diff --git a/clang/lib/DPCT/AnalysisInfo.cpp b/clang/lib/DPCT/AnalysisInfo.cpp index 82100e88d755..cad1ca73d815 100644 --- a/clang/lib/DPCT/AnalysisInfo.cpp +++ b/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(Param); + } } } std::string DeviceFunctionDecl::getExtraParameters(LocInfo LI) { diff --git a/clang/lib/DPCT/AnalysisInfo.h b/clang/lib/DPCT/AnalysisInfo.h index 5e88bee6936e..b6f4b01d8ef8 100644 --- a/clang/lib/DPCT/AnalysisInfo.h +++ b/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; diff --git a/clang/lib/DPCT/RulesLang/APINamesTexture.inc b/clang/lib/DPCT/RulesLang/APINamesTexture.inc index d0ccccd864ee..6ce5365b3aae 100644 --- a/clang/lib/DPCT/RulesLang/APINamesTexture.inc +++ b/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))) diff --git a/clang/lib/DPCT/RulesLang/CallExprRewriterTexture.cpp b/clang/lib/DPCT/RulesLang/CallExprRewriterTexture.cpp index ad0c9bb3f06e..ea93456381ea 100644 --- a/clang/lib/DPCT/RulesLang/CallExprRewriterTexture.cpp +++ b/clang/lib/DPCT/RulesLang/CallExprRewriterTexture.cpp @@ -16,25 +16,25 @@ template class TextureReadRewriterFactory : public CallExprRewriterFactoryBase { std::string Source; int TexType; + std::string TargetName; inline int getDim() const { return TexType & 0x0f; } template std::shared_ptr createRewriter(const CallExpr *C, bool RetAssign, BaseT Base) const { - const static std::string MemberName = "read"; using ReaderPrinter = decltype(makeMemberCallCreator( std::declval>(), false, - MemberName, makeCallArgCreatorWithCall(Idx)...)(C)); + TargetName, makeCallArgCreatorWithCall(Idx)...)(C)); if (RetAssign) { return std::make_shared>>( 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>( - 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>; using ReaderPrinter = @@ -55,11 +53,12 @@ class TextureReadRewriterFactory : public CallExprRewriterFactoryBase { BinaryOperatorPrinter>>( 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>( - 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>; using ReaderPrinter = @@ -82,12 +79,13 @@ class TextureReadRewriterFactory : public CallExprRewriterFactoryBase { BinaryOperatorPrinter>>( 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>( - 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>; using ReaderPrinter = @@ -112,17 +107,17 @@ class TextureReadRewriterFactory : public CallExprRewriterFactoryBase { BinaryOperatorPrinter>>( 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>( - 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 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 create(const CallExpr *Call) const override { const Expr *SourceExpr = Call->getArg(0); @@ -227,14 +222,15 @@ class TextureReadRewriterFactory : public CallExprRewriterFactoryBase { /// original call expr. template inline std::shared_ptr -createTextureReaderRewriterFactory(const std::string &Source, int TextureType) { - return std::make_shared>(Source, - TextureType); +createTextureReaderRewriterFactory(const std::string &Source, int TextureType, + const std::string &TargetName) { + return std::make_shared>( + 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,7 +240,8 @@ createTextureReaderRewriterFactory(const std::string &Source, int TextureType) { REWRITER_FACTORY_ENTRY(FuncName, FuncCallExprRewriterFactory, RewriterName) #define UNSUPPORTED_FACTORY_ENTRY(FuncName, MsgID) \ REWRITER_FACTORY_ENTRY(FuncName, \ - UnsupportFunctionRewriterFactory, MsgID, FuncName) + UnsupportFunctionRewriterFactory, MsgID, \ + FuncName) void CallExprRewriterFactoryBase::initRewriterMapTexture() { RewriterMap->merge( @@ -252,8 +249,8 @@ void CallExprRewriterFactoryBase::initRewriterMapTexture() { std::shared_ptr>({ #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, ...) \ diff --git a/clang/lib/DPCT/RulesLang/RulesLangTexture.cpp b/clang/lib/DPCT/RulesLang/RulesLangTexture.cpp index 29272deda0d6..3ee51c786ece 100644 --- a/clang/lib/DPCT/RulesLang/RulesLangTexture.cpp +++ b/clang/lib/DPCT/RulesLang/RulesLangTexture.cpp @@ -560,6 +560,9 @@ void TextureRule::registerMatcher(MatchFinder &MF) { "tex1Dfetch", "tex1DLayered", "tex2DLayered", + "surf1Dread", + "surf2Dread", + "surf3Dread", "cudaCreateTextureObject", "cudaDestroyTextureObject", "cudaGetTextureObjectResourceDesc", diff --git a/clang/lib/DPCT/SrcAPI/APINames.inc b/clang/lib/DPCT/SrcAPI/APINames.inc index c750e4de889f..49e3a7421f79 100644 --- a/clang/lib/DPCT/SrcAPI/APINames.inc +++ b/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") diff --git a/clang/lib/Headers/__clang_cuda_runtime_wrapper.h b/clang/lib/Headers/__clang_cuda_runtime_wrapper.h index 9c5c09fc405d..f96be6961ece 100644 --- a/clang/lib/Headers/__clang_cuda_runtime_wrapper.h +++ b/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__") diff --git a/clang/runtime/dpct-rt/include/dpct/bindless_images.hpp b/clang/runtime/dpct-rt/include/dpct/bindless_images.hpp index 0bbd4baa6691..dff7e109c297 100644 --- a/clang/runtime/dpct-rt/include/dpct/bindless_images.hpp +++ b/clang/runtime/dpct-rt/include/dpct/bindless_images.hpp @@ -1352,6 +1352,21 @@ static inline void dpct_memcpy(image_mem_wrapper *dest, size_t w_offset_dest, sycl::free(temp, q); } +// A wrapper for sycl sample_image function for the byte addressing image. +template +DataT sample_image_by_byte( + const sycl::ext::oneapi::experimental::sampled_image_handle &imageHandle, + CoordT &&coords) { + if constexpr (std::is_scalar_v) { + return sycl::ext::oneapi::experimental::sample_image( + imageHandle, coords / sizeof(DataT)); + } else { + coords[0] = coords[0] / sizeof(DataT); + return sycl::ext::oneapi::experimental::sample_image( + imageHandle, coords); + } +} + using image_mem_wrapper_ptr = image_mem_wrapper *; #ifdef _WIN32 using external_mem_wrapper_ptr = external_mem_wrapper *; diff --git a/clang/runtime/dpct-rt/include/dpct/image.hpp b/clang/runtime/dpct-rt/include/dpct/image.hpp index 3f81d7c3c9d9..2907423fced7 100644 --- a/clang/runtime/dpct-rt/include/dpct/image.hpp +++ b/clang/runtime/dpct-rt/include/dpct/image.hpp @@ -893,6 +893,36 @@ class image_accessor_ext { typename std::enable_if::type read(CoordT x) { return detail::fetch_data()(_img_acc.read(x, _sampler)); } + + /// Read data from accessor. + template ::value + &&std::is_integral::value + &&std::is_integral::value> + typename std::enable_if::type read_byte(Coord0 x, Coord1 y, + Coord2 z) { + return detail::fetch_data()( + _img_acc.read(sycl::int4(x / sizeof(T), y, z, 0), _sampler)); + } + + /// Read data from accessor. + template ::value + &&std::is_integral::value> + typename std::enable_if::type read_byte(Coord0 x, + Coord1 y) { + return detail::fetch_data()( + _img_acc.read(sycl::int2(x / sizeof(T), y), _sampler)); + } + + /// Read data from accessor. + template ::value> + typename std::enable_if::type read_byte(CoordT x) { + return detail::fetch_data()(_img_acc.read(x / sizeof(T), _sampler)); + } }; template class image_accessor_ext { diff --git a/clang/test/dpct/surface.cu b/clang/test/dpct/surface.cu new file mode 100644 index 000000000000..5e1e5431705b --- /dev/null +++ b/clang/test/dpct/surface.cu @@ -0,0 +1,38 @@ +// RUN: dpct --format-range=none -out-root %T/surface %s --cuda-include-path="%cuda-path/include" -- -x cuda --cuda-host-only -std=c++14 +// RUN: FileCheck --input-file %T/surface/surface.dp.cpp --match-full-lines %s + +// CHECK: template void kernel(dpct::image_accessor_ext surf) { +template __global__ void kernel(cudaSurfaceObject_t surf) { + int i; + float j, k, l, m; + // CHECK: surf.read_byte(i); + surf1Dread(surf, i); + // CHECK: i = surf.read_byte(i); + surf1Dread(&i, surf, i); + // CHECK: surf.read_byte(j, i); + surf2Dread(surf, j, i); + // CHECK: i = surf.read_byte(j, i); + surf2Dread(&i, surf, j, i); + // CHECK: surf.read_byte(k, j, i); + surf3Dread(surf, k, j, i); + // CHECK: i = surf.read_byte(k, j, i); + surf3Dread(&i, surf, k, j, i); +} + +static texture tex21; + +__device__ void device01() { + tex1D(tex21, 1.0f); +} +int main() { + // CHECK: dpct::image_wrapper_base_p surf; + cudaSurfaceObject_t surf; + // CHECK: dpct::image_data resDesc; + cudaResourceDesc resDesc; + // CHECK: surf = dpct::create_image_wrapper(resDesc); + cudaCreateSurfaceObject(&surf, &resDesc); + + kernel<<<1,1>>>(surf); + cudaDestroySurfaceObject(surf); + cudaGetSurfaceObjectResourceDesc(&resDesc, surf); +} diff --git a/clang/test/dpct/texture/surface_object_bindless_image.cu b/clang/test/dpct/texture/surface_object_bindless_image.cu index 6a34429c63c7..70472a11d45f 100644 --- a/clang/test/dpct/texture/surface_object_bindless_image.cu +++ b/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 %} +// CHECK: template void kernel(sycl::ext::oneapi::experimental::sampled_image_handle surf) { +template __global__ void kernel(cudaSurfaceObject_t surf) { + int i; + float j, k, l, m; + // CHECK: dpct::experimental::sample_image_by_byte(surf, float(i)); + surf1Dread(surf, i); + // CHECK: i = dpct::experimental::sample_image_by_byte(surf, float(i)); + surf1Dread(&i, surf, i); + // CHECK: dpct::experimental::sample_image_by_byte(surf, sycl::float2(j, i)); + surf2Dread(surf, j, i); + // CHECK: i = dpct::experimental::sample_image_by_byte(surf, sycl::float2(j, i)); + surf2Dread(&i, surf, j, i); + // CHECK: dpct::experimental::sample_image_by_byte(surf, sycl::float3(k, j, i)); + surf3Dread(surf, k, j, i); + // CHECK: i = dpct::experimental::sample_image_by_byte(surf, sycl::float3(k, j, i)); + surf3Dread(&i, surf, k, j, i); +} int main() { // CHECK: sycl::ext::oneapi::experimental::sampled_image_handle surf; cudaSurfaceObject_t surf;