如何在 OpenCL 中获取“随机”数

Posted

技术标签:

【中文标题】如何在 OpenCL 中获取“随机”数【英文标题】:How to get a "random" number in OpenCL 【发布时间】:2012-04-12 07:38:57 【问题描述】:

我希望在 OpenCL 中获得一个随机数。它不必是真正的随机甚至是随机的。只是简单快捷的事情。

我看到 OpenCL 中有大量真正的随机并行化花式裤子随机算法,就像成千上万行一样。我不需要那样的东西。一个简单的“random()”就可以了,即使很容易看到其中的模式。

我看到有噪声功能?有什么简单的方法可以使用它来获取随机数?

【问题讨论】:

你可以将一个包含一些随机数的数组传递给你的内核,这不是最简单的吗? 是的,但是你在 CPU->GPU 之间来回切换。使用 OpenCL 的全部意义在于保持在 GPU 上。 这取决于应用程序和某些类型的任务(即使在同一个问题中)CPU 可能仍然不错。取决于你有多少时间,在 CPU 上快速生成 rands 用于开发。 其他人注意:这里的许多答案大多已经超越了 OP 希望避免的简单的非平行花式裤子随机算法 - 但是现在存在库来生成随机数所以这样当您面临在 OpenCL 中需要随机数的常见问题时,值得考虑将它们包括在内,或者至少看看他们是如何处理它的。有关详细信息,请参阅下面的帖子。 【参考方案1】:

这是我的 OpenCL 浮点伪随机噪声版本,使用三角函数

//noise values in range if 0.0 to 1.0
static float noise3D(float x, float y, float z) 
    float ptr = 0.0f;
    return fract(sin(x*112.9898f + y*179.233f + z*237.212f) * 43758.5453f, &ptr);


__kernel void fillRandom(float seed, __global float* buffer, int length) 
    int gi = get_global_id(0);
    float fgi = float(gi)/length;
    buffer[gi] = noise3D(fgi, 0.0f, seed);

您可以通过将noise3D 归一化索引坐标作为第一个参数并将随机种子(例如在CPU上生成)作为最后一个参数来生成一维或二维噪声。

以下是使用此内核和不同种子生成的一些噪声图片:

【讨论】:

【参考方案2】:

我目前正在实施实时路径跟踪器。您可能已经知道路径追踪需要许多随机数。 在 GPU 上生成随机数之前,我只是在 CPU 上生成它们(使用 rand(),这很糟糕)并将它们传递给 GPU。 这很快成为了瓶颈。 现在,我正在使用 Park-Miller 伪随机数生成器 (PRNG) 在 GPU 上生成随机数。 实现起来极其简单,效果非常好。 我采集了数千个样本(在 0.0 到 1.0 的范围内)并将它们平均在一起。 结果值非常接近 0.5(这是您所期望的)。在不同的运行之间,与 0.5 的差异约为 0.002。因此它具有非常均匀的分布。 这是描述该算法的论文:http://www.cems.uwe.ac.uk/~irjohnso/coursenotes/ufeen8-15-m/p1192-parkmiller.pdf 这是一篇关于为 CUDA 优化的上述算法的论文(可以很容易地移植到 OpenCL):http://www0.cs.ucl.ac.uk/staff/ucacbbl/ftp/papers/langdon_2009_CIGPU.pdf 这是我如何使用它的示例:

int rand(int* seed) // 1 <= *seed < m

    int const a = 16807; //ie 7**5
    int const m = 2147483647; //ie 2**31-1

    *seed = (long(*seed * a))%m;
    return(*seed);


kernel random_number_kernel(global int* seed_memory)

    int global_id = get_global_id(1) * get_global_size(0) + get_global_id(0); // Get the global id in 1D.

    // Since the Park-Miller PRNG generates a SEQUENCE of random numbers
    // we have to keep track of the previous random number, because the next
    // random number will be generated using the previous one.
    int seed = seed_memory[global_id];

    int random_number = rand(&seed); // Generate the next random number in the sequence.

    seed_memory[global_id] = *seed; // Save the seed for the next time this kernel gets enqueued.

代码仅作为示例。我没有测试过。 在内核第一次执行之前,数组“seed_memory”只被 rand() 填充一次。之后,所有随机数的生成都在 GPU 上进行。我认为也可以简单地使用内核 id 而不是使用 rand() 初始化数组。

【讨论】:

感谢您的链接,我想指出,该论文实际上为 GPU 提出了不同的实现。您发布的只是 Park&Miller 的参考。该论文附带 GPU 的源代码,示例和测试用例的源代码在摘要中链接。 是的,实现有点不同。但结果是一样的。【参考方案3】:

过去几天我一直在解决这个“非随机”问题,我想出了三种不同的方法:

    Xorshift - 我基于这个创建了生成器。您所要做的就是为整个内核提供一个uint2 编号(种子),每个工作项都会计算自己的 rand 编号

    // 'randoms' is uint2 passed to kernel
    uint seed = randoms.x + globalID;
    uint t = seed ^ (seed << 11);  
    uint result = randoms.y ^ (randoms.y >> 19) ^ (t ^ (t >> 8));
    

    Java random - 我使用来自.next(int bits) 方法的代码来生成随机数。这次您必须提供一个ulong 号码作为种子。

    // 'randoms' is ulong passed to kernel
    ulong seed = randoms + globalID;
    seed = (seed * 0x5DEECE66DL + 0xBL) & ((1L << 48) - 1);
    uint result = seed >> 16;
    

    只需在 CPU 上生成所有内容,然后在一个大缓冲区中将其传递给内核。

我在计算图中最小支配集的进化算法中测试了所有三种方法(生成器)。

我喜欢第一个生成的数字,但看起来我的进化算法不喜欢。

第二个生成器生成的数字有一些可见的模式,但我的进化算法无论如何都喜欢这种方式,整个事情的运行速度比第一个生成器快一点。

但第三种方法表明,只提供来自主机 (cpu) 的所有数字绝对没问题。首先,我认为生成(在我的情况下)1536 个 int32 数字并在每个内核调用中将它们传递给 GPU 会太昂贵(计算和传输到 GPU)。但事实证明,它和我之前的尝试一样快。 CPU 负载保持在 5% 以下。

顺便说一句,我也尝试了MWC64X Random,但在我安装了新的 GPU 驱动程序后,函数mul_hi 开始导致构建失败(甚至整个 AMD Kernel Analyer 都崩溃了)。

【讨论】:

应该有什么问题吧,因为这是多线程使用的?更新种子时没有竞争条件吗? ***.com/questions/71155563/… 自从我使用 OpenCL 以来已经 9 年了 :),但如果我没记错的话,每个工人都有一个私有内存,其中保存着所有像这样的原始值 - 所以不会发生竞争条件。然后是整个集群的内存,您需要使用偏移量来读取或写入,然后是全局内存(大内存),您再次需要使用偏移量。请参阅OpenCL Memory Model 了解更多信息。 啊,那么在这个实现中,随机状态是在本地工作空间中吗?我之所以问,是因为我链接的问题是我使用与全局工作区相关的随机逻辑.. @DavidTóth 看着你的问题,我真的不知道。我不再是这方面的专家了,我 9 年前就用过这个(!!!)。无论如何,一般来说,本地创建的变量(如数字)对于每个工人都是私有的。但是,如果要将结果写入全局内存,则每个工作人员都必须使用该内存中的不同位置 - 这就是使用带工作人员 id 的偏移量的原因。如果内存在工作人员之间共享,那么如果任何工作人员都可以修改它,那么总是存在竞争条件。【参考方案4】:

您不能在内核中生成随机数,最好的选择是在主机(CPU)中生成随机数,然后通过缓冲区将其传输到 GPU 并在内核中使用。

【讨论】:

您可以使用任何随机化算法生成随机数,但种子需要包含线程的 global_id,因此每个都是真正随机的 OpenCL 同时使用 GPU 和 CPU!【参考方案5】:

以下是java.util.Random类根据doc使用的算法:

(seed * 0x5DEECE66DL + 0xBL) & ((1L << 48) - 1)

请参阅文档以了解其各种实现。将工作人员的 id 传递给种子并循环几次应该会产生不错的随机性

或者另一种方法是发生一些相当确定会溢出的随机操作:

 long rand= yid*xid*as_float(xid-yid*xid);
 rand*=rand<<32^rand<<16|rand;
 rand*=rand+as_double(rand);

xid=get_global_id(0);yid= get_global_id(1);

【讨论】:

上述(java 随机)在seed[n] 和seed[n-1] 之间具有非常高的相关性。我是在构建骰子模拟器时发现的。取上面的输出,然后做模 6。你永远不会得到两个相邻的 5!【参考方案6】:

OpenCL 似乎不提供这样的功能。但是,some people have done some research 并提供 BSD 许可代码以在 GPU 上生成良好的随机数。

【讨论】:

【参考方案7】:

我遇到了同样的问题。 www.thesalmons.org/john/random123/papers/random123sc11.pdf

您可以在此处找到文档。 http://www.thesalmons.org/john/random123/releases/latest/docs/index.html

您可以在此处下载库: http://www.deshawresearch.com/resources_random123.html

【讨论】:

【参考方案8】:

GPU 没有良好的随机性来源,但可以通过使用来自主机的随机种子为内核播种来轻松克服这一问题。之后,您只需要一个可以处理大量并发线程的算法。

此链接描述了使用 OpenCL 的 Mersenne Twister 实现:Parallel Mersenne Twister。您还可以在 NVIDIA SDK 中找到实现。

【讨论】:

【参考方案9】:

为什么不呢?您可以只编写一个生成随机数的内核,这需要更多内核调用并最终将随机数作为参数传递给需要它们的其他内核

【讨论】:

以上是关于如何在 OpenCL 中获取“随机”数的主要内容,如果未能解决你的问题,请参考以下文章

从 GPU 获取 OpenCL 程序代码

如何使用 PyOpenCL 将带有数组和变量的 C 结构传递给 OpenCL 内核

如何获取 CUDA 内核的汇编代码?

如何在 OpenCL 内核中更新 OpenCL-OpenGL 共享缓冲区数据?

如何在 OpenCL 中使用本地内存?

如何在 OpenCL 中使用固定内存/映射内存