[Обмен опытом] GPU CUDA использует заполнение памяти, чтобы избежать конфликта банков

GPU
[Обмен опытом] GPU CUDA использует заполнение памяти, чтобы избежать конфликта банков

 欢迎关注我的公众号 [极智视界]，回复001获取Google编程规范

 O_o >_<  o_O O_o ~_~ o_O

  В этой статье рассказывается о том, как использовать заполнение памяти в программировании GPU CUDA, чтобы избежать конфликта банков.

1. Общая память

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

  • Канал связи между потоками внутри блока: канал связи между потоками;
  • Кэш, управляемый программой, для данных глобальной памяти: программируемый кеш;
  • Блокнотная память для преобразования данных для улучшения шаблонов доступа к глобальной памяти. Сократите количество обращений к глобальной памяти за счет кэширования данных.

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

__shared__     /// 标识符

__shared__ float tile[_y][_x];     /// 静态的声明了一个2D浮点型数组
    
/// kernel 内声明
extern __shared__ int tile[];

kernel<<<grid, block, isize * sizeof(int)>>>(...);

  Чтобы получить высокую пропускную способность, общая память разделена на 32 карты памяти одинакового размера, соответствующие потокам в варпе, к ним можно обращаться одновременно.

2. Используйте заполнение памяти, чтобы избежать конфликта банков

  Общая память работает так же быстро, как и регистры, если нет конфликтов банков.

  Быстрая ситуация:

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

  Медленный случай:

  • Конфликт банков: несколько потоков в варпе обращаются к одному и тому же банку;
  • Выборки должны быть сериализованы;
  • Максимальное количество потоков, которое несколько потоков могут получить доступ к одному и тому же банку одновременно.

  Примером конфликта банков является общая память:

  В случае отсутствия конфликта банков:

int x_id = blockDim.x * blockIdx.x + threadIdx.x;     // 列坐标
int y_id = blockDim.y * blockIdx.y + threadIdx.y;     // 行坐标
int index = y_id * col + x_id;

__shared__ float sData[BLOCKSIZE][BLOCKSIZE];

if(x_id < col && y_id < row)
{
    sData[threadIdx.y][threadIdx.x] = matrix[index];
    __syncthreads();
    matrixTest[index] = sData[threadIdx.y][threadIdx.x];
}

  В случае конфликта банков:

int x_id = blockDim.x * blockIdx.x + threadIdx.x;     // 列坐标
int y_id = blockDim.y * blockIdx.y + threadIdx.y;     // 行坐标
int index = y_id * col + x_id;

__shared__ float sData[BLOCKSIZE][BLOCKSIZE];

if(x_id < col && y_id < row)
{
    sData[threadIdx.x][threadIdx.y] = matrix[index];
    __syncthreads();
    matrixTest[index] = sData[threadIdx.x][threadIdx.y];
}

  В приведенном выше примере нет конфликта между банками и есть конфликт между банками, но было внесено лишь небольшое изменение. Давайте посмотрим, как решить вышеупомянутый конфликт между банками.

  Взяв приведенный выше пример в качестве примера, конфликта банков можно избежать, просто используя заполнение памяти, как показано на следующем рисунке:

  С точки зрения кода, как улучшить производительность кода с конфликтом банков выше с помощью заполнения памяти:

int x_id = blockDim.x * blockIdx.x + threadIdx.x;     // 列坐标
int y_id = blockDim.y * blockIdx.y + threadIdx.y;     // 行坐标
int index = y_id * col + x_id;

__shared__ float sData[BLOCKSIZE][BLOCKSIZE + 1];     // memory padding

if(x_id < col && y_id < row)
{
    sData[threadIdx.x][threadIdx.y] = matrix[index];
    __syncthreads();
    matrixTest[index] = sData[threadIdx.x][threadIdx.y];
}

  Выше был описан метод использования заполнения памяти, чтобы избежать конфликта банков при программировании GPU CUDA. Я надеюсь, что мой рассказ немного поможет вам в вашем исследовании.

【Передача по общему номеру】 "[Обмен опытом] GPU CUDA использует заполнение памяти, чтобы избежать конфликта банков》

Категории