跨厂商 GPU 可移植性：SYCL 与 HIP 的编程模型与迁移

HIP          与 CUDA 近似同构，迁移成本最低，但基本只覆盖 AMD 与 NVIDIA
SYCL         开放标准，单源 C++，覆盖 Intel、AMD、NVIDIA 与部分 FPGA
Kokkos/RAJA  更高层抽象，用 C++ 模板与执行空间切换后端

warp 宽度      NVIDIA 32，AMD 64，Intel 16 或 32
共享内存大小    每代差异极大，直接影响 tile 策略
原子操作语义    不同实现的内存序与作用域支持不一
硬件特性       张量核、矩阵指令、异步拷贝各有专有扩展

// CUDA
__global__ void saxpy(int n, float a, const float *x, float *y) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) y[i] = a * x[i] + y[i];
}

// HIP：只需把头文件与启动方式换成 HIP 版本
__global__ void saxpy(int n, float a, const float *x, float *y) {
    int i = hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x;
    if (i < n) y[i] = a * x[i] + y[i];
}

#include <hip/hip_runtime.h>

hipLaunchKernelGGL(saxpy, dim3(blocks), dim3(threads), 0, 0, n, a, x, y);
hipDeviceSynchronize();

# 逐文件转换 CUDA 源码到 HIP
hipify-perl saxpy.cu > saxpy.hip.cpp

# 基于 clang 的更精确转换
hipify-clang saxpy.cu -- -I/usr/local/cuda/include

// 危险：写死 32，在 AMD 上 warp 宽度是 64
int lane = threadIdx.x % 32;

// 安全：使用运行时查询
int lane = threadIdx.x % warpSize;

只需要 NVIDIA 与 AMD    HIP 迁移成本最低，hipify 加少量手工修复
需要 Intel 或 FPGA      考虑 SYCL
追求最前沿特性          仍需平台专有代码

#include <sycl/sycl.hpp>
using namespace sycl;

queue q;
buffer<float, 1> buf_x(x, range<1>(n));
buffer<float, 1> buf_y(y, range<1>(n));

q.submit([&](handler &h) {
    auto acc_x = buf_x.get_access<access::mode::read>(h);
    auto acc_y = buf_y.get_access<access::mode::read_write>(h);
    h.parallel_for(range<1>(n), [=](id<1> i) {
        acc_y[i] = a * acc_x[i] + acc_y[i];
    });
});

queue::submit      提交一个命令组
accessor           在命令组中声明对缓冲的访问，依赖由此自动推导
queue::wait        等待队列中所有任务完成
event              细粒度同步句柄

float *x = malloc_shared<float>(n, q);   // 主机与设备均可访问
float *d = malloc_device<float>(n, q);   // 仅设备可访问
q.parallel_for(range<1>(n), [=](id<1> i) { d[i] = x[i] * 2.0f; }).wait();
free(x, q);

性能关键路径      用 malloc_device 加显式 memcpy，行为可预测
复杂指针结构      用 USM，避免访问器与指针混用的复杂性
已有 C 风格代码   USM 改动更小，不必重构成缓冲模型
纯新代码且数据规则 缓冲加访问器更安全，依赖自动推导

// 一维简单并行
q.parallel_for(range<1>(n), [=](id<1> i) { ... });

// 显式工作组
q.parallel_for(nd_range<2>({nx, ny}, {16, 16}), [=](nd_item<2> it) {
    auto gid = it.get_global_id();
    auto lid = it.get_local_id();
    ...
});

// 分层并行
q.parallel_for(nd_range<1>({n}, {64}), [=](nd_item<1> it) { ... });

q.parallel_for(nd_range<1>({n}, {256}), [=](nd_item<1> it) {
    auto sg = it.get_sub_group();
    int width = sg.get_local_range()[0];        // 运行时查询，不写死
    int lane  = sg.get_local_id()[0];
    float v = sycl::shift_group_left(sg, value, 1);
    float total = sycl::reduce_over_group(sg, value, sycl::plus<float>());
});

# Intel oneAPI：为不同后端生成多目标二进制
icpx -fsycl -fsycl-targets=spir64,spir64_gen -O3 -o app app.cpp

# AdaptiveCpp：为 NVIDIA 与 AMD 同时编译
acpp -O3 --acpp-targets='cuda:sm_80;hip:gfx90a' -o app app.cpp

工作组大小   必须是 sub-group 宽度的整数倍，且随平台调整
sub-group 大小 由实现决定，可通过 kernel 属性请求
向量宽度     用 sycl::vec 显式向量化，或依赖编译器自动向量化

// 查询设备能力
auto dev = q.get_device();
auto max_wg = dev.get_info<info::device::max_work_group_size>();
auto sg_sizes = dev.get_info<info::device::sub_group_sizes>();

// 选择 256 与上限的较小值，且保证是 sub-group 宽度的整数倍
size_t wg = std::min<size_t>(256, max_wg);

using Vec4 = sycl::vec<float, 4>;
q.parallel_for(range<1>(n / 4), [=](id<1> i) {
    Vec4 a = in_vec[i];
    Vec4 b = in_vec2[i];
    out_vec[i] = a * b + Vec4(1.0f);
});

q.parallel_for(nd_range<2>({N, N}, {16, 16}), [=](nd_item<2> it) {
    local_accessor<float, 2> tile({16, 16}, it.get_handler());
    tile[it.get_local_id()] = A[it.get_global_id()];
    group_barrier(it.get_group());   // 之后从 tile 读取，减少全局访存
});

CUDA → HIP      hipify 自动转换加少量手工修复，成本最低
CUDA → SYCL     dpct 工具辅助，需要重写启动与内存管理
原生新代码      直接写 SYCL，用 USM 降低学习成本

□ 消除所有写死的 warp 宽度与工作组大小
□ 用组操作替换 __shfl 与 __syncthreads
□ 检查原子操作的内存序与作用域语义
□ 检查依赖平台专有特性的代码（纹理、动态并行）
□ 为每个目标平台单独测性能，不假设自动最优
□ 建立跨平台的数值回归测试

对每个目标平台记录：该内核的峰值占比（如 DRAM 带宽利用率）
可接受标准：所有平台都达到各自峰值的 70% 以上

工具链        固定各平台的编译器与驱动版本矩阵，写进 CI
多目标构建    一份二进制覆盖多个后端，降低运维成本
性能门槛      设成相对该平台峰值的比例，而非绝对时间