diff --git a/catch/unit/kernels/CMakeLists.txt b/catch/ABM/AddKernels/CMakeLists.txt
similarity index 62%
rename from catch/unit/kernels/CMakeLists.txt
rename to catch/ABM/AddKernels/CMakeLists.txt
index 27a95cf29d..543f01d7a7 100644
--- a/catch/unit/kernels/CMakeLists.txt
+++ b/catch/ABM/AddKernels/CMakeLists.txt
@@ -4,7 +4,7 @@ set(TEST_SRC
 )
 
 # Create shared lib of all tests
-add_library(Kernels SHARED EXCLUDE_FROM_ALL ${TEST_SRC})
+add_library(ABMAddKernels SHARED EXCLUDE_FROM_ALL ${TEST_SRC})
 
 # Add dependency on build_tests to build it on this custom target
-add_dependencies(build_tests Kernels)
+add_dependencies(build_tests ABMAddKernels)
diff --git a/catch/unit/kernels/add.cc b/catch/ABM/AddKernels/add.cc
similarity index 92%
rename from catch/unit/kernels/add.cc
rename to catch/ABM/AddKernels/add.cc
index 4f70ffef77..f2be747a77 100644
--- a/catch/unit/kernels/add.cc
+++ b/catch/ABM/AddKernels/add.cc
@@ -6,7 +6,7 @@ template <typename T> __global__ void add(T* a, T* b, T* c, size_t size) {
   if (i < size) c[i] = a[i] + b[i];
 }
 
-TEMPLATE_TEST_CASE("Add Kernel", "[kernel][add]", int, long, float, long long, double) {
+TEMPLATE_TEST_CASE("ABM_AddKernel_MultiTypeMultiSize", "", int, long, float, long long, double) {
   auto size = GENERATE(as<size_t>{}, 100, 500, 1000);
   TestType *d_a, *d_b, *d_c;
   auto res = hipMalloc(&d_a, sizeof(TestType) * size);
diff --git a/catch/ABM/CMakeLists.txt b/catch/ABM/CMakeLists.txt
new file mode 100644
index 0000000000..53fa9e8b0f
--- /dev/null
+++ b/catch/ABM/CMakeLists.txt
@@ -0,0 +1 @@
+add_subdirectory(AddKernels)
diff --git a/catch/CMakeLists.txt b/catch/CMakeLists.txt
index d3b3f028e0..0292f378dc 100644
--- a/catch/CMakeLists.txt
+++ b/catch/CMakeLists.txt
@@ -66,6 +66,7 @@ add_custom_target(build_tests)
 
 # Tests folder
 add_subdirectory(unit)
+add_subdirectory(ABM)
 add_subdirectory(hipTestMain)
 add_subdirectory(stress)
 
diff --git a/catch/README.md b/catch/README.md
index 9d57c8a59d..2f66a1f18f 100644
--- a/catch/README.md
+++ b/catch/README.md
@@ -71,3 +71,18 @@ Catch2 allows multiple ways in which you can debug the test case.
 ## External Libs being used
 - [Catch2](https://github.com/catchorg/Catch2) - Testing framework
 - [picojson](https://github.com/kazuho/picojson) - For config file parsing
+
+# Testing Guidelines
+Tests fall in 5 categories and its file name prefix are as follows:
+ - Unit tests (Prefix: Unit_\*API\*_\*Optional Scenario\*, example : Unit_hipMalloc_Negative or Unit_hipMalloc): Unit Tests are simplest test for an API, the target here is to test the API with different types of input and different ways of calling.
+ - Application Behavior Modelling tests (Prefix: ABM_\*Intent\*_\*Optional Scenario\*, example: ABM_ModuleLoadAndRun): ABM tests are used to model a specific use case of HIP APIs, either seen in a customer app or a general purpose app. It mimics the calling behavior seen in aforementioned app.
+ - Stress/Scale tests (Prefix: Stress_\*API\*_\*Intent\*_\*Optional Scenario\*, example: Stress_hipMemset_ExhaustVRAM): These tests are used to see the behavior of HIP APIs in edge scenarios, for example what happens when we have exhausted vram and do a hipMalloc or run many instances of same API in parallel.
+ - Multi Process tests (Prefix: MultiProc_\*API\*_\*Optional Scenario\*, example: MultiProc_hipIPCMemHandle_GetDataFromProc): These tests are multi process tests and will only run on linux. They are used to test HIP APIs in multi process environment
+ - Performance tests(Prefix: Perf_\*Intent\*_\*Optional Scenario\*, example: Perf_DispatchLatenc  y): Performance tests are used to get results of HIP APIs.
+
+General Guidelines:
+ - Do not use the catch2 tags. Tags wont be used for filtering
+ - Add as many INFO() as you can in tests which prints state of the t est, this will help the debugger when the test fails (INFO macro only prints when the test fails)
+ - Check return of each HIP API and fail whenever there is a misma    tch with hipSuccess or hiprtcSuccess.
+ - Each Category of test will hav e its own exe and catch_discover_test macro will be called on it to discover its tests
+ - Optional Scenario in test names are optional. For example you  can test all Scenarios of hipMalloc API in one file, you can name the file Unit_hipMalloc, if you are having a file just for negative scenarios you can name it as Unit_hipMalloc_Negative.
diff --git a/catch/hipTestMain/CMakeLists.txt b/catch/hipTestMain/CMakeLists.txt
index 3ca07fb7c7..08daceb9e7 100644
--- a/catch/hipTestMain/CMakeLists.txt
+++ b/catch/hipTestMain/CMakeLists.txt
@@ -9,9 +9,8 @@ else()
     target_compile_options(UnitTests PUBLIC -std=c++17)
 endif()
 
-target_link_libraries(UnitTests PRIVATE DeviceLibs
+target_link_libraries(UnitTests PRIVATE UnitDeviceTests
                                         MemoryTest
-                                        Kernels
                                         stdc++fs)
 
 # Add AMD Only Tests
@@ -20,7 +19,19 @@ if(HIP_PLATFORM MATCHES "amd")
 endif()
 
 catch_discover_tests(UnitTests PROPERTIES  SKIP_REGULAR_EXPRESSION "HIP_SKIP_THIS_TEST")
-add_dependencies(build_tests UnitTests)
+
+# ABM exe
+add_executable(ABMTests EXCLUDE_FROM_ALL main.cc hip_test_context.cc)
+if(HIP_PLATFORM MATCHES "amd")
+    set_property(TARGET ABMTests PROPERTY CXX_STANDARD 17)
+else()
+    target_compile_options(ABMTests PUBLIC -std=c++17)
+endif()
+
+target_link_libraries(ABMTests PRIVATE ABMAddKernels)
+catch_discover_tests(ABMTests PROPERTIES  SKIP_REGULAR_EXPRESSION "HIP_SKIP_THIS_TEST")
+
+add_dependencies(build_tests UnitTests ABMTests)
 
 # Add Multiproc tests as seperate binary
 if(UNIX AND HIP_PLATFORM MATCHES "amd")
diff --git a/catch/include/hip_test_checkers.hh b/catch/include/hip_test_checkers.hh
index d0c180c25c..4b2e2f49ba 100644
--- a/catch/include/hip_test_checkers.hh
+++ b/catch/include/hip_test_checkers.hh
@@ -75,13 +75,13 @@ bool initArraysForHost(T** A_h, T** B_h, T** C_h, size_t N, bool usePinnedHost =
 
   if (usePinnedHost) {
     if (A_h) {
-      HIPCHECK(hipHostMalloc((void**)A_h, Nbytes));
+      HIP_CHECK(hipHostMalloc((void**)A_h, Nbytes));
     }
     if (B_h) {
-      HIPCHECK(hipHostMalloc((void**)B_h, Nbytes));
+      HIP_CHECK(hipHostMalloc((void**)B_h, Nbytes));
     }
     if (C_h) {
-      HIPCHECK(hipHostMalloc((void**)C_h, Nbytes));
+      HIP_CHECK(hipHostMalloc((void**)C_h, Nbytes));
     }
   } else {
     if (A_h) {
@@ -110,13 +110,13 @@ bool initArrays(T** A_d, T** B_d, T** C_d, T** A_h, T** B_h, T** C_h, size_t N,
   size_t Nbytes = N * sizeof(T);
 
   if (A_d) {
-    HIPCHECK(hipMalloc(A_d, Nbytes));
+    HIP_CHECK(hipMalloc(A_d, Nbytes));
   }
   if (B_d) {
-    HIPCHECK(hipMalloc(B_d, Nbytes));
+    HIP_CHECK(hipMalloc(B_d, Nbytes));
   }
   if (C_d) {
-    HIPCHECK(hipMalloc(C_d, Nbytes));
+    HIP_CHECK(hipMalloc(C_d, Nbytes));
   }
 
   return initArraysForHost(A_h, B_h, C_h, N, usePinnedHost);
@@ -125,13 +125,13 @@ bool initArrays(T** A_d, T** B_d, T** C_d, T** A_h, T** B_h, T** C_h, size_t N,
 template <typename T> bool freeArraysForHost(T* A_h, T* B_h, T* C_h, bool usePinnedHost) {
   if (usePinnedHost) {
     if (A_h) {
-      HIPCHECK(hipHostFree(A_h));
+      HIP_CHECK(hipHostFree(A_h));
     }
     if (B_h) {
-      HIPCHECK(hipHostFree(B_h));
+      HIP_CHECK(hipHostFree(B_h));
     }
     if (C_h) {
-      HIPCHECK(hipHostFree(C_h));
+      HIP_CHECK(hipHostFree(C_h));
     }
   } else {
     if (A_h) {
@@ -150,13 +150,13 @@ template <typename T> bool freeArraysForHost(T* A_h, T* B_h, T* C_h, bool usePin
 template <typename T>
 bool freeArrays(T* A_d, T* B_d, T* C_d, T* A_h, T* B_h, T* C_h, bool usePinnedHost) {
   if (A_d) {
-    HIPCHECK(hipFree(A_d));
+    HIP_CHECK(hipFree(A_d));
   }
   if (B_d) {
-    HIPCHECK(hipFree(B_d));
+    HIP_CHECK(hipFree(B_d));
   }
   if (C_d) {
-    HIPCHECK(hipFree(C_d));
+    HIP_CHECK(hipFree(C_d));
   }
 
   return freeArraysForHost(A_h, B_h, C_h, usePinnedHost);
diff --git a/catch/include/hip_test_common.hh b/catch/include/hip_test_common.hh
index a6e07973f9..dc2d6980d5 100644
--- a/catch/include/hip_test_common.hh
+++ b/catch/include/hip_test_common.hh
@@ -1,10 +1,32 @@
+/*
+Copyright (c) 2021-present 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_context.hh"
 #include <catch.hpp>
 
 #define HIP_PRINT_STATUS(status) INFO(hipGetErrorName(status) << " at line: " << __LINE__);
 
-#define HIPCHECK(error)                                                                            \
+#define HIP_CHECK(error)                                                                           \
   {                                                                                                \
     hipError_t localError = error;                                                                 \
     if ((localError != hipSuccess) && (localError != hipErrorPeerAccessAlreadyEnabled)) {          \
@@ -14,3 +36,12 @@
     }                                                                                              \
   }
 
+// Although its assert, it will be evaluated at runtime
+#define HIP_ASSERT(x)                                                                              \
+  { REQUIRE((x)); }
+
+
+// Utility Functions
+namespace HipTest {
+int getDeviceCount();
+}
diff --git a/catch/include/hip_test_helper.hh b/catch/include/hip_test_helper.hh
new file mode 100644
index 0000000000..db51ceecdb
--- /dev/null
+++ b/catch/include/hip_test_helper.hh
@@ -0,0 +1,32 @@
+/*
+Copyright (c) 2021-present 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"
+
+namespace HipTest {
+static inline int getGeviceCount() {
+  int dev = 0;
+  HIPCHECK(hipGetDeviceCount(&dev));
+  return dev;
+}
+}  // namespace HipTest
diff --git a/catch/include/hip_test_kernels.hh b/catch/include/hip_test_kernels.hh
index 7196accd97..999cced8c3 100644
--- a/catch/include/hip_test_kernels.hh
+++ b/catch/include/hip_test_kernels.hh
@@ -1,3 +1,25 @@
+/*
+Copyright (c) 2021-present 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/hip_runtime.h>
@@ -59,4 +81,4 @@ template <typename T> __global__ void memsetReverse(T* C_d, T val, int64_t NELEM
     C_d[i] = val;
   }
 }
-}  // namespace HipTest
\ No newline at end of file
+}  // namespace HipTest
diff --git a/catch/multiproc/hipMallocConcurrency.cc b/catch/multiproc/hipMallocConcurrency.cc
index 72d17c26a0..23d28fd40e 100644
--- a/catch/multiproc/hipMallocConcurrency.cc
+++ b/catch/multiproc/hipMallocConcurrency.cc
@@ -20,9 +20,9 @@ unsigned threadsPerBlock = 256;
 
 unsigned setNumBlocks(unsigned blocksPerCU, unsigned threadsPerBlock, size_t N) {
   int device;
-  HIPCHECK(hipGetDevice(&device));
+  HIP_CHECK(hipGetDevice(&device));
   hipDeviceProp_t props;
-  HIPCHECK(hipGetDeviceProperties(&props, device));
+  HIP_CHECK(hipGetDeviceProperties(&props, device));
 
   unsigned blocks = props.multiProcessorCount * blocksPerCU;
   if (blocks * threadsPerBlock > N) {
@@ -43,20 +43,20 @@ bool validateMemoryOnGPU(int gpu, bool concurOnOneGPU = false) {
   size_t prevAvl, prevTot, curAvl, curTot;
   bool TestPassed = true;
 
-  HIPCHECK(hipSetDevice(gpu));
-  HIPCHECK(hipMemGetInfo(&prevAvl, &prevTot));
+  HIP_CHECK(hipSetDevice(gpu));
+  HIP_CHECK(hipMemGetInfo(&prevAvl, &prevTot));
   printf("tgs allocating..\n");
   HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false);
 
   unsigned blocks = setNumBlocks(blocksPerCU, threadsPerBlock, N);
 
-  HIPCHECK(hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice));
-  HIPCHECK(hipMemcpy(B_d, B_h, Nbytes, hipMemcpyHostToDevice));
+  HIP_CHECK(hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice));
+  HIP_CHECK(hipMemcpy(B_d, B_h, Nbytes, hipMemcpyHostToDevice));
 
   hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), 0, 0,
                      static_cast<const int*>(A_d), static_cast<const int*>(B_d), C_d, N);
 
-  HIPCHECK(hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost));
+  HIP_CHECK(hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost));
 
   if (!HipTest::checkVectorADD(A_h, B_h, C_h, N)) {
     printf("Validation PASSED for gpu %d from pid %d\n", gpu, getpid());
@@ -66,7 +66,7 @@ bool validateMemoryOnGPU(int gpu, bool concurOnOneGPU = false) {
   }
 
   HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false);
-  HIPCHECK(hipMemGetInfo(&curAvl, &curTot));
+  HIP_CHECK(hipMemGetInfo(&curAvl, &curTot));
 
   if (!concurOnOneGPU && (prevAvl != curAvl || prevTot != curTot)) {
     // In concurrent calls on one GPU, we cannot verify leaking in this way
@@ -116,7 +116,7 @@ void getDeviceCount1(int* pdevCnt) {
     // writing only, no need for read-descriptor
     close(fd[0]);
 
-    HIPCHECK(hipGetDeviceCount(&devCnt));
+    HIP_CHECK(hipGetDeviceCount(&devCnt));
     // send the value on the write-descriptor:
     write(fd[1], &devCnt, sizeof(devCnt));
 
@@ -129,7 +129,7 @@ void getDeviceCount1(int* pdevCnt) {
   }
 
 #else
-  HIPCHECK(hipGetDeviceCount(pdevCnt));
+  HIP_CHECK(hipGetDeviceCount(pdevCnt));
 #endif
 }
 #endif
diff --git a/catch/unit/CMakeLists.txt b/catch/unit/CMakeLists.txt
index cff55f37a8..323feaa277 100644
--- a/catch/unit/CMakeLists.txt
+++ b/catch/unit/CMakeLists.txt
@@ -1,5 +1,4 @@
 add_subdirectory(memory)
 add_subdirectory(deviceLib)
-add_subdirectory(kernels)
 # Disable Saxpy test temporarily to see if CI Passes
 # add_subdirectory(rtc)
diff --git a/catch/unit/deviceLib/CMakeLists.txt b/catch/unit/deviceLib/CMakeLists.txt
index 22de79c687..357a68e4a3 100644
--- a/catch/unit/deviceLib/CMakeLists.txt
+++ b/catch/unit/deviceLib/CMakeLists.txt
@@ -1,20 +1,34 @@
 # Common Tests - Test independent of all platforms
 set(TEST_SRC
     floatMath.cc
+    anyAll.cc
+    ballot.cc
+    clz.cc
+    ffs.cc
+    funnelshift.cc
+    brev.cc
 )
 
 # AMD only tests
 set(AMD_TEST_SRC
     vectorTypesDevice.cc
+    bitExtract.cc
+    bitInsert.cc
+    floatTM.cc
 )
 
 if(HIP_PLATFORM MATCHES "amd")
     set(TEST_SRC ${TEST_SRC} ${AMD_TEST_SRC})
+    set_source_files_properties(floatTM.cc PROPERTIES COMPILE_FLAGS -std=c++17)
 endif()
 
 
 # Create shared lib of all tests
-add_library(DeviceLibs SHARED EXCLUDE_FROM_ALL ${TEST_SRC})
+add_library(UnitDeviceTests SHARED EXCLUDE_FROM_ALL ${TEST_SRC})
+
+if(HIP_PLATFORM MATCHES "nvidia")
+    target_compile_options(UnitDeviceTests PUBLIC --Wno-deprecated-declarations)
+endif()
 
 # Add dependency on build_tests to build it on this custom target
-add_dependencies(build_tests DeviceLibs)
+add_dependencies(build_tests UnitDeviceTests)
diff --git a/catch/unit/deviceLib/anyAll.cc b/catch/unit/deviceLib/anyAll.cc
new file mode 100644
index 0000000000..d7d5a0d973
--- /dev/null
+++ b/catch/unit/deviceLib/anyAll.cc
@@ -0,0 +1,83 @@
+/*
+Copyright (c) 2021 - present 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>
+
+
+__global__ void warpvote(int* device_any, int* device_all, int pshift) {
+  int tid = threadIdx.x + blockIdx.x * blockDim.x;
+  device_any[threadIdx.x >> pshift] = __any(tid - 77);
+  device_all[threadIdx.x >> pshift] = __all(tid - 77);
+}
+
+
+TEST_CASE("Unit_AnyAll_CompileTest") {
+  int warpSize, pshift;
+  hipDeviceProp_t devProp;
+  hipGetDeviceProperties(&devProp, 0);
+  warpSize = devProp.warpSize;
+
+  int w = warpSize;
+  pshift = 0;
+  while (w >>= 1) ++pshift;
+
+  INFO("WarpSize:: " << warpSize << " pShift: " << pshift);
+
+
+  int anycount = 0;
+  int allcount = 0;
+  int Num_Threads_per_Block = 1024;
+  int Num_Blocks_per_Grid = 1;
+  int Num_Warps_per_Grid = (Num_Threads_per_Block * Num_Blocks_per_Grid) / warpSize;
+
+  int* host_any = (int*)malloc(Num_Warps_per_Grid * sizeof(int));
+  int* host_all = (int*)malloc(Num_Warps_per_Grid * sizeof(int));
+  int* device_any;
+  int* device_all;
+  HIP_CHECK(hipMalloc((void**)&device_any, Num_Warps_per_Grid * sizeof(int)));
+  HIP_CHECK(hipMalloc((void**)&device_all, Num_Warps_per_Grid * sizeof(int)));
+  for (int i = 0; i < Num_Warps_per_Grid; i++) {
+    host_any[i] = 0;
+    host_all[i] = 0;
+  }
+  HIP_CHECK(hipMemcpy(device_any, host_any, sizeof(int), hipMemcpyHostToDevice));
+  HIP_CHECK(hipMemcpy(device_all, host_all, sizeof(int), hipMemcpyHostToDevice));
+
+  hipLaunchKernelGGL(warpvote, dim3(Num_Blocks_per_Grid), dim3(Num_Threads_per_Block), 0, 0,
+                     device_any, device_all, pshift);
+
+  HIP_CHECK(
+      hipMemcpy(host_any, device_any, Num_Warps_per_Grid * sizeof(int), hipMemcpyDeviceToHost));
+  HIP_CHECK(
+      hipMemcpy(host_all, device_all, Num_Warps_per_Grid * sizeof(int), hipMemcpyDeviceToHost));
+  for (int i = 0; i < Num_Warps_per_Grid; i++) {
+    INFO("Warp Number: " << i << " __any: " << host_any[i] << " __all: " << host_all[i]);
+
+    if (host_all[i] != 1) ++allcount;
+    if (host_any[i] != 1) ++anycount;
+  }
+
+  HIP_CHECK(hipFree(device_any));
+  HIP_CHECK(hipFree(device_all));
+  REQUIRE(anycount == 0);
+  REQUIRE(allcount == 1);
+}
diff --git a/catch/unit/deviceLib/ballot.cc b/catch/unit/deviceLib/ballot.cc
new file mode 100644
index 0000000000..0668d51435
--- /dev/null
+++ b/catch/unit/deviceLib/ballot.cc
@@ -0,0 +1,85 @@
+/*
+Copyright (c) 2021 - present 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>
+#include <hip/device_functions.h>
+
+
+__global__ void gpu_ballot(unsigned int* device_ballot, unsigned Num_Warps_per_Block,
+                           unsigned pshift) {
+  int tid = threadIdx.x + blockIdx.x * blockDim.x;
+  const unsigned int warp_num = threadIdx.x >> pshift;
+
+  // TODO try to remove this
+#ifdef __HIP_PLATFORM_AMD__
+  atomicAdd(&device_ballot[warp_num + blockIdx.x * Num_Warps_per_Block],
+            __popcll(__ballot(tid - 245)));
+#else
+  atomicAdd(&device_ballot[warp_num + blockIdx.x * Num_Warps_per_Block],
+            __popc(__ballot(tid - 245)));
+#endif
+}
+
+TEST_CASE("Unit_ballot") {
+  unsigned warpSize, pshift;
+  hipDeviceProp_t devProp;
+  HIP_CHECK(hipGetDeviceProperties(&devProp, 0));
+
+  warpSize = devProp.warpSize;
+
+  int w = warpSize;
+  pshift = 0;
+  while (w >>= 1) ++pshift;
+
+  unsigned int Num_Threads_per_Block = 512;
+  unsigned int Num_Blocks_per_Grid = 1;
+  unsigned int Num_Warps_per_Block = Num_Threads_per_Block / warpSize;
+  unsigned int Num_Warps_per_Grid = (Num_Threads_per_Block * Num_Blocks_per_Grid) / warpSize;
+  unsigned int* host_ballot = (unsigned int*)malloc(Num_Warps_per_Grid * sizeof(unsigned int));
+  unsigned int* device_ballot;
+  HIP_CHECK(hipMalloc((void**)&device_ballot, Num_Warps_per_Grid * sizeof(unsigned int)));
+  int divergent_count = 0;
+  for (unsigned i = 0; i < Num_Warps_per_Grid; i++) host_ballot[i] = 0;
+
+  HIP_CHECK(hipMemcpy(device_ballot, host_ballot, Num_Warps_per_Grid * sizeof(unsigned int),
+                       hipMemcpyHostToDevice));
+
+  hipLaunchKernelGGL(gpu_ballot, dim3(Num_Blocks_per_Grid), dim3(Num_Threads_per_Block), 0, 0,
+                     device_ballot, Num_Warps_per_Block, pshift);
+
+  HIP_CHECK(hipMemcpy(host_ballot, device_ballot, Num_Warps_per_Grid * sizeof(unsigned int),
+                       hipMemcpyDeviceToHost));
+
+  for (unsigned i = 0; i < Num_Warps_per_Grid; i++) {
+    if ((host_ballot[i] == 0) || (host_ballot[i] / warpSize == warpSize)) {
+      INFO("Warp: " << i << " is convergent - predicate true for " << (host_ballot[i] / warpSize)
+                    << " threads");
+    } else {
+      INFO("Warp: " << i << " is divergent - predicate true for " << (host_ballot[i] / warpSize)
+                    << " threads");
+      divergent_count++;
+    }
+  }
+
+  HIP_CHECK(hipFree(device_ballot));
+  REQUIRE(divergent_count == 1);
+}
diff --git a/catch/unit/deviceLib/bitExtract.cc b/catch/unit/deviceLib/bitExtract.cc
new file mode 100644
index 0000000000..249c0949e0
--- /dev/null
+++ b/catch/unit/deviceLib/bitExtract.cc
@@ -0,0 +1,194 @@
+/*
+Copyright (c) 2021 - present 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>
+#include <hip/device_functions.h>
+
+#include <assert.h>
+#include <stdio.h>
+#include <algorithm>
+#include <stdlib.h>
+#include <iostream>
+#include <random>
+
+// CPU implementation of bitextract
+template <typename T> T bit_extract(T src0, unsigned int src1, unsigned int src2) {
+  unsigned int bits = sizeof(T) * 8;
+  T offset = src1 & (bits - 1);
+  T width = src2 & (bits - 1);
+  if (width == 0) {
+    return 0;
+  } else {
+    return (src0 << (bits - width - offset)) >> (bits - width);
+  }
+}
+
+__global__ void HIP_kernel(unsigned int* out32, unsigned int* in32_0, unsigned int* in32_1,
+                           unsigned int* in32_2, unsigned long long int* out64,
+                           unsigned long long int* in64_0, unsigned int* in64_1,
+                           unsigned int* in64_2) {
+  int x = blockDim.x * blockIdx.x + threadIdx.x;
+
+  out32[x] = __bitextract_u32(in32_0[x], in32_1[x], in32_2[x]);
+  out64[x] = __bitextract_u64(in64_0[x], in64_1[x], in64_2[x]);
+}
+
+TEST_CASE("Unit_bitExtract") {
+  using namespace std;
+
+  unsigned int* hostOut32;
+  unsigned int* hostSrc032;
+  unsigned int* hostSrc132;
+  unsigned int* hostSrc232;
+  unsigned long long int* hostOut64;
+  unsigned long long int* hostSrc064;
+  unsigned int* hostSrc164;
+  unsigned int* hostSrc264;
+
+  unsigned int* deviceOut32;
+  unsigned int* deviceSrc032;
+  unsigned int* deviceSrc132;
+  unsigned int* deviceSrc232;
+  unsigned long long int* deviceOut64;
+  unsigned long long int* deviceSrc064;
+  unsigned int* deviceSrc164;
+  unsigned int* deviceSrc264;
+
+  hipDeviceProp_t devProp;
+  HIP_CHECK(hipGetDeviceProperties(&devProp, 0));
+  INFO("System minor : " << devProp.minor);
+  INFO("System major : " << devProp.major);
+  INFO("agent prop name : " << devProp.name);
+
+  INFO("hip Device prop succeeded");
+
+  unsigned int wave_size = devProp.warpSize;
+  unsigned int num_waves_per_block = 2;
+  unsigned int num_threads_per_block = wave_size * num_waves_per_block;
+  unsigned int num_blocks = 2;
+  unsigned int NUM = num_threads_per_block * num_blocks;
+
+  unsigned i;
+  int errors;
+
+  hostOut32 = (unsigned int*)malloc(NUM * sizeof(unsigned int));
+  hostSrc032 = (unsigned int*)malloc(NUM * sizeof(unsigned int));
+  hostSrc132 = (unsigned int*)malloc(NUM * sizeof(unsigned int));
+  hostSrc232 = (unsigned int*)malloc(NUM * sizeof(unsigned int));
+
+  hostOut64 = (unsigned long long int*)malloc(NUM * sizeof(unsigned long long int));
+  hostSrc064 = (unsigned long long int*)malloc(NUM * sizeof(unsigned long long int));
+  hostSrc164 = (unsigned int*)malloc(NUM * sizeof(unsigned int));
+  hostSrc264 = (unsigned int*)malloc(NUM * sizeof(unsigned int));
+
+  // initialize the input data
+  std::random_device rd;
+  std::uniform_int_distribution<uint32_t> uint32_src0_dist;
+  std::uniform_int_distribution<uint32_t> uint32_src12_dist(0, 31);
+  std::uniform_int_distribution<uint64_t> uint64_src0_dist;
+  std::uniform_int_distribution<uint32_t> uint64_src12_dist(0, 63);
+  for (i = 0; i < NUM; i++) {
+    hostOut32[i] = 0;
+    hostSrc032[i] = uint32_src0_dist(rd);
+    hostSrc132[i] = uint32_src12_dist(rd);
+    hostSrc232[i] = uint32_src12_dist(rd);
+    if (hostSrc132[i] + hostSrc232[i] > 32) hostSrc232[i] = 32 - hostSrc132[i];
+    hostOut64[i] = 0;
+    hostSrc064[i] = uint64_src0_dist(rd);
+    hostSrc164[i] = uint64_src12_dist(rd);
+    hostSrc264[i] = uint64_src12_dist(rd);
+  }
+
+  HIP_CHECK(hipMalloc((void**)&deviceOut32, NUM * sizeof(unsigned int)));
+  HIP_CHECK(hipMalloc((void**)&deviceSrc032, NUM * sizeof(unsigned int)));
+  HIP_CHECK(hipMalloc((void**)&deviceSrc132, NUM * sizeof(unsigned int)));
+  HIP_CHECK(hipMalloc((void**)&deviceSrc232, NUM * sizeof(unsigned int)));
+
+  HIP_CHECK(hipMalloc((void**)&deviceOut64, NUM * sizeof(unsigned long long int)));
+  HIP_CHECK(hipMalloc((void**)&deviceSrc064, NUM * sizeof(unsigned long long int)));
+  HIP_CHECK(hipMalloc((void**)&deviceSrc164, NUM * sizeof(unsigned int)));
+  HIP_CHECK(hipMalloc((void**)&deviceSrc264, NUM * sizeof(unsigned int)));
+
+  HIP_CHECK(
+      hipMemcpy(deviceSrc032, hostSrc032, NUM * sizeof(unsigned int), hipMemcpyHostToDevice));
+  HIP_CHECK(
+      hipMemcpy(deviceSrc132, hostSrc132, NUM * sizeof(unsigned int), hipMemcpyHostToDevice));
+  HIP_CHECK(
+      hipMemcpy(deviceSrc232, hostSrc232, NUM * sizeof(unsigned int), hipMemcpyHostToDevice));
+
+  HIP_CHECK(hipMemcpy(deviceSrc064, hostSrc064, NUM * sizeof(unsigned long long int),
+                       hipMemcpyHostToDevice));
+  HIP_CHECK(
+      hipMemcpy(deviceSrc164, hostSrc164, NUM * sizeof(unsigned int), hipMemcpyHostToDevice));
+  HIP_CHECK(
+      hipMemcpy(deviceSrc264, hostSrc264, NUM * sizeof(unsigned int), hipMemcpyHostToDevice));
+
+
+  hipLaunchKernelGGL(HIP_kernel, dim3(num_blocks), dim3(num_threads_per_block), 0, 0, deviceOut32,
+                     deviceSrc032, deviceSrc132, deviceSrc232, deviceOut64, deviceSrc064,
+                     deviceSrc164, deviceSrc264);
+
+
+  HIP_CHECK(hipMemcpy(hostOut32, deviceOut32, NUM * sizeof(unsigned int), hipMemcpyDeviceToHost));
+  HIP_CHECK(hipMemcpy(hostOut64, deviceOut64, NUM * sizeof(unsigned long long int),
+                       hipMemcpyDeviceToHost));
+
+  // verify the results
+  errors = 0;
+  for (i = 0; i < NUM; i++) {
+    if (hostOut32[i] != bit_extract<uint32_t>(hostSrc032[i], hostSrc132[i], hostSrc232[i])) {
+      errors++;
+      INFO("device: " << hostOut32[i] << " host: "
+                      << bit_extract<uint32_t>(hostSrc032[i], hostSrc132[i], hostSrc232[i]) << " "
+                      << hostSrc032[i] << " " << hostSrc132[i] << " " << hostSrc232[i] << "\n");
+    }
+  }
+
+  for (i = 0; i < NUM; i++) {
+    if (hostOut64[i] != bit_extract<uint64_t>(hostSrc064[i], hostSrc164[i], hostSrc264[i])) {
+      errors++;
+      INFO("device: " << hostOut64[i] << " host: "
+                      << bit_extract<uint64_t>(hostSrc064[i], hostSrc164[i], hostSrc264[i]) << " "
+                      << hostSrc064[i] << " " << hostSrc164[i] << " " << hostSrc264[i] << "\n");
+    }
+  }
+
+  HIP_CHECK(hipFree(deviceOut32));
+  HIP_CHECK(hipFree(deviceSrc032));
+  HIP_CHECK(hipFree(deviceSrc132));
+  HIP_CHECK(hipFree(deviceSrc232));
+  HIP_CHECK(hipFree(deviceOut64));
+  HIP_CHECK(hipFree(deviceSrc064));
+  HIP_CHECK(hipFree(deviceSrc164));
+  HIP_CHECK(hipFree(deviceSrc264));
+
+  free(hostOut32);
+  free(hostSrc032);
+  free(hostSrc132);
+  free(hostSrc232);
+  free(hostOut64);
+  free(hostSrc064);
+  free(hostSrc164);
+  free(hostSrc264);
+
+  REQUIRE(errors == 0);
+}
diff --git a/catch/unit/deviceLib/bitInsert.cc b/catch/unit/deviceLib/bitInsert.cc
new file mode 100644
index 0000000000..e8e92f37e8
--- /dev/null
+++ b/catch/unit/deviceLib/bitInsert.cc
@@ -0,0 +1,216 @@
+/*
+Copyright (c) 2021 - present 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>
+#include <hip/device_functions.h>
+
+#include <assert.h>
+#include <stdio.h>
+#include <algorithm>
+#include <stdlib.h>
+#include <iostream>
+#include <random>
+
+
+// CPU implementation of bitinsert
+template <typename T> T bit_insert(T src0, T src1, unsigned int src2, unsigned int src3) {
+  unsigned int bits = sizeof(T) * 8;
+  T offset = src2 & (bits - 1);
+  T width = src3 & (bits - 1);
+  T mask = (((T)1) << width) - 1;
+  return ((src0 & ~(mask << offset)) | ((src1 & mask) << offset));
+}
+
+__global__ void HIP_kernel(unsigned int* out32, unsigned int* in32_0, unsigned int* in32_1,
+                           unsigned int* in32_2, unsigned int* in32_3,
+                           unsigned long long int* out64, unsigned long long int* in64_0,
+                           unsigned long long int* in64_1, unsigned int* in64_2,
+                           unsigned int* in64_3) {
+  int x = blockDim.x * blockIdx.x + threadIdx.x;
+
+  out32[x] = __bitinsert_u32(in32_0[x], in32_1[x], in32_2[x], in32_3[x]);
+  out64[x] = __bitinsert_u64(in64_0[x], in64_1[x], in64_2[x], in64_3[x]);
+}
+
+TEST_CASE("Unit_bitInsert") {
+  using namespace std;
+
+  unsigned int* hostOut32;
+  unsigned int* hostSrc032;
+  unsigned int* hostSrc132;
+  unsigned int* hostSrc232;
+  unsigned int* hostSrc332;
+  unsigned long long int* hostOut64;
+  unsigned long long int* hostSrc064;
+  unsigned long long int* hostSrc164;
+  unsigned int* hostSrc264;
+  unsigned int* hostSrc364;
+
+  unsigned int* deviceOut32;
+  unsigned int* deviceSrc032;
+  unsigned int* deviceSrc132;
+  unsigned int* deviceSrc232;
+  unsigned int* deviceSrc332;
+  unsigned long long int* deviceOut64;
+  unsigned long long int* deviceSrc064;
+  unsigned long long int* deviceSrc164;
+  unsigned int* deviceSrc264;
+  unsigned int* deviceSrc364;
+
+  hipDeviceProp_t devProp;
+  HIP_CHECK(hipGetDeviceProperties(&devProp, 0));
+  INFO("System minor : " << devProp.minor);
+  INFO("System major : " << devProp.major);
+  INFO("agent prop name : " << devProp.name);
+
+  INFO("hip Device prop succeeded");
+
+  unsigned int wave_size = devProp.warpSize;
+  unsigned int num_waves_per_block = 2;
+  unsigned int num_threads_per_block = wave_size * num_waves_per_block;
+  unsigned int num_blocks = 2;
+  unsigned int NUM = num_threads_per_block * num_blocks;
+
+  unsigned i;
+  int errors;
+
+  hostOut32 = (unsigned int*)malloc(NUM * sizeof(unsigned int));
+  hostSrc032 = (unsigned int*)malloc(NUM * sizeof(unsigned int));
+  hostSrc132 = (unsigned int*)malloc(NUM * sizeof(unsigned int));
+  hostSrc232 = (unsigned int*)malloc(NUM * sizeof(unsigned int));
+  hostSrc332 = (unsigned int*)malloc(NUM * sizeof(unsigned int));
+
+  hostOut64 = (unsigned long long int*)malloc(NUM * sizeof(unsigned long long int));
+  hostSrc064 = (unsigned long long int*)malloc(NUM * sizeof(unsigned long long int));
+  hostSrc164 = (unsigned long long int*)malloc(NUM * sizeof(unsigned long long int));
+  hostSrc264 = (unsigned int*)malloc(NUM * sizeof(unsigned int));
+  hostSrc364 = (unsigned int*)malloc(NUM * sizeof(unsigned int));
+
+  // initialize the input data
+  std::random_device rd;
+  std::uniform_int_distribution<uint32_t> uint32_src01_dist;
+  std::uniform_int_distribution<uint32_t> uint32_src23_dist(0, 31);
+  std::uniform_int_distribution<uint64_t> uint64_src01_dist;
+  std::uniform_int_distribution<uint32_t> uint64_src23_dist(0, 63);
+  for (i = 0; i < NUM; i++) {
+    hostOut32[i] = 0;
+    hostSrc032[i] = uint32_src01_dist(rd);
+    hostSrc132[i] = uint32_src01_dist(rd);
+    hostSrc232[i] = uint32_src23_dist(rd);
+    hostSrc232[i] = uint32_src23_dist(rd);
+    hostOut64[i] = 0;
+    hostSrc064[i] = uint64_src01_dist(rd);
+    hostSrc164[i] = uint64_src01_dist(rd);
+    hostSrc264[i] = uint64_src23_dist(rd);
+    hostSrc264[i] = uint64_src23_dist(rd);
+  }
+
+  HIP_CHECK(hipMalloc((void**)&deviceOut32, NUM * sizeof(unsigned int)));
+  HIP_CHECK(hipMalloc((void**)&deviceSrc032, NUM * sizeof(unsigned int)));
+  HIP_CHECK(hipMalloc((void**)&deviceSrc132, NUM * sizeof(unsigned int)));
+  HIP_CHECK(hipMalloc((void**)&deviceSrc232, NUM * sizeof(unsigned int)));
+  HIP_CHECK(hipMalloc((void**)&deviceSrc332, NUM * sizeof(unsigned int)));
+
+  HIP_CHECK(hipMalloc((void**)&deviceOut64, NUM * sizeof(unsigned long long int)));
+  HIP_CHECK(hipMalloc((void**)&deviceSrc064, NUM * sizeof(unsigned long long int)));
+  HIP_CHECK(hipMalloc((void**)&deviceSrc164, NUM * sizeof(unsigned long long int)));
+  HIP_CHECK(hipMalloc((void**)&deviceSrc264, NUM * sizeof(unsigned int)));
+  HIP_CHECK(hipMalloc((void**)&deviceSrc364, NUM * sizeof(unsigned int)));
+
+  HIP_CHECK(
+      hipMemcpy(deviceSrc032, hostSrc032, NUM * sizeof(unsigned int), hipMemcpyHostToDevice));
+  HIP_CHECK(
+      hipMemcpy(deviceSrc132, hostSrc132, NUM * sizeof(unsigned int), hipMemcpyHostToDevice));
+  HIP_CHECK(
+      hipMemcpy(deviceSrc232, hostSrc232, NUM * sizeof(unsigned int), hipMemcpyHostToDevice));
+  HIP_CHECK(
+      hipMemcpy(deviceSrc332, hostSrc332, NUM * sizeof(unsigned int), hipMemcpyHostToDevice));
+
+  HIP_CHECK(hipMemcpy(deviceSrc064, hostSrc064, NUM * sizeof(unsigned long long int),
+                       hipMemcpyHostToDevice));
+  HIP_CHECK(hipMemcpy(deviceSrc164, hostSrc164, NUM * sizeof(unsigned long long int),
+                       hipMemcpyHostToDevice));
+  HIP_CHECK(
+      hipMemcpy(deviceSrc264, hostSrc264, NUM * sizeof(unsigned int), hipMemcpyHostToDevice));
+  HIP_CHECK(
+      hipMemcpy(deviceSrc364, hostSrc364, NUM * sizeof(unsigned int), hipMemcpyHostToDevice));
+
+
+  hipLaunchKernelGGL(HIP_kernel, dim3(num_blocks), dim3(num_threads_per_block), 0, 0, deviceOut32,
+                     deviceSrc032, deviceSrc132, deviceSrc232, deviceSrc332, deviceOut64,
+                     deviceSrc064, deviceSrc164, deviceSrc264, deviceSrc364);
+
+
+  HIP_CHECK(hipMemcpy(hostOut32, deviceOut32, NUM * sizeof(unsigned int), hipMemcpyDeviceToHost));
+  HIP_CHECK(hipMemcpy(hostOut64, deviceOut64, NUM * sizeof(unsigned long long int),
+                       hipMemcpyDeviceToHost));
+
+  // verify the results
+  errors = 0;
+  for (i = 0; i < NUM; i++) {
+    if (hostOut32[i] !=
+        bit_insert<uint32_t>(hostSrc032[i], hostSrc132[i], hostSrc232[i], hostSrc332[i])) {
+      errors++;
+      INFO("device: " << hostOut32[i] << " host: "
+                      << bit_insert<uint32_t>(hostSrc032[i], hostSrc132[i], hostSrc232[i],
+                                              hostSrc332[i])
+                      << " " << hostSrc032[i] << " " << hostSrc132[i] << " " << hostSrc232[i] << " "
+                      << hostSrc332[i] << "\n");
+    }
+  }
+
+  for (i = 0; i < NUM; i++) {
+    if (hostOut64[i] !=
+        bit_insert<uint64_t>(hostSrc064[i], hostSrc164[i], hostSrc264[i], hostSrc364[i])) {
+      errors++;
+      INFO("device: " << hostOut64[i] << " host: "
+                      << bit_insert<uint64_t>(hostSrc064[i], hostSrc164[i], hostSrc264[i],
+                                              hostSrc364[i])
+                      << " " << hostSrc064[i] << " " << hostSrc164[i] << " " << hostSrc264[i] << " "
+                      << hostSrc364[i] << "\n");
+    }
+  }
+
+  HIP_CHECK(hipFree(deviceOut32));
+  HIP_CHECK(hipFree(deviceSrc032));
+  HIP_CHECK(hipFree(deviceSrc132));
+  HIP_CHECK(hipFree(deviceSrc232));
+  HIP_CHECK(hipFree(deviceSrc332));
+  HIP_CHECK(hipFree(deviceOut64));
+  HIP_CHECK(hipFree(deviceSrc064));
+  HIP_CHECK(hipFree(deviceSrc164));
+  HIP_CHECK(hipFree(deviceSrc264));
+  HIP_CHECK(hipFree(deviceSrc364));
+
+  free(hostOut32);
+  free(hostSrc032);
+  free(hostSrc132);
+  free(hostSrc232);
+  free(hostSrc332);
+  free(hostOut64);
+  free(hostSrc064);
+  free(hostSrc164);
+  free(hostSrc264);
+  free(hostSrc364);
+
+  REQUIRE(errors == 0);
+}
diff --git a/catch/unit/deviceLib/brev.cc b/catch/unit/deviceLib/brev.cc
new file mode 100644
index 0000000000..5631f97679
--- /dev/null
+++ b/catch/unit/deviceLib/brev.cc
@@ -0,0 +1,150 @@
+/*
+Copyright (c) 2021 - present 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>
+#include <hip/device_functions.h>
+
+#include <assert.h>
+#include <stdio.h>
+#include <algorithm>
+#include <stdlib.h>
+#include <iostream>
+
+
+#define WIDTH 32
+#define HEIGHT 32
+
+#define NUM (WIDTH * HEIGHT)
+
+#define THREADS_PER_BLOCK_X 8
+#define THREADS_PER_BLOCK_Y 8
+#define THREADS_PER_BLOCK_Z 1
+
+
+// CPU implementation of bitreverse
+template <typename T> T bitreverse(T num) {
+  T count = sizeof(num) * 8 - 1;
+  T reverse_num = num;
+
+  num >>= 1;
+  while (num) {
+    reverse_num <<= 1;
+    reverse_num |= num & 1;
+    num >>= 1;
+    count--;
+  }
+  reverse_num <<= count;
+  return reverse_num;
+}
+
+__global__ void HIP_kernel(unsigned int* a, unsigned int* b, unsigned long long int* c,
+                           unsigned long long int* d, int width, int height) {
+  int x = blockDim.x * blockIdx.x + threadIdx.x;
+  int y = blockDim.y * blockIdx.y + threadIdx.y;
+
+  int i = y * width + x;
+  if (i < (width * height)) {
+    a[i] = __brev(b[i]);
+    c[i] = __brevll(d[i]);
+  }
+}
+
+TEST_CASE("Unit_brev") {
+  using namespace std;
+  unsigned int* hostA;
+  unsigned int* hostB;
+  unsigned long long int* hostC;
+  unsigned long long int* hostD;
+
+  unsigned int* deviceA;
+  unsigned int* deviceB;
+  unsigned long long int* deviceC;
+  unsigned long long int* deviceD;
+
+  hipDeviceProp_t devProp;
+  HIP_CHECK(hipGetDeviceProperties(&devProp, 0));
+  INFO("System minor : " << devProp.minor);
+  INFO("System major : " << devProp.major);
+  INFO("agent prop name : " << devProp.name);
+
+  INFO("hip Device prop succeeded");
+
+  int i;
+  int errors;
+
+  hostA = (unsigned int*)malloc(NUM * sizeof(unsigned int));
+  hostB = (unsigned int*)malloc(NUM * sizeof(unsigned int));
+  hostC = (unsigned long long int*)malloc(NUM * sizeof(unsigned long long int));
+  hostD = (unsigned long long int*)malloc(NUM * sizeof(unsigned long long int));
+
+  // initialize the input data
+  for (i = 0; i < NUM; i++) {
+    hostB[i] = i;
+    hostD[i] = i;
+  }
+
+  HIP_CHECK(hipMalloc((void**)&deviceA, NUM * sizeof(unsigned int)));
+  HIP_CHECK(hipMalloc((void**)&deviceB, NUM * sizeof(unsigned int)));
+  HIP_CHECK(hipMalloc((void**)&deviceC, NUM * sizeof(unsigned long long int)));
+  HIP_CHECK(hipMalloc((void**)&deviceD, NUM * sizeof(unsigned long long int)));
+
+  HIP_CHECK(hipMemcpy(deviceB, hostB, NUM * sizeof(unsigned int), hipMemcpyHostToDevice));
+  HIP_CHECK(
+      hipMemcpy(deviceD, hostD, NUM * sizeof(unsigned long long int), hipMemcpyHostToDevice));
+
+
+  hipLaunchKernelGGL(HIP_kernel, dim3(WIDTH / THREADS_PER_BLOCK_X, HEIGHT / THREADS_PER_BLOCK_Y),
+                     dim3(THREADS_PER_BLOCK_X, THREADS_PER_BLOCK_Y), 0, 0, deviceA, deviceB,
+                     deviceC, deviceD, WIDTH, HEIGHT);
+
+
+  HIP_CHECK(hipMemcpy(hostA, deviceA, NUM * sizeof(unsigned int), hipMemcpyDeviceToHost));
+  HIP_CHECK(
+      hipMemcpy(hostC, deviceC, NUM * sizeof(unsigned long long int), hipMemcpyDeviceToHost));
+
+
+  // verify the results
+  errors = 0;
+  for (i = 0; i < NUM; i++) {
+    if (hostA[i] != bitreverse(hostB[i])) {
+      errors++;
+    }
+  }
+
+  for (i = 0; i < NUM; i++) {
+    if (hostC[i] != bitreverse(hostD[i])) {
+      errors++;
+    }
+  }
+
+  HIP_CHECK(hipFree(deviceA));
+  HIP_CHECK(hipFree(deviceB));
+  HIP_CHECK(hipFree(deviceC));
+  HIP_CHECK(hipFree(deviceD));
+
+  free(hostA);
+  free(hostB);
+  free(hostC);
+  free(hostD);
+
+  REQUIRE(errors == 0);
+}
\ No newline at end of file
diff --git a/catch/unit/deviceLib/clz.cc b/catch/unit/deviceLib/clz.cc
new file mode 100644
index 0000000000..0b5bf04b30
--- /dev/null
+++ b/catch/unit/deviceLib/clz.cc
@@ -0,0 +1,172 @@
+/*
+Copyright (c) 2021 - present 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>
+#include <hip/device_functions.h>
+
+#include <assert.h>
+#include <stdio.h>
+#include <algorithm>
+#include <stdlib.h>
+#include <iostream>
+
+#define WIDTH 8
+#define HEIGHT 8
+#define NUM (WIDTH * HEIGHT)
+
+#define THREADS_PER_BLOCK_X 8
+#define THREADS_PER_BLOCK_Y 8
+#define THREADS_PER_BLOCK_Z 1
+
+unsigned int firstbit_u32(unsigned int a) {
+  if (a == 0) {
+    return 32;
+  }
+  unsigned int pos = 0;
+  while ((int)a > 0) {
+    a <<= 1;
+    pos++;
+  }
+  return pos;
+}
+
+unsigned int firstbit_u64(unsigned long long int a) {
+  if (a == 0) {
+    return 64;
+  }
+  unsigned int pos = 0;
+  while ((long long int)a > 0) {
+    a <<= 1;
+    pos++;
+  }
+  return pos;
+}
+
+// Check implicit conversion will not cause ambiguity.
+__device__ void test_ambiguity() {
+  short s = 1;
+  unsigned short us = 2;
+  float f = 3;
+  int i = 4;
+  unsigned int ui = 5;
+  __clz(f);
+  __clz(s);
+  __clz(us);
+  __clzll(f);
+  __clzll(i);
+  __clzll(ui);
+}
+
+__global__ void clz_HIP_kernel(unsigned int* a, unsigned int* b, unsigned int* c,
+                               unsigned long long int* d, int width, int height) {
+  int x = blockDim.x * blockIdx.x + threadIdx.x;
+  int y = blockDim.y * blockIdx.y + threadIdx.y;
+
+  int i = y * width + x;
+  if (i < (width * height)) {
+    a[i] = __clz(b[i]);
+    c[i] = __clzll(d[i]);
+  }
+}
+
+TEST_CASE("Unit_clz") {
+  using namespace std;
+
+  unsigned int* hostA;
+  unsigned int* hostB;
+  unsigned int* hostC;
+  unsigned long long int* hostD;
+
+  unsigned int* deviceA;
+  unsigned int* deviceB;
+  unsigned int* deviceC;
+  unsigned long long int* deviceD;
+
+  hipDeviceProp_t devProp;
+  HIP_CHECK(hipGetDeviceProperties(&devProp, 0));
+  INFO("System minor : " << devProp.minor);
+  INFO("System major : " << devProp.major);
+  INFO("agent prop name : " << devProp.name);
+
+  INFO("hip Device prop succeeded");
+
+  unsigned int i;
+  int errors;
+
+  hostA = (unsigned int*)malloc(NUM * sizeof(unsigned int));
+  hostB = (unsigned int*)malloc(NUM * sizeof(unsigned int));
+  hostC = (unsigned int*)malloc(NUM * sizeof(unsigned int));
+  hostD = (unsigned long long int*)malloc(NUM * sizeof(unsigned long long int));
+
+  // initialize the input data
+  for (i = 0; i < NUM; i++) {
+    hostB[i] = 419430 * i;
+    hostD[i] = i;
+  }
+
+  HIP_CHECK(hipMalloc((void**)&deviceA, NUM * sizeof(unsigned int)));
+  HIP_CHECK(hipMalloc((void**)&deviceB, NUM * sizeof(unsigned int)));
+  HIP_CHECK(hipMalloc((void**)&deviceC, NUM * sizeof(unsigned int)));
+  HIP_CHECK(hipMalloc((void**)&deviceD, NUM * sizeof(unsigned long long int)));
+
+  HIP_CHECK(hipMemcpy(deviceB, hostB, NUM * sizeof(unsigned int), hipMemcpyHostToDevice));
+  HIP_CHECK(
+      hipMemcpy(deviceD, hostD, NUM * sizeof(unsigned long long int), hipMemcpyHostToDevice));
+
+  hipLaunchKernelGGL(clz_HIP_kernel,
+                     dim3(WIDTH / THREADS_PER_BLOCK_X, HEIGHT / THREADS_PER_BLOCK_Y),
+                     dim3(THREADS_PER_BLOCK_X, THREADS_PER_BLOCK_Y), 0, 0, deviceA, deviceB,
+                     deviceC, deviceD, WIDTH, HEIGHT);
+
+
+  HIP_CHECK(hipMemcpy(hostA, deviceA, NUM * sizeof(unsigned int), hipMemcpyDeviceToHost));
+  HIP_CHECK(hipMemcpy(hostC, deviceC, NUM * sizeof(unsigned int), hipMemcpyDeviceToHost));
+
+  // verify the results
+  errors = 0;
+  for (i = 0; i < NUM; i++) {
+    if (hostA[i] != firstbit_u32(hostB[i])) {
+      INFO("Match Failed: " << hostA[i] << " - " << firstbit_u32(hostB[i]));
+      errors++;
+    }
+  }
+
+  for (i = 0; i < NUM; i++) {
+    if (hostC[i] != firstbit_u64(hostD[i])) {
+      INFO("Match Failed: " << hostC[i] << " - " << firstbit_u32(hostD[i]));
+      errors++;
+    }
+  }
+
+
+  HIP_CHECK(hipFree(deviceA));
+  HIP_CHECK(hipFree(deviceB));
+  HIP_CHECK(hipFree(deviceC));
+  HIP_CHECK(hipFree(deviceD));
+
+  free(hostA);
+  free(hostB);
+  free(hostC);
+  free(hostD);
+
+  REQUIRE(errors == 0);
+}
diff --git a/catch/unit/deviceLib/ffs.cc b/catch/unit/deviceLib/ffs.cc
new file mode 100644
index 0000000000..5c94cc3d20
--- /dev/null
+++ b/catch/unit/deviceLib/ffs.cc
@@ -0,0 +1,146 @@
+/*
+Copyright (c) 2021 - present 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>
+#include <hip/device_functions.h>
+
+#include <assert.h>
+#include <stdio.h>
+#include <algorithm>
+#include <stdlib.h>
+#include <iostream>
+
+
+#define WIDTH 8
+#define HEIGHT 8
+
+#define NUM (WIDTH * HEIGHT)
+
+#define THREADS_PER_BLOCK_X 8
+#define THREADS_PER_BLOCK_Y 8
+#define THREADS_PER_BLOCK_Z 1
+
+template <typename T> unsigned lastbit(T a) {
+  if (a == 0) return 0;
+  unsigned pos = 1;
+  while ((a & 1) != 1) {
+    a >>= 1;
+    pos++;
+  }
+  return pos;
+}
+
+__global__ void ffs_HIP_kernel(unsigned int* a, unsigned int* b, unsigned int* c,
+                               unsigned long long int* d, int width, int height) {
+  int x = blockDim.x * blockIdx.x + threadIdx.x;
+  int y = blockDim.y * blockIdx.y + threadIdx.y;
+
+  int i = y * width + x;
+  if (i < (width * height)) {
+    a[i] = __ffs(b[i]);
+    c[i] = __ffsll(d[i]);
+  }
+}
+
+TEST_CASE("Unit_ffs") {
+  using namespace std;
+
+  unsigned int* hostA;
+  unsigned int* hostB;
+  unsigned int* hostC;
+  unsigned long long int* hostD;
+
+  unsigned int* deviceA;
+  unsigned int* deviceB;
+  unsigned int* deviceC;
+  unsigned long long int* deviceD;
+
+  hipDeviceProp_t devProp;
+  HIP_CHECK(hipGetDeviceProperties(&devProp, 0));
+  INFO("System minor : " << devProp.minor);
+  INFO("System major : " << devProp.major);
+  INFO("agent prop name : " << devProp.name);
+
+  INFO("hip Device prop succeeded");
+
+
+  int i;
+  int errors;
+
+  hostA = (unsigned int*)malloc(NUM * sizeof(unsigned int));
+  hostB = (unsigned int*)malloc(NUM * sizeof(unsigned int));
+  hostC = (unsigned int*)malloc(NUM * sizeof(unsigned int));
+  hostD = (unsigned long long int*)malloc(NUM * sizeof(unsigned long long int));
+
+  // initialize the input data
+  for (i = 0; i < NUM; i++) {
+    hostB[i] = i;
+    hostD[i] = 1099511627776 + i;
+  }
+
+  HIP_CHECK(hipMalloc((void**)&deviceA, NUM * sizeof(unsigned int)));
+  HIP_CHECK(hipMalloc((void**)&deviceB, NUM * sizeof(unsigned int)));
+  HIP_CHECK(hipMalloc((void**)&deviceC, NUM * sizeof(unsigned int)));
+  HIP_CHECK(hipMalloc((void**)&deviceD, NUM * sizeof(unsigned long long int)));
+
+  HIP_CHECK(hipMemcpy(deviceB, hostB, NUM * sizeof(unsigned int), hipMemcpyHostToDevice));
+  HIP_CHECK(
+      hipMemcpy(deviceD, hostD, NUM * sizeof(unsigned long long int), hipMemcpyHostToDevice));
+
+  hipLaunchKernelGGL(ffs_HIP_kernel,
+                     dim3(WIDTH / THREADS_PER_BLOCK_X, HEIGHT / THREADS_PER_BLOCK_Y),
+                     dim3(THREADS_PER_BLOCK_X, THREADS_PER_BLOCK_Y), 0, 0, deviceA, deviceB,
+                     deviceC, deviceD, WIDTH, HEIGHT);
+
+
+  HIP_CHECK(hipMemcpy(hostA, deviceA, NUM * sizeof(unsigned int), hipMemcpyDeviceToHost));
+  HIP_CHECK(hipMemcpy(hostC, deviceC, NUM * sizeof(unsigned int), hipMemcpyDeviceToHost));
+
+  // verify the results
+  errors = 0;
+  for (i = 0; i < NUM; i++) {
+    if (hostA[i] != lastbit(hostB[i])) {
+      INFO("Mismatch: " << hostA[i] << " - " << lastbit(hostB[i]));
+      errors++;
+    }
+  }
+
+  for (i = 0; i < NUM; i++) {
+    if (hostC[i] != lastbit(hostD[i])) {
+      INFO("Mismatch: " << hostC[i] << " - " << lastbit(hostD[i]));
+      errors++;
+    }
+  }
+
+  HIP_CHECK(hipFree(deviceA));
+  HIP_CHECK(hipFree(deviceB));
+  HIP_CHECK(hipFree(deviceC));
+  HIP_CHECK(hipFree(deviceD));
+
+
+  free(hostA);
+  free(hostB);
+  free(hostC);
+  free(hostD);
+
+  REQUIRE(errors == 0);
+}
diff --git a/catch/unit/deviceLib/floatMath.cc b/catch/unit/deviceLib/floatMath.cc
index c163666e83..ee6037635d 100644
--- a/catch/unit/deviceLib/floatMath.cc
+++ b/catch/unit/deviceLib/floatMath.cc
@@ -1,3 +1,25 @@
+/*
+Copyright (c) 2021 - present 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>
 
 #define LEN 512
@@ -27,7 +49,7 @@ __global__ void floatMath(float* In, float* Out) {
   Out[tid] = __tanf(Out[tid]);
 }
 
-TEST_CASE("FloatMathTest") {
+TEST_CASE("Unit_deviceFunctions_CompileTest") {
   float *Ind, *Outd;
   auto res = hipMalloc((void**)&Ind, SIZE);
   REQUIRE(res == hipSuccess);
@@ -38,4 +60,6 @@ TEST_CASE("FloatMathTest") {
   REQUIRE(res == hipSuccess);
   res = hipGetLastError();
   REQUIRE(res == hipSuccess);
+  HIP_CHECK(hipFree(Ind));
+  HIP_CHECK(hipFree(Outd));
 }
diff --git a/catch/unit/deviceLib/floatTM.cc b/catch/unit/deviceLib/floatTM.cc
new file mode 100644
index 0000000000..e4f2a05e2e
--- /dev/null
+++ b/catch/unit/deviceLib/floatTM.cc
@@ -0,0 +1,226 @@
+/*
+Copyright (c) 2021 - present 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>
+
+#include <type_traits>
+#include <random>
+
+
+static std::random_device dev;
+static std::mt19937 rng(dev());
+
+template <typename T, typename M> __host__ __device__ inline constexpr int count() {
+  return sizeof(T) / sizeof(M);
+}
+
+inline float getRandomFloat(float min = 10, float max = 100) {
+  std::uniform_real_distribution<float> gen(min, max);
+  return gen(rng);
+}
+
+template <typename T, typename B> void fillMatrix(T* a, int size) {
+  for (int i = 0; i < size; i++) {
+    T t;
+    t.x = getRandomFloat();
+    if constexpr (count<T, B>() >= 2) t.y = getRandomFloat();
+    if constexpr (count<T, B>() >= 3) t.z = getRandomFloat();
+    if constexpr (count<T, B>() >= 4) t.w = getRandomFloat();
+
+    a[i] = t;
+  }
+}
+
+// Test operations
+template <typename T, typename B> __host__ __device__ void testOperations(T& a, T& b) {
+  a.x += b.x;
+  a.x++;
+  b.x++;
+  if constexpr (count<T, B>() >= 2) {
+    a.y = b.x;
+    a.x = b.y;
+  }
+  if constexpr (count<T, B>() >= 3) {
+    if (a.x > 0) b.x /= a.x;
+    a.x *= b.z;
+    a.y--;
+  }
+  if constexpr (count<T, B>() >= 4) {
+    b.w = a.x;
+    a.w += (-b.y);
+  }
+}
+
+template <typename T, typename B> __global__ void testOperationsGPU(T* d_a, T* d_b, int size) {
+  int id = threadIdx.x;
+  if (id > size) return;
+  T& a = d_a[id];
+  T& b = d_b[id];
+
+  testOperations<T, B>(a, b);
+}
+
+
+template <typename T> void dcopy(T* a, T* b, int size) {
+  for (int i = 0; i < size; i++) {
+    a[i] = b[i];
+  }
+}
+
+template <typename T> bool isEqual(T* a, T* b, int size) {
+  for (int i = 0; i < size; i++) {
+    if (a[i] != b[i]) {
+      return false;
+    }
+  }
+  return true;
+}
+
+// Main function that tests type
+// T = what you want to test
+// D = pack of 1 i.e. float1 int1
+template <typename T, typename D> void testType(int msize) {
+  T *fa, *fb, *fc, *h_fa, *h_fb;
+  fa = new T[msize];
+  fb = new T[msize];
+  fc = new T[msize];
+  h_fa = new T[msize];
+  h_fb = new T[msize];
+
+  T *d_fa, *d_fb;
+
+  constexpr int c = count<T, D>();
+
+  if (c <= 0 || c >= 5) {
+    INFO("It should be between 1 to 5");
+    REQUIRE(false);
+  }
+
+  fillMatrix<T, D>(fa, msize);
+  dcopy(fb, fa, msize);
+  dcopy(h_fa, fa, msize);
+  dcopy(h_fb, fa, msize);
+  for (int i = 0; i < msize; i++) testOperations<T, D>(h_fa[i], h_fb[i]);
+
+  HIP_CHECK(hipMalloc(&d_fa, sizeof(T) * msize));
+  HIP_CHECK(hipMalloc(&d_fb, sizeof(T) * msize));
+
+  HIP_CHECK(hipMemcpy(d_fa, fa, sizeof(T) * msize, hipMemcpyHostToDevice));
+  HIP_CHECK(hipMemcpy(d_fb, fb, sizeof(T) * msize, hipMemcpyHostToDevice));
+
+  auto kernel = testOperationsGPU<T, D>;
+  hipLaunchKernelGGL(kernel, 1, msize, 0, 0, d_fa, d_fb, msize);
+
+  HIP_CHECK(hipMemcpy(fc, d_fa, sizeof(T) * msize, hipMemcpyDeviceToHost));
+
+  bool pass = true;
+  if (!isEqual<T>(h_fa, fc, msize)) {
+    pass = false;
+  }
+
+  delete[] fa;
+  delete[] fb;
+  delete[] fc;
+  delete[] h_fa;
+  delete[] h_fb;
+  HIP_CHECK(hipFree(d_fa));
+  HIP_CHECK(hipFree(d_fb));
+
+  REQUIRE(pass);
+}
+
+
+TEST_CASE("Unit_floatTM") {
+  constexpr int msize = 100;
+
+  // double
+  testType<double1, double1>(msize);
+  testType<double2, double1>(msize);
+  testType<double3, double1>(msize);
+  testType<double4, double1>(msize);
+
+  // floats
+  testType<float1, float1>(msize);
+  testType<float2, float1>(msize);
+  testType<float3, float1>(msize);
+  testType<float4, float1>(msize);
+
+  // ints
+  testType<int1, int1>(msize);
+  testType<int2, int1>(msize);
+  testType<int3, int1>(msize);
+  testType<int4, int1>(msize);
+
+  // chars
+  testType<char1, char1>(msize);
+  testType<char2, char1>(msize);
+  testType<char3, char1>(msize);
+  testType<char4, char1>(msize);
+
+  // long
+  testType<long1, long1>(msize);
+  testType<long2, long1>(msize);
+  testType<long3, long1>(msize);
+  testType<long4, long1>(msize);
+
+  // longlong
+  testType<longlong1, longlong1>(msize);
+  testType<longlong2, longlong1>(msize);
+  testType<longlong3, longlong1>(msize);
+  testType<longlong4, longlong1>(msize);
+
+  // short
+  testType<short1, short1>(msize);
+  testType<short2, short1>(msize);
+  testType<short3, short1>(msize);
+  testType<short4, short1>(msize);
+
+  // uints
+  testType<uint1, uint1>(msize);
+  testType<uint2, uint1>(msize);
+  testType<uint3, uint1>(msize);
+  testType<uint4, uint1>(msize);
+
+  // uchars
+  testType<uchar1, uchar1>(msize);
+  testType<uchar2, uchar1>(msize);
+  testType<uchar3, uchar1>(msize);
+  testType<uchar4, uchar1>(msize);
+
+  // ulong
+  testType<ulong1, ulong1>(msize);
+  testType<ulong2, ulong1>(msize);
+  testType<ulong3, ulong1>(msize);
+  testType<ulong4, ulong1>(msize);
+
+  // ulonglong
+  testType<ulonglong1, ulonglong1>(msize);
+  testType<ulonglong2, ulonglong1>(msize);
+  testType<ulonglong3, ulonglong1>(msize);
+  testType<ulonglong4, ulonglong1>(msize);
+
+  // ushort
+  testType<ushort1, ushort1>(msize);
+  testType<ushort2, ushort1>(msize);
+  testType<ushort3, ushort1>(msize);
+  testType<ushort4, ushort1>(msize);
+}
diff --git a/catch/unit/deviceLib/funnelshift.cc b/catch/unit/deviceLib/funnelshift.cc
new file mode 100644
index 0000000000..ef2a96f73a
--- /dev/null
+++ b/catch/unit/deviceLib/funnelshift.cc
@@ -0,0 +1,208 @@
+/*
+Copyright (c) 2021 - present 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>
+#include <hip/device_functions.h>
+
+
+#include <assert.h>
+#include <stdio.h>
+#include <stdint.h>
+#include <stdlib.h>
+
+
+#define NUM_TESTS 65
+
+#define HI_INT 0xfacefeed
+#define LO_INT 0xdeadbeef
+
+__global__ void funnelshift_kernel(unsigned int* l_out, unsigned int* lc_out, unsigned int* r_out,
+                                   unsigned int* rc_out) {
+  for (int i = 0; i < NUM_TESTS; i++) {
+    l_out[i] = __funnelshift_l(LO_INT, HI_INT, i);
+    lc_out[i] = __funnelshift_lc(LO_INT, HI_INT, i);
+    r_out[i] = __funnelshift_r(LO_INT, HI_INT, i);
+    rc_out[i] = __funnelshift_rc(LO_INT, HI_INT, i);
+  }
+}
+
+static unsigned int cpu_funnelshift_l(unsigned int lo, unsigned int hi, unsigned int shift) {
+  // Concatenate hi:lo
+  uint64_t val = hi;
+  val <<= 32;
+  val |= lo;
+  // left shift by intput & 31
+  val <<= (shift & 31);
+  // pull out upper 32 bits and return them
+  val >>= 32;
+  return val & 0xffffffff;
+}
+
+static unsigned int cpu_funnelshift_lc(unsigned int lo, unsigned int hi, unsigned int shift) {
+  // Concatenate hi:lo
+  uint64_t val = hi;
+  val <<= 32;
+  val |= lo;
+  // left shift by min(input,32)
+  if (shift > 32) shift = 32;
+  val <<= shift;
+  // pull out upper 32 bits and return them
+  val >>= 32;
+  return val & 0xffffffff;
+}
+
+static unsigned int cpu_funnelshift_r(unsigned int lo, unsigned int hi, unsigned int shift) {
+  // Concatenate hi:lo
+  uint64_t val = hi;
+  val <<= 32;
+  val |= lo;
+  // right shift by intput & 31
+  val >>= (shift & 31);
+  // return lower 32 bits
+  return val & 0xffffffff;
+}
+
+static unsigned int cpu_funnelshift_rc(unsigned int lo, unsigned int hi, unsigned int shift) {
+  // Concatenate hi:lo
+  uint64_t val = hi;
+  val <<= 32;
+  val |= lo;
+  // left shift by min(input, 32)
+  if (shift > 32) shift = 32;
+  val >>= shift;
+  // return lower 32 bits
+  return val & 0xffffffff;
+}
+
+TEST_CASE("Unit_funnelshift") {
+  unsigned int* host_l_output;
+  unsigned int* host_lc_output;
+  unsigned int* host_r_output;
+  unsigned int* host_rc_output;
+
+  unsigned int* device_l_output;
+  unsigned int* device_lc_output;
+  unsigned int* device_r_output;
+  unsigned int* device_rc_output;
+
+  unsigned int* golden_l;
+  unsigned int* golden_lc;
+  unsigned int* golden_r;
+  unsigned int* golden_rc;
+
+  hipDeviceProp_t devProp;
+  HIP_CHECK(hipGetDeviceProperties(&devProp, 0));
+  INFO("System minor : " << devProp.minor);
+  INFO("System major : " << devProp.major);
+  INFO("agent prop name : " << devProp.name);
+
+  INFO("hip Device prop succeeded");
+
+
+  int i;
+  int errors;
+
+  host_l_output = (unsigned int*)calloc(NUM_TESTS, sizeof(unsigned int));
+  host_lc_output = (unsigned int*)calloc(NUM_TESTS, sizeof(unsigned int));
+  host_r_output = (unsigned int*)calloc(NUM_TESTS, sizeof(unsigned int));
+  host_rc_output = (unsigned int*)calloc(NUM_TESTS, sizeof(unsigned int));
+
+  golden_l = (unsigned int*)calloc(NUM_TESTS, sizeof(unsigned int));
+  golden_lc = (unsigned int*)calloc(NUM_TESTS, sizeof(unsigned int));
+  golden_r = (unsigned int*)calloc(NUM_TESTS, sizeof(unsigned int));
+  golden_rc = (unsigned int*)calloc(NUM_TESTS, sizeof(unsigned int));
+
+  for (int i = 0; i < NUM_TESTS; i++) {
+    golden_l[i] = cpu_funnelshift_l(LO_INT, HI_INT, i);
+    golden_lc[i] = cpu_funnelshift_lc(LO_INT, HI_INT, i);
+    golden_r[i] = cpu_funnelshift_r(LO_INT, HI_INT, i);
+    golden_rc[i] = cpu_funnelshift_rc(LO_INT, HI_INT, i);
+  }
+
+  HIP_CHECK(hipMalloc((void**)&device_l_output, NUM_TESTS * sizeof(unsigned int)));
+  HIP_CHECK(hipMalloc((void**)&device_lc_output, NUM_TESTS * sizeof(unsigned int)));
+  HIP_CHECK(hipMalloc((void**)&device_r_output, NUM_TESTS * sizeof(unsigned int)));
+  HIP_CHECK(hipMalloc((void**)&device_rc_output, NUM_TESTS * sizeof(unsigned int)));
+
+  hipLaunchKernelGGL(funnelshift_kernel, dim3(1), dim3(1), 0, 0, device_l_output, device_lc_output,
+                     device_r_output, device_rc_output);
+
+  HIP_CHECK(hipMemcpy(host_l_output, device_l_output, NUM_TESTS * sizeof(unsigned int),
+                       hipMemcpyDeviceToHost));
+  HIP_CHECK(hipMemcpy(host_lc_output, device_lc_output, NUM_TESTS * sizeof(unsigned int),
+                       hipMemcpyDeviceToHost));
+  HIP_CHECK(hipMemcpy(host_r_output, device_r_output, NUM_TESTS * sizeof(unsigned int),
+                       hipMemcpyDeviceToHost));
+  HIP_CHECK(hipMemcpy(host_rc_output, device_rc_output, NUM_TESTS * sizeof(unsigned int),
+                       hipMemcpyDeviceToHost));
+
+  // verify the results
+  errors = 0;
+  printf("HI val: 0x%x\n", HI_INT);
+  printf("LO val: 0x%x\n", LO_INT);
+
+  for (i = 0; i < NUM_TESTS; i++) {
+    if (host_l_output[i] != golden_l[i]) {
+      errors++;
+      INFO("Mismatch: " << host_l_output[i] << " - " << golden_l[i]);
+      break;
+    }
+  }
+
+
+  for (i = 0; i < NUM_TESTS && errors == 0; i++) {
+    if (host_lc_output[i] != golden_lc[i]) {
+      errors++;
+      INFO("Mismatch: " << host_lc_output[i] << " - " << golden_lc[i]);
+      break;
+    }
+  }
+
+  errors = 0;
+  for (i = 0; i < NUM_TESTS && errors == 0; i++) {
+    if (host_r_output[i] != golden_r[i]) {
+      errors++;
+      INFO("Mismatch: " << host_r_output[i] << " - " << golden_r[i]);
+      break;
+    }
+  }
+
+  for (i = 0; i < NUM_TESTS && errors == 0; i++) {
+    if (host_rc_output[i] != golden_rc[i]) {
+      errors++;
+      INFO("Mismatch: " << host_rc_output[i] << " - " << golden_rc[i]);
+      break;
+    }
+  }
+
+  HIP_CHECK(hipFree(device_l_output));
+  HIP_CHECK(hipFree(device_lc_output));
+  HIP_CHECK(hipFree(device_r_output));
+  HIP_CHECK(hipFree(device_rc_output));
+
+  free(host_l_output);
+  free(host_lc_output);
+  free(host_r_output);
+  free(host_rc_output);
+
+  REQUIRE(errors == 0);
+}
\ No newline at end of file
diff --git a/catch/unit/deviceLib/vectorTypesDevice.cc b/catch/unit/deviceLib/vectorTypesDevice.cc
index 331766190e..8676a34264 100644
--- a/catch/unit/deviceLib/vectorTypesDevice.cc
+++ b/catch/unit/deviceLib/vectorTypesDevice.cc
@@ -1,3 +1,25 @@
+/*
+Copyright (c) 2021 - present 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>
 #include <type_traits>
 #include <memory>
@@ -218,7 +240,7 @@ template <typename V, typename... Vs> bool run_CheckSharedVectorTypes() {
   return run_CheckSharedVectorType<V>() && run_CheckSharedVectorTypes<Vs...>();
 }
 
-TEST_CASE("VectorTypesTest") {
+TEST_CASE("Unit_vectorTypes_CompileTest") {
   static_assert(sizeof(float1) == 4, "");
   static_assert(sizeof(float2) >= 8, "");
   static_assert(sizeof(float3) >= 12, "");
@@ -242,5 +264,7 @@ TEST_CASE("VectorTypesTest") {
                uint3, uint4, long1, long2, long3, long4, ulong1, ulong2, ulong3, ulong4, longlong1,
                longlong2, longlong3, longlong4, ulonglong1, ulonglong2, ulonglong3, ulonglong4,
                float1, float2, float3, float4, double1, double2, double3, double4>();
+
+  HIP_CHECK(hipFree(ptr));
   REQUIRE(passed == true);
 }
diff --git a/catch/unit/memory/malloc.cc b/catch/unit/memory/malloc.cc
index dcfb6b53f4..24680b3d77 100644
--- a/catch/unit/memory/malloc.cc
+++ b/catch/unit/memory/malloc.cc
@@ -1,6 +1,6 @@
 #include <hip_test_common.hh>
 
-TEST_CASE("HostAllocBasic") {
+TEST_CASE("Unit_hipHostMalloc_4bytes") {
   int* d_a;
   auto res = hipHostMalloc(&d_a, sizeof(int), 0);
   REQUIRE(res == hipSuccess);
@@ -8,7 +8,7 @@ TEST_CASE("HostAllocBasic") {
   REQUIRE(res == hipSuccess);
 }
 
-TEST_CASE("AllocateBasic") {
+TEST_CASE("Unit_hipMalloc_4bytes") {
   int* d_a;
   auto res = hipMalloc(&d_a, sizeof(int)); 
   REQUIRE(res == hipSuccess);
diff --git a/catch/unit/memory/memset.cc b/catch/unit/memory/memset.cc
index 953decde1e..677d8a2d85 100644
--- a/catch/unit/memory/memset.cc
+++ b/catch/unit/memory/memset.cc
@@ -1,6 +1,6 @@
 #include <hip_test_common.hh>
 
-TEST_CASE("MemsetBasic") {
+TEST_CASE("Unit_hipMemset_4bytes") {
   int* d_a;
   auto res = hipMalloc(&d_a, sizeof(int));
   REQUIRE(res == hipSuccess);
@@ -9,7 +9,7 @@ TEST_CASE("MemsetBasic") {
   hipFree(d_a);
 }
 
-TEST_CASE("HostMemsetBasic") {
+TEST_CASE("Unit_hipMemset_4bytes_hostMem") {
   int* d_a;
   auto res = hipHostMalloc(&d_a, sizeof(int), 0);
   REQUIRE(res == hipSuccess);