Jak wybrać wymiary siatki i bloków dla jądra CUDA?

To jest pytanie, jak określić rozmiary siatki, bloków i gwintów CUDA. To jest dodatkowe pytanie do zamieszczonego tutaj .

Pod tym linkiem odpowiedź z talonmies zawiera fragment kodu (patrz poniżej). Nie rozumiem komentarza „wartość zwykle wybierana przez strojenie i ograniczenia sprzętowe”.

Nie znalazłem dobrego wyjaśnienia lub wyjaśnienia, które to wyjaśnia w dokumentacji CUDA. Podsumowując, moje pytanie brzmi jak określić optymalną blocksize(liczbę wątków) biorąc pod uwagę poniższy kod:

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);

Ta odpowiedź składa się z dwóch części (napisałem ją). Jedna część jest łatwa do oszacowania, druga jest bardziej empiryczna.

Ograniczenia sprzętowe:

To jest łatwa do oszacowania część. Dodatek F aktualnego przewodnika programowania CUDA zawiera listę twardych ograniczeń, które ograniczają liczbę wątków na blok, które może mieć uruchomienie jądra. Jeśli przekroczysz którąkolwiek z tych wartości, twoje jądro nigdy nie będzie działać. Można je z grubsza podsumować jako:

1. Każdy blok nie może mieć łącznie więcej niż 512/1024 wątków (odpowiednio Compute Capability 1.x lub 2.xi nowsze)
2. Maksymalne wymiary każdego bloku są ograniczone do [512,512,64] / [1024,1024,64] (Oblicz 1.x / 2.x lub nowszy)
3. Każdy blok nie może zużywać więcej niż 8k / 16k / 32k / 64k / 32k / 64k / 32k / 64k / 32k / 64k łącznie rejestrów (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. Każdy blok nie może zużywać więcej niż 16 kb / 48 kb / 96 kb pamięci współdzielonej (Compute 1.x / 2.x-6.2 / 7.0)

Jeśli pozostaniesz w tych granicach, każde jądro, które uda Ci się pomyślnie skompilować, uruchomi się bez błędów.

Podnoszenie wydajności:

To jest część empiryczna. Liczba wątków na blok, którą wybierzesz w ramach ograniczeń sprzętowych opisanych powyżej, może wpływać i wpływa na wydajność kodu działającego na sprzęcie. Zachowanie każdego kodu będzie inne, a jedynym prawdziwym sposobem jego ilościowego określenia jest staranna analiza porównawcza i profilowanie. Ale znowu, bardzo z grubsza podsumowując:

1. Liczba wątków na blok powinna być zaokrągloną wielokrotnością rozmiaru osnowy, który wynosi 32 na całym obecnym sprzęcie.
2. Każda jednostka wieloprocesorowa do przesyłania strumieniowego w GPU musi mieć wystarczającą liczbę aktywnych wypaczeń, aby w wystarczającym stopniu ukryć całą różną pamięć i opóźnienia potoku instrukcji architektury i osiągnąć maksymalną przepustowość. Ortodoksyjne podejście polega tutaj na próbie uzyskania optymalnego wykorzystania sprzętu (do czego odnosi się odpowiedź Rogera Dahla ).

Druga kwestia to obszerny temat, w którym wątpię, aby ktokolwiek spróbował ująć go w jednej odpowiedzi StackOverflow. Są ludzie, którzy piszą prace doktorskie wokół ilościowej analizy aspektów problemu (zob. Prezentację Wasilija Wołkowa z UC Berkley i ten artykuł Henry'ego Wonga z University of Toronto, gdzie znajdują się przykłady tego, jak skomplikowane jest to pytanie).

Na poziomie początkowym powinieneś być świadomy, że wybrany rozmiar bloku (w zakresie dozwolonych rozmiarów bloków zdefiniowanych przez powyższe ograniczenia) może i ma wpływ na szybkość działania kodu, ale zależy to od sprzętu masz i kod, który uruchamiasz. Porównując testy porównawcze, prawdopodobnie okaże się, że większość nietrywialnych kodów ma „słodki punkt” w zakresie 128-512 wątków na blok, ale będzie to wymagało pewnej analizy z Twojej strony, aby znaleźć to, gdzie to jest. Dobra wiadomość jest taka, że ​​ponieważ pracujesz z wielokrotnościami rozmiaru warp, przestrzeń wyszukiwania jest bardzo ograniczona, a najlepsza konfiguracja dla danego fragmentu kodu jest stosunkowo łatwa do znalezienia.

Powyższe odpowiedzi wskazują, jak rozmiar bloku może wpływać na wydajność i sugerują typową heurystykę jej wyboru opartą na maksymalizacji obłożenia. Nie chcąc zapewnić kryterium wyboru rozmiaru bloku, warto byłoby wspomnieć, że CUDA 6.5 (obecnie w wersji Release Candidate) zawiera kilka nowych funkcji, aby pomóc w czasie wykonywania obliczeń obłożenia i konfiguracji startu, patrz

Wskazówka CUDA Pro: API Occupancy upraszcza konfigurację uruchamiania

Jedną z przydatnych funkcji jest cudaOccupancyMaxPotentialBlockSizeheurystyczna kalkulacja rozmiaru bloku, który zapewnia maksymalną liczbę osób. Wartości dostarczane przez tę funkcję można następnie wykorzystać jako punkt wyjścia do ręcznej optymalizacji parametrów uruchamiania. Poniżej znajduje się mały przykład.

#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");

}

EDYTOWAĆ

Element cudaOccupancyMaxPotentialBlockSizejest zdefiniowany w cuda_runtime.hpliku i jest zdefiniowany w następujący sposób:

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

Znaczenie parametrów jest następujące

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.

Zauważ, że od CUDA 6.5, trzeba obliczyć własne wymiary bloku 2D / 3D z rozmiaru bloku 1D sugerowanego przez API.

Należy również zauważyć, że interfejs API sterownika CUDA zawiera funkcjonalnie równoważne interfejsy API do obliczania zajętości, więc możliwe jest użycie cuOccupancyMaxPotentialBlockSizew kodzie interfejsu API sterownika w taki sam sposób, jak pokazano w przypadku interfejsu API środowiska wykonawczego w powyższym przykładzie.

Rozmiar bloku jest zwykle wybierany w celu maksymalizacji „zajętości”. Wyszukaj w CUDA Occupancy, aby uzyskać więcej informacji. W szczególności zobacz arkusz kalkulacyjny kalkulatora CUDA Occupancy Calculator.