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] субматрица исходного входа AX может быть дополнен нулем при необходимости, например, на контурах входа AX и окно может также быть 1D.

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

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

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

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

Если A 1D вектор-столбец, окном должен быть [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® ядра, которые выполняют фильтрацию изображения при помощи операций шаблона.

Этот пример выполняет среднюю фильтрацию 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

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

Ограничения

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

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

Приложения

Функции

Объекты

Темы

Введенный в R2017b

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

Поддержка

Для просмотра документации необходимо авторизоваться на сайте