Эффективная пропускная способность в Cuda
При вычислении эффективной пропускной способности в Cuda я подсчитываю количество операций чтения / записи в разделяемой памяти. Пример кода приведен ниже.
__global__ void kernel(float *a, float * b, float * c, int num){
int i = threadIdx.x + blockIdx.x*blockDim.x;
__shared__ a_shared[NUM];
__shared__ b_shared[NUM];
if (i < NUM){
a_shared[i] = a[i];
b_shared[i] = b[i]
c[i] = a_shared[i] + b_shared[i];
}
}
Внимательное прочтение этого раздела подсказывает мне, что имеются в виду глобальные загрузки / хранилища, где резервным хранилищем для глобального трафика является DRAM. Следовательно, совместная деятельность не рассматривается и не должна включаться.
Ответы 1
Что касается секция руководства по передовой практике, которое вы указываете в комментариях выше, я бы сказал, что ответ отрицательный, общий трафик не должен включаться.
Откуда нам это знать?
- Основная цель расчета эффективной пропускной способности - сравнить ее с теоретической пропускной способностью:
To measure performance accurately, it is useful to calculate theoretical and effective bandwidth. When the latter is much lower than the former, design or implementation details are likely to reduce bandwidth, and it should be the primary goal of subsequent optimization efforts to increase it.
Однако расчет теоретической пропускной способности включает только глобальный трафик памяти в DRAM:
Using these data items, the peak theoretical memory bandwidth of the NVIDIA Tesla M2090 is 177.6 GB/s:
Это число - пропускная способность DRAM. Он не включает пропускную способность разделяемой памяти.
- Все ссылки на измерения профилировщика относятся к глобальному трафику памяти, а не к разделяемой памяти:
Requested Global Load Throughput
Requested Global Store Throughput
Global Load Throughput
Global Store Throughput
DRAM Read Throughput
DRAM Write Throughput
- Метод расчета теоретической пропускной способности разделяемой памяти нигде не задокументирован в официальной документации CUDA, поэтому его нельзя было включить в расчет теоретической пропускной способности. Поэтому включение измерения пропускной способности разделяемой памяти не имеет смысла для целей сравнения.