 	// time non-streamed execution for reference
 	threads=dim3(512, 1);
 	blocks=dim3(n / threads.x, 1);
+	launch_config = cuda::make_launch_config(blocks, threads);
 	start_event.record();
 
 	for (int k = 0; k < nreps; k++)
 	{
-		init_array<<<blocks, threads>>>(d_a.get(), d_c.get(), niterations);
+		device.launch(init_array, launch_config, d_a.get(), d_c.get(), niterations);
 		cuda::memory::copy(h_a.get(), d_a.get(), nbytes);
 	}
 

	threads=dim3(512, 1);
	assert_(n % threads.x == 0);
	blocks=dim3(n / threads.x, 1);
	start_event.record();
	init_array<<<blocks, threads, 0, streams[0].handle()>>>(d_a.get(), d_c.get(), niterations);
	stop_event.record();
	stop_event.synchronize();
	auto time_kernel = cuda::event::time_elapsed_between(start_event, stop_event);

		// asynchronously launch nstreams kernels, each operating on its own portion of data
		for (int i = 0; i < nstreams; i++)
		{
			streams[i].enqueue.kernel_launch(
				init_array, launch_config, d_a.get() + i *n / nstreams, d_c.get(), niterations);
		}

		// asynchronously launch nstreams memcopies.  Note that memcopy in stream x will only
