Metal - 线程本地的命名空间变量?

Metal - Namespace variable that is local to a thread?

我正在尝试在 Metal 中创建伪随机数生成器 (PRNG),类似于 thrust 的 RNG,每次您在线程中调用 RNG 给定一个特定的种子,它会产生一个不同的随机数,在这种情况下,它将是 thread_position_in_grid。我已经完美地设置了它,并且我现在使用我拥有的代码得到了一张漂亮的均匀随机图片。

但是,我的代码每个线程只工作一次。我想实现一个 next_rng() 函数,该函数 returns 一个新的 rng 使用最后一个结果作为种子。然而,我在存储最后一个结果时遇到了问题,因为我的 RNG 是一个名称 space,而且我只能将最后一个结果存储在 constant 地址 space 中,这是不可更新的。有什么办法绕过 this/restructure 我的代码来避免这种情况吗?

我能在网上找到的最好的帮助是 this SO post 不幸的是,它对我不起作用,因为我不能像他们在他们的解决方案中那样在函数中将变量声明为静态变量。

如果有帮助,我愿意以任何方式重构我的代码。我能不能把它变成静态的 class 或 class 之类的?没有线索,但我只想要一个 PRNG

我的代码(简化版)基本上与 post 的代码具有相同的结构:

namespace metal_rng {

    thread unsigned* last_seed() {
        thread uint last_seed = 1; // Doesn't work because last-seed falls out of scope after the function quits

        // The following comment line doesn't even compile:
        // thread static uint last_seed = 1;

        return &last_seed;
    }

    unsigned TausStep(const unsigned z, const int s1, const int s2, const int s3, const unsigned M)
    {
        unsigned b=(((z << s1) ^ z) >> s2);
        return (((z & M) << s3) ^ b);
    }

    device float rng(const int initial_seed) {
        int seed = initial_seed * 1099087573UL;

        unsigned hashed = (1664525*(*last_seed()) + 1013904223UL);
        *last_seed() = hashed
        return (hashed) * 2.3283064365387e-10;
    }

    device float next_rng() {
          if (*last_seed() == 0) {
              return 0.0;
          } else {
              unsigned hashed = (1664525*(*last_seed()) + 1013904223UL);
              return (hashed) * 2.3283064365387e-10;
          }
    }
}

编辑:

框架在这里完成并整齐地打包:Link

我最终使用了@KenThomases 的建议,我刚刚为 RNG 创建了一个 class。我打包了代码,调用了RNGLokiHere is the link to the repo如果以后有人想用的话