From 7a4b7087d7856d04149ab85cf38f7907b71ec5b6 Mon Sep 17 00:00:00 2001
From: intwanghao <hao3.wang@intel.com>
Date: Thu, 12 Dec 2024 15:29:01 +0800
Subject: [PATCH] [SYCLomatic] Emit msg to suggest user use
 "--use-experimental-features=virtual_mem" to migrate virtual mem API  (#2562)

Signed-off-by: intwanghao <hao3.wang@intel.com>
---
 clang/lib/DPCT/RulesLang/RulesLang.cpp            |  7 +++++++
 clang/test/dpct/driver_device_after11.2.cu        |  8 ++++++--
 .../dpct/driver_device_virtual_mem_after11.2.cu   | 15 +++++++++++++++
 3 files changed, 28 insertions(+), 2 deletions(-)
 create mode 100644 clang/test/dpct/driver_device_virtual_mem_after11.2.cu

diff --git a/clang/lib/DPCT/RulesLang/RulesLang.cpp b/clang/lib/DPCT/RulesLang/RulesLang.cpp
index c2ee40561638..4f016f911617 100644
--- a/clang/lib/DPCT/RulesLang/RulesLang.cpp
+++ b/clang/lib/DPCT/RulesLang/RulesLang.cpp
@@ -8123,6 +8123,13 @@ void DriverDeviceAPIRule::runRule(
     auto SecArg = CE->getArg(1);
     if (auto DRE = dyn_cast<DeclRefExpr>(SecArg)) {
       auto AttributeName = DRE->getNameInfo().getAsString();
+      if (AttributeName ==
+              "CU_DEVICE_ATTRIBUTE_VIRTUAL_MEMORY_MANAGEMENT_SUPPORTED" &&
+          !DpctGlobalInfo::useExpVirtualMemory()) {
+        report(CE->getBeginLoc(), Diagnostics::TRY_EXPERIMENTAL_FEATURE, false,
+               AttributeName, "--use-experimental-features=virtual_mem");
+        return;
+      }
       auto Search = MapNames::EnumNamesMap.find(AttributeName);
       if (Search == MapNames::EnumNamesMap.end()) {
         report(CE->getBeginLoc(), Diagnostics::NOT_SUPPORTED_PARAMETER, false,
diff --git a/clang/test/dpct/driver_device_after11.2.cu b/clang/test/dpct/driver_device_after11.2.cu
index 0fee08ec7d2a..0e2cce5e5875 100644
--- a/clang/test/dpct/driver_device_after11.2.cu
+++ b/clang/test/dpct/driver_device_after11.2.cu
@@ -2,14 +2,18 @@
 // 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 %}
+// RUN: %if build_lit %{icpx -c -fsycl -DBUILD_TEST %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);
+#ifndef BUILD_TEST
+  // CHECK: /*
+  // CHECK-NEXT: DPCT1119:{{[0-9]+}}: Migration of CU_DEVICE_ATTRIBUTE_VIRTUAL_MEMORY_MANAGEMENT_SUPPORTED is not supported, please try to remigrate with option: --use-experimental-features=virtual_mem.
+  // CHECK-NEXT: */
   cuDeviceGetAttribute(&i, CU_DEVICE_ATTRIBUTE_VIRTUAL_MEMORY_MANAGEMENT_SUPPORTED, d);
+#endif
   return 0;
 }
diff --git a/clang/test/dpct/driver_device_virtual_mem_after11.2.cu b/clang/test/dpct/driver_device_virtual_mem_after11.2.cu
new file mode 100644
index 000000000000..02f4ff8fb45c
--- /dev/null
+++ b/clang/test/dpct/driver_device_virtual_mem_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 --use-experimental-features=virtual_mem -out-root %T/driver_device_virtual_mem_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_virtual_mem_after11.2/driver_device_virtual_mem_after11.2.dp.cpp
+// RUN: %if build_lit %{icpx -c -fsycl %T/driver_device_virtual_mem_after11.2/driver_device_virtual_mem_after11.2.dp.cpp -o %T/driver_device_virtual_mem_after11.2/driver_device_virtual_mem_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;
+}