Main Content

このページの翻訳は最新ではありません。ここをクリックして、英語の最新版を参照してください。

gpucoder.stencilKernel

ステンシル関数の CUDA コードの作成

説明

B = gpucoder.stencilKernel(FUN,A,[M N],shape,param1,param2...) は、関数 FUN を、入力 A の各 [M,N] スライディング ウィンドウに適用します。関数 FUN は、A の各 [M,N] 部分行列に対して呼び出され、出力 B の要素を計算します。この要素のインデックスは [M,N] ウィンドウの中心に対応します。

FUN は、入力と同じ型のスカラー出力を返すユーザー定義関数のハンドルです。

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

X は元の入力 A[M,N] 部分行列です。X は、入力 A の境界などで、必要に応じてゼロでパディングできます。X とこのウィンドウは 1 次元にすることもできます。

CFUN のスカラー値の出力です。これは、[M,N] 配列 X の中心要素について計算された出力であり、出力配列 B の対応する要素に代入されます。

param1,param2 はオプション引数です。FUN に、入力ウィンドウのほかに追加のパラメーターが必要な場合、これらの引数を渡します。

ウィンドウ [M,N]A 以下のサイズで、A と同じ形状でなければなりません。

A が 1 次元行ベクトルの場合、ウィンドウは [1,N] でなければなりません。

A が 1 次元列ベクトルの場合、ウィンドウは [N,1] でなければなりません。

shape は出力配列 B のサイズを決定します。これは次の 3 つの値のいずれかを持つことができます。

  • 'same' - A と同じサイズの出力 B が返されます。

  • 'full' - (既定) 完全な出力が返されます。つまり、A のサイズが (x,y) の場合は、B のサイズ > A のサイズ。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 次元イメージの平均値フィルター処理を実行します。1 つのファイル内に、イメージ行列 A を受け入れるエントリポイント関数 test を記述します。3x3 の部分行列の平均値を計算する部分行列 my_mean を作成します。

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 により 3 つのカーネルが作成されます。メモリの初期化用の 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 は、MATLAB® のシミュレーションと数値的に一致しない CUDA コードを生成する可能性があります。このような場合、正確な結果が得られるように、入力サイズを小さくすることを検討してください。

R2017b で導入