バインドされたCUDAテクスチャはゼロを読み取る

Question

テクスチャから値を読み込み、グローバルメモリに書き戻そうとします。 私はカーネルに一定の値を置くことができると私は出力でそれらを見ることができますbeause私は、書込部が動作すると確信しています：

__global__ void 
bartureKernel(float* g_odata, int width, int height) 
{ 
    unsigned int x = blockIdx.x*blockDim.x + threadIdx.x; 
    unsigned int y = blockIdx.y*blockDim.y + threadIdx.y; 

    if(x < width && y < height) { 
      unsigned int idx = (y*width + x); 
      g_odata[idx] = tex2D(texGrad, (float)x, (float)y).x; 

    } 
} 

私が使用したいテクスチャはそう、2つのチャンネルの2Dフロート質感ですトンを設定し、

float* d_data_grad = NULL; 

cudaMalloc((void**) &d_data_grad, gradientSize * sizeof(float)); 
CHECK_CUDA_ERROR; 

texGrad.addressMode[0] = cudaAddressModeClamp; 
texGrad.addressMode[1] = cudaAddressModeClamp; 
texGrad.filterMode = cudaFilterModeLinear; 
texGrad.normalized = false; 

cudaMemset(d_data_grad, 50, gradientSize * sizeof(float)); 
CHECK_CUDA_ERROR; 

cudaBindTexture(NULL, texGrad, d_data_grad, cudaCreateChannelDesc<float2>(), gradientSize * sizeof(float)); 

float* d_data_barture = NULL; 
cudaMalloc((void**) &d_data_barture, outputSize * sizeof(float)); 
CHECK_CUDA_ERROR; 

dim3 dimBlock(8, 8, 1); 
dim3 dimGrid(((width-1)/dimBlock.x)+1, ((height-1)/dimBlock.y)+1, 1); 

bartureKernel<<< dimGrid, dimBlock, 0 >>>(d_data_barture, width, height); 

私が知っている：

texture<float2, 2, cudaReadModeElementType> texGrad; 

とカーネルを呼び出すコードは、いくつかの一定の非ゼロの値を持つテクスチャを初期化します。私は、として定義しエクステントバイトをすべて "50"にすることは、浮動小数点数の文脈ではあまり意味がありませんが、少なくともゼロ以外の値を読み取る必要があります。

私だけ

Answer 1

あなたはcudaMallocによって割り当てられたメモリへのあなたのテクスチャをバインドするcudaBindTextureを使用している...しかしゼロを読み取ることができます。カーネルでは、tex2D関数を使用して、テクスチャから値を読み込みます。それがゼロを読み込んでいる理由です。

cudaBindTextureを使用してテクスチャを線形メモリにバインドする場合は、カーネル内でtex1Dfetchを使用して読み取られます。

tex2Dは機能cudaBindTexture2D、または機能cudaBindTextureToArray

を用いcudaArray に結合され、それらのテクスチャを用いて（cudaMallocPitchによって割り当てられる）ピッチ線形メモリに結合されているもののテクスチャのみから読み取るために使用されていますここで

は、基本的なテーブルであり、あなたがプログラミングガイドから読み取ることができ休ま：

メモリタイプ ----------------- アル所在地を使用して ----------------- バインドを使用して -----------------------

リニアなメモリによってカーネル................... cudaMalloc .................... .... cudaBindTexture ............................. tex1Dfetch

ピッチリニアメモリ........ cudaMallocPitch ............. cudaBindTexture2D ........................ tex2D

cudaArray .... ........................ cudaMallocArray ............. cudaBindTextureToArray ............. tex1Dまたはtex2D

3D cudaArray ............. ......... cudaMalloc3DArray ........ cudaBindTextureToArray ............. tex3D

Answer 2

追加するには、tex1Dfetchを使用したアクセスは整数インデックスに基づいています。 しかし、残りは浮動小数点に基づいてインデックスが付けられており、+ 0.5を加えて必要な正確な値を取得する必要があります。

浮動小数点を作成してfloat2テクスチャにバインドするのはなぜですか？あいまいな結果になることがあります。 float2は2Dフロートテクスチャではありません。実際に複素数の表現に使用できます。

typedef struct {float x; float y;} float2; 

このチュートリアルは、cudaでテクスチャメモリを使用する方法を理解するのに役立つと思います。 http://www.drdobbs.com/parallel/cuda-supercomputing-for-the-masses-part/218100902

あなたが示されたカーネルは、テクスチャを使用してからあまりメリットはありません。しかし、適切に利用されれば、局所性を利用することにより、テクスチャメモリは性能をかなり向上させることができる。また、補間にも役立ちます。