Skip to content

sync

产品支持情况

  • Ascend 950PR/Ascend 950DT:支持
  • Atlas A3 训练系列产品/Atlas A3 推理系列产品:不支持
  • Atlas A2 训练系列产品/Atlas A2 推理系列产品:不支持
  • Atlas 200I/500 A2 推理产品:不支持
  • Atlas 推理系列产品AI Core:不支持
  • Atlas 推理系列产品Vector Core:不支持
  • Atlas 训练系列产品:不支持

功能说明

同步当前Grid内所有SIMT线程。所有线程块中的所有线程都执行到该同步点位置后,Grid内线程才能继续执行。

函数原型

C++
void sync() const

参数说明

返回值说明

约束说明

  • 仅支持SIMT编程场景,不支持SIMD与SIMT混合编程场景。
  • 必须保证Grid内所有线程都能执行到同一个sync()调用,否则是未定义行为。
  • sync()接口仅保证线程指令执行的同步,不保证各线程块的Data Cache和Global Memory间的缓存一致性

调用示例

  • SIMT编程场景:

    C++
    using namespace cooperative_groups;
    __global__ void simt_kernel(uint64_t* output)
    {
        grid_group g = this_grid(); // 获取grid_group对象
        unsigned int block = static_cast<unsigned int>(g.block_rank());
    
        if (threadIdx.x == 0) {
            output[block] = block + 1;
        }
    
        g.sync(); // 同步Grid内所有线程
    
        if (threadIdx.x == 0) {
            uint64_t sum = 0;
            asc_dcci_entire(output); // 刷新Cache
            for (unsigned int i = 0; i < g.num_blocks(); ++i) {
                sum += output[i];
            }
            output[g.num_blocks() + block] = sum;
        }
    }
    

免责声明:本站内容由 asc-devkit 仓 master 分支自动编译生成,属于持续开发版本,可能存在缺陷,仅供预览与参考。如需稳定及商用资料,请查阅官方 昇腾社区