Как можно оптимизировать следующее ядро CUDA? или это уже выбрано для его цели?
Я думал, может быть, я могу использовать __constant__
память в коде хоста для массивов, которые будут установлены со случайными числами. Это возможно? Я знаю, что это только для чтения памяти, поэтому я не понимаю, могу ли я использовать постоянную память вместо __global__
объем памяти.
/*
* CUDA kernel that will execute 100 threads in parallel
* and will populate these parallel arrays with 100 random numbers
* array size = 100.
*/
__global__ void initializeArrays(float* posx, float* posy,float* rayon, float* veloc,
float* opacity ,float* angle, unsigned char* color, int height,
int width, curandState* state, size_t pitch){
int idx = blockIdx.x * blockDim.x + threadIdx.x;
curandState localState = state[idx];
posx[idx] = (float)(curand_normal(&localState)*width);
posy[idx] = (float)(curand_normal(&localState)*height);
rayon[idx] = (float)(10 + curand_normal(&localState)*50);
angle[idx] = (float)(curand_normal(&localState)*360);
veloc[idx] = (float)(curand_uniform(&localState)*20 - 10);
color[idx*pitch] = (unsigned char)(curand_normal(&localState)*255);
color[(idx*pitch)+1] = (unsigned char)(curand_normal(&localState)*255);
color[(idx*pitch)+2] = (unsigned char)(curand_normal(&localState)*255);
opacity[idx] = (float)(0.3f + 1.5f *curand_normal(&localState));
__syncthreads();
}
Я постараюсь сделать блок 2D потоков и сделать так, чтобы каждый поток выполнял только одну операцию.
Рассмотрим ядро, как это:
__global__ void initializeArrays(float* posx, float* posy,float* rayon, float* veloc,
float* opacity ,float* angle, unsigned char* color, int height,
int width, curandState* state, size_t pitch){
int idx = blockIdx.x * blockDim.x + threadIdx.x;
int idy = threadIdx.y;
curandState localState = state[idy][idx];
switch(idy)
{
case 0:
posx[idx] = (float)(curand_normal(&localState)*width);
break;
case 1:
posy[idx] = (float)(curand_normal(&localState)*height);
break;
case 2:
rayon[idx] = (float)(10 + curand_normal(&localState)*50);
break;
case 3:
angle[idx] = (float)(curand_normal(&localState)*360);
break;
case 4:
veloc[idx] = (float)(curand_uniform(&localState)*20 - 10);
break;
case 5:
color[idx*pitch] = (unsigned char)(curand_normal(&localState)*255);
break;
case 6:
color[(idx*pitch)+1] = (unsigned char)(curand_normal(&localState)*255);
break;
case 7:
color[(idx*pitch)+2] = (unsigned char)(curand_normal(&localState)*255);
break;
case 8:
opacity[idx] = (float)(0.3f + 1.5f *curand_normal(&localState));
break;
default:
break;
}
__syncthreads();
}
Это может на самом деле дать вам некоторое ускорение.
Других решений пока нет …