diff --git a/projects/hip-tests/catch/unit/device/chkUUIDInGrandChild_Exe.cc b/projects/hip-tests/catch/unit/device/chkUUIDInGrandChild_Exe.cc
index aedc65c640..2b43ee3f0d 100644
--- a/projects/hip-tests/catch/unit/device/chkUUIDInGrandChild_Exe.cc
+++ b/projects/hip-tests/catch/unit/device/chkUUIDInGrandChild_Exe.cc
@@ -19,6 +19,7 @@ THE SOFTWARE.
 
 #include <hip/hip_runtime.h>
 #include <cstring>
+#include <string>
 int main(int argc, char** argv) {
   if (argc < 0) {
     return -1;
diff --git a/projects/hip-tests/catch/unit/device/multipleUUID_Exe.cc b/projects/hip-tests/catch/unit/device/multipleUUID_Exe.cc
index 60237683f6..9b0d14f409 100644
--- a/projects/hip-tests/catch/unit/device/multipleUUID_Exe.cc
+++ b/projects/hip-tests/catch/unit/device/multipleUUID_Exe.cc
@@ -18,6 +18,7 @@ THE SOFTWARE.
 */
 #include <hip/hip_runtime.h>
 #include <cstring>
+#include <string>
 #include <vector>
 int main(int argc, char** argv) {
   if (argc < 0) {
diff --git a/projects/hip-tests/catch/unit/device/setEnvInChildProc_Exe.cc b/projects/hip-tests/catch/unit/device/setEnvInChildProc_Exe.cc
index 53dfc3bf0e..7242c00a04 100644
--- a/projects/hip-tests/catch/unit/device/setEnvInChildProc_Exe.cc
+++ b/projects/hip-tests/catch/unit/device/setEnvInChildProc_Exe.cc
@@ -17,7 +17,7 @@ OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
 THE SOFTWARE.
 */
 #include <hip/hip_runtime.h>
-#include <cstring>
+#include <string>
 #ifdef _WIN64
 #define setenv(x, y, z) _putenv_s(x, y)
 #define unsetenv(x) _putenv(x)
diff --git a/projects/hip-tests/catch/unit/graph/hipLaunchHostFunc.cc b/projects/hip-tests/catch/unit/graph/hipLaunchHostFunc.cc
index 49ee0c867c..2f5a7432c1 100644
--- a/projects/hip-tests/catch/unit/graph/hipLaunchHostFunc.cc
+++ b/projects/hip-tests/catch/unit/graph/hipLaunchHostFunc.cc
@@ -211,7 +211,8 @@ TEST_CASE("Unit_hipLaunchHostFunc_H2D_Kernel_D2H_Capture") {
   HIP_CHECK(hipStreamBeginCapture(captureStream, hipStreamCaptureModeGlobal));
 
   // Device alloc (async so it belongs to the captured stream)
-  HIP_CHECK(hipMallocAsync(&d_arrayA, arraySize * sizeof(double), captureStream));
+  HIP_CHECK(hipMallocAsync(reinterpret_cast<void**>(&d_arrayA), arraySize * sizeof(double),
+                           captureStream));
 
   // Initialize host data via a host function in the stream
   set_vector_args args{h_array, initValue};
diff --git a/projects/hip-tests/catch/unit/texture/hipTexRefSetGetFilterMode.cc b/projects/hip-tests/catch/unit/texture/hipTexRefSetGetFilterMode.cc
index 6203e2d85e..b336a4df93 100644
--- a/projects/hip-tests/catch/unit/texture/hipTexRefSetGetFilterMode.cc
+++ b/projects/hip-tests/catch/unit/texture/hipTexRefSetGetFilterMode.cc
@@ -18,6 +18,7 @@ THE SOFTWARE.
 */
 
 #include <hip_test_common.hh>
+#if defined(__HIP_PLATFORM_AMD__) || CUDA_VERSION < CUDA_12000
 texture<float, 2, hipReadModeElementType> tex;
 // Test for hipTexRefSetFilterMode and hipTexRefGetFilterMode, including error handling
 TEST_CASE("Unit_hipTexRefSetGetFilterMode") {
@@ -63,3 +64,4 @@ TEST_CASE("Unit_hipTexRefSetGetFilterMode") {
     REQUIRE(errGetMode == hipErrorInvalidValue);
   }
 }
+#endif
diff --git a/projects/hip-tests/catch/unit/texture/hipTexRefSetGetMipmapFilterMode.cc b/projects/hip-tests/catch/unit/texture/hipTexRefSetGetMipmapFilterMode.cc
index 73406b5639..7e208c136a 100644
--- a/projects/hip-tests/catch/unit/texture/hipTexRefSetGetMipmapFilterMode.cc
+++ b/projects/hip-tests/catch/unit/texture/hipTexRefSetGetMipmapFilterMode.cc
@@ -18,6 +18,7 @@ THE SOFTWARE.
 */
 
 #include <hip_test_common.hh>
+#if defined(__HIP_PLATFORM_AMD__) || CUDA_VERSION < CUDA_12000
 texture<float, 2, hipReadModeElementType> tex;
 // Test for hipTexRefSetMipmapFilterMode and hipTexRefGetMipmapFilterMode, including error handling
 TEST_CASE("Unit_hipTexRefSetGetMipmapFilterMode") {
@@ -56,3 +57,4 @@ TEST_CASE("Unit_hipTexRefSetGetMipmapFilterMode") {
     REQUIRE(err == hipErrorInvalidValue);
   }
 }
+#endif
diff --git a/projects/hip-tests/catch/unit/texture/hipTexRefSetGetMipmapLevelBias.cc b/projects/hip-tests/catch/unit/texture/hipTexRefSetGetMipmapLevelBias.cc
index e2e4298d80..8476826e99 100644
--- a/projects/hip-tests/catch/unit/texture/hipTexRefSetGetMipmapLevelBias.cc
+++ b/projects/hip-tests/catch/unit/texture/hipTexRefSetGetMipmapLevelBias.cc
@@ -18,6 +18,7 @@ THE SOFTWARE.
 */
 
 #include <hip_test_common.hh>
+#if defined(__HIP_PLATFORM_AMD__) || CUDA_VERSION < CUDA_12000
 texture<float, 2, hipReadModeElementType> tex;
 // Test for hipTexRefSetMipmapLevelBias and hipTexRefGetMipmapLevelBias, including error handling
 TEST_CASE("Unit_hipTexRefSetGetMipmapLevelBias") {
@@ -50,3 +51,4 @@ TEST_CASE("Unit_hipTexRefSetGetMipmapLevelBias") {
     REQUIRE(err == hipErrorInvalidValue);
   }
 }
+#endif
diff --git a/projects/hip-tests/catch/unit/texture/hipTexRefSetGetMipmapLevelClamp.cc b/projects/hip-tests/catch/unit/texture/hipTexRefSetGetMipmapLevelClamp.cc
index 764464a859..c28fbfe898 100644
--- a/projects/hip-tests/catch/unit/texture/hipTexRefSetGetMipmapLevelClamp.cc
+++ b/projects/hip-tests/catch/unit/texture/hipTexRefSetGetMipmapLevelClamp.cc
@@ -19,6 +19,7 @@ THE SOFTWARE.
 
 #include <hip_test_common.hh>
 #include <float.h>
+#if defined(__HIP_PLATFORM_AMD__) || CUDA_VERSION < CUDA_12000
 texture<float, 2, hipReadModeElementType> tex;
 // Test for hipTexRefSetMipmapLevelClamp and hipTexRefGetMipmapLevelClamp, including error handling
 TEST_CASE("Unit_texRefSetGetMipmapLevelClamp") {
@@ -55,3 +56,4 @@ TEST_CASE("Unit_texRefSetGetMipmapLevelClamp") {
     REQUIRE(err == hipErrorInvalidValue);
   }
 }
+#endif
diff --git a/projects/hip-tests/catch/unit/texture/hipTexRefSetGetMipmappedArray.cc b/projects/hip-tests/catch/unit/texture/hipTexRefSetGetMipmappedArray.cc
index d12bd31595..5f9208c71b 100644
--- a/projects/hip-tests/catch/unit/texture/hipTexRefSetGetMipmappedArray.cc
+++ b/projects/hip-tests/catch/unit/texture/hipTexRefSetGetMipmappedArray.cc
@@ -17,6 +17,7 @@ OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
 THE SOFTWARE.
 */
 #include <hip_test_common.hh>
+#if defined(__HIP_PLATFORM_AMD__) || CUDA_VERSION < CUDA_12000
 texture<float, 2, hipReadModeElementType> tex;
 
 // Test for hipTexRefSetMipmappedArray and hipTexRefGetMipmappedArray, including error handling
@@ -74,3 +75,4 @@ TEST_CASE("Unit_hipTexRefSetGetMipmappedArray") {
     REQUIRE(err == hipErrorInvalidValue);
   }
 }
+#endif
diff --git a/projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemAddressFree.cc b/projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemAddressFree.cc
index fee98051ca..b7c5733146 100644
--- a/projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemAddressFree.cc
+++ b/projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemAddressFree.cc
@@ -62,11 +62,11 @@ TEST_CASE("Unit_hipMemAddressFree_negative") {
   REQUIRE(granularity > 0);
   size_t size_mem = ((granularity + buffer_size - 1) / granularity) * granularity;
   // Allocate virtual address range
-  hipDeviceptr_t ptrA;
+  void* ptrA;
   HIP_CHECK(hipMemAddressReserve(&ptrA, size_mem, 0, 0, 0));
 
   SECTION("nullptr to devptr") {
-    REQUIRE(hipMemAddressFree((hipDeviceptr_t) nullptr, size_mem) == hipErrorInvalidValue);
+    REQUIRE(hipMemAddressFree(nullptr, size_mem) == hipErrorInvalidValue);
   }
 
   SECTION("pass zero to size") { REQUIRE(hipMemAddressFree(ptrA, 0) == hipErrorInvalidValue); }
diff --git a/projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemAddressReserve.cc b/projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemAddressReserve.cc
index df152ad5eb..495bf9b226 100644
--- a/projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemAddressReserve.cc
+++ b/projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemAddressReserve.cc
@@ -68,7 +68,7 @@ TEST_CASE("Unit_hipMemAddressReserve_AlignmentTest") {
   REQUIRE(granularity > 0);
   size_t size_mem = ((granularity + buffer_size - 1) / granularity) * granularity;
   // Allocate virtual address range
-  hipDeviceptr_t ptrA;
+  void* ptrA;
   size_t alignmnt = 1;
   hipMemGenericAllocationHandle_t handle;
   // Allocate physical memory
@@ -93,8 +93,8 @@ TEST_CASE("Unit_hipMemAddressReserve_AlignmentTest") {
     accessDesc.flags = hipMemAccessFlagsProtReadWrite;
     // Make the address accessible to GPU 0
     HIP_CHECK(hipMemSetAccess(ptrA, size_mem, &accessDesc, 1));
-    HIP_CHECK(hipMemcpyHtoD(ptrA, A_h.data(), buffer_size));
-    HIP_CHECK(hipMemcpyDtoH(B_h.data(), ptrA, buffer_size));
+    HIP_CHECK(hipMemcpyHtoD(reinterpret_cast<hipDeviceptr_t>(ptrA), A_h.data(), buffer_size));
+    HIP_CHECK(hipMemcpyDtoH(B_h.data(), reinterpret_cast<hipDeviceptr_t>(ptrA), buffer_size));
     REQUIRE(true == std::equal(B_h.begin(), B_h.end(), A_h.data()));
     HIP_CHECK(hipMemUnmap(ptrA, size_mem));
     HIP_CHECK(hipMemAddressFree(ptrA, size_mem));
@@ -131,7 +131,7 @@ TEST_CASE("Unit_hipMemAddressReserve_Negative") {
   REQUIRE(granularity > 0);
   size_t size_mem = ((granularity + buffer_size - 1) / granularity) * granularity;
   // Allocate virtual address range
-  hipDeviceptr_t ptrA;
+  void* ptrA;
 
   SECTION("Nullptr to ptr") {
     REQUIRE(hipMemAddressReserve(nullptr, size_mem, 0, 0, 0) == hipErrorInvalidValue);
diff --git a/projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemCreate.cc b/projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemCreate.cc
index b60a3100dd..89f141cc3f 100644
--- a/projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemCreate.cc
+++ b/projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemCreate.cc
@@ -138,7 +138,7 @@ TEST_CASE("Unit_hipMemCreate_ChkDev2HstMemcpy_ReleaseHdlPostUnmap") {
   // Allocate physical memory
   HIP_CHECK(hipMemCreate(&handle, size_mem, &prop, 0));
   // Allocate virtual address range
-  hipDeviceptr_t ptrA;
+  void* ptrA;
   HIP_CHECK(hipMemAddressReserve(&ptrA, size_mem, 0, 0, 0));
   HIP_CHECK(hipMemMap(ptrA, size_mem, 0, handle, 0));
   // Set access
@@ -154,8 +154,8 @@ TEST_CASE("Unit_hipMemCreate_ChkDev2HstMemcpy_ReleaseHdlPostUnmap") {
   for (size_t idx = 0; idx < N; idx++) {
     A_h[idx] = idx;
   }
-  HIP_CHECK(hipMemcpyHtoD(ptrA, A_h.data(), buffer_size));
-  HIP_CHECK(hipMemcpyDtoH(B_h.data(), ptrA, buffer_size));
+  HIP_CHECK(hipMemcpyHtoD(reinterpret_cast<hipDeviceptr_t>(ptrA), A_h.data(), buffer_size));
+  HIP_CHECK(hipMemcpyDtoH(B_h.data(), reinterpret_cast<hipDeviceptr_t>(ptrA), buffer_size));
   REQUIRE(true == std::equal(B_h.begin(), B_h.end(), A_h.data()));
   HIP_CHECK(hipMemUnmap(ptrA, size_mem));
   HIP_CHECK(hipMemAddressFree(ptrA, size_mem));
@@ -207,7 +207,7 @@ TEST_CASE("Unit_hipMemCreate_ChkDev2HstMemcpy_ReleaseHdlPreUse") {
   // Allocate physical memory
   HIP_CHECK(hipMemCreate(&handle, size_mem, &prop, 0));
   // Allocate virtual address range
-  hipDeviceptr_t ptrA;
+  void* ptrA;
   HIP_CHECK(hipMemAddressReserve(&ptrA, size_mem, 0, 0, 0));
   HIP_CHECK(hipMemMap(ptrA, size_mem, 0, handle, 0));
   HIP_CHECK(hipMemRelease(handle));
@@ -224,8 +224,8 @@ TEST_CASE("Unit_hipMemCreate_ChkDev2HstMemcpy_ReleaseHdlPreUse") {
   for (size_t idx = 0; idx < N; idx++) {
     A_h[idx] = idx;
   }
-  HIP_CHECK(hipMemcpyHtoD(ptrA, A_h.data(), buffer_size));
-  HIP_CHECK(hipMemcpyDtoH(B_h.data(), ptrA, buffer_size));
+  HIP_CHECK(hipMemcpyHtoD(reinterpret_cast<hipDeviceptr_t>(ptrA), A_h.data(), buffer_size));
+  HIP_CHECK(hipMemcpyDtoH(B_h.data(), reinterpret_cast<hipDeviceptr_t>(ptrA), buffer_size));
   REQUIRE(true == std::equal(B_h.begin(), B_h.end(), A_h.data()));
   HIP_CHECK(hipMemUnmap(ptrA, size_mem));
   HIP_CHECK(hipMemAddressFree(ptrA, size_mem));
@@ -276,7 +276,7 @@ TEST_CASE("Unit_hipMemCreate_ChkWithKerLaunch") {
   // Allocate physical memory
   HIP_CHECK(hipMemCreate(&handle, size_mem, &prop, 0));
   // Allocate virtual address range
-  hipDeviceptr_t ptrA;
+  void* ptrA;
   HIP_CHECK(hipMemAddressReserve(&ptrA, size_mem, 0, 0, 0));
   HIP_CHECK(hipMemMap(ptrA, size_mem, 0, handle, 0));
   HIP_CHECK(hipMemRelease(handle));
@@ -293,11 +293,11 @@ TEST_CASE("Unit_hipMemCreate_ChkWithKerLaunch") {
     A_h[idx] = idx;
     C_h[idx] = idx * idx;
   }
-  HIP_CHECK(hipMemcpyHtoD(ptrA, A_h.data(), buffer_size));
+  HIP_CHECK(hipMemcpyHtoD(reinterpret_cast<hipDeviceptr_t>(ptrA), A_h.data(), buffer_size));
   // Invoke kernel
   hipLaunchKernelGGL(square_kernel, dim3(N / THREADS_PER_BLOCK), dim3(THREADS_PER_BLOCK), 0, 0,
                      reinterpret_cast<int*>(ptrA));
-  HIP_CHECK(hipMemcpyDtoH(B_h.data(), ptrA, buffer_size));
+  HIP_CHECK(hipMemcpyDtoH(B_h.data(), reinterpret_cast<hipDeviceptr_t>(ptrA), buffer_size));
   HIP_CHECK(hipDeviceSynchronize());
   REQUIRE(true == std::equal(B_h.begin(), B_h.end(), C_h.data()));
   HIP_CHECK(hipMemUnmap(ptrA, size_mem));
@@ -352,12 +352,12 @@ TEST_CASE("Unit_hipMemCreate_MapNonContiguousChunks") {
     HIP_CHECK(hipMemCreate(&handle[count], size_mem, &prop, 0));
   }
   // Allocate virtual address range for all the memory chunks
-  hipDeviceptr_t ptrA;
+  void* ptrA;
   HIP_CHECK(hipMemAddressReserve(&ptrA, (numOfBuffers * size_mem), 0, 0, 0));
   for (int idx = 0; idx < numOfBuffers; idx++) {
     unsigned long long uiptr = reinterpret_cast<unsigned long long>(ptrA);
     uiptr = uiptr + idx * size_mem;
-    HIP_CHECK(hipMemMap(reinterpret_cast<hipDeviceptr_t>(uiptr), size_mem, 0, handle[idx], 0));
+    HIP_CHECK(hipMemMap(reinterpret_cast<void*>(uiptr), size_mem, 0, handle[idx], 0));
     HIP_CHECK(hipMemRelease(handle[idx]));
   }
   hipMemAccessDesc accessDesc = {};
@@ -373,18 +373,20 @@ TEST_CASE("Unit_hipMemCreate_MapNonContiguousChunks") {
     A_h[idx] = idx;
     C_h[idx] = idx * idx;
   }
-  HIP_CHECK(hipMemcpyHtoD(ptrA, A_h.data(), numOfBuffers * buffer_size));
+  HIP_CHECK(hipMemcpyHtoD(reinterpret_cast<hipDeviceptr_t>(ptrA), A_h.data(),
+                          numOfBuffers * buffer_size));
   // Launch square kernel
   hipLaunchKernelGGL(square_kernel, dim3((N * numOfBuffers) / THREADS_PER_BLOCK),
                      dim3(THREADS_PER_BLOCK), 0, 0, reinterpret_cast<int*>(ptrA));
-  HIP_CHECK(hipMemcpyDtoH(B_h.data(), ptrA, numOfBuffers * buffer_size));
+  HIP_CHECK(hipMemcpyDtoH(B_h.data(), reinterpret_cast<hipDeviceptr_t>(ptrA),
+                         numOfBuffers * buffer_size));
   HIP_CHECK(hipDeviceSynchronize());
   // Validate Results
   REQUIRE(true == std::equal(B_h.begin(), B_h.end(), C_h.data()));
   for (int idx = 0; idx < numOfBuffers; idx++) {
     unsigned long long uiptr = reinterpret_cast<unsigned long long>(ptrA);
     uiptr = uiptr + idx * size_mem;
-    HIP_CHECK(hipMemUnmap(reinterpret_cast<hipDeviceptr_t>(uiptr), size_mem));
+    HIP_CHECK(hipMemUnmap(reinterpret_cast<void*>(uiptr), size_mem));
   }
   HIP_CHECK(hipMemAddressFree(ptrA, (numOfBuffers * size_mem)));
   CTX_DESTROY();
@@ -434,7 +436,7 @@ TEST_CASE("Unit_hipMemCreate_ChkWithMemset") {
   // Allocate physical memory
   HIP_CHECK(hipMemCreate(&handle, size_mem, &prop, 0));
   // Allocate virtual address range
-  hipDeviceptr_t ptrA;
+  void* ptrA;
   HIP_CHECK(hipMemAddressReserve(&ptrA, size_mem, 0, 0, 0));
   HIP_CHECK(hipMemMap(ptrA, size_mem, 0, handle, 0));
   // Set access
@@ -445,8 +447,8 @@ TEST_CASE("Unit_hipMemCreate_ChkWithMemset") {
   // Make the address accessible to GPU 0
   HIP_CHECK(hipMemSetAccess(ptrA, size_mem, &accessDesc, 1));
   std::vector<int> A_h(N);
-  HIP_CHECK(hipMemset(reinterpret_cast<void*>(ptrA), init_val, buffer_size));
-  HIP_CHECK(hipMemcpyDtoH(A_h.data(), ptrA, buffer_size));
+  HIP_CHECK(hipMemset(ptrA, init_val, buffer_size));
+  HIP_CHECK(hipMemcpyDtoH(A_h.data(), reinterpret_cast<hipDeviceptr_t>(ptrA), buffer_size));
   for (int idx = 0; idx < N; idx++) {
     REQUIRE(A_h[idx] == init_val);
   }
diff --git a/projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemMap.cc b/projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemMap.cc
index 7aa80d1e66..92b7f893cd 100644
--- a/projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemMap.cc
+++ b/projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemMap.cc
@@ -99,7 +99,7 @@ TEST_CASE("Unit_hipMemMap_SameMemoryReuse") {
   // Allocate a physical memory chunk
   HIP_CHECK(hipMemCreate(&handle, size_mem, &prop, 0));
   // Allocate num_buf virtual address ranges
-  hipDeviceptr_t ptrA;
+  void* ptrA;
   HIP_CHECK(hipMemAddressReserve(&ptrA, size_mem, 0, 0, 0));
   hipMemAccessDesc accessDesc = {};
   accessDesc.location.type = hipMemLocationTypeDevice;
@@ -110,12 +110,12 @@ TEST_CASE("Unit_hipMemMap_SameMemoryReuse") {
     HIP_CHECK(hipMemMap(ptrA, size_mem, 0, handle, 0));
     // Set access to GPU 0
     HIP_CHECK(hipMemSetAccess(ptrA, size_mem, &accessDesc, 1));
-    HIP_CHECK(hipMemcpyHtoD(ptrA, A_h.data(), buffer_size));
-    HIP_CHECK(hipMemcpyDtoH(B_h.data(), ptrA, buffer_size));
+    HIP_CHECK(hipMemcpyHtoD(reinterpret_cast<hipDeviceptr_t>(ptrA), A_h.data(), buffer_size));
+    HIP_CHECK(hipMemcpyDtoH(B_h.data(), reinterpret_cast<hipDeviceptr_t>(ptrA), buffer_size));
     REQUIRE(true == std::equal(B_h.begin(), B_h.end(), A_h.data()));
     square_kernel<<<dim3(N / threadsPerBlk), dim3(threadsPerBlk), 0, 0>>>(
         reinterpret_cast<int*>(ptrA));
-    HIP_CHECK(hipMemcpyDtoH(B_h.data(), ptrA, buffer_size));
+    HIP_CHECK(hipMemcpyDtoH(B_h.data(), reinterpret_cast<hipDeviceptr_t>(ptrA), buffer_size));
     HIP_CHECK(hipStreamSynchronize(0));
     REQUIRE(true == std::equal(B_h.begin(), B_h.end(), C_h.data()));
     HIP_CHECK(hipMemUnmap(ptrA, size_mem));
@@ -175,7 +175,7 @@ TEST_CASE("Unit_hipMemMap_PhysicalMemoryReuse_SingleGPU") {
   // Allocate a physical memory chunk
   HIP_CHECK(hipMemCreate(&handle, size_mem, &prop, 0));
   // Allocate num_buf virtual address ranges
-  hipDeviceptr_t ptrA[num_buf];
+  void* ptrA[num_buf];
   for (int buf = 0; buf < num_buf; buf++) {
     HIP_CHECK(hipMemAddressReserve(&ptrA[buf], size_mem, 0, 0, 0));
   }
@@ -188,12 +188,12 @@ TEST_CASE("Unit_hipMemMap_PhysicalMemoryReuse_SingleGPU") {
     HIP_CHECK(hipMemMap(ptrA[buf], size_mem, 0, handle, 0));
     // Set access to GPU 0
     HIP_CHECK(hipMemSetAccess(ptrA[buf], size_mem, &accessDesc, 1));
-    HIP_CHECK(hipMemcpyHtoD(ptrA[buf], A_h.data(), buffer_size));
-    HIP_CHECK(hipMemcpyDtoH(B_h.data(), ptrA[buf], buffer_size));
+    HIP_CHECK(hipMemcpyHtoD(reinterpret_cast<hipDeviceptr_t>(ptrA[buf]), A_h.data(), buffer_size));
+    HIP_CHECK(hipMemcpyDtoH(B_h.data(), reinterpret_cast<hipDeviceptr_t>(ptrA[buf]), buffer_size));
     REQUIRE(true == std::equal(B_h.begin(), B_h.end(), A_h.data()));
     square_kernel<<<dim3(N / threadsPerBlk), dim3(threadsPerBlk), 0, 0>>>(
         reinterpret_cast<int*>(ptrA[buf]));
-    HIP_CHECK(hipMemcpyDtoH(B_h.data(), ptrA[buf], buffer_size));
+    HIP_CHECK(hipMemcpyDtoH(B_h.data(), reinterpret_cast<hipDeviceptr_t>(ptrA[buf]), buffer_size));
     HIP_CHECK(hipStreamSynchronize(0));
     REQUIRE(true == std::equal(B_h.begin(), B_h.end(), C_h.data()));
     HIP_CHECK(hipMemUnmap(ptrA[buf], size_mem));
@@ -255,7 +255,7 @@ TEST_CASE("Unit_hipMemMap_PhysicalMemory_Map2MultVMMs") {
   // Allocate a physical memory chunk
   HIP_CHECK(hipMemCreate(&handle, size_mem, &prop, 0));
   // Allocate num_buf virtual address ranges
-  hipDeviceptr_t ptrA[num_buf];
+  void* ptrA[num_buf];
   for (int buf = 0; buf < num_buf; buf++) {
     HIP_CHECK(hipMemAddressReserve(&ptrA[buf], size_mem, 0, 0, 0));
   }
@@ -271,12 +271,12 @@ TEST_CASE("Unit_hipMemMap_PhysicalMemory_Map2MultVMMs") {
     HIP_CHECK(hipMemSetAccess(ptrA[buf], size_mem, &accessDesc, 1));
   }
   // Copy data to VMM via ptrA[0]
-  HIP_CHECK(hipMemcpyHtoD(ptrA[0], A_h.data(), buffer_size));
+  HIP_CHECK(hipMemcpyHtoD(reinterpret_cast<hipDeviceptr_t>(ptrA[0]), A_h.data(), buffer_size));
   // Validate the data contained in VMM using ptrA[0], ptrA[1],
   // ......, ptrA[num_buf-1]
   for (int buf = 0; buf < num_buf; buf++) {
     std::fill(B_h.begin(), B_h.end(), initializer);
-    HIP_CHECK(hipMemcpyDtoH(B_h.data(), ptrA[buf], buffer_size));
+    HIP_CHECK(hipMemcpyDtoH(B_h.data(), reinterpret_cast<hipDeviceptr_t>(ptrA[buf]), buffer_size));
     REQUIRE(true == std::equal(B_h.begin(), B_h.end(), A_h.data()));
   }
 
@@ -321,7 +321,7 @@ void physicalMemoryReuse_MultiDev (hipMemAllocationProp prop) {
     // Allocate a physical memory chunk
     HIP_CHECK(hipMemCreate(&handle, size_mem, &prop, 0));
     // Allocate devicecount virtual address ranges
-    std::vector<hipDeviceptr_t> ptrA(devicecount);
+    std::vector<void*> ptrA(devicecount);
     for (int devY = 0; devY < devicecount; devY++) {
       HIP_CHECK(hipMemAddressReserve(&ptrA[devY], size_mem, 0, 0, 0));
     }
@@ -337,8 +337,10 @@ void physicalMemoryReuse_MultiDev (hipMemAllocationProp prop) {
       HIP_CHECK(hipMemMap(ptrA[devY], size_mem, 0, handle, 0));
       // Set access to GPU 0
       HIP_CHECK(hipMemSetAccess(ptrA[devY], size_mem, &accessDesc, 1));
-      HIP_CHECK(hipMemcpyHtoD(ptrA[devY], A_h.data(), buffer_size));
-      HIP_CHECK(hipMemcpyDtoH(B_h.data(), ptrA[devY], buffer_size));
+      HIP_CHECK(
+          hipMemcpyHtoD(reinterpret_cast<hipDeviceptr_t>(ptrA[devY]), A_h.data(), buffer_size));
+      HIP_CHECK(
+          hipMemcpyDtoH(B_h.data(), reinterpret_cast<hipDeviceptr_t>(ptrA[devY]), buffer_size));
       REQUIRE(true == std::equal(B_h.begin(), B_h.end(), A_h.data()));
       HIP_CHECK(hipMemUnmap(ptrA[devY], size_mem));
     }
@@ -431,7 +433,7 @@ TEST_CASE("Unit_hipMemMap_VMMMemoryReuse_SingleGPU") {
     HIP_CHECK(hipMemCreate(&handle[buf], size_mem, &prop, 0));
   }
   // Allocate num_buf virtual address ranges
-  hipDeviceptr_t ptrA;
+  void* ptrA;
   HIP_CHECK(hipMemAddressReserve(&ptrA, size_mem, 0, 0, 0));
   hipMemAccessDesc accessDesc = {};
   accessDesc.location.type = hipMemLocationTypeDevice;
@@ -443,13 +445,13 @@ TEST_CASE("Unit_hipMemMap_VMMMemoryReuse_SingleGPU") {
     HIP_CHECK(hipMemMap(ptrA, size_mem, 0, handle[buf], 0));
     // Set access to GPU 0
     HIP_CHECK(hipMemSetAccess(ptrA, size_mem, &accessDesc, 1));
-    HIP_CHECK(hipMemcpyHtoD(ptrA, A_h.data(), buffer_size));
-    HIP_CHECK(hipMemcpyDtoH(B_h.data(), ptrA, buffer_size));
+    HIP_CHECK(hipMemcpyHtoD(reinterpret_cast<hipDeviceptr_t>(ptrA), A_h.data(), buffer_size));
+    HIP_CHECK(hipMemcpyDtoH(B_h.data(), reinterpret_cast<hipDeviceptr_t>(ptrA), buffer_size));
     REQUIRE(true == std::equal(B_h.begin(), B_h.end(), A_h.data()));
 #if HT_NVIDIA
     square_kernel<<<dim3(N / threadsPerBlk), dim3(threadsPerBlk), 0, 0>>>(
         reinterpret_cast<int*>(ptrA));
-    HIP_CHECK(hipMemcpyDtoH(B_h.data(), ptrA, buffer_size));
+    HIP_CHECK(hipMemcpyDtoH(B_h.data(), reinterpret_cast<hipDeviceptr_t>(ptrA), buffer_size));
     HIP_CHECK(hipStreamSynchronize(0));
     REQUIRE(true == std::equal(B_h.begin(), B_h.end(), C_h.data()));
 #endif
@@ -497,7 +499,7 @@ void vMMMemoryReuse_MultiGPU (hipMemAllocationProp prop) {
     HIP_CHECK(hipMemCreate(&handle[dev], size_mem, &prop, 0));
   }
   // Allocate devicecount virtual address ranges
-  hipDeviceptr_t ptrA;
+  void* ptrA;
   HIP_CHECK(hipMemAddressReserve(&ptrA, size_mem, 0, 0, 0));
   // Map ptrA to physical chunk
   SECTION("Set Access of VMM to Different GPU") {
@@ -512,8 +514,8 @@ void vMMMemoryReuse_MultiGPU (hipMemAllocationProp prop) {
       std::fill(B_h.begin(), B_h.end(), initializer);
       HIP_CHECK(hipMemMap(ptrA, size_mem, 0, handle[dev], 0));
       HIP_CHECK(hipMemSetAccess(ptrA, size_mem, &accessDesc, 1));
-      HIP_CHECK(hipMemcpyHtoD(ptrA, A_h.data(), buffer_size));
-      HIP_CHECK(hipMemcpyDtoH(B_h.data(), ptrA, buffer_size));
+      HIP_CHECK(hipMemcpyHtoD(reinterpret_cast<hipDeviceptr_t>(ptrA), A_h.data(), buffer_size));
+      HIP_CHECK(hipMemcpyDtoH(B_h.data(), reinterpret_cast<hipDeviceptr_t>(ptrA), buffer_size));
       HIP_CHECK(hipMemUnmap(ptrA, size_mem));
       REQUIRE(true == std::equal(B_h.begin(), B_h.end(), A_h.data()));
     }
@@ -527,8 +529,8 @@ void vMMMemoryReuse_MultiGPU (hipMemAllocationProp prop) {
       std::fill(B_h.begin(), B_h.end(), initializer);
       HIP_CHECK(hipMemMap(ptrA, size_mem, 0, handle[dev], 0));
       HIP_CHECK(hipMemSetAccess(ptrA, size_mem, &accessDesc, 1));
-      HIP_CHECK(hipMemcpyHtoD(ptrA, A_h.data(), buffer_size));
-      HIP_CHECK(hipMemcpyDtoH(B_h.data(), ptrA, buffer_size));
+      HIP_CHECK(hipMemcpyHtoD(reinterpret_cast<hipDeviceptr_t>(ptrA), A_h.data(), buffer_size));
+      HIP_CHECK(hipMemcpyDtoH(B_h.data(), reinterpret_cast<hipDeviceptr_t>(ptrA), buffer_size));
       HIP_CHECK(hipMemUnmap(ptrA, size_mem));
       REQUIRE(true == std::equal(B_h.begin(), B_h.end(), A_h.data()));
     }
@@ -617,7 +619,7 @@ TEST_CASE("Unit_hipMemMap_MapPartialVMMMem") {
   // Allocate a bigger physical memory chunk of size_mem
   HIP_CHECK(hipMemCreate(&handle, size_mem, &prop, 0));
   // Allocate virtual address range of size twice size_mem
-  hipDeviceptr_t ptrA;
+  void* ptrA;
   HIP_CHECK(hipMemAddressReserve(&ptrA, 2 * size_mem, 0, 0, 0));
   hipMemAccessDesc accessDesc = {};
   accessDesc.location.type = hipMemLocationTypeDevice;
@@ -626,8 +628,8 @@ TEST_CASE("Unit_hipMemMap_MapPartialVMMMem") {
   std::fill(B_h.begin(), B_h.end(), initializer);
   HIP_CHECK(hipMemMap(ptrA, size_mem, 0, handle, 0));
   HIP_CHECK(hipMemSetAccess(ptrA, size_mem, &accessDesc, 1));
-  HIP_CHECK(hipMemcpyHtoD(ptrA, A_h.data(), buffer_size));
-  HIP_CHECK(hipMemcpyDtoH(B_h.data(), ptrA, buffer_size));
+  HIP_CHECK(hipMemcpyHtoD(reinterpret_cast<hipDeviceptr_t>(ptrA), A_h.data(), buffer_size));
+  HIP_CHECK(hipMemcpyDtoH(B_h.data(), reinterpret_cast<hipDeviceptr_t>(ptrA), buffer_size));
   REQUIRE(true == std::equal(B_h.begin(), B_h.end(), A_h.data()));
   HIP_CHECK(hipMemUnmap(ptrA, size_mem));
   // Release resources
@@ -663,14 +665,14 @@ TEST_CASE("Unit_hipMemMap_negative") {
   REQUIRE(granularity > 0);
   size_t size_mem = ((granularity + buffer_size - 1) / granularity) * granularity;
   hipMemGenericAllocationHandle_t handle;
-  hipDeviceptr_t ptrA;
+  void* ptrA;
   // Allocate physical memory
   HIP_CHECK(hipMemCreate(&handle, size_mem, &prop, 0));
   // Allocate virtual address range
   HIP_CHECK(hipMemAddressReserve(&ptrA, size_mem, 0, 0, 0));
 
   SECTION("nullptr to ptrA") {
-    REQUIRE(hipMemMap((hipDeviceptr_t) nullptr, size_mem, 0, handle, 0) == hipErrorInvalidValue);
+    REQUIRE(hipMemMap(nullptr, size_mem, 0, handle, 0) == hipErrorInvalidValue);
   }
 
   SECTION("pass zero to size") {
diff --git a/projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemRetainAllocationHandle.cc b/projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemRetainAllocationHandle.cc
index 28dee8b1b6..ba95ea2f51 100644
--- a/projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemRetainAllocationHandle.cc
+++ b/projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemRetainAllocationHandle.cc
@@ -65,12 +65,12 @@ TEST_CASE("Unit_hipMemRetainAllocationHandle_SetGet") {
   REQUIRE(granularity > 0);
   size_t size_mem = ((granularity + buffer_size - 1) / granularity) * granularity;
   hipMemGenericAllocationHandle_t handle;
-  hipDeviceptr_t ptrA;
+  void* ptrA;
   // Allocate physical memory
   HIP_CHECK(hipMemCreate(&handle, size_mem, &prop, 0));
   // Allocate virtual address range
-  HIP_CHECK(hipMemAddressReserve(&ptrA, size_mem, 0, 0, 0));
-  HIP_CHECK(hipMemMap(ptrA, size_mem, 0, handle, 0));
+  HIP_CHECK(hipMemAddressReserve(reinterpret_cast<void**>(&ptrA), size_mem, 0, 0, 0));
+  HIP_CHECK(hipMemMap((void*)ptrA, size_mem, 0, handle, 0));
   // Test hipMemRetainAllocationHandle
   hipMemGenericAllocationHandle_t gethandle;
   // Check beginning of VMM ptr
@@ -109,7 +109,7 @@ TEST_CASE("Unit_hipMemRetainAllocationHandle_NegTst") {
   REQUIRE(granularity > 0);
   size_t size_mem = ((granularity + buffer_size - 1) / granularity) * granularity;
   hipMemGenericAllocationHandle_t handle;
-  hipDeviceptr_t ptrA;
+  void* ptrA;
   // Allocate physical memory
   HIP_CHECK(hipMemCreate(&handle, size_mem, &prop, 0));
   // Allocate virtual address range
@@ -125,7 +125,7 @@ TEST_CASE("Unit_hipMemRetainAllocationHandle_NegTst") {
     REQUIRE(hipMemRetainAllocationHandle(&gethandle, nullptr) == hipErrorInvalidValue);
   }
   SECTION("not mapped address") {
-    hipDeviceptr_t ptrB;
+    void* ptrB;
     HIP_CHECK(hipMemAddressReserve(&ptrB, size_mem, 0, 0, 0));
     REQUIRE(hipMemRetainAllocationHandle(&gethandle, reinterpret_cast<void*>(ptrB)) ==
             hipErrorInvalidValue);
diff --git a/projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemSetGetAccess.cc b/projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemSetGetAccess.cc
index 7829de69e9..1b83ed3fb7 100644
--- a/projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemSetGetAccess.cc
+++ b/projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemSetGetAccess.cc
@@ -100,7 +100,7 @@ TEST_CASE("Unit_hipMemSetAccess_SetGet") {
   // Allocate physical memory
   HIP_CHECK(hipMemCreate(&handle, size_mem, &prop, 0));
   // Allocate virtual address range
-  hipDeviceptr_t ptrA;
+  void* ptrA;
   HIP_CHECK(hipMemAddressReserve(&ptrA, size_mem, 0, 0, 0));
   HIP_CHECK(hipMemMap(ptrA, size_mem, 0, handle, 0));
   HIP_CHECK(hipMemRelease(handle));
@@ -176,7 +176,7 @@ TEST_CASE("Unit_hipMemSetAccess_MultDevSetGet") {
   // Allocate physical memory
   HIP_CHECK(hipMemCreate(&handle, size_mem, &prop, 0));
   // Allocate virtual address range
-  hipDeviceptr_t ptrA;
+  void* ptrA;
   HIP_CHECK(hipMemAddressReserve(&ptrA, size_mem, 0, 0, 0));
   HIP_CHECK(hipMemMap(ptrA, size_mem, 0, handle, 0));
   HIP_CHECK(hipMemRelease(handle));
@@ -237,7 +237,7 @@ TEST_CASE("Unit_hipMemSetAccess_EntireVMMRangeSetGet") {
   size_t size_mem = ((granularity + buffer_size - 1) / granularity) * granularity;
   // Allocate physical memory
   hipMemGenericAllocationHandle_t handle;
-  hipDeviceptr_t ptrA;
+  void* ptrA;
   HIP_CHECK(hipMemCreate(&handle, size_mem, &prop, 0));
   // Allocate virtual address range
   HIP_CHECK(hipMemAddressReserve(&ptrA, size_mem, 0, 0, 0));
@@ -259,7 +259,7 @@ TEST_CASE("Unit_hipMemSetAccess_EntireVMMRangeSetGet") {
   REQUIRE(flags == hipMemAccessFlagsProtReadWrite);
   unsigned long long uiptr = reinterpret_cast<unsigned long long>(ptrA);
   uiptr += (size_mem - 1);
-  HIP_CHECK(hipMemGetAccess(&flags, &location, reinterpret_cast<hipDeviceptr_t>(uiptr)));
+  HIP_CHECK(hipMemGetAccess(&flags, &location, reinterpret_cast<void*>(uiptr)));
   REQUIRE(flags == hipMemAccessFlagsProtReadWrite);
   HIP_CHECK(hipMemUnmap(ptrA, size_mem));
   HIP_CHECK(hipMemAddressFree(ptrA, size_mem));
@@ -295,7 +295,7 @@ TEST_CASE("Unit_hipMemGetAccess_NegTst") {
   size_t size_mem = ((granularity + buffer_size - 1) / granularity) * granularity;
   // Allocate physical memory
   hipMemGenericAllocationHandle_t handle;
-  hipDeviceptr_t ptrA;
+  void* ptrA;
   HIP_CHECK(hipMemCreate(&handle, size_mem, &prop, 0));
   // Allocate virtual address range
   HIP_CHECK(hipMemAddressReserve(&ptrA, size_mem, 0, 0, 0));
@@ -320,7 +320,7 @@ TEST_CASE("Unit_hipMemGetAccess_NegTst") {
   REQUIRE(status == hipErrorInvalidValue);
   unsigned long long uiptr = reinterpret_cast<unsigned long long>(ptrA);
   uiptr += size_mem;
-  status = hipMemGetAccess(&flags, &location, reinterpret_cast<hipDeviceptr_t>(uiptr));
+  status = hipMemGetAccess(&flags, &location, reinterpret_cast<void*>(uiptr));
   REQUIRE(status == hipErrorInvalidValue);
   HIP_CHECK(hipMemUnmap(ptrA, size_mem));
   HIP_CHECK(hipMemAddressFree(ptrA, size_mem));
@@ -363,7 +363,7 @@ TEST_CASE("Unit_hipMemSetAccess_FuncTstOnMultDev") {
     REQUIRE(granularity > 0);
     size_t size_mem = ((granularity + buffer_size - 1) / granularity) * granularity;
     // Allocate physical memory
-    hipDeviceptr_t ptrA;
+    void* ptrA;
     hipMemGenericAllocationHandle_t handle;
     HIP_CHECK(hipMemCreate(&handle, size_mem, &prop, 0));
     // Allocate virtual address range
@@ -382,7 +382,7 @@ TEST_CASE("Unit_hipMemSetAccess_FuncTstOnMultDev") {
     for (int idx = 0; idx < N; idx++) {
       A_h[idx] = idx;
     }
-    HIP_CHECK(hipMemcpyHtoD(ptrA, A_h.data(), buffer_size));
+    HIP_CHECK(hipMemcpyHtoD(reinterpret_cast<hipDeviceptr_t>(ptrA), A_h.data(), buffer_size));
     // Set the A_h to verify with square kernel.
     for (int idx = 0; idx < N; idx++) {
       A_h[idx] = idx * idx;
@@ -390,7 +390,7 @@ TEST_CASE("Unit_hipMemSetAccess_FuncTstOnMultDev") {
     // Launch square kernel
     hipLaunchKernelGGL(square_kernel, dim3(N / THREADS_PER_BLOCK), dim3(THREADS_PER_BLOCK), 0, 0,
                        reinterpret_cast<int*>(ptrA));
-    HIP_CHECK(hipMemcpyDtoH(B_h.data(), ptrA, buffer_size));
+    HIP_CHECK(hipMemcpyDtoH(B_h.data(), reinterpret_cast<hipDeviceptr_t>(ptrA), buffer_size));
     HIP_CHECK(hipDeviceSynchronize());
     REQUIRE(true == std::equal(B_h.begin(), B_h.end(), A_h.data()));
     HIP_CHECK(hipMemUnmap(ptrA, size_mem));
@@ -438,7 +438,7 @@ TEST_CASE("Unit_hipMemSetAccess_ChangeAccessProp") {
   // Allocate a physical memory chunk
   HIP_CHECK(hipMemCreate(&handle, size_mem, &prop, 0));
   // Allocate num_buf virtual address ranges
-  hipDeviceptr_t ptrA;
+  void* ptrA;
   HIP_CHECK(hipMemAddressReserve(&ptrA, size_mem, 0, 0, 0));
   HIP_CHECK(hipMemMap(ptrA, size_mem, 0, handle, 0));
   hipMemAccessDesc accessDesc = {};
@@ -451,8 +451,8 @@ TEST_CASE("Unit_hipMemSetAccess_ChangeAccessProp") {
     // Change property of virtual memory range to read only
     accessDesc.flags = hipMemAccessFlagsProtReadWrite;
     HIP_CHECK(hipMemSetAccess(ptrA, size_mem, &accessDesc, 1));
-    HIP_CHECK(hipMemcpyHtoD(ptrA, A_h.data(), buffer_size));
-    HIP_CHECK(hipMemcpyDtoH(B_h.data(), ptrA, buffer_size));
+    HIP_CHECK(hipMemcpyHtoD(reinterpret_cast<hipDeviceptr_t>(ptrA), A_h.data(), buffer_size));
+    HIP_CHECK(hipMemcpyDtoH(B_h.data(), reinterpret_cast<hipDeviceptr_t>(ptrA), buffer_size));
     REQUIRE(true == std::equal(B_h.begin(), B_h.end(), A_h.data()));
   }
 
@@ -462,8 +462,8 @@ TEST_CASE("Unit_hipMemSetAccess_ChangeAccessProp") {
     // Change property of virtual memory range to read only
     accessDesc.flags = hipMemAccessFlagsProtReadWrite;
     HIP_CHECK(hipMemSetAccess(ptrA, size_mem, &accessDesc, 1));
-    HIP_CHECK(hipMemcpyHtoD(ptrA, A_h.data(), buffer_size));
-    HIP_CHECK(hipMemcpyDtoH(B_h.data(), ptrA, buffer_size));
+    HIP_CHECK(hipMemcpyHtoD(reinterpret_cast<hipDeviceptr_t>(ptrA), A_h.data(), buffer_size));
+    HIP_CHECK(hipMemcpyDtoH(B_h.data(), reinterpret_cast<hipDeviceptr_t>(ptrA), buffer_size));
     REQUIRE(true == std::equal(B_h.begin(), B_h.end(), A_h.data()));
   }
 
@@ -471,13 +471,15 @@ TEST_CASE("Unit_hipMemSetAccess_ChangeAccessProp") {
   SECTION("Check error while writing on Read-Only memory") {
     accessDesc.flags = hipMemAccessFlagsProtRead;
     HIP_CHECK(hipMemSetAccess(ptrA, size_mem, &accessDesc, 1));
-    REQUIRE(hipErrorInvalidValue == hipMemcpyHtoD(ptrA, A_h.data(), buffer_size));
+    REQUIRE(hipErrorInvalidValue ==
+            hipMemcpyHtoD(reinterpret_cast<hipDeviceptr_t>(ptrA), A_h.data(), buffer_size));
   }
 
   SECTION("Check error while writing on inaccessible memory") {
     accessDesc.flags = hipMemAccessFlagsProtNone;
     HIP_CHECK(hipMemSetAccess(ptrA, size_mem, &accessDesc, 1));
-    REQUIRE(hipErrorInvalidValue == hipMemcpyHtoD(ptrA, A_h.data(), buffer_size));
+    REQUIRE(hipErrorInvalidValue ==
+            hipMemcpyHtoD(reinterpret_cast<hipDeviceptr_t>(ptrA), A_h.data(), buffer_size));
   }
 #endif
 
@@ -524,7 +526,7 @@ TEST_CASE("Unit_hipMemSetAccess_Vmm2UnifiedMemCpy") {
   size_t size_mem = ((granularity + buffer_size - 1) / granularity) * granularity;
   // Allocate physical memory
   hipMemGenericAllocationHandle_t handle;
-  hipDeviceptr_t ptrA, ptrB;
+  void *ptrA, *ptrB;
   HIP_CHECK(hipMemCreate(&handle, size_mem, &prop, 0));
   // Allocate virtual address range
   HIP_CHECK(hipMemAddressReserve(&ptrA, size_mem, 0, 0, 0));
@@ -543,10 +545,11 @@ TEST_CASE("Unit_hipMemSetAccess_Vmm2UnifiedMemCpy") {
   for (int idx = 0; idx < N; idx++) {
     ptrA_h[idx] = idx;
   }
-  HIP_CHECK(hipMemcpyHtoD(ptrA, ptrA_h, buffer_size));
+  HIP_CHECK(hipMemcpyHtoD(reinterpret_cast<hipDeviceptr_t>(ptrA), ptrA_h, buffer_size));
   HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&ptrB), buffer_size));
-  HIP_CHECK(hipMemcpyDtoD(ptrB, ptrA, buffer_size));
-  HIP_CHECK(hipMemcpyDtoH(ptrB_h, ptrB, buffer_size));
+  HIP_CHECK(hipMemcpyDtoD(reinterpret_cast<hipDeviceptr_t>(ptrB),
+                         reinterpret_cast<hipDeviceptr_t>(ptrA), buffer_size));
+  HIP_CHECK(hipMemcpyDtoH(ptrB_h, reinterpret_cast<hipDeviceptr_t>(ptrB), buffer_size));
   bool bPassed = true;
   for (int idx = 0; idx < N; idx++) {
     if (ptrB_h[idx] != idx) {
@@ -594,7 +597,7 @@ TEST_CASE("Unit_hipMemSetAccess_Vmm2DevMemCpy") {
   size_t size_mem = ((granularity + buffer_size - 1) / granularity) * granularity;
   // Allocate physical memory
   hipMemGenericAllocationHandle_t handle;
-  hipDeviceptr_t ptrA, ptrB;
+  void *ptrA, *ptrB;
   HIP_CHECK(hipMemCreate(&handle, size_mem, &prop, 0));
   // Allocate virtual address range
   HIP_CHECK(hipMemAddressReserve(&ptrA, size_mem, 0, 0, 0));
@@ -611,10 +614,11 @@ TEST_CASE("Unit_hipMemSetAccess_Vmm2DevMemCpy") {
   for (int idx = 0; idx < N; idx++) {
     A_h[idx] = idx;
   }
-  HIP_CHECK(hipMemcpyHtoD(ptrA, A_h.data(), buffer_size));
+  HIP_CHECK(hipMemcpyHtoD(reinterpret_cast<hipDeviceptr_t>(ptrA), A_h.data(), buffer_size));
   HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&ptrB), buffer_size));
-  HIP_CHECK(hipMemcpyDtoD(ptrB, ptrA, buffer_size));
-  HIP_CHECK(hipMemcpyDtoH(B_h.data(), ptrB, buffer_size));
+  HIP_CHECK(hipMemcpyDtoD(reinterpret_cast<hipDeviceptr_t>(ptrB),
+                          reinterpret_cast<hipDeviceptr_t>(ptrA), buffer_size));
+  HIP_CHECK(hipMemcpyDtoH(B_h.data(), reinterpret_cast<hipDeviceptr_t>(ptrB), buffer_size));
   REQUIRE(true == std::equal(B_h.begin(), B_h.end(), A_h.data()));
   HIP_CHECK(hipFree(reinterpret_cast<void*>(ptrB)));
   HIP_CHECK(hipMemUnmap(ptrA, size_mem));
@@ -659,7 +663,7 @@ TEST_CASE("Unit_hipMemSetAccess_Vmm2PeerDevMemCpy") {
   size_t size_mem = ((granularity + buffer_size - 1) / granularity) * granularity;
   // Allocate physical memory
   hipMemGenericAllocationHandle_t handle;
-  hipDeviceptr_t ptrA;
+  void* ptrA;
   HIP_CHECK(hipMemCreate(&handle, size_mem, &prop, 0));
   // Allocate virtual address range
   HIP_CHECK(hipMemAddressReserve(&ptrA, size_mem, 0, 0, 0));
@@ -676,7 +680,7 @@ TEST_CASE("Unit_hipMemSetAccess_Vmm2PeerDevMemCpy") {
   for (int idx = 0; idx < N; idx++) {
     A_h[idx] = idx;
   }
-  HIP_CHECK(hipMemcpyHtoD(ptrA, A_h.data(), buffer_size));
+  HIP_CHECK(hipMemcpyHtoD(reinterpret_cast<hipDeviceptr_t>(ptrA), A_h.data(), buffer_size));
   // Check Peer Access
   for (deviceId = 1; deviceId < devicecount; deviceId++) {
     int canAccessPeer = 0;
@@ -703,7 +707,7 @@ TEST_CASE("Unit_hipMemSetAccess_Vmm2PeerDevMemCpy") {
     HIP_CHECK(hipMemSetAccess(ptrA, size_mem, &access, 1));
     hipDeviceptr_t dptr_peer;
     HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&dptr_peer), buffer_size));
-    HIP_CHECK(hipMemcpyDtoD(dptr_peer, ptrA, buffer_size));
+    HIP_CHECK(hipMemcpyDtoD(dptr_peer, reinterpret_cast<hipDeviceptr_t>(ptrA), buffer_size));
     HIP_CHECK(hipMemcpyDtoH(B_h.data(), dptr_peer, buffer_size));
     REQUIRE(true == std::equal(B_h.begin(), B_h.end(), A_h.data()));
     HIP_CHECK(hipFree(reinterpret_cast<void*>(dptr_peer)));
@@ -750,7 +754,7 @@ TEST_CASE("Unit_hipMemSetAccess_Vmm2PeerPeerMemCpy") {
   size_t size_mem = ((granularity + buffer_size - 1) / granularity) * granularity;
   // Allocate physical memory
   hipMemGenericAllocationHandle_t handle;
-  hipDeviceptr_t ptrA;
+  void* ptrA;
   HIP_CHECK(hipMemCreate(&handle, size_mem, &prop, 0));
   // Allocate virtual address range
   HIP_CHECK(hipMemAddressReserve(&ptrA, size_mem, 0, 0, 0));
@@ -767,7 +771,7 @@ TEST_CASE("Unit_hipMemSetAccess_Vmm2PeerPeerMemCpy") {
   for (int idx = 0; idx < N; idx++) {
     A_h[idx] = idx;
   }
-  HIP_CHECK(hipMemcpyHtoD(ptrA, A_h.data(), buffer_size));
+  HIP_CHECK(hipMemcpyHtoD(reinterpret_cast<hipDeviceptr_t>(ptrA), A_h.data(), buffer_size));
   // Check Peer Access
   for (deviceId = 1; deviceId < devicecount; deviceId++) {
     std::fill(B_h.begin(), B_h.end(), initializer);
@@ -843,7 +847,7 @@ TEST_CASE("Unit_hipMemSetAccess_Vmm2VMMMemCpy") {
   HIP_CHECK(hipMemCreate(&handle1, size_mem, &prop, 0));
   HIP_CHECK(hipMemCreate(&handle2, size_mem, &prop, 0));
   // Allocate virtual address range
-  hipDeviceptr_t ptrA, ptrB;
+  void *ptrA, *ptrB;
   HIP_CHECK(hipMemAddressReserve(&ptrA, size_mem, 0, 0, 0));
   HIP_CHECK(hipMemAddressReserve(&ptrB, size_mem, 0, 0, 0));
   HIP_CHECK(hipMemMap(ptrA, size_mem, 0, handle1, 0));
@@ -862,9 +866,10 @@ TEST_CASE("Unit_hipMemSetAccess_Vmm2VMMMemCpy") {
   for (int idx = 0; idx < N; idx++) {
     A_h[idx] = idx;
   }
-  HIP_CHECK(hipMemcpyHtoD(ptrA, A_h.data(), buffer_size));
-  HIP_CHECK(hipMemcpyDtoD(ptrB, ptrA, buffer_size));
-  HIP_CHECK(hipMemcpyDtoH(B_h.data(), ptrB, buffer_size));
+  HIP_CHECK(hipMemcpyHtoD(reinterpret_cast<hipDeviceptr_t>(ptrA), A_h.data(), buffer_size));
+  HIP_CHECK(hipMemcpyDtoD(reinterpret_cast<hipDeviceptr_t>(ptrB),
+                          reinterpret_cast<hipDeviceptr_t>(ptrA), buffer_size));
+  HIP_CHECK(hipMemcpyDtoH(B_h.data(), reinterpret_cast<hipDeviceptr_t>(ptrB), buffer_size));
   REQUIRE(true == std::equal(B_h.begin(), B_h.end(), A_h.data()));
   HIP_CHECK(hipMemUnmap(ptrA, size_mem));
   HIP_CHECK(hipMemUnmap(ptrB, size_mem));
@@ -913,7 +918,7 @@ TEST_CASE("Unit_hipMemSetAccess_Vmm2VMMInterDevMemCpy") {
   // Allocate physical memory
   HIP_CHECK(hipMemCreate(&handle, size_mem, &prop, 0));
   // Allocate virtual address range
-  hipDeviceptr_t ptrA;
+  void* ptrA;
   HIP_CHECK(hipMemAddressReserve(&ptrA, size_mem, 0, 0, 0));
   HIP_CHECK(hipMemMap(ptrA, size_mem, 0, handle, 0));
   HIP_CHECK(hipMemRelease(handle));
@@ -928,7 +933,7 @@ TEST_CASE("Unit_hipMemSetAccess_Vmm2VMMInterDevMemCpy") {
   for (int idx = 0; idx < N; idx++) {
     A_h[idx] = idx;
   }
-  HIP_CHECK(hipMemcpyHtoD(ptrA, A_h.data(), buffer_size));
+  HIP_CHECK(hipMemcpyHtoD(reinterpret_cast<hipDeviceptr_t>(ptrA), A_h.data(), buffer_size));
   for (deviceId = 1; deviceId < devicecount; deviceId++) {
     int canAccessPeer = 0;
     hipDevice_t device_other;
@@ -958,7 +963,7 @@ TEST_CASE("Unit_hipMemSetAccess_Vmm2VMMInterDevMemCpy") {
     // Allocate physical memory
     HIP_CHECK(hipMemCreate(&handle_loc, size_mem_loc, &prop_loc, 0));
     // Allocate virtual address range
-    hipDeviceptr_t ptrB;
+    void* ptrB;
     HIP_CHECK(hipMemAddressReserve(&ptrB, size_mem_loc, 0, 0, 0));
     HIP_CHECK(hipMemMap(ptrB, size_mem_loc, 0, handle_loc, 0));
     HIP_CHECK(hipMemRelease(handle_loc));
@@ -969,9 +974,8 @@ TEST_CASE("Unit_hipMemSetAccess_Vmm2VMMInterDevMemCpy") {
     accessDesc_loc.flags = hipMemAccessFlagsProtReadWrite;
     // Make the address accessible to GPU 0
     HIP_CHECK(hipMemSetAccess(ptrB, size_mem_loc, &accessDesc_loc, 1));
-    HIP_CHECK(hipMemcpyPeer(reinterpret_cast<void*>(ptrB), deviceId, reinterpret_cast<void*>(ptrA),
-                            0, buffer_size));
-    HIP_CHECK(hipMemcpyDtoH(B_h.data(), ptrB, buffer_size));
+    HIP_CHECK(hipMemcpyPeer(ptrB, deviceId, ptrA, 0, buffer_size));
+    HIP_CHECK(hipMemcpyDtoH(B_h.data(), reinterpret_cast<hipDeviceptr_t>(ptrB), buffer_size));
     REQUIRE(true == std::equal(B_h.begin(), B_h.end(), A_h.data()));
     HIP_CHECK(hipMemUnmap(ptrB, size_mem_loc));
     HIP_CHECK(hipMemAddressFree(ptrB, size_mem_loc));
@@ -984,7 +988,7 @@ TEST_CASE("Unit_hipMemSetAccess_Vmm2VMMInterDevMemCpy") {
 class vmm_resize_class {
   size_t current_size_tot;
   size_t current_size_rounded_tot;
-  hipDeviceptr_t ptrVmm;
+  void* ptrVmm;
   std::vector<hipMemGenericAllocationHandle_t> vhandle;
   std::vector<size_t> vsize;
   // allocate initial VMM memory chunk
@@ -1014,7 +1018,7 @@ class vmm_resize_class {
     accessDesc.flags = hipMemAccessFlagsProtReadWrite;
     // Make the address accessible to GPU device
     HIP_CHECK(hipMemSetAccess(ptrVmm, size_rounded, &accessDesc, 1));
-    *ptr = ptrVmm;
+    *ptr = reinterpret_cast<hipDeviceptr_t>(ptrVmm);
     current_size_tot += size;
     current_size_rounded_tot += size_rounded;
     return 0;
@@ -1067,7 +1071,7 @@ class vmm_resize_class {
       } else {
         unsigned long long uiptr = reinterpret_cast<unsigned long long>(ptrVmm);
         uiptr = uiptr + vsize[idx - 1];
-        HIP_CHECK(hipMemMap(reinterpret_cast<hipDeviceptr_t>(uiptr), vsize[idx], 0, myhandle, 0));
+        HIP_CHECK(hipMemMap(reinterpret_cast<void*>(uiptr), vsize[idx], 0, myhandle, 0));
       }
       idx++;
     }
@@ -1078,7 +1082,7 @@ class vmm_resize_class {
     accessDesc.flags = hipMemAccessFlagsProtReadWrite;
     // Make the address accessible to GPU 0
     HIP_CHECK(hipMemSetAccess(ptrVmm, (size_rounded + current_size_rounded_tot), &accessDesc, 1));
-    *ptr = ptrVmm;
+    *ptr = reinterpret_cast<hipDeviceptr_t>(ptrVmm);
     current_size_tot += size;
     current_size_rounded_tot += size_rounded;
     return 0;
@@ -1253,7 +1257,7 @@ TEST_CASE("Unit_hipMemSetAccess_negative") {
   // Allocate physical memory
   HIP_CHECK(hipMemCreate(&handle, size_mem, &prop, 0));
   // Allocate virtual address range
-  hipDeviceptr_t ptrA;
+  void* ptrA;
   HIP_CHECK(hipMemAddressReserve(&ptrA, size_mem, 0, 0, 0));
   HIP_CHECK(hipMemMap(ptrA, size_mem, 0, handle, 0));
   // Set access
@@ -1263,8 +1267,7 @@ TEST_CASE("Unit_hipMemSetAccess_negative") {
   accessDesc.flags = hipMemAccessFlagsProtReadWrite;
 
   SECTION("nullptr to ptrA") {
-    REQUIRE(hipMemSetAccess((hipDeviceptr_t) nullptr, size_mem, &accessDesc, 1) ==
-            hipErrorInvalidValue);
+    REQUIRE(hipMemSetAccess(nullptr, size_mem, &accessDesc, 1) == hipErrorInvalidValue);
   }
 
   SECTION("pass zero to size") {
@@ -1310,7 +1313,7 @@ TEST_CASE("Unit_hipMemSetAccess_negative") {
   }
 
   SECTION("uninitialized virtual memory") {
-    hipDeviceptr_t ptrB;
+    void* ptrB;
     HIP_CHECK(hipMemAddressReserve(&ptrB, size_mem, 0, 0, 0));
     REQUIRE(hipMemSetAccess(ptrB, size_mem, &accessDesc, 1) == hipErrorInvalidValue);
     HIP_CHECK(hipMemAddressFree(ptrB, size_mem));
@@ -1335,7 +1338,7 @@ TEST_CASE("Unit_hipMemSetAccessHostDevice_hostalloc") {
   prop.type = hipMemAllocationTypePinned;       // pinned system RAM
   prop.location.type = hipMemLocationTypeHost;  // generic host
   prop.location.id = 0;                         // host id must be 0
-  prop.requestedHandleType = hipMemHandleTypeNone;
+  prop.requestedHandleTypes = hipMemHandleTypeNone;
 
   constexpr size_t N = 1024;
   constexpr size_t bytes = N * sizeof(int);
@@ -1407,7 +1410,7 @@ TEST_CASE("Unit_hipMemSetAccessHost_devicealloc") {
   prop.type = hipMemAllocationTypePinned;         // pinned system RAM
   prop.location.type = hipMemLocationTypeDevice;  // generic host
   prop.location.id = 0;                           // host id must be 0
-  prop.requestedHandleType = hipMemHandleTypeNone;
+  prop.requestedHandleTypes = hipMemHandleTypeNone;
 
   constexpr size_t N = 1024;
   constexpr size_t bytes = N * sizeof(int);
diff --git a/projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemUnmap.cc b/projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemUnmap.cc
index 03c791fcfa..49cc156dc2 100644
--- a/projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemUnmap.cc
+++ b/projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemUnmap.cc
@@ -66,16 +66,14 @@ TEST_CASE("Unit_hipMemUnmap_negative") {
   size_t size_mem = ((granularity + buffer_size - 1) / granularity) * granularity;
 
   hipMemGenericAllocationHandle_t handle;
-  hipDeviceptr_t ptrA;
+  void* ptrA;
   // Allocate physical memory
   HIP_CHECK(hipMemCreate(&handle, size_mem, &prop, 0));
   // Allocate virtual address range
   HIP_CHECK(hipMemAddressReserve(&ptrA, size_mem, 0, 0, 0));
   HIP_CHECK(hipMemMap(ptrA, size_mem, 0, handle, 0));
 
-  SECTION("nullptr to ptrA") {
-    REQUIRE(hipMemUnmap((hipDeviceptr_t) nullptr, size_mem) == hipErrorInvalidValue);
-  }
+  SECTION("nullptr to ptrA") { REQUIRE(hipMemUnmap(nullptr, size_mem) == hipErrorInvalidValue); }
 
   SECTION("pass zero to size") { REQUIRE(hipMemUnmap(ptrA, 0) == hipErrorInvalidValue); }
 
diff --git a/projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemVmm_old.cc b/projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemVmm_old.cc
index 17303eb035..cc22d56723 100644
--- a/projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemVmm_old.cc
+++ b/projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemVmm_old.cc
@@ -59,13 +59,12 @@ TEST_CASE("Unit_hipMemVmm_Basic") {
 
   size_t size = granularity;
   void* reservedAddress{nullptr};
-  HIP_CHECK(hipMemAddressReserve(reinterpret_cast<hipDeviceptr_t*>(&reservedAddress), size, 0,
-                                 (hipDeviceptr_t) nullptr, 0));
+  HIP_CHECK(hipMemAddressReserve(&reservedAddress, size, 0, nullptr, 0));
 
   hipMemGenericAllocationHandle_t gaHandle;
   HIP_CHECK(hipMemCreate(&gaHandle, size, &memAllocationProp, 0));
 
-  HIP_CHECK(hipMemMap(reinterpret_cast<hipDeviceptr_t>(reservedAddress), size, 0, gaHandle, 0));
+  HIP_CHECK(hipMemMap(reservedAddress, size, 0, gaHandle, 0));
 
   hipDevice_t device;
   HIP_CHECK(hipDeviceGet(&device, 0));
@@ -76,7 +75,7 @@ TEST_CASE("Unit_hipMemVmm_Basic") {
   std::vector<char> values(size);
   const char value = 1;
 
-  HIP_CHECK(hipMemSetAccess(reinterpret_cast<hipDeviceptr_t>(reservedAddress), size, &desc, 1));
+  HIP_CHECK(hipMemSetAccess(reservedAddress, size, &desc, 1));
   HIP_CHECK(hipMemset(reservedAddress, value, size));
   HIP_CHECK(hipMemcpy(&values[0], reservedAddress, size, hipMemcpyDeviceToHost));
 
@@ -84,9 +83,9 @@ TEST_CASE("Unit_hipMemVmm_Basic") {
     REQUIRE(values[i] == value);
   }
 
-  HIP_CHECK(hipMemUnmap(reinterpret_cast<hipDeviceptr_t>(reservedAddress), size));
+  HIP_CHECK(hipMemUnmap(reservedAddress, size));
 
   HIP_CHECK(hipMemRelease(gaHandle));
-  HIP_CHECK(hipMemAddressFree(reinterpret_cast<hipDeviceptr_t>(reservedAddress), size));
+  HIP_CHECK(hipMemAddressFree(reservedAddress, size));
   CTX_DESTROY();
 }
diff --git a/projects/hipother/hipnv/include/hip/nvidia_detail/nvidia_hip_runtime_api.h b/projects/hipother/hipnv/include/hip/nvidia_detail/nvidia_hip_runtime_api.h
index f775ea8868..22c8afee56 100644
--- a/projects/hipother/hipnv/include/hip/nvidia_detail/nvidia_hip_runtime_api.h
+++ b/projects/hipother/hipnv/include/hip/nvidia_detail/nvidia_hip_runtime_api.h
@@ -1930,6 +1930,7 @@ typedef enum CUmemAllocationGranularity_flags_enum hipMemAllocationGranularity_f
 typedef enum cudaMemLocationType hipMemLocationType;
 #define hipMemLocationTypeInvalid cudaMemLocationTypeInvalid
 #define hipMemLocationTypeDevice cudaMemLocationTypeDevice
+#define hipMemLocationTypeHost cudaMemLocationTypeHost
 #define hipMemHandleTypeNone cudaMemHandleTypeNone
 #define hipMemHandleTypePosixFileDescriptor cudaMemHandleTypePosixFileDescriptor
 #define hipMemHandleTypeWin32 cudaMemHandleTypeWin32
@@ -3057,12 +3058,12 @@ inline static hipError_t hipMemRelease(hipMemGenericAllocationHandle_t handle) {
   return hipCUResultTohipError(cuMemRelease(handle));
 }
 inline static hipError_t hipMemAddressFree(void* devPtr, size_t size) {
-  return hipCUResultTohipError(cuMemAddressFree(reinterpret_cast<CUdeviceptr>(devPtr), size));
+  return hipCUResultTohipError(cuMemAddressFree((CUdeviceptr)devPtr, size));
 }
 inline static hipError_t hipMemAddressReserve(void** ptr, size_t size, size_t alignment, void* addr,
                                               unsigned long long flags) {
-  return hipCUResultTohipError(cuMemAddressReserve(reinterpret_cast<CUdeviceptr*>(ptr), size,
-                                                   alignment, reinterpret_cast<CUdeviceptr>(addr),
+  return hipCUResultTohipError(cuMemAddressReserve((CUdeviceptr*)ptr, size,
+                                                   alignment, (CUdeviceptr)addr,
                                                    flags));
 }
 inline static hipError_t hipMemExportToShareableHandle(void* shareableHandle,
@@ -3075,10 +3076,10 @@ inline static hipError_t hipMemExportToShareableHandle(void* shareableHandle,
 inline static hipError_t hipMemGetAccess(unsigned long long* flags, const hipMemLocation* location,
                                          void* ptr) {
   if (location == NULL) {
-    return hipCUResultTohipError(cuMemGetAccess(flags, NULL, reinterpret_cast<CUdeviceptr>(ptr)));
+    return hipCUResultTohipError(cuMemGetAccess(flags, NULL, (CUdeviceptr)ptr));
   } else {
     CUmemLocation loc = hipMemLocationToCUmemLocation(location);
-    return hipCUResultTohipError(cuMemGetAccess(flags, &loc, reinterpret_cast<CUdeviceptr>(ptr)));
+    return hipCUResultTohipError(cuMemGetAccess(flags, &loc, (CUdeviceptr)ptr));
   }
 }
 inline static hipError_t hipMemGetAllocationPropertiesFromHandle(
@@ -3102,7 +3103,7 @@ inline static hipError_t hipMemMap(void* ptr, size_t size, size_t offset,
                                    hipMemGenericAllocationHandle_t handle,
                                    unsigned long long flags) {
   return hipCUResultTohipError(
-      cuMemMap(reinterpret_cast<CUdeviceptr>(ptr), size, offset, handle, flags));
+      cuMemMap((CUdeviceptr)ptr, size, offset, handle, flags));
 }
 inline static hipError_t hipMemMapArrayAsync(hipArrayMapInfo* mapInfoList, unsigned int count,
                                              hipStream_t stream) {
@@ -3116,17 +3117,17 @@ inline static hipError_t hipMemSetAccess(void* ptr, size_t size, const hipMemAcc
                                          size_t count) {
   if (desc == NULL) {
     return hipCUResultTohipError(
-        cuMemSetAccess(reinterpret_cast<CUdeviceptr>(ptr), size, NULL, count));
+        cuMemSetAccess((CUdeviceptr)ptr, size, NULL, count));
   } else {
     CUmemAccessDesc* cuDesc = hipMemAccessDescToCUmemAccessDesc(desc, count);
     auto status = hipCUResultTohipError(
-        cuMemSetAccess(reinterpret_cast<CUdeviceptr>(ptr), size, cuDesc, count));
+        cuMemSetAccess((CUdeviceptr)ptr, size, cuDesc, count));
     free(cuDesc);
     return status;
   }
 }
 inline static hipError_t hipMemUnmap(void* ptr, size_t size) {
-  return hipCUResultTohipError(cuMemUnmap(reinterpret_cast<CUdeviceptr>(ptr), size));
+  return hipCUResultTohipError(cuMemUnmap((CUdeviceptr)ptr, size));
 }
 #endif  // CUDA_VERSION >= CUDA_10020
 
@@ -3464,16 +3465,16 @@ inline static hipError_t hipCtxGetDevice(hipDevice_t* device) {
   return hipCUResultTohipError(cuCtxGetDevice(device));
 }
 
-inline static hipError_t hipCtxGetApiVersion(hipCtx_t ctx, int* apiVersion) {
-  return hipCUResultTohipError(cuCtxGetApiVersion(ctx, (unsigned int*)apiVersion));
+inline static hipError_t hipCtxGetApiVersion(hipCtx_t ctx, unsigned int* apiVersion) {
+  return hipCUResultTohipError(cuCtxGetApiVersion(ctx, apiVersion));
 }
 
-inline static hipError_t hipCtxGetCacheConfig(hipFuncCache* cacheConfig) {
-  return hipCUResultTohipError(cuCtxGetCacheConfig(cacheConfig));
+inline static hipError_t hipCtxGetCacheConfig(hipFuncCache_t* cacheConfig) {
+  return hipCUResultTohipError(cuCtxGetCacheConfig((CUfunc_cache*)cacheConfig));
 }
 
-inline static hipError_t hipCtxSetCacheConfig(hipFuncCache cacheConfig) {
-  return hipCUResultTohipError(cuCtxSetCacheConfig(cacheConfig));
+inline static hipError_t hipCtxSetCacheConfig(hipFuncCache_t cacheConfig) {
+  return hipCUResultTohipError(cuCtxSetCacheConfig((CUfunc_cache)cacheConfig));
 }
 
 inline static hipError_t hipCtxSetSharedMemConfig(hipSharedMemConfig config) {