У меня есть вопросы по поводу объединенных глобальных загрузок памяти в CUDA.В настоящее время мне нужно иметь возможность выполнять на устройстве CUDA с возможностью вычислений CUDA 1.1 или 1.3.
Я пишу функцию ядра CUDA, которая считывает массив типа T из глобальной памяти в разделяемую память, выполняет некоторые вычисления, а затем выпишет массив типа T обратно в глобальную память.Я использую разделяемую память, потому что вычисление для каждого выходного элемента фактически зависит не только от соответствующего входного элемента, но также и от соседних входных элементов.Я хочу загрузить каждый элемент ввода только один раз, поэтому я хочу кэшировать элементы ввода в общей памяти.
Мой план состоит в том, чтобы каждый поток считывал один элемент в общую память, а затем __syncthreads () перед началом вычисления,В этом сценарии каждый поток загружает, вычисляет и сохраняет один элемент (хотя вычисления зависят от элементов, загруженных в общую память другими потоками).
В этом вопросе я хочу сосредоточиться на чтении из глобальной памяти вобщая память.
Предполагая, что в массиве есть N элементов, я настроил CUDA для выполнения в общей сложности N потоков.Для случая, когда sizeof (T) == 4, это должно хорошо слиться в соответствии с моим пониманием CUDA, так как поток K будет читать слово K (где K - индекс потока).
Однако в случаегде sizeof(T) < 4, например, если T = unsigned char или T = short, то я думаю, что это может быть проблемой.В этом случае мой (наивный) план таков:
- Вычислить numElementsPerWord = 4 / sizeof (T)
- если (K% numElementsPerWord == 0), то прочитать есть поток Kпрочитайте следующее полное 32-битное слово
- сохраните 32-битное слово в разделяемой памяти
- после заполнения разделяемой памяти (и вызовите __syncthreads ()), после чего каждый поток K сможет обработать работупри вычислении выходного элемента K
Меня беспокоит, что он не будет объединяться, потому что (например, в случае, когда T = short)
- Поток 0 читает слово 0 изглобальная память
- поток 1 не читает
- поток 2 читает слово 1 из глобальной памяти
- поток 3 не читает
- и т. д ...
Другими словами, поток K читает слово (K / sizeof (T)).Казалось бы, это не слилось должным образом.
Альтернативный подход, который я рассмотрел, был:
- Запуск с числом потоков = (N + 3) / 4, так что каждый поток будетотвечать за загрузку и обработку (4 / sizeof (T)) элементов (каждый поток обрабатывает одно 32-битное слово - возможно, 1, 2 или 4 элемента в зависимости от sizeof (T)).Однако я обеспокоен тем, что этот подход не будет настолько быстрым, насколько это возможно, поскольку каждый поток должен затем в два раза (если T = short) или даже в четыре раза (если T = unsigned char) количество обработки.
Может кто-нибудь сказать мне, если мои предположения о моем плане верны: то есть: он не будет объединяться должным образом?
Можете ли вы прокомментировать мой альтернативный подход?
Можете ли вы порекомендовать более оптимальныйподход, который правильно объединяется?