Вопросы о памяти CUDA

Question

Я довольно новичок в программировании CUDA, и некоторые модели памяти, которые мне не совсем понятны, есть. Как, как это работает? Например, если у меня есть простое ядро ​​

__global__ void kernel(const int* a, int* b){ 
    some computation where different threads in different blocks might  
    write at the same index of b 
} 

Так я представляю a будет находиться в так называемой постоянной памяти. Но как насчет b? Поскольку в нем будут записываться разные потоки в разных блоках, как это будет работать? Я где-то читал, что было гарантировано, что в случае одновременных записей в глобальной памяти разными потоками в одном блоке будет написан хотя бы один, но нет гарантии о других. Должен ли я беспокоиться об этом, то есть, например, каждый поток в блоке записывается в разделяемую память, и как только все это делается, нужно написать все в глобальную память? Или CUDA заботится о нем для меня?

Answer 1

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

Да, aуказатель будет находиться в постоянной памяти, но не потому, что он отмечен const (это полностью ортогональны). bуказатель также находится в постоянной памяти. Все аргументы ядра передаются в постоянной памяти (кроме CC 1.x). Память, на которую ссылаются a и b, теоретически может быть чем угодно (глобальная память устройства, память, закрепленная хостом, все, что может быть адресовано UVA, я считаю). Где он находится, выбран вами, пользователем.

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

Предположим, что ваш код выглядит следующим образом:

b[0] = 10; // Executed by all threads 

Тогда да, это (доброкачественная) состояние гонки, так как все потоки пишут то же самое значение в том же месте. Результат записи определен, однако количество записей не задано, и также есть поток, который выполняет «окончательную» запись. Единственная гарантия заключается в том, что происходит хотя бы одна запись. На практике я считаю, что выдается одна запись для каждого варпа, что является пустой тратой пропускной способности, если в ваших блоках содержится более одной основы (что они должны).

С другой стороны, если ваш код выглядит следующим образом:

b[0] = threadIdx.x; 

Это ясно неопределенное поведение.

Должен ли я беспокоиться об этом, то есть, например, каждый поток в блоке записывается в разделяемую память, и как только все это будет сделано, нужно написать все в глобальную память?

Да, так оно и делается.