从 CUDA 到 SYCL 2020 通过 DPC++ 进行光线追踪
5.00/5 (4投票s)
如何将并行 C++ 光线追踪代码从 CUDA 转换为 SYCL 2020（通过 Intel® DPC++）
- 下载 simple_raytracing-master.zip - 36.9 KB（使用 C++17 parallel for_each 的光线追踪）
- 下载 simple_raytracing_cuda-main.zip - 32.6 KB（需要 NVIDIA CUDA SDK）
- 下载 simple_raytracing_sycl2020-main.zip - 32.9 KB（需要 Intel DPC++/oneAPI SDK）
目录
引言
在本篇文章中，我们将移植 《Weekend 的光线追踪》（已转换为 CUDA 的并行代码），使 CUDA 在 CPU 上运行，然后通过 Intel® Data-Parallel C++ (DPC++) 工具包再次移植到 SYCL 2020。DPC++ 是 ISO C++ 加上 Khronos SYCL，并包含最终将纳入最终标准的社区扩展。在撰写本文时，SYCL 2020 规范已发布供公众审阅，作为临时规范，以便开发者在最终版本发布和批准之前提供宝贵的反馈。
以下是从串行代码到使用 C++17 parallel for_each 算法的多线程代码，对 《Weekend 的光线追踪》 进行的一系列改进。
- 将所有原始内存分配转换为
shared_ptr，以防止内存泄漏。 - 用 C++11 随机数生成器（每个线程拥有其副本）替换不可重入的
rand()，使用thread_local来避免共享和锁定。 - 使用单头文件 stb_image_write.h 库，将纯文本 PPM 图像格式替换为 Portable Network Graphics (PNG) 格式。
- 使用 C++17 parallel
for_each进行并行化，在我的 6 核 CPU 上实现了超过 5 倍的性能提升。
本文分为两个主要部分，第一部分讨论纯 C++ 到 CUDA 的转换，第二部分涵盖 CUDA 到 SYCL 转换所需的工作。对于 SYCL 代码，使用 Intel 的 Data-Parallel C++ 编译器来编译代码。在本文中，**host** 和 **device** 指的是 CPU 和 CPU 以外的设备，例如 GPU 或 FPGA。在 CUDA 的情况下，device 仅指 GPU，因为 CUDA 只能在 GPU 上运行，特别是 NVIDIA 的 GPU。SYCL 代码可以选择在 CPU、GPU 或 FPGA 上执行，通过其选择器。
转换为 CUDA
在 CUDA 和 SYCL 中，缓冲区中的数据结构必须是 C++ 可平凡复制的，这意味着对象可以安全地逐字节复制，而无需调用复制构造函数。此要求排除了多态性的使用。在原始 《Weekend 的光线追踪》 代码中，存在一个纯虚函数 material 类和 3 个派生 material 类，它们的大小不同。我们将它们合并到一个非虚拟类中，以获得统一的大小，从而使 material 数组在 CUDA 端易于分配。另一个原因是，虚拟类对象的内部有一个看不见的 vptr 成员。简而言之，vptr 是一个指向 vtbl 的指针，在多态环境中会被调用。由于 vptr 和 vtbl 未经文档说明，并且实现方式不透明，并且很可能因平台而异。最好不要依赖我们的 CUDA 光线追踪器中的虚函数，因为 material 对象在 host 上实例化，然后传递到 device 端。
class material {
public:
virtual bool scatter(const ray& r_in, const hit_record& rec,
vec3& attenuation, ray& scattered) const = 0;
};
class lambertian : public material {
public:
lambertian(const vec3& a);
virtual bool scatter(const ray& r_in, const hit_record& rec,
vec3& attenuation, ray& scattered) const;
vec3 albedo;
};
class metal : public material {
public:
metal(const vec3& a, float f);
virtual bool scatter(const ray& r_in, const hit_record& rec,
vec3& attenuation, ray& scattered) const;
vec3 albedo;
float fuzz;
};
class dielectric : public material {
public:
dielectric(float ri);
virtual bool scatter(const ray& r_in, const hit_record& rec,
vec3& attenuation, ray& scattered) const;
float ref_idx;
};
这是我们合并的 material 类。我们有一个 material_type 成员来跟踪对象的 material 类型，以便调用 scatter()。scatter() 将其工作委托给其中一个 xxx_scatter() 函数。__device__ 是一个 CUDA 指令，表示该函数可在 CUDA device 上调用。另一方面，__host__ 指示该函数可在 CPU host 上调用。__device__ 和 __host__ 并非互斥：换句话说，一个函数可以同时是 host 和 device 可调用的。
enum class material_type
{
lambertian,
metal,
dielectric
};
class material {
public:
material() = default;
material(material_type mtype, const vec3& a, float f, float ri);
__device__ bool scatter(const ray& r_in, const hit_record& rec,
vec3& attenuation, ray& scattered, const RandAccessor& rand) const
{
switch (mat_type)
{
case material_type::lambertian:
return lambertian_scatter(r_in, rec, attenuation, scattered, rand);
case material_type::metal:
return metal_scatter(r_in, rec, attenuation, scattered, rand);
case material_type::dielectric:
return dielectric_scatter(r_in, rec, attenuation, scattered, rand);
}
return lambertian_scatter(r_in, rec, attenuation, scattered, rand);
}
__device__ bool lambertian_scatter(const ray& r_in, const hit_record& rec,
vec3& attenuation, ray& scattered, const RandAccessor& rand) const;
__device__ bool metal_scatter(const ray& r_in, const hit_record& rec,
vec3& attenuation, ray& scattered, const RandAccessor& rand) const;
__device__ bool dielectric_scatter(const ray& r_in, const hit_record& rec,
vec3& attenuation, ray& scattered, const RandAccessor& rand) const;
material_type mat_type;
vec3 albedo;
float fuzz;
float ref_idx;
};
随机数是使用 PreGenerated 类提前生成的，并使用 RandAccessor 类在 device 端访问它们。
class PreGenerated
{
public:
PreGenerated(size_t size)
{
// Will be used to obtain a seed for the
// random number engine
std::random_device rd;
// Standard mersenne_twister_engine seeded with rd()
std::mt19937 gen(rd());
std::uniform_real_distribution<> dis(0.0, 1.0);
vec.resize(size);
for (size_t i = 0; i < size; ++i)
vec[i] = dis(gen);
}
const std::vector<float>& GetVector() const
{
return vec;
}
private:
std::vector<float> vec;
};
class RandAccessor
{
public:
__device__ RandAccessor(size_t offset_,
const float* arr_, size_t size_)
: offset(offset_)
, arr(arr_)
, size(size_)
, index(0)
{}
__device__ float Get() const
{
size_t i = (offset + index) % size;
++index;
return arr[i];
}
private:
size_t offset;
const float* arr;
size_t size;
mutable size_t index;
};
接下来，我们将尝试编译代码。CUDA 编译器 nvcc.exe 的伪参数非常简单，如下所示。CUDA 源文件必须具有 *.cu 文件扩展名。
nvcc -o executable source.cu
在 Windows 上，这还不够。我必须添加一个参数 -ccbin，并指定 Microsoft Visual C++ 编译器 cl.exe 的文件夹路径。**注意**：您的路径可能与我的不同，请相应地修改。
nvcc -o CUDARayTracer.exe CUDARayTracer.cu -ccbin
"C:\Program Files
(x86)\Microsoft Visual Studio\2019\Community\VC\Tools\MSVC\14.28.29333\bin\Hostx64\x64"
有一个警告，提示 raytrace() 的堆栈大小无法静态确定。我不知道这具体意味着什么，所以我就忽略了它，毕竟它只是一个警告，而不是错误。但后来我发现，这个警告预示着我转换的 CUDA 代码注定会失败。
ptxas warning : Stack size for entry function '_Z8raytracePfPyP6sphereS0_PjS0_P6cameraPiS6_S6_'
cannot be statically determined
运行 CUDA 代码后，cudaDeviceSynchronize() 返回错误号 700，该错误号在 此 CUDA 错误页面 中描述为 device 在无效内存地址上遇到加载或存储指令。这会导致进程处于不一致状态，并且任何进一步的 CUDA 工作都将返回相同的错误。要继续使用 CUDA，必须终止并重新启动进程。 简而言之，我的 CUDA 程序崩溃了，而且我不知道原因。
通过一些逻辑的纯粹决心和猜测，我正确地推断出根本原因是递归的 color()，因为它的递归深度可能高达 50，导致堆栈溢出。与 CPU 进程中每个线程默认的 1MB 堆栈大小不同，在 CUDA 线程上运行的 CUDA 内核的堆栈大小是有限的，因此出现了 ptxas 警告。我的 CUDA 版本 11.1 支持递归，但在撰写本文时，DPC++ 不支持在 device 上进行递归函数调用，因此我决定将 color() 转换为迭代版本，名为 color_loop()。
__device__ vec3 color(const ray& r, sphere* world, size_t world_size,
int depth, const RandAccessor& rand)
{
hit_record rec;
if (world_hit(world, world_size, r, 0.001, FLT_MAX, rec)) {
ray scattered;
vec3 attenuation;
if (depth < 50 && rec.mat->scatter(r, rec, attenuation, scattered, rand)) {
return attenuation * color(scattered, world, world_size, depth + 1, rand);
}
else {
return vec3(0, 0, 0);
}
}
else {
vec3 unit_direction = unit_vector(r.direction());
float t = 0.5 * (unit_direction.y() + 1.0);
return (1.0 - t) * vec3(1.0, 1.0, 1.0) + t * vec3(0.5, 0.7, 1.0);
}
}
通过将 color() 转换为下面的迭代版本 color_loop()，我成功地消除了警告，程序也能顺利运行完成。
__device__ vec3 color_loop(const ray& r, sphere* world,
size_t world_size, const RandAccessor& rand)
{
vec3 attenuation_result(1.0,1.0,1.0);
bool assigned = false;
ray temp = r;
for(int i=0; i<50; ++i)
{
hit_record rec;
if (world_hit(world, world_size, temp, 0.001, FLT_MAX, rec)) {
ray scattered;
vec3 attenuation;
if (rec.mat->scatter(temp, rec, attenuation, scattered, rand)) {
temp = scattered;
if(assigned == false)
{
attenuation_result = attenuation;
assigned = true;
}
else
attenuation_result *= attenuation;
}
else
{
attenuation_result = vec3(0, 0, 0);
break;
}
}
else {
vec3 unit_direction = unit_vector(temp.direction());
float t = 0.5 * (unit_direction.y() + 1.0);
attenuation_result *=
(1.0 - t) * vec3(1.0, 1.0, 1.0) + t * vec3(0.5, 0.7, 1.0);
break;
}
}
return attenuation_result;
}
这是 raytrace 函数。__global__ 指令指示 CUDA 编译器该函数是从 host 调用。它将繁重的工作委托给 raytrace_pixel()。
#ifndef NO_CUDA
__global__ void raytrace(float* dev_arr, size_t* dev_arr_size, sphere* dev_sphere,
size_t* dev_sphere_size, unsigned int* dev_pixelsSrc, size_t* dev_pixelsSrc_size,
camera* dev_camera, int* nx, int* ny, int* ns)
{
int x = blockIdx.x * blockDim.x + threadIdx.x;
int y = blockIdx.y * blockDim.y + threadIdx.y;
if (x >= *nx || y >= *ny)
{
return;
}
int thread_index = y * (*nx) + x;
if (thread_index >= *dev_pixelsSrc_size)
{
return;
}
unsigned int& pixel = dev_pixelsSrc[thread_index];
raytrace_pixel(&pixel, dev_arr, dev_arr_size, dev_sphere,
dev_sphere_size, dev_camera, nx, ny, ns);
}
#endif
raytrace_pixel 函数的定义。
__device__ void raytrace_pixel(unsigned int* dev_pixel, float* dev_arr,
size_t* dev_arr_size, sphere* dev_sphere, size_t* dev_sphere_size,
camera* dev_camera, int* nx, int* ny, int* ns)
{
int j = (*dev_pixel) & 0xffff;
int i = ((*dev_pixel) & 0xffff0000) >> 16;
RandAccessor rand(i+j*(*nx), dev_arr, *dev_arr_size);
vec3 col(0, 0, 0);
for (int s = 0; s < *ns; s++) {
float u = float(i + rand.Get()) / float(*nx);
float v = float(j + rand.Get()) / float(*ny);
ray r = dev_camera->get_ray(u, v, rand);
col += color_loop(r, dev_sphere, *dev_sphere_size, rand);
}
col /= float(*ns);
col = vec3(sqrt(col[0]), sqrt(col[1]), sqrt(col[2]));
int ir = int(255.99 * col[0]);
int ig = int(255.99 * col[1]);
int ib = int(255.99 * col[2]);
*dev_pixel = (0xff000000 | (ib << 16) | (ig << 8) | ir);
}
使用 C++17 Parallel for_each 使 CUDA 代码在 CPU 上运行
大量访问网络且受硬盘 I/O 限制的 CPU 代码无法转换为 GPU 代码。理想情况下，CPU 代码应该是计算密集型的，以最大化转换的成功率。因此，并非所有 CPU 软件都能在 GPU 上运行。GPU 代码可以在 CPU 上运行，前提是源代码可用，这已不是什么秘密。不幸的是，NVIDIA 不会开源其任何 CUDA 库。即使开源了，它们也不是用 CUDA 编写的，而是用流汇编（SASS）编写的，这对我们的转换工作几乎没有用处。这就是为什么随机数生成不依赖于 NVIDIA 的库，而是在 CPU 端预先生成，以便为维护 CUDA 和 CPU 的单一代码库铺平道路。纯 C++ 版本在 lambda 中调用 raytrace_pixel，该 lambda 又由 C++17 parallel for_each 调用。
std::for_each(std::execution::par, pixelsSrc.begin(), pixelsSrc.end(),
[&](unsigned int& pixel)
{
size_t dev_arr_size = vec.size();
size_t dev_sphere_size = world.size();
raytrace_pixel(&pixel, vec.data(), &dev_arr_size, world.data(),
&dev_sphere_size, &cam, &nx, &ny, &ns);
});
在纯 C++ 中，代码必须放在 *.cpp 文件中，所以我们创建一个 TestRayTracer.cpp 文件并添加以下 6 行。所有与 CUDA 相关的代码都由 NO_CUDA 的不存在来保护。定义 NO_CUDA 将排除所有 CUDA 代码和头文件。__global__、__host__ 和 __device__ 是纯粹的 CUDA 指令，在纯 C++ 代码中可以定义为空。然后包含 CUDARayTracer.cu。纯 C++ 版本的代码工作已完成。以下是 RayTracing.cpp 包含的所有代码。
#pragma once
#define NO_CUDA
#define __global__
#define __host__
#define __device__
#include "CUDARayTracer.cu"
为了基准测试，纯 C++ CPU 版本在 6 核 12 线程的 Intel® i7 8700 处理器上运行耗时 90 秒，而 CUDA 版本在 Geforce 1060 上提取了 18 倍的性能。在下一节中，我们将深入探讨如何通过 Intel® DPC++ 将 CUDA 代码转换为 SYCL 2020，然后将源代码上传到 Intel® DevCloud 进行编译和执行。
通过 DPC++ 转换为 SYCL 2020
本节将尝试使用 Data-Parallel C++ (DPC++) 将 CUDA 代码转换为 SYCL 2020。DPC++ 是 Intel 对 Khronos Khronos 内 SYCL 工作组定义的 SYCL 2020 规范的实现。Springer 出版社网站上免费提供了一本出色的 DPC++ 电子书。我强烈建议有兴趣了解 SYCL 2020 的人下载并研究它。SYCL 光线追踪代码与 CUDA 的代码相同，只是删除了 __host__ 和 __device__ 等 CUDA 指令，并将 CUDA 函数调用替换为 SYCL 等效项。
本书的代码示例广泛使用了 C++17 的 类模板参数推导 (CTAD) 来实例化模板，并非所有模板参数都必须指定。在 C++17 之前，即使可以从 initializer_list 中的参数类型推导出向量类型，也必须指定它。
std::vector<int> vec{1, 2, 3, 4};
CTAD 放宽了此要求，允许省略模板类型。
std::vector vec{1, 2, 3, 4};
然而，本节中的 SYCL 代码列表并未对数据缓冲区使用 CTAD，以使 C++ 编译器清楚实例化类型。这只是我的个人偏好。使用一个简单的 SYCL 代码来尝试并熟悉 SYCL 语法/编译器，然后再进行代码转换过程。请阅读代码注释以获得更多解释。该代码没有进行有用的计算，只是将 A 数组的每个元素赋给 B 数组。
#include <CL/sycl.hpp>
#include <array>
#include <iostream>
#include <exception>
using namespace sycl;
constexpr int N = 42;
int main(){
// I use array on the stack because the array is not
// huge, only 42 elements. I would advise to use std::vector
// when the numbers of elements is massive or exceed the stack
// size of 1 MB
std::array<int, N> a, b;
// Initialize a and b array.
for(int i=0; i< N; ++i) {
a[i] = i;
b[i] = 0;
}
try {
queue Q(host_selector{}); // Select to run on the CPU
// Initialize the buffer variables for the device with
// host array.
buffer<int, 1> A{a.data(), range<1>(a.size())};
buffer<int, 1> B{b.data(), range<1>(b.size())};
// submit() takes in a lambda which is invoked on the host
Q.submit([&](handler& h) {
// Device will access the arrays through accA and accB
auto accA = A.template get_access<access::mode::read>(h);
auto accB = B.template get_access<access::mode::write>(h);
// parallel_for() takes in a lambda which is kernel that
// is invoked on the device.
h.parallel_for<class nstream>(
range<1>{N},
[=](id<1> i) { accB[i] = accA[i]; });
});
// Wait for the device code to complete
Q.wait();
// Synchronize the device's B array to host's b array by reading it
// If Q.wait() is not called previously, B.get_access() will call
// it behind the scene before the array synchronize.
B.get_access<access::mode::read>();
for(int i=0; i< N; ++i) {
std::cout << b[i] << " ";
}
}
catch(sycl::exception& ex)
{
std::cerr << "SYCL Exception thrown: "
<< ex.what() << std::endl;
}
catch(std::exception& ex)
{
std::cerr << "std Exception thrown: "
<< ex.what() << std::endl;
}
std::cout << "\nDone!\n";
return 0;
}
sycl::queue 构造函数接受一个 selector，该 selector 指示计算设备。可用的 selector 列表供选择。
sycl::queue q(sycl::default_selector{}); // run on the GPU if present and a CPU otherwise
sycl::queue q(sycl::host_selector{}); // run on the CPU without a runtime
// (i.e., no OpenCL)
sycl::queue q(sycl::cpu_selector{}); // run on the CPU with a runtime (e.g., OpenCL)
sycl::queue q(sycl::gpu_selector{}); // run on the GPU
sycl::queue q(sycl::accelerator_selector{}); // run on an FPGA or other acclerator
在下面的 SYCL 代码中，创建了缓冲区来引用 host 数据，以便稍后将其复制到 device。raytrace_pixel() 的主体被放入 h.parallel_for lambda 中。然后调用 Q.wait() 等待工作完成。之后，调用 dev_pixelsSrc.get_access() 将像素数据从 device 同步到 host 缓冲区。
try {
using namespace sycl;
queue Q(host_selector{}); // for running on cpu without OpenCL installed
//queue Q(gpu_selector{}); // for running on gpu
buffer<float, 1> dev_arr{ vec.data(), range<1>(vec.size()) };
buffer<sphere, 1> dev_sphere{ world.data(), range<1>(world.size()) };
buffer<unsigned int, 1> dev_pixelsSrc{ pixelsSrc.data(), range<1>(pixelsSrc.size()) };
buffer<camera, 1> dev_camera{ &cam, range<1>(1) };
Q.submit([&](handler& h) {
auto acc_dev_arr = dev_arr.template get_access<access::mode::read>(h);
auto acc_dev_sphere = dev_sphere.template get_access<access::mode::read>(h);
auto acc_dev_pixelsSrc =
dev_pixelsSrc.template get_access<access::mode::read_write>(h);
auto acc_dev_camera = dev_camera.template get_access<access::mode::read>(h);
h.parallel_for<class nstream>(range<1>{pixelsSrc.size()}, [=](id<1> index) {
unsigned int& pixel = acc_dev_pixelsSrc[index];
unsigned int* dev_pixel = &pixel;
int j = (*dev_pixel) & 0xffff;
int i = ((*dev_pixel) & 0xffff0000) >> 16;
RandAccessor rand(i + j * nx, acc_dev_arr.get_pointer(), arr_size);
vec3 col(0, 0, 0);
for (int s = 0; s < ns; s++) {
float u = float(i + rand.Get()) / float(nx);
float v = float(j + rand.Get()) / float(ny);
ray r = acc_dev_camera.get_pointer()->get_ray(u, v, rand);
col += color_loop(r, acc_dev_sphere.get_pointer(), sphere_size, rand);
}
col /= float(ns);
col = vec3(sqrt(col[0]), sqrt(col[1]), sqrt(col[2]));
int ir = int(255.99 * col[0]);
int ig = int(255.99 * col[1]);
int ib = int(255.99 * col[2]);
acc_dev_pixelsSrc[index] = (0xff000000 | (ib << 16) | (ig << 8) | ir);
});
});
Q.wait();
// <--- Host Accessor to Synchronize Memory
dev_pixelsSrc.get_access<access::mode::read>();
}
catch (sycl::exception& ex)
{
std::cerr << "SYCL Exception thrown: "
<< ex.what() << std::endl;
}
catch (std::exception& ex)
{
std::cerr << "std Exception thrown: "
<< ex.what() << std::endl;
}
下面显示了使用 DPC++ 编译器编译 SYCL 的命令。您可以下载 Intel® oneAPI 工具包，或在 Intel® DevCloud 上注册一个帐户来构建 SYCL 程序。Intel® DevCloud 是免费在最新的 Intel® 硬件和软件上开发、测试和运行工作负载的推荐方式。
dpcpp -O2 -g -std=c++17 SYCLTracer.cpp
登录 Intel® DevCloud 后，创建一个名为 myjob 的作业脚本，其中包含以下内容。PBS_O_WORKDIR 环境变量保存工作进行的位置。a.out 是 DPC++ 编译器输出的可执行文件的名称，因为我们没有给它命名，默认名称是 a.out。如果您想更改可执行文件的名称，请在上面的命令中添加 -o 开关，后跟您想要的名称。
cd $PBS_O_WORKDIR
./a.out
在下方提交您的 CPU 作业。
qsub myjob
在下方提交您的 GPU 作业。
qsub -l nodes=1:gpu:ppn=2 -d . myjob
**重要提示**：要使您的队列在 GPU 上运行，请记住修改此行
queue Q(host_selector{}); // for running on cpu without OpenCL installed
改为以下行
queue Q(gpu_selector{}); // for running on gpu
并使用 dpcpp 重新编译。
要获取工作的进度，请运行 qstat 命令。它返回空表示工作已完成。
qstat
输出存储在 myjob.oXXXXX 中，同样，错误（如果有）记录在 myjob.eXXXXX 中，其中 XXXXX 是作业号的数字代码。**注意**：如果您的作业脚本未命名为 myjob，则您的输出和错误文件名可能不会以 myjob 开头。在输出中，记录了光线追踪计时。这是 24 核 Xeon 处理器的 CPU 计时。
ray_tracer timing:23017ms
这是 GPU 计时。它比 CPU 计时缩短了 2 秒。GPU 规格未知。
ray_tracer timing:21559ms
这是将光线追踪图像从 Intel® DevCloud 复制到本地文件夹的命令。uXXXXX 是您的 Intel® DevCloud 用户 ID，因此请替换为您自己的 ID。scp 命令将 ray_trace.png 从我的 DevCloud 主文件夹安全地复制到我的 D: 驱动器。
scp devcloud:/home/uXXXXX/ray_trace.png /cygdrive/d/
结论
本文分为两个主要部分。第一部分涵盖了 NVIDIA CUDA 的代码转换。为了使 CUDA 和 CPU 共享同一代码库，需要一些宏来将 CUDA 指令定义为空。为了在 CPU 版本中实现并行化，使用了 C++17 parallel for_each。第二部分涵盖了 SYCL 2020 转换、DPC++ 编译和 Intel® DevCloud 执行。使 SYCL 代码在 CPU 或 GPU 上运行，只需在队列实例化时更改一行代码。如前所述，DPC++ 是 Intel 对 SYCL 2020 的实现。对于 NVIDIA 和 AMD，有 hipSYCL，尽管没有得到两家公司官方支持。SYCL 2020 是一个开放标准，跨架构（CPU/GPU/FPGA）且支持多厂商（Intel/AMD/Nvidia），而 CUDA 则专门针对 NVIDIA 生态系统。每当 SYCL 不支持 GPU 时，总有一个 CPU 回退选项可用。总而言之，SYCL 2020 是目前针对/支持更广泛客户群体的最佳选择。
代码库
- 原始《Weekend 的光线追踪》：书籍系列
- 使用 C++17 Parallel for_each 进行光线追踪
- 《Weekend 的光线追踪》移植到 CUDA 包含 CUDA 代码和 C++17 parallel
for_each()代码（用于调用转换后的 CUDA 代码）。 - 《Weekend 的光线追踪》移植到 SYCL 2020
参考文献
- 免费下载的电子书 Data-Parallel C++：使用 C++ 和 SYCL 精通 DPC++ 异构系统编程
- Jeff Hammond 的 Intel® Data-Parallel C++ 教程
历史
- 2021 年 1 月 17 日：首次发布