CUDA: доступ к произвольным длинным матрицам в обоих измерениях

Question

Привет, в настоящее время я использую потоки, проиндексированные только в одном измерении, чтобы получить доступ ко всем элементам матрицы следующим образом:

// Thread-ID
int idx = blockIdx.x * blockDim.x + threadIdx.x;
// Offset:
int offset = gridDim.x * blockDim.x;

while(idx < MATRIX_ROWS * MATRIX_COLS)
{   
    row = idx % MATRIX_ROWS;
    col = idx / MATRIX_ROWS;    

    matrix[ row ][ col ] = ...;
    idx += offset;
}

Теперь мне стало интересно, как получить доступ к произвольным длинным матрицам с двумерным индексированием. Мне бы хотелось, чтобы один блок всегда имел доступ к отдельным элементам одной строки. Примерно так (x-index ссылается на столбцы, а y-индекс на строки матрицы):

// Thread-IDs
int idx = blockIdx.x * blockDim.x + threadIdx.x;
int idy = blockIdx.y * blockDim.y + threadIdx.y;

// Offset:
int offset = gridDim.x * blockDim.x;

while(idx < MATRIX_COLS)
{   
    matrix[ idy ][ idx ] = ...;
    idx += offset;
}

Теперь давайте предположим, что матрица имеет больше строк, чем я начал блокировать при вызове ядра: при запуске N блоков первые N строк матрицы обрабатываются правильно, но как насчет других строк? Как бы Вы это сделали? Спасибо!

РЕДАКТИРОВАТЬ : у меня возникла идея, но я не знаю, является ли это каким-то «уродливым» кодированием!?

// Thread-IDs
int idx0 = blockIdx.x * blockDim.x + threadIdx.x;
int idx = idx0;
int idy = blockIdx.y * blockDim.y + threadIdx.y;

// Offset:
int offsetx = gridDim.x * blockDim.x;
int offsety = gridDim.y * blockDim.y;

while(idx < MATRIX_COLS && idy < MATRIX_ROWS)
{   
    matrix[ idy ][ idx ] = ...;

    idx += offsetx;
    if(idx > MATRIX_COLS)
    {
        // Jump to nex row and start from 'beginning' concerning columns
        idy += offsety;
        idx = idx0;
    }
}

Answer 1

Возможно, что-то вроде этого - то, что вы ищете?

// Thread-IDs
int idx = blockIdx.x * blockDim.x + threadIdx.x;
int idy = blockIdx.y * blockDim.y + threadIdx.y;

// Offset:
int offsetx = gridDim.x * blockDim.x;
int offsety = gridDim.y * blockDim.y;

for(row = idy; row < MATRIX_ROWS; i+=offsety) {
    float * pos = matrix + row;

#pragma unroll
    for(col = idx; col < MATRIX_COLS; col+=offsetx) {
        pos[col] = .....;
    }
}

Если MATRIX_COLS является определением или константой препроцессора, компилятор может развернуть внутренний цикл и повысить производительность.

РЕДАКТИРОВАТЬ: первая версия кода была написана с порядком основных столбцов, застрявших в моей голове, так что игнорируйте комментарий, который был здесь. Доступ должен быть объединен таким образом.