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

Ответ 1

Есть две части ответа (я написал). Одна часть легко квантифицировать, другая - более эмпирическая.

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

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

• В каждом блоке не может быть больше 512/1024 потоков (Compute Capability 1.x или 2.x-3.x соответственно)
• Максимальные размеры каждого блока ограничены [512,512,64]/[1024,1024,64] (вычислить 1.x/2.x)
• Каждый блок не может потреблять больше, чем 8k/16k/32k регистров (Compute 1.0.1.1/1.2.1.3/2.x)
• Каждый блок не может потреблять более 16 КБ /48 КБ общей памяти (Compute 1.x/2.x)

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

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

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

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

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

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

Ответ 2

В приведенных выше ответах указывается, как размер блока может повлиять на производительность и предложить общую эвристику для ее выбора, основанную на максимизации занятости. Не желая предоставлять критерий выбора размера блока, стоит упомянуть, что 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 для расчета занятости, поэтому в коде API-драйвера можно использовать cuOccupancyMaxPotentialBlockSize аналогично показанному API-интерфейсу выполнения в приведенном выше примере.

Ответ 3

Размер блока обычно выбирается для максимизации "занятости". Поиск на CUDA Occupancy для получения дополнительной информации. В частности, см. Таблицу электронных калькуляций CUDA.