 	stream.enqueue.kernel_launch(kernel, launch_config, bar);
 	stream.synchronize();
 
+#if TEST_COOPERATIVE_GROUPS
 	try {
 		auto can_launch_cooperatively =
-#if __CUDACC_VER_MAJOR__ >= 9
 			(cuda::device::current::get().get_attribute(cudaDevAttrCooperativeLaunch) > 0);
-#else
-			false; // This is not strictly true, since the device might support it, but
-			       // 1. We can't check and
-			       // 2. We don't have a cooperative launch API before CUDA 9
-			       // so "false" is good enough.
-#endif
 		if (can_launch_cooperatively) {
 			auto cooperative_kernel_function = grid_cooperating_foo;
 			auto cooperative_kernel_name = "grid_cooperating_foo";

#include <cuda_runtime_api.h>

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

#include <iostream>


			// Same, but using cuda::enqueue_launch
			std::cout
				<< "Launching kernel " << cooperative_kernel_name
				<< " wrapped in a kernel_t structure,"
				<< " with " << num_blocks << " blocks, using cuda::enqueue_launch(),"
				<< " and allowing thread block cooperation\n"
