Random erase dropout kernel implementation by RooseweltMcW · Pull Request #556 · r-abishek/rpp

RooseweltMcW · 2025-12-29T14:37:53Z

No description provided.

… HOST AVX, SSE

…ents

…all the bitdepths

…performance

…images in the docs

…opout

… i8 variant

indentation modified

…ied QA for logic to implement random seed for non QA

…ropout

HazarathKumarM · 2025-12-30T07:09:52Z

src/modules/tensor/hip/kernel/random_erase.cpp

+#include "hip_tensor_executors.hpp"
+#include <random>
+
+__device__ inline uint generate_seed(uint x, uint y, uint z, int seed)


add comments for all the hardcoded numbers and specify in the comments why those numbers are needed

As per the new logic the LCG method to produce random noise is removed and the noise buffer is passed from the test suite, Done

…c in test suite and pass buffer for random noise

…ropout

HazarathKumarM

@RooseweltMcW please address the comments

HazarathKumarM · 2026-01-13T08:57:52Z

api/rppt_tensor_effects_augmentations.h

 #endif // GPU_SUPPORT

+/*! \brief Random Erase augmentation on HOST backend for a NCHW/NHWC layout tensor
+ * \details This function erases erases random regions from an image and fills with random noise, for a batch of RGB(3 channel) / greyscale(1 channel) images with an NHWC/NCHW tensor layout.<br>


erase is repeated twice. please correct it

HazarathKumarM · 2026-01-13T08:59:23Z

api/rppt_tensor_effects_augmentations.h

+ * \param [in] anchorBoxInfoTensor anchorBoxInfo values of type RpptRoiLtrb for each erase-region inside each image in the batch HOST memory. Restrictions -
+            - 0 <= anchorBoxInfo[i] < respective image width/height
+            - Erase-region anchor boxes on each image given by the user must not overlap
+ * \param [in] numBoxesTensor number of erase-regions per image, for each image in the batch (numBoxesTensor[n] >= 0)


for random erase its always one box right . Do we even need to pass this tensor to kernel?

Removed, Done

HazarathKumarM · 2026-01-13T09:01:36Z

src/modules/tensor/cpu/kernel/random_erase.cpp

+        RpptROIPtr roiPtrInput = &roiTensorPtrSrc[batchCount];
+        compute_roi_validation_host(roiPtrInput, &roi, &roiDefault, roiType);
+
+        Rpp32u numBoxes = numBoxesTensor[batchCount];


as commneted above numBoxes won't be needed I feel

HazarathKumarM · 2026-01-13T09:03:36Z

src/modules/tensor/cpu/kernel/random_erase.cpp

+
+            for(int count = 0; count < numBoxes; count++)
+            {
+                Rpp32u x1 = static_cast<Rpp32u>(RPPPRANGECHECK(anchorBoxInfo[count].lt.x, roi.xywhROI.xy.x, roi.xywhROI.roiWidth));


modify the func name here, it should be RPPRANGECHECK

This change requires modifying every other function which is using the function RPPP to RPP which has to be a global change

HazarathKumarM · 2026-01-13T09:22:55Z

src/modules/tensor/hip/kernel/random_erase.cpp

+    int id_y = hipBlockIdx_y * hipBlockDim_y + hipThreadIdx_y;
+    int id_z = hipBlockIdx_z * hipBlockDim_z + hipThreadIdx_z;
+
+    auto &roi = roiTensorPtrSrc[id_z].xywhROI;


is it required. please modify the below condition like all other kernels

Removed, Done.

HazarathKumarM · 2026-01-13T09:24:28Z

src/modules/tensor/hip/kernel/random_erase.cpp

+    Rpp32u numBoxes = numBoxesTensor[id_z];
+    uint dstIdx = id_z * dstStridesNH.x + id_y * dstStridesNH.y + id_x * 3;
+
+    for (int i = 0; i < numBoxes; i++)


loop inside hip kernel?, can we not make use of the gpu threads and make it parallel

Removed and processed only one box logic

HazarathKumarM · 2026-01-13T09:25:25Z

src/modules/tensor/hip/kernel/random_erase.cpp

+        // if src layout is NHWC, copy src to dst
+        if (srcDescPtr->layout == RpptLayout::NHWC)
+        {
+            hipMemcpyAsync(dstPtr, srcPtr, static_cast<size_t>(srcDescPtr->n * srcDescPtr->strides.nStride * sizeof(T)), hipMemcpyDeviceToDevice, handle.GetStream());


already srcPtr is copied to dstPtr, then why for loop was there in the hip kernels for num boxes

HazarathKumarM · 2026-01-13T09:26:47Z

src/modules/tensor/hip/kernel/random_erase.cpp

+
+    int globalThreads_x = dstDescPtr->w;
+    int globalThreads_y = dstDescPtr->h;
+    int globalThreads_z = handle.GetBatchSize();


discussed to launch kernel per box. Have to replace it with actual box dimensions

* Add Box and Median Filter ROI fixes after minor corrections * Fix source index computation --------- Co-authored-by: Mukesh <mukesh.jayakodi@multicorewareinc.com> Co-authored-by: Kiriti Gowda <kiritigowda@gmail.com>

…ocs/sphinx (r-abishek#556) Bumps [rocm-docs-core[api_reference]](https://github.com/ROCm/rocm-docs-core) from 1.19.1 to 1.20.0. - [Release notes](https://github.com/ROCm/rocm-docs-core/releases) - [Changelog](https://github.com/ROCm/rocm-docs-core/blob/develop/CHANGELOG.md) - [Commits](ROCm/rocm-docs-core@v1.19.1...v1.20.0) --- updated-dependencies: - dependency-name: rocm-docs-core[api_reference] dependency-version: 1.20.0 dependency-type: direct:production update-type: version-update:semver-minor ... Signed-off-by: dependabot[bot] <support@github.com> Co-authored-by: dependabot[bot] <49699333+dependabot[bot]@users.noreply.github.com> Co-authored-by: Kiriti Gowda <kiritigowda@gmail.com>

snehaa8 and others added 30 commits August 22, 2025 07:39

Optimized version of channel dropout HIP backend and working code for…

e32bbdb

… HOST AVX, SSE

Modified name for dropout compute function

a1fdab1

Modified way of AVX and SSE version channel dropout to avoid if statm…

4c7e1a1

…ents

Modified Channel Dropout with generic compute code reused

5747ea0

Parameters and name change for channel dropout

133236a

Modified HIP for better performance

1fb450e

Modified the code and made the channel dropout templated version for …

63491ac

…all the bitdepths

Modified Erase kernel for grid and cutout dropout version for better …

836ff59

…performance

Modified the .h file to have the dropout to effects and added output …

c9babec

…images in the docs

Added output image and modified the .h file to effects for channel dr…

a374034

…opout

Removed space

f308bbe

added space

709809d

Moved grid dropout .h to effects

938d876

Modified channel dropout for I8 variant HOST side

0bd5686

Resolved all review comments and modified code to produce results for…

751cf48

… i8 variant

Removed empty line

d597022

merged the changes of channel dropout commit to grid, cutout dropout

a0d9ff0

Resolved review comments

4038129

Merge remote-tracking branch 'tot/develop' into apr/channel_dropout

24cada0

Modified HOST after merge

24968cf

Made changes after merging and QA passed for dropout

cbc74c8

Channel dropout make_float 4 macro changes

1baa41e

Updated QA with random generator and updated BIN files

ac97a1e

Modified QA name changes

cd5c576

Merge branch 'apr/channel_dropout' into apr/grid_cutout_dropout

bc62008

Modified RandomSeed value passed as parameter to the function call

09f993c

Update rppt_tensor_effects_augmentations.cpp

da78859

indentation modified

Merge branch 'apr/channel_dropout' into apr/grid_cutout_dropout modif…

ed13d97

…ied QA for logic to implement random seed for non QA

Initial modified HIP backend Grid dropout with better performance

3deb6bd

Removed space and review comments resolved

f0e2195

HazarathKumarM added 7 commits December 17, 2025 11:32

Update kernel and removed unwanted functions

613cf77

HOST modifications for randomErase

f37c8b3

Modified randomization in test suite

b8029c3

Merge remote-tracking branch 'origin/develop' into apr/random_erase_d…

95b2de5

…ropout

Removed unwanted files and headers

c4362a5

Resolved review comments

96b71ea

Updated documentation and resolved comments

00bb001

HazarathKumarM reviewed Dec 30, 2025

View reviewed changes

HazarathKumarM added 5 commits December 31, 2025 01:48

updated random noise generation in test suite HOST implementation

1a034ec

Updated kernel files for random erase to remove noise generation logi…

37c7930

…c in test suite and pass buffer for random noise

updated init dropout function

cac47f3

Merge remote-tracking branch 'origin/develop' into apr/random_erase_d…

5947691

…ropout

Added break statement after merge conflicts

66415f5

HazarathKumarM changed the base branch from develop to master January 13, 2026 08:56

HazarathKumarM changed the base branch from master to develop January 13, 2026 08:56

HazarathKumarM reviewed Jan 13, 2026

View reviewed changes

r-abishek and others added 5 commits January 19, 2026 04:34

ROI fixes - Box filter and Median filter (ROCm#652)

ffeea66

* Add Box and Median Filter ROI fixes after minor corrections * Fix source index computation --------- Co-authored-by: Mukesh <mukesh.jayakodi@multicorewareinc.com> Co-authored-by: Kiriti Gowda <kiritigowda@gmail.com>

Removed numBoxTensor parameter and resolved review comments

1e23365

Merge remote-tracking branch 'tot/develop' into apr/random_erase_dropout

0da73dc

Added early condition to check and return the invalid ROI region

e2fe527

Updated kernel to use anchorBox info for random erase kernel

16d878a

HazarathKumarM added 2 commits February 18, 2026 04:11

Merge branch 'develop' into apr/random_erase_dropout

63e3b27

Add unified api for random_erase

620def3

r-abishek approved these changes Feb 18, 2026

View reviewed changes

r-abishek assigned HazarathKumarM and RooseweltMcW Feb 18, 2026

r-abishek added the enhancement New feature or request label Feb 18, 2026

r-abishek changed the base branch from develop to ar/dropout_random_erase February 18, 2026 23:45

r-abishek merged commit e0b5d2b into r-abishek:ar/dropout_random_erase Feb 18, 2026

Conversation

RooseweltMcW commented Dec 29, 2025

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

HazarathKumarM left a comment

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

4 participants