Разъяснение по CUDA

Конфигурация блоков и грида кроится под задачу (problem). Если задача двухмерная, то логично делать двухмерные блоки и двухмерный грид. Тогда два индекса нитей в пределах блока и два индекса блоков в пределах грида удобно использовать как индексы в задаче.

Ну например хочу умножить 100000 чисел, какую конфигурацию юзать?

Disclaimer: я отстал от жизни во времена Fermi где-то.

ЕМНИП, это число зависит от количества регистров, которое твой алгоритм использует, т.к. регистровая память шарится в пределах варпа.

Для таких вещей проще юзать thrust и не парить мозг.

Одномерные блоки по 256. Или просто взять прикладные либы из SDK (NPP, thrust).

Поясни про многомернось плиз. А thrust он для С++, а у мне голый си.

Для голого C там библиотека NPP есть. А что про многомерность пояснять? Грид и блоки максимум трёхмерные, не заблудишься :)

я отстал от жизни во времена Fermi где-то.

А потом на OpenCL перешли или просто на другую работу?

Если ты не просто изучаешь апи, нужно заметить что твоя квадра настолько убогая затычка, что сливает даже просто затычкам в дофига раз: http://gpu.userbenchmark.com/Compare/Nvidia-Quadro-FX-1800-vs-Nvidia-GTX-1050...

Я в курсе, просто это единственное поделит от нвидиа которое я нашел у себя.

просто на другую работу

this

Если задача двухмерная, то логично

Как понять двумерная ли задача? Есть ли метода принятия решения кроме научного тыка: попробовать и то и другое? Например свертка изображения с небольшим ядром (радиус до 3) при реализации с буферами, а не текстурами работает быстрее.

для разных архитектур (обычно мапится на версию CUDA) максимальное значение может быть разным. Например для CUDA до 1.3 (как раз для твоей карточки) максимальное кол-во threads в блоке 512, а блоков 0xffff.

Причем учти - если у тебя размерность задачи будет 2 (максимальная для твоей видяхи), то это значит, что суммарное количество threads не может быть больше 512. т.е. 512х512 не сработает, но 16х32 - ок.

Для определенной карточки можно рассчитать максимальное (при 100% occupancy) кол-во threads в данный момент:

NumCudaCores*WarpSize*MaxWarpsInFlight = 64*32*24 (для твоего случая)

• WarpSize=32 всегда было и есть для нвидии.
• NumCudaCores - количество грубо говоря, SM(streaming multiprocessor) юнитов, (на сколько крута твоя видяха)
• MaxWarpsInFlight также архитектурнозависимо

(Если твои кернелы используют много локальной памяти или регистров, то реальное кол-во warps in-flight < MaxWarpsInFlight. Соответственно страдает и производительность, так как latency hiding работает хуже)

Хочу еще сказать, что кол-во threads в блоке имеет смысл делать только кратным WarpSize, иначе будет не полная утилизация.

Если ты просто хочешь умножить массив чисел (100000 штук), делай одномерную задачу. Кол-во threads, например, выбери 32 => кол-во блоков будет (100000 + 32-1) / 32;

В своем kernel, ставь проверку:

__global__ void DoStuff(const int* array, int* out_array, int N)
{
    int index = blockDim.x * blockIdx.x + threadIdx.x;
    if (index < N)
        out_array[i] = 2 * array[i];
}

Размерность 2 имеет смысл, например, для обработки изображений. Вообще, это сугубо для удобства, чтобы не заморачиваться с вычислением 2D индекса из линейного массива. Размерность принципиально ничего не меняет.

У нвидии есть куча доков, там все очень красиво и понятно расписано.

Кол-во threads, например, выбери 32 => кол-во блоков будет (100000 + 32-1) / 32;

Поясни, 100031/32 Хм... ????

Поясняю. Обычное выравнивание кол-ва блоков.

Исходное кол-во threads N = 100000

по формуле имеем:

(N+32-1) = 100031/32 = 3125 - это необходимое количество блоков. Замечу, что N/32 = 3125 - дает аналогичный результат.

Однако, допустим, что исходное кол-во threads N = 100001. В данном случае, если просто воспольоватося вормулой N/32, получим также 3125. Теперь если мы запустим 3125 блоков и каждый блок содержит 32 threads, то выполнится всего 3125*32 = 100000 threads, что меньше чем исходное кол-во, поэтому, используя вышеприведенную формулу, можно получить минимально необходимое кол-во блоков, а именно: (10001+31)/32 = 3126. Итого, будет запущено 3126*32 = 100032 threads. Добавив проверку if(index < N), мы решим проблему чтения/записи за границы нашего массива. В некоторых случаях сам массив данных выравнивают, чтобы не писать проверку.

Спасибо