[SYCLomatic] Add query api mapping for virtual memory APIs (#2586)

Signed-off-by: intwanghao <[email protected]>

clang/examples/DPCT/Driver/cuMemAddressFree.cu

@@ -0,0 +1,7 @@
+// Option: --use-experimental-features=virtual_mem
+#include <cuda.h>
+void test(CUdeviceptr ptr, size_t size) {
+  // Start
+  cuMemAddressFree(ptr /*CUdeviceptr*/, size /*size_t*/);
+  // End
+}

clang/examples/DPCT/Driver/cuMemAddressReserve.cu

@@ -0,0 +1,10 @@
+// Option: --use-experimental-features=virtual_mem
+#include <cuda.h>
+void test(CUdeviceptr *ptr, size_t size, size_t alignment, CUdeviceptr addr,
+          unsigned long long flags) {
+  // Start
+  cuMemAddressReserve(ptr /*CUdeviceptr **/, size /*size_t*/,
+                      alignment /*size_t*/, addr /*CUdeviceptr*/,
+                      flags /*unsigned long long*/);
+  // End
+}

clang/examples/DPCT/Driver/cuMemCreate.cu

@@ -0,0 +1,9 @@
+// Option: --use-experimental-features=virtual_mem
+#include <cuda.h>
+void test(size_t size, CUmemAllocationProp *prop, unsigned long long flags) {
+  // Start
+  CUmemGenericAllocationHandle *handle;
+  cuMemCreate(handle /*CUmemGenericAllocationHandle **/, size /*size_t*/,
+              prop /*CUmemAllocationProp **/, flags /*unsigned long long*/);
+  // End
+}

clang/examples/DPCT/Driver/cuMemGetAllocationGranularity.cu

@@ -0,0 +1,12 @@
+// Option: --use-experimental-features=virtual_mem
+#include <cuda.h>
+void test(size_t *granularity, CUmemAllocationProp *prop,
+          CUmemAllocationGranularity_flags option) {
+  // Start
+  cuMemGetAllocationGranularity(granularity /*size_t
+                                             **/
+                                ,
+                                prop /*CUmemAllocationProp **/,
+                                option /*CUmemAllocationGranularity_flags*/);
+  // End
+}

clang/examples/DPCT/Driver/cuMemMap.cu

@@ -0,0 +1,11 @@
+// Option: --use-experimental-features=virtual_mem
+#include <cuda.h>
+void test(CUdeviceptr ptr, size_t size, size_t offset,
+          unsigned long long flags) {
+  // Start
+  CUmemGenericAllocationHandle handle;
+  cuMemMap(ptr /*CUdeviceptr*/, size /*size_t*/, offset /*size_t*/,
+           handle /*CUmemGenericAllocationHandle*/,
+           flags /*unsigned long long */);
+  // End
+}

clang/examples/DPCT/Driver/cuMemRelease.cu

@@ -0,0 +1,8 @@
+// Option: --use-experimental-features=virtual_mem
+#include <cuda.h>
+void test() {
+  // Start
+  CUmemGenericAllocationHandle handle;
+  cuMemRelease(handle /*CUmemGenericAllocationHandle */);
+  // End
+}

clang/examples/DPCT/Driver/cuMemSetAccess.cu

@@ -0,0 +1,8 @@
+// Option: --use-experimental-features=virtual_mem
+#include <cuda.h>
+void test(CUdeviceptr ptr, size_t size, CUmemAccessDesc *desc, size_t count) {
+  // Start
+  cuMemSetAccess(ptr /*CUdeviceptr*/, size /*size_t*/,
+                 desc /*CUmemAccessDesc **/, count /*size_t*/);
+  // End
+}

clang/examples/DPCT/Driver/cuMemUnmap.cu

@@ -0,0 +1,7 @@
+// Option: --use-experimental-features=virtual_mem
+#include <cuda.h>
+void test(CUdeviceptr ptr, size_t size) {
+  // Start
+  cuMemUnmap(ptr /*CUdeviceptr*/, size /*size_t*/);
+  // End
+}

clang/lib/DPCT/DPCT.cpp

@@ -968,6 +968,8 @@ int runDPCT(int argc, const char **argv) {
             Experimentals.addValue(ExperimentalFeatures::Exp_LogicalGroup);
           else if (Option.ends_with("root-group"))
             Experimentals.addValue(ExperimentalFeatures::Exp_RootGroup);
+          else if (Option.ends_with("virtual_mem"))
+            Experimentals.addValue(ExperimentalFeatures::Exp_VirtualMemory);
           else if (Option.ends_with("masked-sub-group-operation"))
             Experimentals.addValue(
                 ExperimentalFeatures::Exp_MaskedSubGroupFunction);

clang/test/dpct/query_api_mapping/Driver/test_virtual_memory.cu

@@ -0,0 +1,59 @@
+// UNSUPPORTED: cuda-8.0, cuda-9.0, cuda-9.1, cuda-9.2, cuda-10.0, cuda-10.1, cuda-10.2
+// UNSUPPORTED: v8.0, v9.0, v9.1, v9.2, v10.0, v10.1, v10.2
+/// Virtual Memory Management
+
+// RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=cuMemMap | FileCheck %s -check-prefix=cuMemMap
+// cuMemMap:  CUDA API:
+// cuMemMap-NEXT:    CUmemGenericAllocationHandle handle;
+// cuMemMap-NEXT:    cuMemMap(ptr /*CUdeviceptr*/, size /*size_t*/, offset /*size_t*/,
+// cuMemMap-NEXT:             handle /*CUmemGenericAllocationHandle*/,
+// cuMemMap-NEXT:             flags /*unsigned long long */);
+// cuMemMap-NEXT:  Is migrated to (with the option --use-experimental-features=virtual_mem):
+// cuMemMap-NEXT:    dpct::experimental::physical_mem_ptr handle;
+// cuMemMap-NEXT:    handle->map((uintptr_t)ptr, size, sycl::ext::oneapi::experimental::address_access_mode::read_write, offset);
+
+// RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=cuMemRelease | FileCheck %s -check-prefix=cuMemRelease
+// cuMemRelease: CUDA API:
+// cuMemRelease-NEXT:   CUmemGenericAllocationHandle handle;
+// cuMemRelease-NEXT:   cuMemRelease(handle /*CUmemGenericAllocationHandle */);
+// cuMemRelease-NEXT: Is migrated to (with the option --use-experimental-features=virtual_mem):
+// cuMemRelease-NEXT:   dpct::experimental::physical_mem_ptr handle;
+// cuMemRelease-NEXT:   delete (handle);
+
+// RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=cuMemUnmap | FileCheck %s -check-prefix=cuMemUnmap
+// cuMemUnmap: CUDA API:
+// cuMemUnmap-NEXT:   cuMemUnmap(ptr /*CUdeviceptr*/, size /*size_t*/);
+// cuMemUnmap-NEXT: Is migrated to (with the option --use-experimental-features=virtual_mem):
+// cuMemUnmap-NEXT:   sycl::ext::oneapi::experimental::unmap(ptr, size, dpct::get_current_device().get_context());
+
+// RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=cuMemCreate | FileCheck %s -check-prefix=cuMemCreate
+// cuMemCreate: CUDA API:
+// cuMemCreate-NEXT:   CUmemGenericAllocationHandle *handle;
+// cuMemCreate-NEXT:   cuMemCreate(handle /*CUmemGenericAllocationHandle **/, size /*size_t*/,
+// cuMemCreate-NEXT:               prop /*CUmemAllocationProp **/, flags /*unsigned long long*/);
+// cuMemCreate-NEXT: Is migrated to (with the option --use-experimental-features=virtual_mem):
+// cuMemCreate-NEXT:   dpct::experimental::physical_mem_ptr *handle;
+// cuMemCreate-NEXT:   *handle = new sycl::ext::oneapi::experimental::physical_mem(dpct::get_device(prop->location.id), dpct::get_device(prop->location.id).get_context(), size);
+
+// RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=cuMemAddressFree | FileCheck %s -check-prefix=cuMemAddressFree
+// cuMemAddressFree: CUDA API:
+// cuMemAddressFree-NEXT:   cuMemAddressFree(ptr /*CUdeviceptr*/, size /*size_t*/);
+// cuMemAddressFree-NEXT: Is migrated to (with the option --use-experimental-features=virtual_mem):
+// cuMemAddressFree-NEXT:   sycl::ext::oneapi::experimental::free_virtual_mem((uintptr_t)ptr, size, dpct::get_current_device().get_context());
+
+// RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=cuMemSetAccess | FileCheck %s -check-prefix=cuMemSetAccess
+// cuMemSetAccess: CUDA API:
+// cuMemSetAccess-NEXT:   cuMemSetAccess(ptr /*CUdeviceptr*/, size /*size_t*/,
+// cuMemSetAccess-NEXT:                  desc /*CUmemAccessDesc **/, count /*size_t*/);
+// cuMemSetAccess-NEXT: Is migrated to (with the option --use-experimental-features=virtual_mem):
+// cuMemSetAccess-NEXT:   sycl::ext::oneapi::experimental::set_access_mode(ptr, size, desc->flags, dpct::get_device(desc->location.id).get_context());
+
+// RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=cuMemGetAllocationGranularity | FileCheck %s -check-prefix=cuMemGetAllocationGranularity
+// cuMemGetAllocationGranularity: CUDA API:
+// cuMemGetAllocationGranularity-NEXT:   cuMemGetAllocationGranularity(granularity /*size_t
+// cuMemGetAllocationGranularity-NEXT:                                              **/
+// cuMemGetAllocationGranularity-NEXT:                                 ,
+// cuMemGetAllocationGranularity-NEXT:                                 prop /*CUmemAllocationProp **/,
+// cuMemGetAllocationGranularity-NEXT:                                 option /*CUmemAllocationGranularity_flags*/);
+// cuMemGetAllocationGranularity-NEXT: Is migrated to (with the option --use-experimental-features=virtual_mem):
+// cuMemGetAllocationGranularity-NEXT:   *granularity = sycl::ext::oneapi::experimental::get_mem_granularity(dpct::get_device(prop->location.id), dpct::get_device(prop->location.id).get_context(), option);

clang/test/dpct/query_api_mapping/test_all.cu

@@ -645,20 +645,28 @@
 // CHECK-NEXT: cuGetErrorString
 // CHECK-NEXT: cuInit
 // CHECK-NEXT: cuLaunchKernel
+// CHECK-NEXT: cuMemAddressFree
+// CHECK-NEXT: cuMemAddressReserve
 // CHECK-NEXT: cuMemAdvise
 // CHECK-NEXT: cuMemAlloc
 // CHECK-NEXT: cuMemAllocHost
 // CHECK-NEXT: cuMemAllocManaged
 // CHECK-NEXT: cuMemAllocPitch
+// CHECK-NEXT: cuMemCreate
 // CHECK-NEXT: cuMemFree
 // CHECK-NEXT: cuMemFreeHost
+// CHECK-NEXT: cuMemGetAllocationGranularity
 // CHECK-NEXT: cuMemGetInfo
 // CHECK-NEXT: cuMemHostAlloc
 // CHECK-NEXT: cuMemHostGetDevicePointer
 // CHECK-NEXT: cuMemHostGetFlags
 // CHECK-NEXT: cuMemHostRegister
 // CHECK-NEXT: cuMemHostUnregister
+// CHECK-NEXT: cuMemMap
 // CHECK-NEXT: cuMemPrefetchAsync
+// CHECK-NEXT: cuMemRelease
+// CHECK-NEXT: cuMemSetAccess
+// CHECK-NEXT: cuMemUnmap
 // CHECK-NEXT: cuMemcpy
 // CHECK-NEXT: cuMemcpy2D
 // CHECK-NEXT: cuMemcpy2DAsync