使用CUDA进行矩阵转置

Question

我需要转置一个方阵。我使用matrix：a[i][j] = 0 if i>j, a[i][j] = if i<=j,测试了程序，但结果显示并非所有元素都在正确的位置。

下面是代码(main()除外)：

#include <stdio.h> 
#include <stdlib.h>
__global__ void transpose_kernel (float *a, float *b, int n) {
    unsigned int ax = blockDim.x * blockIdx.x + threadIdx.x;
    unsigned int ay = blockDim.y * blockIdx.y + threadIdx.y;
    unsigned int aIdx = ax + n * ay;
    unsigned int bIdx = ay + n * ax;

    b[bIdx] = a[aIdx];
}

int transpose_host (float *a, float *b, int n) {
    int size = n * n * sizeof (float);
    float *aDev = NULL, *bDev = NULL;

    cudaError_t cuerr = cudaMalloc ((void**)&aDev, size);
    if (cuerr != cudaSuccess) {
        fprintf (stderr, "Cannot allocate GPU memory for aDev: %s\n", cudaGetErrorString (cuerr));
        return (-1);
    }

cuerr = cudaMalloc ((void**)&bDev, size);
if (cuerr != cudaSuccess) {
    fprintf (stderr, "Cannot allocate GPU memory for bDev: %s\n", cudaGetErrorString (cuerr));
    return (-1);
}

dim3 blockSize = dim3 (16, 16, 1);
dim3 gridSize = dim3 (n/16 + 1, n/16 + 1, 1);

cuerr = cudaMemcpy (aDev, a, size, cudaMemcpyHostToDevice);
if (cuerr != cudaSuccess) {
    fprintf (stderr, "Cannot copy data from a to aDev: %s\n", cudaGetErrorString (cuerr));
    return (-1);
}

transpose_kernel <<< gridSize, blockSize >>> (aDev, bDev, n);

cuerr = cudaGetLastError ();
if (cuerr != cudaSuccess) {
    fprintf (stderr, "Cannot launch CUDA kernel: %s\n", cudaGetErrorString (cuerr));
    return (-1);
}

cuerr = cudaDeviceSynchronize ();
if (cuerr != cudaSuccess) {
    fprintf (stderr, "Cannot synchronize CUDA kernel: %s\n", cudaGetErrorString (cuerr));
    return (-1);
}

cuerr = cudaMemcpy (b, bDev, size, cudaMemcpyDeviceToHost);
if (cuerr != cudaSuccess) {
    fprintf (stderr, "Cannot copy data from b to bDev: %s\n", cudaGetErrorString (cuerr));
    return (-1);
}

cudaFree (aDev);
cudaFree (bDev);

    return (0);
}

为什么我的数组不能正确地转置？

Answer 1

问题出在“额外的”线程在分配的数组之外。

当你分配你的网格块时，你会四舍五入(实际上，当所有的东西被平均分配时，也会强制四舍五入到下一个整数：)

dim3 gridSize = dim3 (n/16 + 1, n/16 + 1, 1);

因此总会有一些线程的ax或ay落在[0，n]之外。因此，无论如何，当您将a[aIdx]复制到b[bIdx]中时，您是在将随机数据复制到内存中，实际上，根据调度情况，可能会覆盖“实际”数据。

您可以通过更改内核来修复此问题，以检查以下内容：

if (ax < n && ay < n)
    b[bIdx] = a[aIdx];

您可能想要更改网格大小的舍入，以便在事物均匀分布时不进行舍入：

dim3 gridSize = dim3 ((n+15)/16, (n+15)/16, 1);