coder.gpu.constantMemory

Прагма, которая сопоставляет переменную с постоянной памятью на графическом процессоре

Синтаксис

coder.gpu.constantMemory(v)

Описание

пример

coder.gpu.constantMemory(v) сопоставляет переменную v с постоянным пространством памяти на устройстве графического процессора. Поместите эту прагму в parallelizable цикле. Если GPU Coder™ генерирует ядро для цикла, он загружает v к устройству постоянная переменная memory и заменяет любой доступ к этой переменной в ядре доступом к постоянной переменной memory. Переменная v должна быть только для чтения в ядре, в противном случае GPU Coder игнорирует эту прагму. Используйте эту прагму, когда каждый поток получит доступ к каждому элементу массива параметров / матрица.

Эта функция является функцией генерации кода. Это не имеет никакого эффекта в MATLAB®.

Примеры

свернуть все

Этот пример показывает, как сопоставить вход с постоянным пространством памяти на графическом процессоре при помощи прагмы coder.gpu.constantMemory.

Запишите функции точки входа 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 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];
  }
}

Введенный в R2017b