__device__ __constant__ const

Question

Есть ли разница, и как лучше всего определить константы устройства в программе CUDA?В C ++, программа хост / устройство, если я хочу определить константы в постоянной памяти устройства, я могу сделать либо

__device__ __constant__ float a = 5;
__constant__ float a = 5;

Вопрос 1. На устройствах 2.x и CUDA 4 это так же, как,

__device__ const float a = 5;

Вопрос 2. Почему в PyCUDA SourceModule ("" "..." ""), который компилирует только код устройства, работают даже следующие?

const float a = 5;

Answer 1

В CUDA __constant__ - это спецификатор типа переменной, который указывает, что объявленная переменная должна храниться в постоянной памяти устройства .Раздел B 2.2 Руководства по программированию CUDA

Спецификатор __constant__, необязательно используемый вместе с __device__, объявляет переменную, которая:

• Находится в постоянной памятипробел,
• Имеет время жизни приложения,
• Доступен из всех потоков в сетке и из хоста через библиотеку времени выполнения (cudaGetSymbolAddress() / cudaGetSymbolSize() / cudaMemcpyToSymbol()/ cudaMemcpyFromSymbol() для API времени выполнения и cuModuleGetGlobal() для API драйвера).

В CUDA постоянная память - это выделенная, статическая, глобальная область памяти, доступ к которой осуществляется через кэш (тамэто специальный набор инструкций загрузки PTX для его цели), которые являются унифицированными и доступны только для чтения для всех потоков в работающем ядре.Но содержимое постоянной памяти может быть изменено во время выполнения посредством использования API на стороне хоста, указанных выше.Это отличается от объявления переменной компилятору с помощью объявления const, которое добавляет характеристику, доступную только для чтения, к переменной в области действия объявления.Они совсем не одно и то же.