Рекомендации по постоянной памяти CUDA

Здесь я представляю код

__constant__ int array[1024];

__global__ void kernel1(int *d_dst) {
   int tId = threadIdx.x + blockIdx.x * blockDim.x;
   d_dst[tId] = array[tId];
}

__global__ void kernel2(int *d_dst, int *d_src) {
   int tId = threadIdx.x + blockIdx.x * blockDim.x;
   d_dst[tId] = d_src[tId];
}

int main(int argc, char **argv) {
   int *d_array;
   int *d_src;
   cudaMalloc((void**)&d_array, sizeof(int) * 1024);
   cudaMalloc((void**)&d_src, sizeof(int) * 1024);

   int *test = new int[1024];
   memset(test, 0, sizeof(int) * 1024);

   for (int i = 0; i < 1024; i++) {
     test[i] = 100;
   }

   cudaMemcpyToSymbol(array, test, sizeof(int) * 1024);
   kernel1<<< 1, 1024 >>>(d_array);

   cudaMemcpy(d_src, test, sizeof(int) * 1024, cudaMemcpyHostToDevice);
   kernel2<<<1, 32 >>>(d_array, d_src),

   free(test);
   cudaFree(d_array);
   cudaFree(d_src);

   return 0;
}

Это просто показывает постоянную память и использование глобальной памяти. При его выполнении "kernel2" выполняется примерно в 4 раза быстрее (с точки зрения времени), чем "kernel1"

Я понимаю из руководства по программированию Cuda C, что это связано с тем, что доступ к постоянной памяти становится сериализованным. Это подводит меня к мысли о том, что постоянную память можно наилучшим образом использовать, если warp обращается к одному постоянному значению, такому как integer, float, double и т.д., Но доступ к массиву вообще не выгоден. Другими словами, я могу сказать, что warp должен получить доступ к одному адресу, чтобы получить выгоду от оптимизации/ускорения от постоянного доступа к памяти. Правильно ли это?

Я также хочу знать, если я сохраняю структуру вместо простого типа в моей постоянной памяти. Любой доступ к структуре потоком с деформацией; также рассматривается как доступ к одной памяти или больше? Я имею в виду, что структура может содержать несколько простых типов и массив, например; при доступе к этим простым типам, эти обращения также сериализованы или нет?

Последний вопрос был бы, если у меня есть массив с постоянными значениями, к которому нужно обращаться через разные потоки в основе; для более быстрого доступа он должен храниться в глобальной памяти вместо постоянной памяти. Это правильно?

Любой может передать мне пример кода, в котором показано эффективное использование постоянной памяти.

С уважением,

Ответы

Ответ 1

Я могу сказать, что warp должен получить доступ к одному адресу, чтобы получить выгоду от оптимизации/ускорения от постоянного доступа к памяти. Правильно ли это?

Да, это обычно правильно и является основным намерением использовать постоянный кеш памяти/константы. Постоянный кеш может обслуживать одно 32-разрядное количество за цикл для каждого SM. Поэтому, если каждый поток в warp получает одно и то же значение:

int i = array[20];

тогда у вас будет возможность для хорошей выгоды от постоянного кеша/памяти. Если каждый поток в warp получает уникальное количество:

int i = array[threadIdx.x];

тогда доступ будет сериализован, и постоянное использование данных будет неутешительным, с точки зрения производительности.

Я также хочу знать, если я сохраняю структуру вместо простого типа в моей постоянной памяти. Любой доступ к структуре потоком с деформацией; также рассматривается как доступ к одной памяти или более?

Вы можете, конечно, поместить структуры в постоянную память. Те же правила применяются:

int i = constant_struct_ptr->array[20];

имеет возможность воспользоваться, но

int i = constant_struct_ptr->array[threadIdx.x];

нет. Если вы обращаетесь к одному и тому же элементу структуры простой структуры по потокам, это идеально подходит для постоянного использования кеша.

Последний вопрос был бы, если у меня есть массив с постоянными значениями, к которому нужно обращаться через разные потоки в основе; для более быстрого доступа он должен храниться в глобальной памяти вместо постоянной памяти. Это правильно?

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

Существует множество примеров кода cuda, которые демонстрируют использование данных __constant__. Вот несколько:

и есть другие.

EDIT: отвечает на вопрос в комментариях, если у нас есть такая структура в постоянной памяти:

struct Simple { int a, int b, int c} s;

И мы обращаемся к нему следующим образом:

int p = s.a + s.b + s.c;
          ^     ^     ^
          |     |     |
cycle:    1     2     3

У нас будет хорошее использование постоянной памяти/кеша. Когда код C скомпилируется, под капотом он будет генерировать коды машинного кода, соответствующие 1,2,3 на диаграмме выше. Представьте себе, что первый доступ происходит первым. Поскольку доступ 1 относится к тому же месту памяти, независимо от того, какой поток в warp, в течение цикла 1, все потоки получат значение в s.a, и он будет использовать кеш для наилучшей возможной выгоды. Аналогично для доступа 2 и 3. Если, с другой стороны, мы имели:

struct Simple { int a[32], int b[32], int c[32]} s;
...
int idx = threadIdx.x + blockDim.x * blockIdx.x;
int p = s.a[idx] + s.b[idx] + s.c[idx];

Это не даст хорошего использования постоянной памяти/кеша. Вместо этого, если бы это было типично для нашего доступа к s, мы, вероятно, имели бы лучшую производительность, располагающую s в обычной глобальной памяти.