Управление использованием в CUDA

Question

Доброе утро.

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

http://docs.nvidia.com/gameworks/content/developertools/desktop/analysis/report/cudaexperiments/kernellevel/achievedoccupancy.htm

-warps на SM (многопроцессорные системы) -блоков на СМ -зарегистрироваться в СМ -shared память на SM

Итак, я перехожу к первым вещам и, в зависимости от GPU, я определял размеры ядра в зависимости от максимальных перекосов на SM и блоков на SM. Моя задача состоит в том, чтобы выполнить тысячу миллионов сумм, чтобы определить, какой метод лучше.

То, что я делаю, это цикл for, в котором я запускаю на каждой итерации ядро, максимально увеличивающее заполняемость. Например для NVidia 1080 GPU я прочитал:

int max_blocks = 32; //maximum number of active blocks per SM int max_threads_per_Block = 64; //maximum number of active threads per SM int max_threads = 2048;

Это дает в общей сложности 2048 потоков на СМ и гарантирует максимальную вместимость. Этот GPU может иметь 64 активных искажения, каждый из которых содержит 32 потока. В этом GPU один активный блок имеет 2 искажения, и это означает, что каждый блок может иметь 64 активных потока одновременно. С этим я запустить ядро ​​следующим образом:

dim3 threadsPerBlock(max_threads_per_Block); dim3 numBlocks(max_blocks); VecAdd<<<numBlocks, threadsPerBlock>>>(d_A, d_B, d_C,max_threads);

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

int N = total_ops; //in this case one thousand millions dim3 threadsPerBlock(256); dim3 numBlocks(2*N/threadsPerBlock.x); VecAdd<<<numBlocks, threadsPerBlock>>>(d_A, d_B, d_C,);

Производительность лучше (t потребляется). Я запускаю тот же эксперимент 5 раз в одном и том же исполнении, чтобы избежать выбросов. Мой вопрос: есть ли способ управлять занятием, чтобы иметь лучшие результаты, чем то, что делает API компилятора и времени выполнения ?. Я понимаю, что оптимизация, которую я пытаюсь сделать, уже каким-то образом управляется графическим процессором. Я понимаю, что если есть документ, объясняющий, как мы должны запускать программное обеспечение (ссылка выше) для достижения хорошей производительности, это должен быть способ контролировать это.

Благодаря

Answer 1

В первом примере,

int max_blocks = 32;   //maximum number of active blocks per SM 
int max_threads_per_Block = 64; //maximum number of active threads per SM 
int max_threads = 2048; 

dim3 threadsPerBlock(max_threads_per_Block); 
dim3 numBlocks(max_blocks); 
VecAdd<<<numBlocks, threadsPerBlock>>>(d_A, d_B, d_C,max_threads); 

вы запускаете столько блоков и нитей на блок, как необходимо, чтобы полностью загрузить один SM. Однако ваш GTX 1080 имеет SMs, поэтому ваше размещение составляет только 1/20 = 5%.

Во втором примере,

int N = total_ops;    //in this case one thousand millions 
dim3 threadsPerBlock(256); 
dim3 numBlocks(2*N/threadsPerBlock.x); 
VecAdd<<<numBlocks, threadsPerBlock>>>(d_A, d_B, d_C,); 

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

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

На стороне примечание, добавление вектора в качестве алгоритма привязки к памяти не особенно хорошо подходит для демонстрации влияния занятости. Однако вы по-прежнему видите разницу, так как требуется определенное минимальное количество транзакций памяти в полете для полной загрузки подсистемы памяти (определяется пропускной способностью памяти, равной задержке доступа к памяти), а пример 5% занятости не соответствует этому минимуму ,