Как рассчитать число блоков

Question

Я пишу код CUDA, и я использую графическую карту GForce 9500 GT.

Я пытаюсь обработать массив 20000000 целочисленных элементов и число потоков Я использую 256

Размер основы составляет 32. Способность вычислить 1,1

Это аппаратные средства http://www.geforce.com/hardware/desktop-gpus/geforce-9500-gt/specifications

Теперь номер блока = 20000000/256 = 78125?

Этот звук неправильный. Как рассчитать номер блока? Любая помощь будет оценена по достоинству.

Функция ядра CUDA заключается в следующем. Идея заключается в том, что каждый блок будет вычислять свою сумму, а затем окончательная сумма будет рассчитываться путем суммирования суммы каждого блока.

__global__ static void calculateSum(int * num, int * result, int DATA_SIZE) 
{ 
    extern __shared__ int shared[]; 
    const int tid = threadIdx.x; 
    const int bid = blockIdx.x; 

    shared[tid] = 0; 
    for (int i = bid * THREAD_NUM + tid; i < DATA_SIZE; i += BLOCK_NUM * THREAD_NUM) { 
     shared[tid] += num[i]; 
    } 

    __syncthreads(); 
    int offset = THREAD_NUM/2; 
    while (offset > 0) { 
     if (tid < offset) { 
      shared[tid] += shared[tid + offset]; 
     } 
     offset >>= 1; 
     __syncthreads(); 
    } 

    if (tid == 0) { 
     result[bid] = shared[0]; 

    } 
} 

И я называю эту функцию как

calculateSum <<<BLOCK_NUM, THREAD_NUM, THREAD_NUM * sizeof(int)>>> (gpuarray, result, size); 

Где THREAD_NUM = 256 и массив является GPU размера 20000000.

Здесь я просто используя номер блока, как 16, но не уверен, что если это верно? Как я могу убедиться, что достигается максимальный параллелизм?

Вот результат моего калькулятора занятости CUDA. Он говорит, что у меня будет 100% заполняемость, когда число блоков равно 8. Таким образом, это означает, что я буду иметь максимальную эффективность, если число блоков = 8 и номер потока = 256. Это верно?

Благодаря

Answer 1

Если каждый thred процесс один элемент, и каждый блок имеет 256 потоков, вы должны запустить 20000000 потоков, в результате чего именно 78125 блоков. Это вполне допустимое число.

Однако есть небольшая проблема. У меня нет CC1.1 устройства под руку, но в CC1.3:

Maximum sizes of each dimension of a grid:  65535 x 65535 x 1 

Таким образом, вы должны либо запустить ядро ​​несколько раз для различных частей данных, или сделать 2D-сетку и просто тривиальным tranform 2D адреса от потока к 1D адресу элемента массива.

Answer 2

В вашем случае общее количество потоков (20000000) равномерно делит на количество потоков на блок (256), поэтому вы можете использовать это число (78125). Если числа не делятся равномерно, регулярное целочисленное деление округляет его, и вы получите меньше потоков, чем вам нужно. Таким образом, в этом случае вам нужно округлить результат разрыва с функцией, как в следующем:

int DivUp(int a, int b) { 
    return ((a % b) != 0) ? (a/b + 1) : (a/b); 
} 

Поскольку эта функция может дать вам больше потоков, чем Есть элементы, то вы также должны добавить тест в вашем ядре, чтобы прервать вычисления на последних нескольких потоках:

int i(blockIdx.x * blockDim.x + threadIdx.x); 
if (i >= n_items) { 
    return; 
} 

Однако есть дополнительная загвоздка. Ваше оборудование ограничено максимум 65535 блоками в каждом измерении в сетке и ограничено двумя измерениями (x и y).Поэтому, если после использования DivUp() вы получите счет, который выше, у вас есть два варианта. Вы можете разделить нагрузку и запустить ядро ​​несколько раз, или вы можете использовать два измерения.

Чтобы использовать два измерения, вы выбираете два числа, каждое из которых ниже аппаратного предела и что при умножении становится фактическим количеством блоков, которое вам нужно. Затем вы добавляете код в верхнюю часть ядра, чтобы объединить два измерения (x и y) в один индекс.

Answer 3

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

Как правило, для такого типа ядро ​​требуется столько блоков, сколько одновременно будет работать на одном мультипроцессоре, умноженное на количество многопроцессоров на карте. Первое число можно получить, используя электронную таблицу занятости CUDA, которая поставляется в наборе инструментов CUDA, но верхний предел будет 8 блоков на мультипроцессор, а второй номер будет 4 для устройства, которое у вас есть. Это означает, что для достижения максимально возможного параллелизма потребуется не более 32 блоков, но для ответа в точности требуется доступ к компилятору, которого у меня сейчас нет.

Вы также можете использовать бенчмаркинг для оптимального определения количества блоков экспериментально, используя один из 4,8,12,16,20,24,28 или 32 блоков (кратных 4, поскольку это количество многопроцессоров на вашем карта).

Answer 4

Вы используете только x-мерность сетки в своем ядре. Таким образом, вы ограничены 65535 блоками с использованием cc 1.1.

20000000/256 = 78125 - это правильно!

Таким образом, вам определенно нужно больше 1 блока.

Ядро:

//get unique block index 
const unsigned int blockId = blockIdx.x //1D 
    + blockIdx.y * gridDim.x //2D 

//terminate unnecessary blocks 
if(blockId >= 78124) 
    return; 

//... rest of kernel 

Самый простой подход будет использовать два у-блоки и проверить идентификатор блока в ядре.

dim3 gridDim = dim3(65535, 2); 

это сделало бы более 52945 блоков бесполезно, я не знаю, Что накладные расходы, но заполнение первых х, а затем у и г размера могут создать очень много неиспользуемых блоков, особенно если достижение г измерения!

(Nvidia должны определенно предусмотрели функцию полезности, получая наилучшее использование сетки для уникального использования блока внутри ядра, как его случае здесь)

Для этого простого примера, как об использовании х, у и вычисления корня ,

grid(280, 280) = 78400 blocks //only 275 blocks overhead, less is not possible 

Это одно из больших преимуществ вычислительной способности 3.0. 32-битные диапазоны на каждом блоке облегчают жизнь. Почему это было ограничено до 65535, я никогда не понимал.

Но я по-прежнему предпочитаю пониженную совместимость.

Я также испытал бы изменение @talonmies.