如何在OpenCL内核中对全局内存数组单个元素加锁？

问题描述 #

在OpenCL中，需要实现工作项对全局内存中unsigned short数组的单个元素独占访问：当某个工作项操作数组的某一元素时，其他工作项无法对该元素进行任何操作，但不限制数组其他元素的访问。简化示例代码如下：

__kernel void calculateValues(
    __global unsigned short* values)
{
    // 前面的一些代码...获取i...
    // 已锁定values[i]，但values[i+1]及其他元素未被锁定
    unsigned short current_value = values[i];
    // 对current_value执行一些操作...
    values[i] += current_value;
    // 已解锁values[i]
}

编辑后，更贴近实际场景的内核代码如下：

__kernel void
setNodeStableConnectionWeight(
    volatile __global       ushort* node_signal_transmission_weights,
             __global const ushort* node_signal_transmission_weight_limit,
             __global const ushort* node_signal_transmission_noise_weight,
             __global const uint  * node_signal_receiver_ids)
{
    const uint
        node_id                                                        = get_global_id(0U);
    const ushort
        node_signal_transmission_sender_side_weight_limit              = node_signal_transmission_weight_limit[node_id];
    const ushort
        node_signal_transmission_noise_weight_local_copy               = *node_signal_transmission_noise_weight;
          uint
        node_signal_receiver_id                                        = node_signal_receiver_ids[node_id];
          ushort
        node_signal_transmission_receiver_side_weight;
           bool
        node_connection_correlation_exists                             = true;
    while (node_connection_correlation_exists) {

        // 锁定node_signal_transmission_weights[node_signal_receiver_id]的读写，禁止其他工作项操作！若已被其他工作项锁定则等待。
        node_signal_transmission_receiver_side_weight                  = node_signal_transmission_weights[node_signal_receiver_id];
        if (node_signal_transmission_sender_side_weight_limit          < node_signal_transmission_receiver_side_weight) {
            node_signal_transmission_receiver_side_weight             -= node_signal_transmission_sender_side_weight_limit;
            node_connection_correlation_exists                         = doesNodeConnectionCorrelationExist(
                node_signal_transmission_receiver_side_weight,
                node_signal_transmission_sender_side_weight_limit,
                2U);
            if (node_connection_correlation_exists) {
                if (node_signal_transmission_receiver_side_weight      < node_signal_transmission_noise_weight_local_copy)
                    node_signal_transmission_receiver_side_weight++;
                else if (node_signal_transmission_receiver_side_weight > node_signal_transmission_noise_weight_local_copy)
                    node_signal_transmission_receiver_side_weight--;
            }
            else
                node_connection_correlation_exists                     = true;
        }
        else
            node_signal_transmission_receiver_side_weight             += node_signal_transmission_noise_weight_local_copy;
        node_signal_transmission_weights[node_signal_receiver_id]      = node_signal_transmission_receiver_side_weight;

        // 解锁node_signal_transmission_weights[node_signal_receiver_id]的读写。
        node_signal_receiver_id                                        = node_signal_transmission_receiver_ids[node_signal_receiver_id];
        if (node_id                                                   == node_signal_receiver_id)
            node_connection_correlation_exists                         = false;
    }
}

解决方案：使用原子操作实现单元素独占访问 #

你需要利用OpenCL的原子函数实现"读取-修改-写入"的原子性，确保同一内存地址的操作互斥，不会被其他工作项打断。

核心原理 #

普通内存读写是分离的指令，会导致竞争条件；原子函数将整个读写流程封装为不可分割的硬件指令，保证同一时间只有一个工作项能操作目标元素。针对ushort类型，推荐使用**原子比较交换（Atomic Compare-and-Swap, CAS）**来适配复杂逻辑。

1. 简化示例的实现

如果只是简单的加法操作，可以直接使用原子加法；如果需要自定义逻辑，用CAS循环重试：

__kernel void calculateValues(
    __global unsigned short* values)
{
    // 获取i的逻辑...
    ushort old_val, new_val;
    do {
        // 读取当前值
        old_val = values[i];
        // 执行自定义操作
        new_val = old_val + old_val; // 对应原代码的values[i] += current_value
        // 原子交换：若内存值仍为old_val，则替换为new_val，否则重试
    } while (!atomic_cmpxchg(&values[i], old_val, new_val));
}

2. 实际内核代码的修改

针对你复杂的分支逻辑，用CAS循环包裹整个读写流程，确保操作的原子性：

__kernel void
setNodeStableConnectionWeight(
    volatile __global       ushort* node_signal_transmission_weights,
             __global const ushort* node_signal_transmission_weight_limit,
             __global const ushort* node_signal_transmission_noise_weight,
             __global const uint  * node_signal_receiver_ids)
{
    const uint
        node_id                                                        = get_global_id(0U);
    const ushort
        node_signal_transmission_sender_side_weight_limit              = node_signal_transmission_weight_limit[node_id];
    const ushort
        node_signal_transmission_noise_weight_local_copy               = *node_signal_transmission_noise_weight;
          uint
        node_signal_receiver_id                                        = node_signal_receiver_ids[node_id];
           bool
        node_connection_correlation_exists                             = true;
    while (node_connection_correlation_exists) {
        ushort old_weight, new_weight;
        // CAS循环实现独占访问
        do {
            old_weight = node_signal_transmission_weights[node_signal_receiver_id];
            new_weight = old_weight;
            
            // 原有业务逻辑，基于old_weight计算new_weight
            if (node_signal_transmission_sender_side_weight_limit < new_weight) {
                new_weight -= node_signal_transmission_sender_side_weight_limit;
                bool correlation = doesNodeConnectionCorrelationExist(
                    new_weight,
                    node_signal_transmission_sender_side_weight_limit,
                    2U);
                if (correlation) {
                    if (new_weight < node_signal_transmission_noise_weight_local_copy)
                        new_weight++;
                    else if (new_weight > node_signal_transmission_noise_weight_local_copy)
                        new_weight--;
                } else {
                    // 注意：原逻辑此处存在矛盾，需根据业务需求修正
                    node_connection_correlation_exists = true;
                }
            } else {
                new_weight += node_signal_transmission_noise_weight_local_copy;
            }
            
            // 尝试原子更新：若内存值未被修改则替换，否则重试
        } while (atomic_cmpxchg(&node_signal_transmission_weights[node_signal_receiver_id], old_weight, new_weight) != old_weight);

        // 循环迭代逻辑
        node_signal_receiver_id = node_signal_transmission_receiver_ids[node_signal_receiver_id];
        if (node_id == node_signal_transmission_receiver_id)
            node_connection_correlation_exists = false;
    }
}

注意事项 #

• 原子函数兼容性：确保OpenCL设备支持atomic_cmpxchg针对ushort的版本（OpenCL 1.2+通常支持，部分设备需检查扩展）。若不支持，可将ushort包装为uint后使用对应原子操作。
• 性能优化：原子操作会带来性能开销，工作项可能因重试等待。尽量减少同一元素的竞争频率，比如优化node_signal_receiver_id的分配逻辑。
• 逻辑修正：原代码中else分支强制将node_connection_correlation_exists设为true的逻辑存在矛盾，需根据实际业务需求调整。

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