OOP-Hw:LuisaCompute Code Analysis

Posted on 2024-09-19 Edited on 2024-12-12

这是本学期按面向对象编程课程的大作业“代码分析”报告。第一次，直面奇迹工程。

~~当然，选这个的原因之一也是我和核心开发者私交甚笃~~

1.

1.1. Main Functions and Procedures

$\text{LuisaCompute}$ （以下简称 $\text{LC}$ ），其中 $\text{LUISA}$ 全称为 $\text{Layered and Unified Interfaces on Stream Architectures}$ ，即“流式架构上的分层统一接口“，是一个面向图形等应用的高性能跨平台计算框架。

~~实际上这个全称是硬凑的，取名为Luisa是出于一个很浪漫的原因~~

~~酸死我了~~

$\text{LC}$ 主要功能可分为如下三部分：

一种用于并行核函数编程的嵌入在现代C++中的领域特定语言，利用即时编译技术进行代码生成和编译。
一个统一的运行时库，提供资源封装器用于跨平台资源管理和命令调度。
多个高度优化的并行执行后端，包括CUDA、DirectX、Metal和CPU。

基本流程：

编写核函数，这些函数的编写是嵌入在C++代码中的，例如：

Callable to_srgb = [](Float3 x) {
    $if (x <= 0.00031308f) {
        x = 12.92f * x;
    } $else {
        x = 1.055f * pow(x, 1.f / 2.4f) - .055f;
    };
    return x;
};
Kernel2D fill = [&](ImageFloat image) {
    auto coord = dispatch_id().xy();
    auto size = make_float2(dispatch_size().xy());
    auto rg = make_float2(coord) / size;
    // invoke the callable
    auto srgb = to_srgb(make_float3(rg, 1.f));
    image.write(coord, make_float4(srgb, 1.f));
};

主机端和设备端的代码逻辑从而可以集成在一起。

将 $\text{C++}$ 代码编译，可以得到一个可执行程序。然而我们编写的DSL在此时并没有被执行，它将被编译器进行宏/模板的展开和替换，转化成普通的 $\text{C++}$ 代码，记录下这些代码的结构信息。
$\text{LC}$ 会在运行时解析并生成IR,并交给不同的后端生成代码执行。不同后端被抽象出平台无关的“资源”的概念，从而可以进行统一的编程。

1.1.1. Record the Operations

从高层次来讲， $\text{C++}$ 中的宏，模板等多种语法特性，为开发人员提供了极高的灵活度，使得“在C++代码内实现内嵌的领域特定语言”成为可能。

个人以为，在 $\text{LC}$ 的“语法糖”中十分重要的设计有两点：

通过高超的 $\text{C++}$ 技巧，实现一套完整且灵活的变量及其运算系统。这套系统的特点是：所有编写的代码经过抽象，用户使用起来与一般的C++代码无异，但是在它们的底层，并没有真正地实现运算，它们实际做的事情是记录下进行的运算，以派发给后端进行代码生成。
而从用户的角度看，用户并不知道底层做了什么。
将复杂代码逻辑转化为可调用对象，并且通过灵活组合，使其自动形成 $\text{AST}$ ，易于访问。

下面我们分别讲述这两部分。

1.1.1.1. Variable and Operator System

在LC中，用户编写的代码里会使用到变量等，但是它们并不是真正的变量，而是类似于“占位符”，对这些占位符，可以使用重载运算符的操作，来让它们“看起来”与正常变量一样，而底层实现逻辑得到高度的自定义。

如果我们观察 $\text{LC}$ 里重载运算符的宏：


#define LUISA_MAKE_GLOBAL_DSL_BINARY_OP(op, op_tag_name)                              \
    template<typename Lhs, typename Rhs>                                              \
        requires luisa::compute::any_dsl_v<Lhs, Rhs> &&                               \
                 luisa::is_basic_v<luisa::compute::expr_value_t<Lhs>> &&              \
                 luisa::is_basic_v<luisa::compute::expr_value_t<Rhs>>                 \
    [[nodiscard]] inline auto operator op(Lhs &&lhs, Rhs &&rhs) noexcept {            \
        using R = luisa::compute::detail::dsl_binary_op_return_type<                  \
            luisa::compute::BinaryOp::op_tag_name, Lhs, Rhs>;                         \
        return luisa::compute::dsl::def<R>(                                           \
            luisa::compute::detail::FunctionBuilder::current()->binary(               \
                luisa::compute::Type::of<R>(),                                        \
                luisa::compute::BinaryOp::op_tag_name,                                \
                luisa::compute::detail::extract_expression(std::forward<Lhs>(lhs)),   \
                luisa::compute::detail::extract_expression(std::forward<Rhs>(rhs)))); \
    }

阅读源码可知，def<R>返回的是一个类似于Var<R>的类型（并不严格是）。知道了这一点，我们就会发现这段代码虽然复杂，但是逻辑很好理解：推断出二元运算操作的返回值类型R，然后在当前正在构造的 $\text{LC}$ 函数对象中创建一个二元表达式，将其作为一个具有R类型的临时变量。与此同时，返回这个临时变量的指针，可以使其继续参与运算，从而自然地形成表达式树。

那么到这里，我们要问一个问题： $\text{C++}$ 特有的“模板”系统，对于LC的设计来说真的很重要吗？

一开始，我的看法是：如果没有模板，那么整个 $\text{LC}$ 的抽象便无法成立了。
但是后来我改变了想法。在这个模块中，“模板”只是一种在编译期实现多态的方式，编译器能够看到用户的代码，然后选取合适的模板进行实例化，在实例化的类型和函数中完成分发和记录。
真正核心的操作在于将“记录”封装为“运算”，而不在于多态分发的过程。即使我们不使用模板，我们一样可以编写一套独立的变量系统，通过动态多态的方式来针对不同的运算完成不同的操作记录。
其缺点在于，由于没有模板，编译器无法拿到更多的信息进行优化，也会引入动态多态的分发开销，会造成极大的性能损失，其优点在于可以带来动态的分析和执行。而 $\text{LC}$ 本身的定位是作为源码库引入项目，
它的使用场景只有静态，并且重视性能，因此选择模板实现性能更高的“静态多态”（或者“编译期多态”）更为合适，实现上虽然技巧复杂，但是配合宏的确减少了大量的重复代码。

对此，我总结为： $\text{LC}$ 本身的设计并不依赖于模板系统，即使没有模板，一样可以通过动态多态完成抽象的任务，同样易于拓展。只是考虑到 $\text{LC}$ 本身为了能够在编译期拿到信息实现大量的优化，出于性能考虑选择了模板来实现“编译期多态”。

到此，我们就知道了 $\text{LC}$ 如何记录用户的基本运算操作。

1.1.1.2. Complex Statments

如果不考虑任何高层语法糖的封装，那么一个基本的实现应该是：

用户编写的基本运算被通过某种方式记录下来，这一点我们已经在前面提过，利用的是 $\text{LC}$ 独特的变量系统和精巧的实现。
更高层的语法树节点对象，例如if/else，for等，它们有着不同的代码执行逻辑，它们能够保存不同的可调用对象，将这些可调用对象按照节点自身的执行特点进行组合。

$\text{实现细节}$
在 $\text{LC}$ 中的 $\text{DSL}$ 中，像$if，$else，$return等关键字都是通过宏定义实现的。这些宏定义在include/luisa/dsl/sugar.h中。

理解 $\text{LC}$ 实现的第一步，是从高层理解这些宏做了什么，它们是如何将用户的代码记录为 $\text{AST}$ 的。在include/luisa/dsl中，我们可以看到有关这一部分的实现。

以前面to_rgb里出现的if/else的定义为例：


#define $return(...) ::luisa::compute::return_(__VA_ARGS__)

#define $if(...)                                                                  \
    ::luisa::compute::detail::IfStmtBuilder::create_with_comment(                 \
        ::luisa::compute::dsl_detail::format_source_location(__FILE__, __LINE__), \
        __VA_ARGS__) %                                                            \
        [&]() noexcept
#define $else \
    / [&]() noexcept
#define $elif(...) \
    *([&] { return __VA_ARGS__; }) % [&]() noexcept

乍一看很复杂，但是如果我们把宏替换进去就能够看出 $\text{LC}$ 的实现意图：


Callable to_srgb = [](Float3 x) {
    ::luisa::compute::detail::IfStmtBuilder::create_with_comment(
        ::luisa::compute::dsl_detail::format_source_location(__FILE__, __LINE__), x <= 0.00031308f) %
        [&]() noexcept {
            x = 12.92f * x;
        } / [&]() noexcept {
            x = 1.055f * pow(x, 1.f / 2.4f) - .055f;
        };
    return x;
};

我们可以看出，每个分支内的代码在宏展开以后都替换为了无返回值的lambda表达式。这两个lambda表达式通过重载运算符，和工厂函数IfStmtBuilder::create_with_comment创建出的对象连接在一起，形成了一个IfStmt的结构。我们暂且可以不管这些重载运算符是什么含义，这里就体现了 $\text{LC}$ 里面非常重要的一个想法，就是将代码转化为可调用对象的形式。

实际上，在 $\text{C++}$ 中，类似的设计可谓常见。降低代码耦合的关键，就是提出不变的部分，然后将“可变化”的部分作为依赖注入其中。对于函数式而言，“可变化”的部分就是函数，对于面向对象而言，“可变化”的部分就是虚接口。这种“可调用对象”的想法，在 $\text{C++ STL}$ 的泛型算法库中得到了大量使用，在 $\text{LC}$ 中也不例外，可以认为是一种函数式编程思想的体现（或者也可以说，函数也是对象，这也是面向对象编程的一种体现）。而在这里，如果使用虚接口，那么用户将不得不自己编写派生类实现虚接口，不仅冗长，性能上也会大打折扣。使用lambda表达式，可以直接在原地编写代码，十分方便。

此外，如果没有宏的封装，那么到此，也不过是用户手写几个lambda函数加上一堆又臭又长的调用，是 $\text{C++}$ 项目里的常见写法，但是这一对于用户来说就暴露了大量的细节。宏的包装真正对用户隐藏了他们不需要关心的部分，把核心逻辑的编写留给了用户。

1.1.2. How to Run the Kernels?

在 $\text{1.1.1}$ 中，我们从高层看到了 $\text{LC}$ 的核函数是如何编写并且被保存记录的，接下来我们需要思考，在程序运行的时候，这些核函数是如何被执行的。

这里摘取tests/test_helloworld.cpp中的部分代码：


Context context{argv[0]};
Device device = context.create_device(argv[1]);
Stream stream = device.create_stream();
Image<float> image{device.create_image<float>(PixelStorage::BYTE4, resolution)};
luisa::vector<std::byte> host_image(image.view().size_bytes());
Kernel2D kernel = [&]() {
    ...
};
Shader2D<> shader = device.compile(kernel);
stream << shader().dispatch(resolution) << image.copy_to(host_image.data()) << synchronize();

可以总结为：传入参数选择后端 -> 创建上下文 -> 创建设备 -> 创建指令流 -> 编译核函数并创建着色器 -> 将着色器分发至后端的执行流 -> 从后端收集数据并同步。

1.1.2.1. Context

“上下文”是一个非常常见的概念，但是在不同的应用中，它的含义也多有所差异，但是总体而言是用来代替全局变量，单例等等，管理一些高层次的状态的。“上下文”的主要作用是避免使用全局的状态，可以认为与“单例“模式截然相反。要理解为什么需要使用上下文，就要理解单例模式有哪些问题。

工程中的直觉告诉我们，在程序中保存一个全局状态是非常危险的。在StackOverflow的指引下，我找到了这篇博文：Singleton Anti-Pattern。

具体而言，保存全局状态其实就是单例模式的一种体现。而“单例”————也就是要求整个进程中只能有一个实例————的要求是非苛刻的。考虑到线程安全等等问题，要在现有的程序中保证这个假设就已经很困难了。而对于未来潜在的拓展需求，这种假设更是自缚手脚。因此，无论是为了编码的方便。还是为了程序的可维护性和可拓展性，我们都应该避免使用全局状态。这也就是为什么“上下文”是一个非常好的设计选择。

在 $\text{LC}$ 中，Context保存了运行时目录，校验层等等，也的确是我们所想的“上下文”的概念。

除此以外，在 $\text{LC}$ 的实现中还有个有趣的小模式： $\text{P-Impl}$ 模式。在include/luisa/runtime/context.h中，我们可以看到Context的定义：


class LC_RUNTIME_API Context {

private:
    luisa::shared_ptr<detail::ContextImpl> _impl;

public:
//  methods...
};

这是 $\text{C++}$ 独有的设计模式，有很多的好处。

这样可以在不改变接口的情况下，修改实现细节。
可以分离成员变量的定义，避免引入过多头文件导致的编译时间过长。
Context实例可以简单快速地拷贝，而无需担忧资源管理的问题。实际上这一点也得益于使用了shared_ptr。在我个人的实践中，如果使用 $\text{P-Impl}$ 模式，我更倾向于使用unique_ptr并禁用移动构造和拷贝构造，明确所有权，保证一份ContextImpl实例必然仅被一个Context实例持有。无论如何， $\text{P-Impl}$ 模式都极大地降低了资源管理的负担。
无论我们如何修改实现，Context编译出的二进制接口都是不变的，这样就可以保证二进制兼容性。如果不使用 $\text{P-Impl}$ 模式，那么我们修改成员变量就很可能导致类内存布局发生变化，进而导致二进制不兼容。

1.1.2.2. Device

Device的结构非常简单，它的成员变量只有一个_impl，这里看起来似乎有点儿 $\text{P-Impl}$ 模式的意思，不过据我观察，实际上Device更像是一个包装类。

class LC_RUNTIME_API Device {

public:
    using Deleter = void(DeviceInterface *);
    using Creator = DeviceInterface *(Context && , const DeviceConfig * );
    using Handle = luisa::shared_ptr<DeviceInterface>;

private:
    Handle _impl;   
}

Device本身的职责基本就是：

作为工厂，创建资源。它有不少create_XXX的方法，当然这些最后还是转发给DeviceInterface来分别实现的。
着色器程序的编译，加载和管理。

实际上Device最后还是会把各种方法转交给DeviceInterface来实现。DeviceInterface非常明显，是一个抽象类。诸如 $\text{Cuda, Vulkan, DirectX}$ 等等与显卡交互的接口其实都有一些共性，例如资源管理之类，这些共性可以被提取出来，也就是DeviceInterface类型。

例如：

// ...

[[nodiscard]] virtual BufferCreationInfo create_buffer(const Type *element, size_t elem_count) noexcept = 0;
  [[nodiscard]] virtual BufferCreationInfo create_buffer(const ir::CArc<ir::Type> *element, size_t elem_count) noexcept = 0;
  virtual void destroy_buffer(uint64_t handle) noexcept = 0;

  [[nodiscard]] virtual ResourceCreationInfo create_texture(
      PixelFormat format, uint dimension,
      uint width, uint height, uint depth,
      uint mipmap_levels, bool simultaneous_access) noexcept = 0;
  virtual void destroy_texture(uint64_t handle) noexcept = 0;

  // bindless array
  [[nodiscard]] virtual ResourceCreationInfo create_bindless_array(size_t size) noexcept = 0;
  virtual void destroy_bindless_array(uint64_t handle) noexcept = 0;

 // ...

我们可以猜想到，不同的后端就刚好对应于不同的DeviceInterface的派生类。这样的设计，使得 $\text{LC}$ 可以很方便地拓展到不同的后端，只需要实现一个新的DeviceInterface即可。

$\text{从字符串到后端}$

这里就不免涉及到一个问题：用户输入的是字符串（argv[1]）， $\text{LC}$ 是如何创建具体的后端的？

Context类作为工厂，调用create_device，接受后端的名字（也就是一个字符串），返回为用户选择的后端的Device对象。

Device Context::create_device(luisa::string_view backend_name_in, const DeviceConfig *settings, bool enable_validation) noexcept {
    luisa::string backend_name{backend_name_in};
    for (auto &c : backend_name) { c = static_cast<char>(std::tolower(c)); }
    auto &&m = _impl->load_backend(backend_name);
    auto interface = m.creator(Context{_impl}, settings);
    // ...
}

所以这里的关键是load_backend返回的东西。


struct BackendModule {
    using BackendDeviceNames = void(luisa::vector<luisa::string> &);
    DynamicModule module;
    Device::Creator *creator;
    Device::Deleter *deleter;
    BackendDeviceNames *backend_device_names;
};

[[nodiscard]] const BackendModule &load_backend(const luisa::string &backend_name) noexcept {
    if (std::find(installed_backends.cbegin(),
                    installed_backends.cend(),
                    backend_name) == installed_backends.cend()) {
        LUISA_ERROR_WITH_LOCATION("Backend '{}' is not installed.", backend_name);
    }
    std::scoped_lock lock{module_mutex};
    if (auto iter = loaded_backends.find(backend_name);
        iter != loaded_backends.cend()) { return *iter->second; }
    BackendModule m{.module = DynamicModule::load(
                        runtime_directory,
                        luisa::format("lc-backend-{}", backend_name))};
    LUISA_ASSERT(m.module, "Failed to load backend '{}'.", backend_name);
    m.creator = m.module.function<Device::Creator>("create");
    m.deleter = m.module.function<Device::Deleter>("destroy");
    m.backend_device_names = m.module.function<BackendModule::BackendDeviceNames>("backend_device_names");
    auto pm = loaded_backends
                    .emplace(backend_name, luisa::make_unique<BackendModule>(std::move(m)))
                    .first->second.get();
    return *pm;
}

BackendModule记录了一个后端的基本信息，包括创建和销毁，它们是通过查找动态链接库中的符号create和destroy来实现的。
这个方法会根据后端的名字，查找相应后端编译出的动态库，加载并记录其创建和销毁函数。

这样就可以获取到用户选择的DeviceInterface的工厂函数create（也就是BackendModule的Creator类型），从而创建一个Device对象。

我之前一直不清楚究竟怎么编写多后端支持的程序， $\text{LC}$ 的实现给了我非常大的启发。

1.1.2.3. Stream

我们关注它重载了哪些<<运算：

Delegate operator<<(luisa::unique_ptr<Command> &&cmd) noexcept;
Delegate operator<<(luisa::move_only_function<void()> &&f) noexcept;
template<typename T>
    requires std::is_rvalue_reference_v<T &&> && is_stream_event_v<T>
Stream &operator<<(T &&t) noexcept {
    luisa::invoke(std::forward<T>(t), device(), handle());
    return *this;
}
Stream &operator<<(CommandList::Commit &&commit) noexcept;
Stream &operator<<(Synchronize &&) noexcept;

先看前两个，Stream能够接受一个命令或者函数（注意这里的右值引用，意味着它转移了所有权），并产生一个委托Delegate。

第三个，如果向Stream推入一个事件，则执行该事件。后两个则是提交和同步，这两个类实际上都是空类，这里只是作 $\text{tag dispatch}$ 的作用。

Delegate的定义如下：

class LC_RUNTIME_API Delegate {
    private:
        Stream *_stream;
        CommandList _command_list;
    // ...
}

我们发现Stream可以一个一个地推入命令或者函数，但是Delegate里只有一个CommandList，因此我们猜测这个CommandList其实是个缓冲区，推入的命令不会立即执行。

Stream支持的<<运算，Delegate全部都支持，我们看其中<<(std::unique_ptr<Command> &&cmd)的实现：

1
2
3

if (!_command_list.callbacks().empty()) { _commit(); }
    _command_list.append(std::move(cmd));
    return std::move(*this);

所以这就很清楚了，Stream推入的命令会被缓存得到Delegate，向Delegate继续推入命令则会继续缓存。

Delegate对于推入事件的处理与Stream小有区别：

template<typename T>
            requires std::is_rvalue_reference_v<T &&> && is_stream_event_v<T>
Stream &operator<<(T &&t) && noexcept {
    _commit();
    luisa::invoke(std::forward<T>(t), _stream->device(), _stream->handle());
    return *_stream;
}

它会立刻提交缓存的所有命令，然后执行推入的事件。包括sycnrhonize也会先进行提交。这和一些并行计算API的设计是类似的。它返回的是Stream的引用，这意味着命令缓存被清空了。

$\text{为什么要使用Delegate?}$

为什么不直接给Stream添加一个CommandList缓冲区，而要使用Delegate呢？关于这个问题我问了核心开发者。

Delegate是一个临时对象，它生存到该语句结束，也就是没有新的东西被推入流中，通过给Delegate设计析构函数，可以实现在一条语句结束时自动提交。

1
2
3

Stream::Delegate::~Delegate() noexcept {
    _commit();
}

Delegate对象生存期只持续到一条语句结束，这样结束的时候就可以自动提交，算是一个小技巧。

1.1.2.4 Shader

着色器原义指在显卡上运行的程序。在这里自然也就拓展为一般意义下交给后端执行的程序。

这一部分反而从设计上暂时没有什么好讲的。Device通过compile方法，由具体后端实现将 $\text{LC}$ 的可调用对象编译为Shader，Shader的dispatch(dispatch_sizes)方法返回一个ShaderInvoke对象，可以直接被推入Stream。

1.2. Modules

$\text{LC}$ 库主要可分为如下几个重要模块：

$\text{Core}$ ：核心模块，主要包括一套专用的STL，并行原语，基础数学函数，日志等基本工具。
$\text{AST}$ ：抽象语法树模块，包括用于表示核函数 $\text{AST}$ 的基本类型。
$\text{Runtime}$ ： $\text{LC}$ 的运行时部分。
$\text{Backends}$ ：后端模块，包括 $\text{CUDA}$ ， $\text{DirectX}$ ， $\text{Metal}$ 和 $\text{CPU}$ 等不同的后端。
$\text{DSL}$ ：领域特定语言模块。
$\text{Tests}$ ：功能测试。

1.3. Summary

这个部分简单分析了一下 $\text{LC}$ 的基本工作流程，目前我们暂不关心 $\text{LC}$ 的后端，只是从 $\text{LC}$ 提供的最上层抽象，也就是它的 $\text{DSL}$ 模块提供的编程模型，来分析它的前端原理。

2. Backend and Command Execution

2.1. Backend

接下来我们以 $\text{LC}$ 的 $\text{Cuda}$ 后端为例，来看一下后端的工作流程。

我们能在src/backends/cuda里找到cuda_device.h，这就是 $\text{CUDA}$ 后端，它直接继承自DeviceInterface，并实现了DeviceInterface的所有纯虚函数。

按照前面所说， $\text{LC}$ 编译成着色器的过程，是由具体后端负责的，它调用的是后端的create_shader方法。

这里其实很大一部分都是比较接近于传统编译器里将 $\text{AST}$ 翻译成 $\text{IR}$ 的过程。这里 $\text{LC}$ 后端做的事情其实是把 $\text{LC}$ 的 $\text{AST}$ 翻译成 $\text{Cuda}$ 代码，然后交给 $\text{nvrtc}$ 编译成 $\text{PTX}$ 代码。

我们先看compile的大致流程。它其实是调用了下面的create_shader方法。


ShaderCreationInfo CUDADevice::create_shader(const ShaderOption &option, Function kernel) noexcept {

    // ...

    // codegen
    Clock clk;
    StringScratch scratch;
    CUDACodegenAST codegen{scratch, !_cudadevrt_library.empty()};
    codegen.emit(kernel, _compiler->device_library(), option.native_include);
    LUISA_VERBOSE("Generated CUDA source in {} ms.", clk.toc());

    // process bound arguments
    // ...

    // NVRTC nvrtc_options
    auto sm_option = luisa::format("-arch=compute_{}", _handle.compute_capability());
    auto nvrtc_version_option = luisa::format("-DLC_NVRTC_VERSION={}", _compiler->nvrtc_version());
    auto optix_version_option = luisa::format("-DLC_OPTIX_VERSION={}", optix::VERSION);
    // ...

    // compute hash
    auto src_hash = _compiler->compute_hash(scratch.string(), nvrtc_options);

    // create metadata
    CUDAShaderMetadata metadata{
        // ...
    };
    return _create_shader(option.name, scratch.string(),
                          option, nvrtc_options,
                          metadata, std::move(bound_arguments));
}

2.1.1. Code Generation

所以可以看出来，代码生成的关键就是CUDACodegenAST，而生成代码的核心方法是emit方法。这个类其实就是一个很传统的编译器，用来把 $\text{AST}$ 翻译成 $\text{Cuda}$ 代码。

在编译器，代码分析以及优化中，非常常用的一个设计模式就是 $\text{Visitor}$ 模式。就我自己而言，在此之前已经对这个模式有所使用，同时因为自己也开发过编译器前端（其实就是学校的实验），因此还算比较熟悉。

在编译器以及各种代码优化中，优化 $\text{Pass}$ 等等需要对语法树进行分析遍历的东西就是 $\text{Visitor}$ ，而 $\text{AST}$ 的节点就是 $\text{Acceptor}$ 。 $\text{LC}$ 作为某种意义上的编译系统，会使用这个设计模式自然也是不意外的。

class CUDACodegenAST final : private TypeVisitor, private ExprVisitor, private StmtVisitor {

public:
    friend class CUDAConstantPrinter;
    class RayQueryLowering;

private:
    StringScratch &_scratch;
    Function _function;
    luisa::vector<uint64_t> _generated_functions;
    luisa::vector<uint64_t> _generated_constants;
    luisa::unique_ptr<RayQueryLowering> _ray_query_lowering;
    uint32_t _indent{0u};
    bool _allow_indirect_dispatch;

private:
    const Type *_ray_type;
    const Type *_triangle_hit_type;
    const Type *_procedural_hit_type;
    const Type *_committed_hit_type;
    const Type *_ray_query_all_type;
    const Type *_ray_query_any_type;
    const Type *_indirect_buffer_type;

private:
    void visit(const Type *type) noexcept override;
    void visit(const UnaryExpr *expr) override;
    void visit(const BinaryExpr *expr) override;
    void visit(const MemberExpr *expr) override;
    void visit(const AccessExpr *expr) override;
    void visit(const LiteralExpr *expr) override;
    void visit(const RefExpr *expr) override;
    void visit(const CallExpr *expr) override;
    void visit(const CastExpr *expr) override;
    void visit(const TypeIDExpr *expr) override;
    void visit(const StringIDExpr *expr) override;
    void visit(const BreakStmt *stmt) override;
    void visit(const ContinueStmt *stmt) override;
    void visit(const ReturnStmt *stmt) override;
    void visit(const ScopeStmt *stmt) override;
    void visit(const IfStmt *stmt) override;
    void visit(const LoopStmt *stmt) override;
    void visit(const ExprStmt *stmt) override;
    void visit(const SwitchStmt *stmt) override;
    void visit(const SwitchCaseStmt *stmt) override;
    void visit(const SwitchDefaultStmt *stmt) override;
    void visit(const AssignStmt *stmt) override;
    void visit(const ForStmt *stmt) override;
    void visit(const ConstantExpr *expr) override;
    void visit(const CommentStmt *stmt) override;
    void visit(const RayQueryStmt *stmt) override;
    void visit(const AutoDiffStmt *stmt) override;
    void visit(const CpuCustomOpExpr *expr) override;
    void visit(const GpuCustomOpExpr *expr) override;

private:
    void _emit_type_decl(Function f) noexcept;
    void _emit_variable_decl(Function f, Variable v, bool force_const) noexcept;
    void _emit_type_name(const Type *type) noexcept;
    void _emit_function(Function f) noexcept;
    void _emit_variable_name(Variable v) noexcept;
    void _emit_indent() noexcept;
    void _emit_statements(luisa::span<const Statement *const> stmts) noexcept;
    void _emit_constant(Function::Constant c) noexcept;
    void _emit_variable_declarations(Function f) noexcept;
    void _emit_builtin_variables() noexcept;
    void _emit_access_chain(luisa::span<const Expression *const> chain) noexcept;

public:
    CUDACodegenAST(StringScratch &scratch, bool allow_indirect) noexcept;
    ~CUDACodegenAST() noexcept override;
    void emit(Function f,
              luisa::string_view device_lib,
              luisa::string_view native_include);
};

可以看到，CUDACodegenAST继承自TypeVisitor，ExprVisitor和StmtVisitor，这三个类都是抽象类，继承它们的类都需要实现对应的visit方法。

例如，访问BinaryExpr的visit(const BinaryExpr *expr)方法：

void CUDACodegenAST::visit(const BinaryExpr *expr) {
    _scratch << "(";
    expr->lhs()->accept(*this);
    switch (expr->op()) {
        case BinaryOp::ADD: _scratch << " + "; break;
        case BinaryOp::SUB: _scratch << " - "; break;
        case BinaryOp::MUL: _scratch << " * "; break;
        case BinaryOp::DIV: _scratch << " / "; break;
        case BinaryOp::MOD: _scratch << " % "; break;
        case BinaryOp::BIT_AND: _scratch << " & "; break;
        case BinaryOp::BIT_OR: _scratch << " | "; break;
        case BinaryOp::BIT_XOR: _scratch << " ^ "; break;
        case BinaryOp::SHL: _scratch << " << "; break;
        case BinaryOp::SHR: _scratch << " >> "; break;
        case BinaryOp::AND: _scratch << " && "; break;
        case BinaryOp::OR: _scratch << " || "; break;
        case BinaryOp::LESS: _scratch << " < "; break;
        case BinaryOp::GREATER: _scratch << " > "; break;
        case BinaryOp::LESS_EQUAL: _scratch << " <= "; break;
        case BinaryOp::GREATER_EQUAL: _scratch << " >= "; break;
        case BinaryOp::EQUAL: _scratch << " == "; break;
        case BinaryOp::NOT_EQUAL: _scratch << " != "; break;
    }
    expr->rhs()->accept(*this);
    _scratch << ")";
}

这里就出现了 $\text{Visitor}$ 模式的经典实现，让Acceptor来accept(*this)，“我访问你，就是你来接受我”。我们只知道lhs和rhs是const Expression*的子类类型，但是在编译的时候，编译器不知道它们的具体类型，这样，就没有办法直接选择相应的重载。但是通过accept(*this)，编译器起码能知道访问者一定是CUDACodeGenAST。

我们如果观察所有的Expression类型，它们其实都有一个方法：

1	void accept(ExprVisitor &visitor) const override { visitor.visit(this); }

这样编译器就知道了被访问的Expression具体是什么类型。从而又能够让CUDACodegenAST选择正确的重载。

由于ExprVisitor的visit是个虚接口，在运行的时候，就会调用子类visitor实现的visit方法，在这里也就是CUDACodegenAST的visit方法。

Tip 1: 这里每个Expression的派生类都需要写上这个accept函数，并且它们都长得一模一样，但它不能直接提到Expression基类里！
因为如果我们在基类里实现这个方法，那么*this的类型是Expression&，这样怎么都没法得到Expression子类的类型，也就没办法选择正确的重载！

Tip 2: 如果要减少重复代码，可以使用 $\text{C++}$ 的 $\text{CRTP}$ （奇异递归模板模式）技巧，其操作方式如下：
1
2
3
4
5
6
7
8
9
template<typename Derived>
class Expression {
  Derived& derived() { return static_cast<Derived&>(*this); }
  void accept(ExprVisitor &visitor) const { visitor.visit(derived()); }
};

class BinaryExpr : public Expression<BinaryExpr> {
  // ...
};
不知出于何种原因， $\text{LC}$ 并没有使用这种方法。

$\text{Update: }$ 跟作者聊了这个问题，单纯就是作者没想到。。。

Tip 3: 在 $\text{C++ 23}$ 中，有了新的语法糖“显式对象参数“，可以更加简洁地实现 $\text{CRTP}$ 。

为什么要用这么弯弯绕的方式？

首先注意到一个事实：在Visitor模式中，Acceptor往往是没有什么拓展需求的，因为我们Visitor中实际要把Acceptor的每种派生的visit都实现一遍，假如多加一个，或者删除一个Acceptor的派生，那么所有的Visitor都要改。因此，在Visitor模式设计使用的时候，首先就会把较为稳定的部分作为Acceptor。在编译器里，语法节点的种类往往是比较固定的，而访问AST的方式是多种多样的，因此Visitor模式就很适合作为各种Pass。

如果我们不使用各种各样的Visitor，而是一股脑全塞到 $\text{AST}$ 节点的接口里，那么在拓展的时候就极为麻烦（例如，写新的代码生成，或者优化 $\text{Pass}$ 的时候）并且由于 $\text{AST}$ 节点种类比较固定，而优化 $\text{Pass}$ 经常拓展，这样做是彻底违背开闭原则的。

CUDACodegenAST的emit方法实际上就是先翻译一些宏，或者选项，核函数配置，然后调用_emit_function方法。_emit_function…就是生成核函数的代码了。有了Visitor方法的加持，接下来的事情可以说，跟普通编译器没什么区别了。

$\text{Code Generation}$ 的过程实际上是很容易拓展的，例如我们完全可以把CUDACodegenAST替换成别的什么代码生成 $\text{Pass}$ ，实现定制的需求。例如，就我所知，在现在的 $\text{LC}$ 中，已经在尝试着引入 $\text{LLVM}$ 了，相信在这个框架下，是不需要太多的重构的。

$\text{Update: }$ 现在 $\text{LC}$ 在这一部分已经弃用 $\text{Visitor}$ ，转而改用switch-case了。。。原因似乎是他们短期内并不打算写很多 $\text{Pass}$ ，这样一来还不如直接写switch-case来得方便呢。

2.1.2. CUDA Code to Binary

当然，人类可读的 $\text{Cuda}$ 代码， $\text{GPU}$ 肯定是不认的，总归是得把它翻译成 $\text{GPU}$ 能执行的二进制代码。这就是前面的create_shader方法中，最后返回的_create_shader方法做的事情。

讲道理， $\text{LC}$ 这个命名做的是真不行，create_shader和_create_shader，这两个方法名字差不多，但是做的事情差别也太大了，不去读代码，还真的不容易搞清楚这两个的区别，你说你把名字改成什么compile_cuda_to_ptx不是挺清晰的吗？


ShaderCreationInfo CUDADevice::_create_shader(luisa::string name,
                                              const string &source, const ShaderOption &option,
                                              luisa::span<const char *const> nvrtc_options,
                                              const CUDAShaderMetadata &expected_metadata,
                                              luisa::vector<ShaderDispatchCommand::Argument> bound_arguments) noexcept {

    // generate a default name if not specified
    auto uses_user_path = // ...
    // ...
    auto metadata_name = // ...

    // try disk cache
    auto ptx = [&] {
        luisa::unique_ptr<BinaryStream> ptx_stream;
        luisa::unique_ptr<BinaryStream> metadata_stream;
        if (uses_user_path) {
            ptx_stream = _io->read_shader_bytecode(name);
            metadata_stream = _io->read_shader_bytecode(metadata_name);
        } else if (option.enable_cache) {
            ptx_stream = _io->read_shader_cache(name);
            metadata_stream = _io->read_shader_cache(metadata_name);
        }
        return load_shader_ptx<false>(
            metadata_stream.get(), ptx_stream.get(),
            name, false, expected_metadata);
    }();

    // compile if not found in cache
    if (ptx.empty()) {
        luisa::filesystem::path src_dump_path;
        if (option.enable_debug_info || LUISA_CUDA_DUMP_SOURCE) {
            // ...
            if (uses_user_path) {
                src_dump_path = _io->write_shader_bytecode(src_name, src_data);
            } else if (option.enable_cache) {
                src_dump_path = _io->write_shader_cache(src_name, src_data);
            }
        }
        luisa::string src_filename{src_dump_path.string()};
        ptx = _compiler->compile(source, src_filename, nvrtc_options, &expected_metadata);
        if (!ptx.empty()) {
            // ...
        }
    }

    if (option.compile_only) {// no shader object should be created
        return ShaderCreationInfo::make_invalid();
    }

    // create the shader object
    // ...
#ifndef NDEBUG
    p->set_name(std::move(name));
#endif
    ShaderCreationInfo info{};
    // filling info...
    return info;
}

在最后返回的ShaderCreationInfo里，记录了刚创建的Shader的句柄。不同后端使用的Shader显然是大不一样的，但是使用句柄，就可以很方便地在前端标识和使用了，至于这个句柄是什么。。。后面揭晓。

说实话，这个方法的传入参数太多太杂了，完全应该用一个新的结构体来传递这些参数————万一以后这里要拓展参数呢？

可以看到， $\text{LC}$ 的编译是有磁盘缓存的机制的。如果找不到缓存，那么会进行一次编译，也就是ptx = _compiler->compile(source, src_filename, nvrtc_options, &expected_metadata)这一行。

我们看这个函数的实现。


luisa::string CUDACompiler::compile(const luisa::string &src, const luisa::string &src_filename,
                                    luisa::span<const char *const> options,
                                    const CUDAShaderMetadata *metadata) const noexcept {

    Clock clk;

    auto hash = metadata ? metadata->checksum : compute_hash(src, options);

    if (auto ptx = _cache->fetch(hash)) { return *ptx; }

    nvrtcProgram prog;
    auto filename = src_filename.empty() ? "my_kernel.cu" : src_filename.c_str();
    LUISA_CHECK_NVRTC(nvrtcCreateProgram(
        &prog, src.data(), filename, 0, nullptr, nullptr));
    auto error = nvrtcCompileProgram(prog, static_cast<int>(options.size()), options.data());

    // log this compilation
    // ...

    size_t ptx_size;
    luisa::string ptx;
    // TODO: use OptiX IR for ray tracing shaders
    // ...
    LUISA_CHECK_NVRTC(nvrtcDestroyProgram(&prog));
    LUISA_VERBOSE("CUDACompiler::compile() took {} ms.", clk.toc());
    return ptx;
}

到这里，就没什么好说的了，就是在调用 $\text{nvrtc}$ （ $\text{Cuda}$ 的运行时编译器）将 $\text{Cuda}$ 代码编译成二进制代码。

回到_create_shader，除了编译以外，自然还要处理二进制文件的 $\text{IO}$ 。这个方法里最重要的函数调用就是write_shader_bytecode，负责做这个事情的类型是BinaryIO抽象类。


class BinaryIO {

public:
    virtual ~BinaryIO() noexcept = default;
    [[nodiscard]] virtual luisa::unique_ptr<BinaryStream> read_shader_bytecode(luisa::string_view name) const noexcept = 0;
    // ...
    // returns the path of the written file (if stored on disk, otherwise returns empty path)
    [[nodiscard]] virtual luisa::filesystem::path write_shader_bytecode(luisa::string_view name, luisa::span<std::byte const> data) const noexcept = 0;
    // ...
};

这里我的理解是， $\text{LC}$ 本身并不关心二进制文件的 $\text{IO}$ 细节，使用一个接口更有利于应对潜在的拓展需求：比如如果要从网络上下载二进制文件，或者直接从内存中读取二进制文件，这样的需求都可以通过实现BinaryIO接口来实现。

2.2. Command Execution

接下来我们把视角从后端切回前端，看一下 $\text{LC}$ 具体是如何把前端的命令派发到后端的。

在 $1$ 中，我们已经知道， $\text{LC}$ 通过将命令Command推到Stream中来派发。我们来看Command类的实现。


class Command {

public:
    enum struct Tag {
#define LUISA_MAKE_COMMAND_TAG(Cmd) E##Cmd,
        LUISA_MAP(LUISA_MAKE_COMMAND_TAG, LUISA_COMPUTE_RUNTIME_COMMANDS)
#undef LUISA_MAKE_COMMAND_TAG
    };

private:
    Tag _tag;

public:
    explicit Command(Tag tag) noexcept : _tag(tag) {}
    virtual ~Command() noexcept = default;
    virtual void accept(CommandVisitor &visitor) const noexcept = 0;
    virtual void accept(MutableCommandVisitor &visitor) noexcept = 0;
    [[nodiscard]] auto tag() const noexcept { return _tag; }
    [[nodiscard]] virtual StreamTag stream_tag() const noexcept = 0;
};

可以注意到，又是 $\text{Visitor}$ 模式，在这里就不详述了。要理解 $\text{LC}$ 的派发，我们主要看ShaderInvoke类的dispatch方法。


[[nodiscard]] auto dispatch(uint2 size) && noexcept {
    return this->_parallelize(uint3{size.x, size.y, 1u}).build();
}

_parallelize返回一个ComputeDispatchCmdEncoder，这个类其实是一个比较典型的工厂类，它设置好自己的各种成员变量，最后将这些成员变量作为构造参数，通过一个build方法返回一个ShaderDispatchCommand。


luisa::unique_ptr<ShaderDispatchCommand> ComputeDispatchCmdEncoder::build() && noexcept {
    if (_argument_idx != _argument_count) [[unlikely]] {
        LUISA_ERROR("Required argument count {}. "
                    "Actual argument count {}.",
                    _argument_count, _argument_idx);
    }
    return luisa::make_unique<ShaderDispatchCommand>(
        _handle, std::move(_argument_buffer),
        _argument_count, _dispatch_size);
}

这里的ShaderDispatchCommand正是Command的一个子类，它记录了核函数派发调用的信息。

我们再看Stream的这个方法：


Stream &Stream::operator<<(CommandList::Commit &&commit) noexcept {
    _dispatch(std::move(commit).command_list());
    return *this;
}

也就是说关键是这里的_dispatch方法，它实际上是调用了后端的dispatch方法。所以最后还是要回到后端去。。。

而如果我们看CUDADevice的dispatch，它实际上是在调用CUDAStream的dispatch。


void CUDAStream::dispatch(CommandList &&command_list) noexcept {
    CUDACommandEncoder encoder{this};
    auto commands = command_list.steal_commands();
    auto callbacks = command_list.steal_callbacks();
    {
        std::scoped_lock lock{_dispatch_mutex};
        for (auto &cmd : commands) { cmd->accept(encoder); }
        encoder.commit(std::move(callbacks));
    }
}

所以这里的关键在于Command接受了encoder。encoder是CUDACommandEncoder类型的成员，而CUDACommandEncoder继承了MutableCommandVisitor，这样一切就合情合理了。我们只需要看最后CUDACommandEncoder对于ShaderDispatchCommand的visit方法。


void CUDACommandEncoder::visit(ShaderDispatchCommand *command) noexcept {
    reinterpret_cast<CUDAShader *>(command->handle())->launch(*this, command);
}

command->handle()返回的实际上是一个uint64_t类型，因此我们可以猜出来，这里的handle实际上就是CUDAShader对象的地址，这就解决了前面“返回给前端的句柄到底是什么”的疑问。


void CUDAShader::launch(CUDACommandEncoder &encoder,
                        ShaderDispatchCommand *command) const noexcept {
    // ...
    _launch(encoder, command);
    // ...
}

我发现这个handle的技巧在 $\text{C++}$ 里是很好用的，对于不同类型的对象，它们总归是要放在内存里的，我们直接取它们的地址作为一个统一的UID，也就是handle，这样就可以很方便地对于不同的类型进行统一的处理。

CUDAShader是一个虚类，_launch就是需要派生类实现的虚函数，我们只需要看CUDAShaderNative这个派生就好了，其他的派生类是为了支持 $\text{OptiX}$ 的。


void CUDAShaderNative::_launch(CUDACommandEncoder &encoder, ShaderDispatchCommand *command) const noexcept {

    static thread_local std::array<std::byte, 65536u> argument_buffer;// should be enough

    auto argument_buffer_offset = static_cast<size_t>(0u);
    auto allocate_argument = [&](size_t bytes) noexcept {
       // ... 
    };

    auto encode_argument = [&allocate_argument, command](const auto &arg) noexcept {
        using Tag = ShaderDispatchCommand::Argument::Tag;
        switch (arg.tag) {
            // ...
        }
    };
    for (auto &&arg : _bound_arguments) { encode_argument(arg); }
    for (auto &&arg : command->arguments()) { encode_argument(arg); }
    // launch
    auto cuda_stream = encoder.stream()->handle();
    if (command->is_indirect()) {
        // ...
        // launch built-in indirect dispatch kernel to run the actual kernel
    } else {
        // ...
        for (auto dispatch_size : dispatch_sizes) {
            auto launch_size_and_kernel_id = make_uint4(dispatch_size, 0u);
            // compute
            LUISA_CHECK_CUDA(cuLaunchKernel(
                _function,
                blocks.x, blocks.y, blocks.z,
                block_size.x, block_size.y, block_size.z,
                0u, cuda_stream,
                &arguments, nullptr));
        }
    }
}

这就非常明确了：首先把参数编码到argument_buffer里，然后调用cuLaunchKernel来启动核函数。

2.4. Summary

这次分析了 $\text{LC}$ 的 $\text{CUDA}$ 后端。结合前面的分析，我们就完整理解了 $\text{LC}$ 整个从记录 $\text{AST}$ 到运行时编译，再到运行时派发执行的过程。

3. Resource Management

对于一个 $\text{C++}$ 程序而言，在默认的情况下我们没有垃圾回收之类的机制来“擦屁股”。所以除了软件本身的架构逻辑以外，编写 $\text{C++}$ 程序还有一个很重要的因素要考虑，就是资源管理。

实际上，这两者往往是相辅相成的，一个好的架构逻辑，往往会明确资源的所有权（也就是，将某一个资源的分配和释放具体到一个或者多个对象），从而使得资源管理变得更加容易，同时容易也让人明白类的职责。
除此以外，得益于 $\text{C++}$ 极强的灵活性，我们也可以自由地设计辅助工具来进行辅助。

3.1. Memory Management

$\text{C++}$ 中，内存管理始终是一个非常重要的问题。一方面，在没有垃圾回收的情况下，只有精确的分配释放，才能保证程序持续安全运行，另一方面，内存分配的频率，时机，大小也会直接影响程序的性能。

在 $\text{LC}$ 中，一切的开始都是Context。如前所述，Context使用shared_ptr指向ContextImpl，这样保证了程序在运行时，不管Context怎么拷贝，实际上只会有一个ContextImpl实例。shared_ptr实现上使用了原子性的引用计数，因此只要在不出现循环指向的情况下，最后一个被析构的shared_ptr总会把引用计数置空，并随即把指向的实例析构。这一点上很类似 $\text{Rust}$ 的Arc，不过它是可以对指向的对象进行修改的。

ContextImpl涉及的成员如下：


std::filesystem::path runtime_directory;
luisa::unordered_map<luisa::string, luisa::unique_ptr<BackendModule>> loaded_backends;
luisa::vector<luisa::string> installed_backends;
ValidationLayer validation_layer;
luisa::unordered_map<luisa::string, luisa::unique_ptr<std::filesystem::path>> runtime_subdir_paths;
std::mutex runtime_subdir_mutex;
std::mutex module_mutex;

在ContextImpl被析构时，它的所有成员都会自动析构，而 $\text{STL}$ 容器以及unique_ptr的使用暗示这个类型对
后端信息以及runtime_subdir_paths拥有所有权，当ContextImpl被析构时，这些成员析构函数被调用，会连带着所有后端以及指向的std::filesystem::path被析构。

所以，只要BackendModule和ValidationLayer的内存管理设计妥当，那么整个ContextImpl类都是安全的。而在前面，我们看到BackendModule唯一一个自建类型的成员是DynamicModule，因此我们就要看DynamicModule的设计实现。

class LC_CORE_API DynamicModule : concepts::Noncopyable {

private:
    void *_handle{nullptr};
    // ...
}

首先，DynamicModule是继承自 $\text{LC}$ 中的concepts::Noncopyable，这个类的实现如下：

struct Noncopyable {
    Noncopyable() noexcept = default;
    Noncopyable(const Noncopyable &) noexcept = delete;
    Noncopyable &operator=(const Noncopyable &) noexcept = delete;
    Noncopyable(Noncopyable &&) noexcept = default;
    Noncopyable &operator=(Noncopyable &&) noexcept = default;
};

这是 $\text{C++}$ 中一种常用技巧。因为父类拷贝构造和拷贝赋值被禁用了，所有继承自这个类的类型都不允许再拷贝，保证了资源的唯一性。

需要注意的是，虽然我们能看到ValidationLayer中也有DynamicModule类成员，但是这和BackendModule里的不是一回事，我们如果读代码就会发现它们加载的是两个完全不一样的动态库。ValidationLayer管理的是单独的 $\text{LC}$ 校验层动态库。

DynamicModule对于动态库的管理比较特殊，不同于传统的内存分配/释放，它需要考虑的是动态库的加载和卸载。它的默认构造函数什么也不做，因此默认情况下句柄为空。这个句柄也可用于标识当前DynamicModule是否有效。

初始化一个DynamicModule通过调用其静态方法load来完成：

DynamicModule DynamicModule::load(std::string_view name) noexcept {
    std::lock_guard lock{dynamic_module_search_path_mutex()};
    Clock clock;
    auto &&paths = dynamic_module_search_paths();
    for (auto iter = paths.crbegin(); iter != paths.crend(); iter++) {
        auto p = iter->first / dynamic_module_name(name);
        if (auto handle = dynamic_module_load(p)) {
            LUISA_INFO(
                "Loaded dynamic module '{}' in {} ms.",
                to_string(p), clock.toc());
            return DynamicModule{handle};
        }
    }
    auto module_name = dynamic_module_name(name);
    if (auto handle = dynamic_module_load(module_name)) {
        LUISA_INFO(
            "Loaded dynamic module '{}' in {} ms.",
            module_name, clock.toc());
        return DynamicModule{handle};
    }
    return DynamicModule{nullptr};
}

在析构时，DynamicModule首先判断句柄是否为空，如果不为空，就调用dynamic_module_destroy来卸载动态库。

因此，动态库的管理逻辑非常直观，DynamicModule的核心就是对于动态库句柄的一个 $\text{RAII}$ 封装（在此基础上提供各种功能）。并且通过禁用拷贝，保证任意时刻只有一个DynamicModule是有效的，最后析构自然也能圆满完成动态库的卸载。

这里也是设计 $\text{C++ RAII}$ 类型的一个重要方式。有的对象管理了某种资源，不可被轻易拷贝，然而可以被移动，就必须给类型设计一个“空状态”。在将对象移动到新对象时，旧对象需要相应置为空状态，并在析构函数里进行判断，以避免重复释放资源。

3.2. Cuda Backend Resource Management

到了后端，所有的资源管理最终将会落到具体的DeviceInterface子类实现上。这里仍然以CUDADevice为例。


class CUDADevice final : public DeviceInterface {

    class ContextGuard {

    private:
        CUcontext _ctx;

    public:
        explicit ContextGuard(CUcontext ctx) noexcept : _ctx{ctx} {
            LUISA_CHECK_CUDA(cuCtxPushCurrent(_ctx));
        }
        ~ContextGuard() noexcept {
            CUcontext ctx = nullptr;
            LUISA_CHECK_CUDA(cuCtxPopCurrent(&ctx));
            if (ctx != _ctx) [[unlikely]] {
                LUISA_ERROR_WITH_LOCATION(
                    "Invalid CUDA context {} (expected {}).",
                    fmt::ptr(ctx), fmt::ptr(_ctx));
            }
        }
    };

public:
    /**
     * @brief Device handle of CUDA
     * 
     */
    class Handle {

    private:
        CUcontext _context{nullptr};
        CUdevice _device{0};
        uint32_t _device_index{};
        uint32_t _compute_capability{};
        uint32_t _driver_version{};
        CUuuid _uuid{};
        // will be lazily initialized
        mutable optix::DeviceContext _optix_context{nullptr};
        mutable spin_mutex _mutex{};
    };

private:
    Handle _handle;
    CUmodule _builtin_kernel_module{nullptr};
    CUfunction _accel_update_function{nullptr};
    CUfunction _instance_handle_update_function{nullptr};
    CUfunction _bindless_array_update_function{nullptr};
    luisa::unique_ptr<CUDACompiler> _compiler;
    luisa::unique_ptr<DefaultBinaryIO> _default_io;
    luisa::unique_ptr<CUDAEventManager> _event_manager;
    const BinaryIO *_io{nullptr};
    luisa::string _cudadevrt_library;

private:
    // extensions
    std::mutex _ext_mutex;
    luisa::unique_ptr<CUDADenoiserExt> _denoiser_ext;
    luisa::unique_ptr<CUDADStorageExt> _dstorage_ext;
}

CUDACompiler，CUDADenoiserExt，CUDADStorageExt是较为平凡的类型，它们的大部分成员作为“观察者”存在，即这些类成员能够引用其它的对象，但是不拥有它们，这些对象的生命周期总是不长于它们所观察的对象。因此我们没有很大必要分析它们的资源管理————因为这几个类压根就没有资源需要管理。

内嵌类Handle也是不用太在意的。它很显然是不管理资源的，它负责给外界一个访问CUDADevice具体信息的接口。

剔除这些类型以后，我们看Device到底负责管理了一些什么资源。

class CUDADevice final : public DeviceInterface {

    class ContextGuard {

    private:
        CUcontext _ctx;

    public:
        explicit ContextGuard(CUcontext ctx) noexcept : _ctx{ctx} {
            LUISA_CHECK_CUDA(cuCtxPushCurrent(_ctx));
        }
        ~ContextGuard() noexcept {
            CUcontext ctx = nullptr;
            LUISA_CHECK_CUDA(cuCtxPopCurrent(&ctx));
            if (ctx != _ctx) [[unlikely]] {
                LUISA_ERROR_WITH_LOCATION(
                    "Invalid CUDA context {} (expected {}).",
                    fmt::ptr(ctx), fmt::ptr(_ctx));
            }
        }
    };

private:
    CUmodule _builtin_kernel_module{nullptr};
    CUfunction _accel_update_function{nullptr};
    CUfunction _instance_handle_update_function{nullptr};
    CUfunction _bindless_array_update_function{nullptr};
    luisa::unique_ptr<DefaultBinaryIO> _default_io;
    luisa::unique_ptr<CUDAEventManager> _event_manager;
}

这个ContextGuard其实是 $\text{C++}$ 里一个很典型的设计，对于成对的操作（尤其是那种先“分配”后“释放”的，比如这里的CUcontext），使用一个Guard类来封装，构造时调用“分配”函数，析构时调用“释放”函数，这样通过任何方式离开作用域都会导致资源的释放，从而避免了资源泄漏。

4. Some Bad Smells?

在阅读 $\text{LC}$ 的源码时，我发现了一些可能不太好的地方，这里简单记录一下。

$\text{Builtin Sources}$

$\text{LC Cuda}$ 后端有一些内建的代码，但是这些代码的引入真是非常粗暴。。。

1
2
3

extern "C" const char luisa_cuda_builtin_cuda_builtin_kernels[5368];
extern "C" const char luisa_cuda_builtin_cuda_device_math[487941];
extern "C" const char luisa_cuda_builtin_cuda_device_resource[114333];

这是直接写死的。。。这是否过于粗暴了，万一以后需要修改内建代码，或者需要增添，甚至管理内建的代码呢？

$\text{Bad Naming}$

上面提到的create_shader和_create_shader，这两个方法名字差不多，但是做的事情差别也太大了。 $\text{LC}$ 整体代码注释不多，但是得益于清晰的代码逻辑，我们还是能够理解代码意图。然而有些命名实在是太差了。

$\text{Fixed Operators in DSL}$

$\text{DSL}$ 中，各种运算符的类型直接就写死在enum里了，后续处理也是直接switch case。坦白说，对于一个编程语言来说，可用的运算符确实应该没有太多的修改与扩展需求，然而还是不能彻底排除未来拓展的可能性。 $\text{LC}$ 这么写死了，万一后续需要增加新的运算符，那么就真的需要修改很多地方了。

5. Summary

作为一个比较新的项目， $\text{LuisaCompute}$ 的代码质量还没有经过太多时间的检验，然而，就我的阅读体验来看，代码和设计还是相当优雅的。在阅读过程中，我更是从中学习到了非常多有益的实战经验，包括动态库的管理，多后端的处理，内置代码的 $\text{AST}$ 构建等等。然而， $\text{LuisaCompute}$ 私以为也存在一些代码的坏味道，也希望它能够在未来的发展中不断完善。