Эффективность блоков CUDA Block и Grid

Ответ 1

Это обычно случай установки размера блока для оптимальной производительности и размера сетки в зависимости от общего объема работы. Большинство ядер имеют "сладкое пятно" количества перекосов на Мп, где они работают лучше всего, и вам нужно провести сравнительный анализ/профилирование, чтобы увидеть, где это. Вероятно, вам все еще нужна логика избыточного разлива в ядре, потому что размеры проблем редко бывают краткими размерами блоков.

EDIT: Чтобы дать конкретный пример того, как это можно сделать для простого ядра (в этом случае операция типа dscal типа 1 BLAS уровня, выполняемая как часть факторизации Cholesky упакованных симметричных матриц):

// Fused square root and dscal operation
__global__ 
void cdivkernel(const int n, double *a)
{
    __shared__ double oneondiagv;

    int imin = threadIdx.x + blockDim.x * blockIdx.x;
    int istride = blockDim.x * gridDim.x;

    if (threadIdx.x == 0) {
        oneondiagv = rsqrt( a[0] );
    }
    __syncthreads();

    for(int i=imin; i<n; i+=istride) {
        a[i] *= oneondiagv;
    }
}

Чтобы запустить это ядро, параметры выполнения вычисляются следующим образом:

• Мы разрешаем до 4 искажений на блок (так что 128 потоков). Обычно вы исправляете это при оптимальном числе, но в этом случае ядро ​​часто вызывается на очень маленьких векторах, поэтому имеет значение переменной размер блока. [/Li >
• Затем мы вычисляем количество блоков в соответствии с общим объемом работы, до 112 полных блоков, что эквивалентно 8 блокам на каждый MP на 14 MP Fermi Telsa. Ядро будет выполнять итерацию, если объем работы превышает размер сетки.

Результирующая функция-оболочка, содержащая вычисления параметров выполнения и запуск ядра, выглядит следующим образом:

// Fused the diagonal element root and dscal operation into
// a single "cdiv" operation
void fusedDscal(const int n, double *a)
{
    // The semibandwidth (column length) determines
    // how many warps are required per column of the 
    // matrix.
    const int warpSize = 32;
    const int maxGridSize = 112; // this is 8 blocks per MP for a Telsa C2050

    int warpCount = (n / warpSize) + (((n % warpSize) == 0) ? 0 : 1);
    int warpPerBlock = max(1, min(4, warpCount));

    // For the cdiv kernel, the block size is allowed to grow to
    // four warps per block, and the block count becomes the warp count over four
    // or the GPU "fill" whichever is smaller
    int threadCount = warpSize * warpPerBlock;
    int blockCount = min( maxGridSize, max(1, warpCount/warpPerBlock) );
    dim3 BlockDim = dim3(threadCount, 1, 1);
    dim3 GridDim  = dim3(blockCount, 1, 1);

    cdivkernel<<< GridDim,BlockDim >>>(n,a);
    errchk( cudaPeekAtLastError() );
}

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

Ответ 2

Хорошо, я думаю, мы имеем дело с двумя вопросами здесь.

1) Хороший способ назначить размеры блоков (т.е. количество потоков) Это обычно зависит от типа данных, с которыми вы имеете дело. Вы имеете дело с векторами? Вы имеете дело с матрицами? Предлагаемый способ заключается в том, чтобы количество потоков было кратным 32. Таким образом, при работе с векторами запуск 256 x 1, 512 x 1 блоков может быть прекрасным. Аналогично, имея дело с матрицами, 32 x 8, 32 x 16.

2) Хороший способ назначения размеров сетки (т.е. количества блоков) Здесь немного сложно. Просто запуск 10 000 блоков, потому что мы можем, как правило, не лучший способ делать что-то. Переключение блоков в и из аппаратных средств является дорогостоящим. Две вещи, которые следует учитывать, - это разделяемая память, используемая для каждого блока, и общее количество доступных SP, а также решение для оптимального числа.

Вы можете найти действительно хорошую реализацию, как это сделать из thrust. Это может занять некоторое время, чтобы выяснить, что происходит внутри кода.

Ответ 3

Я думаю, что обычно лучше всего устанавливать размеры блоков и сетки на основе заданной проблемы, особенно для целей оптимизации. Наличие дополнительных потоков, которые ничего не делают, на самом деле не имеет смысла и может ухудшить производительность ваших программ.

Ответ 4

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

Этот сайт имеет несколько эвристик. Некоторые основные моменты:

Выбор блоков в сетке

• Блоки на каждую сетку должны быть >= количество мультипроцессоров.
• Чем больше использования __syncthreads() в ваших ядрах, тем больше блоков (так что один блок может работать, пока другой ждет синхронизации).

Выбор потоков на блок

• Потоки кратных размеру основы (т.е. обычно 32)

• Обычно полезно выбирать количество потоков, так что максимальное количество потоков на блок (на основе аппаратного обеспечения) является кратным количеству потоков. Например. с максимальными потоками 768, используя 256 потоков на блок, будет лучше, чем 512, потому что несколько потоков могут выполняться одновременно на блоке.