CUDA Constant Memory Best Practices

Question

Я представляю здесь некоторый код

__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 я понимаю, что это происходит потому, что доступ к постоянной памяти становится сериализованным. Это подводит меня к мысли, что постоянную память лучше всего использовать, если деформация обращается к единственному постоянному значению, такому как целое число, число с плавающей запятой, двойное число и т. Д., Но доступ к массиву совсем не выгоден. Другими словами, я могу сказать, что деформация должна получить доступ к одному адресу, чтобы иметь какие-либо преимущества оптимизации / ускорения от постоянного доступа к памяти. Это правильно?

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

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

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

С уважением,

Answer 1

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

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

int i = array[20];

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

int i = array[threadIdx.x]; 

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

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

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

int i = constant_struct_ptr->array[20]; 

имеет возможность извлечь выгоду, но

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

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

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

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

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

1. графика громкости рендера
2. двусторонняя фильтрация изображений
3. Свертка изображений
4. финансы MonteCarloGPU

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

РЕДАКТИРОВАТЬ: отвечая на вопрос в комментариях, если у нас есть такая структура в постоянной памяти:

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 происходит первым. Поскольку доступ 1 находится в одной и той же ячейке памяти, независимо от того, какой поток в деформации, во время цикла 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 в обычной глобальной памяти.