0

次の 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();
}
4

1 に答える 1

0

2D スレッド ブロックを作成して、各スレッドが 1 つの操作のみを実行するようにします。次のようなカーネルを考えてみましょう:

__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();
}

これにより、実際に速度が向上する場合があります。

于 2013-04-24T01:31:32.503 に答える