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 all commits
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
6 changes: 3 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,13 @@ 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, ulong prng_seed, ulong 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
16 changes: 6 additions & 10 deletions src/ocl/dropoutocl.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -132,18 +132,16 @@ 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(states_num) + "x" +
std::to_string(sizeof(prngStates)) + "x" +
std::to_string(rng_mode) + "x" + std::to_string(prng_seed) + "x" +
std::to_string(wk_grp_num);
std::string network_config = "initprngs-" + std::to_string(sizeof(prngStates)) + "x" +
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Please use sdt::ostringstream instead of concatenations if there are more than two strings. It's both faster and easier to read.

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Thank you. I would do this, but the rest of the file (as well as other files such as softmaxocl.cpp) use this style of code. Would it be ok to not match the style of the rest of the code?

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Yes, otherwise it would never really change.

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, prng_seed, states_num);
}
else
{
Expand All @@ -152,14 +150,12 @@ 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, prng_seed, states_num);
#if DROPOUT_DEBUG
std::cout << "Succeeded in launching InitPRNGState()." << stateSizeInBytes << std::endl;
#endif
Expand Down