CUDA сообщило о конфликтах в банках с общей памятью

Я работал над оптимизацией некоторого кода и столкнулся с проблемой отчета о конфликте банков совместно используемой памяти с анализом производительности CUDA Nsight. Мне удалось свести его к очень простому коду, который Nsight сообщает о наличии банковского конфликта, когда он, кажется, не должен быть. Ниже находится ядро:

__global__ void conflict() {
    __shared__ double values[33];
    values[threadIdx.x] = threadIdx.x;
    values[threadIdx.x+1] = threadIdx.x;
}

И основная функция для его вызова:

int main() {
    conflict<<<1,32>>>();
}

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

Кто-нибудь еще испытывал проблемы с отчетами Nsight или я просто что-то упустил из-за функционирования банковских конфликтов? Буду признателен за любые отзывы!

Кстати, я запускаю следующую настройку:

  • Windows 8
  • GTX 770
  • Visual Studio Community 2013
  • CUDA 7
  • Nsight Visual studio, версия 4.5
Источник

2 ответа

Решение

Если цель состоит в том, чтобы запустить опубликованный код как есть, с double тип данных, и никаких банковских конфликтов, я считаю, что это возможно при надлежащем использовании cudaDeviceSetSharedMemConfig (на устройствах cc3.x). Вот тестовый пример:

$ cat t750.cu
#include <stdio.h>

typedef double mytype;


template <typename T>
__global__ void conflict() {
    __shared__ T values[33];
    values[threadIdx.x] = threadIdx.x;
    values[threadIdx.x+1] = threadIdx.x;
}

int main(){

#ifdef EBM
  cudaDeviceSetSharedMemConfig(cudaSharedMemBankSizeEightByte);
#endif

  conflict<mytype><<<1,32>>>();
  cudaDeviceSynchronize();
}

$ nvcc -arch=sm_35 -o t750 t750.cu
t750.cu(8): warning: variable "values" was set but never used
          detected during instantiation of "void conflict<T>() [with T=mytype]"
(19): here

$ nvprof --metrics shared_replay_overhead ./t750
==46560== NVPROF is profiling process 46560, command: ./t750
==46560== Profiling application: ./t750
==46560== Profiling result:
==46560== Metric result:
Invocations                               Metric Name                        Metric Description         Min         Max         Avg
Device "Tesla K40c (0)"
 Kernel: void conflict<double>(void)
          1                    shared_replay_overhead             Shared Memory Replay Overhead    0.142857    0.142857    0.142857
$ nvcc -arch=sm_35 -DEBM -o t750 t750.cu
t750.cu(8): warning: variable "values" was set but never used
          detected during instantiation of "void conflict<T>() [with T=mytype]"
(19): here

$ nvprof --metrics shared_replay_overhead ./t750
==46609== NVPROF is profiling process 46609, command: ./t750
==46609== Profiling application: ./t750
==46609== Profiling result:
==46609== Metric result:
Invocations                               Metric Name                        Metric Description         Min         Max         Avg
Device "Tesla K40c (0)"
 Kernel: void conflict<double>(void)
          1                    shared_replay_overhead             Shared Memory Replay Overhead    0.000000    0.000000    0.000000
$

С указанием EightByteModeиздержки воспроизведения совместно используемой памяти равны нулю.

Источник

Оказывается, моя ошибка была с типом данных, который я использовал. Я ошибочно считал само собой разумеющимся, что каждый элемент будет размещен в 1 банке. Однако двойной тип данных составляет 8 байтов, поэтому он охватывает 2 банка общей памяти. Изменение типа данных на float решило эту проблему, и оно правильно отображает 0 банковских конфликтов. Спасибо за отзыв и помощь.

Источник
Другие вопросы по тегам