CUDA 的随机数算法 API
一、具体使用场景
如下是是在 dropout 优化中手写的 uniform_random 的 Kernel：
#include <cuda_runtime.h>
#include <curand_kernel.h>
__device__ inline float cinn_nvgpu_uniform_random_fp32(int seed){
curandStatePhilox4_32_10_t state;
int idx = threadIdx.x + blockIdx.x * blockDim.x;
curand_init(seed, idx, 1, &state);
return curand_uniform(&state);
}
二、API 解析
我们首先来看 curand_init 函数的签名和语义：
__device__ void
curand_init(unsigned long long seed,
unsigned long long subsequence,
unsigned long long offset,
curandStatePhilox4_32_10_t *state)
给定相同的seed、sequence、offset 参数下，curand_init 会保证产生相同的其实状态 state。另外此函数会在调用 2^67 ⋅ sequence + offset次 cu_rand API 之后「重置」为起始状态。关于 sequence 和 offset 的如何生效的机制，参考 StackOverflow。
注意：在 seed 相同、sequence 不同时，一般不会产生有统计学上关联的结果。原文：Sequences generated with the same seed and different sequence numbers will not have statistically correlated values.
另外在CUDA 并行产生随机数实践上，也有一些经验上的建议：
- 若保证高质量的伪随机数，建议使用不同的 seed
- 若是并行在一个「实验」里，建议指定不同的sequence参数，且最好「单调递增」
- 若果线程里的config都是一样，即 state 一样，则可以把「随机状态量」放到 global memory里，以减少 setup 开销
参考原文：For the highest quality parallel pseudorandom number generation, each experiment should be assigned a unique seed. Within an experiment, each thread of computation should be assigned a unique sequence number. If an experiment spans multiple kernel launches, it is recommended that threads between kernel launches be given the same seed, and sequence numbers be assigned in a monotonically increasing way. If the same configuration of threads is launched, random state can be preserved in global memory between launches to avoid state setup time.
然后我们看下Nvidia 主要提供了哪些常用的随机数生成API：
__device__ float
curand_uniform (curandState_t *state); // <---- It may return from 0.0 to 1.0, where 1.0 is included and 0.0 is excluded.
__device__ float
curand_normal (curandState_t *state); // <----- returns a single normally distributed float with mean 0.0 and standard deviation 1.0.
__device__ float
curand_log_normal (curandState_t *state, float mean, float stddev); // <----- returns a single log-normally distributed float based on a normal distribution with the given mean and standard deviation.
// 如下是上述 3 个API 的 double 版本
__device__ double
curand_uniform_double (curandState_t *state)；
__device__ double
curand_normal_double (curandState_t *state)；
__device__ double
curand_log_normal_double (curandState_t *state, double mean, double stddev);
上面的 device API 在每次调用时，只会生成一个 float/double 的随机数。Nvidia 同样提供了一次可以生成 2个或4个 device API：
__device__ uint4
curand4 (curandStatePhilox4_32_10_t *state)；
__device__ float4
curand_uniform4 (curandStatePhilox4_32_10_t *state)；
__device__ float4
curand_normal4 (curandStatePhilox4_32_10_t *state)；
__device__ float4
curand_log_normal4 (curandStatePhilox4_32_10_t *state, float mean, float stddev)；
从上面的函数接口以及 Nvidia 的文档来看，在初始化某种类型的 state 状态量后，每次调用类似 curand() 的 API 后，state 都会自动进行 offset 偏移。
因此，Nvidia 官网上也提供了单独对 state 进行位移的 API，其效果等价于调用多次无返回值的 curand() API，且性能更好：
__device__ void
skipahead(unsigned long long n, curandState_t *state); // <----- == calls n*curand()
__device__ void
skipahead_sequence(unsigned long long n, curandState_t *state); // <----- == calls n*2^67 curand()
三、性能分析
Nvidia 的官网明确指出了存在的性能问题，给开发者实现高性能 Kernel 提供了充分的经验指导：
- curand_init()要比curand()和curand_uniform()慢！
- curand_init()在 offset 比较大时性能也会比小 offset 差！
- save/load操作 state 比每次重复创建起始 state 性能要快很多 ！
原文如下：Calls to curand_init() are slower than calls to curand() or curand_uniform(). Large offsets to curand_init() take more time than smaller offsets. It is much faster to save and restore random generator state than to recalculate the starting state repeatedly.
对于上述第三点，Nvidia 建议可以将 state 存放到 global memory 中，如下是一个样例代码：
__global__ void example(curandState *global_state)
{
curandState local_state;
local_state = global_state[threadIdx.x];
for(int i = 0; i < 10000; i++) {
unsigned int x = curand(&local_state);
...
}
global_state[threadIdx.x] = local_state;
}
从另一个维度来讲，相对于A产生随机数操作的API，初始化 state 会占用更多的「寄存器」和 local memory 资源。因此 nvidia 建议将 curand_init 和 curand() API 拆分放到不同的 Kernel 中，可以获得最大的性能收益；
原文：Initialization of the random generator state generally requires more registers and local memory than random number generation. It may be beneficial to separate calls to curand_init() and curand() into separate kernels for maximum performance.
State setup can be an expensive operation. One way to speed up the setup is to use different seeds for each thread and a constant sequence number of 0. This can be especially helpful if many generators need to be created. While faster to set up, this method provides less guarantees about the mathematical properties of the generated sequences. If there happens to be a bad interaction between the hash function that initializes the generator state from the seed and the periodicity of the generators, there might be threads with highly correlated outputs for some seed values. We do not know of any problem values; if they do exist they are likely to be rare.
Nvidia 提供的 API 样例代码如下：
#include <stdio.h>
#include <stdlib.h>
#include <cuda_runtime.h>
#include <curand_kernel.h>
#define CUDA_CALL(x) do { if((x) != cudaSuccess) { \
printf("Error at %s:%d\n",__FILE__,__LINE__); \
return EXIT_FAILURE;}} while(0)
__global__ void setup_kernel(curandState *state)
{
int id = threadIdx.x + blockIdx.x * blockDim.x;
/* Each thread gets same seed, a different sequence
number, no offset */
curand_init(1234, id, 0, &state[id]);
}
__global__ void generate_uniform_kernel(curandStatePhilox4_32_10_t *state,
int n,
unsigned int *result)
{
int id = threadIdx.x + blockIdx.x * blockDim.x;
unsigned int count = 0;
float x;
/* Copy state to local memory for efficiency */
curandStatePhilox4_32_10_t localState = state[id];
/* Generate pseudo-random uniforms */
for(int i = 0; i < n; i++) {
x = curand_uniform(&localState);
/* Check if > .5 */
if(x > .5) {
count++;
}
}
/* Copy state back to global memory */
state[id] = localState;
/* Store results */
result[id] += count;
}
