Оптимизация ядра CUDA

Как можно оптимизировать следующее ядро ​​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();
}

0

Решение

Я постараюсь сделать блок 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();
}

Это может на самом деле дать вам некоторое ускорение.

0

Другие решения

Других решений пока нет …

По вопросам рекламы [email protected]