优化CUDA内核

Optimizing CUDA Kernel

本文关键字：内核 CUDA 优化      更新时间：2023-10-16

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

相关文章：