 		}
 		cuda::outstanding_error::clear();
 	}
-
+#endif
+	cuda::kernel_t non_cooperative_kernel(device, kernel_function);
+	auto non_cooperative_config = launch_config;
+	non_cooperative_config.block_cooperation = true;
 	std::cout
 		<< "Launching kernel " << kernel_name
 		<< " with " << num_blocks << " blocks, un-cooperatively, using stream.launch()\n" << std::flush;
-	stream.enqueue.kernel_launch(kernel, launch_config, bar);
+	stream.enqueue.kernel_launch(non_cooperative_kernel, non_cooperative_config, bar);
 	stream.synchronize();
 
 	std::cout << "\nSUCCESS\n";

#include <cuda_runtime_api.h>

#if __CUDACC_VER_MAJOR__ >= 9
#include <cooperative_groups.h>
#endif

#include <iostream>

	stream.enqueue.kernel_launch(kernel, launch_config, bar);
	stream.synchronize();

#if TEST_COOPERATIVE_GROUPS
	try {
		auto can_launch_cooperatively =
			(cuda::device::current::get().get_attribute(cudaDevAttrCooperativeLaunch) > 0);
		if (can_launch_cooperatively) {
			auto cooperative_kernel_function = grid_cooperating_foo;
			auto cooperative_kernel_name = "grid_cooperating_foo";
