Bagaimana cara memilih dimensi grid dan blok untuk kernel CUDA?

Ini adalah pertanyaan tentang bagaimana menentukan ukuran grid, blok, dan utas CUDA. Ini adalah pertanyaan tambahan untuk pertanyaan yang diposting di sini .

Mengikuti tautan ini, jawaban dari talonmies berisi potongan kode (lihat di bawah). Saya tidak mengerti komentar "nilai yang biasanya dipilih oleh tuning dan batasan perangkat keras".

Saya belum menemukan penjelasan atau klarifikasi bagus yang menjelaskan hal ini dalam dokumentasi CUDA. Singkatnya, pertanyaan saya adalah bagaimana menentukan blocksize(jumlah utas) optimal dengan kode berikut:

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

Ada dua bagian untuk jawaban itu (saya menulisnya). Satu bagian mudah diukur, bagian lainnya lebih empiris.

Batasan Perangkat Keras:

Ini adalah bagian yang mudah diukur. Lampiran F dari panduan pemrograman CUDA saat ini mencantumkan sejumlah batas keras yang membatasi berapa banyak utas per blok yang dapat dimiliki peluncuran kernel. Jika Anda melebihi salah satu dari ini, kernel Anda tidak akan pernah berjalan. Mereka secara kasar dapat diringkas sebagai:

1. Setiap blok tidak boleh memiliki total lebih dari 512/1024 utas ( Kemampuan Hitung 1.x atau 2.x dan yang lebih baru masing-masing)
2. Dimensi maksimum setiap blok dibatasi hingga [512.512,64] / [1024,1024,64] (Hitung 1.x / 2.x atau lebih baru)
3. Setiap blok tidak dapat mengkonsumsi lebih dari 8k / 16k / 32k / 64k / 32k / 64k / 32k / 64k / 32k / 64k total register (Hitung 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. Setiap blok tidak dapat menggunakan lebih dari 16kb / 48kb / 96kb memori bersama (Hitung 1.x / 2.x-6.2 / 7.0)

Jika Anda tetap berada dalam batasan tersebut, semua kernel yang berhasil Anda kompilasi akan diluncurkan tanpa kesalahan.

Penyetelan Performa:

Ini adalah bagian empirisnya. Jumlah utas per blok yang Anda pilih dalam batasan perangkat keras yang diuraikan di atas dapat dan memang memengaruhi kinerja kode yang berjalan pada perangkat keras. Bagaimana setiap kode berperilaku akan berbeda dan satu-satunya cara nyata untuk mengukurnya adalah dengan pembandingan dan pembuatan profil yang cermat. Tapi sekali lagi, dirangkum dengan sangat kasar:

1. Jumlah utas per blok harus kelipatan bulat dari ukuran warp, yaitu 32 pada semua perangkat keras saat ini.
2. Setiap unit streaming multiprosesor pada GPU harus memiliki cukup aktif warps untuk cukup menyembunyikan semua memori yang berbeda dan latensi pipeline instruksi dari arsitektur dan mencapai throughput maksimum. Pendekatan ortodoks di sini adalah mencoba mencapai penempatan perangkat keras yang optimal (yang dirujuk oleh jawaban Roger Dahl ).

Poin kedua adalah topik besar yang saya ragu ada orang yang akan mencoba dan membahasnya dalam satu jawaban StackOverflow. Ada orang yang menulis tesis PhD seputar analisis kuantitatif aspek masalah (lihat presentasi ini oleh Vasily Volkov dari UC Berkley dan makalah ini oleh Henry Wong dari Universitas Toronto untuk contoh betapa rumitnya pertanyaan tersebut sebenarnya).

Pada level awal, Anda harus menyadari bahwa ukuran blok yang Anda pilih (dalam kisaran ukuran blok legal yang ditentukan oleh batasan di atas) dapat dan memang berdampak pada seberapa cepat kode Anda akan berjalan, tetapi itu tergantung pada perangkat kerasnya. yang Anda miliki dan kode yang Anda jalankan. Dengan pembandingan, Anda mungkin akan menemukan bahwa sebagian besar kode non-sepele memiliki "sweet spot" di 128-512 utas per rentang blok, tetapi akan memerlukan beberapa analisis di pihak Anda untuk menemukan di mana itu. Kabar baiknya adalah karena Anda bekerja dalam kelipatan ukuran warp, ruang pencarian sangat terbatas dan konfigurasi terbaik untuk potongan kode tertentu relatif mudah ditemukan.

Jawaban di atas menunjukkan bagaimana ukuran blok dapat memengaruhi kinerja dan menyarankan heuristik umum untuk pilihannya berdasarkan maksimalisasi hunian. Tanpa ingin memberikan yang kriteria untuk memilih ukuran blok, itu akan menjadi layak disebut bahwa CUDA 6.5 (sekarang dalam rilis versi Candidate) meliputi beberapa fungsi runtime baru untuk membantu dalam perhitungan hunian dan konfigurasi peluncuran, lihat

Tip CUDA Pro: Occupancy API Menyederhanakan Konfigurasi Peluncuran

Salah satu fungsi yang berguna adalah cudaOccupancyMaxPotentialBlockSizeyang secara heuristik menghitung ukuran blok yang mencapai hunian maksimum. Nilai-nilai yang disediakan oleh fungsi tersebut kemudian dapat digunakan sebagai titik awal pengoptimalan manual parameter peluncuran. Di bawah ini adalah contoh kecil.

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

}

EDIT

The cudaOccupancyMaxPotentialBlockSizedidefinisikan dalam cuda_runtime.hberkas dan didefinisikan sebagai berikut:

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

Arti dari parameter adalah sebagai berikut

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.

Perhatikan bahwa, pada CUDA 6.5, seseorang perlu menghitung dimensi blok 2D / 3D miliknya sendiri dari ukuran blok 1D yang disarankan oleh API.

Perhatikan juga bahwa CUDA driver API berisi API yang memiliki fungsi setara untuk penghitungan okupansi, sehingga memungkinkan untuk digunakan cuOccupancyMaxPotentialBlockSizedalam kode API driver dengan cara yang sama seperti yang ditunjukkan untuk API runtime pada contoh di atas.

Ukuran blok biasanya dipilih untuk memaksimalkan "hunian". Cari di CUDA Occupancy untuk informasi lebih lanjut. Secara khusus, lihat spreadsheet Kalkulator Hunian CUDA.