我原本做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);