随机数生成
CUDA 内置目标弃用通知
Numba 内置的 CUDA 目标已弃用,后续开发已移至 NVIDIA numba-cuda 包。请参阅内置 CUDA 目标弃用和维护状态。
Numba 提供了一种可在 GPU 上执行的随机数生成算法。然而,由于 NVIDIA 实现 cuRAND 的技术问题,Numba 的 GPU 随机数生成器并非基于 cuRAND。相反,Numba 的 GPU RNG 是 xoroshiro128+ 算法 的实现。xoroshiro128+ 算法的周期为 2**128 - 1
,这比 cuRAND 默认使用的 XORWOW 算法的周期短,但 xoroshiro128+ 仍然通过了随机数生成器质量的 BigCrush 测试。
在 GPU 上使用任何 RNG 时,务必确保每个线程都有自己的 RNG 状态,并且它们已初始化以生成不重叠的序列。numba.cuda.random
模块提供了实现此功能的主机函数,以及获取均匀分布或正态分布随机数的 CUDA 设备函数。
注意
Numba(与 cuRAND 类似)使用 Box-Muller 变换 <https://en.wikipedia.org/wiki/Box%E2%80%93Muller_transform> 从均匀生成器生成正态分布的随机数。然而,Box-Muller 生成成对的随机数,而当前实现只返回其中一个。因此,生成正态分布值的速度是均匀分布值的一半。
- numba.cuda.random.create_xoroshiro128p_states(n, seed, subsequence_start=0, stream=0)
返回一个为 n 个随机数生成器初始化后的新设备数组。
这将初始化 RNG 状态,使数组中的每个状态在主序列中相互间隔 2**64 步,形成对应的子序列。因此,只要没有 CUDA 线程请求超过 2**64 个随机数,此函数生成的所有 RNG 状态都保证是独立的。
subsequence_start
参数可用于将第一个 RNG 状态前进 2**64 步的倍数。- 参数
n (int) – 要创建的 RNG 状态数
seed (uint64) – 生成器列表的起始种子
subsequence_start (uint64) –
stream (CUDA stream) – 运行初始化内核的流
- numba.cuda.random.init_xoroshiro128p_states(states, seed, subsequence_start=0, stream=0)
在 GPU 上为并行生成器初始化 RNG 状态。
这将初始化 RNG 状态,使数组中的每个状态在主序列中相互间隔 2**64 步,形成对应的子序列。因此,只要没有 CUDA 线程请求超过 2**64 个随机数,此函数生成的所有 RNG 状态都保证是独立的。
subsequence_start
参数可用于将第一个 RNG 状态前进 2**64 步的倍数。- 参数
states (1D DeviceNDArray, dtype=xoroshiro128p_dtype) – RNG 状态数组
seed (uint64) – 生成器列表的起始种子
- numba.cuda.random.xoroshiro128p_normal_float32(states, index)
返回一个正态分布的 float32 值并推进
states[index]
。返回值是使用 Box-Muller 变换从均值=0、标准差=1 的高斯分布中提取的。这会使 RNG 序列前进两步。
- 参数
states (1D array, dtype=xoroshiro128p_dtype) – RNG 状态数组
index (int64) – 要更新的状态偏移量
- 返回类型
float32
- numba.cuda.random.xoroshiro128p_normal_float64(states, index)
返回一个正态分布的 float32 值并推进
states[index]
。返回值是使用 Box-Muller 变换从均值=0、标准差=1 的高斯分布中提取的。这会使 RNG 序列前进两步。
- 参数
states (1D array, dtype=xoroshiro128p_dtype) – RNG 状态数组
index (int64) – 要更新的状态偏移量
- 返回类型
float64
- numba.cuda.random.xoroshiro128p_uniform_float32(states, index)
返回一个 [0.0, 1.0) 范围内的 float32 值并推进
states[index]
。- 参数
states (1D array, dtype=xoroshiro128p_dtype) – RNG 状态数组
index (int64) – 要更新的状态偏移量
- 返回类型
float32
- numba.cuda.random.xoroshiro128p_uniform_float64(states, index)
返回一个 [0.0, 1.0) 范围内的 float64 值并推进
states[index]
。- 参数
states (1D array, dtype=xoroshiro128p_dtype) – RNG 状态数组
index (int64) – 要更新的状态偏移量
- 返回类型
float64
一个简单示例
这是一个使用随机数生成器的示例程序。
from __future__ import print_function, absolute_import
from numba import cuda
from numba.cuda.random import create_xoroshiro128p_states, xoroshiro128p_uniform_float32
import numpy as np
@cuda.jit
def compute_pi(rng_states, iterations, out):
"""Find the maximum value in values and store in result[0]"""
thread_id = cuda.grid(1)
# Compute pi by drawing random (x, y) points and finding what
# fraction lie inside a unit circle
inside = 0
for i in range(iterations):
x = xoroshiro128p_uniform_float32(rng_states, thread_id)
y = xoroshiro128p_uniform_float32(rng_states, thread_id)
if x**2 + y**2 <= 1.0:
inside += 1
out[thread_id] = 4.0 * inside / iterations
threads_per_block = 64
blocks = 24
rng_states = create_xoroshiro128p_states(threads_per_block * blocks, seed=1)
out = np.zeros(threads_per_block * blocks, dtype=np.float32)
compute_pi[blocks, threads_per_block](rng_states, 10000, out)
print('pi:', out.mean())
管理 RNG 状态大小和使用 3D 网格的示例
RNG 状态的数量随使用 RNG 的线程数量而变化,因此通常最好结合 RNG 使用跨步循环,以保持状态大小可管理。
在以下示例中,初始化一个大型 3D 数组时使用随机数,如果每个输出元素使用一个线程,将导致 453,617,100 个 RNG 状态。这将花费很长时间进行初始化,并且 GPU 利用率很低。相反,它使用一个固定大小的 3D 网格,总共有 2,097,152 ((16 ** 3) * (8 ** 3)
) 个线程跨步处理输出数组。3D 线程索引 startx
、starty
和 startz
被线性化为一个 1D 索引 tid
,用于索引到 2,097,152 个 RNG 状态中。
test_ex_3d_grid of ``numba/cuda/tests/doc_example/test_random.py
1from numba import cuda
2from numba.cuda.random import (create_xoroshiro128p_states,
3 xoroshiro128p_uniform_float32)
4import numpy as np
5
6@cuda.jit
7def random_3d(arr, rng_states):
8 # Per-dimension thread indices and strides
9 startx, starty, startz = cuda.grid(3)
10 stridex, stridey, stridez = cuda.gridsize(3)
11
12 # Linearized thread index
13 tid = (startz * stridey * stridex) + (starty * stridex) + startx
14
15 # Use strided loops over the array to assign a random value to each entry
16 for i in range(startz, arr.shape[0], stridez):
17 for j in range(starty, arr.shape[1], stridey):
18 for k in range(startx, arr.shape[2], stridex):
19 arr[i, j, k] = xoroshiro128p_uniform_float32(rng_states, tid)
20
21# Array dimensions
22X, Y, Z = 701, 900, 719
23
24# Block and grid dimensions
25bx, by, bz = 8, 8, 8
26gx, gy, gz = 16, 16, 16
27
28# Total number of threads
29nthreads = bx * by * bz * gx * gy * gz
30
31# Initialize a state for each thread
32rng_states = create_xoroshiro128p_states(nthreads, seed=1)
33
34# Generate random numbers
35arr = cuda.device_array((X, Y, Z), dtype=np.float32)
36random_3d[(gx, gy, gz), (bx, by, bz)](arr, rng_states)