Matrix Transpose (с общей памятью) с произвольным размером на Cuda C - PullRequest
Новая
       68

Matrix Transpose (с общей памятью) с произвольным размером на Cuda C

0 голосов
/ 28 декабря 2018

Я не могу найти способ транспонировать неквадратную матрицу, используя общую память в CUDA C. (Я новичок в CUDA C и C)

На веб-сайте:

https://devblogs.nvidia.com/efficient-matrix-transpose-cuda-cc/

был продемонстрирован эффективный способ транспонирования матрицы (объединенная транспонирование через общую память).Но это работает только для квадратов матриц.

Также код предоставляется на github (так же, как в блоге).

На Stackoverflow существует аналогичный вопрос .Там TILE_DIM = 16 установлено.Но с этой реализацией каждый поток просто копирует один элемент матрицы в матрицу результата.

Это моя текущая реализация: 

__global__ void transpose(double* matIn, double* matTran, int n, int m){
    __shared__ double tile[TILE_DIM][TILE_DIM];
    int i_n = blockIdx.x*TILE_DIM + threadIdx.x;
    int i_m = blockIdx.y*TILE_DIM + threadIdx.y; // <- threadIdx.y only between 0 and 7

    // Load matrix into tile
    // Every Thread loads in this case 4 elements into tile.
    int i;
    for (i = 0; i < TILE_DIM; i += BLOCK_ROWS){
        if(i_n < n  && (i_m+i) < m){
            tile[threadIdx.y+i][threadIdx.x] = matIn[n*(i_m+i) + i_n];
        } else {
            tile[threadIdx.y+i][threadIdx.x] = -1; 
        }
    }
    __syncthreads();

    for (i = 0; i < TILE_DIM; i += BLOCK_ROWS){
        if(tile[threadIdx.x][threadIdx.y+i] != -1){ // <- is there a better way?
            if(true){      // <- what should be checked here?
                matTran[n*(i_m+i) + i_n] = tile[threadIdx.x][threadIdx.y+i];
            } else {
                matTran[m*i_n + (i_m+i)] = tile[threadIdx.x][threadIdx.y+i];
            }
        }
    }
}

, где 4 элемента копируются из потока в плитку,Также четыре элемента из плитки копируются обратно в матрицу результатов.

Здесь Kernel-Configuration <<<a, b>>>: 

where a: (ceil(n/TILE_DIM), ceil(n/TILE_DIM))  (-> is casted to doubles) and 
      b: (TILE_DIM, BLOCK_ROWS) (-> (32, 8))

В настоящее время я использую if(tile[threadIdx.x][threadIdx.y+i] != -1) -статист, чтобы определитькакой поток следует скопировать в матрицу результатов (может быть другой путь).Что касается моих текущих знаний, это ведет себя следующим образом: в блоке ThreadIdx (x, y) копирует данные в плитку, а ThreadIdx (y, x) копирует данные обратно в матрицу результатов.

Я вставил другуюif -статум, чтобы определить, куда копировать данные, так как существует 2 (?) Возможных пункта назначения, в зависимости от ThreadIdx.В настоящее время там вставлено true, но я пробовал много разных вещей.Лучшее, что я мог придумать, было if(threadIdx.x+1 < threadIdx.y+i), которое успешно транспонирует 3x2 -матрицу.

Может кто-нибудь объяснить, чего мне не хватает, записав обратно в матрицу результатов?Очевидно, только один пункт назначения является правильным.Использование

matTran [n * (i_m + i) + i_n] = плитка [threadIdx.x] [threadIdx.y + i];

как в блогеупомянутое должно быть правильным, но я не могу понять, почему это не работает для неквадратных матриц?

1 Ответ

0 голосов
/ 28 декабря 2018

Я слишком усложнил проблему. Здесь , индексы НЕ меняются местами, как я думал.Они пересчитываются с использованием Y- и X-координат потока / блока.Вот фрагмент: 

i_n = blockIdx.y * TILE_DIM + threadIdx.x;  
i_m = blockIdx.x * TILE_DIM + threadIdx.y

Вот исправленный код: 

__global__ void transposeGPUcoalescing(double* matIn, int n, int m, double* matTran){
    __shared__ double tile[TILE_DIM][TILE_DIM];
    int i_n = blockIdx.x * TILE_DIM + threadIdx.x;
    int i_m = blockIdx.y * TILE_DIM + threadIdx.y; // <- threadIdx.y only between 0 and 7

    // Load matrix into tile
    // Every Thread loads in this case 4 elements into tile.
    int i;
    for (i = 0; i < TILE_DIM; i += BLOCK_ROWS){
        if(i_n < n  && (i_m+i) < m){
            tile[threadIdx.y+i][threadIdx.x] = matIn[(i_m+i)*n + i_n];
        }
    }
    __syncthreads();

    i_n = blockIdx.y * TILE_DIM + threadIdx.x; 
    i_m = blockIdx.x * TILE_DIM + threadIdx.y;

    for (i = 0; i < TILE_DIM; i += BLOCK_ROWS){
        if(i_n < m  && (i_m+i) < n){
            matTran[(i_m+i)*m + i_n] = tile[threadIdx.x][threadIdx.y + i]; // <- multiply by m, non-squared!

        }
    }
}

Благодаря этому комментарию за уведомление об ошибке:)

...