CUDA 2次元配列

Question

私はデバイス上に行列を割り当てようとしています。カーネルにいくつかの番号を記入し、それをホストにコピーし直しています。問題は、ホスト上では1行しか満たされていないようだということです。

が、私はこのような何かを得た：

ここ

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のような数量のインデックスの調整としては使用できません（ドキュメントを調べてください）。第2に、ピッチ値は最終的に列インデックスではなく行のインデックスを掛けなければなりません。あなたの問題に関連していない

3. が、あなたの最終printf文は、64ビットプラットフォーム上にある場合、それは%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 
$