Как выбрать размеры сетки и блока для ядер CUDA?

Это вопрос о том, как определить сетку CUDA, размеры блоков и потоков. Это дополнительный вопрос к тому, что размещен здесь .

После этой ссылки ответ от talonmies содержит фрагмент кода (см. Ниже). Мне непонятен комментарий «значение, обычно выбираемое настройкой и аппаратными ограничениями».

Я не нашел хорошего объяснения или разъяснения, объясняющего это, в документации CUDA. Таким образом, мой вопрос заключается в том, как определить оптимальное blocksize(количество потоков) с учетом следующего кода:

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

Этот ответ состоит из двух частей (я его написал). Одну часть легко определить количественно, другую - более эмпирически.

Аппаратные ограничения:

Это легко измерить количественно. В приложении F к текущему руководству по программированию CUDA перечислено несколько жестких ограничений, которые ограничивают количество потоков на блок, которое может иметь запуск ядра. Если вы превысите любое из этих значений, ваше ядро ​​никогда не запустится. Их можно грубо резюмировать следующим образом:

1. В каждом блоке не может быть более 512/1024 потоков ( Compute Capability 1.x или 2.x и новее соответственно)
2. Максимальные размеры каждого блока ограничены [512,512,64] / [1024,1024,64] (для вычислений 1.x / 2.x или новее).
3. Каждый блок не может занимать более 8k / 16k / 32k / 64k / 32k / 64k / 32k / 64k / 32k / 64k регистров в сумме (вычисление 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. Каждый блок не может использовать более 16/48/96 КБ общей памяти (Compute 1.x / 2.x-6.2 / 7.0)

Если вы останетесь в этих пределах, любое ядро, которое вы сможете успешно скомпилировать, запустится без ошибок.

Настройка производительности:

Это эмпирическая часть. Количество потоков на блок, которое вы выбираете в рамках аппаратных ограничений, описанных выше, может и действительно влияет на производительность кода, выполняемого на оборудовании. Каждый код будет вести себя по-разному, и единственный реальный способ его количественно оценить - это тщательный сравнительный анализ и профилирование. Но, опять же, очень грубо резюмировал:

1. Количество потоков на блок должно быть кратным размеру деформации, который равен 32 на всем текущем оборудовании.
2. Каждый потоковый многопроцессорный блок на графическом процессоре должен иметь достаточно активных деформаций, чтобы в достаточной мере скрыть все различные задержки в памяти и конвейере команд архитектуры и достичь максимальной пропускной способности. Ортодоксальный подход здесь - попытаться достичь оптимальной загрузки оборудования (о чем говорит ответ Роджера Даля ).

Второй момент - это огромная тема, и я сомневаюсь, что кто-то попытается охватить ее в одном ответе StackOverflow. Есть люди, пишущие кандидатские диссертации по количественному анализу аспектов проблемы (см. Эту презентацию Василия Волкова из Калифорнийского университета в Беркли и эту статью Генри Вонга из Университета Торонто, где можно увидеть, насколько сложен этот вопрос на самом деле).

На начальном уровне вы должны в основном знать, что выбранный вами размер блока (в пределах диапазона допустимых размеров блока, определенного указанными выше ограничениями) может и действительно влияет на скорость выполнения вашего кода, но это зависит от оборудования. у вас есть и код, который вы используете. Путем тестирования вы, вероятно, обнаружите, что у большинства нетривиальных кодов есть «золотая середина» в диапазоне 128-512 потоков на каждый блок, но с вашей стороны потребуется некоторый анализ, чтобы найти, где это. Хорошая новость заключается в том, что, поскольку вы работаете с размерами, кратными размеру деформации, пространство поиска очень ограничено, и лучшую конфигурацию для данного фрагмента кода относительно легко найти.

В приведенных выше ответах показано, как размер блока может повлиять на производительность, и предлагается общая эвристика для его выбора, основанная на максимизации занятости. Не желая , чтобы обеспечить в критерий выбора размера блока, то стоило бы отметить , что CUDA 6.5 (сейчас в Release Candidate версии) включает в себя несколько новых функций во время выполнения , чтобы помочь в расчетах занятости и конфигурации запуска, см

Совет CUDA Pro: 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");

}

РЕДАКТИРОВАТЬ

Объект cudaOccupancyMaxPotentialBlockSizeопределяется в cuda_runtime.hфайле и определяется следующим образом:

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

Значения параметров следующие

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.

Обратите внимание, что, начиная с CUDA 6.5, необходимо вычислять собственные размеры 2D / 3D-блока из размера 1D-блока, предложенного API.

Также обратите внимание, что API драйвера CUDA содержит функционально эквивалентные API для расчета занятости, поэтому его можно использовать cuOccupancyMaxPotentialBlockSizeв коде API драйвера таким же образом, как показано для API среды выполнения в приведенном выше примере.

Размер блока обычно выбирается так, чтобы максимально «заполнить». Для получения дополнительной информации выполните поиск по CUDA Occupancy. В частности, см. Таблицу CUDA Occupancy Calculator.