cudaBindTexture2D具體實作出現錯誤

我原本做texture 的時候是參酌手冊的使用方法

        cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<float>();
   cudaArray *cu_rho;
   cudaMallocArray(&cu_rho, &channelDesc, N, N);
   
... ... ...

   initialR<<<grid,block>>>(dev_rho); \\給定初始值
   cudaMemcpyToArray( cu_rho, 0, 0, dev_rho, N*N*sizeof(float), cudaMemcpyDeviceToDevice );
   cudaBindTextureToArray( Trho, cu_rho );

   kernel <<< >>>();\\使用到Trho

   cudaUnbindTexture( Trho );
   cudaFreeArray( cu_rho );
   cudaFreeArray( cu_V );


但是我看到cuda by example 中heat_2d.cu的程式碼是用不同的方式來調用texture

cudaMalloc( (void**)&data.dev_inSrc,mageSize )
cudaChannelFormatDesc desc = cudaCreateChannelDesc<float>();
 cudaBindTexture2D( NULL, texConstSrc , data.dev_inSrc , desc , DIM, DIM,sizeof(float) * DIM ) ;
... ... ...
for (int i=0; i<90; i++) {
   float   *in, *out;
   if (dstOut) {
   in  = d->dev_inSrc;
   out = d->dev_outSrc;
   } else {
   out = d->dev_inSrc;
   in  = d->dev_outSrc;
   }
   copy_const_kernel<<<blocks,threads>>>( in );  
   blend_kernel<<<blocks,threads>>>( out, dstOut ); \\綁定
   dstOut = !dstOut;
   }
... ... ...
cudaUnbindTexture( ... );

把範例程式碼編譯過後有看到效果,我自己實作一次的結果卻是錯誤的

texture<float, 2> Trho;
cudaMalloc((void**)&dev_rho,N*N*sizeof(float));
cudaChannelFormatDesc chDesc = cudaCreateChannelDesc<float>();
cudaBindTexture2D( NULL, Trho, dev_rho, chDesc, N, N, N*sizeof(float) );
   
initialR<<<grid,block>>>(dev_rho);  \\給定初始值
\\把綁定移到這行cudaBindTexture2D();
kernel <<< >>> (); \\讀取數據

結果我讀到的值都是0,而不是我指定的值,我也嘗試把綁定放到給好初始值後結果也是0

另外,對於cuda by example 中的程式碼他是用for迴圈一直做迭代

我的解讀是當dstOut=1他在進去kernel前texture會先綁定好texIn然後out寫值進去

當下一輪dstOut=0時texture會再一次重新綁定給定新值得out到texOut上然後in寫值進去重複如此

如果參酌手冊的做法的話應該要每次都重新綁定cuArray,可是我沒有看見有關cuArray的程式碼

反而見到他直接把裝置端的變數直接綁上去,等於說手冊的做法需要透過中介而浪費時間

又如果我參照原本手冊上的方式先綁到cuArray再綁到texture上

想要用cudaMemcpyArrayToArray做非同步的寫定有點疑問,如果有下面的程式碼

那程式會怎麼運作?kernel會等到input綁完才進kernel,然後等kernel算完才會綁上result?

但如果有其他程式碼1不是用到input的話就可以跟綁定input時同時執行?

cudaMemcpyArrayToArray(input);
cudaBindTextureToArray(input);
//其他程式碼1
kernel<<<>>>(result,input);
//其他程式碼2
cudaBindTextureToArray(result);


自問自答,找到原因了!
cudaBindTexture2D( NULL, Trho, dev_rho, chDesc, N, N, N*sizeof(float) );
最後一個參數要對齊乘2的倍數

不對!好像還是怪怪的

楼主您好,您如果要使用linear memory来和texture reference绑定, 请确保:

(1)该内存的起始地址在512B边界。(这点您通过cudaMalloc后,可以保证)
(2)该缓冲区的每行地址在32B的边界(pitch = N * 32B, fermi上,kepler可能不同)。(这点您无法保证)

因为(2)您无法保证(例如,您每行是10个,或者2个,或者30个元素这种情况), 所以您的绑定会以无效参数而失败。(cudaErrorInvalidValue)

建议的解决方案: 使用cudaMallocPitch来确保上述要求,或者使用cudaMallocArray。

感謝版主提醒!重新實作後已經成功。

我也順便練習了cudaMallocPitch 、 cudaMemcpy2D 的實作。

恭喜楼主。服务您是我们的荣幸。