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