同樣的程序放入不同尺寸的數據後Occupancy掉很多

我要做的東西有點類似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_ij
i^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=256
256和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%)? 而且您都说了速度基本符合预期的。:slight_smile:

果断无视它吧。

感谢来访,周末愉快。

後來我參酌NVVP的結果來修改程序,順利的把運行時間降低了!

Nsight的BUG造成虛驚一場!

版主辛苦了!

嗯嗯。根据你的图看,以及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

后者的叫法是和标准手册一致的。

感谢您的来访,夜安。