p47068915
/
mindspore
forked from mindspore-Ecosystem/mindspore
1
0
Fork 0
Code Issues Pull Requests Projects Releases Wiki Activity

!6601 Change the convention that random seed is generated at GPU back-end 

Merge pull request !6601 from peixu_ren/custom_gpu2
This commit is contained in:
mindspore-ci-bot 2020-09-21 09:21:43 +08:00 committed by Gitee
parent aa9fb70f3c 9f1f32aea8
commit 3621bb2348
2 changed files with 7 additions and 3 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

9
mindspore/ccsrc/backend/kernel_compiler/gpu/cuda_impl/random_op_impl.cu
View File

@ -46,12 +46,13 @@ __global__ void UniformRealKernel(int seed, curandState *globalState, T *output,
template <typename T> template <typename T>
void StandardNormal(int seed, int seed2, curandState *globalState, T *output, size_t count, cudaStream_t cuda_stream) { void StandardNormal(int seed, int seed2, curandState *globalState, T *output, size_t count, cudaStream_t cuda_stream) {
int RNG_seed = 0; int RNG_seed = 0;
std::random_device rd;
if (seed2 != 0) { if (seed2 != 0) {
RNG_seed = seed2; RNG_seed = seed2;
} else if (seed != 0) { } else if (seed != 0) {
RNG_seed = seed; RNG_seed = seed;
} else { } else {
RNG_seed = time(NULL); RNG_seed = static_cast<int>(rd());
} }
NormalKernel<<<GET_BLOCKS(count), GET_THREADS, 0, cuda_stream>>>(RNG_seed, globalState, output, count); NormalKernel<<<GET_BLOCKS(count), GET_THREADS, 0, cuda_stream>>>(RNG_seed, globalState, output, count);
return; return;
@ -61,12 +62,13 @@ template <typename T>
void UniformInt(int seed, int seed2, curandState *globalState, T *input1, size_t input_size_1, void UniformInt(int seed, int seed2, curandState *globalState, T *input1, size_t input_size_1,
T *input2, size_t input_size_2, T *output, size_t count, cudaStream_t cuda_stream) { T *input2, size_t input_size_2, T *output, size_t count, cudaStream_t cuda_stream) {
int RNG_seed = 0; int RNG_seed = 0;
std::random_device rd;
if (seed2 != 0) { if (seed2 != 0) {
RNG_seed = seed2; RNG_seed = seed2;
} else if (seed != 0) { } else if (seed != 0) {
RNG_seed = seed; RNG_seed = seed;
} else { } else {
RNG_seed = time(NULL); RNG_seed = static_cast<int>(rd());
} }
UniformIntKernel<<<GET_BLOCKS(count), GET_THREADS, 0, cuda_stream>>> UniformIntKernel<<<GET_BLOCKS(count), GET_THREADS, 0, cuda_stream>>>
(RNG_seed, globalState, input1, input_size_1, input2, input_size_2, output, count); (RNG_seed, globalState, input1, input_size_1, input2, input_size_2, output, count);
@ -76,12 +78,13 @@ void UniformInt(int seed, int seed2, curandState *globalState, T *input1, size_t
template <typename T> template <typename T>
void UniformReal(int seed, int seed2, curandState *globalState, T *output, size_t count, cudaStream_t cuda_stream) { void UniformReal(int seed, int seed2, curandState *globalState, T *output, size_t count, cudaStream_t cuda_stream) {
int RNG_seed = 0; int RNG_seed = 0;
std::random_device rd;
if (seed2 != 0) { if (seed2 != 0) {
RNG_seed = seed2; RNG_seed = seed2;
} else if (seed != 0) { } else if (seed != 0) {
RNG_seed = seed; RNG_seed = seed;
} else { } else {
RNG_seed = time(NULL); RNG_seed = static_cast<int>(rd());
} }
UniformRealKernel<<<GET_BLOCKS(count), GET_THREADS, 0, cuda_stream>>>(RNG_seed, globalState, output, count); UniformRealKernel<<<GET_BLOCKS(count), GET_THREADS, 0, cuda_stream>>>(RNG_seed, globalState, output, count);
return; return;

1
mindspore/ccsrc/backend/kernel_compiler/gpu/cuda_impl/random_op_impl.cuh
View File

@ -18,6 +18,7 @@
#define MINDSPORE_CCSRC_KERNEL_GPU_CUDA_IMPL_RANDOMOPIMPL_H_ #define MINDSPORE_CCSRC_KERNEL_GPU_CUDA_IMPL_RANDOMOPIMPL_H_
#include <curand_kernel.h> #include <curand_kernel.h>
#include <random>
#include "runtime/device/gpu/cuda_common.h" #include "runtime/device/gpu/cuda_common.h"
template <typename T> template <typename T>