為什麼使用shared memory比直接用global存取還要慢?

加速计算 CUDA
1

我要計算的涵式大改長這樣

__global__ void poissonE(double* V,double* VNew,double* rho) {

   int i = blockIdx.x*blockDim.x+threadIdx.x+1;
   int j = blockIdx.y*blockDim.y+threadIdx.y+1;
   int idx = i+j*N;

   float VS,VC=V[idx];

   VS=V[idx+1]+V[idx-1]+V[idx+N]+V[idx-N]+h*h*rho[idx];
   if((i+j)%2 == 0 )
   VNew[idx]=(1 - omega) * VC + omega / 4 * VS;
}

我使用visual profiler觀察資料都是從L2讀取命中率約90%，傳輸速率約50GB/s
但是我把V先寫道shared memory後我發現L2的命中率、讀取速率變慢了！
可是理論上我寫到shared memory可以減少對於global的讀取頻率，應該要變快才是
畢竟shared memory是L1快取，讀取速度應該會更快才是

我shared menmory因為還要處理邊界所以式宣告成[BD+2][BD+2]的矩陣，BD =32;
我編譯下nvcc *.cu -arch=sm_30，使用GTX650

另外，我想請教一下visual profiler我看到有L1快取記憶體，但命中率永遠都是0
是要怎麼樣才能配置L1記憶體進去，在什麼情況下編譯器會把件數配置到L2的快取上?

同樣是L2快取，使用surf或texture會比較快嗎?

謝謝！

2

楼主您好，我阅读了贵代码，我不讨论您的代码是否有正确性方面的问题。然后就题回答您的问题。

因为您需要讨论的问题“为何使用shared memory后会…"，却给出的是没有使用shared memory的代码。和您的问题不对应。所以只能进行大致性讨论。

（1）sm_30的硬件采取的是shared memory和L2(global memory)公用一个控制单元(LSU)，对shared memory的访问和global memory的访问是互斥的。而3.0的L2能1个port对SMX每个周期提供64B的数据, shared memory能提供128B数据(sizeof(float)*32)。但你这里的重复使用率只有4次，而额外加入的读取和写入shared memory的操作就2次了（初始化shared memory)，所以上述2点，带来的好处综合看，可能不大了。

（2）因为您的重复读取都在shared memory里了，减少了从L2的传输，自然命中率和传输率下降了。（注意这个不代表性能会下降，当然也不代表性能会提升，无关的。）

（3）sm_30上。L1 cache不能服务global memory。所以它的global memory hit rate将永远是0. 以及，在3.0上，任何编译参数都无法为global memory使用L1 cache了。请节哀。

（3-1）在sm_30上，因为SM里有了独立的48KB read-only cache, 所以一般情况下是推荐使用它的。但是sm_30和sm_35不同，无法自动使用48KB的read-only cache来缓冲L2/global memory, 所以必须手工使用texture或者surface来利用上这48KB cache (3.5可以使用const __restrict__修饰指针或者__ldg()来自动使用这48KB的cache）。

(3-2)此外，3.0的48KB read-only cache是用texture cache改进来的，可以服务空间性访问，您的每个线程，会读取v[idx], v[idx + 1], v[idx - N], v[idx + N] 是(x,y)附近的临近空间4个点，此时也将享受到福利。

综合(3-1)和(3-2)是建议立刻使用readonly cache的，一般会让您的代码跑的更快。但请注意的是，虽然您的这种代码会因此享受到更大的等效访存带宽，但read-only cache也会引入更大的延迟，请注意掩盖。

感谢您来访CUDAZone China.----------------------------------------------------
代为修正一处笔误。

3

以及，贵卡只能每个周期执行8条double指令(每SMX)， shared memory可以提供128B-256B数据/周期(依然是每SMX)。

而您现在是每个线程将读取7个double(56Bytes), 而进行5条double运算(3条加法，2条乘-加).

所以如果您改善了访存后，很快就要卡死在double计算能力上，而无法提升运行速率很多。

建议您也考虑下购入新的K20显卡。

4

感謝版主的回應。
前面那個程式碼是我真的跑程式時的程式碼用，想說做個簡單的示意

我重新貼一下我的程式碼：
不使用shared memory


 __global__ void kernel(float* V,float* VNew,float* rho) 
   {

   int i = blockIdx.x*blockDim.x+threadIdx.x+1;
   int j = blockIdx.y*blockDim.y+threadIdx.y+1;
   int idx = i+j*N;

   float VS,VC=V[idx];

   VS=V[idx+1]+V[idx-1]+V[idx+N]+V[idx-N]+h*h*rho[idx];
   if((i+j)%2 == 0 )
   VNew[idx]=(1 - omega) * VC + omega / 4 * VS;
   }

使用shared memory


__global__ void kernel(float* V,float* VNew,float* rho) 
{

   int i = blockIdx.x*blockDim.x+threadIdx.x+1;
   int j = blockIdx.y*blockDim.y+threadIdx.y+1;
   int idx = i+j*N;

   float Rho=rho[idx];
   __shared__ float V_s[BD+2][BD+2];

   V_s[threadIdx.y+1][threadIdx.x+1] = V[idx];

   if(threadIdx.y == 0)
   V_s[0][threadIdx.x+1] = V[idx-N];

   if(threadIdx.y == BD-1)
   V_s[BD+1][threadIdx.x+1] = V[idx+N];

   if(threadIdx.x == 0)
   V_s[threadIdx.y+1][0] = V[idx-1];

   if(threadIdx.x == BD-1)
   V_s[threadIdx.y+1][BD+1] = V[idx+1];

   __syncthreads();


   if((i+j)%2 == 0 )
   VNew[idx]=(1 - omega) * V_s[threadIdx.y+1][threadIdx.x+1] + omega / 4 *(V_s[threadIdx.y+1+1][threadIdx.x+1]+V_s[threadIdx.y+1-1][threadIdx.x+1]+V_s[threadIdx.y+1][threadIdx.x+1+1]+V_s[threadIdx.y+1][threadIdx.x+1-1]+h*h*Rho);


   //__syncthreads();

   //VNew[idx] = VN_s[threadIdx.y+1][threadIdx.x+1] ;

}

可能有點混亂，我實際跑程式的時候用double，但是在測試時是用float，最後應該還是會統一改成用float。
sites.google.com/a/ihakka.twbbs.org/home/dang-an-ku/test2.csv?attredirects=0&d=1
sites.google.com/a/ihakka.twbbs.org/home/dang-an-ku/test6.csv?attredirects=0&d=1
上面兩個檔案是我重新分別測是沒有使用、有使用shared memory的測試結果，不使用大概比使用快了1.1倍。

關於版主所言，我可以這樣理解嗎？
不使用shared memory時從global讀取的時間是BDBD(3232)個thread要讀取5個double除以讀取的時間
→3232564/64=5120個週期
使用shared memory所花的時間是從global讀取寫入shared再從shared的總時間
→343464/64+343464/128+3232645/128=4294個週期
兩者所花的時間比為0.839，如果再加上__syncthreads()就更慢了？

texture、surface記憶體的部分我還沒有很熟悉，可能知後練練再上來請教。

關於L1的48K就是寫死了，不能再透過cudaFuncSetCacheConfig( Kernel, cudaFuncCachePreferL1 );調整？

至於K20我買不起…也得等我把cuda上手後才有辦法申請經費

5

楼主您好，读取需要的时间会掩盖在计算里（或者反过来说也成立）。

如果要单独的看，需要传输的次数是总数*sizeof(float) / 64 / 2(GTX650应该有2个L2到SMX的ports的。我没测试）次传输。

以及，楼主应该用平均效率来看，也许更好理解：
例如shared memory在4B * 32读取的时候，可以一次提供128B数据。
在8B * 32读取的时候，可以一次提供256B效率。
而L2到SMX是64B宽度，一次可以传输64B数据(非官网数据，请谨慎使用。官网没公布数据）。
所以如果提前将数据复制到了shared memory里，再多次的读取，才有效率上的提升。

至于楼主为何用shared变慢，准确的直接的解释我不知道哦。（我已经连续在论坛今日工作快15小时了。有点困乏了）

texture可以用3.0新增的texture object访问，在3.0一般情况总是推荐的，简单好用。

至于L1, 可以修改的。在3.0上可以改成16KB(PreferShared), 32KB(PreferEqual), 48KB(PreferL1)。

K20不买也行，可以考虑买入手个Titan.

6

有个问题：
三维surface或者texture有没有针对数据访问局部性做优化？

7

您已经在另外一个帖子里面提出了您的问题，建议新问题单独发帖，以及实在需要回帖提问的话，在内容相关的帖子里面回复，还是勉强可以的，但是像您这样，同样的问题多次回帖提问+发布在内容不相关的主题帖的回帖里面，这样是不可以的，有灌水嫌疑，请您不要再如此操作，切记。

8

[table=525]
[tr][td=93]method\size
[/td][td=72]
64
[/td][td=72]
256
[/td][td=72]
512
[/td][td=72]
1024
[/td][td=72]
2048
[/td][td=72]
4096
[/td][/tr]
[tr][td]no_shared[/td][td]
8.108
[/td][td]
241.608
[/td][td]
307.57
[/td][td]
1218.16
[/td][td]
4853.21
[/td][td]
19540.5
[/td][/tr]
[tr][td]shared_if(x,y)[/td][td]
23.9834
[/td][td]
245.907
[/td][td]
310.588
[/td][td]
1214.63
[/td][td]
4827.86
[/td][td]
19355.8
[/td][/tr]
[tr][td]shared_if(x)[/td][td]
35.3402
[/td][td]
252.532
[/td][td]
309.693
[/td][td]
1208.63
[/td][td]
4770.82
[/td][td]
19068.2
[/td][/tr]
[tr][td]shared_if(y)[/td][td]
39.1761
[/td][td]
242.566
[/td][td]
364.546
[/td][td]
1541.84
[/td][td]
6156.9
[/td][td]
24668.3
[/td][/tr]
[tr][td]shared_and_global
[/td][td]
8.6817
[/td][td]
76.7193
[/td][td]
284.799
[/td][td]
1167.88
[/td][td]
4674.14
[/td][td]
18734.6
[/td][/tr]
[/table]
變數使用float型態，跑step2000步，計時單位為(ms)，機器為GTX650、編譯參數 nvcc kernel.cu -arch=sm_30

原先我猜測是因為在處理邊界時不同warp會產生空等的問題，

所以我做了幾個實驗
no_shared 就是最直接用global的方式下去做
shared_if(x,y) 是上下的邊界給if(y)去做(用idx.x)，左右則是用if(x)去做(用idx.y)
shared_if(x) 就是通通叫idx.y去做
shared_if(y) 則相反
shared_and_global 是上下給shared而左右邊界給global去做

得到以上的計時結果

我推測 no_shared 不慢的原因是因為取上下左右時都滿足合併存取
通通丟給idx.x去做的時候大概是因為不滿足合併存取，讀了很多只有一個有用其他丟掉，
在同一個wrap讀的時候，可能幾個thread就把頻寬占滿了後面的只能排隊
用idx.y去做時估計因為在不同的wrap上雖然不滿足合併存取但可以同是做事
不過這卻不能解是為什麼shared_and_global卻可以這麼快

估計可能shared_memory有使用上的一些限制吧！

我改用texture就直接快了1.7倍了！

9

楼主您好，无法评论您的数据。

如果您非得需要评论，请给出详细的各种代码和各种测试数据集。

以及，在sm_30上使用texture, 特别是您1#中这种有着明显空间关系的数据(-1,+1,-N,+N（N是宽度))，用texture是推荐的（如同我在2#说过的)。您得到1.7倍的提升应该是正常的。如果您认为提升的数据不正确，请依然给出测试代码和数据。

10

横扫版主说：
“在sm_30上，因为SM里有了独立的48KB read-only cache, 所以一般情况下是推荐使用它的。但是sm_30和sm_35不同，无法自动使用48KB的read-only cache来缓冲L2/global memory, 所以必须手工使用texture或者surface来利用上这48KB cache (3.5可以使用constant __restrict__修饰指针或者__ldg()来自动使用这48KB的cache）。”

我想问下，sm_20下的read-only cache有多大，可以采用constant __restrict__修饰指针来自动使用吗？

11

您好：SM2.x和SM 3.X的cache体系和作用划分不同，并无3.x意义上的read-only cache，因而无法直接比较。以及在2.x上及建议直接使用，L2cache要比纹理cache效果更好一些。

大致如此，祝您好运~