Безконфликтный доступ в общей памяти

Question

Я должен использовать разделяемую память размером 64 элемента, вдвое больше, чем 32 банка. Таким образом, дело в том, что количество обращений к памяти вдвое превышает количество потоков в деформации. Как мне обратиться к ним, чтобы обеспечить бесконтактный доступ к банку?

Answer 1

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

__shared__ int shared[32];
int data = shared[base + stride * tid];

там stride нечетно.

Если у вас есть 64-битный доступ, вы можетеиспользуйте такой трюк, как этот:

struct type 
{  
   int x, y, z;
};
__shared__ struct type shared[32];
struct type data = shared[base + tid];

Answer 2

Предположим, вы используете вычислительную возможность 1.x, поэтому ваша общая память имеет 16 банков, и каждый поток должен иметь доступ к 2 элементам в общей памяти.

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

__shared__ int shared[32];
int data = shared[base + stride * tid];
int data = shared[base + stride * tid + 16];

Я использовал этот шаблон для хранения сложных поплавков, но у меня был массив сложных поплавков, поэтому он выглядел как

#define TILE_WIDTH 16

__shared__ float shared[TILE_WIDTH][2*TILE_WIDTH + 1];
float real = shared[base + stride * tid];
float imag = shared[base + stride * tid + TILE_WIDTH];

Где +1 - чтобы избежать сериализации в транспонированных шаблонах доступа.