[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

@@ -1991,11 +1991,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();
@@ -2051,10 +2053,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/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;