3

私はかなり前からCUDAコードを書いてきましたが、テクスチャキャッシュの使い方に慣れてきました。

simpleTextureインスピレーションを得るためにNvidiaSDKの例を使用して、テクスチャキャッシュを使用する簡単な例をコーディングしました。ホストはLenaイメージをGPUにコピーし、テクスチャとしてバインドします。カーネルは、テクスチャキャッシュの内容を出力配列にコピーするだけです。

奇妙なことに、結果（コードの下のすべて灰色の画像を参照）は入力と一致しません。何がうまくいかないかについて何か考えはありますか？

コード（見てくださいtexCache_dummyKernel）：

texture<float, 2, cudaReadModeElementType> tex; //declare texture reference for 2D float texture

//note: tex is global, so no input ptr is needed
__global__ void texCache_dummyKernel(float* out, const int width, const int height){ //copy tex to output
    int x = blockIdx.x*blockDim.x + threadIdx.x; //my index into "big image"
    int y = blockIdx.y*blockDim.y + threadIdx.y;
    int idx = y*width+x;

    if(x < width && y < height)
        out[idx] = tex2D(tex, y, x);
}

int main(int argc, char **argv){        
    cv::Mat img = getRawImage("./Lena.pgm");
    img.convertTo(img, CV_32FC1);
    float* hostImg = (float*)&img.data[0];
    int width = img.cols; int height = img.rows;

    dim3 grid;  dim3 block;
    block.x = 16;  block.y = 16;
    grid.x = width/block.x + 1;          
    grid.y = height/block.y + 1;

    cudaArray *dImg; //cudaArray*, not float*
    cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc(32, 0, 0, 0, cudaChannelFormatKindFloat);        
    CHECK_CUDART(cudaMallocArray(&dImg, &channelDesc, width, height));
    CHECK_CUDART(cudaMemcpyToArray(dImg, 0, 0, hostImg, width*height*sizeof(float), cudaMemcpyHostToDevice));
    setTexCacheParams(); //defined below
    CHECK_CUDART(cudaBindTextureToArray(tex, dImg, channelDesc)); //Bind the array to the texture

    float* dResult; //device memory for output
    CHECK_CUDART(cudaMalloc((void**)&dResult, sizeof(float)*width*height));

    texCache_dummyKernel<<<grid, block>>>(dResult, width, height); //dImg isn't an input param, since 'tex' is a global variable
    CHECK_CUDART(cudaGetLastError()); //make sure kernel didn't crash

    float* hostResult = (float*)malloc(sizeof(float)*width*height);
    CHECK_CUDART(cudaMemcpy(hostResult, dResult, sizeof(float)*width*height, cudaMemcpyDeviceToHost));
    outputProcessedImage(hostResult, width, height, "result.png"); //defined below
}

上記で使用したヘルパー関数をいくつか提供する必要があります。

void setTexCacheParams(){ //configuration directly pulled from simpleTexture in nvidia sdk
    tex.addressMode[0] = cudaAddressModeWrap;
    tex.addressMode[1] = cudaAddressModeWrap;
    tex.filterMode = cudaFilterModeLinear;
    tex.normalized = true;    // access with normalized texture coordinates
}

void outputProcessedImage(float* processedImg, int width, int height, string out_filename){
    cv::Mat img = cv::Mat::zeros(height, width, CV_32FC1);
    for(int i=0; i<height; i++)
        for(int j=0; j<width; j++)
            img.at<float>(i,j) = processedImg[i*width + j]; //just grab the 1st of the 4 pixel spaces in a uchar4

    img.convertTo(img, CV_8UC1); //float to uchar
    vector<int> compression_params;
    compression_params.push_back(CV_IMWRITE_PNG_COMPRESSION);
    compression_params.push_back(9);
    cv::imwrite(out_filename, img, compression_params);
}

入力：

ここに画像の説明を入力してください

出力：

ここに画像の説明を入力してください

  • この投稿がコードの壁になっていることをお詫びします。このようなコードをより簡潔にするためのアドバイスをいただければ幸いです。
  • 上記のファイルI/OにOpenCVを使用しました...これが混乱しないことを願っています。
  • 1Dfloat*配列から入力画像を読み取るようにカーネルを変更し、他のほとんどすべてを同じに保つと、正しい結果が得られます。
4

1 に答える 1

3

元のコードでは、正規化された座標を使用するようにテクスチャを初期化していました。これは、テクスチャが各空間次元の[0,1]でアドレス指定されることを意味します。したがって、カーネルは次のようになります。

__global__ 
void texCache_dummyKernel(float* out, const int width, const int height)
{
    int x = blockIdx.x*blockDim.x + threadIdx.x; //my index into "big image"
    int y = blockIdx.y*blockDim.y + threadIdx.y;
    int idx = y*width+x;

    if(x < width && y < height) {
        float u = float(x)/float(width), v = float(y)/float(height);
        out[idx] = tex2D(tex, u, v);
    }
}

[標準の免責事項：ブラウザで記述され、コンパイルまたはテストされていないため、自己責任で使用してください]

すなわち。tex2D画像の幅と高さで割って正規化された座標を渡す必要があります。

または、発見したように、テクスチャ定義を変更して、normalized=false相対テクスチャ座標ではなく絶対テクスチャ座標でアドレス指定を使用することもできます。それでも、コードで読み取られるテクスチャは次のようになります。

out[idx] = tex2D(tex, float(x)+0.5f, float(y)+0.5f);

テクスチャのアドレス指定は常に浮動小数点座標を使用して行われ、テクスチャデータはボクセル中心であるため、各座標に0.5が追加され、テクスチャ内の各補間領域またはボリュームの重心から読み取りが行われるようになります。

テクスチャフィルタリングとアドレッシングモードの説明、およびそれらが補間に与える影響については、CUDACプログラミングガイドの付録の1つを参照してください。

于 2012-11-24T08:34:07.977 に答える