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