CUDA二维数组[重复]

Question

我正在尝试在设备上分配矩阵，在内核中填充一些数字，然后将其复制回主机。问题是在主机上似乎只有一行被填满。

我得到了这样的东西：

9 9 9 9
-1 -1 -1 -1
-1 -1 -1 -1
-1 -1 -1 -1

这是我的代码：

#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <stdio.h>
#include <stdlib.h>

void check(cudaError x) {
    fprintf(stderr, "%s\n", cudaGetErrorString(x));
}

void showMatrix2(int* v1, int width, int height) {
    printf("---------------------\n");
    for (int i = 0; i < width; i++) {
        for (int j = 0; j < height; j++) {
            printf("%d ", v1[i * width + j]);
        }
        printf("\n");
    }
}

__global__ void kernel(int* tab,int width, int height, int pitch) {

    int row = threadIdx.x + blockIdx.x * blockDim.x;
    int col = threadIdx.y + blockIdx.y * blockDim.y;

    if (row < width && col < height) {
        tab[col * pitch + row] = 9;
    }
}

int main()
{
    int width = 4;
    int height = 4;

    int* d_tab;
    int* h_tab;

    int realSize = width * height* sizeof(int);

    size_t pitch;
    check( cudaMallocPitch(&d_tab, &pitch, width * sizeof(int), height) );
    h_tab = (int*)malloc(realSize);
    check( cudaMemset(d_tab, 0, realSize) );

    dim3 grid(4, 4);
    dim3 block(4, 4);
    kernel <<<grid, block>>>(d_tab, width, height, pitch);

    check( cudaMemcpy2D(h_tab, width*sizeof(int), d_tab, pitch, width*sizeof(int), height, cudaMemcpyDeviceToHost) );

    showMatrix2(h_tab, width, height);
    printf("\nPitch size: %d \n", pitch);
    getchar();
    return 0;
}

Answer 1

1. 任何时候您在使用 CUDA 代码时遇到问题，除了进行错误检查外，请使用 cuda-memcheck 运行您的代码。如果你这样做了，你至少会得到一个关于发生了什么的提示，然后你可以使用像this 这样的技术来继续你自己的调试。即使您无法弄清楚，cuda-memcheck 的输出也会对其他试图帮助您的人有用。

2. 您的内核中有无效的写入。这里有多个错误。为了正确访问内核代码中的倾斜分配，我强烈建议研究the documentation 中针对cudaMallocPitch 给出的示例。简而言之，这种索引生成就是被打破了：

   tab[col * pitch + row]
   

   首先，cudaMallocPitch 返回的pitch 是以字节 为单位的宽度。您不能将其用作对 int 或 float 等数量的索引的调整（研究文档）。其次，音高值最终应该乘以 行 索引，而不是列索引。

3. 与您的问题无关，但如果您在 64 位平台上，您的最终 printf 语句的格式说明符不正确，它应该是 %ld（或者更好，%lu）。

这是修复了索引问题的代码，对我来说似乎可以正常工作：

$ cat t109.cu
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <stdio.h>
#include <stdlib.h>

void check(cudaError x) {
    fprintf(stderr, "%s\n", cudaGetErrorString(x));
}

void showMatrix2(int* v1, int width, int height) {
    printf("---------------------\n");
    for (int i = 0; i < width; i++) {
        for (int j = 0; j < height; j++) {
            printf("%d ", v1[i * width + j]);
        }
        printf("\n");
    }
}

__global__ void kernel(int* tab,int width, int height, int pitch) {

    int row = threadIdx.x + blockIdx.x * blockDim.x;
    int col = threadIdx.y + blockIdx.y * blockDim.y;

    if (row < width && col < height) {
        *( ((int *)(((char *)tab) + (row * pitch))) + col) = 9;
    }
}

int main()
{
    int width = 4;
    int height = 4;

    int* d_tab;
    int* h_tab;

    int realSize = width * height* sizeof(int);

    size_t pitch;
    check( cudaMallocPitch(&d_tab, &pitch, width * sizeof(int), height) );
    h_tab = (int*)malloc(realSize);
    check( cudaMemset(d_tab, 0, realSize) );

    dim3 grid(4, 4);
    dim3 block(4, 4);
    kernel <<<grid, block>>>(d_tab, width, height, pitch);

    check( cudaMemcpy2D(h_tab, width*sizeof(int), d_tab, pitch, width*sizeof(int), height, cudaMemcpyDeviceToHost) );

    showMatrix2(h_tab, width, height);
    printf("\nPitch size: %ld \n", pitch);
    return 0;
}
$ nvcc -arch=sm_61 -o t109 t109.cu
$ cuda-memcheck ./t109
========= CUDA-MEMCHECK
no error
no error
no error
---------------------
9 9 9 9
9 9 9 9
9 9 9 9
9 9 9 9

Pitch size: 512
========= ERROR SUMMARY: 0 errors
$