JCUDA(CUDA) grid, block ?

CUDA를 접하기전에 막연히 CUDA만 설치하고 GPU에 코드를 전달하면 병렬처리가 자동으로 되는지 알았습니다.
만만한 세상은 없나 봅니다. 병렬 처리할 수 있도록 코드를 만든 뒤 GPU로 보내는 작업을 해야한다고 합니다.
CPU로 보낼때는 앞에서도 언급했지만 cu 파일을 ptx로 컴파일 해서 해당 파일을 Driver API를 통해서 실행시키면 됩니다.
그 후 해당 ptx파일이 동시에 실행되어 thread들이 돌기 시작하고 모든 thread들이 종료되면 끝나게 됩니다.

cu(Kernel)에 넣었던 코드들이 동시에 몇개나 어떤 방식으로 실행할지 결정하는게 block, grid라는 개념이 있습니다. 해당 개념을 접하고 이건 뭘까 하던 생각이 납니다. 하루 동안 계속 검색해보며 정체를 찾아봤었습니다. 아직도 이해하기 힘든 부분인것 같네요.

Grid안에는 여러개의 블럭이 존재하고, Block 안에는 여러개의 thread가 존재합니다. (존재할 수 있는 thread갯수는 GPU마다 다릅니다. GTX 750 1024개)
그리고 Grid, Block 모두 3 차원 구조를 외부에서 설정할 수 있습니다. (차원의 최대값도 GPU에 따라 다릅니다.)

Block의 차원을 4*2*1, Grid의 차원을 3*2*1 로 했다면 아래와 같이 표현됩니다.
그림에서 꼬부랑 화살표 하나가 thread가 됩니다.

기존에 있던 예제입니다.

        int blockSizeX = 256;
        int gridSizeX = (int)Math.ceil((double)numElements / blockSizeX);
        cuLaunchKernel(function,
            gridSizeX,  1, 1,      // Grid dimension
            blockSizeX, 5, 1,      // Block dimension
            0, null,               // Shared memory size and stream
            kernelParameters, null // Kernel- and extra parameters
        );
        cuCtxSynchronize();

위와같이 blockSizeX*5*1=256*5*1=1280이 되면 GTX 750에서는 블럭당1024개의 thread만 만들수 있기 때문에 오류가 납니다.

Exception in thread "main" jcuda.CudaException: CUDA_ERROR_INVALID_VALUE
at jcuda.driver.JCudaDriver.checkResult(JCudaDriver.java:288)
at jcuda.driver.JCudaDriver.cuLaunchKernel(JCudaDriver.java:14366)
at CUDAtest.main(CUDAtest.java:90)

예제 프로그램을 바꿔서 grid 차원을 변경하였습니다.

        // Call the kernel function.
        int blockSizeX = 256;
        int gridSizeX = (int)Math.ceil((double)numElements / blockSizeX);
        System.out.println("gridsize:"+gridSizeX+" "+(gridSizeX*1024*100));
        cuLaunchKernel(function,
            gridSizeX,  1024, 100,      // Grid dimension
            blockSizeX, 1, 1,      // Block dimension
            0, null,               // Shared memory size and stream
            kernelParameters, null // Kernel- and extra parameters
        );

결과 그리드 차원이 무려 40,038,400 이 숫자가 결국 block가 되고 block에 다시 블럭 차원 숫자를 곱하면 엄청난 thread수가 실행되는 것인데 GTX 750에서 순식간에 처리가 됩니다. block은 동시에 처리가 되겠지만 grid는 동시에 처리 되지는 않을꺼라고 생각됩니다. h/w적으로 그만한 thread를 실행할 공간이 없기 때문에 처리하고 나서 다음 block을 구동하는 방법을 사용하지 않을까 추측해봅니다.
gridsize:391 40038400
Test PASSED

그러면 병렬 프로그래밍을 어떻게 해야할까요?
답은 cu파일 Kernel에 있습니다. 예제를 보시죠
blockIdx threadIdx blockDim 이런 변수를 이용해서 차원과 현재의 위치를 구해서 연산 결과를 특정 메모리에 넣는 작업을 하게됩니다. 해당 예제는 block,grid 모두 1차원으로 실행시키기 때문에 blockDim.x를 이용해서 block의 크기를 구해서 block index를 곱한후 threadidx를 더해주면 고유한 i위치가 나오게 됩니다.

extern "C"
__global__ void add(int n, float *a, float *b, float *sum)
{
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i<n)
    {
        sum[i] = a[i] + b[i];
    }

}

그리고 CPU를 host라고 표현하고 GPU를 device라고 표현합니다.
메모리 할당을 device쪽에 해서 host와 device간 메모리를 전달할때 cuMemcpyDtoH ,cuMemcpyHtoD 이런 함수를 이용하게 됩니다.