Эффективная пропускная способность памяти при расширенном доступе к памяти

Question

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

__global__ void strideExample (float *outputData, float *inputData, int stride=2) 
{
        int index = (blockIdx.x * blockDim.x + threadIdx.x) * stride;
        outputData[index] = inputData[index]; 
}

Я понимаю, что доступ с размером шага 2 приведет к 50% эффективности загрузки / хранения, поскольку половина элементов, участвующих в транзакции, не используется (что приводит к потере пропускной способности). Как перейти к расчету эффективности загрузки / хранения для больших размеров шага? Заранее спасибо!

Answer 1

В общем:

load efficiency = requested loads / effective loads

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

Идеально слитые доступы имеют эффективность 1.

Ваш код запрашивает точно (blockIdx.x * blockDim.x + threadIdx.x) * sizeof(float) байт. Если предположить, outputData правильно выровнен (как указатели возвращаются cudaMalloc), аппаратное обеспечение придется читать (blockIdx.x * blockDim.x + threadIdx.x) * sizeof(float) * stride байты, округленные до размера транзакции (128 байтов для SM/L1, 32 байта для L1/L2).

Предполагая, что размер вашего блока достаточно велик, округление до размера транзакции становится незначительным, и вы можете упростить уравнение до 1 / stride, что дает в этом случае эффективность нагрузки около ~16,7%.