[SYCLomatic] Support migration of get VIRTUAL_MEMORY_MANAGEMENT_SUPPORTED attribute. (#2380)

Signed-off-by: Tang, Jiajun [email protected]

Author: tangjj11

clang/lib/DPCT/CallExprRewriterCommon.h

@@ -993,8 +993,11 @@ inline std::function<bool(const CallExpr *C)>
 checkIsMigratedToHasAspect(size_t index) {
   return [=](const CallExpr *C) -> bool {
     if (getStmtSpelling(C->getArg(index))
-            .find("CU_DEVICE_ATTRIBUTE_CAN_MAP_HOST_MEMORY") !=
-        std::string::npos) {
+                .find("CU_DEVICE_ATTRIBUTE_CAN_MAP_HOST_MEMORY") !=
+            std::string::npos ||
+        getStmtSpelling(C->getArg(index))
+                .find("CU_DEVICE_ATTRIBUTE_VIRTUAL_MEMORY_MANAGEMENT_"
+                      "SUPPORTED") != std::string::npos) {
       return true;
     }
     return false;

clang/lib/DPCT/MapNames.cpp

@@ -1486,6 +1486,9 @@ void MapNames::setExplicitNamespaceMap(
       {"CU_DEVICE_ATTRIBUTE_MAX_BLOCK_DIM_Z",
        std::make_shared<EnumNameRule>("get_max_work_item_sizes",
                                       HelperFeatureEnum::device_ext)},
+      {"CU_DEVICE_ATTRIBUTE_VIRTUAL_MEMORY_MANAGEMENT_SUPPORTED",
+       std::make_shared<EnumNameRule>(
+           "has(sycl::aspect::ext_oneapi_virtual_mem)")},
       {"CU_CTX_MAP_HOST", std::make_shared<EnumNameRule>("0")},
       {"CU_CTX_SCHED_BLOCKING_SYNC", std::make_shared<EnumNameRule>("0")},
       {"CU_CTX_SCHED_SPIN", std::make_shared<EnumNameRule>("0")},

clang/test/dpct/driver_device_after11.2.cu

@@ -0,0 +1,15 @@
+// UNSUPPORTED: cuda-8.0, cuda-9.0, cuda-9.1, cuda-9.2, cuda-10.0, cuda-10.1, cuda-10.2, cuda-11.0, cuda-11.1
+// UNSUPPORTED: v8.0, v9.0, v9.1, v9.2, v10.0, v10.1, v10.2, v11.0, v11.1
+// RUN: dpct --format-range=none -out-root %T/driver_device_after11.2 %s --cuda-include-path="%cuda-path/include" -- -std=c++14 -x cuda --cuda-host-only
+// RUN: FileCheck %s --match-full-lines --input-file %T/driver_device_after11.2/driver_device_after11.2.dp.cpp
+// RUN: %if build_lit %{icpx -c -fsycl %T/driver_device_after11.2/driver_device_after11.2.dp.cpp -o %T/driver_device_after11.2/driver_device_after11.2.dp.o %}
+
+#include <cuda.h>
+
+int main() {
+  int i;
+  CUdevice d;
+  // CHECK: i = dpct::get_device(d).has(sycl::aspect::ext_oneapi_virtual_mem);
+  cuDeviceGetAttribute(&i, CU_DEVICE_ATTRIBUTE_VIRTUAL_MEMORY_MANAGEMENT_SUPPORTED, d);
+  return 0;
+}