Блог

Вот моя проблема: у меня довольно большой набор двойников (это массив из 77.500 двойников), который должен храниться где-то в cuda. Теперь мне нужен большой набор потоков, чтобы последовательно выполнять кучу операций с этим массивом. Каждый поток должен будет считывать ОДИН и ТОТ ЖЕ элемент этого массива, выполнять задачи, сохранять результаты в общей памяти и считывать следующий элемент массива. Обратите внимание, что каждый поток одновременно должен будет считывать (просто считывать) из одной и той же ячейки памяти. Итак, мне интересно: есть ли какой-нибудь способ транслировать один и тот же double во все потоки только с одним чтением памяти? Многократное чтение было бы совершенно бесполезным… Есть идеи??

Это обычная оптимизация. Идея состоит в том, чтобы заставить каждый поток взаимодействовать со своими соседями по блоку для чтения данных:

 // choose some reasonable block size
const unsigned int block_size = 256;

__global__ void kernel(double *ptr)
{
  __shared__ double window[block_size];

  // cooperate with my block to load block_size elements
  window[threadIdx.x] = ptr[threadIdx.x];

  // wait until the window is full
  __syncthreads();

  // operate on the data
  ...
}
 

Вы можете итеративно «перемещать» окно по элементам массива block_size (или, может быть, на несколько целых чисел больше) за раз, чтобы использовать все это целиком. Тот же метод применяется, когда вы хотите сохранить данные обратно синхронизированным образом.