Dropout: make seed and states_num kernel arguments #2277

tbennun · 2023-07-26T20:59:11Z

Every time a dropout descriptor is created (even if it's with a different seed), the initialization kernel is recompiled. This PR turns the number of states to initialize and random seed into runtime parameters. There seems to be no effect on the runtime of the initialization kernel (which I'm assuming is not performance critical), but fewer JIT compilation calls.

Note: I tried setting prng_seed to be an unsigned long long parameter, but that passed a wrong value onto the kernel itself.

junliume · 2023-07-29T05:50:18Z

@atamazov could you help to take a look?

tbennun · 2023-07-31T16:58:38Z

@JehandadKhan sorry, called clang-format on the code now.

JehandadKhan · 2023-07-31T17:58:31Z

@tbennun The CL file needs to formatted as well.

tbennun · 2023-07-31T18:01:44Z

@JehandadKhan Done now. I thought I had formatted both.

DrizztDoUrden · 2023-08-10T15:29:23Z

src/ocl/dropoutocl.cpp

-                                 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" +


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

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?

Yes, otherwise it would never really change.

DrizztDoUrden · 2023-08-10T15:31:53Z

src/ocl/dropoutocl.cpp


    auto&& kernels = handle.GetKernels(kernel_name, network_config);
    if(!kernels.empty())
    {
-        kernels.front()(prng_states);
+        kernels.front()(prng_states,


Maybe prng_seed and states_num should be declared as unsigned? That would make calls single lines and remove the casts.

Fixed (through a different way). Changing prng_seed to be unsigned would change the MIOpen C API (in which miopenSetDropoutDescriptor uses unsigned long long). I'd rather pass in an unsigned long long value, which I now do.

I have missed that. That's good now.

junliume · 2023-08-17T20:40:25Z

@tbennun could you resolve the review opinions above? We will try to get this merged.

tbennun · 2023-08-17T22:57:39Z

@junliume I now addressed the two comments, see above and new commit.

* 3D group forward convolution solver (#2286) * [HotFix] Fix build issue after #2286 (#2328) * Fix ConvCkIgemmFwdV6r1DlopsNchw solver to reflect that it's not dynamic (#2325) * Remove target IDs from kdb entries (#2309) * Dropout: make seed and states_num kernel arguments (#2277) * [MI250] Adding missing kernel objects (#2329) * Post-merge fixups: Replace environment variable check with problem config check and reduce lambda capture for Invoker obj (#2305) * [HotFix][CI] fix HIP tidy issue from #2277 (#2335) * [HotFix] Update requirements.txt MLIR ignore PATH for Python conda LLD (#2324) * [NFC] Replace long integers with int64_t and size_t for better compatibility with Windows (#2323) * Windows compatibility: replace long integers with int64_t and size_t, replace uint with unsigned int, replace long with long long for numbers, proper casting * Fix formatting * Fix 3d group forward convolution * Resolve review comments * Fix formatting --------- Co-authored-by: Daming Feng <[email protected]> Co-authored-by: Evgenii Averin <[email protected]> Co-authored-by: JD <[email protected]> Co-authored-by: Tal Ben-Nun <[email protected]> Co-authored-by: amberhassaan <[email protected]> Co-authored-by: Jun Liu <[email protected]>

junliume · 2023-09-20T00:56:35Z

@atamazov could you comment on the user's question?

The private_seg_size of 3204 seems to be caused by the precomputed XORWOW matrices, which are placed in OpenCL constant memory and, upon compilation, may be stored in every thread's memory. It might be useful to store one or both matrices in global memory to avoid the private_seg_size pollution.

CC: @zjing14

atamazov · 2023-09-21T15:16:59Z

@junliume I need some clue to the code as I am not confident with dropout. What are the symbolic names of the "the precomputed XORWOW matrices"?

/cc @zjing14

tbennun · 2023-09-21T15:46:14Z

@junliume I need some clue to the code as I am not confident with dropout. What are the symbolic names of the "the precomputed XORWOW matrices"?

/cc @zjing14

@atamazov This refers to the two included files in the Dropout kernel, which initialize large arrays in constant memory:
https://github.com/ROCmSoftwarePlatform/MIOpen/blob/develop/src/kernels/precalc_xorwow_skipahead_sequence_matrices_kernel.h
https://github.com/ROCmSoftwarePlatform/MIOpen/blob/develop/src/kernels/precalc_xorwow_skipahead_matrices_kernel.h

atamazov · 2023-09-21T21:25:02Z

@tbennun @junliume AFAICS the parts of these matrices are copied to the private memory for a purpose, i.e. each workitem modifies its local copy. If we want to put those into global memory, then we need to allocate some workspace (enough to hold data from matrices) and then manually manage it, which doesn't seem easy to me.

But why we may need to save space in private segment? Are we hitting the limit?

/cc @zjing14

junliume · 2023-10-31T20:47:26Z

@atamazov yes we are hitting a limit on constant memory, with this fix it "gives a reduction in initialization time of a Transformer model from 10:22 minutes to 31 seconds." We should take a look at the OCL code to see if we can work on heap memory instead.

CAHEK7 · 2023-11-02T15:45:21Z

PRNG state requires some workspace in any case, but it supposed to be tiny, but it should be double-checked.
The tables, as far as I know, are required to speedup PRNG initialization and warmup (it may require huge numbers of PRNG invocation) and if that's the case - we can hold the tables on CPU side, choose particular entry for the seed and offset, upload it to the state and save a lot of constant\workspace memory.

But to be sure we need to re-investigate our GPU PRNG implementation.

atamazov · 2023-11-02T16:50:10Z

@CAHEK7 Thanks for feedback. Now we at least better understand what to begin with.

/cc @JehandadKhan @junliume

Dropout: make seed and states_num arguments

97af6d0

tbennun changed the title ~~Dropout: make seed and states_num arguments~~ Dropout: make seed and states_num runtime arguments Jul 26, 2023

tbennun changed the title ~~Dropout: make seed and states_num runtime arguments~~ Dropout: make seed and states_num kernel arguments Jul 26, 2023

junliume requested review from DrizztDoUrden and JehandadKhan July 29, 2023 05:50

junliume added the urgency_normal label Jul 29, 2023

JehandadKhan previously approved these changes Jul 31, 2023

View reviewed changes

clang-format

55ad8ee

tbennun dismissed JehandadKhan’s stale review via 55ad8ee July 31, 2023 16:57

tbennun requested a review from JehandadKhan July 31, 2023 16:58

clang-format

108984b

DrizztDoUrden requested changes Aug 10, 2023

View reviewed changes

Remove casts and pass in a 64-bit value for seed

1c15ae0

junliume requested a review from DrizztDoUrden August 17, 2023 23:10

DrizztDoUrden approved these changes Aug 18, 2023

View reviewed changes

junliume added the TESTING_CI_PASSED label Aug 18, 2023

junliume approved these changes Aug 18, 2023

View reviewed changes

junliume merged commit 5ea8aa6 into ROCm:develop Aug 18, 2023

junliume added a commit that referenced this pull request Aug 20, 2023

[HotFix][CI] fix HIP tidy issue from #2277

4a937be

junliume added a commit that referenced this pull request Aug 21, 2023

[HotFix][CI] fix HIP tidy issue from #2277 (#2335)

21a9261

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

Dropout: make seed and states_num kernel arguments #2277

Dropout: make seed and states_num kernel arguments #2277

tbennun commented Jul 26, 2023

junliume commented Jul 29, 2023

tbennun commented Jul 31, 2023

JehandadKhan commented Jul 31, 2023

tbennun commented Jul 31, 2023

DrizztDoUrden Aug 10, 2023

tbennun Aug 17, 2023

DrizztDoUrden Aug 18, 2023

DrizztDoUrden Aug 10, 2023

tbennun Aug 17, 2023

DrizztDoUrden Aug 18, 2023 •

edited

Loading

junliume commented Aug 17, 2023

tbennun commented Aug 17, 2023

junliume commented Sep 20, 2023

atamazov commented Sep 21, 2023

tbennun commented Sep 21, 2023

atamazov commented Sep 21, 2023

junliume commented Oct 31, 2023

CAHEK7 commented Nov 2, 2023 •

edited

Loading

atamazov commented Nov 2, 2023

Dropout: make seed and states_num kernel arguments #2277

Dropout: make seed and states_num kernel arguments #2277

Conversation

tbennun commented Jul 26, 2023

junliume commented Jul 29, 2023

tbennun commented Jul 31, 2023

JehandadKhan commented Jul 31, 2023

tbennun commented Jul 31, 2023

DrizztDoUrden Aug 10, 2023

Choose a reason for hiding this comment

tbennun Aug 17, 2023

Choose a reason for hiding this comment

DrizztDoUrden Aug 18, 2023

Choose a reason for hiding this comment

DrizztDoUrden Aug 10, 2023

Choose a reason for hiding this comment

tbennun Aug 17, 2023

Choose a reason for hiding this comment

DrizztDoUrden Aug 18, 2023 • edited Loading

Choose a reason for hiding this comment

junliume commented Aug 17, 2023

tbennun commented Aug 17, 2023

junliume commented Sep 20, 2023

atamazov commented Sep 21, 2023

tbennun commented Sep 21, 2023

atamazov commented Sep 21, 2023

junliume commented Oct 31, 2023

CAHEK7 commented Nov 2, 2023 • edited Loading

atamazov commented Nov 2, 2023

DrizztDoUrden Aug 18, 2023 •

edited

Loading

CAHEK7 commented Nov 2, 2023 •

edited

Loading