CUDA 커널의 그리드 및 블록 차원을 어떻게 선택합니까?

CUDA 커널의 그리드 및 블록 차원을 어떻게 선택합니까?

---

이것은 CUDA 그리드, 블록 및 스레드 크기를 결정하는 방법에 대한 질문입니다. 여기에 게시 된 질문에 대한 추가 질문입니다.

https://stackoverflow.com/a/5643838/1292251

이 링크를 따라 가면 talonmies의 답변에 코드 조각이 포함되어 있습니다 (아래 참조). 나는 "조율과 하드웨어 제약에 의해 일반적으로 선택되는 값"이라는 코멘트를 이해하지 못한다.

CUDA 문서에서 이것을 설명하는 좋은 설명이나 설명을 찾지 못했습니다. 요약하면, 내 질문은 다음 코드 에서 최적의 블록 크기 (= 스레드 수) 를 결정하는 방법입니다 .

const int n = 128 * 1024;
int blocksize = 512; // value usually chosen by tuning and hardware constraints
int nblocks = n / nthreads; // value determine by block size and total work
madd<<<nblocks,blocksize>>>mAdd(A,B,C,n);

BTW, 나는 부분적으로 내 첫 번째 질문에 대답하기 때문에 위의 링크로 내 질문을 시작했습니다. 이것이 Stack Overflow에서 질문을하는 적절한 방법이 아니라면 실례하거나 조언 해주십시오.

---

그 대답에는 두 부분이 있습니다 (내가 썼습니다). 한 부분은 정량화하기 쉽고 다른 부분은 더 경험적입니다.

하드웨어 제약 :

이것은 정량화하기 쉬운 부분입니다. 현재 CUDA 프로그래밍 가이드의 부록 F에는 커널 시작이 가질 수있는 블록 당 스레드 수를 제한하는 여러 하드 제한이 나열되어 있습니다. 이 중 하나를 초과하면 커널이 실행되지 않습니다. 대략 다음과 같이 요약 할 수 있습니다.

1. 각 블록은 총 512/1024 스레드를 초과 할 수 없습니다 (각각 컴퓨팅 기능 1.x 또는 2.x 이상).
2. 각 블록의 최대 크기는 [512,512,64] / [1024,1024,64] (Compute 1.x / 2.x 이상)로 제한됩니다.
3. 각 블록은 총 8k / 16k / 32k / 64k / 32k / 64k / 32k / 64k / 32k / 64k 레지스터 (Compute 1.0,1.1 / 1.2,1.3 / 2.x- / 3.0 / 3.2 / 3.5-5.2 / 5.3 / 6-6.1 / 6.2 / 7.0)
4. 각 블록은 16kb / 48kb / 96kb 이상의 공유 메모리를 사용할 수 없습니다 (Compute 1.x / 2.x-6.2 / 7.0).

이 한계 내에 머무르면 성공적으로 컴파일 할 수있는 모든 커널이 오류없이 시작됩니다.

성능 조정 :

이것은 경험적인 부분입니다. 위에서 설명한 하드웨어 제약 조건 내에서 선택한 블록 당 스레드 수는 하드웨어에서 실행되는 코드의 성능에 영향을 미칠 수 있으며 영향을 미칩니다. 각 코드의 작동 방식은 다르며이를 정량화하는 유일한 방법은 신중한 벤치마킹 및 프로파일 링입니다. 그러나 다시 대략적으로 요약하면 다음과 같습니다.

1. 블록 당 스레드 수는 모든 현재 하드웨어에서 32 인 워프 크기의 반올림 배수 여야합니다.
2. GPU의 각 스트리밍 멀티 프로세서 장치에는 아키텍처의 모든 다른 메모리 및 명령어 파이프 라인 대기 시간을 충분히 숨기고 최대 처리량을 달성 할 수있는 충분한 활성 워프가 있어야합니다. 여기서 정통 접근 방식은 최적의 하드웨어 점유율을 달성하는 것입니다 ( Roger Dahl의 답변 이 참조하는 것).

두 번째 요점은 누군가가 하나의 StackOverflow 답변에서 그것을 시도하고 다룰 것이라고 의심하는 거대한 주제입니다. 문제 측면의 정량적 분석을 중심으로 박사 학위 논문을 쓰는 사람들이 있습니다 ( 이 질문이 실제로 얼마나 복잡한 지에 대한 예는 UC Berkley의 Vasily Volkov의 프레젠테이션 과 토론토 대학의 Henry Wong 의이 논문 참조).

엔트리 레벨에서는 선택한 블록 크기 (위의 제약 조건에 정의 된 합법적 인 블록 크기 범위 내)가 코드 실행 속도에 영향을 미칠 수 있고 영향을 미칠 수 있지만 하드웨어에 따라 다릅니다. 가지고있는 코드와 실행중인 코드. 벤치마킹을 통해 대부분의 사소하지 않은 코드가 블록 당 128-512 스레드 범위에서 "스위트 스팟"을 가지고 있음을 발견 할 수 있지만 그 위치를 찾으려면 일부 분석이 필요합니다. 좋은 소식은 워프 크기의 배수로 작업하기 때문에 검색 공간이 매우 한정되어 있으며 주어진 코드 조각에 대한 최상의 구성을 비교적 쉽게 찾을 수 있다는 것입니다.

---

위의 답변은 블록 크기가 성능에 미치는 영향을 지적하고 점유 최대화를 기반으로 선택한 공통 휴리스틱을 제안합니다. 블록 크기를 선택 하는 기준 을 제공하고 싶지 는 않지만 CUDA 6.5 (현재 출시 후보 버전)에는 점유 계산 및 시작 구성을 지원하는 몇 가지 새로운 런타임 기능이 포함되어 있다는 점을 언급 할 가치가 있습니다.

CUDA Pro 팁 : Occupancy API로 시작 구성 간소화

유용한 기능 중 하나는 cudaOccupancyMaxPotentialBlockSize최대 점유율을 달성하는 블록 크기를 경험적으로 계산하는 것입니다. 그런 다음 해당 함수에서 제공하는 값을 시작 매개 변수의 수동 최적화의 시작점으로 사용할 수 있습니다. 아래는 약간의 예입니다.

#include <stdio.h>

/************************/
/* TEST KERNEL FUNCTION */
/************************/
__global__ void MyKernel(int *a, int *b, int *c, int N) 
{ 
    int idx = threadIdx.x + blockIdx.x * blockDim.x; 

    if (idx < N) { c[idx] = a[idx] + b[idx]; } 
} 

/********/
/* MAIN */
/********/
void main() 
{ 
    const int N = 1000000;

    int blockSize;      // The launch configurator returned block size 
    int minGridSize;    // The minimum grid size needed to achieve the maximum occupancy for a full device launch 
    int gridSize;       // The actual grid size needed, based on input size 

    int* h_vec1 = (int*) malloc(N*sizeof(int));
    int* h_vec2 = (int*) malloc(N*sizeof(int));
    int* h_vec3 = (int*) malloc(N*sizeof(int));
    int* h_vec4 = (int*) malloc(N*sizeof(int));

    int* d_vec1; cudaMalloc((void**)&d_vec1, N*sizeof(int));
    int* d_vec2; cudaMalloc((void**)&d_vec2, N*sizeof(int));
    int* d_vec3; cudaMalloc((void**)&d_vec3, N*sizeof(int));

    for (int i=0; i<N; i++) {
        h_vec1[i] = 10;
        h_vec2[i] = 20;
        h_vec4[i] = h_vec1[i] + h_vec2[i];
    }

    cudaMemcpy(d_vec1, h_vec1, N*sizeof(int), cudaMemcpyHostToDevice);
    cudaMemcpy(d_vec2, h_vec2, N*sizeof(int), cudaMemcpyHostToDevice);

    float time;
    cudaEvent_t start, stop;
    cudaEventCreate(&start);
    cudaEventCreate(&stop);
    cudaEventRecord(start, 0);

    cudaOccupancyMaxPotentialBlockSize(&minGridSize, &blockSize, MyKernel, 0, N); 

    // Round up according to array size 
    gridSize = (N + blockSize - 1) / blockSize; 

    cudaEventRecord(stop, 0);
    cudaEventSynchronize(stop);
    cudaEventElapsedTime(&time, start, stop);
    printf("Occupancy calculator elapsed time:  %3.3f ms \n", time);

    cudaEventRecord(start, 0);

    MyKernel<<<gridSize, blockSize>>>(d_vec1, d_vec2, d_vec3, N); 

    cudaEventRecord(stop, 0);
    cudaEventSynchronize(stop);
    cudaEventElapsedTime(&time, start, stop);
    printf("Kernel elapsed time:  %3.3f ms \n", time);

    printf("Blocksize %i\n", blockSize);

    cudaMemcpy(h_vec3, d_vec3, N*sizeof(int), cudaMemcpyDeviceToHost);

    for (int i=0; i<N; i++) {
        if (h_vec3[i] != h_vec4[i]) { printf("Error at i = %i! Host = %i; Device = %i\n", i, h_vec4[i], h_vec3[i]); return; };
    }

    printf("Test passed\n");

}

편집하다

The cudaOccupancyMaxPotentialBlockSize is defined in the cuda_runtime.h file and is defined as follows:

template<class T>
__inline__ __host__ CUDART_DEVICE cudaError_t cudaOccupancyMaxPotentialBlockSize(
    int    *minGridSize,
    int    *blockSize,
    T       func,
    size_t  dynamicSMemSize = 0,
    int     blockSizeLimit = 0)
{
    return cudaOccupancyMaxPotentialBlockSizeVariableSMem(minGridSize, blockSize, func, __cudaOccupancyB2DHelper(dynamicSMemSize), blockSizeLimit);
}

The meanings for the parameters is the following

minGridSize     = Suggested min grid size to achieve a full machine launch.
blockSize       = Suggested block size to achieve maximum occupancy.
func            = Kernel function.
dynamicSMemSize = Size of dynamically allocated shared memory. Of course, it is known at runtime before any kernel launch. The size of the statically allocated shared memory is not needed as it is inferred by the properties of func.
blockSizeLimit  = Maximum size for each block. In the case of 1D kernels, it can coincide with the number of input elements.

Note that, as of CUDA 6.5, one needs to compute one's own 2D/3D block dimensions from the 1D block size suggested by the API.

Note also that the CUDA driver API contains functionally equivalent APIs for occupancy calculation, so it is possible to use cuOccupancyMaxPotentialBlockSize in driver API code in the same way shown for the runtime API in the example above.

---

The blocksize is usually selected to maximize the "occupancy". Search on CUDA Occupancy for more information. In particular, see the CUDA Occupancy Calculator spreadsheet.

참고URL : https://stackoverflow.com/questions/9985912/how-do-i-choose-grid-and-block-dimensions-for-cuda-kernels