#c #memory #cuda #shared
#c #память #cuda #общий
Вопрос:
Обратите внимание, что в этот массив общей памяти никогда не записывается, только считывается.
Как у меня есть, моя общая память инициализируется следующим образом:
__shared__ float TMshared[2592];
for (int i = 0; i< 2592; i )
{
TMshared[i] = TM[i];
}
__syncthreads();
(TM передается во все потоки с момента запуска ядра)
Возможно, вы заметили, что это крайне неэффективно, поскольку распараллеливание не происходит, и потоки внутри одного блока записывают данные в одно и то же место.
Может кто-нибудь порекомендовать более эффективный подход / прокомментировать, действительно ли эта проблема нуждается в оптимизации, поскольку рассматриваемый общий массив относительно невелик?
Спасибо!
Ответ №1:
Используйте все потоки для записи независимых местоположений, это, вероятно, будет быстрее.
Пример предполагает 1D threadblock / grid:
#define SSIZE 2592
__shared__ float TMshared[SSIZE];
int lidx = threadIdx.x;
while (lidx < SSIZE){
TMShared[lidx] = TM[lidx];
lidx = blockDim.x;}
__syncthreads();
Комментарии:
1. неплохо. Куда именно идет «#define SSIZE 2592»? В верхней части файла cu, за пределами глобального ядра?
2. Кроме того, какой смысл использовать #define ? Дает ли это преимущество по сравнению с простым явным кодированием числа 2592 в соответствующем месте?
3. Да, определение обычно находится в верхней части файла, хотя я почти уверен, что вы можете поместить его куда угодно (куда угодно, прежде чем оно будет использовано в коде). В define нет явного кода или преимущества в производительности по сравнению с 2592. Однако, если я изменю размер моего массива общей памяти, мне нужно будет изменить его только в одном месте.