cuda：デバイスメモリにグローバル変数を使用する

Question

私はcudaコードで適切な変数を使用する方法に関する質問があります。私のプログラムは、さまざまな関数でアクセスする必要があるたくさんの配列を持っています。それらを渡すのを避け、平面変数1D配列の代わりにグローバル変数と2D mallocpitch配列を使いたいと思います。ですから、私は次のようなことを考えています：

__device__ double * dataPtr ; 
__device__ size_t dataPitch; 
.... 
int main() 
{ 
double * dataPtrLoc; size_t dataPitchLoc; 
cudaMallocPitch((void**) &dataPtrLoc, &dataPitchLoc, width*sizeof(double), height); 
cudaMemcpyToSymbol(dataPtr, &dataPtrLoc, sizeof(dataPtrLoc)); 
cudaMemcpyToSymbol(dataPitch, &dataPitchLoc, sizeof(dataPitchLoc)); 
... 
} 

グローバルな2Dデバイスデータを得るのは良い方法ですか？あなたは示唆を与えることができますか？

編集：私はこのプログラムを作り、それをコンパイルし、正常に動作：

#include <stdio.h> 
__device__ int *d_gridPtr; 
__device__ size_t d_gridPitch; 

__device__ int valij(int ii, int jj) 
{ 
    int* row = (int*)((char*)d_gridPtr + ii * d_gridPitch); 
    return (row[jj]); 
} 

__global__ void printval() 
{ 
    int val0, val1, val2, val3; 
    val0= valij(0,0); 
    val1= valij(0,1); 
    val2= valij(1,0); 
    val3= valij(1,1); 
    printf("%d %d %d %d \n", val0, val1, val2, val3); 
} 

int main() 
{ 
    size_t d_gridPitchLoc; 
    int * d_gridPtrLoc; 
    cudaMallocPitch((void**)&d_gridPtrLoc, &d_gridPitchLoc, 2 * sizeof(int), 2); 
    cudaMemcpyToSymbol(d_gridPtr, & d_gridPtrLoc, sizeof(d_gridPtrLoc)); 
    cudaMemcpyToSymbol(d_gridPitch, &d_gridPitchLoc, sizeof(float)); 

    int h_mem[2*2]={0,1,100,4}; 
    size_t hostpitch = 2* sizeof(int); 
    cudaMemcpy2D(d_gridPtrLoc,d_gridPitchLoc,h_mem,hostpitch,2*sizeof(int),2,cudaMemcpyHostToDevice); 

    printval<<<1,1>>>(); 
    cudaDeviceReset(); 
} 

Answer 1

反りやブロックアクセスのすべてのスレッドが同じ読み取り専用のグローバルメモリアドレス（例えば配列インデックス）場合は、同時に読み込み専用のグローバルデータをメモリ配列__constant__に格納することを検討してください。データに書き込む場合、__constant__を使用することはできません。

アレイが読み取り専用であり、アクセスパターンが（ワープやブロック内で）強い2Dローカリティを持つ場合は、代わりにテクスチャを使用することを検討してください。