gpucoder.stencilKernel

(非推奨) ステンシル関数の CUDA コードの作成

関数 gpucoder.stencilKernel は推奨されません。代わりに stencilfun を使用してください。詳細については、バージョン履歴を参照してください。

構文

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

説明

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 次元にすることもできます。

C は FUN のスカラー値の出力です。これは、[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 で導入

すべて折りたたむ

R2022b: 関数 `gpucoder.stencilKernel` は非推奨

関数 gpucoder.stencilKernel は推奨されません。代わりに stencilfun を使用してください。

R2022b 以降は、関数 stencilfun を使用して、ステンシルのような演算用の CUDA カーネルを生成します。

gpucoder.stencilKernel の代表的な使用法と、代わりに stencilfun を使用するためのコードの更新方法を、次の表に示します。

非推奨推奨

非推奨	推奨
`gpucoder.stencilKernel` を使用した畳み込み: function Out = myconv(In) Out = gpucoder.stencilKernel(@stencilFcn, In, [5 5], 'same'); end function y = stencilFcn(X) W = rand(5); y = 0; for j = 1:5 for i = 1:5 y = y + X(i,j) * W(i,j); end end end	`stencilfun` を使用した畳み込み: function Out = myconv(In) fh = @(X) stencilFcn(X); Out = stencilfun(fh, In, [5 5], Shape = 'same'); end function y = stencilFcn(X) W = rand(5) y = 0; for j = 1:5 for i = 1:5 y = y + X(i,j) * W(i,j); end end end
ステンシル関数に追加の引数を渡します。 weights = rand(5); In = rand(100); Out = gpucoder.stencilKernel(@stencilFcn, In, [5 5],'same',weights); function y = stencilFcn(X, weights) y = 0; for i = 1 : 5 for j = 1 : 5 y = y + X(j,i) * weights(j,i); end end end	無名関数を使用して、追加の引数をステンシル関数に渡します。 weights = rand(5); fh = @(X) stencilFcn(X, weights); In = rand(100); Out = stencilfun(fh, In, [5 5],Shape='same'); function y = stencilFcn(X, weights) y = 0; for i = 1 : 5 for j = 1 : 5 y = y + X(j,i) * weights(j,i); end end end

gpucoder.stencilKernel を使用した畳み込み:

function Out = myconv(In)
Out = gpucoder.stencilKernel(@stencilFcn, In, [5 5], 'same');
end

function y = stencilFcn(X)
W = rand(5);
y = 0;
for j = 1:5
    for i = 1:5
        y = y + X(i,j) * W(i,j);
    end
end
end

stencilfun を使用した畳み込み:

function Out = myconv(In)
fh = @(X) stencilFcn(X);
Out = stencilfun(fh, In, [5 5], Shape = 'same');
end

function y = stencilFcn(X)
W = rand(5)
y = 0;
for j = 1:5
    for i = 1:5
        y = y + X(i,j) * W(i,j);
    end
end
end

ステンシル関数に追加の引数を渡します。

weights = rand(5);
In = rand(100);
Out = gpucoder.stencilKernel(@stencilFcn, In, [5 5],'same',weights);
 
function y = stencilFcn(X, weights)
    y = 0;
    for i = 1 : 5
        for j = 1 : 5
            y = y + X(j,i) * weights(j,i);
        end
    end
end

無名関数を使用して、追加の引数をステンシル関数に渡します。

weights = rand(5);
fh = @(X) stencilFcn(X, weights);
In = rand(100);
Out = stencilfun(fh, In, [5 5],Shape='same');
 
function y = stencilFcn(X, weights)
    y = 0;
    for i = 1 : 5
        for j = 1 : 5
            y = y + X(j,i) * weights(j,i);
        end
    end
end

参考

gpucoder.stencilKernel

構文

説明

例

ステンシル カーネルを使用した平均値フィルター

制限

バージョン履歴

R2022b: 関数 gpucoder.stencilKernel は非推奨

参考

アプリ

関数

オブジェクト

トピック

ステンシルカーネルを使用した平均値フィルター

R2022b: 関数 `gpucoder.stencilKernel` は非推奨