我要做的東西有點類似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
后者的叫法是和标准手册一致的。
感谢您的来访，夜安。