OpenCL内存合并（Coalescence）疑问：为何示例Kernel函数合并性差？

咱们先从你贴的代码里的明显问题说起：你重复声明了__global float* input参数，应该是一个input、一个output吧？先把这个语法错误修正，不然代码根本跑不起来。假设修正后的代码大概是这样：

__kernel void mapSin( __global float* input, __global float* output, const unsigned int count) {
    int x = get_global_id(0);
    if(x < count) {
        output[x] = sin(input[x]);
    }
}

如果是这个版本，其实内存合并性是不错的，但你说它合并性差，那大概率是你的实际代码里有跨步访问或者不连续的内存寻址逻辑，我来给你拆解核心原因：

内存合并的核心逻辑先回顾下 #

GPU的内存控制器是按「内存事务」干活的——一次会读取连续的一大块内存（比如64/128字节，具体看GPU架构）。如果同一个工作组里的线程，访问的是连续的内存地址（比如线程0读input[0]，线程1读input[1]，...，线程63读input[63]），这些零散的访问会被合并成少量甚至单个事务，带宽利用率拉满。但如果线程访问的地址是跳着来的，控制器就得发起N多个独立事务，完全浪费带宽，这就是合并性差的本质。

你的函数大概率踩了这些坑 #

• 跨步访问内存
  比如你的线程ID不是直接对应数组下标，而是做了乘法偏移：

  // 比如每个线程跳过3个元素，访问0、4、8...这种地址
  int idx = get_global_id(0) * 4;
  

  这种情况下，同一工作组的线程访问的地址间隔很大，完全没法合并成连续事务，每一个线程的访问都要单独触发一次读取，效率极低。
• 二维维度下的存储顺序不匹配
  如果你的代码是处理二维数组，比如数组是行主序存储（C语言默认），但你却让线程按列的顺序访问：

  int x = get_global_id(0);
  int y = get_global_id(1);
  // 行主序下，列访问的跨步是整个宽度，地址完全不连续
  int idx = y + x * width;
  

  比如width是1024，那同一x的线程访问的地址差了1024个float，这种跨步完全破坏了内存合并。
• 全局ID的维度用错了
  比如你启动的是1D全局工作项，但代码里却用了get_global_id(1)来做下标，这会导致大量线程访问同一个地址或者完全离散的地址，直接废掉内存合并。

给你个内存合并优秀的正面例子 #

__kernel void mapGoodSin( __global float* input, __global float* output, const unsigned int count) {
    // 用全局ID直接作为数组下标，确保连续访问
    int global_idx = get_global_id(0);
    if(global_idx < count) {
        output[global_idx] = sin(input[global_idx]);
    }
}

这个版本里，一个工作组的64个线程会连续访问input[0]到input[63]，内存控制器一次就能把这些数据读进来，完美合并，带宽利用率拉满。

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