diff --git a/projects/hip-tests/catch/unit/atomics/CMakeLists.txt b/projects/hip-tests/catch/unit/atomics/CMakeLists.txt
index 2aef527440..fa77a87a99 100644
--- a/projects/hip-tests/catch/unit/atomics/CMakeLists.txt
+++ b/projects/hip-tests/catch/unit/atomics/CMakeLists.txt
@@ -19,23 +19,50 @@
 # THE SOFTWARE.
 
 set(TEST_SRC
+    atomicAnd.cc
+    atomicAnd_system.cc
+    atomicOr.cc
+    atomicOr_system.cc
+    atomicXor.cc
+    atomicXor_system.cc
     atomicExch.cc
     atomicExch_system.cc
 )
 
 if(HIP_PLATFORM MATCHES "nvidia")
+    set_source_files_properties(atomicExch_system.cc PROPERTIES COMPILE_FLAGS "-gencode arch=compute_60,code=sm_60 -gencode arch=compute_70,code=sm_70 -gencode arch=compute_80,code=sm_80")
+    set_source_files_properties(atomicAnd_system.cc PROPERTIES COMPILE_FLAGS "-gencode arch=compute_60,code=sm_60 -gencode arch=compute_70,code=sm_70 -gencode arch=compute_80,code=sm_80")
+    set_source_files_properties(atomicOr_system.cc PROPERTIES COMPILE_FLAGS "-gencode arch=compute_60,code=sm_60 -gencode arch=compute_70,code=sm_70 -gencode arch=compute_80,code=sm_80")
+    set_source_files_properties(atomicXor_system.cc PROPERTIES COMPILE_FLAGS "-gencode arch=compute_60,code=sm_60 -gencode arch=compute_70,code=sm_70 -gencode arch=compute_80,code=sm_80")
     set_source_files_properties(atomicExch_system.cc PROPERTIES COMPILE_FLAGS "-gencode arch=compute_60,code=sm_60 -gencode arch=compute_70,code=sm_70 -gencode arch=compute_80,code=sm_80")
     hip_add_exe_to_target(NAME AtomicsTest
-                        TEST_SRC ${TEST_SRC}
-                        TEST_TARGET_NAME build_tests
-                        LINKER_LIBS "nvrtc -gencode arch=compute_60,code=sm_60 -gencode arch=compute_70,code=sm_70 -gencode arch=compute_80,code=sm_80")
+                          TEST_SRC ${TEST_SRC}
+                          TEST_TARGET_NAME build_tests
+                          LINKER_LIBS "nvrtc -gencode arch=compute_60,code=sm_60 -gencode arch=compute_70,code=sm_70 -gencode arch=compute_80,code=sm_80")
+    set(EXPECTED_ERRORS 36) # EXSWHTEC-278
 elseif(HIP_PLATFORM MATCHES "amd")
     hip_add_exe_to_target(NAME AtomicsTest
-                        TEST_SRC ${TEST_SRC}
-                        TEST_TARGET_NAME build_tests
-                        LINKER_LIBS hiprtc)
+                          TEST_SRC ${TEST_SRC}
+                          TEST_TARGET_NAME build_tests
+                          LINKER_LIBS hiprtc)
+    set(EXPECTED_ERRORS 40)
 endif()
 
+add_test(NAME Unit_atomicAnd_Negative_Parameters
+         COMMAND python3 ${CMAKE_CURRENT_SOURCE_DIR}/../compileAndCaptureOutput.py
+         ${CMAKE_CURRENT_SOURCE_DIR} ${HIP_PLATFORM} ${HIP_PATH}
+         atomicAnd_negative_kernels.cc ${EXPECTED_ERRORS})
+
+add_test(NAME Unit_atomicOr_Negative_Parameters
+         COMMAND python3 ${CMAKE_CURRENT_SOURCE_DIR}/../compileAndCaptureOutput.py
+         ${CMAKE_CURRENT_SOURCE_DIR} ${HIP_PLATFORM} ${HIP_PATH}
+         atomicOr_negative_kernels.cc ${EXPECTED_ERRORS})
+
+add_test(NAME Unit_atomicXor_Negative_Parameters
+         COMMAND python3 ${CMAKE_CURRENT_SOURCE_DIR}/../compileAndCaptureOutput.py
+         ${CMAKE_CURRENT_SOURCE_DIR} ${HIP_PLATFORM} ${HIP_PATH}
+         atomicXor_negative_kernels.cc ${EXPECTED_ERRORS})
+
 # SWDEV-435667: Below 2 tests failed in stress test on 01/12/23
 #add_test(NAME Unit_atomicExch_Negative_Parameters
 #         COMMAND python3 ${CMAKE_CURRENT_SOURCE_DIR}/../compileAndCaptureOutput.py
diff --git a/projects/hip-tests/catch/unit/atomics/atomicAnd.cc b/projects/hip-tests/catch/unit/atomics/atomicAnd.cc
new file mode 100644
index 0000000000..756526a31c
--- /dev/null
+++ b/projects/hip-tests/catch/unit/atomics/atomicAnd.cc
@@ -0,0 +1,222 @@
+/*
+Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved.
+
+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.
+*/
+
+#include "atomicAnd_negative_kernels_rtc.hh"
+#include "bitwise_common.hh"
+
+#include <hip_test_common.hh>
+
+/**
+ * @addtogroup atomicAnd atomicAnd
+ * @{
+ * @ingroup AtomicsTest
+ * `atomicAnd(TestType* address, TestType* val)` -
+ * performs atomic bitwise AND between address and val, returns old value.
+ */
+
+/**
+ * Test Description
+ * ------------------------
+ *  - Performs atomicAnd from multiple threads on the same address.
+ *  - Uses only one device and launches one kernel.
+ * Test source
+ * ------------------------
+ *  - unit/atomics/atomicAnd.cc
+ * Test requirements
+ * ------------------------
+ *  - HIP_VERSION >= 5.2
+ */
+TEMPLATE_TEST_CASE("Unit_atomicAnd_Positive_SameAddress", "", int, unsigned int, unsigned long,
+                   unsigned long long) {
+  for (auto current = 0; current < cmd_options.iterations; ++current) {
+    DYNAMIC_SECTION("Same address " << current) {
+      Bitwise::SingleDeviceSingleKernelTest<TestType, Bitwise::AtomicOperation::kAnd>(
+          1, sizeof(TestType));
+    }
+  }
+}
+
+/**
+ * Test Description
+ * ------------------------
+ *  - Performs atomicAnd from multiple threads on adjacent addresses.
+ *  - Uses only one device and launches one kernel.
+ * Test source
+ * ------------------------
+ *  - unit/atomics/atomicAnd.cc
+ * Test requirements
+ * ------------------------
+ *  - HIP_VERSION >= 5.2
+ */
+TEMPLATE_TEST_CASE("Unit_atomicAnd_Positive_Adjacent_Addresses", "", int, unsigned int,
+                   unsigned long, unsigned long long) {
+  int warp_size = 0;
+  HIP_CHECK(hipDeviceGetAttribute(&warp_size, hipDeviceAttributeWarpSize, 0));
+
+  for (auto current = 0; current < cmd_options.iterations; ++current) {
+    DYNAMIC_SECTION("Adjacent address " << current) {
+      Bitwise::SingleDeviceSingleKernelTest<TestType, Bitwise::AtomicOperation::kAnd>(
+          warp_size, sizeof(TestType));
+    }
+  }
+}
+
+/**
+ * Test Description
+ * ------------------------
+ *  - Performs atomicAnd from multiple threads on the scattered addresses.
+ *  - Uses only one device and launches one kernel.
+ * Test source
+ * ------------------------
+ *  - unit/atomics/atomicAnd.cc
+ * Test requirements
+ * ------------------------
+ *  - HIP_VERSION >= 5.2
+ */
+TEMPLATE_TEST_CASE("Unit_atomicAnd_Positive_Scattered_Addresses", "", int, unsigned int,
+                   unsigned long, unsigned long long) {
+  int warp_size = 0;
+  HIP_CHECK(hipDeviceGetAttribute(&warp_size, hipDeviceAttributeWarpSize, 0));
+  const auto cache_line_size = 128u;
+
+  for (auto current = 0; current < cmd_options.iterations; ++current) {
+    DYNAMIC_SECTION("Scattered address " << current) {
+      Bitwise::SingleDeviceSingleKernelTest<TestType, Bitwise::AtomicOperation::kAnd>(
+          warp_size, cache_line_size);
+    }
+  }
+}
+
+/**
+ * Test Description
+ * ------------------------
+ *  - Performs atomicAnd from multiple threads on the same address.
+ *  - Uses only one device and launches multiple kernels.
+ * Test source
+ * ------------------------
+ *  - unit/atomics/atomicAnd.cc
+ * Test requirements
+ * ------------------------
+ *  - HIP_VERSION >= 5.2
+ */
+TEMPLATE_TEST_CASE("Unit_atomicAnd_Positive_Multi_Kernel_Same_Address", "", int, unsigned int,
+                   unsigned long, unsigned long long) {
+  for (auto current = 0; current < cmd_options.iterations; ++current) {
+    DYNAMIC_SECTION("Same address " << current) {
+      Bitwise::SingleDeviceMultipleKernelTest<TestType, Bitwise::AtomicOperation::kAnd>(
+          2, 1, sizeof(TestType));
+    }
+  }
+}
+
+/**
+ * Test Description
+ * ------------------------
+ *  - Performs atomicAnd from multiple threads on adjacent addresses.
+ *  - Uses only one device and launches multiple kernels.
+ * Test source
+ * ------------------------
+ *  - unit/atomics/atomicAnd.cc
+ * Test requirements
+ * ------------------------
+ *  - HIP_VERSION >= 5.2
+ */
+TEMPLATE_TEST_CASE("Unit_atomicAnd_Positive_Multi_Kernel_Adjacent_Addresses", "", int, unsigned int,
+                   unsigned long, unsigned long long) {
+  int warp_size = 0;
+  HIP_CHECK(hipDeviceGetAttribute(&warp_size, hipDeviceAttributeWarpSize, 0));
+
+  for (auto current = 0; current < cmd_options.iterations; ++current) {
+    DYNAMIC_SECTION("Adjacent address " << current) {
+      Bitwise::SingleDeviceMultipleKernelTest<TestType, Bitwise::AtomicOperation::kAnd>(
+          2, warp_size, sizeof(TestType));
+    }
+  }
+}
+
+/**
+ * Test Description
+ * ------------------------
+ *  - Performs atomicAnd from multiple threads on the scattered addresses.
+ *  - Uses only one device and launches multiple kernels.
+ * Test source
+ * ------------------------
+ *  - unit/atomics/atomicAnd.cc
+ * Test requirements
+ * ------------------------
+ *  - HIP_VERSION >= 5.2
+ */
+TEMPLATE_TEST_CASE("Unit_atomicAnd_Positive_Multi_Kernel_Scattered_Addresses", "", int,
+                   unsigned int, unsigned long, unsigned long long) {
+  int warp_size = 0;
+  HIP_CHECK(hipDeviceGetAttribute(&warp_size, hipDeviceAttributeWarpSize, 0));
+  const auto cache_line_size = 128u;
+
+  for (auto current = 0; current < cmd_options.iterations; ++current) {
+    DYNAMIC_SECTION("Scattered address " << current) {
+      Bitwise::SingleDeviceMultipleKernelTest<TestType, Bitwise::AtomicOperation::kAnd>(
+          2, warp_size, cache_line_size);
+    }
+  }
+}
+
+/**
+ * Test Description
+ * ------------------------
+ *  - Compiles atomicAnd with invalid parameters.
+ *  - Compiles the source with RTC.
+ * Test source
+ * ------------------------
+ *  - unit/atomics/atomicAnd.cc
+ * Test requirements
+ * ------------------------
+ *  - HIP_VERSION >= 5.2
+ */
+TEST_CASE("Unit_atomicAnd_Negative_Parameters_RTC") {
+  hiprtcProgram program{};
+
+  const auto program_source =
+      GENERATE(kAtomicAnd_int, kAtomicAnd_uint, kAtomicAnd_ulong, kAtomicAnd_ulonglong);
+  HIPRTC_CHECK(
+      hiprtcCreateProgram(&program, program_source, "atomicAnd_negative.cc", 0, nullptr, nullptr));
+  hiprtcResult result{hiprtcCompileProgram(program, 0, nullptr)};
+
+  // Get the compile log and count compiler error messages
+  size_t log_size{};
+  HIPRTC_CHECK(hiprtcGetProgramLogSize(program, &log_size));
+  std::string log(log_size, ' ');
+  HIPRTC_CHECK(hiprtcGetProgramLog(program, log.data()));
+  int error_count{0};
+  // Please check the content of negative_kernels_rtc.hh
+  int expected_error_count{9};
+  std::string error_message{"error:"};
+
+  size_t n_pos = log.find(error_message, 0);
+  while (n_pos != std::string::npos) {
+    ++error_count;
+    n_pos = log.find(error_message, n_pos + 1);
+  }
+
+  HIPRTC_CHECK(hiprtcDestroyProgram(&program));
+  HIPRTC_CHECK_ERROR(result, HIPRTC_ERROR_COMPILATION);
+  REQUIRE(error_count == expected_error_count);
+}
diff --git a/projects/hip-tests/catch/unit/atomics/atomicAnd_negative_kernels.cc b/projects/hip-tests/catch/unit/atomics/atomicAnd_negative_kernels.cc
new file mode 100644
index 0000000000..593399b45e
--- /dev/null
+++ b/projects/hip-tests/catch/unit/atomics/atomicAnd_negative_kernels.cc
@@ -0,0 +1,185 @@
+/*
+Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved.
+
+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.
+*/
+
+#include <hip_test_common.hh>
+
+class Dummy {
+ public:
+  __device__ Dummy() {}
+  __device__ ~Dummy() {}
+};
+
+/* int atomicAnd(int* address, int val) */
+__global__ void atomicAnd_int_v1(int* address, int* result) { *result = atomicAnd(&address, 1234); }
+
+__global__ void atomicAnd_int_v2(int* address, int* result) {
+  *result = atomicAnd(address, address);
+}
+
+__global__ void atomicAnd_int_v3(int* address, int* result) { *result = atomicAnd(1234, 1234); }
+
+__global__ void atomicAnd_int_v4(Dummy* address, int* result) {
+  *result = atomicAnd(address, 1234);
+}
+
+__global__ void atomicAnd_int_v5(char* address, int* result) { *result = atomicAnd(address, 1234); }
+
+__global__ void atomicAnd_int_v6(short* address, int* result) {
+  *result = atomicAnd(address, 1234);
+}
+
+__global__ void atomicAnd_int_v7(long* address, int* result) { *result = atomicAnd(address, 1234); }
+
+__global__ void atomicAnd_int_v8(long long* address, int* result) {
+  *result = atomicAnd(address, 1234);
+}
+
+__global__ void atomicAnd_int_v9(float* address, int* result) {
+  *result = atomicAnd(address, 1234);
+}
+
+__global__ void atomicAnd_int_v10(double* address, int* result) {
+  *result = atomicAnd(address, 1234);
+}
+
+/* unsigned int atomicAnd(unsigned int* address, unsigned int val) */
+__global__ void atomicAnd_uint_v1(unsigned int* address, unsigned int* result) {
+  *result = atomicAnd(&address, 1234);
+}
+
+__global__ void atomicAnd_uint_v2(unsigned int* address, unsigned int* result) {
+  *result = atomicAnd(address, address);
+}
+
+__global__ void atomicAnd_uint_v3(unsigned int* address, unsigned int* result) {
+  *result = atomicAnd(1234, 1234);
+}
+
+__global__ void atomicAnd_uint_v4(Dummy* address, unsigned int* result) {
+  *result = atomicAnd(address, 1234);
+}
+
+__global__ void atomicAnd_uint_v5(char* address, unsigned int* result) {
+  *result = atomicAnd(address, 1234);
+}
+
+__global__ void atomicAnd_uint_v6(short* address, unsigned int* result) {
+  *result = atomicAnd(address, 1234);
+}
+
+__global__ void atomicAnd_uint_v7(long* address, unsigned int* result) {
+  *result = atomicAnd(address, 1234);
+}
+
+__global__ void atomicAnd_uint_v8(long long* address, unsigned int* result) {
+  *result = atomicAnd(address, 1234);
+}
+
+__global__ void atomicAnd_uint_v9(float* address, unsigned int* result) {
+  *result = atomicAnd(address, 1234);
+}
+
+__global__ void atomicAnd_uint_v10(double* address, unsigned int* result) {
+  *result = atomicAnd(address, 1234);
+}
+
+/* atomicAnd(unsigned long* address, unsigned long val) */
+__global__ void atomicAnd_ulong_v1(unsigned long* address, unsigned long* result) {
+  *result = atomicAnd(&address, 1234);
+}
+
+__global__ void atomicAnd_ulong_v2(unsigned long* address, unsigned long* result) {
+  *result = atomicAnd(address, address);
+}
+
+__global__ void atomicAnd_ulong_v3(unsigned long* address, unsigned long* result) {
+  *result = atomicAnd(1234, 1234);
+}
+
+__global__ void atomicAnd_ulong_v4(Dummy* address, unsigned long* result) {
+  *result = atomicAnd(address, 1234);
+}
+
+__global__ void atomicAnd_ulong_v5(char* address, unsigned long* result) {
+  *result = atomicAnd(address, 1234);
+}
+
+__global__ void atomicAnd_ulong_v6(short* address, unsigned long* result) {
+  *result = atomicAnd(address, 1234);
+}
+
+__global__ void atomicAnd_ulong_v7(long* address, unsigned long* result) {
+  *result = atomicAnd(address, 1234);
+}
+
+__global__ void atomicAnd_ulong_v8(long long* address, unsigned long* result) {
+  *result = atomicAnd(address, 1234);
+}
+
+__global__ void atomicAnd_ulong_v9(float* address, unsigned long* result) {
+  *result = atomicAnd(address, 1234);
+}
+
+__global__ void atomicAnd_ulong_v10(double* address, unsigned long* result) {
+  *result = atomicAnd(address, 1234);
+}
+
+/* atomicAnd(unsigned long long* address, unsigned long long val) */
+__global__ void atomicAnd_ulonglong_v1(unsigned long long* address, unsigned long long* result) {
+  *result = atomicAnd(&address, 1234);
+}
+
+__global__ void atomicAnd_ulonglong_v2(unsigned long long* address, unsigned long long* result) {
+  *result = atomicAnd(address, address);
+}
+
+__global__ void atomicAnd_ulonglong_v3(unsigned long long* address, unsigned long long* result) {
+  *result = atomicAnd(1234, 1234);
+}
+
+__global__ void atomicAnd_ulonglong_v4(Dummy* address, unsigned long long* result) {
+  *result = atomicAnd(address, 1234);
+}
+
+__global__ void atomicAnd_ulonglong_v5(char* address, unsigned long long* result) {
+  *result = atomicAnd(address, 1234);
+}
+
+__global__ void atomicAnd_ulonglong_v6(short* address, unsigned long long* result) {
+  *result = atomicAnd(address, 1234);
+}
+
+__global__ void atomicAnd_ulonglong_v7(long* address, unsigned long long* result) {
+  *result = atomicAnd(address, 1234);
+}
+
+__global__ void atomicAnd_ulonglong_v8(long long* address, unsigned long long* result) {
+  *result = atomicAnd(address, 1234);
+}
+
+__global__ void atomicAnd_ulonglong_v9(float* address, unsigned long long* result) {
+  *result = atomicAnd(address, 1234);
+}
+
+__global__ void atomicAnd_ulonglong_v10(double* address, unsigned long long* result) {
+  *result = atomicAnd(address, 1234);
+}
diff --git a/projects/hip-tests/catch/unit/atomics/atomicAnd_negative_kernels_rtc.hh b/projects/hip-tests/catch/unit/atomics/atomicAnd_negative_kernels_rtc.hh
new file mode 100644
index 0000000000..d637feb9fe
--- /dev/null
+++ b/projects/hip-tests/catch/unit/atomics/atomicAnd_negative_kernels_rtc.hh
@@ -0,0 +1,223 @@
+/*
+Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved.
+
+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.
+*/
+
+#pragma once
+
+/*
+Negative kernels used for the atomics negative Test Cases that are using RTC.
+*/
+
+static constexpr auto kAtomicAnd_int{
+    R"(
+    __global__ void atomicAnd_int_v1(int* address, int* result) {
+      *result = atomicAnd(&address, 1234);
+    }
+
+    __global__ void atomicAnd_int_v2(int* address, int* result) {
+      *result = atomicAnd(address, address);
+    }
+
+    __global__ void atomicAnd_int_v3(int* address, int* result) {
+      *result = atomicAnd(1234, 1234);
+    }
+
+    class Dummy {
+     public:
+      __device__ Dummy() {}
+      __device__ ~Dummy() {}
+    };
+
+    __global__ void atomicAnd_int_v4(Dummy* address, int* result) {
+      *result = atomicAnd(address, 1234);
+    }
+
+    __global__ void atomicAnd_int_v5(char* address, int* result) {
+      *result = atomicAnd(address, 1234);
+    }
+
+    __global__ void atomicAnd_int_v6(short* address, int* result) {
+      *result = atomicAnd(address, 1234);
+    }
+
+    __global__ void atomicAnd_int_v7(long* address, int* result) {
+      *result = atomicAnd(address, 1234);
+    }
+
+    __global__ void atomicAnd_int_v8(long long* address, int* result) {
+      *result = atomicAnd(address, 1234);
+    }
+
+    __global__ void atomicAnd_int_v9(float* address, int* result) {
+      *result = atomicAnd(address, 1234);
+    }
+
+    __global__ void atomicAnd_int_v10(double* address, int* result) {
+      *result = atomicAnd(address, 1234);
+    }
+  )"};
+
+static constexpr auto kAtomicAnd_uint{
+    R"(
+    __global__ void atomicAnd_uint_v1(unsigned int* address, unsigned int* result) {
+      *result = atomicAnd(&address, 1234);
+    }
+
+    __global__ void atomicAnd_uint_v2(unsigned int* address, unsigned int* result) {
+      *result = atomicAnd(address, address);
+    }
+
+    __global__ void atomicAnd_uint_v3(unsigned int* address, unsigned int* result) {
+      *result = atomicAnd(1234, 1234);
+    }
+
+    class Dummy {
+     public:
+      __device__ Dummy() {}
+      __device__ ~Dummy() {}
+    };
+
+    __global__ void atomicAnd_uint_v4(Dummy* address, unsigned int* result) {
+      *result = atomicAnd(address, 1234);
+    }
+
+    __global__ void atomicAnd_uint_v5(char* address, unsigned int* result) {
+      *result = atomicAnd(address, 1234);
+    }
+
+    __global__ void atomicAnd_uint_v6(short* address, unsigned int* result) {
+      *result = atomicAnd(address, 1234);
+    }
+
+    __global__ void atomicAnd_uint_v7(long* address, unsigned int* result) {
+      *result = atomicAnd(address, 1234);
+    }
+
+    __global__ void atomicAnd_uint_v8(long long* address, unsigned int* result) {
+      *result = atomicAnd(address, 1234);
+    }
+
+    __global__ void atomicAnd_uint_v9(float* address, unsigned int* result) {
+      *result = atomicAnd(address, 1234);
+    }
+
+    __global__ void atomicAnd_uint_v10(double* address, unsigned int* result) {
+      *result = atomicAnd(address, 1234);
+    }
+  )"};
+
+static constexpr auto kAtomicAnd_ulong{
+    R"(
+    __global__ void atomicAnd_ulong_v1(unsigned long* address, unsigned long* result) {
+      *result = atomicAnd(&address, 1234);
+    }
+
+    __global__ void atomicAnd_ulong_v2(unsigned long* address, unsigned long* result) {
+      *result = atomicAnd(address, address);
+    }
+
+    __global__ void atomicAnd_ulong_v3(unsigned long* address, unsigned long* result) {
+      *result = atomicAnd(1234, 1234);
+    }
+
+    class Dummy {
+     public:
+      __device__ Dummy() {}
+      __device__ ~Dummy() {}
+    };
+
+    __global__ void atomicAnd_ulong_v4(Dummy* address, unsigned long* result) {
+      *result = atomicAnd(address, 1234);
+    }
+
+    __global__ void atomicAnd_ulong_v5(char* address, unsigned long* result) {
+      *result = atomicAnd(address, 1234);
+    }
+
+    __global__ void atomicAnd_ulong_v6(short* address, unsigned long* result) {
+      *result = atomicAnd(address, 1234);
+    }
+
+    __global__ void atomicAnd_ulong_v7(long* address, unsigned long* result) {
+      *result = atomicAnd(address, 1234);
+    }
+
+    __global__ void atomicAnd_ulong_v8(long long* address, unsigned long* result) {
+      *result = atomicAnd(address, 1234);
+    }
+
+    __global__ void atomicAnd_ulong_v9(float* address, unsigned long* result) {
+      *result = atomicAnd(address, 1234);
+    }
+
+    __global__ void atomicAnd_ulong_v10(double* address, unsigned long* result) {
+      *result = atomicAnd(address, 1234);
+    }
+  )"};
+
+static constexpr auto kAtomicAnd_ulonglong{
+    R"(
+    __global__ void atomicAnd_ulonglong_v1(unsigned long long* address, unsigned long long* result) {
+      *result = atomicAnd(&address, 1234);
+    }
+
+    __global__ void atomicAnd_ulonglong_v2(unsigned long long* address, unsigned long long* result) {
+      *result = atomicAnd(address, address);
+    }
+
+    __global__ void atomicAnd_ulonglong_v3(unsigned long long* address, unsigned long long* result) {
+      *result = atomicAnd(1234, 1234);
+    }
+
+    class Dummy {
+     public:
+      __device__ Dummy() {}
+      __device__ ~Dummy() {}
+    };
+
+    __global__ void atomicAnd_ulonglong_v4(Dummy* address, unsigned long long* result) {
+      *result = atomicAnd(address, 1234);
+    }
+
+    __global__ void atomicAnd_ulonglong_v5(char* address, unsigned long long* result) {
+      *result = atomicAnd(address, 1234);
+    }
+
+    __global__ void atomicAnd_ulonglong_v6(short* address, unsigned long long* result) {
+      *result = atomicAnd(address, 1234);
+    }
+
+    __global__ void atomicAnd_ulonglong_v7(long* address, unsigned long long* result) {
+      *result = atomicAnd(address, 1234);
+    }
+
+    __global__ void atomicAnd_ulonglong_v8(long long* address, unsigned long long* result) {
+      *result = atomicAnd(address, 1234);
+    }
+
+    __global__ void atomicAnd_ulonglong_v9(float* address, unsigned long long* result) {
+      *result = atomicAnd(address, 1234);
+    }
+
+    __global__ void atomicAnd_ulonglong_v10(double* address, unsigned long long* result) {
+      *result = atomicAnd(address, 1234);
+    }
+  )"};
diff --git a/projects/hip-tests/catch/unit/atomics/atomicAnd_system.cc b/projects/hip-tests/catch/unit/atomics/atomicAnd_system.cc
new file mode 100644
index 0000000000..e696a8ac26
--- /dev/null
+++ b/projects/hip-tests/catch/unit/atomics/atomicAnd_system.cc
@@ -0,0 +1,109 @@
+/*
+Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved.
+
+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.
+*/
+
+#include "bitwise_common.hh"
+
+#include <hip_test_common.hh>
+
+/**
+ * @addtogroup atomicAnd_system atomicAnd_system
+ * @{
+ * @ingroup AtomicsTest
+ * `atomicAnd_system(TestType* address, TestType* val)` -
+ * performs system-wide atomic bitwise AND between address and val, returns old value.
+ */
+
+/**
+ * Test Description
+ * ------------------------
+ *  - Performs atomicAnd_system from multiple threads on the same address.
+ *  - Uses multiple devices and launches multiple kernels.
+ * Test source
+ * ------------------------
+ *  - unit/atomics/atomicAnd_system.cc
+ * Test requirements
+ * ------------------------
+ *  - Multi-device
+ *  - HIP_VERSION >= 5.2
+ */
+TEMPLATE_TEST_CASE("Unit_atomicAnd_system_Positive_Peer_GPUs_Same_Address", "", int, unsigned int,
+                   unsigned long, unsigned long long) {
+  for (auto current = 0; current < cmd_options.iterations; ++current) {
+    DYNAMIC_SECTION("Same address " << current) {
+      Bitwise::MultipleDeviceMultipleKernelTest<TestType, Bitwise::AtomicOperation::kAndSystem>(
+          2, 2, 1, sizeof(TestType));
+    }
+  }
+}
+
+/**
+ * Test Description
+ * ------------------------
+ *  - Performs atomicAnd_system from multiple threads on adjacent addresses.
+ *  - Uses multiple devices and launches multiple kernels.
+ * Test source
+ * ------------------------
+ *  - unit/atomics/atomicAnd_system.cc
+ * Test requirements
+ * ------------------------
+ *  - Multi-device
+ *  - HIP_VERSION >= 5.2
+ */
+TEMPLATE_TEST_CASE("Unit_atomicAnd_system_Positive_Peer_GPUs_Adjacent_Addresses", "", int,
+                   unsigned int, unsigned long, unsigned long long) {
+  int warp_size = 0;
+  HIP_CHECK(hipDeviceGetAttribute(&warp_size, hipDeviceAttributeWarpSize, 0));
+
+  for (auto current = 0; current < cmd_options.iterations; ++current) {
+    DYNAMIC_SECTION("Adjacent address " << current) {
+      Bitwise::MultipleDeviceMultipleKernelTest<TestType, Bitwise::AtomicOperation::kAndSystem>(
+          2, 2, warp_size, sizeof(TestType));
+    }
+  }
+}
+
+/**
+ * Test Description
+ * ------------------------
+ *  - Performs atomicAnd_system from multiple threads on scattered addresses.
+ *  - Uses multiple devices and launches multiple kernels.
+ * Test source
+ * ------------------------
+ *  - unit/atomics/atomicAnd_system.cc
+ * Test requirements
+ * ------------------------
+ *  - Multi-device
+ *  - HIP_VERSION >= 5.2
+ */
+TEMPLATE_TEST_CASE("Unit_atomicAnd_system_Positive_Peer_GPUs_Scattered_Addresses", "", int,
+                   unsigned int, unsigned long, unsigned long long) {
+  int warp_size = 0;
+  HIP_CHECK(hipDeviceGetAttribute(&warp_size, hipDeviceAttributeWarpSize, 0));
+  const auto cache_line_size = 128u;
+
+  for (auto current = 0; current < cmd_options.iterations; ++current) {
+    DYNAMIC_SECTION("Scattered address " << current) {
+      Bitwise::MultipleDeviceMultipleKernelTest<TestType, Bitwise::AtomicOperation::kAndSystem>(
+          2, 2, warp_size, cache_line_size);
+    }
+  }
+}
diff --git a/projects/hip-tests/catch/unit/atomics/atomicOr.cc b/projects/hip-tests/catch/unit/atomics/atomicOr.cc
new file mode 100644
index 0000000000..e2ae9c6825
--- /dev/null
+++ b/projects/hip-tests/catch/unit/atomics/atomicOr.cc
@@ -0,0 +1,222 @@
+/*
+Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved.
+
+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.
+*/
+
+#include "atomicOr_negative_kernels_rtc.hh"
+#include "bitwise_common.hh"
+
+#include <hip_test_common.hh>
+
+/**
+ * @addtogroup atomicOr atomicOr
+ * @{
+ * @ingroup AtomicsTest
+ * `atomicOr(TestType* address, TestType* val)` -
+ * performs atomic bitwise OR between address and val, returns old value.
+ */
+
+/**
+ * Test Description
+ * ------------------------
+ *  - Performs atomicOr from multiple threads on the same address.
+ *  - Uses only one device and launches one kernel.
+ * Test source
+ * ------------------------
+ *  - unit/atomics/atomicOr.cc
+ * Test requirements
+ * ------------------------
+ *  - HIP_VERSION >= 5.2
+ */
+TEMPLATE_TEST_CASE("Unit_atomicOr_Positive_SameAddress", "", int, unsigned int, unsigned long,
+                   unsigned long long) {
+  for (auto current = 0; current < cmd_options.iterations; ++current) {
+    DYNAMIC_SECTION("Same address " << current) {
+      Bitwise::SingleDeviceSingleKernelTest<TestType, Bitwise::AtomicOperation::kOr>(
+          1, sizeof(TestType));
+    }
+  }
+}
+
+/**
+ * Test Description
+ * ------------------------
+ *  - Performs atomicOr from multiple threads on adjacent addresses.
+ *  - Uses only one device and launches one kernel.
+ * Test source
+ * ------------------------
+ *  - unit/atomics/atomicOr.cc
+ * Test requirements
+ * ------------------------
+ *  - HIP_VERSION >= 5.2
+ */
+TEMPLATE_TEST_CASE("Unit_atomicOr_Positive_Adjacent_Addresses", "", int, unsigned int,
+                   unsigned long, unsigned long long) {
+  int warp_size = 0;
+  HIP_CHECK(hipDeviceGetAttribute(&warp_size, hipDeviceAttributeWarpSize, 0));
+
+  for (auto current = 0; current < cmd_options.iterations; ++current) {
+    DYNAMIC_SECTION("Adjacent address " << current) {
+      Bitwise::SingleDeviceSingleKernelTest<TestType, Bitwise::AtomicOperation::kOr>(
+          warp_size, sizeof(TestType));
+    }
+  }
+}
+
+/**
+ * Test Description
+ * ------------------------
+ *  - Performs atomicOr from multiple threads on the scattered addresses.
+ *  - Uses only one device and launches one kernel.
+ * Test source
+ * ------------------------
+ *  - unit/atomics/atomicOr.cc
+ * Test requirements
+ * ------------------------
+ *  - HIP_VERSION >= 5.2
+ */
+TEMPLATE_TEST_CASE("Unit_atomicOr_Positive_Scattered_Addresses", "", int, unsigned int,
+                   unsigned long, unsigned long long) {
+  int warp_size = 0;
+  HIP_CHECK(hipDeviceGetAttribute(&warp_size, hipDeviceAttributeWarpSize, 0));
+  const auto cache_line_size = 128u;
+
+  for (auto current = 0; current < cmd_options.iterations; ++current) {
+    DYNAMIC_SECTION("Scattered address " << current) {
+      Bitwise::SingleDeviceSingleKernelTest<TestType, Bitwise::AtomicOperation::kOr>(
+          warp_size, cache_line_size);
+    }
+  }
+}
+
+/**
+ * Test Description
+ * ------------------------
+ *  - Performs atomicOr from multiple threads on the same address.
+ *  - Uses only one device and launches multiple kernels.
+ * Test source
+ * ------------------------
+ *  - unit/atomics/atomicOr.cc
+ * Test requirements
+ * ------------------------
+ *  - HIP_VERSION >= 5.2
+ */
+TEMPLATE_TEST_CASE("Unit_atomicOr_Positive_Multi_Kernel_Same_Address", "", int, unsigned int,
+                   unsigned long, unsigned long long) {
+  for (auto current = 0; current < cmd_options.iterations; ++current) {
+    DYNAMIC_SECTION("Same address " << current) {
+      Bitwise::SingleDeviceMultipleKernelTest<TestType, Bitwise::AtomicOperation::kOr>(
+          2, 1, sizeof(TestType));
+    }
+  }
+}
+
+/**
+ * Test Description
+ * ------------------------
+ *  - Performs atomicOr from multiple threads on adjacent addresses.
+ *  - Uses only one device and launches multiple kernels.
+ * Test source
+ * ------------------------
+ *  - unit/atomics/atomicOr.cc
+ * Test requirements
+ * ------------------------
+ *  - HIP_VERSION >= 5.2
+ */
+TEMPLATE_TEST_CASE("Unit_atomicOr_Positive_Multi_Kernel_Adjacent_Addresses", "", int, unsigned int,
+                   unsigned long, unsigned long long) {
+  int warp_size = 0;
+  HIP_CHECK(hipDeviceGetAttribute(&warp_size, hipDeviceAttributeWarpSize, 0));
+
+  for (auto current = 0; current < cmd_options.iterations; ++current) {
+    DYNAMIC_SECTION("Adjacent address " << current) {
+      Bitwise::SingleDeviceMultipleKernelTest<TestType, Bitwise::AtomicOperation::kOr>(
+          2, warp_size, sizeof(TestType));
+    }
+  }
+}
+
+/**
+ * Test Description
+ * ------------------------
+ *  - Performs atomicOr from multiple threads on the scattered addresses.
+ *  - Uses only one device and launches multiple kernels.
+ * Test source
+ * ------------------------
+ *  - unit/atomics/atomicOr.cc
+ * Test requirements
+ * ------------------------
+ *  - HIP_VERSION >= 5.2
+ */
+TEMPLATE_TEST_CASE("Unit_atomicOr_Positive_Multi_Kernel_Scattered_Addresses", "", int, unsigned int,
+                   unsigned long, unsigned long long) {
+  int warp_size = 0;
+  HIP_CHECK(hipDeviceGetAttribute(&warp_size, hipDeviceAttributeWarpSize, 0));
+  const auto cache_line_size = 128u;
+
+  for (auto current = 0; current < cmd_options.iterations; ++current) {
+    DYNAMIC_SECTION("Scattered address " << current) {
+      Bitwise::SingleDeviceMultipleKernelTest<TestType, Bitwise::AtomicOperation::kOr>(
+          2, warp_size, cache_line_size);
+    }
+  }
+}
+
+/**
+ * Test Description
+ * ------------------------
+ *  - Compiles atomicAnd with invalid parameters.
+ *  - Compiles the source with RTC.
+ * Test source
+ * ------------------------
+ *  - unit/atomics/atomicOr.cc
+ * Test requirements
+ * ------------------------
+ *  - HIP_VERSION >= 5.2
+ */
+TEST_CASE("Unit_atomicOr_Negative_Parameters_RTC") {
+  hiprtcProgram program{};
+
+  const auto program_source =
+      GENERATE(kAtomicOr_int, kAtomicOr_uint, kAtomicOr_ulong, kAtomicOr_ulonglong);
+  HIPRTC_CHECK(
+      hiprtcCreateProgram(&program, program_source, "atomicOr_negative.cc", 0, nullptr, nullptr));
+  hiprtcResult result{hiprtcCompileProgram(program, 0, nullptr)};
+
+  // Get the compile log and count compiler error messages
+  size_t log_size{};
+  HIPRTC_CHECK(hiprtcGetProgramLogSize(program, &log_size));
+  std::string log(log_size, ' ');
+  HIPRTC_CHECK(hiprtcGetProgramLog(program, log.data()));
+  int error_count{0};
+  // Please check the content of negative_kernels_rtc.hh
+  int expected_error_count{9};
+  std::string error_message{"error:"};
+
+  size_t n_pos = log.find(error_message, 0);
+  while (n_pos != std::string::npos) {
+    ++error_count;
+    n_pos = log.find(error_message, n_pos + 1);
+  }
+
+  HIPRTC_CHECK(hiprtcDestroyProgram(&program));
+  HIPRTC_CHECK_ERROR(result, HIPRTC_ERROR_COMPILATION);
+  REQUIRE(error_count == expected_error_count);
+}
diff --git a/projects/hip-tests/catch/unit/atomics/atomicOr_negative_kernels.cc b/projects/hip-tests/catch/unit/atomics/atomicOr_negative_kernels.cc
new file mode 100644
index 0000000000..47a56e3a0d
--- /dev/null
+++ b/projects/hip-tests/catch/unit/atomics/atomicOr_negative_kernels.cc
@@ -0,0 +1,177 @@
+/*
+Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved.
+
+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.
+*/
+
+#include <hip_test_common.hh>
+
+class Dummy {
+ public:
+  __device__ Dummy() {}
+  __device__ ~Dummy() {}
+};
+
+/* int atomicOr(int* address, int val) */
+__global__ void atomicOr_int_v1(int* address, int* result) { *result = atomicOr(&address, 1234); }
+
+__global__ void atomicOr_int_v2(int* address, int* result) { *result = atomicOr(address, address); }
+
+__global__ void atomicOr_int_v3(int* address, int* result) { *result = atomicOr(1234, 1234); }
+
+__global__ void atomicOr_int_v4(Dummy* address, int* result) { *result = atomicOr(address, 1234); }
+
+__global__ void atomicOr_int_v5(char* address, int* result) { *result = atomicOr(address, 1234); }
+
+__global__ void atomicOr_int_v6(short* address, int* result) { *result = atomicOr(address, 1234); }
+
+__global__ void atomicOr_int_v7(long* address, int* result) { *result = atomicOr(address, 1234); }
+
+__global__ void atomicOr_int_v8(long long* address, int* result) {
+  *result = atomicOr(address, 1234);
+}
+
+__global__ void atomicOr_int_v9(float* address, int* result) { *result = atomicOr(address, 1234); }
+
+__global__ void atomicOr_int_v10(double* address, int* result) {
+  *result = atomicOr(address, 1234);
+}
+
+/* unsigned int atomicOr(unsigned int* address, unsigned int val) */
+__global__ void atomicOr_uint_v1(unsigned int* address, unsigned int* result) {
+  *result = atomicOr(&address, 1234);
+}
+
+__global__ void atomicOr_uint_v2(unsigned int* address, unsigned int* result) {
+  *result = atomicOr(address, address);
+}
+
+__global__ void atomicOr_uint_v3(unsigned int* address, unsigned int* result) {
+  *result = atomicOr(1234, 1234);
+}
+
+__global__ void atomicOr_uint_v4(Dummy* address, unsigned int* result) {
+  *result = atomicOr(address, 1234);
+}
+
+__global__ void atomicOr_uint_v5(char* address, unsigned int* result) {
+  *result = atomicOr(address, 1234);
+}
+
+__global__ void atomicOr_uint_v6(short* address, unsigned int* result) {
+  *result = atomicOr(address, 1234);
+}
+
+__global__ void atomicOr_uint_v7(long* address, unsigned int* result) {
+  *result = atomicOr(address, 1234);
+}
+
+__global__ void atomicOr_uint_v8(long long* address, unsigned int* result) {
+  *result = atomicOr(address, 1234);
+}
+
+__global__ void atomicOr_uint_v9(float* address, unsigned int* result) {
+  *result = atomicOr(address, 1234);
+}
+
+__global__ void atomicOr_uint_v10(double* address, unsigned int* result) {
+  *result = atomicOr(address, 1234);
+}
+
+/* atomicOr(unsigned long* address, unsigned long val) */
+__global__ void atomicOr_ulong_v1(unsigned long* address, unsigned long* result) {
+  *result = atomicOr(&address, 1234);
+}
+
+__global__ void atomicOr_ulong_v2(unsigned long* address, unsigned long* result) {
+  *result = atomicOr(address, address);
+}
+
+__global__ void atomicOr_ulong_v3(unsigned long* address, unsigned long* result) {
+  *result = atomicOr(1234, 1234);
+}
+
+__global__ void atomicOr_ulong_v4(Dummy* address, unsigned long* result) {
+  *result = atomicOr(address, 1234);
+}
+
+__global__ void atomicOr_ulong_v5(char* address, unsigned long* result) {
+  *result = atomicOr(address, 1234);
+}
+
+__global__ void atomicOr_ulong_v6(short* address, unsigned long* result) {
+  *result = atomicOr(address, 1234);
+}
+
+__global__ void atomicOr_ulong_v7(long* address, unsigned long* result) {
+  *result = atomicOr(address, 1234);
+}
+
+__global__ void atomicOr_ulong_v8(long long* address, unsigned long* result) {
+  *result = atomicOr(address, 1234);
+}
+
+__global__ void atomicOr_ulong_v9(float* address, unsigned long* result) {
+  *result = atomicOr(address, 1234);
+}
+
+__global__ void atomicOr_ulong_v10(double* address, unsigned long* result) {
+  *result = atomicOr(address, 1234);
+}
+
+/* atomicOr(unsigned long long* address, unsigned long long val) */
+__global__ void atomicOr_ulonglong_v1(unsigned long long* address, unsigned long long* result) {
+  *result = atomicOr(&address, 1234);
+}
+
+__global__ void atomicOr_ulonglong_v2(unsigned long long* address, unsigned long long* result) {
+  *result = atomicOr(address, address);
+}
+
+__global__ void atomicOr_ulonglong_v3(unsigned long long* address, unsigned long long* result) {
+  *result = atomicOr(1234, 1234);
+}
+
+__global__ void atomicOr_ulonglong_v4(Dummy* address, unsigned long long* result) {
+  *result = atomicOr(address, 1234);
+}
+
+__global__ void atomicOr_ulonglong_v5(char* address, unsigned long long* result) {
+  *result = atomicOr(address, 1234);
+}
+
+__global__ void atomicOr_ulonglong_v6(short* address, unsigned long long* result) {
+  *result = atomicOr(address, 1234);
+}
+
+__global__ void atomicOr_ulonglong_v7(long* address, unsigned long long* result) {
+  *result = atomicOr(address, 1234);
+}
+
+__global__ void atomicOr_ulonglong_v8(long long* address, unsigned long long* result) {
+  *result = atomicOr(address, 1234);
+}
+
+__global__ void atomicOr_ulonglong_v9(float* address, unsigned long long* result) {
+  *result = atomicOr(address, 1234);
+}
+
+__global__ void atomicOr_ulonglong_v10(double* address, unsigned long long* result) {
+  *result = atomicOr(address, 1234);
+}
diff --git a/projects/hip-tests/catch/unit/atomics/atomicOr_negative_kernels_rtc.hh b/projects/hip-tests/catch/unit/atomics/atomicOr_negative_kernels_rtc.hh
new file mode 100644
index 0000000000..dd4117e704
--- /dev/null
+++ b/projects/hip-tests/catch/unit/atomics/atomicOr_negative_kernels_rtc.hh
@@ -0,0 +1,223 @@
+/*
+Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved.
+
+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.
+*/
+
+#pragma once
+
+/*
+Negative kernels used for the atomics negative Test Cases that are using RTC.
+*/
+
+static constexpr auto kAtomicOr_int{
+    R"(
+    __global__ void atomicOr_int_v1(int* address, int* result) {
+      *result = atomicOr(&address, 1234);
+    }
+
+    __global__ void atomicOr_int_v2(int* address, int* result) {
+      *result = atomicOr(address, address);
+    }
+
+    __global__ void atomicOr_int_v3(int* address, int* result) {
+      *result = atomicOr(1234, 1234);
+    }
+
+    class Dummy {
+     public:
+      __device__ Dummy() {}
+      __device__ ~Dummy() {}
+    };
+
+    __global__ void atomicOr_int_v4(Dummy* address, int* result) {
+      *result = atomicOr(address, 1234);
+    }
+
+    __global__ void atomicOr_int_v5(char* address, int* result) {
+      *result = atomicOr(address, 1234);
+    }
+
+    __global__ void atomicOr_int_v6(short* address, int* result) {
+      *result = atomicOr(address, 1234);
+    }
+
+    __global__ void atomicOr_int_v7(long* address, int* result) {
+      *result = atomicOr(address, 1234);
+    }
+
+    __global__ void atomicOr_int_v8(long long* address, int* result) {
+      *result = atomicOr(address, 1234);
+    }
+
+    __global__ void atomicOr_int_v9(float* address, int* result) {
+      *result = atomicOr(address, 1234);
+    }
+
+    __global__ void atomicOr_int_v10(double* address, int* result) {
+      *result = atomicOr(address, 1234);
+    }
+  )"};
+
+static constexpr auto kAtomicOr_uint{
+    R"(
+    __global__ void atomicOr_uint_v1(unsigned int* address, unsigned int* result) {
+      *result = atomicOr(&address, 1234);
+    }
+
+    __global__ void atomicOr_uint_v2(unsigned int* address, unsigned int* result) {
+      *result = atomicOr(address, address);
+    }
+
+    __global__ void atomicOr_uint_v3(unsigned int* address, unsigned int* result) {
+      *result = atomicOr(1234, 1234);
+    }
+
+    class Dummy {
+     public:
+      __device__ Dummy() {}
+      __device__ ~Dummy() {}
+    };
+
+    __global__ void atomicOr_uint_v4(Dummy* address, unsigned int* result) {
+      *result = atomicOr(address, 1234);
+    }
+
+    __global__ void atomicOr_uint_v5(char* address, unsigned int* result) {
+      *result = atomicOr(address, 1234);
+    }
+
+    __global__ void atomicOr_uint_v6(short* address, unsigned int* result) {
+      *result = atomicOr(address, 1234);
+    }
+
+    __global__ void atomicOr_uint_v7(long* address, unsigned int* result) {
+      *result = atomicOr(address, 1234);
+    }
+
+    __global__ void atomicOr_uint_v8(long long* address, unsigned int* result) {
+      *result = atomicOr(address, 1234);
+    }
+
+    __global__ void atomicOr_uint_v9(float* address, unsigned int* result) {
+      *result = atomicOr(address, 1234);
+    }
+
+    __global__ void atomicOr_uint_v10(double* address, unsigned int* result) {
+      *result = atomicOr(address, 1234);
+    }
+  )"};
+
+static constexpr auto kAtomicOr_ulong{
+    R"(
+    __global__ void atomicOr_ulong_v1(unsigned long* address, unsigned long* result) {
+      *result = atomicOr(&address, 1234);
+    }
+
+    __global__ void atomicOr_ulong_v2(unsigned long* address, unsigned long* result) {
+      *result = atomicOr(address, address);
+    }
+
+    __global__ void atomicOr_ulong_v3(unsigned long* address, unsigned long* result) {
+      *result = atomicOr(1234, 1234);
+    }
+
+    class Dummy {
+     public:
+      __device__ Dummy() {}
+      __device__ ~Dummy() {}
+    };
+
+    __global__ void atomicOr_ulong_v4(Dummy* address, unsigned long* result) {
+      *result = atomicOr(address, 1234);
+    }
+
+    __global__ void atomicOr_ulong_v5(char* address, unsigned long* result) {
+      *result = atomicOr(address, 1234);
+    }
+
+    __global__ void atomicOr_ulong_v6(short* address, unsigned long* result) {
+      *result = atomicOr(address, 1234);
+    }
+
+    __global__ void atomicOr_ulong_v7(long* address, unsigned long* result) {
+      *result = atomicOr(address, 1234);
+    }
+
+    __global__ void atomicOr_ulong_v8(long long* address, unsigned long* result) {
+      *result = atomicOr(address, 1234);
+    }
+
+    __global__ void atomicOr_ulong_v9(float* address, unsigned long* result) {
+      *result = atomicOr(address, 1234);
+    }
+
+    __global__ void atomicOr_ulong_v10(double* address, unsigned long* result) {
+      *result = atomicOr(address, 1234);
+    }
+  )"};
+
+static constexpr auto kAtomicOr_ulonglong{
+    R"(
+    __global__ void atomicOr_ulonglong_v1(unsigned long long* address, unsigned long long* result) {
+      *result = atomicOr(&address, 1234);
+    }
+
+    __global__ void atomicOr_ulonglong_v2(unsigned long long* address, unsigned long long* result) {
+      *result = atomicOr(address, address);
+    }
+
+    __global__ void atomicOr_ulonglong_v3(unsigned long long* address, unsigned long long* result) {
+      *result = atomicOr(1234, 1234);
+    }
+
+    class Dummy {
+     public:
+      __device__ Dummy() {}
+      __device__ ~Dummy() {}
+    };
+
+    __global__ void atomicOr_ulonglong_v4(Dummy* address, unsigned long long* result) {
+      *result = atomicOr(address, 1234);
+    }
+
+    __global__ void atomicOr_ulonglong_v5(char* address, unsigned long long* result) {
+      *result = atomicOr(address, 1234);
+    }
+
+    __global__ void atomicOr_ulonglong_v6(short* address, unsigned long long* result) {
+      *result = atomicOr(address, 1234);
+    }
+
+    __global__ void atomicOr_ulonglong_v7(long* address, unsigned long long* result) {
+      *result = atomicOr(address, 1234);
+    }
+
+    __global__ void atomicOr_ulonglong_v8(long long* address, unsigned long long* result) {
+      *result = atomicOr(address, 1234);
+    }
+
+    __global__ void atomicOr_ulonglong_v9(float* address, unsigned long long* result) {
+      *result = atomicOr(address, 1234);
+    }
+
+    __global__ void atomicOr_ulonglong_v10(double* address, unsigned long long* result) {
+      *result = atomicOr(address, 1234);
+    }
+  )"};
diff --git a/projects/hip-tests/catch/unit/atomics/atomicOr_system.cc b/projects/hip-tests/catch/unit/atomics/atomicOr_system.cc
new file mode 100644
index 0000000000..0239056e3f
--- /dev/null
+++ b/projects/hip-tests/catch/unit/atomics/atomicOr_system.cc
@@ -0,0 +1,109 @@
+/*
+Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved.
+
+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.
+*/
+
+#include "bitwise_common.hh"
+
+#include <hip_test_common.hh>
+
+/**
+ * @addtogroup atomicOr_system atomicOr_system
+ * @{
+ * @ingroup AtomicsTest
+ * `atomicOr_system(TestType* address, TestType* val)` -
+ * performs system-wide atomic bitwise OR between address and val, returns old value.
+ */
+
+/**
+ * Test Description
+ * ------------------------
+ *  - Performs atomicOr_system from multiple threads on the same address.
+ *  - Uses multiple devices and launches multiple kernels.
+ * Test source
+ * ------------------------
+ *  - unit/atomics/atomicOr_system.cc
+ * Test requirements
+ * ------------------------
+ *  - Multi-device
+ *  - HIP_VERSION >= 5.2
+ */
+TEMPLATE_TEST_CASE("Unit_atomicOr_system_Positive_Peer_GPUs_Same_Address", "", int, unsigned int,
+                   unsigned long, unsigned long long) {
+  for (auto current = 0; current < cmd_options.iterations; ++current) {
+    DYNAMIC_SECTION("Same address " << current) {
+      Bitwise::MultipleDeviceMultipleKernelTest<TestType, Bitwise::AtomicOperation::kOrSystem>(
+          2, 2, 1, sizeof(TestType));
+    }
+  }
+}
+
+/**
+ * Test Description
+ * ------------------------
+ *  - Performs atomicOr_system from multiple threads on adjacent addresses.
+ *  - Uses multiple devices and launches multiple kernels.
+ * Test source
+ * ------------------------
+ *  - unit/atomics/atomicOr_system.cc
+ * Test requirements
+ * ------------------------
+ *  - Multi-device
+ *  - HIP_VERSION >= 5.2
+ */
+TEMPLATE_TEST_CASE("Unit_atomicOr_system_Positive_Peer_GPUs_Adjacent_Addresses", "", int,
+                   unsigned int, unsigned long, unsigned long long) {
+  int warp_size = 0;
+  HIP_CHECK(hipDeviceGetAttribute(&warp_size, hipDeviceAttributeWarpSize, 0));
+
+  for (auto current = 0; current < cmd_options.iterations; ++current) {
+    DYNAMIC_SECTION("Adjacent address " << current) {
+      Bitwise::MultipleDeviceMultipleKernelTest<TestType, Bitwise::AtomicOperation::kOrSystem>(
+          2, 2, warp_size, sizeof(TestType));
+    }
+  }
+}
+
+/**
+ * Test Description
+ * ------------------------
+ *  - Performs atomicOr_system from multiple threads on scattered addresses.
+ *  - Uses multiple devices and launches multiple kernels.
+ * Test source
+ * ------------------------
+ *  - unit/atomics/atomicOr_system.cc
+ * Test requirements
+ * ------------------------
+ *  - Multi-device
+ *  - HIP_VERSION >= 5.2
+ */
+TEMPLATE_TEST_CASE("Unit_atomicOr_system_Positive_Peer_GPUs_Scattered_Addresses", "", int,
+                   unsigned int, unsigned long, unsigned long long) {
+  int warp_size = 0;
+  HIP_CHECK(hipDeviceGetAttribute(&warp_size, hipDeviceAttributeWarpSize, 0));
+  const auto cache_line_size = 128u;
+
+  for (auto current = 0; current < cmd_options.iterations; ++current) {
+    DYNAMIC_SECTION("Scattered address " << current) {
+      Bitwise::MultipleDeviceMultipleKernelTest<TestType, Bitwise::AtomicOperation::kOrSystem>(
+          2, 2, warp_size, cache_line_size);
+    }
+  }
+}
diff --git a/projects/hip-tests/catch/unit/atomics/atomicXor.cc b/projects/hip-tests/catch/unit/atomics/atomicXor.cc
new file mode 100644
index 0000000000..0fb31252c5
--- /dev/null
+++ b/projects/hip-tests/catch/unit/atomics/atomicXor.cc
@@ -0,0 +1,222 @@
+/*
+Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved.
+
+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.
+*/
+
+#include "atomicXor_negative_kernels_rtc.hh"
+#include "bitwise_common.hh"
+
+#include <hip_test_common.hh>
+
+/**
+ * @addtogroup atomicXor atomicXor
+ * @{
+ * @ingroup AtomicsTest
+ * `atomicXor(TestType* address, TestType* val)` -
+ * performs atomic bitwise XOR between address and val, returns old value.
+ */
+
+/**
+ * Test Description
+ * ------------------------
+ *  - Performs atomicXor from multiple threads on the same address.
+ *  - Uses only one device and launches one kernel.
+ * Test source
+ * ------------------------
+ *  - unit/atomics/atomicXor.cc
+ * Test requirements
+ * ------------------------
+ *  - HIP_VERSION >= 5.2
+ */
+TEMPLATE_TEST_CASE("Unit_atomicXor_Positive_SameAddress", "", int, unsigned int, unsigned long,
+                   unsigned long long) {
+  for (auto current = 0; current < cmd_options.iterations; ++current) {
+    DYNAMIC_SECTION("Same address " << current) {
+      Bitwise::SingleDeviceSingleKernelTest<TestType, Bitwise::AtomicOperation::kXor>(
+          1, sizeof(TestType));
+    }
+  }
+}
+
+/**
+ * Test Description
+ * ------------------------
+ *  - Performs atomicXor from multiple threads on adjacent addresses.
+ *  - Uses only one device and launches one kernel.
+ * Test source
+ * ------------------------
+ *  - unit/atomics/atomicXor.cc
+ * Test requirements
+ * ------------------------
+ *  - HIP_VERSION >= 5.2
+ */
+TEMPLATE_TEST_CASE("Unit_atomicXor_Positive_Adjacent_Addresses", "", int, unsigned int,
+                   unsigned long, unsigned long long) {
+  int warp_size = 0;
+  HIP_CHECK(hipDeviceGetAttribute(&warp_size, hipDeviceAttributeWarpSize, 0));
+
+  for (auto current = 0; current < cmd_options.iterations; ++current) {
+    DYNAMIC_SECTION("Adjacent address " << current) {
+      Bitwise::SingleDeviceSingleKernelTest<TestType, Bitwise::AtomicOperation::kXor>(
+          warp_size, sizeof(TestType));
+    }
+  }
+}
+
+/**
+ * Test Description
+ * ------------------------
+ *  - Performs atomicXor from multiple threads on the scattered addresses.
+ *  - Uses only one device and launches one kernel.
+ * Test source
+ * ------------------------
+ *  - unit/atomics/atomicXor.cc
+ * Test requirements
+ * ------------------------
+ *  - HIP_VERSION >= 5.2
+ */
+TEMPLATE_TEST_CASE("Unit_atomicXor_Positive_Scattered_Addresses", "", int, unsigned int,
+                   unsigned long, unsigned long long) {
+  int warp_size = 0;
+  HIP_CHECK(hipDeviceGetAttribute(&warp_size, hipDeviceAttributeWarpSize, 0));
+  const auto cache_line_size = 128u;
+
+  for (auto current = 0; current < cmd_options.iterations; ++current) {
+    DYNAMIC_SECTION("Scattered address " << current) {
+      Bitwise::SingleDeviceSingleKernelTest<TestType, Bitwise::AtomicOperation::kXor>(
+          warp_size, cache_line_size);
+    }
+  }
+}
+
+/**
+ * Test Description
+ * ------------------------
+ *  - Performs atomicXor from multiple threads on the same address.
+ *  - Uses only one device and launches multiple kernels.
+ * Test source
+ * ------------------------
+ *  - unit/atomics/atomicXor.cc
+ * Test requirements
+ * ------------------------
+ *  - HIP_VERSION >= 5.2
+ */
+TEMPLATE_TEST_CASE("Unit_atomicXor_Positive_Multi_Kernel_Same_Address", "", int, unsigned int,
+                   unsigned long, unsigned long long) {
+  for (auto current = 0; current < cmd_options.iterations; ++current) {
+    DYNAMIC_SECTION("Same address " << current) {
+      Bitwise::SingleDeviceMultipleKernelTest<TestType, Bitwise::AtomicOperation::kXor>(
+          2, 1, sizeof(TestType));
+    }
+  }
+}
+
+/**
+ * Test Description
+ * ------------------------
+ *  - Performs atomicXor from multiple threads on adjacent addresses.
+ *  - Uses only one device and launches multiple kernels.
+ * Test source
+ * ------------------------
+ *  - unit/atomics/atomicXor.cc
+ * Test requirements
+ * ------------------------
+ *  - HIP_VERSION >= 5.2
+ */
+TEMPLATE_TEST_CASE("Unit_atomicXor_Positive_Multi_Kernel_Adjacent_Addresses", "", int, unsigned int,
+                   unsigned long, unsigned long long) {
+  int warp_size = 0;
+  HIP_CHECK(hipDeviceGetAttribute(&warp_size, hipDeviceAttributeWarpSize, 0));
+
+  for (auto current = 0; current < cmd_options.iterations; ++current) {
+    DYNAMIC_SECTION("Adjacent address " << current) {
+      Bitwise::SingleDeviceMultipleKernelTest<TestType, Bitwise::AtomicOperation::kXor>(
+          2, warp_size - 1, sizeof(TestType));
+    }
+  }
+}
+
+/**
+ * Test Description
+ * ------------------------
+ *  - Performs atomicXor from multiple threads on the scattered addresses.
+ *  - Uses only one device and launches multiple kernels.
+ * Test source
+ * ------------------------
+ *  - unit/atomics/atomicXor.cc
+ * Test requirements
+ * ------------------------
+ *  - HIP_VERSION >= 5.2
+ */
+TEMPLATE_TEST_CASE("Unit_atomicXor_Positive_Multi_Kernel_Scattered_Addresses", "", int,
+                   unsigned int, unsigned long, unsigned long long) {
+  int warp_size = 0;
+  HIP_CHECK(hipDeviceGetAttribute(&warp_size, hipDeviceAttributeWarpSize, 0));
+  const auto cache_line_size = 128u;
+
+  for (auto current = 0; current < cmd_options.iterations; ++current) {
+    DYNAMIC_SECTION("Scattered address " << current) {
+      Bitwise::SingleDeviceMultipleKernelTest<TestType, Bitwise::AtomicOperation::kXor>(
+          2, warp_size - 1, cache_line_size);
+    }
+  }
+}
+
+/**
+ * Test Description
+ * ------------------------
+ *  - Compiles atomicXor with invalid parameters.
+ *  - Compiles the source with RTC.
+ * Test source
+ * ------------------------
+ *  - unit/atomics/atomicXor.cc
+ * Test requirements
+ * ------------------------
+ *  - HIP_VERSION >= 5.2
+ */
+TEST_CASE("Unit_atomicXor_Negative_Parameters_RTC") {
+  hiprtcProgram program{};
+
+  const auto program_source =
+      GENERATE(kAtomicXor_int, kAtomicXor_uint, kAtomicXor_ulong, kAtomicXor_ulonglong);
+  HIPRTC_CHECK(
+      hiprtcCreateProgram(&program, program_source, "atomicXor_negative.cc", 0, nullptr, nullptr));
+  hiprtcResult result{hiprtcCompileProgram(program, 0, nullptr)};
+
+  // Get the compile log and count compiler error messages
+  size_t log_size{};
+  HIPRTC_CHECK(hiprtcGetProgramLogSize(program, &log_size));
+  std::string log(log_size, ' ');
+  HIPRTC_CHECK(hiprtcGetProgramLog(program, log.data()));
+  int error_count{0};
+  // Please check the content of negative_kernels_rtc.hh
+  int expected_error_count{9};
+  std::string error_message{"error:"};
+
+  size_t n_pos = log.find(error_message, 0);
+  while (n_pos != std::string::npos) {
+    ++error_count;
+    n_pos = log.find(error_message, n_pos + 1);
+  }
+
+  HIPRTC_CHECK(hiprtcDestroyProgram(&program));
+  HIPRTC_CHECK_ERROR(result, HIPRTC_ERROR_COMPILATION);
+  REQUIRE(error_count == expected_error_count);
+}
diff --git a/projects/hip-tests/catch/unit/atomics/atomicXor_negative_kernels.cc b/projects/hip-tests/catch/unit/atomics/atomicXor_negative_kernels.cc
new file mode 100644
index 0000000000..a180afd6db
--- /dev/null
+++ b/projects/hip-tests/catch/unit/atomics/atomicXor_negative_kernels.cc
@@ -0,0 +1,185 @@
+/*
+Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved.
+
+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.
+*/
+
+#include <hip_test_common.hh>
+
+class Dummy {
+ public:
+  __device__ Dummy() {}
+  __device__ ~Dummy() {}
+};
+
+/* int atomicXor(int* address, int val) */
+__global__ void atomicXor_int_v1(int* address, int* result) { *result = atomicXor(&address, 1234); }
+
+__global__ void atomicXor_int_v2(int* address, int* result) {
+  *result = atomicXor(address, address);
+}
+
+__global__ void atomicXor_int_v3(int* address, int* result) { *result = atomicXor(1234, 1234); }
+
+__global__ void atomicXor_int_v4(Dummy* address, int* result) {
+  *result = atomicXor(address, 1234);
+}
+
+__global__ void atomicXor_int_v5(char* address, int* result) { *result = atomicXor(address, 1234); }
+
+__global__ void atomicXor_int_v6(short* address, int* result) {
+  *result = atomicXor(address, 1234);
+}
+
+__global__ void atomicXor_int_v7(long* address, int* result) { *result = atomicXor(address, 1234); }
+
+__global__ void atomicXor_int_v8(long long* address, int* result) {
+  *result = atomicXor(address, 1234);
+}
+
+__global__ void atomicXor_int_v9(float* address, int* result) {
+  *result = atomicXor(address, 1234);
+}
+
+__global__ void atomicXor_int_v10(double* address, int* result) {
+  *result = atomicXor(address, 1234);
+}
+
+/* unsigned int atomicXor(unsigned int* address, unsigned int val) */
+__global__ void atomicXor_uint_v1(unsigned int* address, unsigned int* result) {
+  *result = atomicXor(&address, 1234);
+}
+
+__global__ void atomicXor_uint_v2(unsigned int* address, unsigned int* result) {
+  *result = atomicXor(address, address);
+}
+
+__global__ void atomicXor_uint_v3(unsigned int* address, unsigned int* result) {
+  *result = atomicXor(1234, 1234);
+}
+
+__global__ void atomicXor_uint_v4(Dummy* address, unsigned int* result) {
+  *result = atomicXor(address, 1234);
+}
+
+__global__ void atomicXor_uint_v5(char* address, unsigned int* result) {
+  *result = atomicXor(address, 1234);
+}
+
+__global__ void atomicXor_uint_v6(short* address, unsigned int* result) {
+  *result = atomicXor(address, 1234);
+}
+
+__global__ void atomicXor_uint_v7(long* address, unsigned int* result) {
+  *result = atomicXor(address, 1234);
+}
+
+__global__ void atomicXor_uint_v8(long long* address, unsigned int* result) {
+  *result = atomicXor(address, 1234);
+}
+
+__global__ void atomicXor_int_v9(float* address, unsigned int* result) {
+  *result = atomicXor(address, 1234);
+}
+
+__global__ void atomicXor_int_v10(double* address, unsigned int* result) {
+  *result = atomicXor(address, 1234);
+}
+
+/* atomicXor(unsigned long* address, unsigned long val) */
+__global__ void atomicXor_ulong_v1(unsigned long* address, unsigned long* result) {
+  *result = atomicXor(&address, 1234);
+}
+
+__global__ void atomicXor_ulong_v2(unsigned long* address, unsigned long* result) {
+  *result = atomicXor(address, address);
+}
+
+__global__ void atomicXor_ulong_v3(unsigned long* address, unsigned long* result) {
+  *result = atomicXor(1234, 1234);
+}
+
+__global__ void atomicXor_ulong_v4(Dummy* address, unsigned long* result) {
+  *result = atomicXor(address, 1234);
+}
+
+__global__ void atomicXor_ulong_v5(char* address, unsigned long* result) {
+  *result = atomicXor(address, 1234);
+}
+
+__global__ void atomicXor_ulong_v6(short* address, unsigned long* result) {
+  *result = atomicXor(address, 1234);
+}
+
+__global__ void atomicXor_ulong_v7(long* address, unsigned long* result) {
+  *result = atomicXor(address, 1234);
+}
+
+__global__ void atomicXor_ulong_v8(long long* address, unsigned long* result) {
+  *result = atomicXor(address, 1234);
+}
+
+__global__ void atomicXor_ulong_v9(float* address, unsigned long* result) {
+  *result = atomicOr(address, 1234);
+}
+
+__global__ void atomicXor_ulong_v10(double* address, unsigned long* result) {
+  *result = atomicOr(address, 1234);
+}
+
+/* atomicXor(unsigned long long* address, unsigned long long val) */
+__global__ void atomicXor_ulonglong_v1(unsigned long long* address, unsigned long long* result) {
+  *result = atomicXor(&address, 1234);
+}
+
+__global__ void atomicXor_ulonglong_v2(unsigned long long* address, unsigned long long* result) {
+  *result = atomicXor(address, address);
+}
+
+__global__ void atomicXor_ulonglong_v3(unsigned long long* address, unsigned long long* result) {
+  *result = atomicXor(1234, 1234);
+}
+
+__global__ void atomicXor_ulonglong_v4(Dummy* address, unsigned long long* result) {
+  *result = atomicXor(address, 1234);
+}
+
+__global__ void atomicXor_ulonglong_v5(char* address, unsigned long long* result) {
+  *result = atomicXor(address, 1234);
+}
+
+__global__ void atomicXor_ulonglong_v6(short* address, unsigned long long* result) {
+  *result = atomicXor(address, 1234);
+}
+
+__global__ void atomicXor_ulonglong_v7(long* address, unsigned long long* result) {
+  *result = atomicXor(address, 1234);
+}
+
+__global__ void atomicXor_ulonglong_v8(long long* address, unsigned long long* result) {
+  *result = atomicXor(address, 1234);
+}
+
+__global__ void atomicOr_ulonglong_v9(float* address, unsigned long long* result) {
+  *result = atomicOr(address, 1234);
+}
+
+__global__ void atomicOr_ulonglong_v10(double* address, unsigned long long* result) {
+  *result = atomicOr(address, 1234);
+}
diff --git a/projects/hip-tests/catch/unit/atomics/atomicXor_negative_kernels_rtc.hh b/projects/hip-tests/catch/unit/atomics/atomicXor_negative_kernels_rtc.hh
new file mode 100644
index 0000000000..3d4e19c7e7
--- /dev/null
+++ b/projects/hip-tests/catch/unit/atomics/atomicXor_negative_kernels_rtc.hh
@@ -0,0 +1,223 @@
+/*
+Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved.
+
+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.
+*/
+
+#pragma once
+
+/*
+Negative kernels used for the atomics negative Test Cases that are using RTC.
+*/
+
+static constexpr auto kAtomicXor_int{
+    R"(
+    __global__ void atomicXor_int_v1(int* address, int* result) {
+      *result = atomicXor(&address, 1234);
+    }
+
+    __global__ void atomicXor_int_v2(int* address, int* result) {
+      *result = atomicXor(address, address);
+    }
+
+    __global__ void atomicXor_int_v3(int* address, int* result) {
+      *result = atomicXor(1234, 1234);
+    }
+
+    class Dummy {
+     public:
+      __device__ Dummy() {}
+      __device__ ~Dummy() {}
+    };
+
+    __global__ void atomicXor_int_v4(Dummy* address, int* result) {
+      *result = atomicXor(address, 1234);
+    }
+
+    __global__ void atomicXor_int_v5(char* address, int* result) {
+      *result = atomicXor(address, 1234);
+    }
+
+    __global__ void atomicXor_int_v6(short* address, int* result) {
+      *result = atomicXor(address, 1234);
+    }
+
+    __global__ void atomicXor_int_v7(long* address, int* result) {
+      *result = atomicXor(address, 1234);
+    }
+
+    __global__ void atomicXor_int_v8(long long* address, int* result) {
+      *result = atomicXor(address, 1234);
+    }
+
+    __global__ void atomicXor_int_v9(float* address, int* result) {
+      *result = atomicXor(address, 1234);
+    }
+
+    __global__ void atomicXor_int_v10(double* address, int* result) {
+      *result = atomicXor(address, 1234);
+    }
+  )"};
+
+static constexpr auto kAtomicXor_uint{
+    R"(
+    __global__ void atomicXor_uint_v1(unsigned int* address, unsigned int* result) {
+      *result = atomicXor(&address, 1234);
+    }
+
+    __global__ void atomicXor_uint_v2(unsigned int* address, unsigned int* result) {
+      *result = atomicXor(address, address);
+    }
+
+    __global__ void atomicXor_uint_v3(unsigned int* address, unsigned int* result) {
+      *result = atomicXor(1234, 1234);
+    }
+
+    class Dummy {
+     public:
+      __device__ Dummy() {}
+      __device__ ~Dummy() {}
+    };
+
+    __global__ void atomicXor_uint_v4(Dummy* address, unsigned int* result) {
+      *result = atomicXor(address, 1234);
+    }
+
+    __global__ void atomicXor_uint_v5(char* address, unsigned int* result) {
+      *result = atomicXor(address, 1234);
+    }
+
+    __global__ void atomicXor_uint_v6(short* address, unsigned int* result) {
+      *result = atomicXor(address, 1234);
+    }
+
+    __global__ void atomicXor_uint_v7(long* address, unsigned int* result) {
+      *result = atomicXor(address, 1234);
+    }
+
+    __global__ void atomicXor_uint_v8(long long* address, unsigned int* result) {
+      *result = atomicXor(address, 1234);
+    }
+
+    __global__ void atomicXor_uint_v9(float* address, unsigned int* result) {
+      *result = atomicXor(address, 1234);
+    }
+
+    __global__ void atomicXor_uint_v10(double* address, unsigned int* result) {
+      *result = atomicXor(address, 1234);
+    }
+  )"};
+
+static constexpr auto kAtomicXor_ulong{
+    R"(
+    __global__ void atomicXor_ulong_v1(unsigned long* address, unsigned long* result) {
+      *result = atomicXor(&address, 1234);
+    }
+
+    __global__ void atomicXor_ulong_v2(unsigned long* address, unsigned long* result) {
+      *result = atomicXor(address, address);
+    }
+
+    __global__ void atomicXor_ulong_v3(unsigned long* address, unsigned long* result) {
+      *result = atomicXor(1234, 1234);
+    }
+
+    class Dummy {
+     public:
+      __device__ Dummy() {}
+      __device__ ~Dummy() {}
+    };
+
+    __global__ void atomicXor_ulong_v4(Dummy* address, unsigned long* result) {
+      *result = atomicXor(address, 1234);
+    }
+
+    __global__ void atomicXor_ulong_v5(char* address, unsigned long* result) {
+      *result = atomicXor(address, 1234);
+    }
+
+    __global__ void atomicXor_ulong_v6(short* address, unsigned long* result) {
+      *result = atomicXor(address, 1234);
+    }
+
+    __global__ void atomicXor_ulong_v7(long* address, unsigned long* result) {
+      *result = atomicXor(address, 1234);
+    }
+
+    __global__ void atomicXor_ulong_v8(long long* address, unsigned long* result) {
+      *result = atomicXor(address, 1234);
+    }
+
+    __global__ void atomicXor_ulong_v9(float* address, unsigned long* result) {
+      *result = atomicXor(address, 1234);
+    }
+
+    __global__ void atomicXor_ulong_v10(double* address, unsigned long* result) {
+      *result = atomicXor(address, 1234);
+    }
+  )"};
+
+static constexpr auto kAtomicXor_ulonglong{
+    R"(
+    __global__ void atomicXor_ulonglong_v1(unsigned long long* address, unsigned long long* result) {
+      *result = atomicXor(&address, 1234);
+    }
+
+    __global__ void atomicXor_ulonglong_v2(unsigned long long* address, unsigned long long* result) {
+      *result = atomicXor(address, address);
+    }
+
+    __global__ void atomicXor_ulonglong_v3(unsigned long long* address, unsigned long long* result) {
+      *result = atomicXor(1234, 1234);
+    }
+
+    class Dummy {
+     public:
+      __device__ Dummy() {}
+      __device__ ~Dummy() {}
+    };
+
+    __global__ void atomicXor_ulonglong_v4(Dummy* address, unsigned long long* result) {
+      *result = atomicXor(address, 1234);
+    }
+
+    __global__ void atomicXor_ulonglong_v5(char* address, unsigned long long* result) {
+      *result = atomicXor(address, 1234);
+    }
+
+    __global__ void atomicXor_ulonglong_v6(short* address, unsigned long long* result) {
+      *result = atomicXor(address, 1234);
+    }
+
+    __global__ void atomicXor_ulonglong_v7(long* address, unsigned long long* result) {
+      *result = atomicXor(address, 1234);
+    }
+
+    __global__ void atomicXor_ulonglong_v8(long long* address, unsigned long long* result) {
+      *result = atomicXor(address, 1234);
+    }
+
+    __global__ void atomicXor_ulonglong_v9(float* address, unsigned long long* result) {
+      *result = atomicXor(address, 1234);
+    }
+
+    __global__ void atomicXor_ulonglong_v10(double* address, unsigned long long* result) {
+      *result = atomicXor(address, 1234);
+    }
+  )"};
diff --git a/projects/hip-tests/catch/unit/atomics/atomicXor_system.cc b/projects/hip-tests/catch/unit/atomics/atomicXor_system.cc
new file mode 100644
index 0000000000..fbfb82d36d
--- /dev/null
+++ b/projects/hip-tests/catch/unit/atomics/atomicXor_system.cc
@@ -0,0 +1,109 @@
+/*
+Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved.
+
+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.
+*/
+
+#include "bitwise_common.hh"
+
+#include <hip_test_common.hh>
+
+/**
+ * @addtogroup atomicXor_system atomicXor_system
+ * @{
+ * @ingroup AtomicsTest
+ * `atomicXor_system(TestType* address, TestType* val)` -
+ * performs system-wide atomic bitwise XOR between address and val, returns old value.
+ */
+
+/**
+ * Test Description
+ * ------------------------
+ *  - Performs atomicXor_system from multiple threads on the same address.
+ *  - Uses multiple devices and launches multiple kernels.
+ * Test source
+ * ------------------------
+ *  - unit/atomics/atomicXor_system.cc
+ * Test requirements
+ * ------------------------
+ *  - Multi-device
+ *  - HIP_VERSION >= 5.2
+ */
+TEMPLATE_TEST_CASE("Unit_atomicXor_system_Positive_Peer_GPUs_Same_Address", "", int, unsigned int,
+                   unsigned long, unsigned long long) {
+  for (auto current = 0; current < cmd_options.iterations; ++current) {
+    DYNAMIC_SECTION("Same address " << current) {
+      Bitwise::MultipleDeviceMultipleKernelTest<TestType, Bitwise::AtomicOperation::kXorSystem>(
+          2, 2, 1, sizeof(TestType));
+    }
+  }
+}
+
+/**
+ * Test Description
+ * ------------------------
+ *  - Performs atomicXor_system from multiple threads on adjacent addresses.
+ *  - Uses multiple devices and launches multiple kernels.
+ * Test source
+ * ------------------------
+ *  - unit/atomics/atomicXor_system.cc
+ * Test requirements
+ * ------------------------
+ *  - Multi-device
+ *  - HIP_VERSION >= 5.2
+ */
+TEMPLATE_TEST_CASE("Unit_atomicXor_system_Positive_Peer_GPUs_Adjacent_Addresses", "", int,
+                   unsigned int, unsigned long, unsigned long long) {
+  int warp_size = 0;
+  HIP_CHECK(hipDeviceGetAttribute(&warp_size, hipDeviceAttributeWarpSize, 0));
+
+  for (auto current = 0; current < cmd_options.iterations; ++current) {
+    DYNAMIC_SECTION("Adjacent address " << current) {
+      Bitwise::MultipleDeviceMultipleKernelTest<TestType, Bitwise::AtomicOperation::kXorSystem>(
+          2, 2, warp_size, sizeof(TestType));
+    }
+  }
+}
+
+/**
+ * Test Description
+ * ------------------------
+ *  - Performs atomicXor_system from multiple threads on scattered addresses.
+ *  - Uses multiple devices and launches multiple kernels.
+ * Test source
+ * ------------------------
+ *  - unit/atomics/atomicXor_system.cc
+ * Test requirements
+ * ------------------------
+ *  - Multi-device
+ *  - HIP_VERSION >= 5.2
+ */
+TEMPLATE_TEST_CASE("Unit_atomicXor_system_Positive_Peer_GPUs_Scattered_Addresses", "", int,
+                   unsigned int, unsigned long, unsigned long long) {
+  int warp_size = 0;
+  HIP_CHECK(hipDeviceGetAttribute(&warp_size, hipDeviceAttributeWarpSize, 0));
+  const auto cache_line_size = 128u;
+
+  for (auto current = 0; current < cmd_options.iterations; ++current) {
+    DYNAMIC_SECTION("Scattered address " << current) {
+      Bitwise::MultipleDeviceMultipleKernelTest<TestType, Bitwise::AtomicOperation::kXorSystem>(
+          2, 2, warp_size, cache_line_size);
+    }
+  }
+}
diff --git a/projects/hip-tests/catch/unit/atomics/bitwise_common.hh b/projects/hip-tests/catch/unit/atomics/bitwise_common.hh
new file mode 100644
index 0000000000..9e71c99cb0
--- /dev/null
+++ b/projects/hip-tests/catch/unit/atomics/bitwise_common.hh
@@ -0,0 +1,345 @@
+/*
+Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved.
+
+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.
+*/
+
+#pragma once
+
+#include <hip_test_common.hh>
+#include <hip/hip_cooperative_groups.h>
+#include <resource_guards.hh>
+#include <cmd_options.hh>
+
+namespace cg = cooperative_groups;
+
+namespace Bitwise {
+enum class AtomicOperation {
+  kAnd = 0,
+  kAndSystem,
+  kOr,
+  kOrSystem,
+  kXor,
+  kXorSystem,
+};
+
+constexpr auto kMask = 0xAAAA;
+constexpr auto kTestValue = 0x4545;
+constexpr auto kAndTestValue = 0xFFFF;
+
+template <typename TestType, AtomicOperation operation>
+__host__ __device__ TestType GetTestValue() {
+  if constexpr (operation == AtomicOperation::kAnd || operation == AtomicOperation::kAndSystem) {
+    return kAndTestValue;
+  }
+
+  return kTestValue;
+}
+
+template <typename TestType, AtomicOperation operation>
+__device__ TestType PerformAtomicOperation(TestType* const mem) {
+  const auto mask = kMask;
+
+  if constexpr (operation == AtomicOperation::kAnd) {
+    return atomicAnd(mem, mask);
+  } else if constexpr (operation == AtomicOperation::kAndSystem) {
+    return atomicAnd_system(mem, mask);
+  } else if constexpr (operation == AtomicOperation::kOr) {
+    return atomicOr(mem, mask);
+  } else if constexpr (operation == AtomicOperation::kOrSystem) {
+    return atomicOr_system(mem, mask);
+  } else if constexpr (operation == AtomicOperation::kXor) {
+    return atomicXor(mem, mask);
+  } else if constexpr (operation == AtomicOperation::kXorSystem) {
+    return atomicXor_system(mem, mask);
+  }
+}
+
+template <typename TestType, AtomicOperation operation, bool use_shared_mem>
+__global__ void TestKernel(TestType* const global_mem, TestType* const old_vals) {
+  __shared__ TestType shared_mem;
+
+  const auto tid = cg::this_grid().thread_rank();
+
+  TestType* const mem = use_shared_mem ? &shared_mem : global_mem;
+
+  if constexpr (use_shared_mem) {
+    if (tid == 0) mem[0] = global_mem[0];
+    __syncthreads();
+  }
+
+  old_vals[tid] = PerformAtomicOperation<TestType, operation>(mem);
+
+  if constexpr (use_shared_mem) {
+    __syncthreads();
+    if (tid == 0) global_mem[0] = mem[0];
+  }
+}
+
+template <typename TestType>
+__host__ __device__ TestType* PitchedOffset(TestType* const ptr, const unsigned int pitch,
+                                            const unsigned int idx) {
+  const auto byte_ptr = reinterpret_cast<uint8_t*>(ptr);
+  return reinterpret_cast<TestType*>(byte_ptr + idx * pitch);
+}
+
+template <typename TestType, AtomicOperation operation, bool use_shared_mem>
+__global__ void TestKernel(TestType* const global_mem, TestType* const old_vals,
+                           const unsigned int width, const unsigned pitch) {
+  extern __shared__ uint8_t shared_mem[];
+
+  const auto tid = cg::this_grid().thread_rank();
+
+  TestType* const mem = use_shared_mem ? reinterpret_cast<TestType*>(shared_mem) : global_mem;
+
+  if constexpr (use_shared_mem) {
+    if (tid < width) {
+      const auto target = PitchedOffset(mem, pitch, tid);
+      *target = *PitchedOffset(global_mem, pitch, tid);
+    };
+    __syncthreads();
+  }
+
+  old_vals[tid] =
+      PerformAtomicOperation<TestType, operation>(PitchedOffset(mem, pitch, tid % width));
+
+  if constexpr (use_shared_mem) {
+    __syncthreads();
+    if (tid < width) {
+      const auto target = PitchedOffset(global_mem, pitch, tid);
+      *target = *PitchedOffset(mem, pitch, tid);
+    };
+  }
+}
+
+struct TestParams {
+  auto ThreadCount() const {
+    return blocks.x * blocks.y * blocks.z * threads.x * threads.y * threads.z;
+  }
+
+  dim3 blocks;
+  dim3 threads;
+  unsigned int num_devices = 1u;
+  unsigned int kernel_count = 1u;
+  unsigned int width = 1u;
+  unsigned int pitch = 0u;
+  unsigned int host_thread_count = 0u;
+  LinearAllocs alloc_type;
+};
+
+template <typename TestType, AtomicOperation operation>
+std::tuple<std::vector<TestType>, std::vector<TestType>> TestKernelHostRef(const TestParams& p) {
+  const auto thread_count = p.num_devices * p.kernel_count * p.ThreadCount();
+
+  TestType test_value = GetTestValue<TestType, operation>();
+  const auto mask = kMask;
+  std::vector<TestType> res_vals(p.width, test_value);
+  std::vector<TestType> old_vals;
+  old_vals.reserve(thread_count);
+
+  for (auto tid = 0u; tid < thread_count; ++tid) {
+    auto& res = res_vals[tid % p.width];
+    old_vals.push_back(res);
+
+    if constexpr (operation == AtomicOperation::kAnd || operation == AtomicOperation::kAndSystem) {
+      res = res & mask;
+    } else if constexpr (operation == AtomicOperation::kOr ||
+                         operation == AtomicOperation::kOrSystem) {
+      res = res | mask;
+    } else if constexpr (operation == AtomicOperation::kXor ||
+                         operation == AtomicOperation::kXorSystem) {
+      res = res ^ mask;
+    }
+  }
+
+  return {res_vals, old_vals};
+}
+
+template <typename TestType, AtomicOperation operation>
+void Verify(const TestParams& p, std::vector<TestType>& res_vals, std::vector<TestType>& old_vals) {
+  auto [expected_res_vals, expected_old_vals] = TestKernelHostRef<TestType, operation>(p);
+
+  for (auto i = 0u; i < res_vals.size(); ++i) {
+    INFO("Results index: " << i);
+    REQUIRE(expected_res_vals[i] == res_vals[i]);
+  }
+
+  std::sort(begin(old_vals), end(old_vals));
+  std::sort(begin(expected_old_vals), end(expected_old_vals));
+  for (auto i = 0u; i < old_vals.size(); ++i) {
+    INFO("Old values index: " << i);
+    REQUIRE(expected_old_vals[i] == old_vals[i]);
+  }
+}
+
+template <typename TestType, AtomicOperation operation, bool use_shared_mem>
+void LaunchKernel(const TestParams& p, hipStream_t stream, TestType* const mem_ptr,
+                  TestType* const old_vals) {
+  const auto shared_mem_size = use_shared_mem ? p.width * p.pitch : 0u;
+  if (p.width == 1 && p.pitch == sizeof(TestType))
+    TestKernel<TestType, operation, use_shared_mem>
+        <<<p.blocks, p.threads, shared_mem_size, stream>>>(mem_ptr, old_vals);
+  else
+    TestKernel<TestType, operation, use_shared_mem>
+        <<<p.blocks, p.threads, shared_mem_size, stream>>>(mem_ptr, old_vals, p.width, p.pitch);
+}
+
+template <typename TestType, AtomicOperation operation, bool use_shared_mem>
+void TestCore(const TestParams& p) {
+  const auto old_vals_alloc_size = p.kernel_count * p.ThreadCount() * sizeof(TestType);
+  std::vector<LinearAllocGuard<TestType>> old_vals_devs;
+  std::vector<StreamGuard> streams;
+  for (auto i = 0; i < p.num_devices; ++i) {
+    HIP_CHECK(hipSetDevice(i));
+    old_vals_devs.emplace_back(LinearAllocs::hipMalloc, old_vals_alloc_size);
+    for (auto j = 0; j < p.kernel_count; ++j) {
+      streams.emplace_back(Streams::created);
+    }
+  }
+
+  const auto mem_alloc_size = p.width * p.pitch;
+  LinearAllocGuard<TestType> mem_dev(p.alloc_type, mem_alloc_size);
+
+  std::vector<TestType> old_vals(p.num_devices * p.kernel_count * p.ThreadCount());
+  std::vector<TestType> res_vals(p.width);
+
+  TestType* const mem_ptr =
+      p.alloc_type == LinearAllocs::hipMalloc ? mem_dev.ptr() : mem_dev.host_ptr();
+
+  TestType test_value = GetTestValue<TestType, operation>();
+  HIP_CHECK(hipMemset(mem_ptr, 0, mem_alloc_size));
+  for (int i = 0; i < p.width * p.pitch / sizeof(TestType); ++i) {
+    HIP_CHECK(hipMemcpy(&mem_ptr[i], &test_value, sizeof(TestType), hipMemcpyHostToDevice));
+  }
+
+  for (auto i = 0u; i < p.num_devices; ++i) {
+    for (auto j = 0u; j < p.kernel_count; ++j) {
+      const auto& stream = streams[i * p.kernel_count + j].stream();
+      const auto old_vals = old_vals_devs[i].ptr() + j * p.ThreadCount();
+      LaunchKernel<TestType, operation, use_shared_mem>(p, stream, mem_dev.ptr(), old_vals);
+    }
+  }
+
+  for (auto i = 0u; i < p.num_devices; ++i) {
+    const auto device_offset = i * p.kernel_count * p.ThreadCount();
+    HIP_CHECK(hipMemcpy(old_vals.data() + device_offset, old_vals_devs[i].ptr(),
+                        old_vals_alloc_size, hipMemcpyDeviceToHost));
+  }
+  HIP_CHECK(hipMemcpy2D(res_vals.data(), sizeof(TestType), mem_ptr, p.pitch, sizeof(TestType),
+                        p.width, hipMemcpyDeviceToHost));
+
+  Verify<TestType, operation>(p, res_vals, old_vals);
+}
+
+template <typename TestType, AtomicOperation operation>
+void SingleDeviceSingleKernelTest(const unsigned int width, const unsigned int pitch) {
+  TestParams params;
+  params.num_devices = 1;
+  params.kernel_count = 1;
+  params.threads = GENERATE(dim3(1023));
+  params.width = width;
+  params.pitch = pitch;
+
+  SECTION("Global memory") {
+    params.blocks = GENERATE(dim3(3));
+    using LA = LinearAllocs;
+    for (const auto alloc_type :
+         {LA::hipMalloc, LA::hipHostMalloc, LA::hipMallocManaged, LA::mallocAndRegister}) {
+      params.alloc_type = alloc_type;
+      DYNAMIC_SECTION("Allocation type: " << to_string(alloc_type)) {
+        TestCore<TestType, operation, false>(params);
+      }
+    }
+  }
+
+  SECTION("Shared memory") {
+    params.blocks = dim3(1);
+    params.alloc_type = LinearAllocs::hipMalloc;
+    TestCore<TestType, operation, true>(params);
+  }
+}
+
+template <typename TestType, AtomicOperation operation>
+void SingleDeviceMultipleKernelTest(const unsigned int kernel_count, const unsigned int width,
+                                    const unsigned int pitch) {
+  int concurrent_kernels = 0;
+  HIP_CHECK(hipDeviceGetAttribute(&concurrent_kernels, hipDeviceAttributeConcurrentKernels, 0));
+  if (!concurrent_kernels) {
+    HipTest::HIP_SKIP_TEST("Test requires support for concurrent kernel execution");
+    return;
+  }
+
+  TestParams params;
+  params.num_devices = 1;
+  params.kernel_count = kernel_count;
+  params.blocks = GENERATE(dim3(3));
+  params.threads = GENERATE(dim3(1023));
+  params.width = width;
+  params.pitch = pitch;
+
+  using LA = LinearAllocs;
+  for (const auto alloc_type :
+       {LA::hipMalloc, LA::hipHostMalloc, LA::hipMallocManaged, LA::mallocAndRegister}) {
+    params.alloc_type = alloc_type;
+    DYNAMIC_SECTION("Allocation type: " << to_string(alloc_type)) {
+      TestCore<TestType, operation, false>(params);
+    }
+  }
+}
+
+template <typename TestType, AtomicOperation operation>
+void MultipleDeviceMultipleKernelTest(const unsigned int num_devices,
+                                      const unsigned int kernel_count, const unsigned int width,
+                                      const unsigned int pitch) {
+  if (num_devices > 1) {
+    if (HipTest::getDeviceCount() < num_devices) {
+      std::string msg = std::to_string(num_devices) + " devices are required";
+      HipTest::HIP_SKIP_TEST(msg.c_str());
+      return;
+    }
+  }
+
+  if (kernel_count > 1) {
+    for (auto i = 0u; i < num_devices; ++i) {
+      int concurrent_kernels = 0;
+      HIP_CHECK(hipDeviceGetAttribute(&concurrent_kernels, hipDeviceAttributeConcurrentKernels, i));
+      if (!concurrent_kernels) {
+        HipTest::HIP_SKIP_TEST("Test requires support for concurrent kernel execution");
+        return;
+      }
+    }
+  }
+
+  TestParams params;
+  params.num_devices = num_devices;
+  params.kernel_count = kernel_count;
+  params.blocks = GENERATE(dim3(3));
+  params.threads = GENERATE(dim3(1023));
+  params.width = width;
+  params.pitch = pitch;
+
+  using LA = LinearAllocs;
+  for (const auto alloc_type : {LA::hipHostMalloc, LA::hipMallocManaged, LA::mallocAndRegister}) {
+    params.alloc_type = alloc_type;
+    DYNAMIC_SECTION("Allocation type: " << to_string(alloc_type)) {
+      TestCore<TestType, operation, false>(params);
+    }
+  }
+}
+}  // namespace Bitwise