 
   cuda_timer(cudaStream_t stream = 0) : stream_(stream), start_{0}, end_{0}
   {
-    checkCuda(cudaEventCreate(&start_));
-    checkCuda(cudaEventCreate(&end_));
+    cudaEventCreate(&start_);
+    cudaEventCreate(&end_);
     reset();
   }
 
   void reset()
   {
-    checkCuda(cudaEventRecord(start_, stream_));
+    cudaEventRecord(start_, stream_);
   }
 
   float elapsed_milliseconds() const
   {
-    checkCuda(cudaEventRecord(end_, stream_));
-    checkCuda(cudaEventSynchronize(end_));
+    cudaEventRecord(end_, stream_);
+    cudaEventSynchronize(end_);
 
     float result = 0;
-    checkCuda(cudaEventElapsedTime(&result, start_, end_));
+    cudaEventElapsedTime(&result, start_, end_);
     return result;
   }
 
   ~cuda_timer()
   {
-    checkCuda(cudaEventDestroy(start_));
-    checkCuda(cudaEventDestroy(end_));
+    cudaEventDestroy(start_);
+    cudaEventDestroy(end_);
   }
 };
 

{
  const int nx = 1024;
  const int ny = 1024;
  const int mem_size = nx*ny*sizeof(float);

  dim3 dimGrid(nx/TILE_DIM, ny/TILE_DIM, 1);
  dim3 dimBlock(TILE_DIM, BLOCK_ROWS, 1);

#include <cstdio>
#include <cassert>
#include <algorithm>
#include <thrust/device_vector.h>
#include <thrust/host_vector.h>
#include <agency/cuda/execution_policy.hpp>

const int TILE_DIM = 32;
const int BLOCK_ROWS = 8;
const int NUM_REPS = 100;

// Check errors and print GB/s
void postprocess(const thrust::host_vector<float>& ref, const thrust::host_vector<float>& res, float ms)
{
  auto mismatch = std::mismatch(ref.begin(), ref.end(), res.begin());
  if(mismatch.first != ref.end())
