生成コードからのカスタム CUDA デバイス関数の呼び出し

生成コードに組み込もうとしている特定のサブ関数について高度に最適化された CUDA^® コードがある場合、GPU Coder™ は coder.ceval 機能を拡張することでこの目標達成を後押しします。

外部の CUDA 関数を GPU デバイス上で実行するには、__device__ 修飾子を使用しなければなりません。これらのデバイス関数は、他のデバイス関数またはグローバル関数からのみ呼び出すことができるという点で、グローバル関数 (カーネル) とは異なります。そのため、デバイス関数の coder.ceval 呼び出しは、カーネルにマッピングされるループ内に配置しなければなりません。CUDA カーネルと生成コードの統合の詳細については、生成コードからのカスタム CUDA カーネルの呼び出しを参照してください。

メモ

coder.ceval 呼び出しが含まれたループをカーネルにマッピングできない場合、コード生成は失敗します。カーネルの作成を妨げている問題および推奨される回避方法を確認するには、GPU Coder ドキュメンテーション内のトラブルシューティングのトピックを参照してください。MATLAB^® コードセクション内にサポートされない関数が含まれている場合、そのようなセクションから coder.ceval 呼び出しを削除しなければなりません。

CUDA デバイス関数 `__usad4_wrap` の呼び出し

ステレオ視差の例では、ステレオペアの左右のイメージ内にある 2 つの対応する点間の距離を測定します。エントリポイント関数 stereoDisparity_cuda_sample は、関数 coder.ceval を使用して外部デバイス関数 __usad4_wrap を呼び出します。

%% modified algorithm for stereo disparity block matching
% In this implementation instead of finding shifted image ,indices are mapped 
% accordingly to save memory and some processing RGBA column major packed 
% data is used as input for compatibility with CUDA intrinsics. Convolution
% is performed using separable filters (Horizontal and then Vertical)

function [out_disp] = stereoDisparity_cuda_sample(img0,img1)
coder.cinclude('cuda_intrinsic.h');

% gpu code generation pragma
coder.gpu.kernelfun;

%% Stereo disparity Parameters
% WIN_RAD is the radius of the window to be operated,min_disparity is the 
% minimum disparity level the search continues for, max_disparity is the maximum
% disparity level the search continues for.
WIN_RAD = 8;
min_disparity = -16;
max_disparity = 0;

%% Image dimensions for loop control
% The number of channels packed are 4 (RGBA) so as nChannels are 4
[imgHeight,imgWidth]=size(img0);
nChannels = 4;
imgHeight = imgHeight/nChannels;

%% To store the raw differences
diff_img = zeros([imgHeight+2*WIN_RAD,imgWidth+2*WIN_RAD],'int32');

%To store the minimum cost
min_cost = zeros([imgHeight,imgWidth],'int32');
min_cost(:,:) = 99999999;

% Store the final disparity
out_disp = zeros([imgHeight,imgWidth],'int16');

%% Filters for aggregating the differences
% filter_h is the horizontal filter used in separable convolution
% filter_v is the vertical filter used in separable convolution which
% operates on the output of the row convolution
filt_h = ones([1 17],'int32');
filt_v = ones([17 1],'int32');

%% Main Loop that runs for all the disparity levels. This loop is currently
% expected to run on CPU.
for d=min_disparity:max_disparity
    
    % Find the difference matrix for the current disparity level. Expect
    % this to generate a Kernel function.
    coder.gpu.kernel;
    for colIdx=1:imgWidth+2*WIN_RAD
        coder.gpu.kernel;
        for rowIdx=1:imgHeight+2*WIN_RAD
            % Row index calculation
            ind_h = rowIdx - WIN_RAD;
            
            % Column indices calculation for left image
            ind_w1 = colIdx - WIN_RAD;
            
            % Row indices calculation for right image
            ind_w2 = colIdx + d - WIN_RAD;
            
            % Border clamping for row Indices
            if ind_h <= 0
                ind_h = 1;
            end
            if ind_h > imgHeight
                ind_h = imgHeight;
            end
            
            % Border clamping for column indices for left image
            if ind_w1 <= 0
                ind_w1 = 1;
            end
            if ind_w1 > imgWidth
                ind_w1 = imgWidth;
            end
            
            % Border clamping for column indices for right image
            if ind_w2 <= 0
                ind_w2 = 1;
            end
            if ind_w2 > imgWidth
                ind_w2 = imgWidth;
            end
            
            % In this step, Sum of absolute Differences is performed
            % across Four channels. This piece of code is suitable 
            % for replacement with SAD intrinsics
            tDiff = int32(0);
            tDiff = coder.ceval('-gpudevicefcn', '__usad4_wrap', 
                    coder.rref(img0((ind_h-1)*(nChannels)+1,ind_w1)), 
                    coder.rref(img1((ind_h-1)*(nChannels)+1,ind_w2)));
            
            %Store the SAD cost into a matrix
            diff_img(rowIdx,colIdx) = tDiff;
        end
    end
    
    % Aggregating the differences using separable convolution. Expect this 
    % to generate two Kernel using shared memory.The first kernel is the 
    % convolution with the horizontal kernel and second kernel operates on
    % its output the column wise convolution.
    cost_v = conv2(diff_img,filt_h,'valid');
    cost = conv2(cost_v,filt_v,'valid');
    
    % This part updates the min_cost matrix with by comparing the values
    % with current disparity level. Expect to generate a Kernel for this.
    for ll=1:imgWidth
        for kk=1:imgHeight
            % load the cost
            temp_cost = int32(cost(kk,ll));
            
            % compare against the minimum cost available and store the
            % disparity value
            if min_cost(kk,ll) > temp_cost
                min_cost(kk,ll) = temp_cost;
                out_disp(kk,ll) = abs(d) + 8;
            end
            
        end
    end
    
end
end

__usad4_wrap の定義は外部ファイル cuda_intrinsic.h に記述されています。このファイルはエントリポイント関数と同じフォルダーにあります。

__device__ unsigned int __usad4(unsigned int A, unsigned int B, unsigned int C=0)
{
    unsigned int result;
#if (__CUDA_ARCH__ >= 300) // Kepler (SM 3.x) supports a 4 vector SAD SIMD
    asm("vabsdiff4.u32.u32.u32.add" " %0, %1, %2, %3;": "=r"(result):"r"(A),
    "r"(B), "r"(C));
#else // SM 2.0            // Fermi  (SM 2.x) supports only 1 SAD SIMD, 
                           // so there are 4 instructions
    asm("vabsdiff.u32.u32.u32.add" " %0, %1.b0, %2.b0, %3;": 
         "=r"(result):"r"(A), "r"(B), "r"(C));
    asm("vabsdiff.u32.u32.u32.add" " %0, %1.b1, %2.b1, %3;": 
         "=r"(result):"r"(A), "r"(B), "r"(result));
    asm("vabsdiff.u32.u32.u32.add" " %0, %1.b2, %2.b2, %3;": 
         "=r"(result):"r"(A), "r"(B), "r"(result));
    asm("vabsdiff.u32.u32.u32.add" " %0, %1.b3, %2.b3, %3;": 
         "=r"(result):"r"(A), "r"(B), "r"(result));
#endif
    return result;
}

__device__ unsigned int packBytes(const uint8_T *inBytes)
{
    unsigned int packed = inBytes[0] | (inBytes[1] << 8) | 
                    (inBytes[2] << 16) | (inBytes[3] << 24);
    return packed;
}

__device__ unsigned int __usad4_wrap(const uint8_T *A, const uint8_T *B)
{
    unsigned int x = packBytes(A);
    unsigned int y = packBytes(B);

    return __usad4(x, y);
}

CUDA コードの生成

コード生成構成オブジェクトを作成して CUDA コードを生成します。構成オブジェクトのカスタムコードプロパティ (CustomInclude) を設定して、カスタム C ファイルの場所を指定します。以下は、cuda_intrinsic.h ファイルの場所を指すコード生成スクリプトの例です。

cfg = coder.gpuConfig('mex');
cfg.CustomInclude = pwd;

codegen -config cfg -args {imgRGB0, imgRGB1} stereoDisparity_cuda_sample_intrinsic;