Как лучше организовать чтение в общую память вот в таком случае:
Есть матрица t[nx][ny] заведомо больше размера блока(16х16), Каждому потоку требеются
текущее значение t[x][y] и 4 соседа t[x][y+1], t[x][y-1], t[x+1][y], t[x-1][y]. Т.е. в идеале каждый поток читает свое значение в общую память потом делаем синхронизацию блока(те все как в примере с матрицой из sdk). Однако тут получается еще нужны "краевые" значения для блока
Можете что нить лучше предложить такого варианта
- __shared__ float sTimecc[BLOCK_SIDE+2][BLOCK_SIDE+2];
- const int threadX = threadIdx.x+1;
- const int threadY = threadIdx.y+1;
- sTimecc[threadX][threadY]=t[(y)*nz + (z)];
- if(threadX == blockDim.x){
- sTimecc[threadX+1][threadY] = t[(y+1)*nz + (z)];
- }
- if(threadX == 1){
- sTimecc[threadX-1][threadY] = t[(y-1)*nz + (z)];
- }
- if(threadY == blockDim.y){
- sTimecc[threadX][threadY+1] = t[(y)*nz + (z+1)];
- }
- if(threadY == 1){
- sTimecc[threadX][threadY-1] = t[(y)*nz + (z-1)];
- }
- __syncthreads();
Comments
Я не понимаю как (z, y) связаны с (threadX, threadY). Если никак -- использовать broadcast, прочитать всеми нитями краевые значения. Когда все нити читают одно слово, то оно не читается за один запрос, а не за количество нитей.
А вообще тут проблема в branching-е, а не в памяти. Надо свести if'ы к вычислению.
Мне не представляется возможным здесь избавится от ветвления за счет вычислений ибо чтение идет из глобальной памяти уже определенных величин.
Я вот на схемке зарисовал что я хочу прочитать блоком потоков)) :
Где "0" - это данные которые читает эта строчка
sTimecc[threadX][threadY]=t[(y)*nz + (z)];"1" - это то что мне надо еще прочитать. "-" - неиспользуемые данные в данном блокеПС надеюсь стало яснее )))
Думаю, что тут нельзя ничего "подправить", чтобы стало эффективно. Тут не-coalescing чтение из глобальной, потому что не последовательные слова читаются из памяти, плюс heavy-branching, убивающий производительность в разы.
Предлагаю изменить порядок загрузки данных нитями. Сделать, чтобы нить читала блок 32x32, а не 18x18. То есть чтобы каждая нить из блока 16x16 грузила 4 значения.
1 1 2 2 3 3 4 4 ... 16 16
1 1 2 2 3 3 4 4 ... 16 16
17 17 18 18 ...
17 17 18 18 ...
...
Тут не нужно if'ов и можно сделать coalescing чтение. Не знаю, правда, как это усложнит обработку.
При этом как Вы собираетесь использовать эту шаред-память тоже интересно, потому что при блоке 18x18, конфликты банков второго порядка приобритете -- к бабке не ходи.
Спасибо будем пробовать )
Ну а вообще я их дальше как то так использую. Вроде конфликтов банков второго рода не должно быть.......или я чего-то не понимаю
Ну, тут надо брать листочек бумаги и считать :)
Каждая строка Вашего блока нитей (16 нитей) объединяются в half-warp (в полном warp'е 32 нити), half-warp одновременно обращается к памяти. Шаред-память разбивается на 16 банков. 16 подрядидущих слов попадают в разные банки. Таким образом, если две нити из одной строки (при данном блоке нитей 16x16 это half-warp) обращаются к элементам, разложенным в памяти на расстоянии sizeof(float)*16, происходит конфликт и это обращение происходит не за один такт, а за два. Если больше нитей конфликтуют, то еще больше.
Но это все ерунда, конечно. Всё, что я писал, довольно вяло отзывалось на оптимизации, убирающие конфликты банков до 16 включительно. Вот устранение бранчингов приводит к фантастическим изменениям, да.
В общем я избавился от ветвлений таким образом:
То есть увеличение количества чтений из глобальной памяти примерно в 3,9 раз (вы предлогали в 3,2 раза увеличить). Это привело к 17% увеличению времени исполнения......
Я понимаю это не то что вы предлогали - в данном случае одно и тоже значение из глобальной памяти читается 4 различными тредами. Не думаю что ваш вариант даст прирост производительности :-(.
Да я вроде не предлагал увеличивать количество чтений :) Вы читали блоками 16х16 нитей матрицу 16х16, или, если учитывать пересекающиеся края, 18x18. Я предлагал блоком 16x16 нитей читать матрицу 30x30 с пересекающимися краями, то есть 32x32. Число чтений меньше, чем в Вашем исходном варианте.
Единственное что непонятно при этом, важно ли Вашему алгоритму, что блок обрабатывает блок именно 16х16. По косвенным признакам похоже, что все-таки блок 16x16 памяти это важно. Если это действительно так, и нельзя никак изменить обработку блока 16x16 на обработку блока 30x30, то тогда имеет смысл все пересекающиеся края блоков вынести в отдельные массивы.
Вы правильно заметели что блок размером 16х16 обрабатывает блок данных именно 16х16. Вот на счет выносить края в отдельный массив это уже интересно там есть конечно ряд тонкостей и ограничений.....но надо пробовать. Спасибо
Ну вoобще как то так связаны )))