强制支持ICE!!
以及强制赞美那些，伟大、英明、睿智、可爱的，用.x分量做y下标，用.y分量做x坐标的人！
（我不是说楼主！）
我也应该见ICE的帖子而自勉！
感谢答复！
那么，我可以这么下结论：一个block里处理一列数据，这个方法没有办法提高访存效率。
接着考虑另一个方法，考虑运算前，把行列转置一下。
如果数据已经是放在显存里，逐行保存的。下面这个内核函数有瓶颈吗？
读取的时候，访问数据可以被合并，写入数据能被合并吗？
global void transposeM(const int *a, int *b, int w, int h)
{
int x = blockIdx.x * blockDim.x + threadIdx.x;
int y = blockIdx.y * blockDim.y + threadIdx.y;
//check border
…
b[x * h + y] = a[y * w + x];
}
谢谢！
我提出这个问题，是在实际应用碰到的问题：要对已经放置在显存里的逐行保存数据，做列方向的运算。
看起来这个问题没有更好的办法了
LZ您好：
我表示无法认同您在22#定下的结论。
也许随便一个单一block处理一列数据的算法我还不能随便下结论，但是就您给出的具体例子，已经在19#给出了详细分析，以及给出了一个合并访问的实现的核心代码。此代码经测试，可以比您之前您的测试结果的访存速度高出一个量级。
按照举一反例即可推翻命题的逻辑原则，我有充分把握和理由不认同您的结论。
以及鉴于我无法干涉您的人身自由（包括认识自由），我谨建议您重新观察和实验下19#的实现方法，但并无法强迫您接受我的观点，或者放弃您的固有观点。
2：无法完整回答您的核函数瓶颈如何，因为您没有给出完整的代码。
以及关于此写入是否能合并访问，请自行判断。您可以参考programming guide的内容，以及本版的其他讨论帖。
如果您不能自己做出结论，我将不再回答您的问题。
因为最近几天以来，我们花费了大量的时间在和您讨论访存问题，本帖已经写到第三页，累计回答数千字以上，而您如果是在不明白基本概念的前提下一直询问的话，这真心让人伤心的，希望不是这样。
祝您好运~
LZ您好，请重新思考19#相关内容，此建议已经说过多次，这是最后一次了。
您再发表类似的观点，将被我无视的。
祝您好运~
其实楼主死活不看ICE帖子，真心不对。
ICE已经说了很明白了，您的y和x颠倒定义即可高速。。。真心无奈楼主你。
你想想是这个道理不？
谢谢ice！
您在19#提到的方法，我已理解。
事实上，在我们的应用中，对列方向的运算比测试代码里上下三个元素做平均要复杂的非常非常非常多！所以解决这个问题，我们把重新编写算法作为最后一个办法。
anyway，非常感谢！
LZ您好：
我无法判断您要实现的目标代码是否具备更好的访存方式，以及我尊重您对于您具体问题解决策略的优先级判断，因为这个涉及开发成本，需要您承担而不是我。
需要指出的是，前面的所有回帖内容都只针对本帖出现的内容，我无法脑补出您没有给出的实际问题，并作出相关建议的。
以及，一般地，除了事先做好一份转置数据以外，您也可以在kernel内部将局部需要使用的数据，以合并访问的方式读入shared memory，然后再从shared memory读取，最后的回写也是一样。
整个kernel的流程为 合并访问读入shared memory——访问shared memory进行计算并回写到shared memory——合并访问的方式回写global memory。这样虽然增加了环节，但是由于global 访存都是合并的，也是划算的，以及或许您可以用得上。
大致如此了，祝您编码顺利~
。。。。。明明可以充分合并的，你却拼命用threadIdx.x做y
（你要知道，线程组合成warp是先threadIdx.x, 再.y，再.z的)
（你这样是故意让一个warp的32个线程间访存拉开巨大步长。。。）
（和你说了还用“算法复杂”拒绝接受。）
建议已经给出。既然你拒绝接受。我以后将不再为你服务。
[
谢谢！
shared memory的方法也已测试过，可以提高一些速度，但仍不够高。
[
解决此问题，在实践中，有3个方法。
在本贴中，我一直想寻求确认：cuda是否有自身特性方法可解决标题所述问题。
如果只是改个下标就能解决问题，我就根本不会来问这个问题。
[
一个col-based的算法改成line-based或者elem-based是足够麻烦的！
[
一个col-based的算法改成line-based或者elem-based是足够麻烦的！
LZ您好：
收到您在您的实现中对shared memory的评价，以及无法确定：
1：是否有其他提高效率的做法；
2：使用shared memory后速度仍不够高是由于算法和硬件极限限制还是由于您的具体实现限制。
因此只能谨慎地对您的评价采取无法肯定也无法否定的态度，敬请谅解。
祝您好运~
LZ您好：
您的标题涉及范围宽广，无法给您一个肯定或者否定的结论。
但就您给出的具体实例而言，符合cuda自身特性的可解决问题的方法恰恰就是您所不屑的交换坐标的方法。这是最直接的方法，也是效率相当不错的解决方案。在不进行其他细致的优化时，仅仅做这样的交换，就可以有10X以上的访存效能提升。
如果您的实例能较好地代表您的实际的复杂问题的行为，那么依然建议您多少参考下这个方法。
如果您的实例并不能代表您的实际问题的行为，那么讨论这个实例未免有南辕北辙之嫌，而直接对没有讨论过的您的实际问题下结论，却又流于空泛。
您说是这样么？
祝您好运~
LZ您好：
就本帖讨论的实例而言，对于第一个kernel，在算法层面是完全没有改变的，依然是基于列上的处理，我们只是就CUDA的具体实现特点做了不同的安排而已。
此外，一般而言，这世上是没有免费的午餐的，code fast和fast code一般总是一对矛盾的。因此在有必要的情况下，修改实现，修改算法，甚至修改算法框架都是可选的手段。
注意到您是手机发帖，颇为感动，来日方长，论坛随时欢迎您来讨论问题的~
祝您好运~
楼主：“如果只是改个下标就能解决问题，我就根本不会来问这个问题。”
我：既然不相信，我将无视您。信任是彼此的。