diff --git a/projects/hip-tests/catch/unit/memory/hipMemCreate.cc b/projects/hip-tests/catch/unit/memory/hipMemCreate.cc
index c6823505ac..83f6ff7c01 100644
--- a/projects/hip-tests/catch/unit/memory/hipMemCreate.cc
+++ b/projects/hip-tests/catch/unit/memory/hipMemCreate.cc
@@ -37,6 +37,15 @@ THE SOFTWARE.
 #define NUM_OF_BUFFERS 3
 #define DATA_SIZE (1 << 13)
 
+/**
+ Kernel to perform Square of input data.
+ */
+static __global__ void square_kernel(int* Buff) {
+  int i = threadIdx.x + blockDim.x * blockIdx.x;
+  int temp = Buff[i] * Buff[i];
+  Buff[i] = temp;
+}
+
 /**
  * Test Description
  * ------------------------
@@ -201,9 +210,7 @@ TEST_CASE("Unit_hipMemCreate_ChkDev2HstMemcpy_ReleaseHdlPreUse") {
  * ------------------------
  *    - HIP_VERSION >= 6.1
  */
-#if HT_NVIDIA
-// This test is disabled. Will be enabled once VMM feature is fully
-// available
+
 TEST_CASE("Unit_hipMemCreate_ChkWithKerLaunch") {
   size_t granularity = 0;
   constexpr int N = DATA_SIZE;
@@ -253,7 +260,6 @@ TEST_CASE("Unit_hipMemCreate_ChkWithKerLaunch") {
   HIP_CHECK(hipMemUnmap(ptrA, size_mem));
   HIP_CHECK(hipMemAddressFree(ptrA, size_mem));
 }
-#endif
 
 /**
  * Test Description
@@ -268,9 +274,6 @@ TEST_CASE("Unit_hipMemCreate_ChkWithKerLaunch") {
  * ------------------------
  *    - HIP_VERSION >= 6.1
  */
-#if HT_NVIDIA
-// This test is disabled. Will be enabled once VMM feature is fully
-// available
 TEST_CASE("Unit_hipMemCreate_MapNonContiguousChunks") {
   size_t granularity = 0;
   constexpr int numOfBuffers = NUM_OF_BUFFERS;
@@ -333,7 +336,6 @@ TEST_CASE("Unit_hipMemCreate_MapNonContiguousChunks") {
   }
   HIP_CHECK(hipMemAddressFree(ptrA, (numOfBuffers * size_mem)));
 }
-#endif
 
 /**
  * Test Description
diff --git a/projects/hip-tests/catch/unit/memory/hipMemSetGetAccess.cc b/projects/hip-tests/catch/unit/memory/hipMemSetGetAccess.cc
index 220c2c11ba..20910f9803 100644
--- a/projects/hip-tests/catch/unit/memory/hipMemSetGetAccess.cc
+++ b/projects/hip-tests/catch/unit/memory/hipMemSetGetAccess.cc
@@ -44,6 +44,16 @@ THE SOFTWARE.
 #define DATA_SIZE (1 << 13)
 #define NEW_DATA_SIZE (2*DATA_SIZE)
 constexpr int initializer = 0;
+
+/**
+ Kernel to perform Square of input data.
+ */
+static __global__ void square_kernel(int* Buff) {
+  int i = threadIdx.x + blockDim.x * blockIdx.x;
+  int temp = Buff[i] * Buff[i];
+  Buff[i] = temp;
+}
+
 /**
  * Test Description
  * ------------------------
@@ -58,9 +68,6 @@ constexpr int initializer = 0;
  * ------------------------
  *    - HIP_VERSION >= 6.1
  */
-#if HT_NVIDIA
-// This test is disabled. Will be enabled once VMM feature is fully
-// available
 TEST_CASE("Unit_hipMemSetAccess_SetGet") {
   size_t granularity = 0;
   constexpr int N = DATA_SIZE;
@@ -134,7 +141,7 @@ TEST_CASE("Unit_hipMemSetAccess_MultDevSetGet") {
   size_t granularity = 0;
   constexpr int N = DATA_SIZE;
   size_t buffer_size = N * sizeof(int);
-  int deviceId = 0, value = 0, device_count = 0;
+  int deviceId = 0, device_count = 0;
   hipDevice_t device0, device1;
   HIP_CHECK(hipGetDeviceCount(&device_count));
   if (device_count < 2) {
@@ -188,7 +195,6 @@ TEST_CASE("Unit_hipMemSetAccess_MultDevSetGet") {
   HIP_CHECK(hipMemUnmap(ptrA, size_mem));
   HIP_CHECK(hipMemAddressFree(ptrA, size_mem));
 }
-#endif
 
 /**
  * Test Description
@@ -325,9 +331,6 @@ TEST_CASE("Unit_hipMemGetAccess_NegTst") {
  * ------------------------
  *    - HIP_VERSION >= 6.1
  */
-#if HT_NVIDIA
-// This test is disabled. Will be enabled once VMM feature is fully
-// available
 TEST_CASE("Unit_hipMemSetAccess_FuncTstOnMultDev") {
   size_t granularity = 0;
   constexpr int N = DATA_SIZE;
@@ -378,7 +381,6 @@ TEST_CASE("Unit_hipMemSetAccess_FuncTstOnMultDev") {
     HIP_CHECK(hipMemAddressFree(ptrA, size_mem));
   }
 }
-#endif
 
 /**
  * Test Description
@@ -1018,9 +1020,6 @@ TEST_CASE("Unit_hipMemSetAccess_Vmm2VMMInterDevMemCpy") {
  * ------------------------
  *    - HIP_VERSION >= 6.1
  */
-#if HT_NVIDIA
-// This test is disabled. Will be enabled once VMM feature is fully
-// available
 TEST_CASE("Unit_hipMemSetAccess_MapPhysChksFromMulDev") {
   int devicecount = 0;
   HIP_CHECK(hipGetDeviceCount(&devicecount));
@@ -1097,7 +1096,6 @@ TEST_CASE("Unit_hipMemSetAccess_MapPhysChksFromMulDev") {
   free(handle);
   free(size_mem);
 }
-#endif
 
 /**
  * Test Description
@@ -1225,9 +1223,7 @@ class vmm_resize_class {
     return 0;
   }
 };
-#if HT_NVIDIA
-// This test is disabled. Will be enabled once VMM feature is fully
-// available
+
 TEST_CASE("Unit_hipMemSetAccess_GrowVMM") {
   hipDeviceptr_t ptr;
   constexpr int N = DATA_SIZE;
@@ -1278,7 +1274,6 @@ TEST_CASE("Unit_hipMemSetAccess_GrowVMM") {
   free(ptrA_h);
   resizeobj.free_vmm();
 }
-#endif
 
 /**
  * Test Description
diff --git a/projects/hip-tests/catch/unit/memory/hip_vmm_common.hh b/projects/hip-tests/catch/unit/memory/hip_vmm_common.hh
index 83d209ac97..24ecb6408a 100644
--- a/projects/hip-tests/catch/unit/memory/hip_vmm_common.hh
+++ b/projects/hip-tests/catch/unit/memory/hip_vmm_common.hh
@@ -22,6 +22,8 @@ THE SOFTWARE.
 
 #pragma once
 
+#include "hip_test_context.hh"
+
 #define checkVMMSupported(device) {\
   int value = 0;\
   hipDeviceAttribute_t\
@@ -46,14 +48,4 @@ THE SOFTWARE.
   }\
 }
 
-#if HT_NVIDIA
-constexpr int threadsPerBlk = 64;
-/**
- Kernel to perform Square of input data.
- */
-static __global__ void square_kernel(int* Buff) {
-  int i = threadIdx.x + blockDim.x * blockIdx.x;
-  int temp = Buff[i] * Buff[i];
-  Buff[i] = temp;
-}
-#endif
+constexpr int threadsPerBlk = 64;
\ No newline at end of file