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