 
   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_);
   }
 };
 

  if (argc > 1) devId = atoi(argv[1]);

  cudaDeviceProp prop;
  checkCuda( cudaGetDeviceProperties(&prop, devId));
  printf("\nDevice : %s\n", prop.name);
  printf("Matrix size: %d %d, Block size: %d %d, Tile size: %d %d\n", 
         nx, ny, TILE_DIM, BLOCK_ROWS, TILE_DIM, TILE_DIM);
  printf("dimGrid: %d %d %d. dimBlock: %d %d %d\n",
         dimGrid.x, dimGrid.y, dimGrid.z, dimBlock.x, dimBlock.y, dimBlock.z);
  
  checkCuda( cudaSetDevice(devId) );

  std::vector<float> h_idata(nx * ny);
  std::vector<float> h_cdata(nx * ny);
  std::vector<float> h_tdata(nx * ny);
  std::vector<float> gold(nx * ny);
  
  float *d_idata, *d_cdata, *d_tdata;
  checkCuda( cudaMalloc(&d_idata, mem_size) );
  checkCuda( cudaMalloc(&d_cdata, mem_size) );
  checkCuda( cudaMalloc(&d_tdata, mem_size) );

  // check parameters and calculate execution configuration
  if (nx % TILE_DIM || ny % TILE_DIM) {
    throw std::logic_error("nx and ny must be a multiple of TILE_DIM");

#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())
