[SYCLomatic] Support migration of 3 surface sample APIs.

clang/lib/DPCT/AnalysisInfo.cpp

@@ -3600,8 +3600,12 @@ const std::string MemVarInfo::ExternVariableName = "dpct_local";
 std::unordered_map<std::string, int> MemVarInfo::AnonymousTypeDeclStmtMap;
 ///// class TextureTypeInfo /////
 TextureTypeInfo::TextureTypeInfo(std::string &&DataType, int TexType) {
+  TypeLength = DataType.length();
   setDataTypeAndTexType(std::move(DataType), TexType);
 }
+int TextureTypeInfo::getTypeLength() {
+  return TypeLength;
+}
 void TextureTypeInfo::setDataTypeAndTexType(std::string &&Type, int TexType) {
   DataType = std::move(Type);
   IsArray = TexType & 0xF0;
@@ -3791,7 +3795,7 @@ void TextureObjectInfo::merge(std::shared_ptr<TextureObjectInfo> Target) {
 void TextureObjectInfo::addParamDeclReplacement() {
   if (Type) {
     DpctGlobalInfo::getInstance().addReplacement(
-        std::make_shared<ExtReplacement>(FilePath, Offset, ReplaceTypeLength,
+        std::make_shared<ExtReplacement>(FilePath, Offset, Type->getTypeLength(),
                                          getParamDeclType(), nullptr));
   }
 }
@@ -4924,7 +4928,7 @@ void DeviceFunctionDecl::emplaceReplacement() {
         DpctGlobalInfo::getInstance().addReplacement(
             std::make_shared<ExtReplacement>(
                 Obj->getFilePath(), Obj->getOffset(),
-                strlen("cudaTextureObject_t"),
+                Obj->getType()->getTypeLength(),
                 MapNames::getClNamespace() +
                     "ext::oneapi::experimental::sampled_image_handle",
                 nullptr));
@@ -5034,9 +5038,12 @@ void DeviceFunctionDecl::buildTextureObjectParamsInfo(
     return;
   for (unsigned Idx = 0; Idx < Parms.size(); ++Idx) {
     auto Param = Parms[Idx];
-    if (DpctGlobalInfo::getUnqualifiedTypeName(Param->getType()) ==
-        "cudaTextureObject_t")
+    std::string ParamName =
+        DpctGlobalInfo::getUnqualifiedTypeName(Param->getType());
+    if (ParamName == "cudaTextureObject_t" ||
+        ParamName == "cudaSurfaceObject_t") {
       TextureObjectList[Idx] = std::make_shared<TextureObjectInfo>(Param);
+    }
   }
 }
 std::string DeviceFunctionDecl::getExtraParameters(LocInfo LI) {
@@ -6583,8 +6590,6 @@ std::string CudaMallocInfo::getAssignArgs(const std::string &TypeName) {
 ///// end /////
 int HostDeviceFuncInfo::MaxId = 0;
 
-const int TextureObjectInfo::ReplaceTypeLength = strlen("cudaTextureObject_t");
-
 #define TYPE_CAST(qual_type, type) dyn_cast<type>(qual_type)
 #define ARG_TYPE_CAST(type) TYPE_CAST(ArgType, type)
 #define PARM_TYPE_CAST(type) TYPE_CAST(ParmType, type)

clang/lib/DPCT/AnalysisInfo.h

@@ -1989,11 +1989,13 @@ class MemVarInfo : public VarInfo {
 
 class TextureTypeInfo {
   std::string DataType;
+  int TypeLength;
   int Dimension;
   bool IsArray;
 
 public:
   TextureTypeInfo(std::string &&DataType, int TexType);
+  int getTypeLength();
   void setDataTypeAndTexType(std::string &&Type, int TexType);
   void prepareForImage();
   void endForImage();
@@ -2049,10 +2051,8 @@ class TextureInfo {
   bool isUseHelperFunc() { return true; }
 };
 
-// texture handle info
+// texture object info can be used for CUDA texture and suface objects.
 class TextureObjectInfo : public TextureInfo {
-  static const int ReplaceTypeLength;
-
   // If it is a parameter in the function, it is the parameter index, either it
   // is 0.
   unsigned ParamIdx;

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)
 
@@ -261,3 +287,10 @@ FEATURE_REQUEST_FACTORY(
         "cuTexRefSetFormat", true, false, true, true,
         MEMBER_CALL(ARG(0), true, "set_channel_type", ARG(1)),
         MEMBER_CALL(ARG(0), true, "set_channel_num", ARG(2))))
+
+FEATURE_REQUEST_FACTORY(HelperFeatureEnum::device_ext,
+                        ENTRY_TEXTURE("surf1Dread", 0x01, 1))
+FEATURE_REQUEST_FACTORY(HelperFeatureEnum::device_ext,
+                        ENTRY_TEXTURE("surf2Dread", 0x02, 1, 2))
+FEATURE_REQUEST_FACTORY(HelperFeatureEnum::device_ext,
+                        ENTRY_TEXTURE("surf3Dread", 0x03, 1, 2, 3))

clang/lib/DPCT/RulesLang/RulesLang.cpp

@@ -11957,10 +11957,13 @@ void TextureRule::registerMatcher(MatchFinder &MF) {
           )
           .bind("tex"),
       this);
-  MF.addMatcher(typeLoc(loc(qualType(hasDeclaration(typedefDecl(hasAnyName(
-                            "cudaTextureObject_t", "CUtexObject"))))))
-                    .bind("texObj"),
-                this);
+  MF.addMatcher(
+      typeLoc(
+          loc(qualType(hasDeclaration(typedefDecl(hasAnyName(
+              "cudaTextureObject_t", "cudaSurfaceObject_t", "CUtexObject"))))))
+          .bind("texObj"),
+      this);
+
   MF.addMatcher(
       memberExpr(hasObjectExpression(hasType(
           type(hasUnqualifiedDesugaredType(recordType(hasDeclaration(
@@ -12015,11 +12018,17 @@ void TextureRule::registerMatcher(MatchFinder &MF) {
       "tex1Dfetch",
       "tex1DLayered",
       "tex2DLayered",
+      "surf1Dread",
+      "surf2Dread",
+      "surf3Dread",
       "cudaCreateTextureObject",
       "cudaDestroyTextureObject",
       "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/lib/Headers/__clang_cuda_runtime_wrapper.h

@@ -425,6 +425,7 @@ __host__ __device__ void __nv_tex_surf_handler(const char *name, T *ptr,
 #endif // __cplusplus >= 201103L && CUDA_VERSION >= 9000
 #include "texture_fetch_functions.h"
 #include "texture_indirect_functions.h"
+#include "surface_indirect_functions.h"
 
 // Restore state of __CUDA_ARCH__ and __THROW we had on entry.
 #pragma pop_macro("__CUDA_ARCH__")

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,34 @@
+// 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 %}
+
+
+template<typename T> __global__ void kernel(cudaSurfaceObject_t surf) {
+  int i;
+  float j, k, l, m;
+  // CHECK: sycl::ext::oneapi::experimental::sample_image<T>(surf, float(i));
+  surf1Dread<T>(surf, i);
+  // CHECK: i = sycl::ext::oneapi::experimental::sample_image<T>(surf, float(i));
+  surf1Dread<T>(&i, surf, i);
+  // CHECK: i = sycl::ext::oneapi::experimental::sample_image<T>(surf, sycl::float2(j, i));
+  surf2Dread<T>(surf, j, i);
+  // CHECK: sycl::ext::oneapi::experimental::sample_image<T>(surf, sycl::float2(j, i));
+  surf2Dread<T>(&i, surf, j, i);
+  // CHECK: sycl::ext::oneapi::experimental::sample_image<T>(surf, sycl::float3(k, j, i));
+  surf3Dread<T>(surf, k, j, i);
+  // CHECK: i = sycl::ext::oneapi::experimental::sample_image<T>(surf, sycl::float3(k, j, i));
+  surf3Dread<T>(&i, surf, k, j, i);
+}
+
+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);
+}