 
     for(int j = 0; j < TILE_DIM; j += BLOCK_ROWS)
     {
-      odata[(y+j)*width + x] = tile[self.inner().index()[0]][self.inner().index()[1] + j];
-    }
-  }
+      odata[(y+j)*width + x] = tile.data[self.inner().index()[0]][self.inner().index()[1] + j];
+    }
+  }
+};
+
+
+struct conflict_free_tile2d_t
+{
+  float data[TILE_DIM][TILE_DIM+1];
 };
 
 
 struct transpose_no_bank_conflicts
 {
-  __device__ void operator()(cuda_thread_2d& self, float* odata, const float* idata)
-  {
-    __shared__ float tile[TILE_DIM][TILE_DIM+1];
-      
+  __device__ void operator()(cuda_thread_2d& self, float* odata, const float* idata, conflict_free_tile2d_t& tile)
+  {
     int x = blockIdx.x * TILE_DIM + threadIdx.x;
     int y = blockIdx.y * TILE_DIM + threadIdx.y;
     int width = gridDim.x * TILE_DIM;
 
     for(int j = 0; j < TILE_DIM; j += BLOCK_ROWS)
     {
-      tile[threadIdx.y+j][threadIdx.x] = idata[(y+j)*width + x];
+      tile.data[threadIdx.y+j][threadIdx.x] = idata[(y+j)*width + x];
     }
 
     self.inner().wait();

};


struct transpose_coalesced
{
  __device__ void operator()(cuda_thread_2d& self, float* odata, const float* idata)
  {
    __shared__ float tile[TILE_DIM][TILE_DIM];
      
    int x = self.outer().index()[0] * TILE_DIM + self.inner().index()[0];
    int y = self.outer().index()[1] * TILE_DIM + self.inner().index()[1];
    int width = self.outer().group_shape()[0] * TILE_DIM;

    for(int j = 0; j < TILE_DIM; j += BLOCK_ROWS)
    {
      tile[self.inner().index()[1]+j][self.inner().index()[0]] = idata[(y+j)*width + x];
    }

    self.inner().wait();


  ms = time_invocation([&]
  {
    agency::cuda::bulk_async(grid(dim_grid, dim_block), copy_shared_mem{}, raw_pointer_cast(d_cdata.data()), raw_pointer_cast(d_idata.data()), agency::share<1,tile1d_t>());
  });

  h_cdata = d_cdata;
