优化CUDA内核

Question

如何进一步优化以下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();
}

Answer 1

我将尝试使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();
}

这实际上可能会给你一些速度。