我要計算的涵式大改長這樣
__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會比較快嗎?
謝謝!
楼主您好,我阅读了贵代码,我不讨论您的代码是否有正确性方面的问题。然后就题回答您的问题。
因为您需要讨论的问题“为何使用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.----------------------------------------------------
代为修正一处笔误。
以及,贵卡只能每个周期执行8条double指令(每SMX), shared memory可以提供128B-256B数据/周期(依然是每SMX)。
而您现在是每个线程将读取7个double(56Bytes), 而进行5条double运算(3条加法,2条乘-加).
所以如果您改善了访存后,很快就要卡死在double计算能力上,而无法提升运行速率很多。
建议您也考虑下购入新的K20显卡。
感謝版主的回應。
前面那個程式碼是我真的跑程式時的程式碼用,想說做個簡單的示意
我重新貼一下我的程式碼:
不使用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上手後才有辦法申請經費
楼主您好,读取需要的时间会掩盖在计算里(或者反过来说也成立)。
如果要单独的看,需要传输的次数是总数*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.
有个问题:
三维surface或者texture有没有针对数据访问局部性做优化?
您已经在另外一个帖子里面提出了您的问题,建议新问题单独发帖,以及实在需要回帖提问的话,在内容相关的帖子里面回复,还是勉强可以的,但是像您这样,同样的问题多次回帖提问+发布在内容不相关的主题帖的回帖里面,这样是不可以的,有灌水嫌疑,请您不要再如此操作,切记。
[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倍了!
楼主您好,无法评论您的数据。
如果您非得需要评论,请给出详细的各种代码和各种测试数据集。
以及,在sm_30上使用texture, 特别是您1#中这种有着明显空间关系的数据(-1,+1,-N,+N(N是宽度)),用texture是推荐的(如同我在2#说过的)。您得到1.7倍的提升应该是正常的。如果您认为提升的数据不正确,请依然给出测试代码和数据。
system
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__修饰指针来自动使用吗?
system
11
您好:SM2.x和SM 3.X的cache体系和作用划分不同,并无3.x意义上的read-only cache,因而无法直接比较。以及在2.x上及建议直接使用,L2cache要比纹理cache效果更好一些。
大致如此,祝您好运~