[SYCLomatic] Support migration of 3 surface sample APIs.

clang/lib/DPCT/AnalysisInfo.cpp

@@ -5036,9 +5036,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

@@ -2054,7 +2054,7 @@ class TextureInfo {
   bool isUseHelperFunc() { return true; }
 };
 
-// texture handle info
+// texture object info can be used for CUDA texture and surface objects.
 class TextureObjectInfo : public TextureInfo {
   static const int ReplaceTypeLength;
 

clang/lib/DPCT/RulesLang/APINamesTexture.inc

@@ -120,48 +120,101 @@ CONDITIONAL_FACTORY_ENTRY(
                             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))
+CONDITIONAL_FACTORY_ENTRY(
+    UseExtBindlessImages,
+    FEATURE_REQUEST_FACTORY(
+        HelperFeatureEnum::device_ext,
+        ENTRY_TEXTURE("tex1D", 0x01,
+                      MapNames::getClNamespace() +
+                          "ext::oneapi::experimental::sample_image",
+                      1)),
+    FEATURE_REQUEST_FACTORY(HelperFeatureEnum::device_ext,
+                            ENTRY_TEXTURE("tex1D", 0x01, "read", 1)))
 
 CONDITIONAL_FACTORY_ENTRY(
-    UseExtBindlessImages, ENTRY_TEXTURE("tex1DLod", 0x11, 1),
+    UseExtBindlessImages,
+    ENTRY_TEXTURE("tex1DLod", 0x11,
+                  MapNames::getClNamespace() +
+                      "ext::oneapi::experimental::sample_mipmap",
+                  1),
     UNSUPPORT_FACTORY_ENTRY("tex1DLod", Diagnostics::TRY_EXPERIMENTAL_FEATURE,
                             ARG("tex1DLod"),
                             ARG("--use-experimental-features=bindless_images")))
-
-FEATURE_REQUEST_FACTORY(HelperFeatureEnum::device_ext,
-                        ENTRY_TEXTURE("tex2D", 0x02, 1, 2))
+CONDITIONAL_FACTORY_ENTRY(
+    UseExtBindlessImages,
+    FEATURE_REQUEST_FACTORY(
+        HelperFeatureEnum::device_ext,
+        ENTRY_TEXTURE("tex2D", 0x02,
+                      MapNames::getClNamespace() +
+                          "ext::oneapi::experimental::sample_image",
+                      1, 2)),
+    FEATURE_REQUEST_FACTORY(HelperFeatureEnum::device_ext,
+                            ENTRY_TEXTURE("tex2D", 0x02, "read", 1, 2)))
 
 CONDITIONAL_FACTORY_ENTRY(
-    UseExtBindlessImages, ENTRY_TEXTURE("tex2DLod", 0x12, 1, 2),
+    UseExtBindlessImages,
+    ENTRY_TEXTURE("tex2DLod", 0x12,
+                  MapNames::getClNamespace() +
+                      "ext::oneapi::experimental::sample_mipmap",
+                  1, 2),
     UNSUPPORT_FACTORY_ENTRY("tex2DLod", Diagnostics::TRY_EXPERIMENTAL_FEATURE,
                             ARG("tex2DLod"),
                             ARG("--use-experimental-features=bindless_images")))
 
-FEATURE_REQUEST_FACTORY(HelperFeatureEnum::device_ext,
-                        ENTRY_TEXTURE("tex3D", 0x03, 1, 2, 3))
+FEATURE_REQUEST_FACTORY(
+    HelperFeatureEnum::device_ext,
+    ENTRY_TEXTURE("tex3D", 0x03,
+                  MapNames::getClNamespace() +
+                      "ext::oneapi::experimental::sample_image",
+                  1, 2, 3))
 
 CONDITIONAL_FACTORY_ENTRY(
-    UseExtBindlessImages, ENTRY_TEXTURE("tex3DLod", 0x13, 1, 2, 3),
+    UseExtBindlessImages,
+    ENTRY_TEXTURE("tex3DLod", 0x13,
+                  MapNames::getClNamespace() +
+                      "ext::oneapi::experimental::sample_mipmap",
+                  1, 2, 3),
     UNSUPPORT_FACTORY_ENTRY("tex3DLod", Diagnostics::TRY_EXPERIMENTAL_FEATURE,
                             ARG("tex3DLod"),
                             ARG("--use-experimental-features=bindless_images")))
 
-FEATURE_REQUEST_FACTORY(
-    HelperFeatureEnum::device_ext,
-    WARNING_FACTORY_ENTRY("tex1Dfetch", ENTRY_TEXTURE("tex1Dfetch", 0x01, 1),
-                          Diagnostics::TEX_FETCH))
-CONDITIONAL_FACTORY_ENTRY(UseExtBindlessImages,
-                          ENTRY_TEXTURE("tex1DLayered", 0xF1, 1),
-                          FEATURE_REQUEST_FACTORY(HelperFeatureEnum::device_ext,
-                                                  ENTRY_TEXTURE("tex1DLayered",
-                                                                0xF1, 2, 1)))
-CONDITIONAL_FACTORY_ENTRY(UseExtBindlessImages,
-                          ENTRY_TEXTURE("tex2DLayered", 0xF2, 1, 2),
-                          FEATURE_REQUEST_FACTORY(HelperFeatureEnum::device_ext,
-                                                  ENTRY_TEXTURE("tex2DLayered",
-                                                                0xF2, 3, 1, 2)))
+CONDITIONAL_FACTORY_ENTRY(
+    UseExtBindlessImages,
+    FEATURE_REQUEST_FACTORY(
+        HelperFeatureEnum::device_ext,
+        WARNING_FACTORY_ENTRY(
+            "tex1Dfetch",
+            ENTRY_TEXTURE("tex1Dfetch", 0x01,
+                          MapNames::getClNamespace() +
+                              "ext::oneapi::experimental::sample_image",
+                          1),
+            Diagnostics::TEX_FETCH)),
+    FEATURE_REQUEST_FACTORY(
+        HelperFeatureEnum::device_ext,
+        WARNING_FACTORY_ENTRY("tex1Dfetch",
+                              ENTRY_TEXTURE("tex1Dfetch", 0x01, "read", 1),
+                              Diagnostics::TEX_FETCH)))
+
+CONDITIONAL_FACTORY_ENTRY(
+    UseExtBindlessImages,
+    FEATURE_REQUEST_FACTORY(
+        HelperFeatureEnum::device_ext,
+        ENTRY_TEXTURE("tex1DLayered", 0xF1,
+                      MapNames::getClNamespace() +
+                          "ext::oneapi::experimental::sample_image_array",
+                      1)),
+    FEATURE_REQUEST_FACTORY(HelperFeatureEnum::device_ext,
+                            ENTRY_TEXTURE("tex1DLayered", 0xF1, "read", 2, 1)))
+
+CONDITIONAL_FACTORY_ENTRY(
+    UseExtBindlessImages,
+    ENTRY_TEXTURE("tex2DLayered", 0xF2,
+                  MapNames::getClNamespace() +
+                      "ext::oneapi::experimental::sample_image_array",
+                  1, 2),
+    FEATURE_REQUEST_FACTORY(HelperFeatureEnum::device_ext,
+                            ENTRY_TEXTURE("tex2DLayered", 0xF2, "read", 3, 1,
+                                          2)))
 
 FEATURE_REQUEST_FACTORY(HelperFeatureEnum::device_ext,
                         ASSIGNABLE_FACTORY(ENTRY_BIND("cudaBindTexture",
@@ -287,3 +340,37 @@ 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))))
+
+CONDITIONAL_FACTORY_ENTRY(
+    UseExtBindlessImages,
+    FEATURE_REQUEST_FACTORY(
+        HelperFeatureEnum::device_ext,
+        ENTRY_TEXTURE("surf1Dread", 0x01,
+                      MapNames::getLibraryHelperNamespace() +
+                          "experimental::sample_image_by_byte",
+                      1)),
+    FEATURE_REQUEST_FACTORY(HelperFeatureEnum::device_ext,
+                            ENTRY_TEXTURE("surf1Dread", 0x01, "read_byte", 1)))
+
+CONDITIONAL_FACTORY_ENTRY(
+    UseExtBindlessImages,
+    FEATURE_REQUEST_FACTORY(
+        HelperFeatureEnum::device_ext,
+        ENTRY_TEXTURE("surf2Dread", 0x02,
+                      MapNames::getLibraryHelperNamespace() +
+                          "experimental::sample_image_by_byte",
+                      1, 2)),
+    FEATURE_REQUEST_FACTORY(HelperFeatureEnum::device_ext,
+                            ENTRY_TEXTURE("surf2Dread", 0x02, "read_byte", 1,
+                                          2)))
+CONDITIONAL_FACTORY_ENTRY(
+    UseExtBindlessImages,
+    FEATURE_REQUEST_FACTORY(
+        HelperFeatureEnum::device_ext,
+        ENTRY_TEXTURE("surf3Dread", 0x03,
+                      MapNames::getLibraryHelperNamespace() +
+                          "experimental::sample_image_by_byte",
+                      1, 2, 3)),
+    FEATURE_REQUEST_FACTORY(HelperFeatureEnum::device_ext,
+                            ENTRY_TEXTURE("surf3Dread", 0x03, "read_byte", 1, 2,
+                                          3)))

clang/lib/DPCT/RulesLang/CallExprRewriterTexture.cpp

@@ -16,25 +16,25 @@ template <size_t... Idx>
 class TextureReadRewriterFactory : public CallExprRewriterFactoryBase {
   std::string Source;
   int TexType;
+  std::string TargetName;
 
   inline int getDim() const { return TexType & 0x0f; }
 
   template <class BaseT>
   std::shared_ptr<CallExprRewriter>
   createRewriter(const CallExpr *C, bool RetAssign, BaseT Base) const {
-    const static std::string MemberName = "read";
     using ReaderPrinter = decltype(makeMemberCallCreator<false>(
         std::declval<std::function<BaseT(const CallExpr *)>>(), false,
-        MemberName, makeCallArgCreatorWithCall(Idx)...)(C));
+        TargetName, makeCallArgCreatorWithCall(Idx)...)(C));
     if (RetAssign) {
       return std::make_shared<PrinterRewriter<
           BinaryOperatorPrinter<BO_Assign, DerefExpr, ReaderPrinter>>>(
           C, Source, DerefExpr(C->getArg(0), C),
-          ReaderPrinter(std::move(Base), false, MemberName,
+          ReaderPrinter(std::move(Base), false, TargetName,
                         std::make_pair(C, C->getArg(Idx + 1))...));
     }
     return std::make_shared<PrinterRewriter<ReaderPrinter>>(
-        C, Source, Base, false, MemberName,
+        C, Source, Base, false, TargetName,
         std::make_pair(C, C->getArg(Idx))...);
   }
 
@@ -43,8 +43,6 @@ class TextureReadRewriterFactory : public CallExprRewriterFactoryBase {
   createbindlessRewriterNormal(const CallExpr *C, bool RetAssign,
                                const TemplateArgumentInfo &TAI,
                                const std::string &VecTypeName) const {
-    const static std::string FuncName =
-        MapNames::getClNamespace() + "ext::oneapi::experimental::sample_image";
     using FuncNamePrinter =
         TemplatedNamePrinter<StringRef, std::vector<TemplateArgumentInfo>>;
     using ReaderPrinter =
@@ -55,11 +53,12 @@ class TextureReadRewriterFactory : public CallExprRewriterFactoryBase {
           BinaryOperatorPrinter<BO_Assign, DerefExpr, ReaderPrinter>>>(
           C, Source, DerefExpr(C->getArg(0), C),
           ReaderPrinter(
-              FuncNamePrinter(FuncName, {TAI}), std::make_pair(C, C->getArg(1)),
+              FuncNamePrinter(TargetName, {TAI}),
+              std::make_pair(C, C->getArg(1)),
               VecType(VecTypeName, std::make_pair(C, C->getArg(Idx + 1))...)));
     }
     return std::make_shared<PrinterRewriter<ReaderPrinter>>(
-        C, Source, FuncNamePrinter(FuncName, {TAI}),
+        C, Source, FuncNamePrinter(TargetName, {TAI}),
         std::make_pair(C, C->getArg(0)),
         VecType(VecTypeName, std::make_pair(C, C->getArg(Idx))...));
   }
@@ -69,8 +68,6 @@ class TextureReadRewriterFactory : public CallExprRewriterFactoryBase {
   createbindlessRewriterLod(const CallExpr *C, bool RetAssign,
                             const TemplateArgumentInfo &TAI,
                             const std::string &VecTypeName) const {
-    const static std::string FuncName =
-        MapNames::getClNamespace() + "ext::oneapi::experimental::sample_mipmap";
     using FuncNamePrinter =
         TemplatedNamePrinter<StringRef, std::vector<TemplateArgumentInfo>>;
     using ReaderPrinter =
@@ -82,12 +79,13 @@ class TextureReadRewriterFactory : public CallExprRewriterFactoryBase {
           BinaryOperatorPrinter<BO_Assign, DerefExpr, ReaderPrinter>>>(
           C, Source, DerefExpr(C->getArg(0), C),
           ReaderPrinter(
-              FuncNamePrinter(FuncName, {TAI}), std::make_pair(C, C->getArg(1)),
+              FuncNamePrinter(TargetName, {TAI}),
+              std::make_pair(C, C->getArg(1)),
               VecType(VecTypeName, std::make_pair(C, C->getArg(Idx + 1))...),
               std::make_pair(C, C->getArg(C->getNumArgs() - 1))));
     }
     return std::make_shared<PrinterRewriter<ReaderPrinter>>(
-        C, Source, FuncNamePrinter(FuncName, {TAI}),
+        C, Source, FuncNamePrinter(TargetName, {TAI}),
         std::make_pair(C, C->getArg(0)),
         VecType(VecTypeName, std::make_pair(C, C->getArg(Idx))...),
         std::make_pair(C, C->getArg(C->getNumArgs() - 1)));
@@ -98,9 +96,6 @@ class TextureReadRewriterFactory : public CallExprRewriterFactoryBase {
   createbindlessRewriterLayered(const CallExpr *C, bool RetAssign,
                                 const TemplateArgumentInfo &TAI,
                                 const std::string &VecTypeName) const {
-    const static std::string FuncName =
-        MapNames::getClNamespace() +
-        "ext::oneapi::experimental::sample_image_array";
     using FuncNamePrinter =
         TemplatedNamePrinter<StringRef, std::vector<TemplateArgumentInfo>>;
     using ReaderPrinter =
@@ -112,17 +107,17 @@ class TextureReadRewriterFactory : public CallExprRewriterFactoryBase {
           BinaryOperatorPrinter<BO_Assign, DerefExpr, ReaderPrinter>>>(
           C, Source, DerefExpr(C->getArg(0), C),
           ReaderPrinter(
-              FuncNamePrinter(FuncName, {TAI}), std::make_pair(C, C->getArg(1)),
+              FuncNamePrinter(TargetName, {TAI}),
+              std::make_pair(C, C->getArg(1)),
               VecType(VecTypeName, std::make_pair(C, C->getArg(Idx + 1))...),
               std::make_pair(C, C->getArg(C->getNumArgs() - 1))));
     }
     return std::make_shared<PrinterRewriter<ReaderPrinter>>(
-        C, Source, FuncNamePrinter(FuncName, {TAI}),
+        C, Source, FuncNamePrinter(TargetName, {TAI}),
         std::make_pair(C, C->getArg(0)),
         VecType(VecTypeName, std::make_pair(C, C->getArg(Idx))...),
         std::make_pair(C, C->getArg(C->getNumArgs() - 1)));
   }
-
   std::shared_ptr<CallExprRewriter>
   createbindlessRewriter(const CallExpr *C, bool RetAssign,
                          QualType TargetType) const {
@@ -160,8 +155,8 @@ class TextureReadRewriterFactory : public CallExprRewriterFactoryBase {
   }
 
 public:
-  TextureReadRewriterFactory(std::string Name, int Tex)
-      : Source(std::move(Name)), TexType(Tex) {}
+  TextureReadRewriterFactory(std::string Name, int Tex, std::string TargetName)
+      : Source(std::move(Name)), TexType(Tex), TargetName(TargetName) {}
   std::shared_ptr<CallExprRewriter>
   create(const CallExpr *Call) const override {
     const Expr *SourceExpr = Call->getArg(0);
@@ -227,14 +222,15 @@ class TextureReadRewriterFactory : public CallExprRewriterFactoryBase {
 /// original call expr.
 template <size_t... Idx>
 inline std::shared_ptr<CallExprRewriterFactoryBase>
-createTextureReaderRewriterFactory(const std::string &Source, int TextureType) {
-  return std::make_shared<TextureReadRewriterFactory<Idx...>>(Source,
-                                                              TextureType);
+createTextureReaderRewriterFactory(const std::string &Source, int TextureType,
+                                   const std::string &TargetName) {
+  return std::make_shared<TextureReadRewriterFactory<Idx...>>(
+      Source, TextureType, TargetName);
 }
 
-#define TEX_FUNCTION_FACTORY_ENTRY(FuncName, TexType, ...)                     \
-  {FuncName,                                                                   \
-   createTextureReaderRewriterFactory<__VA_ARGS__>(FuncName, TexType)},
+#define TEX_FUNCTION_FACTORY_ENTRY(FuncName, TexType, TargetName, ...)         \
+  {FuncName, createTextureReaderRewriterFactory<__VA_ARGS__>(                  \
+                 FuncName, TexType, TargetName)},
 #define BIND_TEXTURE_FACTORY_ENTRY(FuncName, ...)                              \
   {FuncName, createBindTextureRewriterFactory<__VA_ARGS__>(FuncName)},
 
@@ -244,16 +240,17 @@ createTextureReaderRewriterFactory(const std::string &Source, int TextureType) {
   REWRITER_FACTORY_ENTRY(FuncName, FuncCallExprRewriterFactory, RewriterName)
 #define UNSUPPORTED_FACTORY_ENTRY(FuncName, MsgID)                             \
   REWRITER_FACTORY_ENTRY(FuncName,                                             \
-      UnsupportFunctionRewriterFactory<std::string>, MsgID, FuncName)
+                         UnsupportFunctionRewriterFactory<std::string>, MsgID, \
+                         FuncName)
 
 void CallExprRewriterFactoryBase::initRewriterMapTexture() {
   RewriterMap->merge(
       std::unordered_map<std::string,
                          std::shared_ptr<CallExprRewriterFactoryBase>>({
 #define ENTRY_RENAMED(SOURCEAPINAME, TARGETAPINAME)                            \
   FUNC_NAME_FACTORY_ENTRY(SOURCEAPINAME, TARGETAPINAME)
-#define ENTRY_TEXTURE(SOURCEAPINAME, TEXTYPE, ...)                             \
-  TEX_FUNCTION_FACTORY_ENTRY(SOURCEAPINAME, TEXTYPE, __VA_ARGS__)
+#define ENTRY_TEXTURE(SOURCEAPINAME, TEXTYPE, TARGETAPINAME, ...)              \
+  TEX_FUNCTION_FACTORY_ENTRY(SOURCEAPINAME, TEXTYPE, TARGETAPINAME, __VA_ARGS__)
 #define ENTRY_UNSUPPORTED(SOURCEAPINAME, MSGID)                                \
   UNSUPPORTED_FACTORY_ENTRY(SOURCEAPINAME, MSGID)
 #define ENTRY_BIND(SOURCEAPINAME, ...)                                         \

clang/lib/DPCT/RulesLang/RulesLangTexture.cpp

@@ -560,6 +560,9 @@ void TextureRule::registerMatcher(MatchFinder &MF) {
       "tex1Dfetch",
       "tex1DLayered",
       "tex2DLayered",
+      "surf1Dread",
+      "surf2Dread",
+      "surf3Dread",
       "cudaCreateTextureObject",
       "cudaDestroyTextureObject",
       "cudaGetTextureObjectResourceDesc",

clang/lib/DPCT/SrcAPI/APINames.inc

@@ -1484,11 +1484,11 @@ ENTRY(texCubemapLod, texCubemapLod, false, NO_FLAG, P4, "comment")
 ENTRY(texCubemapLayered, texCubemapLayered, false, NO_FLAG, P4, "comment")
 ENTRY(texCubemapLayeredLod, texCubemapLayeredLod, false, NO_FLAG, P4, "comment")
 ENTRY(tex2Dgather, tex2Dgather, false, NO_FLAG, P0, "comment")
-ENTRY(surf1Dread, surf1Dread, false, NO_FLAG, P4, "comment")
+ENTRY(surf1Dread, surf1Dread, true, NO_FLAG, P4, "Successful")
 ENTRY(surf1Dwrite, surf1Dwrite, false, NO_FLAG, P0, "comment")
-ENTRY(surf2Dread, surf2Dread, false, NO_FLAG, P4, "comment")
+ENTRY(surf2Dread, surf2Dread, true, NO_FLAG, P4, "Successful")
 ENTRY(surf2Dwrite, surf2Dwrite, false, NO_FLAG, P0, "comment")
-ENTRY(surf3Dread, surf3Dread, false, NO_FLAG, P4, "comment")
+ENTRY(surf3Dread, surf3Dread, true, NO_FLAG, P4, "Successful")
 ENTRY(surf3Dwrite, surf3Dwrite, false, NO_FLAG, P0, "comment")
 ENTRY(surf1DLayeredread, surf1DLayeredread, false, NO_FLAG, P4, "comment")
 ENTRY(surf1DLayeredwrite, surf1DLayeredwrite, false, NO_FLAG, P4, "comment")

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__")