CODEWITHSUNDEEP
arrayscuda2doffset
pQB
pQB

Reputation: 3127

Know the Block ID in CUDA from a given 2D offset

i've trying to calculate the blockIdx.x and blockIdx.y from a given offset in CUDA but i'm totally mind-blocked. The idea is read data from shared memory when possible and from global memory in other case.

In example, if I've a 1D array of 64 elements and I configure a kernel with 16x1 threads (4 blocks in total) each thread can access to a position using:

int idx = blockDim.x*blockIdx.x + threadIdx.x

and i can easily get the blockIdx.x of a given index value from the idx as 

int blockNumber = idx / blockDim.x;

but in a 2D scenario with 8x8 elements and a kernel configuration of 4x4 threads (2x2 blocks in total) each thread accesses to a position using:

int x = threadIdx.x + blockIdx.x * blockDim.x;
int y = threadIdx.y + blockIdx.y * blockDim.y;
int pitch = blockDim.x * gridDim.x;
int idx = x + y * pitch;

int sharedMemIndex = threadIdx.x+threadIdx.y+BLOCK_DIM_X;
__shared_block[sharedMemIndex] = fromGlobalMemory[idx];
__syncthreads();

// ... some operations

int unknow_index = __shared_block[sharedMemIndex];

if ( unknow_index within this block? )
    // ... read from shared memory
else
    // ... read from global memory

How can i know the Block ID.x and ID.y at a given idx? i.e. index 34 and 35 are in block (0, 1) and index 36 in block (1, 1). So, if a thread in block (0, 1) read a value of index 35, that thread will know that the value is within its block and will read it from shared memory. The index 35 value will be in stored in the position 11 of the shared memory of the block (0. 1).

Thanks in advance!

Upvotes: 1

Views: 10987

Answers (3)

talonmies
talonmies

Reputation: 72349

In practice, I really can't think of a good reason why this is ever necessary, but you can compute the result like this, for an arbitrary index value idx(assuming column ordered indexing):

int pitch = blockDim.x * gridDim.x;
int tidy = idx / pitch; // div(idx,pitch)
int tidx = idx - (pitch * tidy); // mod(idx,pitch)
int bidx = idx / blockDim.x;
int bidy = idy / blockDim.y;

that should give you the block coordinates of the index in bidx and bidy.

Upvotes: 1

jwdmsd
jwdmsd

Reputation: 2127

There's no need to apply math on Idx to find out the X and Y blocks or go backwards from Idx to find the block index. For every thread (Idx) you can find out the Y and X blocks simply by calling the blockIdx.x and blockIdx.y.

at any point in kernel:

int x = blockIdx.x // will give you X block Index at that particular thread
int y = blockIdx.y // will give you Y block Index at that particular thread.

Update: If you're dead set on the reverse operation, you need to know the value of pitch and block dimensions

   int currentRow = idx/pitch;
   int currentCol = idx%pitch;

   int block_idx_x = currentCol/blockDim.x;
   int block_idx_y = currentRow/blockDim.y;

Upvotes: 1

Pavan Yalamanchili
Pavan Yalamanchili

Reputation: 12099

You are performing unnecessary calculations.

idx / blockDim.x
-->(blockDim.x * blockIdx.x + threadIdx.x)/blockDim.x
-->(blockIdx.x  + threadIdx.x/blockDim.x)
--> blockIdx.x + 0 (threadIdx.x always less than blockDim.x)

You can just use blockIdx.x instead of the convoluted calculation. The same is true for 2D grids (blockIdx.x and blockIdx.y).

Upvotes: 0

Related Questions