В этом примере показано, как сопоставить вход с постоянным пространством памяти на графическом процессоре при помощи 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];
}
}