Запись функции точки входа myFun который принимает два входа a размера 256x256 и постоянной k размера 1x3. Функция имеет вложенный for- контуры, добавляющие константы к каждому элементу a. Чтобы создать ядро, разместите coder.gpu.kernel() прагматика за пределами вложенного for-луп. coder.gpu.constantMemory(k) помещает вход только для чтения k в постоянную память ГПУ.

function b = myFun(a,k)
  b = coder.nullcopy(zeros(size(a)));
  coder.gpu.kernel();
    for j = 1:256
      for i = 1:256
        coder.gpu.constantMemory(k);  
        b(i,j) = a(i,j) + k(1) + k(2) + k(3);
      end
    end
end

Создайте объект конфигурации для создания кода MEX.

cfg = coder.gpuConfig('mex');

Определение массива ячеек input который объявляет размер и тип данных входных данных a,k к функции myFun.

input = {ones(256),ones(1,3)}

Создание функции MEX myFun_mex с помощью -config, -args, и -report параметры для задания конфигурации, предоставления входных аргументов и создания отчета о создании кода.

codegen -config cfg -args input -report myFun

В отчете на вкладке Код C щелкните 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];
  }
}