   if (argc > 1) devId = atoi(argv[1]);
 
   cudaDeviceProp prop;
-  checkCuda( cudaGetDeviceProperties(&prop, devId));
+  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) );
-
+  cudaSetDevice(devId);
+
+  thrust::host_vector<float> h_idata(nx * ny);
+  thrust::host_vector<float> h_cdata(nx * ny);
+  thrust::host_vector<float> h_tdata(nx * ny);
+  thrust::host_vector<float> gold(nx * ny);
+
+  thrust::device_vector<float> d_idata(nx * ny);
+  thrust::device_vector<float> d_cdata(nx * ny);
+  thrust::device_vector<float> d_tdata(nx * ny);
+  
   // 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");

  // copy 
  // ----
  printf("%25s", "copy");
  checkCuda( cudaMemset(d_cdata, 0, mem_size) );
  // warm up
  copy<<<dimGrid, dimBlock>>>(d_cdata, d_idata);
  timer.reset();
  for(int i = 0; i < NUM_REPS; i++)
  {
     copy<<<dimGrid, dimBlock>>>(d_cdata, d_idata);
  }
  ms = timer.elapsed_milliseconds();
  checkCuda( cudaMemcpy(h_cdata.data(), d_cdata, mem_size, cudaMemcpyDeviceToHost) );
  postprocess(h_idata, h_cdata, ms);

  // -------------
  // copySharedMem 
  // -------------
  printf("%25s", "shared memory copy");
  checkCuda( cudaMemset(d_cdata, 0, mem_size) );
  // warm up
  copySharedMem<<<dimGrid, dimBlock>>>(d_cdata, d_idata);
  timer.reset();
  for(int i = 0; i < NUM_REPS; i++)
  {
     copySharedMem<<<dimGrid, dimBlock>>>(d_cdata, d_idata);
  }
  ms = timer.elapsed_milliseconds();
  checkCuda( cudaMemcpy(h_cdata.data(), d_cdata, mem_size, cudaMemcpyDeviceToHost) );
  postprocess(h_idata, h_cdata, ms);

  // --------------
  // transposeNaive 
  // --------------
  printf("%25s", "naive transpose");
  checkCuda( cudaMemset(d_tdata, 0, mem_size) );
  // warmup
  agency::cuda::bulk_async(grid({dimGrid.x,dimGrid.y}, {dimBlock.x,dimBlock.y}), transpose_naive{}, d_tdata, d_idata);
  timer.reset();
  for(int i = 0; i < NUM_REPS; i++)
  {
    agency::cuda::bulk_async(grid({dimGrid.x,dimGrid.y}, {dimBlock.x,dimBlock.y}), transpose_naive{}, d_tdata, d_idata);
  }
  ms = timer.elapsed_milliseconds();
  checkCuda( cudaMemcpy(h_tdata.data(), d_tdata, mem_size, cudaMemcpyDeviceToHost) );
  postprocess(gold, h_tdata, ms);

  // ------------------
  // transposeCoalesced 
  // ------------------
  printf("%25s", "coalesced transpose");
  checkCuda( cudaMemset(d_tdata, 0, mem_size) );
  // warmup
  agency::cuda::bulk_async(grid({dimGrid.x,dimGrid.y}, {dimBlock.x,dimBlock.y}), transpose_coalesced{}, d_tdata, d_idata);
  timer.reset();
  for(int i = 0; i < NUM_REPS; i++)
  {
    agency::cuda::bulk_async(grid({dimGrid.x,dimGrid.y}, {dimBlock.x,dimBlock.y}), transpose_coalesced{}, d_tdata, d_idata);
  }
  ms = timer.elapsed_milliseconds();
  checkCuda( cudaMemcpy(h_tdata.data(), d_tdata, mem_size, cudaMemcpyDeviceToHost) );
  postprocess(gold, h_tdata, ms);

  // ------------------------
  // transposeNoBankConflicts
  // ------------------------
  printf("%25s", "conflict-free transpose");
  checkCuda( cudaMemset(d_tdata, 0, mem_size) );
  // warmup
  agency::cuda::bulk_async(grid({dimGrid.x,dimGrid.y}, {dimBlock.x,dimBlock.y}), transpose_no_bank_conflicts{}, d_tdata, d_idata);
  timer.reset();
  for(int i = 0; i < NUM_REPS; i++)
  {
    agency::cuda::bulk_async(grid({dimGrid.x,dimGrid.y}, {dimBlock.x,dimBlock.y}), transpose_no_bank_conflicts{}, d_tdata, d_idata);
  }
  ms = timer.elapsed_milliseconds();
  checkCuda( cudaMemcpy(h_tdata.data(), d_tdata, mem_size, cudaMemcpyDeviceToHost) );
  postprocess(gold, h_tdata, ms);
}


    for (int i = 0; i < nx; i++)
      gold[j*nx + i] = h_idata[i*nx + j];
  
  // copy input to device
  d_idata = h_idata;
  
  float ms;
  cuda_timer timer;
