为什么CUDA向量化内存访问内核的性能比朴素实现更慢？

• 线程网格配置不匹配，产生大量无用调度开销
  你调用两个内核时使用了完全相同的grid配置，该配置是按照朴素内核需要处理的全量size个元素计算的。但向量化内核每个线程处理2个int元素，仅需要处理size/2个逻辑元素，当前配置下近一半线程进入循环后会直接触发i < size/2不成立的逻辑直接退出，额外的线程启动、调度开销抵消了向量化访存的收益。
• 编译器自动优化已经实现了朴素内核的向量化
  NVCC在默认优化等级下会自动对规则的连续全局内存访问做向量化合并，你的朴素拷贝内核访存模式非常简单规整，编译器已经自动将其优化为向量访存指令，手动强转int2做向量化没有带来额外的访存效率提升。
• 向量化访存的事务开销更高
  朴素拷贝场景中，单个warp的32个线程单次读取32个int，总大小为128字节，刚好匹配GPU全局内存/L1缓存的128字节标准访存事务粒度，仅需要1个访存事务就能完成整个warp的读取请求。而int2向量化拷贝每个warp单次读取32个int2，总大小为256字节，需要拆分为2个128字节的访存事务，额外增加了访存开销。

如果要验证问题，只需要将向量化内核的调用grid调整为适配size/2的尺寸：

dim3 grid_vec((size/2 - 1)/block.x + 1, 1);
vecCopy<<<grid_vec, block, 0, stream>>>(d_in, d_out, size);

调整后即可观测到向量化内核的性能符合预期。

内容的提问来源于stack exchange，提问作者kingwales