 #include <numeric>
 #include <limits>
 
-using std::size_t;
+using cuda::size_t;
 
 namespace kernels {
 

	cuda::memory::copy(ptr_in.get(), other_arr);

	check_output_is_iota("copy from (managed) global memory into a 2D array", ptr_in.get(), arr.size());
	
	// also asynchronously
	auto stream = device.create_stream(cuda::stream::async);
	cuda::memory::async::copy(other_arr, ptr_out.get(), stream);
	cuda::memory::async::copy(ptr_in.get(), other_arr, stream);
	device.synchronize();

		threadIdx.x + blockIdx.x * blockDim.x,
		threadIdx.y + blockIdx.y * blockDim.y
	};
	const auto gtid_serialized = gtid.x + gtid.y * static_cast<unsigned>(w);

	if (gtid.x < w && gtid.y < h) {
		const float x = tex2D<float>(texture_source, gtid.x, gtid.y);
//		printf("Thread %u %u, reading value %.4g, and writing to index %3u\n", gtid.x, gtid.y, x, gtid_serliazed);
		destination[gtid_serialized] = x;
	}
}

