diff --git a/projects/hip-tests/catch/hipTestMain/config/config_amd_linux b/projects/hip-tests/catch/hipTestMain/config/config_amd_linux
index 0f6f626367..a02d218e60 100644
--- a/projects/hip-tests/catch/hipTestMain/config/config_amd_linux
+++ b/projects/hip-tests/catch/hipTestMain/config/config_amd_linux
@@ -141,6 +141,11 @@
         "Unit_Device___hbgtu2_Accuracy_Positive",
         "Unit_Device___hbleu2_Accuracy_Positive",
         "Unit_Device___hbltu2_Accuracy_Positive",
+        "=== Below 4 tests are disable due to defect EXSWHTEC-355 ===",
+        "Unit_Device___hadd_Sanity_Positive",
+        "Unit_Device___uhadd_Sanity_Positive",
+        "Unit_Device___rhadd_Sanity_Positive",
+        "Unit_Device___urhadd_Sanity_Positive",
         "=== Patch which removes the typetraits implementation from std namespace in hiprtc is reverted ===",
         "Unit_hiprtc_stdheaders",
         "Unit_hipGraphAddMemcpyNode_Negative_Parameters",
diff --git a/projects/hip-tests/catch/hipTestMain/config/config_amd_windows b/projects/hip-tests/catch/hipTestMain/config/config_amd_windows
index 1522525fea..8dbe04c1b0 100644
--- a/projects/hip-tests/catch/hipTestMain/config/config_amd_windows
+++ b/projects/hip-tests/catch/hipTestMain/config/config_amd_windows
@@ -466,6 +466,11 @@
         "Unit_Device___hbgtu2_Accuracy_Positive",
         "Unit_Device___hbleu2_Accuracy_Positive",
         "Unit_Device___hbltu2_Accuracy_Positive",
+        "=== Below 4 tests are disable due to defect EXSWHTEC-355 ===",
+        "Unit_Device___hadd_Sanity_Positive",
+        "Unit_Device___uhadd_Sanity_Positive",
+        "Unit_Device___rhadd_Sanity_Positive",
+        "Unit_Device___urhadd_Sanity_Positive",
     #endif
         "End of json"
     ]
diff --git a/projects/hip-tests/catch/unit/math/integer_intrinsics.cc b/projects/hip-tests/catch/unit/math/integer_intrinsics.cc
index d851577831..ee9cd760a8 100644
--- a/projects/hip-tests/catch/unit/math/integer_intrinsics.cc
+++ b/projects/hip-tests/catch/unit/math/integer_intrinsics.cc
@@ -317,4 +317,478 @@ TEST_CASE("Unit_Device___umul24_Sanity_Positive") {
   HIP_CHECK(hipDeviceSynchronize());
 
   REQUIRE(y.ptr()[0] == x1 * x2);
+}
+
+__global__ void __funnelshift_l_kernel(unsigned int* y, unsigned int lo, unsigned int hi,
+                                       unsigned int shift) {
+  y[0] = __funnelshift_l(lo, hi, shift);
+}
+
+/**
+ * Test Description
+ * ------------------------
+ *    - Sanity test for `__funnelshift_l(lo,hi,shift)`.
+ *
+ * Test source
+ * ------------------------
+ *    - unit/math/integer_intrinsics.cc
+ * Test requirements
+ * ------------------------
+ *    - HIP_VERSION >= 5.2
+ */
+TEST_CASE("Unit_Device___funnelshift_l_Sanity_Positive") {
+  LinearAllocGuard<unsigned int> y(LinearAllocs::hipMallocManaged, sizeof(unsigned int));
+
+  const unsigned int lo = 0xAAAAAAAA, hi = 0xBBBBBBBB;
+  const unsigned long long hi_lo = (static_cast<unsigned long long>(hi) << 32) | lo;
+
+  for (unsigned int shift = 0; shift < 64; ++shift) {
+    __funnelshift_l_kernel<<<1, 1>>>(y.ptr(), lo, hi, shift);
+    HIP_CHECK(hipDeviceSynchronize());
+
+    INFO("shift: " << shift);
+    REQUIRE(y.ptr()[0] == static_cast<unsigned int>((hi_lo << (shift & 31)) >> 32));
+  }
+}
+
+__global__ void __funnelshift_lc_kernel(unsigned int* y, unsigned int lo, unsigned int hi,
+                                        unsigned int shift) {
+  y[0] = __funnelshift_lc(lo, hi, shift);
+}
+
+/**
+ * Test Description
+ * ------------------------
+ *    - Sanity test for `__funnelshift_lc(lo,hi,shift)`.
+ *
+ * Test source
+ * ------------------------
+ *    - unit/math/integer_intrinsics.cc
+ * Test requirements
+ * ------------------------
+ *    - HIP_VERSION >= 5.2
+ */
+TEST_CASE("Unit_Device___funnelshift_lc_Sanity_Positive") {
+  LinearAllocGuard<unsigned int> y(LinearAllocs::hipMallocManaged, sizeof(unsigned int));
+
+  const unsigned int lo = 0xAAAAAAAA, hi = 0xBBBBBBBB;
+  const unsigned long long hi_lo = (static_cast<unsigned long long>(hi) << 32) | lo;
+
+  for (unsigned int shift = 0; shift < 64; ++shift) {
+    __funnelshift_lc_kernel<<<1, 1>>>(y.ptr(), lo, hi, shift);
+    HIP_CHECK(hipDeviceSynchronize());
+
+    INFO("shift: " << shift);
+    REQUIRE(y.ptr()[0] == static_cast<unsigned int>((hi_lo << std::min(shift, 32u)) >> 32));
+  }
+}
+
+__global__ void __funnelshift_r_kernel(unsigned int* y, unsigned int lo, unsigned int hi,
+                                       unsigned int shift) {
+  y[0] = __funnelshift_r(lo, hi, shift);
+}
+
+/**
+ * Test Description
+ * ------------------------
+ *    - Sanity test for `__funnelshift_r(lo,hi,shift)`.
+ *
+ * Test source
+ * ------------------------
+ *    - unit/math/integer_intrinsics.cc
+ * Test requirements
+ * ------------------------
+ *    - HIP_VERSION >= 5.2
+ */
+TEST_CASE("Unit_Device___funnelshift_r_Sanity_Positive") {
+  LinearAllocGuard<unsigned int> y(LinearAllocs::hipMallocManaged, sizeof(unsigned int));
+
+  const unsigned int lo = 0xAAAAAAAA, hi = 0xBBBBBBBB;
+  const unsigned long long hi_lo = (static_cast<unsigned long long>(hi) << 32) | lo;
+
+  for (unsigned int shift = 0; shift < 64; ++shift) {
+    __funnelshift_r_kernel<<<1, 1>>>(y.ptr(), lo, hi, shift);
+    HIP_CHECK(hipDeviceSynchronize());
+
+    INFO("shift: " << shift);
+    REQUIRE(y.ptr()[0] == static_cast<unsigned int>(hi_lo >> (shift & 31)));
+  }
+}
+
+__global__ void __funnelshift_rc_kernel(unsigned int* y, unsigned int lo, unsigned int hi,
+                                        unsigned int shift) {
+  y[0] = __funnelshift_rc(lo, hi, shift);
+}
+
+/**
+ * Test Description
+ * ------------------------
+ *    - Sanity test for `__funnelshift_rc(lo,hi,shift)`.
+ *
+ * Test source
+ * ------------------------
+ *    - unit/math/integer_intrinsics.cc
+ * Test requirements
+ * ------------------------
+ *    - HIP_VERSION >= 5.2
+ */
+TEST_CASE("Unit_Device___funnelshift_rc_Sanity_Positive") {
+  LinearAllocGuard<unsigned int> y(LinearAllocs::hipMallocManaged, sizeof(unsigned int));
+
+  const unsigned int lo = 0xAAAAAAAA, hi = 0xBBBBBBBB;
+  const unsigned long long hi_lo = (static_cast<unsigned long long>(hi) << 32) | lo;
+
+  for (unsigned int shift = 0; shift < 64; ++shift) {
+    __funnelshift_rc_kernel<<<1, 1>>>(y.ptr(), lo, hi, shift);
+    HIP_CHECK(hipDeviceSynchronize());
+
+    INFO("shift: " << shift);
+    REQUIRE(y.ptr()[0] == static_cast<unsigned int>(hi_lo >> std::min(shift, 32u)));
+  }
+}
+
+__global__ void __hadd_kernel(int* y, int x1, int x2) { y[0] = __hadd(x1, x2); }
+
+/**
+ * Test Description
+ * ------------------------
+ *    - Sanity test for `__hadd(x,y)`.
+ *
+ * Test source
+ * ------------------------
+ *    - unit/math/integer_intrinsics.cc
+ * Test requirements
+ * ------------------------
+ *    - HIP_VERSION >= 5.2
+ */
+TEST_CASE("Unit_Device___hadd_Sanity_Positive") {
+  LinearAllocGuard<int> y(LinearAllocs::hipMallocManaged, sizeof(int));
+
+  int x1 = GENERATE(0, -42, 42, 0xFFFFFFFF);
+  int x2 = GENERATE(0, -42, 42, 0xFFFFFFFF);
+
+  __hadd_kernel<<<1, 1>>>(y.ptr(), x1, x2);
+  HIP_CHECK(hipDeviceSynchronize());
+
+  INFO("x1: " << x1);
+  INFO("x2: " << x2);
+  REQUIRE(y.ptr()[0] == static_cast<int>((static_cast<long long>(x1) + x2) >> 1));
+}
+
+__global__ void __uhadd_kernel(unsigned int* y, unsigned int x1, unsigned int x2) {
+  y[0] = __uhadd(x1, x2);
+}
+
+/**
+ * Test Description
+ * ------------------------
+ *    - Sanity test for `__uhadd(x,y)`.
+ *
+ * Test source
+ * ------------------------
+ *    - unit/math/integer_intrinsics.cc
+ * Test requirements
+ * ------------------------
+ *    - HIP_VERSION >= 5.2
+ */
+TEST_CASE("Unit_Device___uhadd_Sanity_Positive") {
+  LinearAllocGuard<unsigned int> y(LinearAllocs::hipMallocManaged, sizeof(unsigned int));
+
+  unsigned int x1 = GENERATE(0, 42, 0xFFFFFFFF);
+  unsigned int x2 = GENERATE(0, 42, 0xFFFFFFFF);
+
+  __uhadd_kernel<<<1, 1>>>(y.ptr(), x1, x2);
+  HIP_CHECK(hipDeviceSynchronize());
+
+  INFO("x1: " << x1);
+  INFO("x2: " << x2);
+  REQUIRE(y.ptr()[0] == static_cast<unsigned int>((static_cast<unsigned long long>(x1) + x2) >> 1));
+}
+
+__global__ void __rhadd_kernel(int* y, int x1, int x2) { y[0] = __rhadd(x1, x2); }
+
+/**
+ * Test Description
+ * ------------------------
+ *    - Sanity test for `__rhadd(x,y)`.
+ *
+ * Test source
+ * ------------------------
+ *    - unit/math/integer_intrinsics.cc
+ * Test requirements
+ * ------------------------
+ *    - HIP_VERSION >= 5.2
+ */
+TEST_CASE("Unit_Device___rhadd_Sanity_Positive") {
+  LinearAllocGuard<int> y(LinearAllocs::hipMallocManaged, sizeof(int));
+
+  int x1 = GENERATE(0, -42, 42, 0xFFFFFFFF);
+  int x2 = GENERATE(0, -42, 42, 0xFFFFFFFF);
+
+  __rhadd_kernel<<<1, 1>>>(y.ptr(), x1, x2);
+  HIP_CHECK(hipDeviceSynchronize());
+
+  INFO("x1: " << x1);
+  INFO("x2: " << x2);
+  REQUIRE(y.ptr()[0] == static_cast<int>((static_cast<long long>(x1) + x2 + 1) >> 1));
+}
+
+__global__ void __urhadd_kernel(unsigned int* y, unsigned int x1, unsigned int x2) {
+  y[0] = __urhadd(x1, x2);
+}
+
+/**
+ * Test Description
+ * ------------------------
+ *    - Sanity test for `__urhadd(x,y)`.
+ *
+ * Test source
+ * ------------------------
+ *    - unit/math/integer_intrinsics.cc
+ * Test requirements
+ * ------------------------
+ *    - HIP_VERSION >= 5.2
+ */
+TEST_CASE("Unit_Device___urhadd_Sanity_Positive") {
+  LinearAllocGuard<unsigned int> y(LinearAllocs::hipMallocManaged, sizeof(unsigned int));
+
+  unsigned int x1 = GENERATE(0, 42, 0xFFFFFFFF);
+  unsigned int x2 = GENERATE(0, 42, 0xFFFFFFFF);
+
+  __urhadd_kernel<<<1, 1>>>(y.ptr(), x1, x2);
+  HIP_CHECK(hipDeviceSynchronize());
+
+  INFO("x1: " << x1);
+  INFO("x2: " << x2);
+  REQUIRE(y.ptr()[0] ==
+          static_cast<unsigned int>((static_cast<unsigned long long>(x1) + x2 + 1) >> 1));
+}
+
+__global__ void __mulhi_kernel(int* y, int x1, int x2) { y[0] = __mulhi(x1, x2); }
+
+/**
+ * Test Description
+ * ------------------------
+ *    - Sanity test for `__mulhi(x,y)`.
+ *
+ * Test source
+ * ------------------------
+ *    - unit/math/integer_intrinsics.cc
+ * Test requirements
+ * ------------------------
+ *    - HIP_VERSION >= 5.2
+ */
+TEST_CASE("Unit_Device___mulhi_Sanity_Positive") {
+  LinearAllocGuard<int> y(LinearAllocs::hipMallocManaged, sizeof(int));
+
+  int x1 = GENERATE(0, -42, 42, 0xFFFFFFFF);
+  int x2 = GENERATE(0, -42, 42, 0xFFFFFFFF);
+
+  __mulhi_kernel<<<1, 1>>>(y.ptr(), x1, x2);
+  HIP_CHECK(hipDeviceSynchronize());
+
+  INFO("x1: " << x1);
+  INFO("x2: " << x2);
+  REQUIRE(y.ptr()[0] ==
+          static_cast<int>((static_cast<long long>(x1) * static_cast<long long>(x2)) >> 32));
+}
+
+__global__ void __umulhi_kernel(unsigned int* y, unsigned int x1, unsigned int x2) {
+  y[0] = __umulhi(x1, x2);
+}
+
+/**
+ * Test Description
+ * ------------------------
+ *    - Sanity test for `__umulhi(x,y)`.
+ *
+ * Test source
+ * ------------------------
+ *    - unit/math/integer_intrinsics.cc
+ * Test requirements
+ * ------------------------
+ *    - HIP_VERSION >= 5.2
+ */
+TEST_CASE("Unit_Device___umulhi_Sanity_Positive") {
+  LinearAllocGuard<unsigned int> y(LinearAllocs::hipMallocManaged, sizeof(unsigned int));
+
+  unsigned int x1 = GENERATE(0, 42, 0xFFFFFFFF);
+  unsigned int x2 = GENERATE(0, 42, 0xFFFFFFFF);
+
+  __umulhi_kernel<<<1, 1>>>(y.ptr(), x1, x2);
+  HIP_CHECK(hipDeviceSynchronize());
+
+  INFO("x1: " << x1);
+  INFO("x2: " << x2);
+  REQUIRE(y.ptr()[0] ==
+          static_cast<unsigned int>((static_cast<unsigned long long>(x1) * x2) >> 32));
+}
+
+__global__ void __mul64hi_kernel(long long* y, long long x1, long long x2) {
+  y[0] = __mul64hi(x1, x2);
+}
+
+/**
+ * Test Description
+ * ------------------------
+ *    - Sanity test for `__mul64hi(x,y)`.
+ *
+ * Test source
+ * ------------------------
+ *    - unit/math/integer_intrinsics.cc
+ * Test requirements
+ * ------------------------
+ *    - HIP_VERSION >= 5.2
+ */
+TEST_CASE("Unit_Device___mul64hi_Sanity_Positive") {
+  LinearAllocGuard<long long> y(LinearAllocs::hipMallocManaged, sizeof(long long));
+
+  long long x1 = GENERATE(0, -42, 42, 0xFFFFFFFF);
+  long long x2 = GENERATE(0, -42, 42, 0xFFFFFFFF);
+
+  __mul64hi_kernel<<<1, 1>>>(y.ptr(), x1, x2);
+  HIP_CHECK(hipDeviceSynchronize());
+
+  INFO("x1: " << x1);
+  INFO("x2: " << x2);
+  REQUIRE(
+      y.ptr()[0] ==
+      static_cast<long long>((static_cast<__int128_t>(x1) * static_cast<__int128_t>(x2)) >> 64));
+}
+
+__global__ void __umul64hi_kernel(unsigned long long* y, unsigned long long x1,
+                                  unsigned long long x2) {
+  y[0] = __umul64hi(x1, x2);
+}
+
+/**
+ * Test Description
+ * ------------------------
+ *    - Sanity test for `__umul64hi(x,y)`.
+ *
+ * Test source
+ * ------------------------
+ *    - unit/math/integer_intrinsics.cc
+ * Test requirements
+ * ------------------------
+ *    - HIP_VERSION >= 5.2
+ */
+TEST_CASE("Unit_Device___umul64hi_Sanity_Positive") {
+  LinearAllocGuard<unsigned long long> y(LinearAllocs::hipMallocManaged,
+                                         sizeof(unsigned long long));
+
+  unsigned long long x1 = GENERATE(0, 42, 0xFFFFFFFF);
+  unsigned long long x2 = GENERATE(0, 42, 0xFFFFFFFF);
+
+  __umul64hi_kernel<<<1, 1>>>(y.ptr(), x1, x2);
+  HIP_CHECK(hipDeviceSynchronize());
+
+  INFO("x1: " << x1);
+  INFO("x2: " << x2);
+  REQUIRE(y.ptr()[0] ==
+          static_cast<unsigned long long>(
+              (static_cast<__uint128_t>(x1) * static_cast<__uint128_t>(x2)) >> 64));
+}
+
+__global__ void __sad_kernel(unsigned int* y, int x1, int x2, unsigned int x3) {
+  y[0] = __sad(x1, x2, x3);
+}
+
+/**
+ * Test Description
+ * ------------------------
+ *    - Sanity test for `__sad(x,y,z)`.
+ *
+ * Test source
+ * ------------------------
+ *    - unit/math/integer_intrinsics.cc
+ * Test requirements
+ * ------------------------
+ *    - HIP_VERSION >= 5.2
+ */
+TEST_CASE("Unit_Device___sad_Sanity_Positive") {
+  LinearAllocGuard<unsigned int> y(LinearAllocs::hipMallocManaged, sizeof(unsigned int));
+
+  int x1 = GENERATE(0, -42, 42, 0xFFFFFFFF);
+  int x2 = GENERATE(0, -42, 42, 0xFFFFFFFF);
+  unsigned int x3 = GENERATE(0, 42, 0xFFFFFFFF);
+
+  __sad_kernel<<<1, 1>>>(y.ptr(), x1, x2, x3);
+  HIP_CHECK(hipDeviceSynchronize());
+
+  INFO("x1: " << x1);
+  INFO("x2: " << x2);
+  REQUIRE(y.ptr()[0] == (static_cast<unsigned int>(std::abs(x1 - x2)) + x3));
+}
+
+__global__ void __usad_kernel(unsigned int* y, unsigned int x1, unsigned int x2, unsigned int x3) {
+  y[0] = __usad(x1, x2, x3);
+}
+
+/**
+ * Test Description
+ * ------------------------
+ *    - Sanity test for `__usad(x,y,z)`.
+ *
+ * Test source
+ * ------------------------
+ *    - unit/math/integer_intrinsics.cc
+ * Test requirements
+ * ------------------------
+ *    - HIP_VERSION >= 5.2
+ */
+TEST_CASE("Unit_Device___usad_Sanity_Positive") {
+  LinearAllocGuard<unsigned int> y(LinearAllocs::hipMallocManaged, sizeof(unsigned int));
+
+  unsigned int x1 = GENERATE(0, 42, 0xFFFFFFFF);
+  unsigned int x2 = GENERATE(0, 42, 0xFFFFFFFF);
+  unsigned int x3 = GENERATE(0, 42, 0xFFFFFFFF);
+
+  __usad_kernel<<<1, 1>>>(y.ptr(), x1, x2, x3);
+  HIP_CHECK(hipDeviceSynchronize());
+
+  INFO("x1: " << x1);
+  INFO("x2: " << x2);
+  REQUIRE(y.ptr()[0] ==
+          (static_cast<unsigned int>(
+               std::abs(static_cast<long long>(x1) - static_cast<long long>(x2))) +
+           x3));
+}
+
+__global__ void __byte_perm(unsigned int* y, unsigned int x1, unsigned int x2, unsigned int s) {
+  y[0] = __byte_perm(x1, x2, s);
+}
+
+/**
+ * Test Description
+ * ------------------------
+ *    - Sanity test for `__byte_perm(x,y,s)`.
+ *
+ * Test source
+ * ------------------------
+ *    - unit/math/integer_intrinsics.cc
+ * Test requirements
+ * ------------------------
+ *    - HIP_VERSION >= 5.2
+ */
+TEST_CASE("Unit_Device___byte_perm_Sanity_Positive") {
+  LinearAllocGuard<unsigned int> y(LinearAllocs::hipMallocManaged, sizeof(unsigned int));
+
+  unsigned int bytes[] = {0x88, 0x99, 0xAA, 0xBB, 0xCC, 0xDD, 0xEE, 0xFF};
+
+  unsigned int x1 = (bytes[3] << 24) | (bytes[2] << 16) | (bytes[1] << 8) | bytes[0];
+  unsigned int x2 = (bytes[7] << 24) | (bytes[6] << 16) | (bytes[5] << 8) | bytes[4];
+
+  unsigned int s0 = GENERATE(0, 1);
+  unsigned int s1 = GENERATE(2, 3);
+  unsigned int s2 = GENERATE(4, 5);
+  unsigned int s3 = GENERATE(6, 7);
+
+  unsigned int s = (s3 << 12) | (s2 << 8) | (s1 << 4) | s0;
+
+  __byte_perm<<<1, 1>>>(y.ptr(), x1, x2, s);
+  HIP_CHECK(hipDeviceSynchronize());
+
+  unsigned int expected = (bytes[s3] << 24) | (bytes[s2] << 16) | (bytes[s1] << 8) | bytes[s0];
+  REQUIRE(y.ptr()[0] == expected);
 }
\ No newline at end of file