Как выбрать размеры сетки и блока для ядер CUDA? - PullRequest
94 голосов
/ 03 апреля 2012

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

https://stackoverflow.com/a/5643838/1292251

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

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

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

Кстати, я начал свой вопрос с ссылки выше, потому что эточастично отвечает на мой первый вопрос.Если это не правильный способ задать вопросы о переполнении стека, извините или сообщите мне.

Ответы [ 3 ]

136 голосов
/ 03 апреля 2012

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

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

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

  1. В каждом блоке может быть не более 512/1024 потоков ( Compute Capability 1.x или 2.x и более поздних версий соответственно)
  2. Максимальные размеры каждого блока ограничены [512,512,64] / [1024,1024,64] (вычисление 1.x / 2.x или более поздняя)
  3. В каждом блоке не может быть больше регистров 8К / 16К / 32К / 64К / 32К / 64К / 32К / 64К / 32К / 64К (Вычислить 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 потоках на диапазон блоков, но с вашей стороны потребуется некоторый анализ, чтобы найти, где это находится. Хорошая новость заключается в том, что, поскольку вы работаете с кратными размерами деформации, пространство поиска очень ограничено, и лучшую конфигурацию для данного фрагмента кода относительно легко найти.

33 голосов
/ 29 июля 2014

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

Совет CUDA Pro: API Occupancy упрощает настройку запуска

Одна из полезных функций - 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 среды выполнения.в приведенном выше примере.

10 голосов
/ 03 апреля 2012

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

...