Как перенести огромную произвольную матрицу в cuda, используя разделяемую память? - PullRequest
Новая
       65

Как перенести огромную произвольную матрицу в cuda, используя разделяемую память?

0 голосов
/ 17 мая 2019

У меня есть задача транспонировать матрицу в CUDA, используя разделяемую память без конфликтов банков Ограничения: при * высоте <= 10 ^ 8. Ключевые размеры теста: 1x10 ^ 8, 10 ^ 4x10 ^ 4, 10 ^ 8 * 1. </p>

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

Я попытался создать цикл, который помогает работать с огромной матрицей: 

const int BLOCK_DIM = 32;
__global__ void transposeMatrixFast(double* inputMatrix, double* outputMatrix, int width, int height)
{
    __shared__ double temp[BLOCK_DIM][BLOCK_DIM+1];

    int xIndex = blockIdx.x * blockDim.x + threadIdx.x;
    int yIndex = blockIdx.y * blockDim.y + threadIdx.y;
    int offsetx = gridDim.x * blockDim.x;
    int offsety = gridDim.y * blockDim.y;

    for (int y = yIndex; y < height; y += offsety)
    {
        for (int x = xIndex; x < width; x += offsetx)
        {
            if ((xIndex < width) && (yIndex < height))
            {
                int idx = y * width + x;

                temp[threadIdx.y][threadIdx.x] = inputMatrix[idx];
            }

            __syncthreads();


            if ((x < width) && (y < height))
            {
                int idx = x * height + y;

                outputMatrix[idx] = temp[threadIdx.y][threadIdx.x];

            }
        }
    }
}

Теперь на тестовом сервере появляется ошибка «Превышено ограничение по времени». Причина в том, что я не могу использовать преимущества общей памяти в этой строке: outputMatrix[idx] = temp[threadIdx.x][threadIdx.y]; и мой кернер замедляется. Я думаю, что есть другой способ организовать мой цикл, но я не знаю как.

1 Ответ

0 голосов
/ 22 мая 2019

Я нашел другой способ организовать мой цикл, и теперь я могу транспонировать матрицы ЛЮБОГО размера: 

const int BLOCK_SIZE = 32;
__global__ void matrixTransposeSolveBankConflicts(const double *d_a, double *d_b, const unsigned long rows, const unsigned long cols) {

    __shared__ double mat[BLOCK_SIZE][BLOCK_SIZE + 1];

    unsigned long bh = ceil((double)rows / BLOCK_SIZE);
    unsigned long bw = ceil((double)cols / BLOCK_SIZE);

    for (unsigned long blocky = blockIdx.y; blocky < bh; blocky += gridDim.y) {
        for (unsigned long blockx = blockIdx.x; blockx < bw; blockx += gridDim.x) {
            unsigned long bx = blockx * BLOCK_SIZE;
            unsigned long by = blocky * BLOCK_SIZE;

            unsigned long i = by + threadIdx.y;
            unsigned long j = bx + threadIdx.x;

            if (i < rows && j < cols)
            {
                mat[threadIdx.x][threadIdx.y] = d_a[i*cols + j];
            }

            __syncthreads();

            unsigned long ti = bx + threadIdx.y;
            unsigned long tj = by + threadIdx.x;

            if (tj < rows && ti < cols)
            {
                d_b[ti*rows + tj] = mat[threadIdx.y][threadIdx.x];
            }

            __syncthreads();
        }
    }
}
...