[SYCLomatic] Add 11 cuda driver API query Mappings.

clang/examples/DPCT/Driver/cuMemGetAllocationPropertiesFromHandle.cu

@@ -0,0 +1,9 @@
+// Option: --use-experimental-features=virtual_mem
+#include <cuda.h>
+void test() {
+  CUmemAllocationProp prop = {};
+  CUmemGenericAllocationHandle allocHandle;
+  // Start
+  cuMemGetAllocationPropertiesFromHandle(&prop/*CUmemAllocationProp **/, allocHandle/*CUmemGenericAllocationHandle*/);
+  // End
+}

clang/examples/DPCT/Driver/cuMipmappedArrayDestroy.cu

@@ -0,0 +1,7 @@
+// Option: --use-experimental-features=bindless_images
+void test() {
+    CUmipmappedArray mmArray;
+    // Start
+    cuMipmappedArrayDestroy(mmArray/*CUmipmappedArray*/);
+    // End
+}

clang/examples/DPCT/Driver/cuMipmappedArrayGetLevel.cu

@@ -0,0 +1,8 @@
+// Option: --use-experimental-features=bindless_images
+void test() {
+  CUmipmappedArray mmArray;
+  CUarray level_arr;
+  // Start
+  cuMipmappedArrayGetLevel(&level_arr/*CUarray **/, mmArray/*CUmipmappedArray*/, 1/*unsigned int*/);
+  // End
+}

clang/examples/DPCT/Driver/cuPointerGetAttribute.cu

@@ -0,0 +1,9 @@
+// Option: --usm-level=none
+#include <cuda.h>
+void test() {
+  void *base_ptr;
+  void *ptr;
+  // Start
+  cuPointerGetAttribute(base_ptr/*void **/, CU_POINTER_ATTRIBUTE_RANGE_START_ADDR/*CUpointer_attribute*/, (CUdeviceptr)ptr/*CUdeviceptr*/);
+  // End
+}

clang/examples/DPCT/Driver/cuStreamGetCtx.cu

@@ -0,0 +1,6 @@
+#include <cuda.h>
+void test(CUstream stream, CUcontext& context) {
+  // Start
+  cuStreamGetCtx(stream/*CUstream*/, &context/*CUcontext **/);
+  // End
+}

clang/examples/DPCT/Driver/cuTexRefCreate.cu

@@ -0,0 +1,8 @@
+// Option: --use-experimental-features=bindless_images
+#include <cuda.h>
+void test() {
+  CUtexref r;
+  // Start
+  cuTexRefCreate(&r/*CUtexref **/);
+  // End
+}

clang/examples/DPCT/Driver/cuTexRefDestroy.cu

@@ -0,0 +1,8 @@
+// Option: --use-experimental-features=bindless_images
+#include <cuda.h>
+void test() {
+  CUtexref r;
+  // Start
+  cuTexRefDestroy(r/*CUtexref*/);
+  // End
+}

clang/examples/DPCT/Driver/cuTexRefGetArray.cu

@@ -0,0 +1,9 @@
+// Option: --use-experimental-features=bindless_images
+#include <cuda.h>
+void test() {
+  CUtexref r;
+  CUarray a;
+  // Start
+  cuTexRefGetArray(&a/*CUarray **/, r/*CUtexref*/);
+  // End
+}

clang/examples/DPCT/Driver/cuTexRefGetFlags.cu

@@ -1,3 +1,4 @@
+#include <cuda.h>
 void test(unsigned int *pu) {
   // Start
   CUtexref t;

clang/examples/DPCT/Driver/cuTexRefGetMipmapFilterMode.cu

@@ -0,0 +1,9 @@
+// Option: --use-experimental-features=bindless_images
+#include <cuda.h>
+void test() {
+  CUfilter_mode fm = CU_TR_FILTER_MODE_POINT;
+  CUtexref texRef;
+  // Start
+  cuTexRefGetMipmapFilterMode(&fm /*CUfilter_mode **/, texRef /*CUtexref*/);
+  // End
+}

clang/examples/DPCT/Driver/cuTexRefGetMipmapLevelClamp.cu

@@ -0,0 +1,10 @@
+// Option: --use-experimental-features=bindless_images
+#include <cuda.h>
+
+void test() {
+  CUtexref texRef;
+  float min_clamp, max_clamp;
+  // Start
+  cuTexRefGetMipmapLevelClamp(&min_clamp/*float **/, &max_clamp/*float **/, texRef/*CUtexref*/);
+  // End
+}

clang/lib/DPCT/DPCT.cpp

@@ -1065,6 +1065,9 @@ int runDPCT(int argc, const char **argv) {
           NoDRYPattern.setValue(true);
         } else if (Option == "--enable-profiling") {
           EnablepProfiling.setValue(true);
+        } else if (Option.starts_with("--usm-level=")) {
+          if (Option.ends_with("none"))
+            USMLevel.setValue(UsmLevel::UL_None);
         }
         // Need add more option.
       }

clang/lib/DPCT/RulesLang/APINamesTexture.inc

@@ -467,6 +467,17 @@ CONDITIONAL_FACTORY_ENTRY(
                             Diagnostics::TRY_EXPERIMENTAL_FEATURE,
                             ARG("cuSurfRefGetArray"),
                             ARG("--use-experimental-features=bindless_images")))
+CONDITIONAL_FACTORY_ENTRY(
+    UseExtBindlessImages,
+    ASSIGNABLE_FACTORY(ASSIGN_FACTORY_ENTRY(
+        "cuTexRefGetArray", DEREF(0),
+        CALL(MapNames::getDpctNamespace() + "experimental::get_img_mem",
+             ARG(1)))),
+    UNSUPPORT_FACTORY_ENTRY("cuTexRefGetArray",
+                            Diagnostics::TRY_EXPERIMENTAL_FEATURE,
+                            ARG("cuTexRefGetArray"),
+                            ARG("--use-experimental-features=bindless_images")))
+
 
 CONDITIONAL_FACTORY_ENTRY(
     UseExtBindlessImages,

clang/lib/DPCT/RulesLang/RulesLangTexture.cpp

@@ -589,6 +589,7 @@ void TextureRule::registerMatcher(MatchFinder &MF) {
       "cuTexRefDestroy",
       "cuSurfRefSetArray",
       "cuSurfRefGetArray",
+      "cuTexRefGetArray",
       "cuTexRefSetArray",
       "cuTexRefSetFormat",
       "cuTexRefSetAddressMode",

clang/test/dpct/query_api_mapping/Driver/test-after9.2.cu

@@ -8,3 +8,30 @@
 // CUDEVICEGETUUID-NEXT:   cuDeviceGetUuid(pu /*CUuuid **/, d /*CUdevice*/);
 // CUDEVICEGETUUID-NEXT: Is migrated to:
 // CUDEVICEGETUUID-NEXT:   *pu = dpct::get_device(d).get_device_info().get_uuid();
+
+// RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=cuMemAddressReserve | FileCheck %s -check-prefix=CUMEMADDRESSRESERVE
+// CUMEMADDRESSRESERVE: CUDA API:
+// CUMEMADDRESSRESERVE-NEXT:   cuMemAddressReserve(ptr /*CUdeviceptr **/, size /*size_t*/,
+// CUMEMADDRESSRESERVE-NEXT:                      alignment /*size_t*/, addr /*CUdeviceptr*/,
+// CUMEMADDRESSRESERVE-NEXT:                      flags /*unsigned long long*/);
+// CUMEMADDRESSRESERVE-NEXT: Is migrated to (with the option --use-experimental-features=virtual_mem):
+// CUMEMADDRESSRESERVE-NEXT:   *ptr = (dpct::device_ptr)sycl::ext::oneapi::experimental::reserve_virtual_mem((uintptr_t)addr, size, dpct::get_current_device().get_context());
+
+
+// RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=cuMemGetAllocationPropertiesFromHandle | FileCheck %s -check-prefix=CUMEMGETALLOCATIONPROPERTIESFROMHANDLE
+// CUMEMGETALLOCATIONPROPERTIESFROMHANDLE: CUDA API:
+// CUMEMGETALLOCATIONPROPERTIESFROMHANDLE-NEXT:   cuMemGetAllocationPropertiesFromHandle(&prop/*CUmemAllocationProp **/, allocHandle/*CUmemGenericAllocationHandle*/);
+// CUMEMGETALLOCATIONPROPERTIESFROMHANDLE-NEXT: Is migrated to (with the option --use-experimental-features=virtual_mem):
+// CUMEMGETALLOCATIONPROPERTIESFROMHANDLE-NEXT:    prop.location.id = dpct::get_device_id(allocHandle->get_device());
+
+// RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=cuPointerGetAttribute | FileCheck %s -check-prefix=CUPOINTERGETATTRIBUTE
+// CUPOINTERGETATTRIBUTE: CUDA API:
+// CUPOINTERGETATTRIBUTE-NEXT:   cuPointerGetAttribute(base_ptr/*void **/, CU_POINTER_ATTRIBUTE_RANGE_START_ADDR/*CUpointer_attribute*/, (CUdeviceptr)ptr/*CUdeviceptr*/);
+// CUPOINTERGETATTRIBUTE-NEXT: Is migrated to (with the option --usm-level=none):
+// CUPOINTERGETATTRIBUTE-NEXT:   base_ptr = dpct::get_base_addr((dpct::device_ptr)ptr);
+
+// RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=cuStreamGetCtx | FileCheck %s -check-prefix=CUSTREAMGETCTX
+// CUSTREAMGETCTX: CUDA API:
+// CUSTREAMGETCTX-NEXT:   cuStreamGetCtx(stream/*CUstream*/, &context/*CUcontext **/);
+// CUSTREAMGETCTX-NEXT: Is migrated to:
+// CUSTREAMGETCTX-NEXT:   context = dpct::get_device_id(stream->get_device());

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

@@ -561,3 +561,48 @@
 // CUDEVICECANACCESSPEER-NEXT:   cuDeviceCanAccessPeer(pi /*int **/, d1 /*CUdevice*/, d2 /*CUdevice*/);
 // CUDEVICECANACCESSPEER-NEXT: Is migrated to:
 // CUDEVICECANACCESSPEER-NEXT:   *pi = dpct::get_device(d1).ext_oneapi_can_access_peer(dpct::get_device(d2));
+
+
+/// Bindless Image
+
+// RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=cuMipmappedArrayDestroy | FileCheck %s -check-prefix=CUMIPMAPPEDARRAYDESTROY
+// CUMIPMAPPEDARRAYDESTROY: CUDA API:
+// CUMIPMAPPEDARRAYDESTROY-NEXT:   cuMipmappedArrayDestroy(mmArray/*CUmipmappedArray*/);
+// CUMIPMAPPEDARRAYDESTROY-NEXT: Is migrated to (with the option --use-experimental-features=bindless_images):
+// CUMIPMAPPEDARRAYDESTROY-NEXT:   delete mmArray;
+
+// RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=cuMipmappedArrayGetLevel | FileCheck %s -check-prefix=CUMIPMAPPEDARRAYGETLEVEL
+// CUMIPMAPPEDARRAYGETLEVEL: CUDA API:
+// CUMIPMAPPEDARRAYGETLEVEL-NEXT:   cuMipmappedArrayGetLevel(&level_arr/*CUarray*/, mmArray/*CUmipmappedArray*/, 1/*level*/);
+// CUMIPMAPPEDARRAYGETLEVEL-NEXT: Is migrated to (with the option --use-experimental-features=bindless_images):
+// CUMIPMAPPEDARRAYGETLEVEL-NEXT:   level_arr = mmArray->get_mip_level(1);
+
+// RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=cuTexRefCreate | FileCheck %s -check-prefix=CUTEXREFCREATE
+// CUTEXREFCREATE: CUDA API:
+// CUTEXREFCREATE-NEXT:   cuTexRefCreate(&r/*CUtexref **/);
+// CUTEXREFCREATE-NEXT: Is migrated to (with the option --use-experimental-features=bindless_images):
+// CUTEXREFCREATE-NEXT:   r = new dpct::experimental::bindless_image_wrapper_base();
+
+// RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=cuTexRefDestroy | FileCheck %s -check-prefix=CUTEXREFDESTROY
+// CUTEXREFDESTROY: CUDA API:
+// CUTEXREFDESTROY-NEXT:   cuTexRefDestroy(r/*CUtexref*/);
+// CUTEXREFDESTROY-NEXT: Is migrated to (with the option --use-experimental-features=bindless_images):
+// CUTEXREFDESTROY-NEXT:   delete r;
+
+// RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=cuTexRefGetArray | FileCheck %s -check-prefix=CUTEXREFGETARRAY
+// CUTEXREFGETARRAY: CUDA API:
+// CUTEXREFGETARRAY-NEXT:   cuTexRefGetArray(phArray /*CUarray **/, hTexRef /*CUtexref*/);
+// CUTEXREFGETARRAY-NEXT: Is migrated to (with the option --use-experimental-features=bindless_images):
+// CUTEXREFGETARRAY-NEXT:   a = dpct::experimental::get_img_mem(r);
+
+// RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=cuTexRefGetMipmapFilterMode | FileCheck %s -check-prefix=CUTEXREFGETMIPMAPFILTERMODE
+// CUTEXREFGETMIPMAPFILTERMODE: CUDA API:
+// CUTEXREFGETMIPMAPFILTERMODE-NEXT:   cuTexRefGetMipmapFilterMode(&fm /*CUfilter_mode **/, texRef /*CUtexref*/);
+// CUTEXREFGETMIPMAPFILTERMODE-NEXT: Is migrated to (with the option --use-experimental-features=bindless_images):
+// CUTEXREFGETMIPMAPFILTERMODE-NEXT:   fm = texRef->get_mip_filtering_mode();
+
+// RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=cuTexRefGetMipmapLevelClamp | FileCheck %s -check-prefix=CUTEXREFGETMIPMAPLEVELCLAMP
+// CUTEXREFGETMIPMAPLEVELCLAMP: CUDA API:
+// CUTEXREFGETMIPMAPLEVELCLAMP-NEXT:   cuTexRefGetMipmapLevelClamp(&min_clamp/*float **/, &max_clamp/*float **/, texRef/*CUtexref*/);
+// CUTEXREFGETMIPMAPLEVELCLAMP-NEXT: Is migrated to (with the option --use-experimental-features=bindless_images):
+// CUTEXREFGETMIPMAPLEVELCLAMP-NEXT:   texRef->get_mip_level_clamp(&min_clamp, &max_clamp);

clang/test/dpct/query_api_mapping/test_all.cu

@@ -757,6 +757,7 @@
 // CHECK-NEXT: cuMemFree
 // CHECK-NEXT: cuMemFreeHost
 // CHECK-NEXT: cuMemGetAllocationGranularity
+// CHECK-NEXT: cuMemGetAllocationPropertiesFromHandle
 // CHECK-NEXT: cuMemGetInfo
 // CHECK-NEXT: cuMemHostAlloc
 // CHECK-NEXT: cuMemHostGetDevicePointer
@@ -804,28 +805,37 @@
 // CHECK-NEXT: cuMemsetD32Async
 // CHECK-NEXT: cuMemsetD8
 // CHECK-NEXT: cuMemsetD8Async
+// CHECK-NEXT: cuMipmappedArrayDestroy
+// CHECK-NEXT: cuMipmappedArrayGetLevel
 // CHECK-NEXT: cuModuleGetFunction
 // CHECK-NEXT: cuModuleGetTexRef
 // CHECK-NEXT: cuModuleLoad
 // CHECK-NEXT: cuModuleLoadData
 // CHECK-NEXT: cuModuleLoadDataEx
 // CHECK-NEXT: cuModuleUnload
 // CHECK-NEXT: cuOccupancyMaxActiveBlocksPerMultiprocessor
+// CHECK-NEXT: cuPointerGetAttribute
 // CHECK-NEXT: cuPointerGetAttributes
 // CHECK-NEXT: cuStreamAddCallback
 // CHECK-NEXT: cuStreamAttachMemAsync
 // CHECK-NEXT: cuStreamCreate
 // CHECK-NEXT: cuStreamDestroy
+// CHECK-NEXT: cuStreamGetCtx
 // CHECK-NEXT: cuStreamQuery
 // CHECK-NEXT: cuStreamSynchronize
 // CHECK-NEXT: cuStreamWaitEvent
 // CHECK-NEXT: cuTexObjectCreate
 // CHECK-NEXT: cuTexObjectDestroy
 // CHECK-NEXT: cuTexObjectGetResourceDesc
 // CHECK-NEXT: cuTexObjectGetTextureDesc
+// CHECK-NEXT: cuTexRefCreate
+// CHECK-NEXT: cuTexRefDestroy
 // CHECK-NEXT: cuTexRefGetAddressMode
+// CHECK-NEXT: cuTexRefGetArray
 // CHECK-NEXT: cuTexRefGetFilterMode
 // CHECK-NEXT: cuTexRefGetFlags
+// CHECK-NEXT: cuTexRefGetMipmapFilterMode
+// CHECK-NEXT: cuTexRefGetMipmapLevelClamp
 // CHECK-NEXT: cuTexRefSetAddress
 // CHECK-NEXT: cuTexRefSetAddress2D
 // CHECK-NEXT: cuTexRefSetAddressMode

clang/test/dpct/texture/texture_reference_bindless_image.cu

@@ -75,6 +75,8 @@ void driverTextureReferenceManagement() {
   // CHECK: r->set_channel_type(f);
   // CHECK-NEXT: r->set_channel_num(i);
   cuTexRefSetFormat(r, f, i);
+  // CHECK: a = dpct::experimental::get_img_mem(r);
+  cuTexRefGetArray(&a, r);
 }
 void test_surf_ref() {
   // CHECK: dpct::experimental::image_mem_wrapper_ptr arr;