Как лучше организовать чтение в общую память вот в таком случае:

Как лучше организовать чтение в общую память вот в таком случае:
Есть матрица 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). Однако тут получается еще нужны "краевые" значения для блока

Можете что нить лучше предложить такого варианта

  1. __shared__  float sTimecc[BLOCK_SIDE+2][BLOCK_SIDE+2];
  2. const int threadX = threadIdx.x+1;
  3. const int threadY = threadIdx.y+1;
  4.  
  5. sTimecc[threadX][threadY]=t[(y)*nz + (z)];
  6.             if(threadX == blockDim.x){
  7.                 sTimecc[threadX+1][threadY] = t[(y+1)*nz + (z)];
  8.             }
  9.             if(threadX == 1){
  10.                 sTimecc[threadX-1][threadY] = t[(y-1)*nz + (z)];
  11.             }
  12.             if(threadY == blockDim.y){
  13.                 sTimecc[threadX][threadY+1] = t[(y)*nz + (z+1)];
  14.             }
  15.             if(threadY == 1){
  16.                 sTimecc[threadX][threadY-1] = t[(y)*nz + (z-1)];
  17.             }
  18. __syncthreads();

Forums: 

Я не понимаю как (z, y)

Я не понимаю как (z, y) связаны с (threadX, threadY). Если никак -- использовать broadcast, прочитать всеми нитями краевые значения. Когда все нити читают одно слово, то оно не читается за один запрос, а не за количество нитей.

А вообще тут проблема в branching-е, а не в памяти. Надо свести if'ы к вычислению.

Ну вoобще как то так связаны

Ну вoобще как то так связаны )))

  1.                 const int y =  (blockIdx.x)*blockDim.x + threadY-1;
  2.                 const int z = (blockIdx.y)*blockDim.y + threadZ-1

Мне не представляется

Мне не представляется возможным здесь избавится от ветвления за счет вычислений ибо чтение идет из глобальной памяти уже определенных величин.
Я вот на схемке зарисовал что я хочу прочитать блоком потоков)) :

  1. -111111111-
  2. 10000000001
  3. 10000000001
  4. 10000000001
  5. 10000000001
  6. 10000000001
  7. -111111111-

Где "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, конфликты банков второго порядка приобритете -- к бабке не ходи.

Спасибо будем пробовать ) Ну

Спасибо будем пробовать )

Ну а вообще я их дальше как то так использую. Вроде конфликтов банков второго рода не должно быть.......или я чего-то не понимаю

  1. sTimecf[threadIdx.x][threadIdx.y] = sTimecc[threadX][threadY+1];
  2. sTimefc[threadIdx.x][threadIdx.y] = sTimecc[threadX+1][threadY];
  3. __syncthreads();
  4. foo(sTimecc[threadX][threadY],sTimefc[threadIdx.x][threadIdx.y],sTimecf[threadIdx.x][threadIdx.y]);

Ну, тут надо брать листочек

Ну, тут надо брать листочек бумаги и считать :)

Каждая строка Вашего блока нитей (16 нитей) объединяются в half-warp (в полном warp'е 32 нити), half-warp одновременно обращается к памяти. Шаред-память разбивается на 16 банков. 16 подрядидущих слов попадают в разные банки. Таким образом, если две нити из одной строки (при данном блоке нитей 16x16 это half-warp) обращаются к элементам, разложенным в памяти на расстоянии sizeof(float)*16, происходит конфликт и это обращение происходит не за один такт, а за два. Если больше нитей конфликтуют, то еще больше.

Но это все ерунда, конечно. Всё, что я писал, довольно вяло отзывалось на оптимизации, убирающие конфликты банков до 16 включительно. Вот устранение бранчингов приводит к фантастическим изменениям, да.

В общем я избавился от

В общем я избавился от ветвлений таким образом:

  1.             sTimecc[threadIdx.x][threadIdx.y]=t[(y)*nz + (z)];
  2.                 sTimecf[threadIdx.x][threadIdx.y] = t[(y)*nz + (z+1)];
  3.                 sTimecb[threadIdx.x][threadIdx.y] = t[(y)*nz + (z-1)];
  4.                 sTimefc[threadIdx.x][threadIdx.y] = t[(y+1)*nz + (z)];
  5.                 sTimebc[threadIdx.x][threadIdx.y] = t[(y-1)*nz + (z)];

То есть увеличение количества чтений из глобальной памяти примерно в 3,9 раз (вы предлогали в 3,2 раза увеличить). Это привело к 17% увеличению времени исполнения......
Я понимаю это не то что вы предлогали - в данном случае одно и тоже значение из глобальной памяти читается 4 различными тредами. Не думаю что ваш вариант даст прирост производительности :-(.

Да я вроде не предлагал

Да я вроде не предлагал увеличивать количество чтений :) Вы читали блоками 16х16 нитей матрицу 16х16, или, если учитывать пересекающиеся края, 18x18. Я предлагал блоком 16x16 нитей читать матрицу 30x30 с пересекающимися краями, то есть 32x32. Число чтений меньше, чем в Вашем исходном варианте.

Единственное что непонятно при этом, важно ли Вашему алгоритму, что блок обрабатывает блок именно 16х16. По косвенным признакам похоже, что все-таки блок 16x16 памяти это важно. Если это действительно так, и нельзя никак изменить обработку блока 16x16 на обработку блока 30x30, то тогда имеет смысл все пересекающиеся края блоков вынести в отдельные массивы.

Вы правильно заметели что

Вы правильно заметели что блок размером 16х16 обрабатывает блок данных именно 16х16. Вот на счет выносить края в отдельный массив это уже интересно там есть конечно ряд тонкостей и ограничений.....но надо пробовать. Спасибо