Прежде чем я углублюсь, это общая идея того, что происходит:
Общая идея заключается в том, что у меня есть x массивов с плавающей точкой, и я хочу добавить каждый из них последовательно в другой массив (скалярное добавление):
t = массив;
a = массив массива;
t = нули
t + = a [0]
t + = a [1]
...
t + = a [N]
где + = представляет скалярное сложение.
Этопрямо вперед.Я пытался сжать код, я должен быть как можно более компактным и сохранить функциональность.Проблема здесь заключается в том, что для массивов определенного размера - я вижу проблемы с размером, превышающим 128 x 128 x 108. По сути, суммация памяти, скопированной обратно на хост, не совпадает с тем, что я рассчитал.Я застрял на этом весь день, поэтому я перестану тратить свое время.Я действительно не могу объяснить, почему это происходит.Я рассуждаю так:
- Использование слишком большого количества постоянного пространства (без использования)
- Использование слишком большого количества регистров (нет)
- Неправильное условие в ядре для проверки, если idx, idy, idz находятся в пределах (это все еще может быть)
- Что-то смешное с gpu (пробовал на gt280 и tesla C1060 и C2060)
- Неверный формат printf (надеюсь, этоэто) * ...
Этот список можно продолжить.Спасибо за просмотр этого, если у вас есть время.Кажется, проблема почти связана с памятью (т. Е. Размеры памяти> 128 * 128 * 108 не работают. Поэтому 64 * 128 * 256 будет работать или любая их перестановка).
Вот полный исходный код (который должен быть скомпилирован с nvcc):
#include <cuda.h>
#include <iostream>
#include <stdio.h>
#include <assert.h>
#define BSIZE 8
void cudaCheckError(cudaError_t e,const char * msg) {
if (e != cudaSuccess){
printf("Error number: %d\n",e);
printf("%s\n",msg);
}
};
__global__ void accumulate(float * in,float * out, int3 gdims, int zlevel) {
int idx = blockIdx.x*blockDim.x + threadIdx.x;
int idy = blockIdx.y*blockDim.y + threadIdx.y;
int idz = threadIdx.z;
long int index = (zlevel*((int)BSIZE)+idz)*gdims.x*gdims.y+ \
idy*gdims.x+ \
idx;
if ( idx < gdims.x && idy < gdims.y && (idz + zlevel*(int)BSIZE) < gdims.z) {
out[index] += in[index];
}
};
int main(int argc, char * argv[]) {
int width,
height,
depth;
if (argc != 4) {
printf("Must have 3 inputs: width height depth\n");
exit(0);
}
float tempsum;
int count =0;
width = atoi(argv[1]);
height = atoi(argv[2]);
depth = atoi(argv[3]);
printf("Dimensions (%d,%d,%d)\n",width,height,depth);
int3 dFull;
dFull.x = width+2;
dFull.y = height+2;
dFull.z = depth+2;
printf("Dimensions (%d,%d,%d)\n",dFull.x,dFull.y,dFull.z);
int fMemSize=dFull.x*dFull.y*dFull.z;
int nHostF=9;
float * f_hostZero;
float ** f_dev;
float * f_temp_host;
float * f_temp_dev;
dim3 grid( dFull.x/(int)BSIZE+1, dFull.y/(int)BSIZE + 1);
dim3 threads((int)BSIZE,(int)BSIZE,(int)BSIZE);
printf("Threads (x,y) : (%d,%d)\nGrid (x,y) : (%d,%d)\n",threads.x,threads.y,grid.x,grid.y);
int num_zsteps=dFull.z/(int)BSIZE + 1;
printf("Number of z steps to take : %d\n",num_zsteps);
// Host array allocation
f_temp_host = new float[fMemSize];
f_hostZero = new float[fMemSize];
// Allocate nHostF address on host
f_dev = new float*[nHostF];
// Host array assignment
for(int i=0; i < fMemSize; i++){
f_temp_host[i] = 1.0;
f_hostZero[i] = 0.0;
}
// Device allocations - allocated for array size + 2
for(int i=0; i<nHostF; i++){
cudaMalloc((void**)&f_dev[i],sizeof(float)*fMemSize);
}
// Allocate the decive pointer
cudaMalloc( (void**)&f_temp_dev, sizeof(float)*fMemSize);
cudaCheckError(cudaMemcpy((void *)f_temp_dev,(const void *)f_hostZero,
sizeof(float)*fMemSize,cudaMemcpyHostToDevice),"At first mem copy");
printf("Memory regions allocated\n");
// Copy memory to each array
for(int i=0; i<nHostF; i++){
cudaCheckError(cudaMemcpy((void *)(f_dev[i]),(const void *)f_temp_host,
sizeof(float)*fMemSize,cudaMemcpyHostToDevice),"At first mem copy");
}
// Add value 1.0 (from each array n f_dev[i]) to f_temp_dev
for (int i=0; i<nHostF; i++){
for (int zLevel=0; zLevel<num_zsteps; zLevel++){
accumulate<<<grid,threads>>>(f_dev[i],f_temp_dev,dFull,zLevel);
cudaThreadSynchronize();
}
cudaCheckError(cudaMemcpy((void *)f_temp_host,(const void *)f_temp_dev,
sizeof(float)*fMemSize,cudaMemcpyDeviceToHost),"At mem copy back");
tempsum=0.f;
count =0;
for(int k = 0 ; k< fMemSize; k++){
tempsum += f_temp_host[k];
assert ( (int)f_temp_host[k] == (i+1) );
if ( f_temp_host[k] !=(float)(i+1) ) {
printf("Found invalid return value\n");
exit(0);
}
count++;
}
printf("Total Count: %d\n",count);
printf("Real Array sum: %18f\nTotal values counted : %d\n",tempsum,count*(i+1));
printf("Calculated Array sum: %ld\n\n",(i+1)*fMemSize );
}
for(int i=0; i<nHostF; i++){
cudaFree(f_dev[i]);
}
cudaFree(f_temp_dev);
printf("Memory free. Program successfully complete\n");
delete f_dev;
delete f_temp_host;
}