identifier cudaSurfaceObject is undefined

我參照 Nvidia 手冊使用 surface memory 但是發生錯誤
[
我從網站上複製貼上,結果還是一樣,手冊說CC 2.0以上都可以使用,我的卡是GTX650 下的指令是 -arch=sm_30

以下是我的程式碼

#include <stdio.h>
#include <iostream>
#include <fstream>
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
using namespace std;

__global__ void kernel_surf(cudaSurfaceObject V, cudaSurfaceObject VN, cudaSurfaceObject rho)
{

   ... ... ...
   surf2Dread(&r,rho,i,j, cudaBoundaryModeTrap);
   temp= 0.25* (VR + VD + VL + VU - h*h*r);

   surf2Dwrite(temp, VN, i,j, cudaBoundaryModeTrap);
}

int main()
{

   float *V, *dev_V, *dev_VNew, *dev_rho;

   cudaHostAlloc((void**)&V,N*N*sizeof(float),cudaHostAllocDefault);        
   cudaMalloc((void**)&dev_V,N*N*sizeof(float));
   cudaMalloc((void**)&dev_VNew,N*N*sizeof(float));
   cudaMalloc((void**)&dev_rho,N*N*sizeof(float));

   cudaChannelFormatDesc chDesc = cudaCreateChannelDesc<float>();
   cudaArray *cu_rho,*cu_V,*cu_VN;
   cudaMallocArray(&cu_rho, &chDesc, N, N, cudaArraySurfaceLoadStore);
   cudaMallocArray(&cu_VN, &chDesc, N, N, cudaArraySurfaceLoadStore);
   cudaMallocArray(&cu_V, &chDesc, N, N, cudaArraySurfaceLoadStore);
   cudaMemcpyToArray( cu_rho, 0, 0, dev_rho, N*N*sizeof(float),cudaMemcpyDeviceToDevice );
   cudaMemcpyToArray( cu_VN, 0, 0, dev_VNew, N*N*sizeof(float),cudaMemcpyDeviceToDevice );
   cudaMemcpyToArray( cu_V, 0, 0, dev_V, N*N*sizeof(float),cudaMemcpyDeviceToDevice );

   // Create the cuda resource description
   struct cudaResourceDesc resDesc;
   memset(&resDesc, 0, sizeof(resoureDescription));
   resDesc.resType = cudaResourceTypeArray;    // be sure to set the resource type to cudaResourceTypeArray
 
   // Create the surface object
   resDesc.res.array.array = cu_rho;    // this is the important bit
   cudaSurfaceObject_t su_rho = 0;
   cudaCreateSurfaceObject(&su_rho, &resDesc);
   resDesc.res.array.array = cu_VN;    // this is the important bit
   cudaSurfaceObject_t su_VN = 0;
   cudaCreateSurfaceObject(&su_VN, &resDesc);
   resDesc.res.array.array = cu_V;    // this is the important bit
   cudaSurfaceObject_t su_V = 0;
   cudaCreateSurfaceObject(&su_V, &resDesc);

   ... ... ...

   if(tf)
   {kernel_surf<<<grid,block>>>(su_V,su_VN,su_rho);}
   else
   {kernel_surf<<<grid,block>>>(su_VN,su_V,su_rho);}
   
   ... ... ...

   // Destroy surface objects
   cudaDestroySurfaceObject(su_rho);
   cudaDestroySurfaceObject(su_VN);
   cudaDestroySurfaceObject(su_V);

   // Free device memory
   cudaFreeArray(cu_rho);
   cudaFreeArray(cu_VN);
   cudaFreeArray(cu_V);

   cudaFree(dev_V);
   cudaFree(dev_VNew);
   cudaFree(dev_rho);

   cudaFreeHost(V);

   cudaEventDestroy(start);
   cudaEventDestroy(stop);

   cudaThreadExit();


}

楼主您好,
请使用cudaSurfaceObject_t代替cudaSurfaceObject
请使用sizeof(struct cudaResourceDesc)代替sizeof(resourceDescription)

第一个是手册的错误,请原谅手册编写组。

第二个是您的错误,请自我原谅。

感謝版主的回應,編譯過了!

不講手冊有錯,我大概一輩子也找不到吧!

不過跑出來的結果怪怪的!

原本的計算時間大約要20秒,現在變成0秒,我在研究看看

再次感謝版主!

您客气了。服务您是我的荣幸。

任何东西都无法强求十全十美,但是我们可以尽量辩证的看。

感谢您的来访。

我找出問題了!就自己程式碼打錯!
surf2Dwrite定位時x方向要乘上4倍這跟texture不一樣

感谢楼主,因为我用2.x一直,很少关注surface.

未能之前知道surface使用的坐标是字节的,而不是元素的。

的确和texture不同,需要sizeof(元素), 例如您这里的X需要4。

以及,我刚才也实践了下,1D需要4, 而2D的需要X4即可, Y的确不需要变化(可能是因为内部Y*宽度了,于是texture和surface对Y都一样)。

通过这个,我得到了教育,感谢楼主!!

以及,我为我之前未能向您指出您代码里的这个BUG表示歉意。

大家互相切磋切磋才會進步嘛!

不過用surface後我以為不用複製來複製去的會比texture還要快可惜沒有!

詳細原因可能還要再研究

另外補充一下
surf2Dwrite(temp, VN, i,j, cudaBoundaryModeTrap);
寫入的值要很小心,surf宣告成float 那寫入的值就一定要是float型態。
如果要放其他型態的值進去要先強制轉換才不會出錯。

最後,我實驗的結果發現surface Obj雖然提供寫入的功能
但是依然沒有比tex Obj直接綁定變數來讀/寫還要快
目前的結論是不用surf改用tex Obj綁定變數
kernel重新啟動後tex會再更新值,這其實就跟surf寫入功能是一樣的
只是不用再透過cuArray中介讀取

(1)
楼主道:
“另外補充一下
surf2Dwrite(temp, VN, i,j, cudaBoundaryModeTrap);
寫入的值要很小心,surf宣告成float 那寫入的值就一定要是float型態。
如果要放其他型態的值進去要先強制轉換才不會出錯。”
----这个我不能赞同。
(1A) surface只有简单的访存功能(以及利用texture/readonly cache的功能), 而不涉及具体的数据的。所以您可以对一个surface写入任何种类的数据的。
(2A)surface本身不能有数据类型。但是它绑定的cudaArray中的每一个元素具有类型(4个分量的大小,是否是float还是整数)。但此类型会被surface无视的。
以上2点,应该不存在“宣告成float的surface和只能写入float"的问题。

(2)
楼主继续写道:“
最後,我實驗的結果發現surface Obj雖然提供寫入的功能
但是依然沒有比tex Obj直接綁定變數來讀/寫還要快

----您确定您的texture能写入,而且比您的surface更快么???
因为我不认为texture可以直接像surface一样的方便的在kernel进行写入,所以我不认为您的“甚至texture写入得比surface"更快的说法成立。

(3)
楼主还写道:“

  • 目前的結論是不用surf改用tex Obj綁定變數
    kernel重新啟動後tex會再更新值,這其實就跟surf寫入功能是一樣的
    只是不用再透過cuArray中介讀取”
    –用texture代替surface来写入,上文说了,这个恐怕不行,前者没写入功能。
    –以及,就算您是考虑直接写它后备的cudaArray, 但是此类型依然是用户透明的类型,您无法直接在kernel里写入一个cudaArray的(除了通过surface).

楼主三思????

我講一下測試的結果,主程式宣告的部分(正確的來講應該是說cuArray是宣告成float的形式)

float  *dev_V ;
   size_t pitch_V;

   //Malloc memory
   cudaMallocPitch((void**) &dev_V  , &pitch_V  , sizeof(float) * N, N);

   cudaChannelFormatDesc chDesc = cudaCreateChannelDesc<float>();
   cudaArray *cu_V;
   cudaMallocArray(&cu_V, &chDesc, N, N, cudaArraySurfaceLoadStore);
   cudaMemcpy2DToArray( cu_V  , 0, 0, dev_V  , pitch_V  , N*sizeof(float), N , cudaMemcpyDeviceToDevice );

   // Create the cuda resource description
   struct cudaResourceDesc resDesc;
   memset(&resDesc, 0, sizeof(resDesc));
   resDesc.resType = cudaResourceTypeArray;    // be sure to set the resource type to cudaResourceTypeArray
   
   // Create the surface object
   resDesc.res.array.array = cu_V;    // this is the important bit
   cudaSurfaceObject_t su_V = 0;
   cudaCreateSurfaceObject(&su_V, &resDesc);

   initial<<<grid,block>>>(su_V);
cudaMemcpy2DFromArray (    V, N*sizeof(float) , cu_V ,0 ,0 ,N*sizeof(float) ,N ,cudaMemcpyDeviceToHost);
//write file


測試的kernel部分

__global__ void initialR(cudaSurfaceObject_t R) 
{

   int i = blockIdx.x*blockDim.x+threadIdx.x+1;
   int j = blockIdx.y*blockDim.y+threadIdx.y+1;
//surf2Dread
//float
   surf2Dwrite(float(-8), R, 4*i,j, cudaBoundaryModeTrap);
//int
surf2Dwrite(-7, R, 4*i,j, cudaBoundaryModeTrap);
   
}

我先測試過把d_V都設定成-1才綁到cuArray上,不做任何事情直接寫檔案,確認讀到的值都是-1,因此我推斷在宣告、複製回主機等程式碼是沒問題的,之後我才把有關於surf2Dread/write的kernel加入。

然後我在測試直接在kernel中用surf2Dread來讀值,發現正確無誤,最後我把值用surf2Dwrite寫進去。一開始我沒有轉換型態,一直得到錯誤的結果,後來我在想究竟是哪裡錯誤,我想想可能是型態的問題,我做了float()轉換後發現結果就對了

我這邊敘述的測試可能跟上面主程式端的程式碼有一點點差異,但是基本上主程式端的部分是沒問題的。

關於tex寫入的部分,參考cuda by example看到他直接把tex綁到變數上不經過cuArray,然後重複啟動kernel就更新tex值

遂我做了以下實驗

__global__ void kernel_tex( float *V, size_t pitch ,cudaTextureObject_t TV  )
{

   int i = blockIdx.x*blockDim.x+threadIdx.x+1;
   int j = blockIdx.y*blockDim.y+threadIdx.y+1;
   int idx = i+j*pitch;
   
if(idx == 0)
{
   printf("before Write\n");
   V[idx] = V[idx] +1;
   printf("after Write\n");
}
   
}


int main()
{

   float  *dev_V ;
   size_t pitch_V, pitch_rho;

   //Malloc memory
   cudaMallocPitch((void**) &dev_V  , &pitch_V  , sizeof(float) * N, N);
   
   //TextureSet
   cudaResourceDesc resDesc;
   memset(&resDesc, 0, sizeof(resDesc));
   resDesc.res.pitch2D.desc = cudaCreateChannelDesc<float>();
   resDesc.resType = cudaResourceTypePitch2D;
   resDesc.res.pitch2D.width = N;
   resDesc.res.pitch2D.height = N;

   cudaTextureDesc texDesc;
   memset(&texDesc, 0, sizeof(texDesc));
   texDesc.readMode = cudaReadModeElementType;

   //Texture Obj
   resDesc.res.pitch2D.devPtr = dev_V;
   resDesc.res.pitch2D.pitchInBytes =  pitch_V;
   cudaCreateTextureObject(&TV, &resDesc, &texDesc, NULL);
   
   //initial
   initial <<<grid,block>>>(dev_V   , pitch_V   / sizeof(float) );

   for(int step=0;step<2;step=step+1)
   { kernel_tex<<<grid,block>>>(dev_V , pitch_V / sizeof(float),TV); }

   float elapsedTime;
   cudaEventRecord(stop, 0);
   cudaEventSynchronize(stop);
   cudaEventElapsedTime(&elapsedTime, start, stop);
   cout << " GPU time = " << elapsedTime  << " (ms)" << endl;

   cudaMemcpy2D( V ,N*sizeof(float), dev_V, pitch_V, N*sizeof(float), N ,cudaMemcpyDeviceToHost);

... .... ....

}

一開始我用initial 的kernel寫入 -9 再綁到texture上,關鍵在於for迴圈的kernel
我先用tex2d讀出值來,然後把他+1寫回綁定的變數上去,在讀值出來,會發現讀出來的值都是進去kernel的值
但是到下一個迴圈後讀出來的值就是我寫進去的值。

(這邊我所描述的部分也跟我付的主程式有些差異。)

所以我得到的結論是如果texture直接綁變數的會其實可以寫入值進去,只是他不是直接寫到texture裡面
我用了這個方式做了其他程式得到的結果跟用global結果看起來是滿一致的

最後我把結果實作到我的程式,計時結果如下:可以看到用sufObj居然還比global還要慢
size 64 256 512 1024 2048 4096
GM 08.1080 241.608 307.570 1218.16 4853.21 19540.5
TexObj 8.60365 58.2406 198.489 754.691 2966.99 11905.5
SufObj 16.9739 168.491 644.159 2567.96 10242.8 40978.7

以上是我的實驗,或許可能有些細節的部分做錯導致錯誤的結果,還請版主指教

您应该直接指出能证明您的3个观点的语句段落。而不是让我阅读您的大量代码,替您总结。

实际上我在前楼已经表达过我的观点了。
我不能立刻转换角色,放弃我的说话,化身为您,来替您证明。

我希望您能直接自证。

抱歉!程式碼只是示意的部分,我主要想把具體的宣告寫出來參考。

容我花幾些分鐘重新編輯,用更多的文字來說明我的實驗。

我要先睡了,就不等你了。我先说几个事情吧,都是事实:

(1)surface本身无类型,可以写入任何类型的数据。
(2)surface的写入最坏要下次kernel启动才能生效。
(3)cudaArray不能在kernel里直接写入。
(4)surface和texture的后备的线性内存(即非cudaArray)可以在kernel里写入,但最坏要下次启动才能生效。

以及内容基本都可见于手册。

先講surface的問題由於實作的時候得到錯誤的結果,所以我先把有關surf的kernel的刪掉,只剩下宣告、寫值、綁定到surf、回傳讀值確定沒錯。

然後我把surf的kernel寫入,我先做surf2Dread看看讀出來的有沒有錯結果沒錯,所以我猜測錯誤的關鍵在於寫入值的部分,我嘗試把寫入的值加上float()轉換發現結果正確、拿掉就錯了,詳細的宣告可以參考前面的那篇程式碼,關鍵的程式碼:

    //float
surf2Dwrite(float(-8), R, 4*i,j, cudaBoundaryModeTrap);  //與cuArray宣告的float型態一致,結果正確
   //int
   surf2Dwrite(-7, R, 4*i,j, cudaBoundaryModeTrap); //與cuArray宣告型態不一致,結果錯誤
float temp = -6;
   surf2Dwrite(-7, R, 4*i,j, cudaBoundaryModeTrap); //結果正確

這部分我認為也有待商榷,可能我明天再做一次實驗來確定。


關於tex寫入的部分,參考cuda by example看到他直接把tex綁到變數上不經過cuArray,每次都是把更新值寫到綁定的變數上,又不用重新綁定tex,所以我就推斷"如果tex不經過cuArray直接綁變數的話每次kernel啟動會更新值"

遂我做了以下實驗,我把關鍵的程式碼列出來

//在kernel理的關鍵3行
 printf("before Write %f\n",tex2Dread);
   V[idx] = V[idx] +1;
   printf("after Write %f\n",tex2Dread);

//在主程式裡的關鍵2行
for(int step=0;step<2;step=step+1)
   { kernel_tex<<<>>>(dev_V , pitch_V / sizeof(float),TV); }

我的Tex是直接綁上dev_V的,kernel做的事其就是先重tex中讀值,然後把值寫入tex綁定的變數上,再把值讀出來
我測試印出來的結果是
before Write -7
after Write -7
before Write -6
after Write -6
看起來似乎就是kernel重新啟動後就會重新更新值一樣,這樣就實現了"tex寫入"的功能,或許這邊我用的詞會誤解

最後基於以上兩個實驗,我把surfObj、"tex寫入"的程式碼應用到我另外的程式上去檢查,發現跟global的結果是一致
正確一點講是跑出來的圖形一樣,數量級也一樣但我沒有仔細分析細部的誤差

最後我做了計時,結果如下:可以看到用sufObj居然還比global還要慢
size 64 256 512 1024 2048 4096
GM 08.1080 241.608 307.570 1218.16 4853.21 19540.5
TexObj 8.60365 58.2406 198.489 754.691 2966.99 11905.5
SufObj 16.9739 168.491 644.159 2567.96 10242.8 40978.7

以上僅供版主參考,或許還有不清楚、錯誤的地方。


剛剛看了版主的話,的確
(2)surface的写入最坏要下次kernel启动才能生效。
(3)cudaArray不能在kernel里直接写入。
(4)surface和texture的后备的线性内存(即非cudaArray)可以在kernel里写入,但最坏要下次启动才能生效。

這些我也同意,只是我可能用詞不好造成誤解,真抱歉,"tex寫入"其實我想表達的就是第4點

上面的說明都是想表達版主說的這幾件事情

關於(1)surface本身无类型,可以写入任何类型的数据。

我想我明天在做實驗看看究竟是不是我出錯

只是在效率上我計時的結果確實是surf比較慢

嗯嗯。实际上一个写入是直接越过了texture unit, 直接写显存,一个是通过surface写到cudaArray(然后其后备的显存)。

您完全可以用更快的直接写显存,但这样需要使用线性内存,无法使用cudaArray了。

以及在都使用cudaArray的时候,texture是无法写入的。

大致如此吧。

以及,您的确可以直接读写texture的后备的线性内存(即直接读写普通的global memory)。

但你这还是“surface比texture么”,不如直接说直接比较surface读取和普通读取吧。

以及,直接写入global memory可能上去一时快,但可能会带来之后的读取不利的代价(普通线性显存VS cuda array)。

以及,楼主您的“surface写入的速率慢的问题”也可能是错觉。因为这可以容易得到一个结论:用surface写入永远无法到达显卡的最大写入带宽(因为您说了还有更大速率的直接写入方式存在)。

但实际上用surface写入也能达到最大写入带宽的。您的实验可能只是无法展现出这个最大值来。(例如无法掩盖的延迟)

每个warp在结束的时候,都要等待自己的写入完成到某个阶段,然后才能结束。您只有一次写入的surface, 可能会造成大量warps无法及时结束而无法展现带宽,导致您认为不如直接写入了。

您可以重复改成每个线程4x份量的surface写入,看看有无效果。

確實版主說的"tex讀取global寫入"會比我的用詞還要精準

也確實聽版主這麼說可能是改用surface後kernel的策略要調整才能達到比較好的效果

這幾天可能先忙我的模擬順便想想怎麼樣使用surface會比較好,到時候在說說我做的結果