在openCL中的多个设备中维护RNG状态的最佳方式

Question

所以我正试图为openCL使用这个自定义RNG库：http://cas.ee.ic.ac.uk/people/dt10/research/rngs-gpu-mwc64x.html

库定义了一个状态结构：

//! Represents the state of a particular generator
typedef struct{ uint x; uint c; } mwc64x_state_t;

为了生成一个随机uint，您可以将状态传递给以下函数：

uint MWC64X_NextUint(mwc64x_state_t *s)

它会更新状态，这样当你再次将其传递到函数中时，序列中的下一个"随机"数就会生成。

对于我正在创建的项目，我不仅需要能够在不同的工作组/项目中生成随机数，还需要同时在多个设备上生成随机数。我很难找到最好的设计方法。比如，我应该为每个设备/命令队列创建1个mwc64x_state_t对象，并将该状态作为全局变量传入吗？或者可以同时为所有设备创建一个状态对象吗？或者，我甚至不将其作为全局变量传入，并在每个内核函数中本地声明一个新状态？

该库还具有以下功能：

void MWC64X_SeedStreams(mwc64x_state_t *s, ulong baseOffset, ulong perStreamOffset)

据推测，这应该将RNG分成多个"流"，但将其包含在我的内核中会使其速度慢得令人难以置信。例如，如果我做一些非常简单的事情，比如：

__kernel void myKernel()
{
    mwc64x_state_t rng;
    MWC64X_SeedStreams(&rng, 0, 10000);
}

然后内核调用的速度变慢了大约40倍。

该库确实提供了一些源代码作为示例用法，但示例代码有点有限，似乎没有那么大帮助。

因此，如果有人熟悉openCL中的RNG，或者你以前使用过这个特定的库，我将非常感谢你的建议。

Answer 1

MWC64X_SeedStreams函数确实相对较慢，至少相比之下到MWC64X_NextUint调用，但对于尝试将大型全局流拆分为多个子流，这些子流可用于平行的假设您将多次致电NextUint在内核中（例如一百个或更多），但SeedStreams仅位于顶部。

这是EstimatePi示例的注释版本，随附使用库（mwc64x/test/estimate_pi.cpp和mwc64x/trest/test_mwc64x.cl）。

__kernel void EstimatePi(ulong n, ulong baseOffset, __global ulong *acc)
{
    // One RNG state per work-item
    mwc64x_state_t rng;
    // This calculates the number of samples that each work-item uses
    ulong samplesPerStream=n/get_global_size(0);
    // And then skip each work-item to their part of the stream, which
    // will from stream offset:
    //   baseOffset+2*samplesPerStream*get_global_id(0)
    // up to (but not including):
    //   baseOffset+2*samplesPerStream*(get_global_id(0)+1)
    //
    MWC64X_SeedStreams(&rng, baseOffset, 2*samplesPerStream);

    // Now use the numbers
    uint count=0;
    for(uint i=0;i<samplesPerStream;i++){
        ulong x=MWC64X_NextUint(&rng);
        ulong y=MWC64X_NextUint(&rng);
        ulong x2=x*x;
        ulong y2=y*y;
        if(x2+y2 >= x2)
            count++;
    }
    acc[get_global_id(0)] = count;
}

因此，目的是n应该很大，并随着数字的增长而增长工作项的数量增加，因此samplesPerStream仍然存在一百个或更多。

如果您希望在多个设备上使用多个内核，那么需要向流分割添加另一级别的层次结构，例如，如果你有：

• K：设备数量（可能在并行机上）
• W：每个设备的工作项数
• C：每个工作项调用NextUint的次数

您最终会得到N=KWC所有对NextUint的调用总数工作项。如果您的设备被标识为k=0..（k-1），然后在每个内核中，您将执行以下操作：

MWC64X_SeedStreams(&rng, W*C*k, C);

那么流中的索引将是：

[0             .. N ) : Parts of stream used across all devices
[k*(W*C)       .. (k+1)*(W*C) )    : Used within device k
[k*(W*C)+(i*C) .. (k*W*C)+(i+1)*C ) : Used by work-item i in device k.

如果每个内核使用的样本少于C，则可以如有必要，估计过高。

（我是图书馆的作者）。