Знать идентификатор блока в CUDA по заданному 2D смещению

Question

Я пытаюсь вычислить blockIdx.x и blockIdx.y из заданного смещения в CUDA, но я полностью заблокирован. Идея заключается в чтении данных из общей памяти, когда это возможно, и из глобальной памяти в другом случае.

Например, если у меня есть одномерный массив из 64 элементов, и я настраиваю ядро ​​с потоками 16x1 (всего 4 блока), каждый поток может получить доступ к позиции, используя:

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

и я легко могу получить blockIdx.x заданного значения индекса из idx как

int blockNumber = idx / blockDim.x; 

, но в 2D-сценарии с элементами 8x8 и конфигурацией ядра с потоками 4x4 (всего 2x2 блока) каждый поток получает доступ к позиции, используя:

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

Как я могу узнать идентификатор блока.х и идентификатор.й для данного идентификатора? то есть индексы 34 и 35 находятся в блоке (0, 1) и индексы 36 в блоке (1, 1). Таким образом, если поток в блоке (0, 1) считывает значение индекса 35, этот поток будет знать, что значение находится внутри своего блока, и будет читать его из общей памяти. Значение индекса 35 будет храниться в позиции 11 общей памяти блока (0. 1).

Заранее спасибо!

Answer 1

На практике я действительно не могу придумать вескую причину, почему это когда-либо необходимо, но вы можете вычислить результат следующим образом для произвольного значения индекса idx (при условии индексации по порядку столбцов):

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;

, который должен дать вам координаты блока индекса в bidx и bidy.

Answer 2

Нет необходимости применять математику к Idx , чтобы найти блоки X и Y или перейти назад от Idx, чтобы найти индекс блока.Для каждого потока (Idx) вы можете узнать блоки Y и X, просто вызвав blockIdx.x и blockIdx.y .

в любой точке ядра.:

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. 

Обновление: Если вы полностью настроены на обратную операцию, вам нужно знать значение шага и размеры блока

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

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

Answer 3

Вы выполняете ненужные вычисления.

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)

Вы можете просто использовать blockIdx.x вместо сложного вычисления.То же самое верно для двумерных сеток (blockIdx.x и blockIdx.y).