[SYCLomatic] Support migration of 3 surface sample APIs.

clang/lib/DPCT/AnalysisInfo.cpp

@@ -5034,9 +5034,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) {

clang/lib/DPCT/AnalysisInfo.h

@@ -2051,7 +2051,7 @@ 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;
 

clang/lib/DPCT/RulesLang/APINamesTexture.inc

@@ -287,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

@@ -9706,10 +9706,6 @@ void TextureRule::registerMatcher(MatchFinder &MF) {
               "cudaTextureObject_t", "cudaSurfaceObject_t", "CUtexObject"))))))
           .bind("texObj"),
       this);
-  MF.addMatcher(typeLoc(loc(qualType(hasDeclaration(typedefDecl(hasAnyName(
-                            "cudaTextureObject_t", "CUtexObject"))))))
-                    .bind("texObj"),
-                this);
   MF.addMatcher(
       memberExpr(hasObjectExpression(hasType(
           type(hasUnqualifiedDesugaredType(recordType(hasDeclaration(
@@ -9764,6 +9760,9 @@ void TextureRule::registerMatcher(MatchFinder &MF) {
       "tex1Dfetch",
       "tex1DLayered",
       "tex2DLayered",
+      "surf1Dread",
+      "surf2Dread",
+      "surf3Dread",
       "cudaCreateTextureObject",
       "cudaDestroyTextureObject",
       "cudaGetTextureObjectResourceDesc",

clang/lib/Headers/__clang_cuda_runtime_wrapper.h

@@ -425,6 +425,9 @@ __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"
+#ifdef SYCLomatic_CUSTOMIZATION
+#include "surface_indirect_functions.h"
+#endif // SYCLomatic_CUSTOMIZATION
 
 // Restore state of __CUDA_ARCH__ and __THROW we had on entry.
 #pragma pop_macro("__CUDA_ARCH__")

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

@@ -893,6 +893,51 @@ class image_accessor_ext {
   typename std::enable_if<Available, data_t>::type read(CoordT x) {
     return detail::fetch_data<T>()(_img_acc.read(x, _sampler));
   }
+
+  /// Read data from accessor.
+  template <bool Available = dimensions == 3>
+  typename std::enable_if<Available, data_t>::type read_byte(float x, float y,
+                                                        float z) {
+    return detail::fetch_data<T>()(
+        _img_acc.read(sycl::float4(x / sizeof(T), y, z, 0), _sampler));
+  }
+  /// Read data from accessor.
+  template <class Coord0, class Coord1, class Coord2,
+            bool Available = dimensions == 3 &&
+                             std::is_integral<Coord0>::value
+                                 &&std::is_integral<Coord1>::value
+                                     &&std::is_integral<Coord2>::value>
+  typename std::enable_if<Available, data_t>::type read_byte(Coord0 x, Coord1 y,
+                                                        Coord2 z) {
+    return detail::fetch_data<T>()(
+        _img_acc.read(sycl::int4(x / sizeof(T), y, z, 0), _sampler));
+  }
+  /// Read data from accessor.
+  template <bool Available = dimensions == 2>
+  typename std::enable_if<Available, data_t>::type read_byte(float x, float y) {
+    return detail::fetch_data<T>()(
+        _img_acc.read(sycl::float2(x / sizeof(T), y), _sampler));
+  }
+  /// Read data from accessor.
+  template <class Coord0, class Coord1,
+            bool Available = dimensions == 2 &&
+                             std::is_integral<Coord0>::value
+                                 &&std::is_integral<Coord1>::value>
+  typename std::enable_if<Available, data_t>::type read_byte(Coord0 x, Coord1 y) {
+    return detail::fetch_data<T>()(
+        _img_acc.read(sycl::int2(x / sizeof(T), y), _sampler));
+  }
+  /// Read data from accessor.
+  template <bool Available = dimensions == 1>
+  typename std::enable_if<Available, data_t>::type read_byte(float x) {
+    return detail::fetch_data<T>()(_img_acc.read(x / sizeof(T), _sampler));
+  }
+  /// Read data from accessor.
+  template <class CoordT,
+            bool Available = dimensions == 1 && std::is_integral<CoordT>::value>
+  typename std::enable_if<Available, data_t>::type read_byte(CoordT x) {
+    return detail::fetch_data<T>()(_img_acc.read(x / sizeof(T), _sampler));
+  }
 };
 
 template <class T, int dimensions> class image_accessor_ext<T, dimensions, true> {

clang/test/dpct/texture/surface_object_bindless_image.cu

@@ -2,6 +2,23 @@
 // 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;