В этом примере показано, как сопоставить вход в постоянное пространство памяти на графическом процессоре с помощью coder.gpu.constantMemory прагма.
Написание функции точки входа myFun который принимает два входов a размера 256x256 и постоянные k размера 1x3. Функция имеет вложенное for-циклы, который добавляет константы к каждому элементу a. Чтобы создать ядро, разместите coder.gpu.kernel() прагма вне вложенного for-цикл. coder.gpu.constantMemory(k) помещает вход только для чтения k в постоянную память графический процессор.
Создайте объект строения для генерации кода MEX.
Задайте массив ячеек input который объявляет размер и тип данных входов a,k в myFun функций.
Сгенерируйте MEX-функцию myFun_mex при помощи -config, -args, и -report опции, чтобы задать строение, предоставить входные параметры и сгенерировать отчет генерации кода.
В отчете на вкладке C code нажмите myFun.cu.
Переменная только для чтения k объявлен как const_k при помощи __constant__ как показано на фрагменте кода.
/* Variable Definitions */
__constant__ real_T const_k[3];
cudaMemcpyToSymbol вызов копирует значение k от хоста к устройству постоянной памяти const_k.
cudaMemcpyToSymbol(const_k, k, 24U, 0U, cudaMemcpyHostToDevice);
cudaMemcpy(gpu_a, a, 524288U, cudaMemcpyHostToDevice);
myFun_kernel1<<<dim3(128U, 1U, 1U), dim3(512U, 1U, 1U)>>>(gpu_a, gpu_b);
cudaMemcpy(b, gpu_b, 524288U, cudaMemcpyDeviceToHost);
Тело ядра обращается к постоянной const_k и добавляет его к каждому элементу a
static __global__ __launch_bounds__(512, 1) void myFun_kernel1(const real_T *a,
real_T *b)
{
int32_T i;
int32_T j;
int32_T threadIdX;
threadIdX = (int32_T)(blockDim.x * blockIdx.x + threadIdx.x);
i = threadIdX / 256;
j = threadIdX - i * 256;
if ((!(j >= 256)) && (!(i >= 256))) {
b[i + (j << 8)] = ((a[i + (j << 8)] + const_k[0]) + const_k[1]) + const_k[2];
}
}