[SYCLomatic] Support migration of 3 runtime Surface APIs.

Signed-off-by: Tang, Jiajun [email protected]

Author: tangjj11

clang/lib/DPCT/RuleInfra/MapNames.cpp

@@ -554,6 +554,13 @@ void MapNames::setExplicitNamespaceMap(
                      "ext::oneapi::experimental::sampled_image_handle"
                : getDpctNamespace() + "image_wrapper_base_p",
            HelperFeatureEnum::device_ext)},
+      {"cudaSurfaceObject_t",
+       std::make_shared<TypeNameRule>(
+           DpctGlobalInfo::useExtBindlessImages()
+               ? getClNamespace() +
+                     "ext::oneapi::experimental::sampled_image_handle"
+               : getDpctNamespace() + "image_wrapper_base_p",
+           HelperFeatureEnum::device_ext)},
       {"textureReference",
        std::make_shared<TypeNameRule>(getDpctNamespace() + "image_wrapper_base",
                                       HelperFeatureEnum::device_ext)},
@@ -868,6 +875,7 @@ void MapNames::setExplicitNamespaceMap(
       "cudaTextureDesc",
       "cudaResourceDesc",
       "cudaTextureObject_t",
+      "cudaSurfaceObject_t",
       "textureReference",
       "cudaTextureAddressMode",
       "cudaTextureFilterMode",

clang/lib/DPCT/RulesLang/APINamesTexture.inc

@@ -64,6 +64,13 @@ ASSIGNABLE_FACTORY(CONDITIONAL_FACTORY_ENTRY(
                                 "experimental::destroy_bindless_image",
                             ARG(0), QUEUESTR)),
     DELETER_FACTORY_ENTRY("cuTexObjectDestroy", ARG(0))))
+ASSIGNABLE_FACTORY(CONDITIONAL_FACTORY_ENTRY(
+    UseExtBindlessImages,
+    CALL_FACTORY_ENTRY("cudaDestroySurfaceObject",
+                       CALL(MapNames::getDpctNamespace() +
+                                "experimental::destroy_bindless_image",
+                            ARG(0), QUEUESTR)),
+    DELETER_FACTORY_ENTRY("cudaDestroySurfaceObject", ARG(0))))
 
 CONDITIONAL_FACTORY_ENTRY(
     UseExtBindlessImages,
@@ -104,6 +111,15 @@ CONDITIONAL_FACTORY_ENTRY(
                                 "cuTexObjectGetTextureDesc", DEREF(0),
                                 MEMBER_CALL(ARG(1), true,
                                             "get_sampling_info")))))
+CONDITIONAL_FACTORY_ENTRY(
+    UseExtBindlessImages,
+    ASSIGNABLE_FACTORY(ASSIGN_FACTORY_ENTRY(
+        "cudaGetSurfaceObjectResourceDesc", DEREF(0),
+        CALL(MapNames::getDpctNamespace() + "experimental::get_data", ARG(1)))),
+    FEATURE_REQUEST_FACTORY(HelperFeatureEnum::device_ext,
+                            ASSIGNABLE_FACTORY(ASSIGN_FACTORY_ENTRY(
+                                "cudaGetSurfaceObjectResourceDesc", DEREF(0),
+                                MEMBER_CALL(ARG(1), true, "get_data")))))
 
 FEATURE_REQUEST_FACTORY(HelperFeatureEnum::device_ext,
                         ENTRY_TEXTURE("tex1D", 0x01, 1))
@@ -184,6 +200,16 @@ FEATURE_REQUEST_FACTORY(HelperFeatureEnum::device_ext,
                                            "create_image_wrapper",
                                  DEREF(1), DEREF(2)))))
 
+FEATURE_REQUEST_FACTORY(HelperFeatureEnum::device_ext,
+                        ASSIGNABLE_FACTORY(ASSIGN_FACTORY_ENTRY(
+                            "cudaCreateSurfaceObject", DEREF(0),
+                            CALL(DpctGlobalInfo::useExtBindlessImages()
+                                     ? MapNames::getDpctNamespace() +
+                                           "experimental::create_bindless_image"
+                                     : MapNames::getDpctNamespace() +
+                                           "create_image_wrapper",
+                                 DEREF(1)))))
+
 ENTRY_UNSUPPORTED("cudaGetTextureObjectResourceViewDesc",
                   Diagnostics::API_NOT_MIGRATED)
 

clang/lib/DPCT/RulesLang/RulesLang.cpp

@@ -9698,6 +9698,12 @@ void TextureRule::registerMatcher(MatchFinder &MF) {
           )
           .bind("tex"),
       this);
+  MF.addMatcher(
+      typeLoc(
+          loc(qualType(hasDeclaration(typedefDecl(hasAnyName(
+              "cudaTextureObject_t", "cudaSurfaceObject_t", "CUtexObject"))))))
+          .bind("texObj"),
+      this);
   MF.addMatcher(typeLoc(loc(qualType(hasDeclaration(typedefDecl(hasAnyName(
                             "cudaTextureObject_t", "CUtexObject"))))))
                     .bind("texObj"),
@@ -9761,6 +9767,9 @@ void TextureRule::registerMatcher(MatchFinder &MF) {
       "cudaGetTextureObjectResourceDesc",
       "cudaGetTextureObjectTextureDesc",
       "cudaGetTextureObjectResourceViewDesc",
+      "cudaCreateSurfaceObject",
+      "cudaDestroySurfaceObject",
+      "cudaGetSurfaceObjectResourceDesc",
       "cuArray3DCreate_v2",
       "cuArrayCreate_v2",
       "cuArrayDestroy",

clang/lib/DPCT/SrcAPI/APINames.inc

@@ -379,9 +379,9 @@ ENTRY(cudaGetTextureObjectResourceViewDesc, cudaGetTextureObjectResourceViewDesc
 ENTRY(cudaGetTextureObjectTextureDesc, cudaGetTextureObjectTextureDesc, true, NO_FLAG, P4, "Successful")
 
 // low level texture surface management functions of runtime API
-ENTRY(cudaCreateSurfaceObject, cudaCreateSurfaceObject, false, NO_FLAG, P0, "comment")
-ENTRY(cudaDestroySurfaceObject, cudaDestroySurfaceObject, false, NO_FLAG, P0, "comment")
-ENTRY(cudaGetSurfaceObjectResourceDesc, cudaGetSurfaceObjectResourceDesc, false, NO_FLAG, P4, "comment")
+ENTRY(cudaCreateSurfaceObject, cudaCreateSurfaceObject, true, NO_FLAG, P0, "Successful")
+ENTRY(cudaDestroySurfaceObject, cudaDestroySurfaceObject, true, NO_FLAG, P0, "Successful")
+ENTRY(cudaGetSurfaceObjectResourceDesc, cudaGetSurfaceObjectResourceDesc, true, NO_FLAG, P4, "Successful")
 
 // Version Management
 ENTRY(cudaDriverGetVersion, cudaDriverGetVersion, true, NO_FLAG, P0, "DPCT1043")

clang/runtime/dpct-rt/include/dpct/bindless_images.hpp

@@ -844,7 +844,7 @@ inline void unmap_resources(int count, external_mem_wrapper **handles,
 /// \param [in] q The queue where the image creation be executed.
 /// \returns The sampled image handle of created bindless image.
 static inline sycl::ext::oneapi::experimental::sampled_image_handle
-create_bindless_image(image_data data, sampling_info info,
+create_bindless_image(image_data data, sampling_info info = {},
                       sycl::queue q = get_default_queue()) {
   auto samp = sycl::ext::oneapi::experimental::bindless_image_sampler(
       info.get_addressing_mode(), info.get_coordinate_normalization_mode(),

clang/test/dpct/texture/surface_object_bindless_image.cu

@@ -0,0 +1,16 @@
+// RUN: dpct --format-range=none --use-experimental-features=bindless_images -out-root %T/texture/surface_object_bindless_image %s --cuda-include-path="%cuda-path/include" -- -x cuda --cuda-host-only -std=c++14
+// RUN: FileCheck --input-file %T/texture/surface_object_bindless_image/surface_object_bindless_image.dp.cpp --match-full-lines %s
+// RUN: %if build_lit %{icpx -c -fsycl %T/texture/surface_object_bindless_image/surface_object_bindless_image.dp.cpp -o %T/texture/surface_object_bindless_image/surface_object_bindless_image.dp.o %}
+
+int main() {
+  // CHECK: sycl::ext::oneapi::experimental::sampled_image_handle surf;
+  cudaSurfaceObject_t surf;
+  // CHECK: dpct::image_data resDesc;
+  cudaResourceDesc resDesc;
+  // CHECK: surf = dpct::experimental::create_bindless_image(resDesc);
+  cudaCreateSurfaceObject(&surf, &resDesc);
+  // CHECK: dpct::experimental::destroy_bindless_image(surf, dpct::get_in_order_queue());
+  cudaDestroySurfaceObject(surf);
+  // CHECK: resDesc = dpct::experimental::get_data(surf);
+  cudaGetSurfaceObjectResourceDesc(&resDesc, surf);
+}