exponenta event banner

gpucoder.stencilKernel

Создание кода CUDA для функций набора элементов

Описание

пример

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 и окно также может быть 1-D.

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

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

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

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

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

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

  • 'same' - Возвращает выходные данные B это тот же размер, что и A.

  • 'full' - (по умолчанию) Возвращает полный вывод. Размер 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 ®, которые выполняют фильтрацию изображения с помощью операций трафарета.

В этом примере выполняется средняя фильтрация 2-D изображения. В одном файле запишите функцию точки входа 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

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

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

Используйте codegen для создания функции CUDA MEX.

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__ квалификатор) для улучшения пропускной способности памяти и локализации данных.

Ограничения

  • Для очень больших входных размеров gpucoder.stencilKernel функция может создавать код CUDA, не соответствующий числовому значению моделирования MATLAB ®. В таких случаях следует рассмотреть возможность уменьшения объема вводимых данных для получения точных результатов.

Представлен в R2017b