id="chapter-1-lc-programming">Chapter

programming

内线程必须同步执行。

分支会导致部分线程被掩码（Masked），执行串行化。

🔥

[[unroll]]、分支预测或通过逻辑运算消除

个线程。

这些线程在硬件上是“步调一致”的。

Branch

switch-case，会对性能产生什么影响？

Warp

Serialization（串行化）。

虽然逻辑上线程是并行的，但物理上

个线程会轮流执行每一个分支逻辑，有效计算吞吐量可能下降到原来的

技巧：尽可能

color="blue">使用

lerp、clamp

或逻辑位运算代替

id="12-延迟掩盖-latency-hiding-与-occupancy">1.2

延迟掩盖

慢在“记性不好”（显存延迟高），快在“人多”（并行度高）。

申请了太多临时变量，每个线程分配到的寄存器就多，能并行的线程就少。

里的寄存器使用，性能反而没提升，甚至下降了？

Occupancy

已经足够掩盖内存延迟，再往上提高也不会带来收益。

反而，如果为了降低寄存器使用而增加了额外的指令（如反复计算），总执行时间反而会增加。

color="blue">当寄存器实在不够用时，编译器会将变量存在“本地内存”中。

虽然叫本地，其实是在显存里，

color="blue">读写极慢！

这是性能优化的红线。

对应硬件的什么？

它为什么比全局变量快？

color="blue">片上存储（On-chip），距离

ALU

color="blue">不走全局内存总线。

(平铺优化)：在处理大规模数据（如多模态模型中的矩阵乘法）时，先将数据

color="blue">从

Memory，在内部完成计算再写回。

这能把访存延迟降低一个数量级，减少HBM

后：

color="blue">数据被缓存在片上，对显存的读取次数减少了

BlockSize

color="blue">LuisaCompute简化了编程，运行时自动选择最优

block

考点预警：对于CUDA，需要自行配置block_size

Warp：硬件执行的灵魂。

Block

内部每

id="b-存储线材质与位置-sram-vs-dram">B.

存储线：材质与位置

特点：极快，但贵，占用面积大。

物理实现：Register、Shared

Memory、L1

特点：容量大，相对便宜，但慢（延迟高）。

物理实现：Global

color="blue">芯片外面（Off-chip）。

GDDR

id="c-逻辑空间global-memory-的真面目">C.

逻辑空间：Global

┌─────────────────────────────────────────────────────────┐
GPU
┌─────────────────────────────────────────────────┐
Register
└─────────────────────────────────────────────────┘
┌─────────────────────────────────────────────────┐
Shared
└─────────────────────────────────────────────────┘
┌─────────────────────────────────────────────────┐
Cache（二级缓存）
└─────────────────────────────────────────────────┘
┌─────────────────────────────────────────────────┐
HBM
└─────────────────────────────────────────────────┘
┌─────────────────────────────────────────────────┐
Host
└─────────────────────────────────────────────────┘
└─────────────────────────────────────────────────────────┘

[32][33]，错开地址。

atomicAdd、__syncthreads()、__threadfence()。

🔥

warp_reduce，最后由一个线程写回全局。

id="21-合并访问-coalesced-access">2.1

合并访问

Buffer<Particle>)：[x,

b],

sizeof(Particle)）。

这会导致大量带宽浪费在不需要的数据（y,

pos_y)：[x,

x...],

color="blue">地址是完美连续的，只需极少的事务就能填满

Buffer<float>

Buffer<float4>。

对齐陷阱：如果你非要用

16

float4）。

float3

字节，这在某些架构上会导致跨

id="22-shared-memory-与-bank-conflict">2.2

Shared

这是面试中最高频的“手撕代码优化点”，特别是涉及矩阵运算或卷积时。

color="red">Bank（通常每

字节一个

color="blue">理想情况：Warp

个线程同时访问

B。

硬件必须将这两个请求串行化（先给

0，再给

32\)。

读取是按行读（连续，无冲突）。

但转置后写回是按列写，Shared

Memory

Data[i+1][0]。

如果宽度正好是

恰好都在

color="blue">32和33互质）

Memory定义数组维度时增加一个偏移量。

32和33互质，取模运算能直接避免Bank

Conflict。

compile_kernel(.,

KernelOption::DetectBankConflict)。

但实际开发中一般不用

compile_kernel，用异步编译

compiler.compile<>()

更灵活，然后使用Nsight

atomic_fetch_add

Cache）。

它会阻塞内存控制器，直到“读-改-写”完成。

10,000

local_hist[256]。

局部统计：Block

全局汇总：Block

结束后，再将

id="三luisacompute-核心机制">三、LuisaCompute

核心机制

color="red">eDSL)。

color="red">AST（抽象语法树）。

JIT

float？

Var<float>（或

Expr<float>）是因为我们要

color="blue">捕获的是变量的计算逻辑而非当前的数值。

只有这样，LC

Kernel。

必须在初始化阶段完成编译并缓存。

C++

Binding：Shader

Slot（如

t1），每画一个物体都要重新切换绑定。

这在

Bindless：将所有资源（Buffer,

Texture）放在一个全局的大数组（Heap）里。

Shader

只持有一个索引（int

里，我们可以通过一个动态索引直接访问数组中的任意资源：images.read2d(index,

痛点：解决了大规模场景（如

AIGC

动态更新：Bindless

允许在不重新记录

上的索引表。

这对于多模态模型中频繁变动的数据采样非常高效。

device.compile_kernel([&](ImageFloat

gBuffer,

shadowMap.sample(shadowCoord);//

...

device.create_bindless_array(MAX_TEXTURES);//

注册资源到数组

device.compile_kernel([&](BindlessVar

bindless,

bindless.tex2d(texId).sample(uv);//

nextTex

bindless.tex2d(nextTex).sample(uv);

});//

uv_buffer).dispatch(pixel_count);

GetNativeResource()

RDG

import_buffer。

Non_Pixel_Shader_Resource；计算完回传给

生命周期对齐：Luisa

Context

static_cast<ID3D12DynamicRHI*>(GDynamicRHI);ID3D12Device*

NativeDevice

DX12RHI->RHIGetNativeDevice();IDXGIAdapter1*

NativeAdapter

DX12RHI->RHIGetAdapter();static

context{

};luisa::compute::DirectXDeviceConfig

NativeAdapter;dx_config.headless

true;static

device.create_stream(luisa::compute::StreamTag::COMPUTE);

color="blue">ID3D12Device

管理接口"，luisa::compute::Device

Luisa

做的"套壳"，保存前者的指针，以便Luisa自动进行RAII管理。

color="blue">Adapter

指代物理上的显卡硬件。

有的电脑可能有两块显卡：一个是

Intel

：把一堆任务打包，放在一条传送带（Stream）上，送进显卡排队执行。

视频演示：https://www.bilibili.com/video/BV1KrzrYeEDw/

Radix

Sort的流程：基数为2时，uint32需要进行32轮，

color="blue">实际工程中可能取Radix为16，只需要处理8轮。

Sum)：对桶进行scan。

确定元素在全局数组中的起始存放位置。

Heuristic)：一种更优的分割策略，但

GPU

时，最大的性能杀手是什么？

如何缓解？

杀手：分支分歧

(Branch

显存跳跃访问。

相邻的光线（同一个

缓解：光线重排序

(Ray

Sorting/Reordering)。

在着色前，根据光线的方向或击中点的材质

对光线进行排序，让同一个

(Autodiff)：LuisaCompute

支持类似

(Reparameterization)：在光栅化或光追中，几何边缘的导数通常是不连续的。

需要了解如何处理这种不连续性（如

Edge

的顶点位置，使其渲染结果接近目标图片，LuisaCompute

Autodiff

Pass：记录计算图（Tape）或直接生成伴随代码。

反向

(__shfl_sync)、Reduce。

Reduce

warp_shuffle_down(var,

delta):

warp_shuffle_xor(var,

laneMask):

color="blue">经典的面试题：“如何实现一个高性能的并行求和？

”.

Warp

device.compile_kernel([&](BufferFloat

input,

device.compile_kernel([&](Buffer<uint>

out_count

__popc(warp_ballot(is_leader));}

});

Zero-Copy：中间数据（Per-thread

Ops：Warp

减少指令数：省去了大量的

Load/Store

算完数据，不写回显存，直接在寄存器里交给

Reduce

id="52-warp-specialization">5.2

Warp

它们不需要分配用于矩阵乘法的寄存器。

它们不需要分配用于全局内存地址计算的寄存器。

数据好了mbarrier.arrive(transaction_count);

else

(Overlap)wmma_gemm(s_mem_ptr[step

2]);

Isolation)：这是最关键的。

传统模式下，一个线程既要存

Global

Accumulator，大家都能“轻装上阵”，可以用更少的资源跑更多的

color="blue">把数据锁在寄存器里”，这是

LuisaCompute

总是好的吗？

什么时候融合反而会变慢？

Spilling)：如果融合的操作太多，导致单个

Kernel

并行度不匹配：如果前面的操作是

Reduce

Element-wise，它们的线程映射逻辑不同（一个需要

Block

小算子地狱：如果是大规模训练，每个

Kernel

(实例化)：驱动将这一整张图编译成一个“超级指令包”。

Launch

cudaGraphLaunch。

GPU

(GSP)

ExecuteIndirect。

预先录制好

Command

LuisaCompute：

color="blue">LC

本身就是一种录制机制。

你可以构建一个包含成百上千个

Dispatch

stream.submit(list)。

虽然底层实现依赖于后端，但在

DX12

id="chapter-2-profilling">Chapter

Profilling

Unit)，比标准函数快得多，但精度略低。

atan2）计算太慢，且输入范围有限，

Buffer

合并访问：确保内存事务是对齐的（

向量化读写：

color="blue">在

float。

这样可以触发硬件的

指令，一次读取

位宽压缩：如果精度允许，一个uint的32bit存放多个数据。

访存量直接减半，带宽利用率翻倍。