Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Dropout: make seed and states_num kernel arguments #2277

Merged
merged 4 commits into from
Aug 18, 2023
Merged

Dropout: make seed and states_num kernel arguments #2277

Show file tree
Hide file tree
Changes from 1 commit
Commits
Show all changes
4 commits
Select commit Hold shift + click to select a range
97af6d0
Dropout: make seed and states_num arguments
tbennun Jul 26, 2023
55ad8ee
clang-format
tbennun Jul 31, 2023
108984b
clang-format
tbennun Jul 31, 2023
1c15ae0
Remove casts and pass in a 64-bit value for seed
tbennun Aug 17, 2023
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
7 changes: 4 additions & 3 deletions src/kernels/MIOpenDropout.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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);

Expand Down
12 changes: 6 additions & 6 deletions src/ocl/dropoutocl.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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<unsigned int>(prng_seed),
static_cast<unsigned int>(states_num));
}
else
{
Expand All @@ -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<unsigned int>(prng_seed),
static_cast<unsigned int>(states_num));
#if DROPOUT_DEBUG
std::cout << "Succeeded in launching InitPRNGState()." << stateSizeInBytes << std::endl;
#endif
Expand Down