Exponenta Banner
exponenta event banner

gpucoder.stencilKernel

Создайте код CUDA для функций шаблона

свернуть все на странице

Синтаксис

B = gpucoder.stencilKernel(FUN,A,[M N],shape,param1,param2...)

Описание

пример

B = gpucoder.stencilKernel(FUN,A,[M N],shape,param1,param2...) применяет функциональный FUN к каждому раздвижному окну [M,N] входа A. Функциональный FUN называется для каждой субматрицы [M,N] A и вычисляет элемент вывода B. Индекс этого элемента соответствует центру окна [M,N].

FUN является указателем на пользовательскую функцию, которая возвращает скалярный выходной параметр того же типа как вход.

C= FUN(X,param1,param2, ...)

X является субматрицей [M,N] исходного входа A. X может быть дополнен нулем при необходимости, например, на контурах входа A. X и окно могут также быть 1D.

C является скаляром, оцененным вывод FUN. Это - вывод, вычисленный для центрального элемента массива [M,N] X, и присвоено соответствующему элементу выходного массива B.

param1,param2 является дополнительными аргументами. Передайте эти аргументы, если FUN требует каких-либо дополнительных параметров в дополнение к входному окну.

Окно [M,N] должно быть меньше чем или равно размеру A с той же формой как A.

Если A является 1D вектором - строкой, окном должен быть [1,N].

Если A является 1D вектор-столбцом, окном должен быть [N,1].

shape определяет размер выходного массива B. Это может иметь одно из трех возможных значений:

  • B вывода 'same' - Returns, который одного размера как A.

  • полный- (значение по умолчанию) Возвращает полный выходной параметр. Размер B> размер A, то есть, если A имеет размер (x, y). Размер B = [x + floor(M/2), y + floor(N/2)]

  • 'valid' - Возвращает только те части вывода, которые вычисляются без дополненных нулем ребер A. Размер B = [x - floor(M/2), y - floor(N/2)]

Вход A должен быть вектором или матрицей с числовым типом, поддержанным FUN. Класс B совпадает с классом A.

Генерация кода поддерживается только для фиксированного размера выходные параметры. Форма и окно должны быть константами времени компиляции, потому что они определяют размер вывода.

Примеры

свернуть все

Этот пример показывает, как использовать gpucoder.stencilKernel и сгенерировать ядра CUDA®, которые выполняют фильтрацию изображения при помощи операций шаблона.

Этот пример выполняет среднюю фильтрацию 2D изображения. В одном файле запишите функции точки входа test, который принимает матрицу изображений A. Создайте подфункциональный my_mean, который вычисляет среднее значение субматрицы 3x3.

function B = meanImgFilt(A)  %#codegen
  B = gpucoder.stencilKernel(@my_mean,A,[3 3],'same');
  
  function out = my_mean(A)
    out = cast(mean(A(:)), class(A));
  end
end

Setup изображение тестового воздействия для функции meanImgFilt.

inImage = im2double(imread('cameraman.tif'));

Используйте функцию codegen, чтобы сгенерировать MEX-функцию CUDA.

codegen -config coder.gpuConfig('mex') -args {inImage} -report meanImgFilt

GPU Coder создает три ядра: meanImgFilt_kernel1 для инициализации памяти, meanImgFilt_kernel2 для оптимизации входной структуры памяти и meanImgFilt_kernel3 для средней операции фильтрации. Следующее является отрывком сгенерированного кода.

  cudaMalloc(&gpu_B, 524288ULL);
  cudaMalloc(&gpu_A, 524288ULL);
  cudaMalloc(&gpu_expanded, 532512ULL);
  meanImgFilt_kernel1<<<dim3(131U, 1U, 1U), dim3(512U, 1U, 1U)>>>(gpu_expanded);
  cudaMemcpy((void *)gpu_A, (void *)&A[0], 524288ULL, cudaMemcpyHostToDevice);
  meanImgFilt_kernel2<<<dim3(128U, 1U, 1U), dim3(512U, 1U, 1U)>>>(gpu_A,
    gpu_expanded);
  meanImgFilt_kernel3<<<dim3(8U, 8U, 1U), dim3(32U, 32U, 1U)>>>(gpu_expanded,
    gpu_B);
  cudaMemcpy((void *)&B[0], (void *)gpu_B, 524288ULL, cudaMemcpyDeviceToHost);

meanImgFilt_kernel3 использует общую память (спецификатор __shared__), чтобы улучшить пропускную способность памяти и местность данных.

Смотрите также

| | |

Темы

Введенный в R2017b

Документация GPU Coder
Поддержка