2

正方行列を転置する必要があります。マトリックスを使用してプログラムをテストします。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);
}

配列が正しく転置されないのはなぜですか？

4

1 に答える 1

2

問題は、割り当てられた配列の外に出る「余分な」スレッドにあります。

グリッド ブロックを割り当てると、切り上げられます (実際、均等に分割された場合でも、次の整数への丸めが強制されます:)

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

ax または y が [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);
于 2012-11-09T21:01:18.053 に答える