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

Ограничения

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

См. также

Приложения

Функции

Объекты

Темы

Введенный в R2017b

Документация по GPU Coder

Поддержка