 
     for(int j = 0; j < TILE_DIM; j += BLOCK_ROWS)
     {
-      odata[(y+j)*width + x] = tile[threadIdx.x][threadIdx.y + j];
+      odata[(y+j)*width + x] = tile.data[threadIdx.x][threadIdx.y + j];
     }
   }
 };

};


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), transpose_coalesced{}, raw_pointer_cast(d_tdata.data()), raw_pointer_cast(d_idata.data()), agency::share<1,tile2d_t>());
  });

  h_tdata = d_tdata;
