From af325b6b397d4e219293c0e78e29cbfa379b9f18 Mon Sep 17 00:00:00 2001 From: tomflinda Date: Thu, 17 Nov 2022 11:52:54 +0800 Subject: [PATCH] [SYCLomatic] Refine migration of cudaMemGetInfo() (#228) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit 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 --- clang/include/clang/DPCT/DPCTOptions.inc | 6 +- clang/lib/DPCT/APINames.inc | 2 +- clang/lib/DPCT/ASTTraversal.cpp | 93 ++++++++++++------- clang/lib/DPCT/AnalysisInfo.h | 4 + .../lib/DPCT/IncrementalMigrationUtility.cpp | 3 + clang/lib/DPCT/ValidateArguments.h | 1 + clang/runtime/dpct-rt/include/device.hpp.inc | 32 +++++++ .../dpct/helper_files_ref/include/device.hpp | 23 +++++ clang/test/dpct/memory_management.cu | 2 +- .../test/dpct/memory_management_restricted.cu | 11 +++ .../dpct/test_api_level/Device/api_test13.cu | 2 +- .../dpct/test_api_level/Device/api_test29.cu | 2 +- .../dpct/test_api_level/Device/api_test31.cu | 13 +++ 13 files changed, 153 insertions(+), 41 deletions(-) create mode 100644 clang/test/dpct/test_api_level/Device/api_test31.cu diff --git a/clang/include/clang/DPCT/DPCTOptions.inc b/clang/include/clang/DPCT/DPCTOptions.inc index 069a1ff53d10..c78fe9785167 100644 --- 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::listgetArg(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 index a67258ddd5c2..5ce697eb1e07 100644 --- 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 index f5ba12837d8d..4d39f1c78d2b 100644 --- a/clang/lib/DPCT/IncrementalMigrationUtility.cpp +++ b/clang/lib/DPCT/IncrementalMigrationUtility.cpp @@ -199,6 +199,9 @@ bool printOptions( if (!(UValue & static_cast(DPCPPExtensionsDefaultEnabled::ExtDE_EnqueueBarrier))) Str = Str + "enqueued_barriers,"; + if (!(UValue & + static_cast(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 index c7fb2f092eab..1bf89534943b 100644 --- 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 index 551fae0e8f63..ef4d7f78918b 100644 --- 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(); + } +#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 index aa35311f9567..76c4ac12af24 100644 --- 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(); + } +#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().c_str()); diff --git a/clang/test/dpct/memory_management.cu b/clang/test/dpct/memory_management.cu index 3da1163fab97..5423e7cd04e4 100644 --- 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 diff --git a/clang/test/dpct/memory_management_restricted.cu b/clang/test/dpct/memory_management_restricted.cu index f69556630622..1c5699977cf2 100644 --- 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(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 diff --git a/clang/test/dpct/test_api_level/Device/api_test13.cu b/clang/test/dpct/test_api_level/Device/api_test13.cu index 12f6de197f8b..fac0cd4bab63 100644 --- 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 index cb787eb43a51..5aed853f87bf 100644 --- 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 new file mode 100644 index 000000000000..88ea2af7868d --- /dev/null +++ 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; +}