第一次发帖。
接触GPU编程一年多，最近在CUDA程序优化过程中出现了一个奇怪的问题，跟各位讨论一下，
硬件环境是NVS4200M，ComputeCapability: ‘2.1’ , 软件环境为win7 x64 + vs2010 + CUDA 5.0。
kernel的功能是进行矩阵向量乘法，即
rst = [A] x， A为NN矩阵，rst 与 x为N1的向量
由于矩阵[A]规模很大，不能直接存储，所以A中的元素需要每次进行计算。故无法使用cublas库中的矩阵向量乘法函数。同时，host与device之间的数据传输量很小，与N同阶量级（矩阵为N*N大小）。在kernel中使用单精度浮点运算。
kernel的输入量是长度为N的坐标向量，即N组（float x, float y, float z），为了方便描述，记为pos。矩阵中的元素a_ij由pos[i]与pos[j]确定。
算法大致如下
kernel(float *pos, float *x, float *rst) // x为要乘的向量，rst为存储矩阵向量乘法结果向量的数组
-
i = threadIdx.x + blockDim.x*(blockIdx.x + gridDim.x*blockIdx.y);
-
x = pos[i]
-
pr = 0
-
for j = 0 ~ N
-
y = pos[j]
-
a = f(x,y) // a为矩阵中的元素a_ij, 为了表述方便，这里省略了a的具体计算过程（比较长，只与x,y坐标有关）
-
pr = pr + a*x[j]
- rst[i] =pr
此kernel在memory上进行了比较充分的优化，目前使用global memory的版本已经优化达到了近似shared memory版本程序的速度。在进一步的优化中，发现了如下的问题
现象：假设上述的kernel运行一次时间为120ms，如果把第7行注释掉，运行时间立即降到20ms左右。
需要说明，在for j = 0~N循环中，使用的中间变量都是寄存器，其中第6行实际上是很多行代码，包含了主要的数学计算量，而第7行的计算量很小。
如果把第7行改成：
temp = pr + a*x[j]
pr = temp
程序执行时间为120ms，而注释掉pr = temp后，执行时间降到20ms。
通过这个对比不难发现，使程序性能下降的主要原因并不是第7行中的数学运算，那么是什么原因导致的呢？
欢迎大家讨论，如果有不清楚的地方请指出，我补充说明。
[/i][/i][/i]
一段没有任何输出的计算，一般认为是无意义的，一般会被compiler整块整块的干掉。所以你感觉是一条语句，却加速了6倍。是因为输出被干掉了，大量相关的语句整体被干掉。
如同人脑的神经元。如果没有任何输出，该细胞将死亡。
kernel(float *pos, float *x, float *rst) // x为要乘的向量，rst为存储矩阵向量乘法结果向量的数组
-
i = threadIdx.x + blockDim.x*(blockIdx.x + gridDim.x*blockIdx.y);
-
x = pos[i]
-
pr = 0
-
for j = 0 ~ N
-
y = pos[j]
-
a = f(x,y) // a为矩阵中的元素a_ij, 为了表述方便，这里省略了a的具体计算过程（比较长，只与x,y坐标有关）
-
pr = pr + a*x[j]
- rst[i] =pr
您去掉第七句后，中间过程将整体被干掉。
您的kernel实际变成：
-
i = threadIdx.x + blockDim.x*(blockIdx.x + gridDim.x*blockIdx.y);
-
x = pos[i]
-
pr = 0
- rst[i] = pr
回答的好快，谢谢。
你的意思是，如果我不把结果输出的话，实际上中间的所有计算都没有进行，对吗？
应该是和该输出相关的所有部分会被拿掉，因为编译器认为这个是无用功。
我之前也猜测过可能是这个原因，但是因为没有相关的资料，所以不能确定。
不是计算机专业的，不太懂编译器，请问一下这个问题是nvcc编译器，还以一般编译器都会有的一种“优化”或“清理”措施吗？
谢谢您的回复，我明白了。还好在这里问了一下，否则不知道自己还要在这里面钻多久。
system
12
非常感谢。
nv官方的bbs之前是不是关闭过一次？我记得以前注册过，不过是英文的，当时有一个上海的Andrew Shao还邮件回复过我问题，后来bbs和他的电邮地址全都失效了。
forums.developer.nvidia.com网站是这个地址
system
13
您好，NV的英文版论坛曾经关闭过一段时间，现在已经重新开放了，依然是可以访问的。那个论坛隶属于美国NV总部。
您现在访问的cuda zone中文版是NV中国分公司负责支持的中文版开发者论坛。
两者都是官方提供的（以及其他国家还有其他语言的cuda zone）。
您可以自由选择在这些论坛里面参与讨论，但他们相互之间是分别维护的，帖子内容也不相同。
关于上海的Andrew Shao，您可以向本论坛常驻的NV官方支持人员询问下，看看是不是他们的同事。
总之，欢迎您莅临cuda zone CHINA，祝您在论坛开心~
system
14
这种知识在什么课程里可以获得呢？有什么推荐阅读的书籍吗？谢谢
system
17
听说有一种叫做EXP的东西可以兑换为各种知识，以及您多来论坛潜水/提问可以有效提升EXP数量，这多少是一个可行的方法。
祝您好运~
