BitwiseAND and BitwiseOR on HOST and HIP #230

snehaa8 · 2024-02-06T06:39:58Z

No description provided.

Also includes fixing f16 and f32 datatype of HOST

Update Copywrite

…sn/bitwise_OR

Includes reference output

r-abishek

@snehaa8 Pls address comments. Replicate changes mentioned in either bitwiseOr/AND for both.
Lets combine BitwiseAND + BitwiseOR in this Internal PR and close the other one for ease. You can work on this sn/bitwise_OR branch for both.

r-abishek · 2024-02-07T03:02:00Z

src/modules/cpu/kernel/bitwise_and.hpp

@@ -0,0 +1,957 @@
+/*
+MIT License


Replace exact text with the blank lines as in LICENSE file outside:

MIT License Copyright (c) 2019 - 2024 Advanced Micro Devices, Inc. Permission is hereby granted, free of charge, to any person obtaining a copy of this software and associated documentation files (the "Software"), to deal in the Software without restriction, including without limitation the rights to use, copy, modify, merge, publish, distribute, sublicense, and/or sell copies of the Software, and to permit persons to whom the Software is furnished to do so, subject to the following conditions: The above copyright notice and this permission notice shall be included in all copies or substantial portions of the Software. THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE.

r-abishek · 2024-02-07T03:05:47Z

src/include/cpu/rpp_cpu_simd.hpp

+{
+    __m128i pxSrc[8];
+    __m128i pxMask = _mm_setr_epi8(0, 3, 6, 9, 1, 4, 7, 10, 2, 5, 8, 11, 12, 13, 14, 15);
+    __m128i pxMaskRGB = _mm_setr_epi8(0, 4, 8, 12, 2, 6, 10, 14, 1, 5, 9, 13, 3, 7, 11, 15);


Don't we have the 0,3,6,9 mask or the 0,4,8,12 mask pre-allocated outside of the runtime execution path somewhere since they are common?

Checked this while implementing, didn't find any in the fashion I needed.

Should i add this at the start of rpp_cpu_simd where other common constants are defined?

r-abishek · 2024-02-07T03:18:00Z

src/modules/cpu/kernel/bitwise_and.hpp

+#endif
+                for (; vectorLoopCount < bufferLength; vectorLoopCount += 3)
+                {
+                    *dstPtrTempR++ = RPPPIXELCHECKF32((float)((uint)(srcPtr1Temp[0] * 255) & (uint)(srcPtr2Temp[0] * 255)) / 255);


Please add a comment like below on HOST/HIP for bitwiseAND and bitwiseOR to clarify to the reader. Either the link you pointed to before, or establish validity with openCV or other lib.

// BitwiseAND / BitwiseOR are logical operations only on U8/I8 types. For a float / half precision image (pixel values from 0-1), the BitwiseAND / BitwiseOR is applied on a 0-255 range-translated approximation, of the original 0-1 decimal-range image

Added comments and link as pointed to before

r-abishek · 2024-02-07T03:19:35Z

src/modules/cpu/kernel/bitwise_or.hpp

@@ -0,0 +1,957 @@
+/*
+MIT License


Same comment

r-abishek · 2024-02-07T03:20:02Z

src/modules/cpu/kernel/bitwise_or.hpp

+#endif
+                for (; vectorLoopCount < bufferLength; vectorLoopCount += 3)
+                {
+                    *dstPtrTempR++ = RPPPIXELCHECKF32((float)((uint)(srcPtr1Temp[0] * 255) | (uint)(srcPtr2Temp[0] * 255)) / 255);


Same comment for Bitwise OR

r-abishek · 2024-02-07T03:26:51Z

src/modules/hip/kernel/bitwise_or.hpp

+#include "rpp_hip_common.hpp"
+
+template <typename T>
+__device__ void bitwise_or_hip_compute(T *srcPtr, d_float8 *src1_f8, d_float8 *src2_f8, d_float8 *dst_f8)


template <typename T> __device__ void bitwise_or_hip_compute(T *srcPtr, d_float8 *src1_f8, d_float8 *src2_f8, d_float8 *dst_f8) { if constexpr ((std::is_same<T, float>::value) || (std::is_same<T, half>::value)) { rpp_hip_math_multiply8_const(src1_f8, src1_f8, (float4)255); rpp_hip_math_multiply8_const(src2_f8, src2_f8, (float4)255); rpp_hip_math_bitwiseOr8(src1_f8, src2_f8, dst_f8); rpp_hip_math_multiply8_const(dst_f8, dst_f8, (float4)ONE_OVER_255); } else if constexpr (std::is_same<T, signed char>::value) { rpp_hip_math_add8_const(src1_f8, src1_f8, (float4)128); rpp_hip_math_add8_const(src2_f8, src2_f8, (float4)128); rpp_hip_math_bitwiseOr8(src1_f8, src2_f8, dst_f8); rpp_hip_math_subtract8_const(dst_f8, dst_f8, (float4)128); } }

Modified it a little more like

template <typename T> __device__ void bitwise_and_hip_compute(T *srcPtr, d_float8 *src1_f8, d_float8 *src2_f8, d_float8 *dst_f8) { if constexpr ((std::is_same<T, float>::value) || (std::is_same<T, half>::value)) { rpp_hip_math_multiply8_const(src1_f8, src1_f8, (float4)255); rpp_hip_math_multiply8_const(src2_f8, src2_f8, (float4)255); rpp_hip_math_bitwiseAnd8(src1_f8, src2_f8, dst_f8); rpp_hip_math_multiply8_const(dst_f8, dst_f8, (float4)ONE_OVER_255); } else if constexpr (std::is_same<T, signed char>::value) { rpp_hip_math_add8_const(src1_f8, src1_f8, (float4)128); rpp_hip_math_add8_const(src2_f8, src2_f8, (float4)128); rpp_hip_math_bitwiseAnd8(src1_f8, src2_f8, dst_f8); rpp_hip_math_subtract8_const(dst_f8, dst_f8, (float4)128); } else rpp_hip_math_bitwiseAnd8(src1_f8, src2_f8, dst_f8); }

r-abishek · 2024-02-07T03:27:51Z

src/modules/hip/kernel/bitwise_or.hpp

+template <typename T>
+__device__ void bitwise_or_hip_compute(T *srcPtr, d_float8 *src1_f8, d_float8 *src2_f8, d_float8 *dst_f8)
+{
+    float4 adjustment_f4;


Remove unused variable

r-abishek · 2024-02-07T03:31:56Z

utilities/test_suite/rpp_test_suite_common.h

@@ -97,7 +98,8 @@ std::map<int, string> augmentationMap =
    {84, "spatter"},
    {85, "swap_channels"},
    {86, "color_to_greyscale"},
-    {87, "tensor_sum"}
+    {87, "tensor_sum"},
+    {92, "bitwise_or"}


Why is the case number for bitwiseOR this far apart from bitwiseAND? Pls check BatchPD case numbers.

Modified testCase of Bitwise OR to 68 to match with Inclusive OR of BatchPD.

snehaa8 · 2024-02-08T14:24:38Z

Please take a final look

r-abishek · 2024-02-20T22:59:42Z

@snehaa8 Conflicts. Pull upstream develop into your branch.

snehaa8 · 2024-02-21T05:51:34Z

Resolved merge conflicts

r-abishek

@snehaa8 I made some minor formatting changes on your branch. Please also move the header files as in comment.

r-abishek · 2024-02-21T19:19:51Z

src/modules/hip/kernel/bitwise_and.hpp

+#include <hip/hip_runtime.h>
+#include "rpp_hip_common.hpp"
+
+/* BitwiseAND is logical operation only on U8/I8 types.


We need to create a "logical_operations" header for bitwise ops. Currently everything is under arithmetic.
Test suite grouping/classification, external .h include, and internal .hpp includes need to change.

Done, please recheck.
Confirmed QA test pass too.

snehaa8 added 13 commits January 5, 2024 06:54

Initial commit - Bitwise AND HOST Tensor

3debd2f

Match u8 and i8 outputs with BatchPD variant

569ad14

Fix i8 PKD3 -> PLN3

8c68c8a

Initial commit - Bitwise AND HIP Tensor

562c336

Also includes fixing f16 and f32 datatype of HOST

Add reference outputs

d03f07a

Merge branch 'master' of https://github.com/ROCm/rpp into sn/bitwise_AND

d12870a

Modify reference outputs

5e524f7

Update Copywrite

Combine templated functions to support all datatypes

5e4b736

Initial commit - Bitwise OR HOST

ac88212

Fix GPU kernel details

97b9aa1

Merge branch 'sn/bitwise_AND' of https://github.com/snehaa8/rpp into …

557d8d4

…sn/bitwise_OR

Fix case number for HOST testsuite

226bf9d

Initial commit - Bitwise OR HIP

3df2c11

Includes reference output

r-abishek requested changes Feb 7, 2024

View reviewed changes

r-abishek assigned snehaa8 Feb 7, 2024

r-abishek added the enhancement New feature or request label Feb 7, 2024

r-abishek added this to the sow10ms3 milestone Feb 7, 2024

r-abishek changed the base branch from master to ar/opt_bitwise_and_or February 7, 2024 03:35

r-abishek mentioned this pull request Feb 7, 2024

Bitwise AND Tensor #218

Closed

r-abishek changed the title ~~Bitwise OR Kernel~~ BitwiseAND and BitwiseOR on HOST and HIP Feb 7, 2024

Address review comments

1590f34

r-abishek approved these changes Feb 20, 2024

View reviewed changes

Merge branch 'ar/opt_bitwise_and_or' into sn/bitwise_OR

3246085

r-abishek added 2 commits February 21, 2024 11:04

Update rppt_tensor_arithmetic_operations.h

860b853

Update rppt_tensor_arithmetic_operations.h

304fc09

r-abishek requested changes Feb 21, 2024

View reviewed changes

Merge branch 'develop' of https://github.com/ROCm/rpp into sn/bitwise_OR

ea89580

Move bitwise operations into under logical ops

e4c6388

r-abishek changed the base branch from ar/opt_bitwise_and_or to develop February 27, 2024 23:52

r-abishek changed the base branch from develop to ar/opt_bitwise_and_or February 27, 2024 23:53

r-abishek approved these changes Mar 6, 2024

View reviewed changes

r-abishek merged commit 0aa9f07 into r-abishek:ar/opt_bitwise_and_or Mar 6, 2024

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

BitwiseAND and BitwiseOR on HOST and HIP #230

BitwiseAND and BitwiseOR on HOST and HIP #230

snehaa8 commented Feb 6, 2024

r-abishek left a comment

r-abishek Feb 7, 2024

r-abishek Feb 7, 2024

snehaa8 Feb 8, 2024

snehaa8 Feb 8, 2024

r-abishek Feb 7, 2024

snehaa8 Feb 8, 2024

r-abishek Feb 7, 2024

r-abishek Feb 7, 2024

r-abishek Feb 7, 2024

snehaa8 Feb 8, 2024

r-abishek Feb 7, 2024

r-abishek Feb 7, 2024

snehaa8 Feb 8, 2024

snehaa8 commented Feb 8, 2024

r-abishek commented Feb 20, 2024

snehaa8 commented Feb 21, 2024

r-abishek left a comment

r-abishek Feb 21, 2024

snehaa8 Feb 27, 2024

BitwiseAND and BitwiseOR on HOST and HIP #230

BitwiseAND and BitwiseOR on HOST and HIP #230

Conversation

snehaa8 commented Feb 6, 2024

r-abishek left a comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

snehaa8 commented Feb 8, 2024

r-abishek commented Feb 20, 2024

snehaa8 commented Feb 21, 2024

r-abishek left a comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment