Отправка трехмерного массива в ядро CUDA
Я взял код, приведенный в качестве ответа на вопрос Как я могу сложить два 2d (разбитых) массива, используя вложенные циклы for? и попытался использовать его для 3D вместо 2D и немного изменил другие части, теперь это выглядит следующим образом:
__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*** d_c;
GPUerrchk(cudaMalloc((void****)&d_c,2*sizeof(int**)));
GPUerrchk(cudaMemcpy(d_c,h_c,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]);
}
В приведенном выше коде я использую 2 в качестве размеров для каждого из измерений h_c, в реальной реализации я буду иметь эти размеры в очень больших числах и в разных для каждой части подмассивов "int***" или более измерений, У меня проблема с частью после вызова ядра, где я пытаюсь скопировать результаты обратно в массив res. Можете ли вы помочь мне решить проблему? Пожалуйста, вы можете показать решение так, как я пишу выше. Спасибо!
1 ответ
Прежде всего, я думаю, что talonmies, когда он разместил ответ на предыдущий вопрос, который вы упомянули, не намеревался, чтобы это было представлением хорошего кодирования. Поэтому выяснение того, как распространить его на 3D, может быть не лучшим использованием вашего времени. Например, почему мы хотим писать программы, которые используют ровно один поток? Хотя такое ядро может быть законно использовано, это не одно из них. Ваше ядро имеет возможность выполнять кучу независимых работ параллельно, но вместо этого вы вынуждаете все это в один поток и сериализуете его. Определение параллельной работы:
a[i][j][k]=i+j+k;
Давайте разберемся, как обрабатывать это параллельно на GPU.
Еще одно вступительное замечание, которое я хотел бы сделать, заключается в том, что, поскольку мы имеем дело с проблемами, размеры которых известны заранее, давайте используем C для их решения с максимальной пользой, которую мы можем получить от языка. В некоторых случаях могут понадобиться вложенные циклы для выполнения cudaMalloc, но я не думаю, что это один из них.
Вот код, который выполняет работу параллельно:
#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 - многомерный массив указателей (на хосте), на одно измерение меньше размера проблемы, причем последнее измерение представляет собой набор указателей на регионы cudaMalloc, расположенные на устройстве, а не на хосте.
- создайте другой многомерный массив указателей того же класса, что и созданный на предыдущем шаге, но на одно измерение меньше, чем созданный на предыдущем шаге. этот массив также должен иметь свои последние оценки cudaMalloc на устройстве.
- скопируйте последний набор указателей хоста из второго предыдущего шага в область cudaMalloced на устройстве на предыдущем шаге.
- повторяйте шаги 2-3 до тех пор, пока мы не получим один (хост) указатель, указывающий на многомерный массив указателей, которые теперь все находятся на устройстве.