[SYCLomatic] Added migration support for Semaphore APIs (#2695)

TejaX-Alaghari · web-flow · commit f5621e288c45 · 2025-03-26T13:04:59.000+08:00
This PR adds support for -

4 External Semaphore Resource APIs
cudaImportExternalSemaphore
cudaSignalExternalSemaphoresAsync
cudaWaitExternalSemaphoresAsync
cudaDestroyExternalSemaphore
4 Data Types
cudaExternalSemaphore_t
cudaExternalSemaphoreHandleDesc
cudaExternalSemaphoreSignalParams
cudaExternalSemaphoreWaitParams
3 Semaphore Types
opaqueFd
win32_nt_handle
win32_nt_dx12_resource
Supported Libraries
DirectX12
Vulkan
diff --git a/clang/lib/DPCT/Diagnostics/Diagnostics.inc b/clang/lib/DPCT/Diagnostics/Diagnostics.inc
@@ -296,8 +296,8 @@ DEF_WARNING(SPARSE_NNZ, 1134, MEDIUM_LEVEL, "The tool cannot deduce the consumer
 DEF_COMMENT(SPARSE_NNZ, 1134, MEDIUM_LEVEL, "The tool cannot deduce the consumer API (\"dpct::sparse::csrgemm\") of this API, and this API has 2 arguments depending on the 8th and the 12th parameters of the consumer API. Please replace the 2 arguments tagged as \"dpct_placeholder\" with the corresponding value.")
 DEF_WARNING(JOINT_MATRIX_SHAPE, 1135, HIGH_LEVEL, "Please check if joint_matrix implementations support the combination of data type and matrix shape type in the target hardware.")
 DEF_COMMENT(JOINT_MATRIX_SHAPE, 1135, HIGH_LEVEL, "Please check if joint_matrix implementations support the combination of data type and matrix shape type in the target hardware.")
-DEF_WARNING(UNSUPPORTED_EXTMEM_WIN_HANDLE, 1136, HIGH_LEVEL, "SYCL Bindless Images extension only supports importing external resource memory using NT handle on Windows. If assert(%0.get_win32_handle()) fails, you may need to adjust the code to use (%0.get_win32_handle()).")
-DEF_COMMENT(UNSUPPORTED_EXTMEM_WIN_HANDLE, 1136, HIGH_LEVEL, "SYCL Bindless Images extension only supports importing external resource memory using NT handle on Windows. If assert({0}.get_win32_handle()) fails, you may need to adjust the code to use ({0}.get_win32_handle()).")
+DEF_WARNING(UNSUPPORTED_EXTMEM_WIN_HANDLE, 1136, HIGH_LEVEL, "SYCL Bindless Images extension only supports importing external resource using NT handle on Windows. If assert(%0.get_win32_handle()) fails, you may need to adjust the code to use (%0.get_win32_handle()).")
+DEF_COMMENT(UNSUPPORTED_EXTMEM_WIN_HANDLE, 1136, HIGH_LEVEL, "SYCL Bindless Images extension only supports importing external resource using NT handle on Windows. If assert({0}.get_win32_handle()) fails, you may need to adjust the code to use ({0}.get_win32_handle()).")
 DEF_WARNING(ASYNC_COPY_DEVICE_WARN, 1137, LOW_LEVEL, "ASM instruction \"%0\" is asynchronous copy, current it is migrated to synchronous copy operation. You may need to adjust the code to tune the performance.")
 DEF_COMMENT(ASYNC_COPY_DEVICE_WARN, 1137, LOW_LEVEL, "ASM instruction \"{0}\" is asynchronous copy, current it is migrated to synchronous copy operation. You may need to adjust the code to tune the performance.")
 
diff --git a/clang/lib/DPCT/RuleInfra/APINamesTemplateType.inc b/clang/lib/DPCT/RuleInfra/APINamesTemplateType.inc
@@ -593,3 +593,60 @@ TYPE_REWRITE_ENTRY(
             WARNING_FACTORY(
                 Diagnostics::TRY_EXPERIMENTAL_FEATURE, TYPESTR,
                 STR("--use-experimental-features=bindless_images")))))
+
+// External Semaphore Handle Descs
+TYPE_REWRITE_ENTRY(
+    "cudaExternalSemaphore_t",
+    TYPE_CONDITIONAL_FACTORY(
+        [](const TypeLoc) { return DpctGlobalInfo::useSYCLCompat(); },
+        WARNING_FACTORY(Diagnostics::UNSUPPORT_SYCLCOMPAT,
+                        STR("cudaExternalSemaphore_t")),
+        TYPE_CONDITIONAL_FACTORY(
+            checkEnableBindlessImagesForType(),
+            TYPE_FACTORY(STR(MapNames::getDpctNamespace() +
+                             "experimental::external_sem_wrapper_ptr")),
+            WARNING_FACTORY(
+                Diagnostics::TRY_EXPERIMENTAL_FEATURE, TYPESTR,
+                STR("--use-experimental-features=bindless_images")))))
+
+TYPE_REWRITE_ENTRY(
+    "cudaExternalSemaphoreHandleDesc",
+    TYPE_CONDITIONAL_FACTORY(
+        [](const TypeLoc) { return DpctGlobalInfo::useSYCLCompat(); },
+        WARNING_FACTORY(Diagnostics::UNSUPPORT_SYCLCOMPAT,
+                        STR("cudaExternalSemaphoreHandleDesc")),
+        TYPE_CONDITIONAL_FACTORY(
+            checkEnableBindlessImagesForType(),
+            TYPE_FACTORY(STR(MapNames::getDpctNamespace() +
+                             "experimental::external_sem_handle_desc")),
+            WARNING_FACTORY(
+                Diagnostics::TRY_EXPERIMENTAL_FEATURE, TYPESTR,
+                STR("--use-experimental-features=bindless_images")))))
+
+TYPE_REWRITE_ENTRY(
+    "cudaExternalSemaphoreSignalParams",
+    TYPE_CONDITIONAL_FACTORY(
+        [](const TypeLoc) { return DpctGlobalInfo::useSYCLCompat(); },
+        WARNING_FACTORY(Diagnostics::UNSUPPORT_SYCLCOMPAT,
+                        STR("cudaExternalSemaphoreSignalParams")),
+        TYPE_CONDITIONAL_FACTORY(
+            checkEnableBindlessImagesForType(),
+            TYPE_FACTORY(STR(MapNames::getDpctNamespace() +
+                             "experimental::external_sem_params")),
+            WARNING_FACTORY(
+                Diagnostics::TRY_EXPERIMENTAL_FEATURE, TYPESTR,
+                STR("--use-experimental-features=bindless_images")))))
+
+TYPE_REWRITE_ENTRY(
+    "cudaExternalSemaphoreWaitParams",
+    TYPE_CONDITIONAL_FACTORY(
+        [](const TypeLoc) { return DpctGlobalInfo::useSYCLCompat(); },
+        WARNING_FACTORY(Diagnostics::UNSUPPORT_SYCLCOMPAT,
+                        STR("cudaExternalSemaphoreWaitParams")),
+        TYPE_CONDITIONAL_FACTORY(
+            checkEnableBindlessImagesForType(),
+            TYPE_FACTORY(STR(MapNames::getDpctNamespace() +
+                             "experimental::external_sem_params")),
+            WARNING_FACTORY(
+                Diagnostics::TRY_EXPERIMENTAL_FEATURE, TYPESTR,
+                STR("--use-experimental-features=bindless_images")))))
diff --git a/clang/lib/DPCT/RuleInfra/MapNames.cpp b/clang/lib/DPCT/RuleInfra/MapNames.cpp
@@ -260,6 +260,13 @@ void MapNames::setExplicitNamespaceMap(
       {"cudaExternalMemoryDedicated",
        MacroMigrationRule("dpct_build_in_macro_rule", RulePriority::Fallback,
                           "cudaExternalMemoryDedicated", "0")},
+      {"cudaExternalSemaphoreSignalSkipNvSciBufMemSync",
+       MacroMigrationRule("dpct_build_in_macro_rule", RulePriority::Fallback,
+                          "cudaExternalSemaphoreSignalSkipNvSciBufMemSync",
+                          "0")},
+      {"cudaExternalSemaphoreWaitSkipNvSciBufMemSync",
+       MacroMigrationRule("dpct_build_in_macro_rule", RulePriority::Fallback,
+                          "cudaExternalSemaphoreWaitSkipNvSciBufMemSync", "0")},
       {"cudaArrayDefault",
        MacroMigrationRule("dpct_build_in_macro_rule", RulePriority::Fallback,
                           "cudaArrayDefault",
@@ -331,7 +338,8 @@ void MapNames::setExplicitNamespaceMap(
       {"__half2", std::make_shared<TypeNameRule>(getClNamespace() + "half2")},
       {"half", std::make_shared<TypeNameRule>(getClNamespace() + "half")},
       {"half2", std::make_shared<TypeNameRule>(getClNamespace() + "half2")},
-      {"__nv_half2", std::make_shared<TypeNameRule>(getClNamespace() + "half2")},
+      {"__nv_half2",
+       std::make_shared<TypeNameRule>(getClNamespace() + "half2")},
       {"__nv_half", std::make_shared<TypeNameRule>(getClNamespace() + "half")},
       {"cudaEvent_t",
        std::make_shared<TypeNameRule>(getDpctNamespace() + "event_ptr",
@@ -886,6 +894,9 @@ void MapNames::setExplicitNamespaceMap(
       {"cudaExternalMemoryHandleType",
        std::make_shared<TypeNameRule>(getExpNamespace() +
                                       "external_mem_handle_type")},
+      {"cudaExternalSemaphoreHandleType",
+       std::make_shared<TypeNameRule>(getExpNamespace() +
+                                      "external_semaphore_handle_type")},
       // ...
   };
   // SYCLcompat unsupport types
@@ -1521,6 +1532,24 @@ void MapNames::setExplicitNamespaceMap(
                ? getExpNamespace() +
                      "external_mem_handle_type::win32_nt_dx12_resource"
                : "cudaExternalMemoryHandleTypeD3D12Resource")},
+      // enum cudaExternalSemaphoreHandleType
+      {"cudaExternalSemaphoreHandleTypeOpaqueFd",
+       std::make_shared<EnumNameRule>(
+           DpctGlobalInfo::useExtBindlessImages()
+               ? getExpNamespace() + "external_semaphore_handle_type::opaque_fd"
+               : "cudaExternalSemaphoreHandleTypeOpaqueFd")},
+      {"cudaExternalSemaphoreHandleTypeOpaqueWin32",
+       std::make_shared<EnumNameRule>(
+           DpctGlobalInfo::useExtBindlessImages()
+               ? getExpNamespace() +
+                     "external_semaphore_handle_type::win32_nt_handle"
+               : "cudaExternalSemaphoreHandleTypeOpaqueWin32")},
+      {"cudaExternalSemaphoreHandleTypeD3D12Fence",
+       std::make_shared<EnumNameRule>(
+           DpctGlobalInfo::useExtBindlessImages()
+               ? getExpNamespace() +
+                     "external_semaphore_handle_type::win32_nt_dx12_fence"
+               : "cudaExternalSemaphoreHandleTypeD3D12Fence")},
       // ...
   };
 
diff --git a/clang/lib/DPCT/RuleInfra/TypeLocRewriters.cpp b/clang/lib/DPCT/RuleInfra/TypeLocRewriters.cpp
@@ -360,6 +360,10 @@ void initTypeLocSYCLCompatRewriterMap(
   SYCLCOMPAT_UNSUPPORT("cudaExternalMemoryHandleDesc")
   SYCLCOMPAT_UNSUPPORT("cudaExternalMemoryMipmappedArrayDesc")
   SYCLCOMPAT_UNSUPPORT("cudaExternalMemoryBufferDesc")
+  SYCLCOMPAT_UNSUPPORT("cudaExternalSemaphore_t")
+  SYCLCOMPAT_UNSUPPORT("cudaExternalSemaphoreHandleDesc")
+  SYCLCOMPAT_UNSUPPORT("cudaExternalSemaphoreSignalParams")
+  SYCLCOMPAT_UNSUPPORT("cudaExternalSemaphoreWaitParams")
   SYCLCOMPAT_UNSUPPORT("thrust::system::cuda::experimental::pinned_allocator")
   SYCLCOMPAT_UNSUPPORT("thrust::cuda::experimental::pinned_allocator")
   SYCLCOMPAT_UNSUPPORT("thrust::device_allocator")
diff --git a/clang/lib/DPCT/RulesLang/APINamesGraphicsInterop.inc b/clang/lib/DPCT/RulesLang/APINamesGraphicsInterop.inc
@@ -184,4 +184,54 @@ CONDITIONAL_FACTORY_ENTRY(
     UNSUPPORT_FACTORY_ENTRY("cudaDestroyExternalMemory", Diagnostics::TRY_EXPERIMENTAL_FEATURE,
                             ARG("cudaDestroyExternalMemory"),
                             ARG("--use-experimental-features=bindless_images")))
+
+CONDITIONAL_FACTORY_ENTRY(
+    UseExtBindlessImages,
+    ASSIGNABLE_FACTORY(CALL_FACTORY_ENTRY("cudaImportExternalSemaphore",
+                           CALL(MapNames::getDpctNamespace() +
+                                    "experimental::import_external_semaphore",
+                                ARG(0), ARG(1)))),
+    UNSUPPORT_FACTORY_ENTRY("cudaImportExternalSemaphore", Diagnostics::TRY_EXPERIMENTAL_FEATURE,
+                            ARG("cudaImportExternalSemaphore"),
+                            ARG("--use-experimental-features=bindless_images")))
+
+CONDITIONAL_FACTORY_ENTRY(
+    UseExtBindlessImages,
+    CONDITIONAL_FACTORY_ENTRY(
+        makeCheckNot(CheckArgIsDefaultCudaStream(3)),
+        ASSIGNABLE_FACTORY(CALL_FACTORY_ENTRY("cudaSignalExternalSemaphoresAsync_v2",
+                           CALL(MapNames::getDpctNamespace() +
+                                    "experimental::signal_external_semaphore",
+                                ARG(0), ARG(1), ARG(2), ARG(3)))),
+        ASSIGNABLE_FACTORY(CALL_FACTORY_ENTRY("cudaSignalExternalSemaphoresAsync_v2",
+                           CALL(MapNames::getDpctNamespace() +
+                                    "experimental::signal_external_semaphore",
+                                ARG(0), ARG(1), ARG(2))))),
+    UNSUPPORT_FACTORY_ENTRY("cudaSignalExternalSemaphoresAsync_v2", Diagnostics::TRY_EXPERIMENTAL_FEATURE,
+                            ARG("cudaSignalExternalSemaphoresAsync_v2"),
+                            ARG("--use-experimental-features=bindless_images")))
+
+CONDITIONAL_FACTORY_ENTRY(
+    UseExtBindlessImages,
+    CONDITIONAL_FACTORY_ENTRY(
+        makeCheckNot(CheckArgIsDefaultCudaStream(3)),
+        ASSIGNABLE_FACTORY(CALL_FACTORY_ENTRY("cudaWaitExternalSemaphoresAsync_v2",
+                           CALL(MapNames::getDpctNamespace() +
+                                    "experimental::wait_external_semaphore",
+                                ARG(0), ARG(1), ARG(2), ARG(3)))),
+        ASSIGNABLE_FACTORY(CALL_FACTORY_ENTRY("cudaWaitExternalSemaphoresAsync_v2",
+                           CALL(MapNames::getDpctNamespace() +
+                                    "experimental::wait_external_semaphore",
+                                ARG(0), ARG(1), ARG(2))))),
+    UNSUPPORT_FACTORY_ENTRY("cudaWaitExternalSemaphoresAsync_v2", Diagnostics::TRY_EXPERIMENTAL_FEATURE,
+                            ARG("cudaWaitExternalSemaphoresAsync_v2"),
+                            ARG("--use-experimental-features=bindless_images")))
+
+ASSIGNABLE_FACTORY(CONDITIONAL_FACTORY_ENTRY(
+    UseExtBindlessImages,
+    DELETER_FACTORY_ENTRY("cudaDestroyExternalSemaphore", ARG(0)),
+    UNSUPPORT_FACTORY_ENTRY(
+        "cudaDestroyExternalSemaphore", Diagnostics::TRY_EXPERIMENTAL_FEATURE,
+        ARG("cudaDestroyExternalSemaphore"),
+        ARG("--use-experimental-features=bindless_images"))))
 // clang-format on
diff --git a/clang/lib/DPCT/RulesLang/MapNamesLang.cpp b/clang/lib/DPCT/RulesLang/MapNamesLang.cpp
@@ -235,13 +235,18 @@ MapNamesLang::MapTy TextureRule::ResourceTypeNames{
     {"format", "channel_type"},
     {"numChannels", "channel_num"}};
 
-// External memory handle descs fields mapping
-MapNamesLang::MapTy GraphicsInteropRule::ExtMemHandleDescNames{
+// External resource mem handle descs fields mapping
+MapNamesLang::MapTy GraphicsInteropRule::ExtResMemHandleDescNames{
     {"fd", "fd_handle"},        {"handle", "win32_handle"},
     {"name", "win32_obj_name"}, {"size", "res_size"},
     {"type", "handle_type"},    {"flags", "flags"},
     {"offset", "mem_offset"},   {"numLevels", "num_levels"},
-    {"extent", "size"},         {"formatDesc", "image_channel"}};
+    {"extent", "size"},         {"formatDesc", "image_channel"},
+    {"value", "value"}};
+
+// External resource sem param descs fields mapping
+MapNamesLang::MapTy GraphicsInteropRule::ExtResSemParamsNames{
+    {"value", "value"}};
 
 const MapNamesLang::MapTy MemoryDataTypeRule::PitchMemberNames{
     {"pitch", "pitch"}, {"ptr", "data_ptr"}, {"xsize", "x"}, {"ysize", "y"}};
diff --git a/clang/lib/DPCT/RulesLang/RulesLang.cpp b/clang/lib/DPCT/RulesLang/RulesLang.cpp
@@ -346,8 +346,8 @@ void TypeInDeclRule::registerMatcher(MatchFinder &MF) {
               "cublasLtMatmulHeuristicResult_t", "CUjit_target",
               "cublasLtMatrixTransformDesc_t", "cudaGraphicsMapFlags",
               "cudaGraphicsRegisterFlags", "cudaExternalMemoryHandleType",
-              "CUstreamCallback", "cudaHostFn_t", "__nv_half2",
-              "__nv_half"))))))
+              "cudaExternalSemaphoreHandleType", "CUstreamCallback",
+              "cudaHostFn_t", "__nv_half2", "__nv_half"))))))
           .bind("cudaTypeDef"),
       this);
 
@@ -360,7 +360,10 @@ void TypeInDeclRule::registerMatcher(MatchFinder &MF) {
                   "cudaGraphicsResource_t", "CUgraphicsResource",
                   "cudaExternalMemory_t", "cudaExternalMemoryHandleDesc",
                   "cudaExternalMemoryMipmappedArrayDesc",
-                  "cudaExternalMemoryBufferDesc"))))))
+                  "cudaExternalMemoryBufferDesc", "cudaExternalSemaphore_t",
+                  "cudaExternalSemaphoreHandleDesc",
+                  "cudaExternalSemaphoreSignalParams",
+                  "cudaExternalSemaphoreWaitParams"))))))
           .bind("cudaTypeDefEA"),
       this);
   MF.addMatcher(varDecl(hasType(classTemplateSpecializationDecl(
@@ -1981,7 +1984,14 @@ void EnumConstantRule::runRule(const MatchFinder::MatchResult &Result) {
       EnumName == "cudaExternalMemoryHandleTypeD3D12Heap" ||
       EnumName == "cudaExternalMemoryHandleTypeD3D11Resource" ||
       EnumName == "cudaExternalMemoryHandleTypeD3D11ResourceKmt" ||
-      EnumName == "cudaExternalMemoryHandleTypeNvSciBuf") {
+      EnumName == "cudaExternalMemoryHandleTypeNvSciBuf" ||
+      EnumName == "cudaExternalSemaphoreHandleTypeOpaqueWin32Kmt" ||
+      EnumName == "cudaExternalSemaphoreHandleTypeD3D11Fence" ||
+      EnumName == "cudaExternalSemaphoreHandleTypeNvSciSync" ||
+      EnumName == "cudaExternalSemaphoreHandleTypeKeyedMutex" ||
+      EnumName == "cudaExternalSemaphoreHandleTypeKeyedMutexKmt" ||
+      EnumName == "cudaExternalSemaphoreHandleTypeTimelineSemaphoreFd" ||
+      EnumName == "cudaExternalSemaphoreHandleTypeTimelineSemaphoreWin32") {
     report(E->getBeginLoc(), Diagnostics::API_NOT_MIGRATED, false, EnumName);
     return;
   } else if (EnumName == "cudaComputeModeDefault" ||
@@ -2007,7 +2017,10 @@ void EnumConstantRule::runRule(const MatchFinder::MatchResult &Result) {
               EnumName == "cudaGraphicsMapFlagsWriteDiscard" ||
               EnumName == "cudaExternalMemoryHandleTypeOpaqueFd" ||
               EnumName == "cudaExternalMemoryHandleTypeOpaqueWin32" ||
-              EnumName == "cudaExternalMemoryHandleTypeD3D12Resource")) {
+              EnumName == "cudaExternalMemoryHandleTypeD3D12Resource" ||
+              EnumName == "cudaExternalSemaphoreHandleTypeOpaqueFd" ||
+              EnumName == "cudaExternalSemaphoreHandleTypeOpaqueWin32" ||
+              EnumName == "cudaExternalSemaphoreHandleTypeD3D12Fence")) {
     report(E->getBeginLoc(), Diagnostics::TRY_EXPERIMENTAL_FEATURE, false,
            EnumName, "--use-experimental-features=bindless_images");
     return;
diff --git a/clang/lib/DPCT/RulesLang/RulesLang.h b/clang/lib/DPCT/RulesLang/RulesLang.h
@@ -1011,11 +1011,14 @@ class AssertRule : public NamedMigrationRule<AssertRule> {
 };
 
 class GraphicsInteropRule : public NamedMigrationRule<GraphicsInteropRule> {
-  static MapNames::MapTy ExtMemHandleDescNames;
+  static MapNames::MapTy ExtResMemHandleDescNames, ExtResSemParamsNames;
 
   const Expr *getAssignedBO(const Expr *E, ASTContext &Context);
   const Expr *getParentAsAssignedBO(const Expr *E, ASTContext &Context);
-  void replaceExtMemHandleDataExpr(const MemberExpr *ME, ASTContext &Context);
+  void replaceExtResMemHandleDataExpr(const MemberExpr *ME,
+                                      ASTContext &Context);
+  void replaceExtResSemParamsDataExpr(const MemberExpr *ME,
+                                      ASTContext &Context);
   inline const MemberExpr *getParentMemberExpr(const Stmt *S) {
     return DpctGlobalInfo::findParent<MemberExpr>(S);
   }
diff --git a/clang/lib/DPCT/RulesLang/RulesLangGraphicsInterop.cpp b/clang/lib/DPCT/RulesLang/RulesLangGraphicsInterop.cpp
diff --git a/clang/lib/DPCT/SrcAPI/APINames.inc b/clang/lib/DPCT/SrcAPI/APINames.inc
diff --git a/clang/runtime/dpct-rt/include/dpct/bindless_images.hpp b/clang/runtime/dpct-rt/include/dpct/bindless_images.hpp
diff --git a/clang/test/dpct/externalResInterop.cu b/clang/test/dpct/externalResInterop.cu
