关于面向GPU优化的并行化

我这里有段自适应采样的简要代码:

@ti.kernel
def pathtrace():
    for i, j in image_pixels:
        diff = diff_pixels[i, j]  # for self-adaptive sampling
        if diff > NOISE_THRESHOLD:
            sample(i, j)

对于并行此任务,随着 sample 执行次数增加,diff 逐渐减小。随着 kernel 调用次数越来越多,大多数子任务会提前结束掉,而 sample 是一个计算量非常大的函数。

这会造成线程发散吗?这种情况应该如何优化?能否增加一个前置 kernel 和一个 buffer,将 diff < NOISE_THRESHOLD(i,j) 映射到 diff > NOISE_THRESHOLD(i',j') ,这样就可以让这个kernel不需要 if 来决定是否 sample

  • 但是这样的话可能会增加(随机)访存?
  • 而且还要一个算法找到映射的方法,我想了一下这个是可以并行化的,而且计算复杂度应该比较小。
@ti.kernel
def build_map():
    x, y = find_index(diff_pixels, i, j, NOISE_THRESHOLD)
    # find_index 方法是适合并行的,它不一次性构建整个 map
    # 而是随着 kernel 逐渐调用而修改 map
    index_map[i, j] = (x, y)
@ti.kernel
def pathtrace():
    for i, j in index_map:
        x, y = index_map[i, j]
        sample(x, y)

我理解你的思路应该是弄一个buffer把要算的那些idx存起来,然后遍历这个index来跑。这里面发散的访存不用过于担心,因为30系GPU上有5MB左右的L2 data cache,4090甚至有个64MB的巨大的cache,这个kenrel局部性还不错,是个递增的index而不是随机走, fetch进cache的那些数据总能hit一些。这个角度来说访存会有问题也不致于太离谱。相反,如果一个warp中只有一个线程在狂奔,那实际上的ALU的利用率会非常低,综上我觉得你的思路是可行的。

1 个赞

感觉类似于“着色器重排序”这样的技术,taichi 有方便的实现吗?

这个我简单看了下是40系显卡的新API,支持dx12 / vulkan,似乎主要针对ray tracing shaders做的扩展,通用计算可能匹配起来有点儿麻烦。Vulkan的vendor extension暂时还没有,得等nv升级驱动。

你的代码用Vulkan后端跑性能咋样?另外现在用的是什么显卡呀?

纯软件算法层面应该也是可以重排序的,一次性把统一任务提交给GPU

没太看懂~你说的这个感觉更像是通过warp scheduling来掩盖latency?

纯软件这里得确定你能走哪个分支,这个是运行时信息,硬件如果没支持的话软件优化空间其实不是很大的样子