GPU编程PPT4

SYCL 编程模型笔记（COMP52315 GPU Programming）

📌 课程信息

• 课程名称：GPU Programming
• 讲师：Christopher Marcotte
• 邮箱：christopher.marcotte@durham.ac.uk

---

🎯 SYCL 目标

• 用于异构系统（CPU/GPU/FPGA）的 高层次单源编程模型
• 使用现代 ISO C++ 编写
• 三大扩展：
  1. 设备发现与信息获取
  2. 设备控制（计算内核 + 内存管理）
  3. 依赖管理与调度

---

📖 SYCL 规范

• 由 Khronos Group 维护
• 当前版本：SYCL 2020
• SYCL 规范 PDF
• SYCL 语法速查表
• C++ Execution Library (C++26)

---

🧩 SYCL 实现

常用编译器：

• clang++（LLVM）
• dpcpp / icpx（Intel）

在 NCC 上使用方法：

module load llvm-clang
module load cuda/11.5

clang++ -fsycl -fsycl-targets=nvptx64-nvidia-cuda my_code.cpp -o my_executable

CUDA版本要求 ≤ 11.5

---

🔍 设备发现 & 信息查询

auto devices = sycl::device::get_devices();
for (auto d : devices)
    std::cout << d.get_info<sycl::info::device::name>() << "
";

快速查看当前设备：

sycl::queue q;
std::cout << "Device: " 
          << q.get_device().get_info<sycl::info::device::name>() << "
";

指定设备（如 cpu, gpu）：

sycl::gpu_selector gpu;
sycl::queue q(gpu);

---

🧠 SYCL 程序结构

Lambda 表达式

[capture](params) { body }

• [=]：默认值捕获（by copy）
• [&]：引用捕获（by reference）

---

SYCL 示例程序

#include <sycl/sycl.hpp>
sycl::queue q;
std::vector<double> x(N, 1.0);

{
    sycl::buffer<double, 1> x_buf(x.data(), x.size());
    q.submit([&](sycl::handler& cgh) {
        auto x_d = x_buf.get_access<sycl::access::mode::read_write>(cgh);
        cgh.single_task([=]() {
            for (int i = 0; i < N; i++)
                x_d[i] = 2.0;
        });
    });
}
for (auto val : x) std::cout << val << " ";

---

🚀 SYCL 执行模型

• 源于 OpenCL
• ND-range：由 work-group 组成，work-group 包含多个 work-item
• 可为 1D、2D、3D 空间定义：
  • Global range: 总 work-items 数量
  • Local range: 每组 work-items 的数量

---

🧵 内核函数（Kernel）API

Single Task

q.submit([&](handler& cgh){
    cgh.single_task([=]() {
        /* device code */
    });
}).wait();

Parallel For（高层版本）

cgh.parallel_for(sycl::range<1>(N), [=](sycl::id<1> i) {
    x_d[i] += 1;
});

ND-Range 版本

sycl::range global{2, 2, 2};
sycl::range local{2, 2, 2};

q.submit([&](sycl::handler& cgh) {
    cgh.parallel_for(sycl::nd_range{global, local},
        [=](sycl::nd_item<3> i) {
            printf("My global id in z is %lu\n", i.get_global_id(2));
        });
}).wait();

---

⚠️ SYCL 内核限制

在内核中 不允许：

• 动态内存分配
• 递归
• 多态/虚函数
• 函数指针
• std:: 标准库（如 std::vector, std::string 等）

---

🧠 内存管理方式

Buffer-Accessor 模型

• 使用 sycl::buffer 封装数据
• 在 queue.submit() 内用 accessor 访问
• buffer 生命周期结束时会自动同步数据

Unified Shared Memory (USM)

int* data = malloc_shared<int>(N, q);
q.parallel_for(sycl::range<1>(N), [=](sycl::id<1> i) {
    data[i] *= 2;
}).wait();

---

📚 更多学习资源

• SYCL API 文档: SYCL Doxygen
• CUDA 到 SYCL 迁移指南:
  Codeplay Guide
• SYCL Academy 教程:
  GitHub
• DPC++ 互动教学平台：
  CodinGame Playground
• 免费书籍：
  Mastering DPC++ (Springer)