From f6f7c51c268b1e64d8fd25b644c94f939433ffd6 Mon Sep 17 00:00:00 2001 From: "chenwei.sun" Date: Thu, 3 Jul 2025 14:06:41 +0800 Subject: [PATCH 1/2] [SYCLomatic] Add query-api-mapping for 1 runtime API `cudaEventRecordWithFlags` Signed-off-by: chenwei.sun --- .../DPCT/Runtime/cudaEventRecordWithFlags.cu | 6 ++++++ clang/lib/DPCT/RulesLang/RulesLang.cpp | 13 +++++++++++++ clang/test/dpct/query_api_mapping/Runtime/test.cu | 8 ++++++-- clang/test/dpct/query_api_mapping/test_all.cu | 1 + 4 files changed, 26 insertions(+), 2 deletions(-) create mode 100644 clang/examples/DPCT/Runtime/cudaEventRecordWithFlags.cu diff --git a/clang/examples/DPCT/Runtime/cudaEventRecordWithFlags.cu b/clang/examples/DPCT/Runtime/cudaEventRecordWithFlags.cu new file mode 100644 index 000000000000..a45a6544d7cf --- /dev/null +++ b/clang/examples/DPCT/Runtime/cudaEventRecordWithFlags.cu @@ -0,0 +1,6 @@ +void test(cudaEvent_t event, cudaStream_t stream) { + // Start + cudaEventRecordWithFlags(event /*cudaEvent_t*/, stream /*cudaStream_t*/, + cudaEventRecordDefault /*unsigned int*/); + // End +} diff --git a/clang/lib/DPCT/RulesLang/RulesLang.cpp b/clang/lib/DPCT/RulesLang/RulesLang.cpp index 99b55946bf96..4cf0cca92245 100644 --- a/clang/lib/DPCT/RulesLang/RulesLang.cpp +++ b/clang/lib/DPCT/RulesLang/RulesLang.cpp @@ -3276,6 +3276,19 @@ void EventAPICallRule::handleEventRecordWithProfilingDisabled( static std::set> DeclDupFilter; auto &SM = DpctGlobalInfo::getSourceManager(); + int NumArgs = CE->getNumArgs(); + if (NumArgs == 3) { // Special process for cudaEventRecordWithFlags(). + auto APIName = CE->getDirectCallee()->getNameInfo().getName().getAsString(); + const Expr *SecArg = CE->getArg(2); + ExprAnalysis Arg2EA(SecArg); + auto Arg2Name = Arg2EA.getReplacedString(); + if (Arg2Name != "cudaEventRecordDefault") { + report(CE->getBeginLoc(), Diagnostics::NOT_SUPPORTED_PARAMETER, false, + APIName, "parameter " + Arg2Name + " is unsupported"); + return; + } + } + const ValueDecl *MD = nullptr; if ((MD = getDecl(CE->getArg(0))) == nullptr) return; diff --git a/clang/test/dpct/query_api_mapping/Runtime/test.cu b/clang/test/dpct/query_api_mapping/Runtime/test.cu index 1664739882ab..112429e95a0b 100644 --- a/clang/test/dpct/query_api_mapping/Runtime/test.cu +++ b/clang/test/dpct/query_api_mapping/Runtime/test.cu @@ -520,8 +520,6 @@ // CUDARUNTIMEGETVERSION-NEXT: Is migrated to: // CUDARUNTIMEGETVERSION-NEXT: *pi = dpct::get_major_version(dpct::get_current_device()); - - // RUN: dpct --cuda-include-path="%cuda-path/include" -query-api-mapping=cudaMemcpy3DPeer | FileCheck %s -check-prefix=CUDAMEMCPY3DPEER // CUDAMEMCPY3DPEER: CUDA API: // CUDAMEMCPY3DPEER-NEXT: cudaMemcpy3DPeer(p/*const cudaMemcpy3DPeerParms**/); @@ -557,3 +555,9 @@ // CUDAPROFILERSTOP-NEXT: cudaProfilerStop(); // CUDAPROFILERSTOP-NEXT: The API is Removed. // CUDAPROFILERSTOP-EMPTY: +// RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=cudaEventRecordWithFlags | FileCheck %s -check-prefix=cudaEventRecordWithFlags +// cudaEventRecordWithFlags: CUDA API: +// cudaEventRecordWithFlags-NEXT: cudaEventRecordWithFlags(event /*cudaEvent_t*/, stream /*cudaStream_t*/, +// cudaEventRecordWithFlags-NEXT: cudaEventRecordDefault /*unsigned int*/); +// cudaEventRecordWithFlags-NEXT: Is migrated to: +// cudaEventRecordWithFlags-NEXT: *event = dpct::get_in_order_queue().ext_oneapi_submit_barrier(); diff --git a/clang/test/dpct/query_api_mapping/test_all.cu b/clang/test/dpct/query_api_mapping/test_all.cu index 6352b08ca822..d1af2a07f629 100644 --- a/clang/test/dpct/query_api_mapping/test_all.cu +++ b/clang/test/dpct/query_api_mapping/test_all.cu @@ -1447,6 +1447,7 @@ // CHECK-NEXT: cudaEventElapsedTime // CHECK-NEXT: cudaEventQuery // CHECK-NEXT: cudaEventRecord +// CHECK-NEXT: cudaEventRecordWithFlags // CHECK-NEXT: cudaEventSynchronize // CHECK-NEXT: cudaExternalMemoryGetMappedBuffer // CHECK-NEXT: cudaExternalMemoryGetMappedMipmappedArray From 6bc1d808617a2e7773fb37ce13ad21ca5bc8eda5 Mon Sep 17 00:00:00 2001 From: "chenwei.sun" Date: Fri, 4 Jul 2025 10:01:35 +0800 Subject: [PATCH 2/2] Refine migration command with profiling enabled Signed-off-by: chenwei.sun --- .../DPCT/Runtime/cudaEventRecordWithFlags.cu | 2 ++ clang/lib/DPCT/DPCT.cpp | 2 ++ clang/lib/DPCT/RulesLang/RulesLang.cpp | 13 ------------- clang/test/dpct/query_api_mapping/Runtime/test.cu | 4 ++-- 4 files changed, 6 insertions(+), 15 deletions(-) diff --git a/clang/examples/DPCT/Runtime/cudaEventRecordWithFlags.cu b/clang/examples/DPCT/Runtime/cudaEventRecordWithFlags.cu index a45a6544d7cf..501a9f738bf4 100644 --- a/clang/examples/DPCT/Runtime/cudaEventRecordWithFlags.cu +++ b/clang/examples/DPCT/Runtime/cudaEventRecordWithFlags.cu @@ -1,3 +1,5 @@ +// Option: --enable-profiling + void test(cudaEvent_t event, cudaStream_t stream) { // Start cudaEventRecordWithFlags(event /*cudaEvent_t*/, stream /*cudaStream_t*/, diff --git a/clang/lib/DPCT/DPCT.cpp b/clang/lib/DPCT/DPCT.cpp index 2f06c278c1ea..3d4df482bc8f 100644 --- a/clang/lib/DPCT/DPCT.cpp +++ b/clang/lib/DPCT/DPCT.cpp @@ -1072,6 +1072,8 @@ int runDPCT(int argc, const char **argv) { Experimentals.addValue(ExperimentalFeatures::Exp_NonUniformGroups); } else if (Option == "--no-dry-pattern") { NoDRYPattern.setValue(true); + } else if (Option == "--enable-profiling") { + EnablepProfiling.setValue(true); } // Need add more option. } diff --git a/clang/lib/DPCT/RulesLang/RulesLang.cpp b/clang/lib/DPCT/RulesLang/RulesLang.cpp index 4cf0cca92245..99b55946bf96 100644 --- a/clang/lib/DPCT/RulesLang/RulesLang.cpp +++ b/clang/lib/DPCT/RulesLang/RulesLang.cpp @@ -3276,19 +3276,6 @@ void EventAPICallRule::handleEventRecordWithProfilingDisabled( static std::set> DeclDupFilter; auto &SM = DpctGlobalInfo::getSourceManager(); - int NumArgs = CE->getNumArgs(); - if (NumArgs == 3) { // Special process for cudaEventRecordWithFlags(). - auto APIName = CE->getDirectCallee()->getNameInfo().getName().getAsString(); - const Expr *SecArg = CE->getArg(2); - ExprAnalysis Arg2EA(SecArg); - auto Arg2Name = Arg2EA.getReplacedString(); - if (Arg2Name != "cudaEventRecordDefault") { - report(CE->getBeginLoc(), Diagnostics::NOT_SUPPORTED_PARAMETER, false, - APIName, "parameter " + Arg2Name + " is unsupported"); - return; - } - } - const ValueDecl *MD = nullptr; if ((MD = getDecl(CE->getArg(0))) == nullptr) return; diff --git a/clang/test/dpct/query_api_mapping/Runtime/test.cu b/clang/test/dpct/query_api_mapping/Runtime/test.cu index 112429e95a0b..6db94acbb8a7 100644 --- a/clang/test/dpct/query_api_mapping/Runtime/test.cu +++ b/clang/test/dpct/query_api_mapping/Runtime/test.cu @@ -559,5 +559,5 @@ // cudaEventRecordWithFlags: CUDA API: // cudaEventRecordWithFlags-NEXT: cudaEventRecordWithFlags(event /*cudaEvent_t*/, stream /*cudaStream_t*/, // cudaEventRecordWithFlags-NEXT: cudaEventRecordDefault /*unsigned int*/); -// cudaEventRecordWithFlags-NEXT: Is migrated to: -// cudaEventRecordWithFlags-NEXT: *event = dpct::get_in_order_queue().ext_oneapi_submit_barrier(); +// cudaEventRecordWithFlags-NEXT: Is migrated to (with the option --enable-profiling): +// cudaEventRecordWithFlags-NEXT: dpct::sync_barrier(event /*cudaEvent_t*/, stream /*unsigned int*/);