оптимизация кода CUDA с помощью веток
У меня есть код CUDA, который я хочу оптимизировать. Мое ядро работает с dim3 grid=(35,48) а также dim3 threads=(18,18), Прежде всего, каждый блок выполняет независимые 290 векторных вычислений, где каждый поток выполняет 1 векторное вычисление (что составляет 1024 сложения-умножения).
Однако входные данные для первых 17 * 17 = 289 этих вычислений сохраняются в общем массиве im1, а данные для последних сохраняются в im2 (также выходные массивы отличаются). После этого я использую все полученные данные для дальнейших расчетов.
Я реализовал это следующим образом:
if ((threadIdx.x < 17) && (threadIdx.y < 17)){
**instructions for 289s vector calculations**
}
else if ((threadIdx.x == 17) && (threadIdx.y == 17)){
**instruction for 290 vector calculation**
}
__syncthreads();
***further calculations***
Итак, если я правильно понимаю, мои первые 289 следуют за 1 веткой, а поток # 324 следует за другой. Пока первая группа потоков находится в основе # 0,1, .., 10, а поток # 324 находится в основе # 11, расходящихся ветвей нет. Тем не менее, я читал, что, как правило, лучше избегать любых if операторы в таких ядрах и замените их на пошаговый индекс или что-то в этом роде. Итак, я могу как-то улучшить этот код?
Мой графический процессор GTX 980 с cc 5.2, я использую VS2013 для кодирования.
Спасибо Михаил
Решение
Рассмотрим блок из 18 * 18 потоков, пронумерованных от 0 (0, 0) до 323 (17, 17).
So, if I understand correct, my first 289 follow 1 branch [...]
Если под «первыми 289» вы ссылаетесь на потоки, пронумерованные от 0 (0, 0) до 288 (16, 16), то нет, не все они принимают первую ветвь. Например, поток 17 (0, 17) не принимает ветвь (см. Иллюстрацию ниже). Однако за промежуток блока 289 потоков действительно принимают эту ветвь.
[...] and thread #324 follows another
Это верно, поток 323 (17, 17) занимает вторую ветвь.
Темы 17 (0, 17), 35 (1, 17) … 305 (16, 17) и 306 (17, 0), 307 (17, 1) … 322 (17, 16) (всего 35 тем ) не бери какую-то ветку и впустую. С точки зрения производительности это плохо, но и не очень катастрофично.
Но рассмотрим следующую схему того, что вы делаете:
0 1 2 … 15 16 17
0 * * * * * * - * represents a thread that takes branch 1
1 * * * * * * - X represents a thread that takes branch 2
2 * * * * * * - - represents a thread that takes no branch
… * * * * * * -
15 * * * * * * -
16 * * * * * * -
17 - - - - - - X
Помните, что основы состоят из 32 потоков. Таким образом, потоки 0..31, 32..63 и т. Д. Выполняются в режиме ожидания. Как вы можете заметить на схеме выше, у вас есть один неактивный поток на каждые 18 потоков. Другими словами, это означает все ваши перекосы расходятся.
Это, вероятно, не огромный удар по производительности (если вообще), потому что одна из веток всегда «ничего не делает». При этом я определенно рекомендую вам исправить свой дизайн, и я уверен, что вы могли бы заметить улучшение производительности (хотя, скорее, из-за паттернов доступа к памяти, чем из-за расхождения).
Очевидным решением было бы запустить только 290 потоков вместо 324 и выполнить сопоставление с координатами x и y самостоятельно, но тогда ваш последний перекос будет заметно отличаться.
Другое решение состоит в том, чтобы запустить достаточно перекосов, чтобы покрыть первые 289 потоков (это означает, что 10 перекосов, а последний потратил впустую 31 поток), и запустить дополнительный перекос, в котором вы используете поток для второй ветви (например, последней) , Так что это будет 11 перекосов, 352 нитей, 62 впустую. Это может показаться хуже с точки зрения эффективности, но на самом деле это сложнее, чем из-за шаблонов доступа к памяти, так что попробуйте.
Также обратите внимание, что если тела if/else операторы на самом деле не различаются по коду, но по данным (как вы вроде бы подразумеваете …), тогда использование ветки бессмысленно. Просто играй с указателями. Могут возникнуть другие проблемы (связанные с объединением доступа к памяти), но не будет расхождения потока кода.
Я хотел бы предложить больше улучшений, но не видя ваш код или не зная, как ваши данные размещены, это своего рода выстрел в темноте. В комментариях вы говорите, что не можете заставить NSIGHT работать: я настоятельно рекомендую вам сделать это приоритетом.
Другие решения