Смещение столбцов в матрице строк-мажоров

Question

Учитывая жирную матрицу A=[[1,2,3,4,5],[6,7,8,9,10],[11,12,13,14,15]], я пытаюсь сместить (переставить) строки, т.е. B=[[11,12,13,14,15],[1,2,3,4,5],[6,7,8,9,10]], и сместить столбцы, например, C=[[5,1,2,3,4],[10,6,7,8,9],[15,11,12,13,14]]

Я сделал аналогичное Ядро как смещающиеся столбцы:

// A->C
__global__ void column_shift(int* mat, int row, int col) {
    int row_num = blockDim.x * blockIdx.x + threadIdx.x;
    if (row_num < row) {
        int a = mat[row_num * col];
        for (int i = 0; i < col - 1; ++i) {
            mat[row_num * col + i] = mat[row_num * col + i + 1];
        }
        mat[row_num * col + (col - 1)] = a;
    }
}

// A->B
__global__ void row_shift(int* mat, int row, int col) {
    int col_num = blockDim.x * blockIdx.x + threadIdx.x;
    if (col_num < col) {
        int a = mat[(row - 1) * col + col_num];
        for (int i = row - 1; i > 0; i--) {
            mat[i * col + col_num] = mat[(i - 1) * col + col_num];
        }
        mat[col_num] = a;
    }
}

Однако, по сравнению с row_shift, column_shift работает хуже. Я думаю, это связано с объединением памяти. Есть ли эффективный способ повысить производительность column_shift?

Answer 1

Самая очевидная проблема с производительностью column_shift - это отсутствие объединения памяти. Это может быть исправлено с помощью перекоса потоков, выполняющих сдвиг данных строки, а не одного потока.

Рассмотрим следующий пример (обратите внимание, я переписал ваши ядра, чтобы использовать простой вспомогательный класс, который упрощает ядро). код, и снижает риск индексации ошибок вычислений (как было в случае по крайней мере в одном из ядер, которые вы первоначально разместили):

#include <thrust/device_vector.h>
#include <thrust/iterator/counting_iterator.h>
#include <vector>

struct stride2D
{
   int* p;
   int s0;
   __host__ __device__
   stride2D(int* _p, int _s0) : p(_p), s0(_s0) {};
   __host__ __device__
   int operator  () (int x, int y) const { return p[x*s0 + y]; };
   __host__ __device__
   int& operator () (int x, int y) { return p[x*s0 + y]; };
};

__global__ void column_shift2(int* mat, int row, int col)
{
    int row_num = blockDim.x * blockIdx.x + threadIdx.x;
    stride2D m(mat, col);   

    if (row_num < row) {
        int a = m(row_num, 0);
        for (int i = 0; i < col-1; i++) {
            m(row_num, i) = m(row_num, i+1);
        }
        m(row_num, col-1) = a;
    }
}

__global__ void column_shift3(int* mat, int row, int col)
{
    int row_num = blockDim.y * blockIdx.y + threadIdx.y;
    stride2D m(mat, col);   
    if (row_num < row) {
        int a = m(row_num, 0);
        for (int i = threadIdx.x; i < col-1; i += warpSize) {
            m(row_num, i) = m(row_num, i+1);
        }
        if (threadIdx.x == 0) m(row_num, col-1) = a;
    }
}

__global__ void row_shift2(int* mat, int row, int col) {
    int col_num = blockDim.x * blockIdx.x + threadIdx.x;
    stride2D m(mat, col);   
    if (col_num < col) {
        int a = m(row-1, col_num);
        for (int i = row - 1; i > 0; i--) {
            m(i, col_num) = m(i-1, col_num);
        }
        m(0, col_num) = a;
    }
}

int main()
{
    const int r = 300, c = 900, n = r * c;

    {
        std::vector<int> idata(n);
        thrust::counting_iterator<int> first(1);
        thrust::copy(first, first+n, idata.begin());

        thrust::device_vector<int> ddata(idata);
        int* d = thrust::raw_pointer_cast(ddata.data());

        int bsize = 256, nblocks = (c / bsize) + (c % bsize > 0) ? 1 : 0;
        row_shift2<<<nblocks, bsize>>>(d, r, c);
        cudaDeviceSynchronize();

        std::vector<int> odata(n);
        thrust::copy(ddata.begin(), ddata.end(), odata.begin());
    }

    {
        std::vector<int> idata(n);
        thrust::counting_iterator<int> first(1);
        thrust::copy(first, first+n, idata.begin());

        thrust::device_vector<int> ddata(idata);
        int* d = thrust::raw_pointer_cast(ddata.data());

        int bsize = 256, nblocks = (r / bsize) + (r % bsize > 0) ? 1 : 0;
        column_shift2<<<nblocks, bsize>>>(d, r, c);
        cudaDeviceSynchronize();

        std::vector<int> odata(n);
        thrust::copy(ddata.begin(), ddata.end(), odata.begin());
    }

    {
        std::vector<int> idata(n);
        thrust::counting_iterator<int> first(1);
        thrust::copy(first, first+n, idata.begin());

        thrust::device_vector<int> ddata(idata);
        int* d = thrust::raw_pointer_cast(ddata.data());

        const int bwidth = 32;
        dim3 bsize(bwidth, 1024/bwidth);
        int nblocks = (r / bsize.y) + (r % bsize.y > 0) ? 1 : 0;
        column_shift3<<<nblocks, bsize>>>(d, r, c);
        cudaDeviceSynchronize();

        std::vector<int> odata(n);
        thrust::copy(ddata.begin(), ddata.end(), odata.begin());
    }

    cudaDeviceReset();

    return 0;
}

Единственное реальное необходимое изменение - это внутреннее копирование l oop в пределах операция column_shift:

    for (int i = threadIdx.x; i < col-1; i += warpSize) {
        m(row_num, i) = m(row_num, i+1);
    }

Теперь мы используем l oop с перекосом (для корректности его необходимо запустить с blockDim.x = 32). Профилирование этого кода показывает это:

nvprof ./permute
==13687== NVPROF is profiling process 13687, command: ./permute
==13687== Profiling application: ./permute
==13687== Profiling result:
            Type  Time(%)      Time     Calls       Avg       Min       Max  Name
 GPU activities:   51.24%  643.80us         1  643.80us  643.80us  643.80us  column_shift2(int*, int, int)
                   21.36%  268.41us         3  89.471us  89.087us  89.887us  [CUDA memcpy HtoD]
                   21.06%  264.57us         3  88.191us  87.647us  89.023us  [CUDA memcpy DtoH]
                    5.54%  69.631us         1  69.631us  69.631us  69.631us  row_shift2(int*, int, int)
                    0.81%  10.144us         1  10.144us  10.144us  10.144us  column_shift3(int*, int, int)
      API calls:   68.19%  114.44ms         3  38.148ms  78.552us  114.28ms  cudaMalloc
                   30.00%  50.352ms         1  50.352ms  50.352ms  50.352ms  cudaDeviceReset
                    0.65%  1.0974ms         6  182.89us  102.55us  246.46us  cudaMemcpyAsync
                    0.44%  732.75us         3  244.25us  13.565us  646.95us  cudaDeviceSynchronize
                    0.21%  348.53us        97  3.5930us     263ns  197.14us  cuDeviceGetAttribute
                    0.17%  290.47us         1  290.47us  290.47us  290.47us  cuDeviceTotalMem
                    0.16%  266.04us         6  44.339us  2.3170us  87.602us  cudaStreamSynchronize
                    0.11%  184.85us         3  61.616us  53.903us  71.672us  cudaFree
                    0.03%  54.650us         3  18.216us  13.862us  25.133us  cudaLaunchKernel
                    0.03%  51.108us         1  51.108us  51.108us  51.108us  cuDeviceGetName
                    0.00%  4.0760us         3  1.3580us     408ns  3.1910us  cuDeviceGetCount
                    0.00%  3.4620us         1  3.4620us  3.4620us  3.4620us  cuDeviceGetPCIBusId
                    0.00%  1.6850us         2     842ns     248ns  1.4370us  cuDeviceGet
                    0.00%     585ns         1     585ns     585ns     585ns  cuDeviceGetUuid

т. Е. Копия с выделением основы примерно в 60 раз быстрее, чем ваша оригинальная реализация.

[Обратите внимание, что весь код чрезвычайно проверен незначительно, и никаких гарантий правильности или оптимальности не делается и не подразумевается]