Можно ли выделить часть общей памяти для каждого потока? - PullRequest
Новая
       21

Можно ли выделить часть общей памяти для каждого потока?

0 голосов
/ 14 октября 2019

Я использую CUDA 10.1 в Ubuntu 18.04, и я хотел бы знать, можно ли выделить часть общей памяти для каждого потока. Под этим я подразумеваю, скажем, я хочу иметь относительно большой массив, приватный для каждого потока, который не помещается в регистры, и помещение массива в разделяемую память вводит условия гонки, поскольку каждый поток в блоке будет иметь доступ кЭто. Поэтому мне очень интересно использовать часть разделяемой памяти в качестве «регистра», чтобы избежать гонок. Я понимаю, что это ограничит мою занятость, ограничив число потоков, которые я мог бы использовать в SM, но компромисс стоит в моем случае. Любая помощь приветствуется!

1 Ответ

1 голос
/ 15 октября 2019

Нет, невозможно выделить выделение общей памяти (статической или динамической) определенному потоку. Совместно используемая память имеет только область видимости блока.

Однако вполне возможно разработать схему индексирования, при которой потоки внутри блока используют исключительно уникальные местоположения в распределении общей памяти области действия блока. Например: 

template<int nthreads, int words_per_thread>
__global__
void kernel(..)
{
    __shared__ int buffer[nthreads * words_per_thread];

    int* localbuff = &buffer[threadIdx.x * words_per_thread];

    // localbuff is now safely indexed from [0] to [words_per_thread-1]

}

Другой возможный метод будет выглядеть примерно так: 

#include <stdio.h>
template<typename T>
class sr_
{
  T *sp;
  public:
  __device__
  sr_(T *_sd) { sp = _sd;}
  __device__
  T &operator[](int idx){return sp[blockDim.x*idx+threadIdx.x];}
};
// for statically allocated shared memory
#define SREG(size,type,block_size) \
  __shared__ type SREG_sdata[size*block_size]; \
  typedef type SREG_type; \
  sr_<SREG_type> sreg(SREG_sdata);
// for dynamically allocated shared memory
#define DSREG(type) \
  __shared__ type SREG_sdata[]; \
  typedef type SREG_type; \
  sr_<SREG_type> sreg(SREG_sdata);

const int BS = 8;

__global__ void k2(){
  SREG(8,float,BS)
  sreg[0] = 1.0f;
  printf("%f\n", sreg[0]);
}

int main(){

  k2<<<1,BS>>>();
  cudaDeviceSynchronize();
}

Преимущества этого метода в том, что у одного потока нет возможности индексировать пространство другого потока, итакже, что не может быть никаких банковских конфликтов. Обратите внимание, что это не будет обрабатывать все варианты использования. Его необходимо изменить, например, если бы в одном модуле было несколько ядер, которые использовали разные типы данных для совместной памяти.

...