OpenCLでlanczos3を使った画像の拡大処理(2)

1: imgSrcの画像を水平方向に補完してimgTmpに書込む。
2: imgTmpの画像を垂直方向に補完してimgDstに書込む

imgTmp.width = imgDst.width
imgTmp.height = imgSrc.height
とする。

//////// begin:
#define GAMMA_CORRECT 0
__kernel void lanczos3_hmag_vmag_passFirst(__write_only image2d_t imgTmp, __read_only image2d_t imgSrc)
{
    const sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE| CLK_ADDRESS_MIRRORED_REPEAT| CLK_FILTER_NEAREST;

    int2 coordDst = (int2)(get_global_id(0), get_global_id(1));
    float widthSrc = get_image_width(imgSrc);
    float widthDst = get_image_width(imgTmp);
    float heightDst = get_image_height(imgTmp);


    if (widthDst<=coordDst.x || heightDst<=coordDst.y)
        return;

    float center = (coordDst.x+0.5f) * (widthSrc/widthDst);
    float beginSrc = trunc(center-3.0f);
    float distance = (beginSrc+0.5f)-center;
    float2 coordSample = (float2)(beginSrc, coordDst.y);
    float sumWeight = 0.0f;
    float4 sumColor = 0.0f;

    for (int i=0; i<6; i++)
    {
        float weight = lanczos3(distance);
        float4 color = read_imagef(imgSrc, sampler, coordSample);
#if GAMMA_CORRECT
        color = pow(color, 2.2);
#endif
        sumColor += color * weight;
        sumWeight += weight;
        coordSample.x += 1.0f;
        distance += 1.0f;
    }

    write_imagef(imgTmp, coordDst, sumColor / sumWeight);
}

__kernel void lanczos3_hmag_vmag_passSecond(__write_only image2d_t imgDst, __read_only image2d_t imgTmp)
{
    const sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE| CLK_ADDRESS_MIRRORED_REPEAT| CLK_FILTER_NEAREST;

    int2 coordDst = (int2)(get_global_id(0), get_global_id(1));
    float heightSrc = get_image_height(imgTmp);
    float heightDst = get_image_height(imgDst);
    float widthDst = get_image_width(imgTmp);


    if (widthDst<=coordDst.x || heightDst<=coordDst.y)
        return;

    float center = (coordDst.y+0.5f) * (heightSrc/heightDst);
    float beginSrc = trunc(center-3.0f);
    float distance = (beginSrc+0.5f)-center;
    float2 coordSample = (float2)(coordDst.x, beginSrc);
    float sumWeight = 0.0f;
    float4 sumColor = 0.0f;

    for (int i=0; i<6; i++)
    {
        float weight = lanczos3(distance);
        sumColor += weight * read_imagef(imgTmp, sampler, coordSample);
        sumWeight += weight;
        coordSample.y += 1.0f;
        distance += 1.0f;
    }

#if GAMMA_CORRECT
    write_imagef(imgDst, coordDst, pow(sumColor / sumWeight, 1.0f/2.2f));
#else
    write_imagef(imgDst, coordDst, sumColor / sumWeight);
#endif
}
//////// end:
遅いカーネルでは平均38.7[ms]、今回のカーネルでは平均8.69[ms]といった具合。
処理点数は36+6から6+6への減少なのにスループットは4倍以上になっているが、原因は分からない。

今回のカーネルでは重みを毎回計算しているが、事前にテーブルを作成しておき、そこから引けば速くなりそうだ。
次回はテーブルの作成のみ書く。