Cuda gridDim i blockDim

Parafrazowane z CUDA Programming Guide :

GridDim: ta zmienna zawiera wymiary siatki.

BlockIdx: ta zmienna zawiera indeks bloków w siatce.

BlockDim: ta zmienna i zawiera wymiary bloku.

ThreadIdx: ta zmienna zawiera indeks wątku w bloku.

Wydaje ci się, że jesteś trochę zdezorientowany wątkiem hierachy, który CUDA ma; w w skrócie, dla jądra będzie 1 Siatka (którą zawsze wizualizuję jako trójwymiarową kostkę). Każdy z jego elementów jest blokiem, takim, że siatka zadeklarowana jako dim3 grid(10, 10, 2); miałaby 10*10*2 całkowita liczba bloków. Z kolei każdy blok jest trójwymiarową sześcianem wątków.

Z tym, że powiedział, To jest powszechne używać tylko x-wymiar bloków i siatek, co wygląda jak kod w twoim pytaniu robi. Jest to szczególnie przydatne w przypadku pracy z tablicami 1D. W takim przypadku twój tid+=blockDim.x * gridDim.x linia będzie w efekcie unikalnym indeksem każdego wątku w Twojej siatce. To dlatego, że blockDim.x będzie wielkością każdego bloku, a gridDim.x będzie całkowitą liczbą bloków.

Więc jeśli uruchomisz jądro z parametrami

dim3 block_dim(128,1,1);
dim3 grid_dim(10,1,1);
kernel<<<grid_dim,block_dim>>>(...);

Wtedy w jądrze miałbyś threadIdx.x + blockIdx.x*blockDim.x będziesz miał:

threadIdx.x range from [0 ~ 128)

blockIdx.x range from [0 ~ 10)

blockDim.x equal to 128

gridDim.x equal to 10

Stąd przy obliczaniu threadIdx.x + blockIdx.x*blockDim.x mielibyśmy wartości w zakresie określonym przez: [0, 128) + 128 * [1, 10), co oznaczałoby, że twoje wartości tid wahają się od {0, 1, 2,..., 1279}. Jest to przydatne, gdy chcesz mapować wątki do zadań, ponieważ zapewnia to unikalny identyfikator dla wszystkich wątków w jądrze.

Jednakże, jeśli masz

int tid = threadIdx.x + blockIdx.x * blockDim.x;
tid += blockDim.x * gridDim.x;

Wtedy będziesz miał: tid = [0, 128) + 128 * [1, 10) + (128 * 10), a twoje wartości tid będą wynosić od {1280, 1281,..., 2559} Nie jestem pewien, gdzie byłoby to istotne, ale wszystko zależy od aplikacji i sposobu mapowania wątków do danych. To mapowanie jest dość centralne dla każdego uruchomienia jądra i to Ty decydujesz, jak to zrobić. Kiedy uruchamiasz jądro, określasz wymiary siatki i bloku i to Ty musisz wymusić mapowanie do swoich danych wewnątrz jądra. Jeśli nie przekroczysz limitów sprzętowych (w przypadku nowoczesnych kart możesz mieć maksymalnie 2^10 wątków na blok i 2^16-1 bloki na siatkę) {]}

W tym kodzie źródłowym mamy nawet 4 thredy, funkcja jądra może uzyskać dostęp do wszystkich 10 tablic. Jak?

#define N 10 //(33*1024)

__global__ void add(int *c){
    int tid = threadIdx.x + blockIdx.x * gridDim.x;

    if(tid < N)
        c[tid] = 1;

    while( tid < N)
    {
        c[tid] = 1;
        tid += blockDim.x * gridDim.x;
    }
}

int main(void)
{
    int c[N];
    int *dev_c;
    cudaMalloc( (void**)&dev_c, N*sizeof(int) );

    for(int i=0; i<N; ++i)
    {
        c[i] = -1;
    }

    cudaMemcpy(dev_c, c, N*sizeof(int), cudaMemcpyHostToDevice);

    add<<< 2, 2>>>(dev_c);
    cudaMemcpy(c, dev_c, N*sizeof(int), cudaMemcpyDeviceToHost );

    for(int i=0; i< N; ++i)
    {
        printf("c[%d] = %d \n" ,i, c[i] );
    }

    cudaFree( dev_c );
}

Dlaczego nie tworzymy 10 wątków ex) add<<<2,5>>> or add<5,2>>> Ponieważ musimy utworzyć dość małą liczbę wątków, jeśli N jest większe niż 10 ex) 33*1024.

Ten kod źródłowy jest przykładem tego przypadku. macierze są 10, wątki cuda są 4. Jak uzyskać dostęp do wszystkich 10 tablic tylko przez 4 wątki.

Zobacz stronę o znaczeniu threadIdx, blockIdx, blockDim, gridDim w cuda szczegóły.

W tym kodzie źródłowym,

gridDim.x : 2    this means number of block of x

gridDim.y : 1    this means number of block of y

blockDim.x : 2   this means number of thread of x in a block

blockDim.y : 1   this means number of thread of y in a block

Nasza liczba wątków to 4, Ponieważ 2*2(bloki * wątek).

W funkcji add kernel możemy uzyskać dostęp do indeksu 0, 1, 2, 3 wątku

->tid = threadIdx.x + blockIdx.x * blockDim.x

①0+0*2=0

②1+0*2=1

③0+1*2=2

④1+1*2=3

Jak uzyskać dostęp do reszty indeksu 4, 5, 6, 7, 8, 9. W pętli while

tid += blockDim.x + gridDim.x in while

** pierwsze wywołanie jądra * *

-1 pętla: 0+2*2=4

-2 pętla: 4+2*2=8

-3 pętla: 8+2*2=12 (ale ta wartość jest fałszywa, podczas gdy out!)

** drugie wywołanie jądra * *

-1 pętla: 1+2*2=5

-2 pętla: 5+2*2=9

-3 pętla: 9+2*2=13 (ale ta wartość jest fałszywa, podczas gdy out!)

** third call of kernel * *

-1 pętla: 2+2*2=6

-2 pętla: 6+2*2=10 (ale ta wartość jest fałszywa, podczas gdy out!)

** czwarte wywołanie jądra * *

-1 pętla: 3+2*2=7

-2 pętla: 7+2*2=11 (ale ta wartość jest fałszywa, while out!)

Więc wszystkie indeksy 0, 1, 2, 3, 4, 5, 6, 7, 8, 9 można uzyskać dostęp przez wartość tid.

Zobacz tę stronę. http://study.marearts.com/2015/03/to-process-all-arrays-by-reasonably.html Nie mogę przesłać obrazu, ponieważ niska reputacja.

Najpierw patrz rysunek Siatka bloków wątków z oficjalnego dokumentu CUDA

Zwykle używamy jądra w następujący sposób:

__global__ void kernelname(...){
    const id_x = blockDim.x * blockIdx.x + threadIdx.x;
    const id_y = blockDim.y * blockIdx.y + threadIdx.y;
    ...
}

// invoke kernel
// assume we have assigned the proper gridsize and blocksize
kernelname<<<gridsize, blocksize>>>(...)

Znaczenie niektórych zmiennych:

gridsize Liczba bloków na siatkę, odpowiadająca gridDim

blocksize Liczba wątków w bloku, odpowiadająca blockDim

threadIdx.x waha się w [0, blockDim.x)

blockIdx.x zmienia się w [0, gridDim.x)

Spróbujmy więc obliczyć indeks w kierunku x , gdy mamy threadIdx.x i blockIdx.x. Według rysunku , blockIdx.x określa, którym blokiem jesteś, a threadIdx.x określa, którym wątkiem jesteś, gdy podano lokalizację bloku. Stąd mamy:

which_blk = blockDim.x * blockIdx.x; // which block you are
final_index_x = which_blk + threadIdx.x; // based on the given block, we can have the final location by adding the threadIdx.x

Czyli:

final_index_x = blockDim.x * blockIdx.x + threadIdx.x;

Który jest taki sam jak powyższy przykładowy kod.

Podobnie możemy uzyskać indeks odpowiednio w kierunku y lub z .

Jak widzimy, zwykle nie używamy gridDim w naszym kodzie, ponieważ to informacje są wykonywane jako zakres blockIdx. Przeciwnie, musimy użyć blockDim, chociaż ta informacja jest wykonywana jako zakres threadIdx. Powód pokazałem powyżej Krok po kroku.

Mam nadzieję, że ta odpowiedź pomoże rozwiązać Twoje zamieszanie.