我的数据量已经很大，一次性传到显存里面，处理后一次性传回主机，数据传输的延迟已经完全覆盖。
算法等价的情况下，还有什么原因可能造成GPU比CPU还慢。
我有一些理解不知道对不对，请指出：
1、如果我的核函数里充满了BANK CONFLICT,在最坏的情况下，造成所有的数据都排队串行执行，这相当于CPU运行，运行的时间应与CPU相当。
2、GPU每一个核都相当于一个独立的CPU,我现在把每一个核都当成一个独立的CPU来使用，也就是每个线程都独立跑一个程序副本，线程之间无合作。假如我现在有384个线程，我知道GPU运行时间不可能达到CPU的384倍，但至少不会比CPU慢。
3、GPU存储资源使用率很高甚至是满载的时候，与资源使用率非常低的时候，会不会对GPU运行速度造成很大影响。
LZ您好：
您的1#中充斥着理解错误，下面将一一为您指出：
0：“我的数据量已经很大，一次性传到显存里面，处理后一次性传回主机，数据传输的延迟已经完全覆盖。”——数据量大小暂且不论，一次性copy，计算一次，然后整体回拷，这本身是完全没有用计算来掩盖传输的做法，不知LZ的“数据传输的延迟已经完全覆盖”这种说法从何而来？
事实上，如果需要整体数据搬移，然后相应的计算量很少，这并不适合在GPU上计算，有可能CPU直接算都比数据copy的时间短。
如果copy到device端的数据是被反复利用的，或者有相对密集的复杂的计算，那么可以发挥出GPU计算吞吐量大的优势。
以及，您可以拆分到多个stream里面，用计算和数据传输相互掩盖，这样能带来一些好处。
以上是对您1#第一句话自我下的结论中错误之处进行的说明。
未完待续。
0.5：至于“还有什么原因可能造成GPU比CPU还慢”，实际上造成GPU比CPU慢的原因非常广泛，而且GPU比CPU慢的情况也很常见。
首先必须强调的是，任何硬件都有其适合的用途，也有其不适合的用途；任何硬件都有适合其自身特点的用法和不适合其自身特点的用法。从来就没什么保证GPU一定比CPU快，如果真的能这样，那么CPU久不用生产了，intel，AMD等一票CPU厂家立即倒闭去卖烤肉就行了。
常见的导致GPU计算很慢的情况包括但不限于以下几点：
a：算法不适合并行计算，前后有很强的依赖性，而每一步又无法并行处理。
b：需要频繁和host端交换大量数据，此时时间基本上都消耗在pci-e总线传输上了。
c：不合适的算法实现，写的不好的算法会极大地降低GPU的计算效能。
d：不合适的线程组织，也将极大地降低GPU的计算效能。
所以，需要选择合适的问题和算法，辅以精心撰写的适合GPU的代码，才能发挥出GPU的效能。
但即便如此，也不会随便几百上千倍的加速。
待续——
以下依次回答您1#的三个问题：
1：“如果我的核函数里充满了BANK CONFLICT,在最坏的情况下，造成所有的数据都排队串行执行，这相当于CPU运行，运行的时间应与CPU相当。”
首先指出，“bank conflict”是使用shared memory的时候，可能出现的情况。如果不使用shared memory，将不会出现bank conflict，以及也无法理解“核函数里充满bank conflict”是什么情况。
其次，kernel中的访存不仅有shared memory的访问，还有其他的访存路径，因此，bank conflict并不会造成所有的访存都串行化。
再次，有bank conflict或者访存串行化，这并不“相当于CPU执行”，我无法理解是如何“相当”的。CPU执行特点和GPU执行特点差别甚远，你并不能说GPU在某种特定情况下就和CPU一样。
以及最后，完全没有理由表明“运行的时间应与CPU相当”，GPU的执行时间，就某些具体问题而言，可能比CPU快，可能慢，也可能相当。但是您之前给出的因果理由的逻辑，是完全没有依据的。
总之，您这样简单地推断，是不可取的。
2：
首先“GPU每一个核都相当于一个独立的CPU”，这个说法是不正确的。在硬件结构上，两者相差很远，CUDA CORE简单得多，硬件规模也小得多，功能上也远远没有CPU完整，多个cuda core配合其他一些硬件在一起组成一个SM，才是相对完整的处理结构，此外还要配合GPU全局的L2和显存控制器，pci-e控制器等，才能组成完整的GPU硬件。
或许从软件层面上讲，将每个线程都看做是运行在一个完整功能的微小的CPU上，这是一种可取的理解，这也正是SIMT比SIMD灵活的地方，但是您的这种“GPU每一个核都相当于一个独立的CPU”观点，是不正确的。
其次，“我现在把每一个核都当成一个独立的CPU来使用，也就是每个线程都独立跑一个程序副本，线程之间无合作。”
a)前面说了一个CUDA CORE不能当做一个CPU看待；
b)您这里成功地混淆了CUDA CORE这个硬件概念和线程这个软件概念；
c)从线程的角度讲，每个线程都会跑相同的kernel代码，如果没有block内的线程合作，那么他们确实可以看作是相互独立的。但是，如果直接拿CPU的代码实现来跑，实现中的一些访存方式是极为不适合GPU的，将极大地影响GPU的效率。或者如果里面有大量的随机的跳转，与warp非对齐的分支，这些都会影响GPU的效率。
最后，“假如我现在有384个线程，我知道GPU运行时间不可能达到CPU的384倍，但至少不会比CPU慢。”
a)您这里依然混淆了线程和CUDA CORE的概念，简单地说，CUDA CORE的数量（以及硬件的频率等参数）决定了GPU硬件本身能达到的峰值计算速度，而要达到这个速度或者要达到一个较好的执行速度，需要比cuda core数量多很多的线程数量，以及需要良好优化的代码。
b)即便您这里的“线程”换成CUDA CORE，这个结论也依然是错误的，因为前面已经说明，无法用一个CUDA CORE去和一个CPU核心相互比较的，CPU核心是高频率，低延迟，多个运算单元，完整功能的一个庞大的核心，而CUDA CORE仅仅是十分简单的运算单元而已，频率也低得多。
或者我们把这个说法简单推广一下，就很容易看出来问题。
这个说法等效于，即使我的GPU只有一个核心（或者只跑一个线程），也和一个CPU核心一样快，这显然是荒谬的。一枚中高档的CPU的晶体管规模和一枚中档的GPU的晶体管规模相当，CPU只有4个核心，GPU动辄好几百个核心，如果还能1个CPU核心对一个GPU核心效能相当，那CPU的实现效率就低到一定境界了。
综上，您提供的各项前提条件实际上都不能支持您最后的结论“但至少不会比CPU慢”。
顺便说一下，“GPU运行时间不可能达到CPU的384倍”这句话，按照您的意思以及正确的中文习惯表达应为“GPU运行时间不可能达到CPU的384分之一（那么低）”
3：“GPU存储资源使用率很高甚至是满载的时候，与资源使用率非常低的时候，会不会对GPU运行速度造成很大影响。”——访存对GPU计算效能的影响，只取决于你使用的那一部分的使用情况，比如是否合并访问等等。而与您不使用的那些部分无关。
只要您cudaMalloc()（或者其他方法）申请成功的显存，都可以正常使用，是无差别的，不会因为这块存储空间实际在某个DRAM的某个位置，而与其他地方不同。
请您主要考虑修正您之前估计问题时的错误观点，以及主要考虑优化您的算法实现（当然前提是您的算法是适合GPU实现的），而不要疑神疑鬼怀疑多用一点DRAM空间就会突然变得非常慢。
您的1#的所有问题回答如上，供您参考。
祝您编码顺利~
谢谢您这么耐心为我讲解。
我曾经试过“将每个CUDA CORE当成一个独立CPU来使用”，发现确实有加速作用，但是加速比很低，1～2的样子。
现在依然有一个核函数是“将每个CUDA CORE当成一个独立CPU来使用”，因为我暂时还没有找出很好并行化的方法，但是又不想将一堆数据传回CPU处理再传回显存。
从您上面的分析看来，“将每个CUDA CORE当成一个独立CPU来使用”是非常低效的，我要寻找新的并行实现算法。
LZ您好：
1：前面5#已经具体说过，一个CUDA CORE不能看做是一个CPU的核心，以及一个线程也并不是一直在一个CUDA CORE上执行，以及一个线程在执行kernel里面的多个指令的时候，也并不都在CUDA CORE中执行，比如访存，比如使用SFU的指令等。
2：前面说过您这个只是“将一个线程看做是在一个完整功能的小型CPU上运作”，但必须再次指出的是，这个“完整功能”是整个GPU提供的，这里的“小型CPU”是一个抽象出来的概念，而不是指CUDA CORE。
3：同时，“将一个线程看做是在一个完整功能的小型CPU上运作”并不是说有多少CUDA CORE就上多少线程，实际上一般应该上多得多的线程，才能比较好的发挥出GPU的性能（也有用较少线程，但是精心安排任务以达到较好效能的例子）。
综合上述三点，您“将每个CUDA CORE当成一个独立CPU来使用”（并暗示有多少CUDA CORE就上多少线程），这是一种不适合GPU特性的做法，一般无法发挥出GPU的效能。
以及，在这里，我不评价您的“加速比”的情况。因为加速比可以有宽泛的不同的内涵。
比如算法加速比，A算法在改进之后比B算法快了多少。
比如并行加速比，C算法中的某部分经过并行化之后快了多少。
比如设备对比的加速比，使用CPU实现某算法与使用GPU实现某算法相比，前者或者后者快了多少。以及设备加速比一则取决于设备自身的能力，二则也受算法实现优劣的影响。
以及您后面的内容：
1：“现在依然有一个核函数是“将每个CUDA CORE当成一个独立CPU来使用”，因为我暂时还没有找出很好并行化的方法，但是又不想将一堆数据传回CPU处理再传回显存。”
是否将数据传回去CPU处理这是一个问题，而将数据留在GPU上如何处理，这是另外的一个问题。留在GPU上就地处理，或许还有好的实现方法。
2：“从您上面的分析看来，“将每个CUDA CORE当成一个独立CPU来使用”是非常低效的，我要寻找新的并行实现算法。”
我没有这样说过，实际上我一直在说，a)不能这么类比，硬件结构差别很大，b)一般需要比cuda core数量更多的线程才能实现较高效率，c)从线程的视角看，是可以看做运行在一个全功能的处理器上，但这个处理器不是CUDA CORE。
因为您前提“将每个CUDA CORE当成一个独立CPU来使用”如上面多楼叙述，是不正确的，硬件方面差别很大，软件方面又不是那么回事。所以在这个错误前提下的任何结论和推断都是无意义的，包络您对效率的推断。您使用一种本身很明显效率就很低的方式去实现您的算法，然后说，这个效率果然很低，这有什么实际意义呢？
望您深思~
难道不是一个CORE提供一个线程吗，我一直是这么认为的。。。。。
您说要上“比CUDA CORE多得多的线程才能较好地发挥GPU性能”，这句话好难理解，能再解释一下吗？
明天见，晚安。
LZ您好：
1：“难道不是一个CORE提供一个线程吗，我一直是这么认为的”，不是的，对于NVIDIA 的GPU来说不是这样的，前面已经多次指出了，现在我为您再次确认这一点。
以及，您这个是CPU上的观念，GPU上有所不同。
以及，如果您打算继续使用您之前固有的想法，那么请把一个SM/SMX看做是一个“core”，它具备大量的SIMD的执行单元（即SP，cuda core）以及其他辅助硬件单元；它是一个具备超线程能力的处理器，相较于intel的CPU，每个CPU核心可以有2个线程，一个SM/SMX可以resident 一两千个threads。
2：“您说要上“比CUDA CORE多得多的线程才能较好地发挥GPU性能”，这句话好难理解，能再解释一下吗？”，这个本来就是GPU基本的执行行为特点的。resident在SM上的threads，以warp为单位接收scheduler调度，或者去访存，或者去计算。
当一个warp的threads将当前的指令灌入执行单元之后，在等待结果的这段延迟中，将被切换出来，换另外一个就绪的warp继续执行。所有的执行单元都是流水线化的。
GPU上就是依赖大量线程的切换来交替向执行单元灌入指令，以及掩盖各种延迟的。
一般来说，如果不是精心设计的kernel，都需要有大量的线程通过这种自动切换机制，来保证GPU有较好的执行效率。以及，这种自动的线程切换的方式，也是一种较为方便的使用方式。线程数太少的时候，将不会有足够的线程来掩盖延迟，执行单元将会等待，执行效率会下降。
如果您的某个特定的任务太小，只需要较少的线程，那么请直接这样使用。
而如果您的任务有一定规模，请尽量多上一些线程，起到“自动优化”的效果。
以及，法无定法，当您十分熟悉CUDA编程的时候，也可以安排使用中等数量的线程，并给每个线程安排较多的工作，这样或许能有其他进一步的改善，或许什么改善也没有。
大致如此，感谢您深夜来访。
祝您晚安~
学习了！
版主，牛叉。
楼上两位网友请注意，本版禁水，请勿使用与讨论内容不相关的回复顶贴，否则将视为灌水处理。
良好的讨论环境需要大家共同维护。
上次讨论过一个小波变换程序，我对程序进行了优化，从第一次写出来时的9秒加速到现在的1秒。现存几点疑问。
优化描述：
算法中存在线程活动数量越来越少的情况，从第一次循环使用64个线程开始，活动线程数逐渐减少，到第六次循环只使用了2个线程。这意味着每个block安排的SP越多，后面造成的线程浪费就会越严重。我起初每个BLOCK有128个SP,程序运行时间9秒（debug模式，下同）。于是我尝试每个block安排32个SP,程序运行用时果然大大减少，为1.6秒。我又得寸进尺，将每个block的sp减小到16，还将一些循环作了展开，代码变得非常冗长。运行时间1秒。
问题就来了：
1、在上述描述中可能依然存在软硬件概念混淆，但您肯定知道我的意思。
2、一个warp为32线程，GPU的计算以warp为单位进行调度。在我准备给每个block中安排16个sp之前，我犹豫了,因为我在论文资料中看到描述说“最好保证每个block的sp数量为warp即32的整数倍”，这是否意味着我把一个warp分散在两个block里面？我的做法是否合理？但实验发现，每个block中安排16个sp依然能使程序进一步加速。
3、虽然程序运行了1秒，但依然没赶上cpu运行时间560ms，从上面的分析我认为，如果我把每个block的sp减为8个，可能还会使gpu运行时间减少一点点，但不可能少于560ms.并且这样做会使的我程序代码变得比较冗长，我不愿意这么做了。GPU赶不上cpu，令我感到非常难过。我一时依然无法确定是算法的并行性不够，还是程序本身写的不合理。
4、我将编绎模式换成release，发现刚才使用了1秒的程序时间减到25ms !这让我很吃惊，DEBUG模式跟release模式的区别有这么大吗？？？我把纯C语言版本的程序也改成release,发现时间依然是560ms左右，没有明显的提升。这是否意味着我的程序已经成功提速，并且加速比为560/25=22.4 ？
最近问题比较多，再次感谢您的耐心和细心，并对您精湛的专业知识表示佩服！
LZ您好：
请先允许我为自己在本帖中回复的数千字以及所花费的时间默哀…
我为您解释问题写了数千字的内容，您还是要坚持自己的固有的错误的概念和想法，我表示，我也不能拿您怎么样。
在您使用正规的概念和术语描述您的问题之前，我拒绝继续再和您详细讨论，同时拒绝通过反猜的方法来主动纠正您的表述，以免和您的主旨不同造成更大的混乱。
下面会简要指出您的问题答案：
0：请保证您在这样删减线程规模的时候，原问题依然能被正确解决，否则讨论时间是无意义的。以及一个block安排如此少的线程数一般而言在任何GPU架构上都是低效的。
1：我肯定不知道您的意思，并拒绝猜测您的意思。
2：warp是不能跨block的。
3：（请允许我不看这一条，真的看不下去了。）
4：请不要使用debug模式测试时间，这个无意义。以及关于C语言在CPU上的编译问题不在本版讨论范围内。以及我不评价您所谓的“加速比”。
最近您的问题比较多，但请您认真查看回帖和相关资料，以正确的方式理解CUDA的相关概念和操作，而不是固守您固有的错误思维，错误概念，基础问题尚未理解清楚，就蛮干并提出各种不靠谱的结论。
古人云：磨刀不误砍柴工，学习来不得半点浮躁，愿共勉。
您每次生气，都会让我更清醒。在些作自我批评；
在知道了CORE与threads并非等同的关系之后，我顿时想起为什么CPU有双核四线程，线程比核多！我明白了为什么plofiler测试说我的mp “mostly idle"，一切都源于我认为线程数量不能超过threads的错误理解，因为我把线程当成了core。我将重点学习基础概念，在我的柴刀没磨好之前，不再向您请教算法与优化问题。
晚安。
LZ您好：
看到您下决心彻底理解并攻克CUDA编程，我感到十分欣慰。
引用主席诗词一句，共勉：“红军不怕远征难，万水千山只等闲。”
（以及，您文中“一切都源于我认为线程数量不能超过threads的错误理解”文中的“threads”应为“cores”才符合上下文。）
祝您晚安~
学习总结，请指正。
1、线程与核心
线程thread是软件概念，CUDA CORE是硬件概念，两者不是等同的关系，但有重要联系：线程要在cuda core上运行，cuda core是支持线程运算的硬件基础。两者也不一定是一一对应的关系，可以是一对多，也可以是多对一。一般情况下，线程数量要比CUDA core多得多的时候才能较好地发挥出GPU的性能。至此，我已抛弃“一个CORE就是一个线程”的错误观点。
2、块和线程束
block是软件概念，多个线程threads可以组成一个块block。SM是硬件概念，SM由sp组成，block要发射到SM上运行。我在16#第（2）小点中认为block由sp（即CUDA CORE)组成是极端错误的，应该受到人民的批判。32个线程组成一个线程束warp，SM以warp为单位对线程进行调度，warp中的32个线程使用SM中的32个sp来执行。因此block中的线程数（而不是sp)最好为32的整数倍，我在16#中在一个block中放入16线程，连一个warp都未能组成，这意味着SM调用该block的线程时将不能取得高效率（但是我想请问这种情况下SM将如何调取这16个线程）。
3、线程切换问题
您在12#中说“当一个warp的threads将当前的指令灌入执行单元之后，在等待结果的这段延迟中，将被切换出来，换另外一个就绪的warp继续执行。所有的执行单元都是流水线化的。”
我理解了大概的意思，就是说要充分利用每个threads的等待时间来做有用功，防止这些处于等待状态的线程占着茅坑不拉屎。
但仔细一想，又产生了疑问：您说的执行单元应该是由sp吧，或者说是由sp组成的SM。当一个warp的threads灌入SM之后，这个warp中的线程就要等待处理结果，此时SM正在忙碌着处理这些数据，那么其它的warp又将自己的线程灌到哪里去呢？您说的“切换出来”是什么意思？
4、线程数与数据量
在常见的CUDA程序例子里，数据量都非常大，远远超过了CUDA CORE的数量。因此每个BLOCK可以安排512，1024甚至更多个线程（受硬件限制），处理过程中“每个线程对应处理一个元素”，用多个BLOCK的时候这些线程可以将所有的数据一次性覆盖。线程数与数据量成1：1的关系。但是如果我的数量很小，还不如CUDA core多。根据您12#的回复，“如果您的某个特定的任务太小，只需要较少的线程，那么请直接这样使用。”这就是说这种情况下无法发挥任何“线程切换”带来的好处。
5、BLOCK的切换
正如上面第4点所说，我的数据量很小，例如128个元素的序列，但CUDA CORE有384个。按照“一个线程对应处理一个数据”，我最多也只能在一个block中安排128个线程，无法发挥“线程切换”的自动优化效能。但是我现在要对大量的序列进行处理，每个序列都是128个元素，每个block处理一个序列。由于block的数量很多，他们在被投射到SM上执行时应该也是“流水线化”的，这就会有一个“切换”作用。“block的切换”这个说法是否准确？既然BLOCK由线程组成，那么它与线程切换是不是一回事？当一个block里的线程数量多于core的数量，在把这个BLOCK投射到SM的时候，就会产生“线程切换”。当一个BLOCK时的线程数量少于或等于core的数量，把这个block投身到SM的时候，就无需切换，只需占着茅坑不拉屎的等待结果。当一个SM能同时处理多个BLOCK时，在BLOCK数量很大的情况下，就会产生“BLOCK的切换”，从而避免占着茅坑不拉屎的现象。
6、debug与release
我之前一直用debug进行程序调试并测时，发现cuda运行时间无法超越c语言程序。根据您17#提示“请不要使用debug模式测试时间，这个无意义。”我才知道，虽然在debug模式下，在一个block下分配16线程所用时间最短，但这个时间毫无意义。我切换到release或profile模式，并测时，发现每个block里分配16线程实际上是延长了运行时间。
请指正！