我要做的東西有點類似nbody的程序,要處理 \sum_{j=0}^N K_iji^sin(theta_i - theta_j)
N 是我的數據的尺寸
K_ij 是一個係數矩陣,對應到程式是 f[]
程式主要就是依序給f[j]去找對應的theta_j
然後去計算\sum_{j=0}^N K_iji^sin(theta_i - theta_j)
第三層迴圈是把sum展開來
__global__ void cuda2( cudaTextureObject_t tex, const float* __restrict__ theta ,float *s)
{
const int i = blockIdx.x*blockDim.x+threadIdx.x;
const int xi = i / sqrtN;
const int yi = i % sqrtN;
const int mid = (sqrtN-1)/2;
float f;
float local_s[4];
for(int xj=0; xj<sqrtN; xj++)
{
for(int yj=0; yj<sqrtN; yj=yj+4)
{
#pragma unroll 4
for(int count=0;count<4;count++)
{
if(xj==mid && yj==mid);
else
{
const int j = sqrtN*xj+yj+count;
//transform (mid,mid)->(xi,yi)
int xj2 = xj+(xi-mid); cuda3(xj2);
int yj2 = yj+count+(yi-mid); cuda3(yj2);
int j2 = sqrtN*xj2+yj2;
//the connection of "i" to "j2" is "d"
f = tex1Dfetch<float>(tex, j);
local_s[count]+=f*sin(theta[j2]-theta[i]);
}
}
}
}
s[i] = local_s[0] + local_s[1] + local_s[2] + local_s[3];
}
當N 的尺寸是 128 * 128 時 Occupancy 約 96% ,kernel 執行的時間約 0137 ms
當N 的尺寸是 256 * 256 時 Occupancy 剩 <5% ,kernel 執行的時間約 3000 ms
時間差將近22倍,若再把尺寸加大的話Occupancy 還是很低
然後為什麼Warp Issue Efficiency會破100%
進去看Issue stall 所占的比例成份差異很大,我不太曉得是什麼原因造成這樣的差異
楼主您好,
您并没有报告您的block size以及sqrtN以及cuda3()等是什么。
以及,您的.rar附件里啥都没有的。
看残缺的代码看不出什么的,
所以建议您补充下信息?
附加檔案裡面有兩個檔案N_128_128.nvreport 、 N_256_256.nvreport
是用CUDA5.5、最新版Nsight3.0.0.13150生成的報告,拉到VS去應該可以打開吧!
如果不能開我晚點用貼圖的好了
N = sqrtN * sqrtN ,sqrtN就是二維數據的邊長
cuda2<<< N / 512 , 512 >>>()
blocksize 為512 ,grid = N / 512
cuda3 是處理計算j2時超出邊界時的情況
device forceinline int cuda3(int x)
{return ( (x % sqrtN) + sqrtN ) % sqrtN ;}
您的2个文件里无任何有用信息的,不信您可以直接用notepad打开。
以及,不看这个,根据您提供的信息,您的代码无法理解为何会突然降速几十倍。(从100多ms到3000ms) 这个真心不知道。
以及,您补充的( (x % sqrtN) + sqrtN ) % sqrtN 等价于式子: x % sqrtN, 您确定您没写错?
感谢您的来访,周末愉快。
如果非要找个理由的话,
这里:
local_s[count]+=fsin(theta[j2]-theta[ i]);
您的j2在之前N=256256和128*128的情况下,范围从~64KB增加到了~256KB, 从而正好导致L2抖动,从而正好性能急剧下降?按理说不应该啊。
建议楼主给出给出profiler的数据报告,否则肉眼看代码真心看不出什么。(请注意,这并不表示其他会员、版主、NVIDIA官方支持、总版主肉眼看不出什么,只是说我自己)。
楼主既然用的nsight里的小profiler, 你可以看下“Issue Stall Reasons"图,如果方便,请发上来,这样看看能否看出什么。
那東西確實不能用notepad開要在VS裡面開,不過不管他,這晚點再貼圖處理好了!
( (x % sqrtN) + sqrtN ) % sqrtN寫這樣是因為負數取餘數會有負數,但最終餘數是要正數,所以才這樣寫
我仔細想了一下我的程序的複雜度是O(N^2),所以N加大4倍消耗時間要增加16倍,從結果看大概是差21倍,感覺還算合理。
後來我用visual profiler跑一次結果後發現Occupancy依然是90%以上
如果這樣來看的話似乎是Nsight裡面出了些問題,否則Occupancy掉這麼多的話執行時間應該會拉成很多。
晚點再貼profier、nsight跑的結果
(1)我知道是用VS打开的,我让你用notepad的目的是让你看看里面到底有没有东西!
(话说你给的2个文件,里面真心没任何有效信息的,你可能没发全)
(2)嗯嗯。你说的对。我以为您的x始终都是正数呢。这是我考虑不周全。
(3)是的。这里我刚才也欠考虑了,您的一个线程是O(N)的(2个sqrt(n)的循环), 然后线程数目也增加了N倍,的确应该至少16倍的。(主要我刚才没想这点,光看你说occupancy下降几十倍)
(4)关于为何nsight自带的小profiler报告achieved occupancy < 5%,而大的nvvp报告>90%, 这个真心不知道。可能是它的BUG?
既然如此,那楼主速度下降还不多,一切都只是虚惊一场?
我把Nsight跑出來的結果放上來了,
上面的是N = 256256,下面的N = 128128
跑出來的結果就差很多
[attach]3262[/attach]
[attach]3261[/attach]
大概不曉得Nsight哪裡出錯了吧!害我不曉得該如何改進程序
版主也辛苦了!
建议无视nsight的profiling结果,而是使用visual profiler的为准:
(1)首先我们不知道nsight的achieved occupancy是否和visual profiler的一样(即:block里提前死去的warp会拖累occupancy)。不知道具体这个1%的occupancy是怎么计算出来的。
(2)既然visual profiler报告的achieved occupancy在90%,而且速度没有下降,那么果断建议无视nsight里的小profiler结果,而是相信visual profiler的。
有人报喜,有人报忧,为何不使用报喜的结果呢(NVVP的90%)? 而且您都说了速度基本符合预期的。
果断无视它吧。
感谢来访,周末愉快。
system
10
後來我參酌NVVP的結果來修改程序,順利的把運行時間降低了!
Nsight的BUG造成虛驚一場!
版主辛苦了!
system
11
嗯嗯。根据你的图看,以及nsight的说法,这个应该是nsight的BUG.
你的图报告的active warps只有0.7(根据nsight手册说法,这个实际上是resident warps)。
所以必然是BUG的。
以及,一般我们常用的是visual profiler, 因为visual profiler里的用词和programming guide之类的手册比较一致,推荐用这个。
当你从nsight转向独立的visual profiler后,请注意一些名词变化:
nsight的active warps → nvvp的resident warps
nsight的eligible warps → nvvp的active warps
后者的叫法是和标准手册一致的。
感谢您的来访,夜安。