まず第一に、あなたが言及した前の質問への回答を彼が投稿したとき、タロンミーズはそれが良いコーディングの代表であることを意図していなかったと思います. したがって、それを 3D に拡張する方法を考え出すことは、あなたの時間の最善の使い方ではないかもしれません。たとえば、厳密に 1 つのスレッドを使用するプログラムを作成する必要があるのはなぜでしょうか。このようなカーネルの正当な使用法はあるかもしれませんが、これはその 1 つではありません。カーネルには、一連の独立した作業を並行して実行する可能性がありますが、代わりに、すべてを 1 つのスレッドに強制してシリアル化しています。並行作業の定義は次のとおりです。
a[i][j][k]=i+j+k;
それを GPU で並列処理する方法を考えてみましょう。
もう 1 つ紹介しておきたいのは、サイズが事前にわかっている問題を扱っているので、言語から得られるメリットを最大限に活用して C を使用して問題に取り組みましょうということです。cudaMalloc を実行するためのネストされたループが必要になる場合もありますが、これはその 1 つではないと思います。
作業を並行して実行するコードは次のとおりです。
#include <stdio.h>
#include <stdlib.h>
// set a 3D volume
// To compile it with nvcc execute: nvcc -O2 -o set3d set3d.cu
//define the data set size (cubic volume)
#define DATAXSIZE 100
#define DATAYSIZE 100
#define DATAZSIZE 20
//define the chunk sizes that each threadblock will work on
#define BLKXSIZE 32
#define BLKYSIZE 4
#define BLKZSIZE 4
// for cuda error checking
#define cudaCheckErrors(msg) \
do { \
cudaError_t __err = cudaGetLastError(); \
if (__err != cudaSuccess) { \
fprintf(stderr, "Fatal error: %s (%s at %s:%d)\n", \
msg, cudaGetErrorString(__err), \
__FILE__, __LINE__); \
fprintf(stderr, "*** FAILED - ABORTING\n"); \
return 1; \
} \
} while (0)
// device function to set the 3D volume
__global__ void set(int a[][DATAYSIZE][DATAXSIZE])
{
unsigned idx = blockIdx.x*blockDim.x + threadIdx.x;
unsigned idy = blockIdx.y*blockDim.y + threadIdx.y;
unsigned idz = blockIdx.z*blockDim.z + threadIdx.z;
if ((idx < (DATAXSIZE)) && (idy < (DATAYSIZE)) && (idz < (DATAZSIZE))){
a[idz][idy][idx] = idz+idy+idx;
}
}
int main(int argc, char *argv[])
{
typedef int nRarray[DATAYSIZE][DATAXSIZE];
const dim3 blockSize(BLKXSIZE, BLKYSIZE, BLKZSIZE);
const dim3 gridSize(((DATAXSIZE+BLKXSIZE-1)/BLKXSIZE), ((DATAYSIZE+BLKYSIZE-1)/BLKYSIZE), ((DATAZSIZE+BLKZSIZE-1)/BLKZSIZE));
// overall data set sizes
const int nx = DATAXSIZE;
const int ny = DATAYSIZE;
const int nz = DATAZSIZE;
// pointers for data set storage via malloc
nRarray *c; // storage for result stored on host
nRarray *d_c; // storage for result computed on device
// allocate storage for data set
if ((c = (nRarray *)malloc((nx*ny*nz)*sizeof(int))) == 0) {fprintf(stderr,"malloc1 Fail \n"); return 1;}
// allocate GPU device buffers
cudaMalloc((void **) &d_c, (nx*ny*nz)*sizeof(int));
cudaCheckErrors("Failed to allocate device buffer");
// compute result
set<<<gridSize,blockSize>>>(d_c);
cudaCheckErrors("Kernel launch failure");
// copy output data back to host
cudaMemcpy(c, d_c, ((nx*ny*nz)*sizeof(int)), cudaMemcpyDeviceToHost);
cudaCheckErrors("CUDA memcpy failure");
// and check for accuracy
for (unsigned i=0; i<nz; i++)
for (unsigned j=0; j<ny; j++)
for (unsigned k=0; k<nx; k++)
if (c[i][j][k] != (i+j+k)) {
printf("Mismatch at x= %d, y= %d, z= %d Host= %d, Device = %d\n", i, j, k, (i+j+k), c[i][j][k]);
return 1;
}
printf("Results check!\n");
free(c);
cudaFree(d_c);
cudaCheckErrors("cudaFree fail");
return 0;
}
あなたがコメントでそれを求めたので、それを機能させるためにあなたのコードに加えることができる最小限の変更をここに示します。また、参照した前の質問からのタロンミーのコメントのいくつかを思い出してください。
「コードの複雑さとパフォーマンス上の理由から、実際にはそうしたくありません。CUDA コードでポインターの配列を使用することは、線形メモリを使用する代替手段よりも難しく、遅くなります。」
「線形メモリを使用することに比べて、それは非常に悪い考えです。」
すべてのポインタのコピーが正しいことを確認するために、これを紙に図解する必要がありました。
#include <cstdio>
inline void GPUassert(cudaError_t code, char * file, int line, bool Abort=true)
{
if (code != 0) {
fprintf(stderr, "GPUassert: %s %s %d\n", cudaGetErrorString(code),file,line);
if (Abort) exit(code);
}
}
#define GPUerrchk(ans) { GPUassert((ans), __FILE__, __LINE__); }
__global__ void doSmth(int*** a) {
for(int i=0; i<2; i++)
for(int j=0; j<2; j++)
for(int k=0; k<2; k++)
a[i][j][k]=i+j+k;
}
int main() {
int*** h_c = (int***) malloc(2*sizeof(int**));
for(int i=0; i<2; i++) {
h_c[i] = (int**) malloc(2*sizeof(int*));
for(int j=0; j<2; j++)
GPUerrchk(cudaMalloc((void**)&h_c[i][j],2*sizeof(int)));
}
int ***h_c1 = (int ***) malloc(2*sizeof(int **));
for (int i=0; i<2; i++){
GPUerrchk(cudaMalloc((void***)&(h_c1[i]), 2*sizeof(int*)));
GPUerrchk(cudaMemcpy(h_c1[i], h_c[i], 2*sizeof(int*), cudaMemcpyHostToDevice));
}
int*** d_c;
GPUerrchk(cudaMalloc((void****)&d_c,2*sizeof(int**)));
GPUerrchk(cudaMemcpy(d_c,h_c1,2*sizeof(int**),cudaMemcpyHostToDevice));
doSmth<<<1,1>>>(d_c);
GPUerrchk(cudaPeekAtLastError());
int res[2][2][2];
for(int i=0; i<2; i++)
for(int j=0; j<2; j++)
GPUerrchk(cudaMemcpy(&res[i][j][0], h_c[i][j],2*sizeof(int),cudaMemcpyDeviceToHost));
for(int i=0; i<2; i++)
for(int j=0; j<2; j++)
for(int k=0; k<2; k++)
printf("[%d][%d][%d]=%d\n",i,j,k,res[i][j][k]);
}
一言で言えば、次の一連のシーケンスを実行する必要があります。
- (ホスト上の) ポインターの多次元配列を malloc します。問題のサイズよりも 1 次元小さく、最後の次元は、ホストではなくデバイス上で cudaMalloc された領域へのポインターのセットです。
- 前のステップで作成されたものと同じクラスですが、前のステップで作成されたものよりも 1 次元小さいポインタの別の多次元配列を作成します。この配列には、デバイス上で cudaMalloc された最終ランクも必要です。
- 前の 2 番目の手順からホスト ポインターの最後のセットを、前の手順でデバイス上の領域 cudaMalloced にコピーします。
- ポインターの多次元配列を指す単一の (ホスト) ポインターになるまで、手順 2 から 3 を繰り返します。これらはすべてデバイス上に常駐しています。