From 97af6d014112c554225e23992c13e250549ea5e6 Mon Sep 17 00:00:00 2001 From: Tal Ben-Nun Date: Wed, 26 Jul 2023 13:50:30 -0700 Subject: [PATCH 1/4] Dropout: make seed and states_num arguments --- src/kernels/MIOpenDropout.cl | 7 ++++--- src/ocl/dropoutocl.cpp | 12 ++++++------ 2 files changed, 10 insertions(+), 9 deletions(-) diff --git a/src/kernels/MIOpenDropout.cl b/src/kernels/MIOpenDropout.cl index 75e29fc4fb..b7410558a7 100644 --- a/src/kernels/MIOpenDropout.cl +++ b/src/kernels/MIOpenDropout.cl @@ -273,13 +273,14 @@ void xorwow_lite_init(prngStates* cur_state, cur_state->d += (uint)(offset)*362437; } -__kernel void InitKernelState(__global prngStates* state) +__kernel void InitKernelState(__global prngStates* state, uint prng_seed, + uint states_num) { - for(uint gid = get_global_id(0); gid < STATES_NUM; gid += get_global_size(0)) + for(uint gid = get_global_id(0); gid < states_num; gid += get_global_size(0)) { prngStates state_gid; xorwow_lite_init(&state_gid, - (unsigned long long)PRNG_SEED, + (unsigned long long)prng_seed, (unsigned long long)gid, (unsigned long long)0); diff --git a/src/ocl/dropoutocl.cpp b/src/ocl/dropoutocl.cpp index c52f21da49..9ae3ae7099 100644 --- a/src/ocl/dropoutocl.cpp +++ b/src/ocl/dropoutocl.cpp @@ -135,15 +135,16 @@ void DropoutDescriptor::InitPRNGState(Handle& handle, size_t states_num = prng_stateSizeInBytes / sizeof(prngStates); size_t wk_grp_num = std::min(size_t(MAX_PRNG_STATE / 256), (states_num + 255) / 256); - std::string network_config = "initprngs-" + std::to_string(states_num) + "x" + + std::string network_config = "initprngs-" + std::to_string(sizeof(prngStates)) + "x" + - std::to_string(rng_mode) + "x" + std::to_string(prng_seed) + "x" + + std::to_string(rng_mode) + "x" + std::to_string(wk_grp_num); auto&& kernels = handle.GetKernels(kernel_name, network_config); if(!kernels.empty()) { - kernels.front()(prng_states); + kernels.front()(prng_states, static_cast(prng_seed), + static_cast(states_num)); } else { @@ -152,14 +153,13 @@ void DropoutDescriptor::InitPRNGState(Handle& handle, std::string params; params += " -DRUN_INIT_PRNG=1"; - params += " -DPRNG_SEED=" + std::to_string(prng_seed); - params += " -DSTATES_NUM=" + std::to_string(states_num); #if DROPOUT_DEBUG std::cout << "Threads allocated for PRNG states: " << vgd[0] << std::endl; std::cout << "Memory allocated for PRNG states: " << stateSizeInBytes << std::endl; #endif handle.AddKernel(kernel_name, network_config, program_name, kernel_name, vld, vgd, params)( - prng_states); + prng_states, static_cast(prng_seed), + static_cast(states_num)); #if DROPOUT_DEBUG std::cout << "Succeeded in launching InitPRNGState()." << stateSizeInBytes << std::endl; #endif From 55ad8ee766e0e33f05b935cb10e3661740f487f4 Mon Sep 17 00:00:00 2001 From: Tal Ben-Nun Date: Mon, 31 Jul 2023 09:57:27 -0700 Subject: [PATCH 2/4] clang-format --- src/ocl/dropoutocl.cpp | 12 ++++++------ 1 file changed, 6 insertions(+), 6 deletions(-) diff --git a/src/ocl/dropoutocl.cpp b/src/ocl/dropoutocl.cpp index 9ae3ae7099..446dcd0e0a 100644 --- a/src/ocl/dropoutocl.cpp +++ b/src/ocl/dropoutocl.cpp @@ -135,15 +135,14 @@ void DropoutDescriptor::InitPRNGState(Handle& handle, size_t states_num = prng_stateSizeInBytes / sizeof(prngStates); size_t wk_grp_num = std::min(size_t(MAX_PRNG_STATE / 256), (states_num + 255) / 256); - std::string network_config = "initprngs-" + - std::to_string(sizeof(prngStates)) + "x" + - std::to_string(rng_mode) + "x" + - std::to_string(wk_grp_num); + std::string network_config = "initprngs-" + std::to_string(sizeof(prngStates)) + "x" + + std::to_string(rng_mode) + "x" + std::to_string(wk_grp_num); auto&& kernels = handle.GetKernels(kernel_name, network_config); if(!kernels.empty()) { - kernels.front()(prng_states, static_cast(prng_seed), + kernels.front()(prng_states, + static_cast(prng_seed), static_cast(states_num)); } else @@ -158,7 +157,8 @@ void DropoutDescriptor::InitPRNGState(Handle& handle, std::cout << "Memory allocated for PRNG states: " << stateSizeInBytes << std::endl; #endif handle.AddKernel(kernel_name, network_config, program_name, kernel_name, vld, vgd, params)( - prng_states, static_cast(prng_seed), + prng_states, + static_cast(prng_seed), static_cast(states_num)); #if DROPOUT_DEBUG std::cout << "Succeeded in launching InitPRNGState()." << stateSizeInBytes << std::endl; From 108984b343b2c84f0f99e2767207d53d5723a17e Mon Sep 17 00:00:00 2001 From: Tal Ben-Nun Date: Mon, 31 Jul 2023 11:00:49 -0700 Subject: [PATCH 3/4] clang-format --- src/kernels/MIOpenDropout.cl | 3 +-- 1 file changed, 1 insertion(+), 2 deletions(-) diff --git a/src/kernels/MIOpenDropout.cl b/src/kernels/MIOpenDropout.cl index b7410558a7..bbeb6f6767 100644 --- a/src/kernels/MIOpenDropout.cl +++ b/src/kernels/MIOpenDropout.cl @@ -273,8 +273,7 @@ void xorwow_lite_init(prngStates* cur_state, cur_state->d += (uint)(offset)*362437; } -__kernel void InitKernelState(__global prngStates* state, uint prng_seed, - uint states_num) +__kernel void InitKernelState(__global prngStates* state, uint prng_seed, uint states_num) { for(uint gid = get_global_id(0); gid < states_num; gid += get_global_size(0)) { From 1c15ae07316fdd8afbb9f429f315e8720a4b32b9 Mon Sep 17 00:00:00 2001 From: Tal Ben-Nun Date: Thu, 17 Aug 2023 15:56:04 -0700 Subject: [PATCH 4/4] Remove casts and pass in a 64-bit value for seed --- src/kernels/MIOpenDropout.cl | 2 +- src/ocl/dropoutocl.cpp | 12 ++++-------- 2 files changed, 5 insertions(+), 9 deletions(-) diff --git a/src/kernels/MIOpenDropout.cl b/src/kernels/MIOpenDropout.cl index bbeb6f6767..596423418f 100644 --- a/src/kernels/MIOpenDropout.cl +++ b/src/kernels/MIOpenDropout.cl @@ -273,7 +273,7 @@ void xorwow_lite_init(prngStates* cur_state, cur_state->d += (uint)(offset)*362437; } -__kernel void InitKernelState(__global prngStates* state, uint prng_seed, uint states_num) +__kernel void InitKernelState(__global prngStates* state, ulong prng_seed, ulong states_num) { for(uint gid = get_global_id(0); gid < states_num; gid += get_global_size(0)) { diff --git a/src/ocl/dropoutocl.cpp b/src/ocl/dropoutocl.cpp index 446dcd0e0a..caf44d9569 100644 --- a/src/ocl/dropoutocl.cpp +++ b/src/ocl/dropoutocl.cpp @@ -132,8 +132,8 @@ void DropoutDescriptor::InitPRNGState(Handle& handle, MIOPEN_THROW("PRNG state size should not exceed system maximum memory allocation size."); } - size_t states_num = prng_stateSizeInBytes / sizeof(prngStates); - size_t wk_grp_num = std::min(size_t(MAX_PRNG_STATE / 256), (states_num + 255) / 256); + unsigned long long states_num = prng_stateSizeInBytes / sizeof(prngStates); + size_t wk_grp_num = std::min(MAX_PRNG_STATE / 256ULL, (states_num + 255) / 256); std::string network_config = "initprngs-" + std::to_string(sizeof(prngStates)) + "x" + std::to_string(rng_mode) + "x" + std::to_string(wk_grp_num); @@ -141,9 +141,7 @@ void DropoutDescriptor::InitPRNGState(Handle& handle, auto&& kernels = handle.GetKernels(kernel_name, network_config); if(!kernels.empty()) { - kernels.front()(prng_states, - static_cast(prng_seed), - static_cast(states_num)); + kernels.front()(prng_states, prng_seed, states_num); } else { @@ -157,9 +155,7 @@ void DropoutDescriptor::InitPRNGState(Handle& handle, std::cout << "Memory allocated for PRNG states: " << stateSizeInBytes << std::endl; #endif handle.AddKernel(kernel_name, network_config, program_name, kernel_name, vld, vgd, params)( - prng_states, - static_cast(prng_seed), - static_cast(states_num)); + prng_states, prng_seed, states_num); #if DROPOUT_DEBUG std::cout << "Succeeded in launching InitPRNGState()." << stateSizeInBytes << std::endl; #endif