[SYCLomatic] Refine migration of cudaMemGetInfo() (#228)

1. Add a new value "device_info" for option “--no-dpcpp-extensions" to guide migration of cudaMemGetInfo(). 2. Add a new helper function get_memory_info() in helper file device.hpp to get the number of bytes of free and total memory of the SYCL device Signed-off-by: chenwei.sun <chenwei.sun@intel.com>
oneapi-src · Nov 17, 2022 · af325b6 · af325b6
1 parent a0758d6
commit af325b6
Show file tree

Hide file tree

Showing 13 changed files with 153 additions and 41 deletions.
diff --git a/clang/include/clang/DPCT/DPCTOptions.inc b/clang/include/clang/DPCT/DPCTOptions.inc
@@ -320,7 +320,11 @@ DPCT_ENUM_OPTION(DPCT_OPT_TYPE(static llvm::cl::list<DPCPPExtensionsDefaultEnabl
              DPCT_OPT_ENUM("enqueued_barriers",
                                      int(DPCPPExtensionsDefaultEnabled::ExtDE_EnqueueBarrier),
                                      "Enqueued barriers extension.",
-                                     false)
+                                     false),
+             DPCT_OPT_ENUM("device_info",
+                                     int(DPCPPExtensionsDefaultEnabled::ExtDE_DeviceInfo),
+                                     "Intel's Extensions for Device Information.",
+                                     false)                            
            ),
            llvm::cl::desc("Comma separated list of extensions not to be used in migrated "
            "code.\n"

diff --git a/clang/lib/DPCT/APINames.inc b/clang/lib/DPCT/APINames.inc
@@ -154,7 +154,7 @@ ENTRY(cudaMallocManaged, cudaMallocManaged, true, NO_FLAG, P0, "Successful")
 ENTRY(cudaMallocMipmappedArray, cudaMallocMipmappedArray, false, NO_FLAG, P0, "comment")
 ENTRY(cudaMallocPitch, cudaMallocPitch, true, NO_FLAG, P0, "Successful")
 ENTRY(cudaMemAdvise, cudaMemAdvise, true, NO_FLAG, P0, "Successful for restricted USM")
-ENTRY(cudaMemGetInfo, cudaMemGetInfo, true, API_CALL_UNSUPPORTED, P0, "Migration rate 0.5, emit DPCT1072 for unsupported part.")
+ENTRY(cudaMemGetInfo, cudaMemGetInfo, true, API_CALL_UNSUPPORTED, P0, "Successful")
 ENTRY(cudaMemPrefetchAsync, cudaMemPrefetchAsync, true, NO_FLAG, P0, "Successful: USM only")
 ENTRY(cudaMemRangeGetAttribute, cudaMemRangeGetAttribute, false, NO_FLAG, P4, "comment")
 ENTRY(cudaMemRangeGetAttributes, cudaMemRangeGetAttributes, false, NO_FLAG, P4, "comment")

diff --git a/clang/lib/DPCT/ASTTraversal.cpp b/clang/lib/DPCT/ASTTraversal.cpp
@@ -11516,47 +11516,68 @@ void MemoryMigrationRule::miscMigration(const MatchFinder::MatchResult &Result,
     emplaceTransformation(new ReplaceStmt(C, OS.str()));
     requestFeature(HelperFeatureEnum::Image_image_matrix_get_channel, C);
   } else if (Name == "cuMemGetInfo_v2" || Name == "cudaMemGetInfo") {
-    auto &SM = DpctGlobalInfo::getSourceManager();
-    std::ostringstream OS;
-    if (IsAssigned)
+    if (DpctGlobalInfo::useDeviceInfo()) {
+      std::ostringstream OS;
+      if (IsAssigned)
+        OS << "(";
+      OS << MapNames::getDpctNamespace() + "get_current_device().get_memory_info";
       OS << "(";
-    auto SecArg = C->getArg(1);
-    printDerefOp(OS, SecArg);
-    OS << " = " << MapNames::getDpctNamespace()
-       << "get_current_device().get_device_info()"
-          ".get_global_mem_size()";
-    requestFeature(HelperFeatureEnum::Device_get_current_device, C);
-    requestFeature(
-        HelperFeatureEnum::Device_device_ext_get_device_info_return_info, C);
-    requestFeature(HelperFeatureEnum::Device_device_info_get_global_mem_size,
-                   C);
-    if (IsAssigned) {
-      OS << ", 0)";
-      report(C->getBeginLoc(), Diagnostics::NOERROR_RETURN_COMMA_OP, false);
-    }
-    SourceLocation CallBegin(C->getBeginLoc());
-    SourceLocation CallEnd(C->getEndLoc());
+      printDerefOp(OS, C->getArg(0));
+      OS << ", ";
+      printDerefOp(OS, C->getArg(1));
+      OS << ")";
 
-    bool IsMacroArg =
-        SM.isMacroArgExpansion(CallBegin) && SM.isMacroArgExpansion(CallEnd);
+      emplaceTransformation(new ReplaceStmt(C, OS.str()));
+      if (IsAssigned) {
+        OS << ", 0)";
+        report(C->getBeginLoc(), Diagnostics::NOERROR_RETURN_COMMA_OP, false);
+      }
+      emplaceTransformation(new ReplaceStmt(C, OS.str()));
+      requestFeature(HelperFeatureEnum::Device_device_ext_get_memory_info, C);
+    } else {
+      auto &SM = DpctGlobalInfo::getSourceManager();
+      std::ostringstream OS;
+      if (IsAssigned)
+        OS << "(";
 
-    if (CallBegin.isMacroID() && IsMacroArg) {
-      CallBegin = SM.getImmediateSpellingLoc(CallBegin);
-      CallBegin = SM.getExpansionLoc(CallBegin);
-    } else if (CallBegin.isMacroID()) {
-      CallBegin = SM.getExpansionLoc(CallBegin);
-    }
+      auto SecondArg = C->getArg(1);
+      printDerefOp(OS, SecondArg);
+      OS << " = " << MapNames::getDpctNamespace()
+         << "get_current_device().get_device_info()"
+            ".get_global_mem_size()";
+      requestFeature(HelperFeatureEnum::Device_get_current_device, C);
+      requestFeature(
+          HelperFeatureEnum::Device_device_ext_get_device_info_return_info, C);
+      requestFeature(HelperFeatureEnum::Device_device_info_get_global_mem_size,
+                     C);
+      if (IsAssigned) {
+        OS << ", 0)";
+        report(C->getBeginLoc(), Diagnostics::NOERROR_RETURN_COMMA_OP, false);
+      }
+      SourceLocation CallBegin(C->getBeginLoc());
+      SourceLocation CallEnd(C->getEndLoc());
 
-    if (CallEnd.isMacroID() && IsMacroArg) {
-      CallEnd = SM.getImmediateSpellingLoc(CallEnd);
-      CallEnd = SM.getExpansionLoc(CallEnd);
-    } else if (CallEnd.isMacroID()) {
-      CallEnd = SM.getExpansionLoc(CallEnd);
-    }
-    CallEnd = CallEnd.getLocWithOffset(1);
+      bool IsMacroArg =
+          SM.isMacroArgExpansion(CallBegin) && SM.isMacroArgExpansion(CallEnd);
 
-    emplaceTransformation(replaceText(CallBegin, CallEnd, OS.str(), SM));
-    report(C->getBeginLoc(), Diagnostics::UNSUPPORT_FREE_MEMORY_SIZE, false);
+      if (CallBegin.isMacroID() && IsMacroArg) {
+        CallBegin = SM.getImmediateSpellingLoc(CallBegin);
+        CallBegin = SM.getExpansionLoc(CallBegin);
+      } else if (CallBegin.isMacroID()) {
+        CallBegin = SM.getExpansionLoc(CallBegin);
+      }
+
+      if (CallEnd.isMacroID() && IsMacroArg) {
+        CallEnd = SM.getImmediateSpellingLoc(CallEnd);
+        CallEnd = SM.getExpansionLoc(CallEnd);
+      } else if (CallEnd.isMacroID()) {
+        CallEnd = SM.getExpansionLoc(CallEnd);
+      }
+      CallEnd = CallEnd.getLocWithOffset(1);
+
+      emplaceTransformation(replaceText(CallBegin, CallEnd, OS.str(), SM));
+      report(C->getBeginLoc(), Diagnostics::UNSUPPORT_FREE_MEMORY_SIZE, false);
+    }
   } else {
     auto Itr = CallExprRewriterFactoryBase::RewriterMap->find(Name);
     if (Itr != CallExprRewriterFactoryBase::RewriterMap->end()) {

diff --git a/clang/lib/DPCT/AnalysisInfo.h b/clang/lib/DPCT/AnalysisInfo.h
@@ -1864,6 +1864,10 @@ class DpctGlobalInfo {
     return getUsingExtensionDD(DPCPPExtensionsDefaultDisabled::ExtDD_CCXXStandardLibrary);
   }
 
+  static bool useDeviceInfo() {
+    return getUsingExtensionDE(DPCPPExtensionsDefaultEnabled::ExtDE_DeviceInfo);
+  }
+
   static bool getSpBLASUnsupportedMatrixTypeFlag() {
     return SpBLASUnsupportedMatrixTypeFlag;
   }

diff --git a/clang/lib/DPCT/IncrementalMigrationUtility.cpp b/clang/lib/DPCT/IncrementalMigrationUtility.cpp
@@ -199,6 +199,9 @@ bool printOptions(
         if (!(UValue &
               static_cast<unsigned>(DPCPPExtensionsDefaultEnabled::ExtDE_EnqueueBarrier)))
           Str = Str + "enqueued_barriers,";
+        if (!(UValue &
+              static_cast<unsigned>(DPCPPExtensionsDefaultEnabled::ExtDE_DeviceInfo)))
+          Str += "device_info,";
       }
       if (!Str.empty()) {
         Str = "--no-dpcpp-extensions=" + Str;

diff --git a/clang/lib/DPCT/ValidateArguments.h b/clang/lib/DPCT/ValidateArguments.h
@@ -58,6 +58,7 @@ enum class ExplicitNamespace : unsigned int {
 };
 enum class DPCPPExtensionsDefaultEnabled : unsigned int {
   ExtDE_EnqueueBarrier = 0x01,
+  ExtDE_DeviceInfo = 0x02,
   ExtDE_DPCPPExtensionsDefaultEnabledEnumSize
 };
 enum class DPCPPExtensionsDefaultDisabled : unsigned int {

diff --git a/clang/runtime/dpct-rt/include/device.hpp.inc b/clang/runtime/dpct-rt/include/device.hpp.inc
@@ -547,6 +547,38 @@ public:
   int get_integrated() const { return get_device_info().get_integrated(); }
 // DPCT_LABEL_END
 
+// DPCT_LABEL_BEGIN|device_ext_get_memory_info|dpct
+// DPCT_PARENT_FEATURE|device_ext
+// DPCT_DEPENDENCY_BEGIN
+// Device|device_ext
+// Device|device_ext_get_device_info_return_info
+// Device|device_info_get_global_mem_size
+// DPCT_DEPENDENCY_END
+// DPCT_CODE
+  /// Get the number of bytes of free and total memory on the SYCL device.
+  /// \param [out] free_memory The number of bytes of free memory on the SYCL device.
+  /// \param [out] total_memory The number of bytes of total memory on the SYCL device.
+  void get_memory_info(size_t &free_memory, size_t &total_memory) {
+#if (defined(__SYCL_COMPILER_VERSION) && __SYCL_COMPILER_VERSION >= 20221105)
+    if (!has(sycl::aspect::ext_intel_free_memory)) {
+      std::cerr << "get_memory_info: ext_intel_free_memory is not supported." << std::endl;
+      free_memory = 0;
+    } else {
+      free_memory = get_info<sycl::ext::intel::info::device::free_memory>();
+    }
+#else
+    std::cerr << "get_memory_info: ext_intel_free_memory is not supported." << std::endl;
+    free_memory = 0;
+#if defined(_MSC_VER) && !defined(__clang__)
+#pragma message("Querying the number of bytes of free memory is not supported")
+#else
+#warning "Querying the number of bytes of free memory is not supported"
+#endif
+#endif
+    total_memory = get_device_info().get_global_mem_size();
+  }
+// DPCT_LABEL_END
+
 // DPCT_LABEL_BEGIN|device_ext_get_device_info_return_void|dpct
 // DPCT_PARENT_FEATURE|device_ext
 // DPCT_DEPENDENCY_BEGIN

diff --git a/clang/test/dpct/helper_files_ref/include/device.hpp b/clang/test/dpct/helper_files_ref/include/device.hpp
@@ -229,6 +229,29 @@ class device_ext : public sycl::device {
 
   int get_integrated() const { return get_device_info().get_integrated(); }
 
+  /// Get the number of bytes of free and total memory on the SYCL device.
+  /// \param [out] free_memory The number of bytes of free memory on the SYCL device.
+  /// \param [out] total_memory The number of bytes of total memory on the SYCL device.
+  void get_memory_info(size_t &free_memory, size_t &total_memory) {
+#if (defined(__SYCL_COMPILER_VERSION) && __SYCL_COMPILER_VERSION >= 20221105)
+    if (!has(sycl::aspect::ext_intel_free_memory)) {
+      std::cerr << "get_memory_info: ext_intel_free_memory is not supported." << std::endl;
+      free_memory = 0;
+    } else {
+      free_memory = get_info<sycl::ext::intel::info::device::free_memory>();
+    }
+#else
+    std::cerr << "get_memory_info: ext_intel_free_memory is not supported." << std::endl;
+    free_memory = 0;
+#if defined(_MSC_VER) && !defined(__clang__)
+#pragma message("Querying the number of bytes of free memory is not supported")
+#else
+#warning "Querying the number of bytes of free memory is not supported"
+#endif
+#endif
+    total_memory = get_device_info().get_global_mem_size();
+  }
+
   void get_device_info(device_info &out) const {
     device_info prop;
     prop.set_name(get_info<sycl::info::device::name>().c_str());

diff --git a/clang/test/dpct/memory_management.cu b/clang/test/dpct/memory_management.cu
@@ -1,6 +1,6 @@
 // FIXME
 // UNSUPPORTED: -windows-
-// RUN: dpct --format-range=none --usm-level=none -out-root %T/memory_management %s --cuda-include-path="%cuda-path/include" -output-file=memory_management_outputfile.txt -- -x cuda --cuda-host-only
+// RUN: dpct --no-dpcpp-extensions=device_info --format-range=none --usm-level=none -out-root %T/memory_management %s --cuda-include-path="%cuda-path/include" -output-file=memory_management_outputfile.txt -- -x cuda --cuda-host-only
 // RUN: FileCheck --match-full-lines --input-file %T/memory_management/memory_management.dp.cpp %s
 
 #include <cuda_runtime.h>

diff --git a/clang/test/dpct/memory_management_restricted.cu b/clang/test/dpct/memory_management_restricted.cu
@@ -377,6 +377,17 @@ void foobar() {
 
   //CHECK: MY_ERROR_CHECKER((d_Output = sycl::malloc_device<float>(1, q_ct1), 0));
   MY_ERROR_CHECKER(cudaMalloc((void **)&d_Output, sizeof(float)));
+
+  size_t free_mem, total_mem;
+
+  //CHECK: dpct::get_current_device().get_memory_info(free_mem, total_mem);
+  cudaMemGetInfo(&free_mem, &total_mem);
+  //CHECK: MY_ERROR_CHECKER((dpct::get_current_device().get_memory_info(free_mem, total_mem), 0));
+  MY_ERROR_CHECKER(cudaMemGetInfo(&free_mem, &total_mem));
+  //CHECK: dpct::get_current_device().get_memory_info(free_mem, total_mem);
+  cuMemGetInfo(&free_mem, &total_mem);
+  //CHECK: MY_ERROR_CHECKER((dpct::get_current_device().get_memory_info(free_mem, total_mem), 0));
+  MY_ERROR_CHECKER(cuMemGetInfo(&free_mem, &total_mem));
 }
 
 template <typename T>

diff --git a/clang/test/dpct/test_api_level/Device/api_test13.cu b/clang/test/dpct/test_api_level/Device/api_test13.cu
@@ -1,4 +1,4 @@
-// RUN: dpct --format-range=none  --usm-level=none  --use-custom-helper=api -out-root %T/Device/api_test13_out %s --cuda-include-path="%cuda-path/include" -- -x cuda --cuda-host-only
+// RUN: dpct --format-range=none --no-dpcpp-extensions=device_info --usm-level=none  --use-custom-helper=api -out-root %T/Device/api_test13_out %s --cuda-include-path="%cuda-path/include" -- -x cuda --cuda-host-only
 // RUN: grep "IsCalled" %T/Device/api_test13_out/MainSourceFiles.yaml | wc -l > %T/Device/api_test13_out/count.txt
 // RUN: FileCheck --input-file %T/Device/api_test13_out/count.txt --match-full-lines %s
 // RUN: rm -rf %T/Device/api_test13_out

diff --git a/clang/test/dpct/test_api_level/Device/api_test29.cu b/clang/test/dpct/test_api_level/Device/api_test29.cu
@@ -1,4 +1,4 @@
-// RUN: dpct --format-range=none  --usm-level=none  --use-custom-helper=api -out-root %T/Device/api_test29_out %s --cuda-include-path="%cuda-path/include" -- -x cuda --cuda-host-only
+// RUN: dpct --format-range=none --usm-level=none  --use-custom-helper=api -out-root %T/Device/api_test29_out %s --cuda-include-path="%cuda-path/include" -- -x cuda --cuda-host-only
 // RUN: grep "IsCalled" %T/Device/api_test29_out/MainSourceFiles.yaml | wc -l > %T/Device/api_test29_out/count.txt
 // RUN: FileCheck --input-file %T/Device/api_test29_out/count.txt --match-full-lines %s
 // RUN: rm -rf %T/Device/api_test29_out

diff --git a/clang/test/dpct/test_api_level/Device/api_test31.cu b/clang/test/dpct/test_api_level/Device/api_test31.cu
@@ -0,0 +1,13 @@
+// RUN: dpct --format-range=none --usm-level=none  --use-custom-helper=api -out-root %T/Device/api_test31_out %s --cuda-include-path="%cuda-path/include" -- -x cuda --cuda-host-only
+// RUN: grep "IsCalled" %T/Device/api_test31_out/MainSourceFiles.yaml | wc -l > %T/Device/api_test31_out/count.txt
+// RUN: FileCheck --input-file %T/Device/api_test31_out/count.txt --match-full-lines %s
+// RUN: rm -rf %T/Device/api_test31_out
+
+// CHECK: 25
+// TEST_FEATURE: Device_device_ext_get_memory_info
+
+int main() {
+  size_t result1, result2;
+  cuMemGetInfo(&result1, &result2);
+  return 0;
+}