From 5ded68707789ff817400da26eea42a5735970ab0 Mon Sep 17 00:00:00 2001 From: "Tang, Jiajun" Date: Wed, 16 Oct 2024 10:05:51 +0800 Subject: [PATCH 01/21] [SYCLomatic] Support migration of 3 runtime Surface APIs. Signed-off-by: Tang, Jiajun jiajun.tang@intel.com --- clang/lib/DPCT/RuleInfra/MapNames.cpp | 8 ++++++ clang/lib/DPCT/RulesLang/APINamesTexture.inc | 26 +++++++++++++++++++ clang/lib/DPCT/RulesLang/RulesLang.cpp | 9 +++++++ clang/lib/DPCT/SrcAPI/APINames.inc | 6 ++--- .../dpct-rt/include/dpct/bindless_images.hpp | 2 +- .../texture/surface_object_bindless_image.cu | 16 ++++++++++++ 6 files changed, 63 insertions(+), 4 deletions(-) create mode 100644 clang/test/dpct/texture/surface_object_bindless_image.cu diff --git a/clang/lib/DPCT/RuleInfra/MapNames.cpp b/clang/lib/DPCT/RuleInfra/MapNames.cpp index a08b1fedac39..2b84f371c87f 100644 --- a/clang/lib/DPCT/RuleInfra/MapNames.cpp +++ b/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( + DpctGlobalInfo::useExtBindlessImages() + ? getClNamespace() + + "ext::oneapi::experimental::sampled_image_handle" + : getDpctNamespace() + "image_wrapper_base_p", + HelperFeatureEnum::device_ext)}, {"textureReference", std::make_shared(getDpctNamespace() + "image_wrapper_base", HelperFeatureEnum::device_ext)}, @@ -868,6 +875,7 @@ void MapNames::setExplicitNamespaceMap( "cudaTextureDesc", "cudaResourceDesc", "cudaTextureObject_t", + "cudaSurfaceObject_t", "textureReference", "cudaTextureAddressMode", "cudaTextureFilterMode", diff --git a/clang/lib/DPCT/RulesLang/APINamesTexture.inc b/clang/lib/DPCT/RulesLang/APINamesTexture.inc index ec2139dee9b4..d0ccccd864ee 100644 --- a/clang/lib/DPCT/RulesLang/APINamesTexture.inc +++ b/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) diff --git a/clang/lib/DPCT/RulesLang/RulesLang.cpp b/clang/lib/DPCT/RulesLang/RulesLang.cpp index 2dee0e04406f..b7d6c1ce41f5 100644 --- a/clang/lib/DPCT/RulesLang/RulesLang.cpp +++ b/clang/lib/DPCT/RulesLang/RulesLang.cpp @@ -11957,6 +11957,12 @@ void TextureRule::registerMatcher(MatchFinder &MF) { ) .bind("tex"), this); + MF.addMatcher( + typeLoc( + loc(qualType(hasDeclaration(typedefDecl(hasAnyName( + "cudaTextureObject_t", "cudaSurfaceObject_t", "CUtexObject")))))) + .bind("texObj"), + this); MF.addMatcher(typeLoc(loc(qualType(hasDeclaration(typedefDecl(hasAnyName( "cudaTextureObject_t", "CUtexObject")))))) .bind("texObj"), @@ -12020,6 +12026,9 @@ void TextureRule::registerMatcher(MatchFinder &MF) { "cudaGetTextureObjectResourceDesc", "cudaGetTextureObjectTextureDesc", "cudaGetTextureObjectResourceViewDesc", + "cudaCreateSurfaceObject", + "cudaDestroySurfaceObject", + "cudaGetSurfaceObjectResourceDesc", "cuArray3DCreate_v2", "cuArrayCreate_v2", "cuArrayDestroy", diff --git a/clang/lib/DPCT/SrcAPI/APINames.inc b/clang/lib/DPCT/SrcAPI/APINames.inc index f2fa3be225ca..93cccc99fbde 100644 --- a/clang/lib/DPCT/SrcAPI/APINames.inc +++ b/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") diff --git a/clang/runtime/dpct-rt/include/dpct/bindless_images.hpp b/clang/runtime/dpct-rt/include/dpct/bindless_images.hpp index 0df49d437e6a..0bbd4baa6691 100644 --- a/clang/runtime/dpct-rt/include/dpct/bindless_images.hpp +++ b/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(), diff --git a/clang/test/dpct/texture/surface_object_bindless_image.cu b/clang/test/dpct/texture/surface_object_bindless_image.cu new file mode 100644 index 000000000000..6a34429c63c7 --- /dev/null +++ b/clang/test/dpct/texture/surface_object_bindless_image.cu @@ -0,0 +1,16 @@ +// 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 %} + +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); +} From 1b8d981698958b0b612dc8a82c7f48aa286db82e Mon Sep 17 00:00:00 2001 From: "Chen, Sheng S" Date: Tue, 12 Nov 2024 16:44:56 +0800 Subject: [PATCH 02/21] update Signed-off-by: Chen, Sheng S --- clang/lib/DPCT/AnalysisInfo.cpp | 17 +++++++++++------ clang/lib/DPCT/AnalysisInfo.h | 6 +++--- clang/lib/DPCT/RulesLang/APINamesTexture.inc | 7 +++++++ clang/lib/DPCT/RulesLang/RulesLang.cpp | 8 ++++---- .../lib/Headers/__clang_cuda_runtime_wrapper.h | 1 + .../texture/surface_object_bindless_image.cu | 18 ++++++++++++++++++ 6 files changed, 44 insertions(+), 13 deletions(-) diff --git a/clang/lib/DPCT/AnalysisInfo.cpp b/clang/lib/DPCT/AnalysisInfo.cpp index 32f5754c2116..7f8c5b3abff5 100644 --- a/clang/lib/DPCT/AnalysisInfo.cpp +++ b/clang/lib/DPCT/AnalysisInfo.cpp @@ -3600,8 +3600,12 @@ const std::string MemVarInfo::ExternVariableName = "dpct_local"; std::unordered_map 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 Target) { void TextureObjectInfo::addParamDeclReplacement() { if (Type) { DpctGlobalInfo::getInstance().addReplacement( - std::make_shared(FilePath, Offset, ReplaceTypeLength, + std::make_shared(FilePath, Offset, Type->getTypeLength(), getParamDeclType(), nullptr)); } } @@ -4924,7 +4928,7 @@ void DeviceFunctionDecl::emplaceReplacement() { DpctGlobalInfo::getInstance().addReplacement( std::make_shared( 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(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(qual_type) #define ARG_TYPE_CAST(type) TYPE_CAST(ArgType, type) #define PARM_TYPE_CAST(type) TYPE_CAST(ParmType, type) diff --git a/clang/lib/DPCT/AnalysisInfo.h b/clang/lib/DPCT/AnalysisInfo.h index 9c516406208a..4e9d58aa9393 100644 --- a/clang/lib/DPCT/AnalysisInfo.h +++ b/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; diff --git a/clang/lib/DPCT/RulesLang/APINamesTexture.inc b/clang/lib/DPCT/RulesLang/APINamesTexture.inc index d0ccccd864ee..c6c2023120c0 100644 --- a/clang/lib/DPCT/RulesLang/APINamesTexture.inc +++ b/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)) diff --git a/clang/lib/DPCT/RulesLang/RulesLang.cpp b/clang/lib/DPCT/RulesLang/RulesLang.cpp index b7d6c1ce41f5..81fc3ca48e61 100644 --- a/clang/lib/DPCT/RulesLang/RulesLang.cpp +++ b/clang/lib/DPCT/RulesLang/RulesLang.cpp @@ -11963,10 +11963,7 @@ 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( @@ -12021,6 +12018,9 @@ void TextureRule::registerMatcher(MatchFinder &MF) { "tex1Dfetch", "tex1DLayered", "tex2DLayered", + "surf1Dread", + "surf2Dread", + "surf3Dread", "cudaCreateTextureObject", "cudaDestroyTextureObject", "cudaGetTextureObjectResourceDesc", diff --git a/clang/lib/Headers/__clang_cuda_runtime_wrapper.h b/clang/lib/Headers/__clang_cuda_runtime_wrapper.h index 9c5c09fc405d..97b790db822a 100644 --- a/clang/lib/Headers/__clang_cuda_runtime_wrapper.h +++ b/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__") diff --git a/clang/test/dpct/texture/surface_object_bindless_image.cu b/clang/test/dpct/texture/surface_object_bindless_image.cu index 6a34429c63c7..9bb3309076b0 100644 --- a/clang/test/dpct/texture/surface_object_bindless_image.cu +++ b/clang/test/dpct/texture/surface_object_bindless_image.cu @@ -2,6 +2,24 @@ // 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 __global__ void kernel(cudaSurfaceObject_t surf) { + int i; + float j, k, l, m; + // CHECK: sycl::ext::oneapi::experimental::sample_image(surf, float(i)); + surf1Dread(surf, i); + // CHECK: i = sycl::ext::oneapi::experimental::sample_image(surf, float(i)); + surf1Dread(&i, surf, i); + // CHECK: i = sycl::ext::oneapi::experimental::sample_image(surf, sycl::float2(j, i)); + surf2Dread(surf, j, i); + // CHECK: sycl::ext::oneapi::experimental::sample_image(surf, sycl::float2(j, i)); + surf2Dread(&i, surf, j, i); + // CHECK: sycl::ext::oneapi::experimental::sample_image(surf, sycl::float3(k, j, i)); + surf3Dread(surf, k, j, i); + // CHECK: i = sycl::ext::oneapi::experimental::sample_image(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; From 778e2c8dae573ba18970443e7c6169ce1fb53124 Mon Sep 17 00:00:00 2001 From: "Chen, Sheng S" Date: Tue, 12 Nov 2024 21:12:14 +0800 Subject: [PATCH 03/21] update Signed-off-by: Chen, Sheng S --- clang/lib/Headers/__clang_cuda_runtime_wrapper.h | 2 ++ 1 file changed, 2 insertions(+) diff --git a/clang/lib/Headers/__clang_cuda_runtime_wrapper.h b/clang/lib/Headers/__clang_cuda_runtime_wrapper.h index 97b790db822a..f96be6961ece 100644 --- a/clang/lib/Headers/__clang_cuda_runtime_wrapper.h +++ b/clang/lib/Headers/__clang_cuda_runtime_wrapper.h @@ -425,7 +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__") From bbeb07fc1fc0f97f49a08dc916a8ef455741a903 Mon Sep 17 00:00:00 2001 From: "Chen, Sheng S" Date: Mon, 18 Nov 2024 16:46:26 +0800 Subject: [PATCH 04/21] update Signed-off-by: Chen, Sheng S --- clang/lib/DPCT/AnalysisInfo.cpp | 10 ++++------ clang/lib/DPCT/AnalysisInfo.h | 4 ++-- 2 files changed, 6 insertions(+), 8 deletions(-) diff --git a/clang/lib/DPCT/AnalysisInfo.cpp b/clang/lib/DPCT/AnalysisInfo.cpp index c378629ec4b1..ab82e1bf7306 100644 --- a/clang/lib/DPCT/AnalysisInfo.cpp +++ b/clang/lib/DPCT/AnalysisInfo.cpp @@ -3600,12 +3600,8 @@ const std::string MemVarInfo::ExternVariableName = "dpct_local"; std::unordered_map 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; @@ -3795,7 +3791,7 @@ void TextureObjectInfo::merge(std::shared_ptr Target) { void TextureObjectInfo::addParamDeclReplacement() { if (Type) { DpctGlobalInfo::getInstance().addReplacement( - std::make_shared(FilePath, Offset, Type->getTypeLength(), + std::make_shared(FilePath, Offset, ReplaceTypeLength, getParamDeclType(), nullptr)); } } @@ -4928,7 +4924,7 @@ void DeviceFunctionDecl::emplaceReplacement() { DpctGlobalInfo::getInstance().addReplacement( std::make_shared( Obj->getFilePath(), Obj->getOffset(), - Obj->getType()->getTypeLength(), + strlen("cudaTextureObject_t"), MapNames::getClNamespace() + "ext::oneapi::experimental::sampled_image_handle", nullptr)); @@ -6590,6 +6586,8 @@ 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(qual_type) #define ARG_TYPE_CAST(type) TYPE_CAST(ArgType, type) #define PARM_TYPE_CAST(type) TYPE_CAST(ParmType, type) diff --git a/clang/lib/DPCT/AnalysisInfo.h b/clang/lib/DPCT/AnalysisInfo.h index aef688e3ee38..33f0d9596b3b 100644 --- a/clang/lib/DPCT/AnalysisInfo.h +++ b/clang/lib/DPCT/AnalysisInfo.h @@ -1991,13 +1991,11 @@ 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(); @@ -2055,6 +2053,8 @@ class TextureInfo { // 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; From 1876e89b844b49e82682b9f2fce33253f0fc2f8c Mon Sep 17 00:00:00 2001 From: "Chen, Sheng S" Date: Tue, 19 Nov 2024 13:41:29 +0800 Subject: [PATCH 05/21] up Signed-off-by: Chen, Sheng S --- clang/lib/DPCT/AnalysisInfo.h | 2 +- .../RulesLang/CallExprRewriterTexture.cpp | 40 ++++++++++++------- clang/lib/DPCT/SrcAPI/APINames.inc | 6 +-- clang/test/dpct/surface.cu | 34 ++++++++++++++++ .../texture/surface_object_bindless_image.cu | 13 +++--- 5 files changed, 70 insertions(+), 25 deletions(-) create mode 100644 clang/test/dpct/surface.cu diff --git a/clang/lib/DPCT/AnalysisInfo.h b/clang/lib/DPCT/AnalysisInfo.h index 33f0d9596b3b..a404a814bd97 100644 --- a/clang/lib/DPCT/AnalysisInfo.h +++ b/clang/lib/DPCT/AnalysisInfo.h @@ -2051,7 +2051,7 @@ class TextureInfo { bool isUseHelperFunc() { return true; } }; -// texture object info can be used for CUDA texture and suface objects. +// 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/CallExprRewriterTexture.cpp b/clang/lib/DPCT/RulesLang/CallExprRewriterTexture.cpp index ad0c9bb3f06e..5a8790cd6ec9 100644 --- a/clang/lib/DPCT/RulesLang/CallExprRewriterTexture.cpp +++ b/clang/lib/DPCT/RulesLang/CallExprRewriterTexture.cpp @@ -21,8 +21,10 @@ class TextureReadRewriterFactory : public CallExprRewriterFactoryBase { template std::shared_ptr - createRewriter(const CallExpr *C, bool RetAssign, BaseT Base) const { - const static std::string MemberName = "read"; + createRewriter(const CallExpr *C, bool RetAssign, bool IsSurfAPI, + BaseT Base) const { + const static std::string MemberName = IsSurfAPI ? "read_byte" : "read"; + using ReaderPrinter = decltype(makeMemberCallCreator( std::declval>(), false, MemberName, makeCallArgCreatorWithCall(Idx)...)(C)); @@ -41,10 +43,14 @@ class TextureReadRewriterFactory : public CallExprRewriterFactoryBase { template std::shared_ptr createbindlessRewriterNormal(const CallExpr *C, bool RetAssign, - const TemplateArgumentInfo &TAI, + bool IsSurfAPI, const TemplateArgumentInfo &TAI, const std::string &VecTypeName) const { - const static std::string FuncName = - MapNames::getClNamespace() + "ext::oneapi::experimental::sample_image"; + const static std::string FuncName = [=]() -> std::string { + if (IsSurfAPI) + return "syclcompat::experimental::sample_image"; + return MapNames::getClNamespace() + + "ext::oneapi::experimental::sample_image"; + }(); using FuncNamePrinter = TemplatedNamePrinter>; using ReaderPrinter = @@ -66,7 +72,7 @@ class TextureReadRewriterFactory : public CallExprRewriterFactoryBase { template std::shared_ptr - createbindlessRewriterLod(const CallExpr *C, bool RetAssign, + createbindlessRewriterLod(const CallExpr *C, bool RetAssign, bool IsSurfAPI, const TemplateArgumentInfo &TAI, const std::string &VecTypeName) const { const static std::string FuncName = @@ -96,7 +102,7 @@ class TextureReadRewriterFactory : public CallExprRewriterFactoryBase { template std::shared_ptr createbindlessRewriterLayered(const CallExpr *C, bool RetAssign, - const TemplateArgumentInfo &TAI, + bool IsSurfAPI, const TemplateArgumentInfo &TAI, const std::string &VecTypeName) const { const static std::string FuncName = MapNames::getClNamespace() + @@ -124,7 +130,7 @@ class TextureReadRewriterFactory : public CallExprRewriterFactoryBase { } std::shared_ptr - createbindlessRewriter(const CallExpr *C, bool RetAssign, + createbindlessRewriter(const CallExpr *C, bool RetAssign, bool IsSurfAPI, QualType TargetType) const { TemplateArgumentInfo TAI; auto TAL = getTemplateArgsList(C); @@ -151,11 +157,12 @@ class TextureReadRewriterFactory : public CallExprRewriterFactoryBase { CallExprPrintergetArg(Idx)))...>; if ((TexType & 0xf0) == 0x10) - return createbindlessRewriterLod(C, RetAssign, TAI, VecTypeName); + return createbindlessRewriterLod(C, RetAssign, IsSurfAPI, TAI, + VecTypeName); if ((TexType & 0xf0) == 0xf0) - return createbindlessRewriterLayered(C, RetAssign, TAI, - VecTypeName); - return createbindlessRewriterNormal(C, RetAssign, TAI, + return createbindlessRewriterLayered(C, RetAssign, IsSurfAPI, + TAI, VecTypeName); + return createbindlessRewriterNormal(C, RetAssign, IsSurfAPI, TAI, VecTypeName); } @@ -169,6 +176,9 @@ class TextureReadRewriterFactory : public CallExprRewriterFactoryBase { QualType TargetType = Call->getType(); StringRef SourceName; bool RetAssign = false; + bool IsSurfAPI = false; + if (Source.find("surf") != std::string::npos) + IsSurfAPI = true; if (SourceExpr->getType()->isPointerType()) { TargetType = SourceExpr->getType()->getPointeeType(); SourceExpr = Call->getArg(1); @@ -181,7 +191,7 @@ class TextureReadRewriterFactory : public CallExprRewriterFactoryBase { } } if (DpctGlobalInfo::useExtBindlessImages()) { - return createbindlessRewriter(Call, RetAssign, TargetType); + return createbindlessRewriter(Call, RetAssign, IsSurfAPI, TargetType); } SourceExpr = SourceExpr->IgnoreImpCasts(); if (auto FD = DpctGlobalInfo::getParentFunction(Call)) { @@ -196,7 +206,7 @@ class TextureReadRewriterFactory : public CallExprRewriterFactoryBase { MemberInfo->setType( DpctGlobalInfo::getUnqualifiedTypeName(TargetType), TexType); SourceName = MemberInfo->getName(); - return createRewriter(Call, RetAssign, SourceName); + return createRewriter(Call, RetAssign, IsSurfAPI, SourceName); } } else if (auto DRE = dyn_cast(SourceExpr)) { auto CallDefRange = @@ -213,7 +223,7 @@ class TextureReadRewriterFactory : public CallExprRewriterFactoryBase { } } - return createRewriter(Call, RetAssign, + return createRewriter(Call, RetAssign, IsSurfAPI, std::make_pair(Call, Call->getArg(RetAssign & 0x01))); } }; 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/test/dpct/surface.cu b/clang/test/dpct/surface.cu new file mode 100644 index 000000000000..a868027a2418 --- /dev/null +++ b/clang/test/dpct/surface.cu @@ -0,0 +1,34 @@ +// 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 +// RUN: %if build_lit %{icpx -c -fsycl %T/surface/surface.dp.cpp -o %T/surface/surface.dp.o %} + +// 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); +} + +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 5dc83bdbb8e0..790cb33a07ca 100644 --- a/clang/test/dpct/texture/surface_object_bindless_image.cu +++ b/clang/test/dpct/texture/surface_object_bindless_image.cu @@ -2,20 +2,21 @@ // 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: sycl::ext::oneapi::experimental::sample_image(surf, float(i)); + // CHECK: syclcompat::experimental::sample_image(surf, float(i)); surf1Dread(surf, i); - // CHECK: i = sycl::ext::oneapi::experimental::sample_image(surf, float(i)); + // CHECK: i = syclcompat::experimental::sample_image(surf, float(i)); surf1Dread(&i, surf, i); - // CHECK: i = sycl::ext::oneapi::experimental::sample_image(surf, sycl::float2(j, i)); + // CHECK: syclcompat::experimental::sample_image(surf, sycl::float2(j, i)); surf2Dread(surf, j, i); - // CHECK: sycl::ext::oneapi::experimental::sample_image(surf, sycl::float2(j, i)); + // CHECK: i = syclcompat::experimental::sample_image(surf, sycl::float2(j, i)); surf2Dread(&i, surf, j, i); - // CHECK: sycl::ext::oneapi::experimental::sample_image(surf, sycl::float3(k, j, i)); + // CHECK: syclcompat::experimental::sample_image(surf, sycl::float3(k, j, i)); surf3Dread(surf, k, j, i); - // CHECK: i = sycl::ext::oneapi::experimental::sample_image(surf, sycl::float3(k, j, i)); + // CHECK: i = syclcompat::experimental::sample_image(surf, sycl::float3(k, j, i)); surf3Dread(&i, surf, k, j, i); } From 0fb7b46d7a4b9c13c42fb3d03291fd1611f934e5 Mon Sep 17 00:00:00 2001 From: "Chen, Sheng S" Date: Tue, 19 Nov 2024 16:50:42 +0800 Subject: [PATCH 06/21] up Signed-off-by: Chen, Sheng S --- .../DPCT/RulesLang/CallExprRewriterTexture.cpp | 2 +- .../dpct-rt/include/dpct/bindless_images.hpp | 16 +++++++++++++++- 2 files changed, 16 insertions(+), 2 deletions(-) diff --git a/clang/lib/DPCT/RulesLang/CallExprRewriterTexture.cpp b/clang/lib/DPCT/RulesLang/CallExprRewriterTexture.cpp index 5a8790cd6ec9..1accba0d0b79 100644 --- a/clang/lib/DPCT/RulesLang/CallExprRewriterTexture.cpp +++ b/clang/lib/DPCT/RulesLang/CallExprRewriterTexture.cpp @@ -47,7 +47,7 @@ class TextureReadRewriterFactory : public CallExprRewriterFactoryBase { const std::string &VecTypeName) const { const static std::string FuncName = [=]() -> std::string { if (IsSurfAPI) - return "syclcompat::experimental::sample_image"; + return MapNames::getDpctNamespace() + "experimental::sample_image"; return MapNames::getClNamespace() + "ext::oneapi::experimental::sample_image"; }(); diff --git a/clang/runtime/dpct-rt/include/dpct/bindless_images.hpp b/clang/runtime/dpct-rt/include/dpct/bindless_images.hpp index 0bbd4baa6691..fb1f8b2ed91c 100644 --- a/clang/runtime/dpct-rt/include/dpct/bindless_images.hpp +++ b/clang/runtime/dpct-rt/include/dpct/bindless_images.hpp @@ -19,7 +19,7 @@ namespace dpct { namespace experimental { -#ifdef SYCL_EXT_ONEAPI_BINDLESS_IMAGES +#ifndef SYCL_EXT_ONEAPI_BINDLESS_IMAGES /// The wrapper class of bindless image memory handle. class image_mem_wrapper { @@ -1352,6 +1352,20 @@ static inline void dpct_memcpy(image_mem_wrapper *dest, size_t w_offset_dest, sycl::free(temp, q); } +template +DataT sample_image( + 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 *; From 24e392a952f72f20b9f58722f15e9bef842bf5c9 Mon Sep 17 00:00:00 2001 From: "Chen, Sheng S" Date: Wed, 20 Nov 2024 08:56:29 +0800 Subject: [PATCH 07/21] add comment Signed-off-by: Chen, Sheng S --- clang/runtime/dpct-rt/include/dpct/bindless_images.hpp | 1 + 1 file changed, 1 insertion(+) diff --git a/clang/runtime/dpct-rt/include/dpct/bindless_images.hpp b/clang/runtime/dpct-rt/include/dpct/bindless_images.hpp index fb1f8b2ed91c..83c80710bd7e 100644 --- a/clang/runtime/dpct-rt/include/dpct/bindless_images.hpp +++ b/clang/runtime/dpct-rt/include/dpct/bindless_images.hpp @@ -1352,6 +1352,7 @@ 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( const sycl::ext::oneapi::experimental::sampled_image_handle &imageHandle, From f3c9c4f7b2653b6cdbe3fc00d4a8e3ae5585b88c Mon Sep 17 00:00:00 2001 From: "Chen, Sheng S" Date: Thu, 21 Nov 2024 10:36:12 +0800 Subject: [PATCH 08/21] resolve the comment. Signed-off-by: Chen, Sheng S --- clang/lib/DPCT/RulesLang/CallExprRewriterTexture.cpp | 2 +- .../dpct/texture/surface_object_bindless_image.cu | 12 ++++++------ 2 files changed, 7 insertions(+), 7 deletions(-) diff --git a/clang/lib/DPCT/RulesLang/CallExprRewriterTexture.cpp b/clang/lib/DPCT/RulesLang/CallExprRewriterTexture.cpp index 1accba0d0b79..feb76d5e01a9 100644 --- a/clang/lib/DPCT/RulesLang/CallExprRewriterTexture.cpp +++ b/clang/lib/DPCT/RulesLang/CallExprRewriterTexture.cpp @@ -47,7 +47,7 @@ class TextureReadRewriterFactory : public CallExprRewriterFactoryBase { const std::string &VecTypeName) const { const static std::string FuncName = [=]() -> std::string { if (IsSurfAPI) - return MapNames::getDpctNamespace() + "experimental::sample_image"; + return MapNames::getLibraryHelperNamespace() + "experimental::sample_image"; return MapNames::getClNamespace() + "ext::oneapi::experimental::sample_image"; }(); diff --git a/clang/test/dpct/texture/surface_object_bindless_image.cu b/clang/test/dpct/texture/surface_object_bindless_image.cu index 790cb33a07ca..56ea491f31e1 100644 --- a/clang/test/dpct/texture/surface_object_bindless_image.cu +++ b/clang/test/dpct/texture/surface_object_bindless_image.cu @@ -6,17 +6,17 @@ template __global__ void kernel(cudaSurfaceObject_t surf) { int i; float j, k, l, m; - // CHECK: syclcompat::experimental::sample_image(surf, float(i)); + // CHECK: dpct::experimental::sample_image(surf, float(i)); surf1Dread(surf, i); - // CHECK: i = syclcompat::experimental::sample_image(surf, float(i)); + // CHECK: i = dpct::experimental::sample_image(surf, float(i)); surf1Dread(&i, surf, i); - // CHECK: syclcompat::experimental::sample_image(surf, sycl::float2(j, i)); + // CHECK: dpct::experimental::sample_image(surf, sycl::float2(j, i)); surf2Dread(surf, j, i); - // CHECK: i = syclcompat::experimental::sample_image(surf, sycl::float2(j, i)); + // CHECK: i = dpct::experimental::sample_image(surf, sycl::float2(j, i)); surf2Dread(&i, surf, j, i); - // CHECK: syclcompat::experimental::sample_image(surf, sycl::float3(k, j, i)); + // CHECK: dpct::experimental::sample_image(surf, sycl::float3(k, j, i)); surf3Dread(surf, k, j, i); - // CHECK: i = syclcompat::experimental::sample_image(surf, sycl::float3(k, j, i)); + // CHECK: i = dpct::experimental::sample_image(surf, sycl::float3(k, j, i)); surf3Dread(&i, surf, k, j, i); } From dedb20be3a76b8e4ef75eeaab73ddd01026a38cb Mon Sep 17 00:00:00 2001 From: "Chen, Sheng S" Date: Fri, 22 Nov 2024 09:07:12 +0800 Subject: [PATCH 09/21] fix comment Signed-off-by: Chen, Sheng S --- .../RulesLang/CallExprRewriterTexture.cpp | 35 +++++++++---------- .../dpct-rt/include/dpct/bindless_images.hpp | 2 +- clang/runtime/dpct-rt/include/dpct/image.hpp | 20 ++--------- 3 files changed, 20 insertions(+), 37 deletions(-) diff --git a/clang/lib/DPCT/RulesLang/CallExprRewriterTexture.cpp b/clang/lib/DPCT/RulesLang/CallExprRewriterTexture.cpp index feb76d5e01a9..d4892a882643 100644 --- a/clang/lib/DPCT/RulesLang/CallExprRewriterTexture.cpp +++ b/clang/lib/DPCT/RulesLang/CallExprRewriterTexture.cpp @@ -19,10 +19,9 @@ class TextureReadRewriterFactory : public CallExprRewriterFactoryBase { inline int getDim() const { return TexType & 0x0f; } - template + template std::shared_ptr - createRewriter(const CallExpr *C, bool RetAssign, bool IsSurfAPI, - BaseT Base) const { + createRewriter(const CallExpr *C, bool RetAssign, BaseT Base) const { const static std::string MemberName = IsSurfAPI ? "read_byte" : "read"; using ReaderPrinter = decltype(makeMemberCallCreator( @@ -40,10 +39,10 @@ class TextureReadRewriterFactory : public CallExprRewriterFactoryBase { std::make_pair(C, C->getArg(Idx))...); } - template + template std::shared_ptr createbindlessRewriterNormal(const CallExpr *C, bool RetAssign, - bool IsSurfAPI, const TemplateArgumentInfo &TAI, + const TemplateArgumentInfo &TAI, const std::string &VecTypeName) const { const static std::string FuncName = [=]() -> std::string { if (IsSurfAPI) @@ -70,9 +69,9 @@ class TextureReadRewriterFactory : public CallExprRewriterFactoryBase { VecType(VecTypeName, std::make_pair(C, C->getArg(Idx))...)); } - template + template std::shared_ptr - createbindlessRewriterLod(const CallExpr *C, bool RetAssign, bool IsSurfAPI, + createbindlessRewriterLod(const CallExpr *C, bool RetAssign, const TemplateArgumentInfo &TAI, const std::string &VecTypeName) const { const static std::string FuncName = @@ -99,10 +98,10 @@ class TextureReadRewriterFactory : public CallExprRewriterFactoryBase { std::make_pair(C, C->getArg(C->getNumArgs() - 1))); } - template + template std::shared_ptr createbindlessRewriterLayered(const CallExpr *C, bool RetAssign, - bool IsSurfAPI, const TemplateArgumentInfo &TAI, + const TemplateArgumentInfo &TAI, const std::string &VecTypeName) const { const static std::string FuncName = MapNames::getClNamespace() + @@ -128,9 +127,9 @@ class TextureReadRewriterFactory : public CallExprRewriterFactoryBase { VecType(VecTypeName, std::make_pair(C, C->getArg(Idx))...), std::make_pair(C, C->getArg(C->getNumArgs() - 1))); } - + template std::shared_ptr - createbindlessRewriter(const CallExpr *C, bool RetAssign, bool IsSurfAPI, + createbindlessRewriter(const CallExpr *C, bool RetAssign, QualType TargetType) const { TemplateArgumentInfo TAI; auto TAL = getTemplateArgsList(C); @@ -157,12 +156,12 @@ class TextureReadRewriterFactory : public CallExprRewriterFactoryBase { CallExprPrintergetArg(Idx)))...>; if ((TexType & 0xf0) == 0x10) - return createbindlessRewriterLod(C, RetAssign, IsSurfAPI, TAI, + return createbindlessRewriterLod(C, RetAssign, TAI, VecTypeName); if ((TexType & 0xf0) == 0xf0) - return createbindlessRewriterLayered(C, RetAssign, IsSurfAPI, + return createbindlessRewriterLayered(C, RetAssign, TAI, VecTypeName); - return createbindlessRewriterNormal(C, RetAssign, IsSurfAPI, TAI, + return createbindlessRewriterNormal(C, RetAssign, TAI, VecTypeName); } @@ -191,7 +190,7 @@ class TextureReadRewriterFactory : public CallExprRewriterFactoryBase { } } if (DpctGlobalInfo::useExtBindlessImages()) { - return createbindlessRewriter(Call, RetAssign, IsSurfAPI, TargetType); + return createbindlessRewriter(Call, RetAssign, TargetType); } SourceExpr = SourceExpr->IgnoreImpCasts(); if (auto FD = DpctGlobalInfo::getParentFunction(Call)) { @@ -206,7 +205,7 @@ class TextureReadRewriterFactory : public CallExprRewriterFactoryBase { MemberInfo->setType( DpctGlobalInfo::getUnqualifiedTypeName(TargetType), TexType); SourceName = MemberInfo->getName(); - return createRewriter(Call, RetAssign, IsSurfAPI, SourceName); + return createRewriter(Call, RetAssign, SourceName); } } else if (auto DRE = dyn_cast(SourceExpr)) { auto CallDefRange = @@ -223,8 +222,8 @@ class TextureReadRewriterFactory : public CallExprRewriterFactoryBase { } } - return createRewriter(Call, RetAssign, IsSurfAPI, - std::make_pair(Call, Call->getArg(RetAssign & 0x01))); + return createRewriter( + Call, RetAssign, std::make_pair(Call, Call->getArg(RetAssign & 0x01))); } }; diff --git a/clang/runtime/dpct-rt/include/dpct/bindless_images.hpp b/clang/runtime/dpct-rt/include/dpct/bindless_images.hpp index 83c80710bd7e..9d485dc82175 100644 --- a/clang/runtime/dpct-rt/include/dpct/bindless_images.hpp +++ b/clang/runtime/dpct-rt/include/dpct/bindless_images.hpp @@ -1354,7 +1354,7 @@ static inline void dpct_memcpy(image_mem_wrapper *dest, size_t w_offset_dest, // A wrapper for sycl sample_image function for the byte addressing image. template -DataT sample_image( +DataT sample_image_by_byte( const sycl::ext::oneapi::experimental::sampled_image_handle &imageHandle, CoordT &&coords) { if constexpr (std::is_scalar_v) { diff --git a/clang/runtime/dpct-rt/include/dpct/image.hpp b/clang/runtime/dpct-rt/include/dpct/image.hpp index 9d0820c5f3db..1c58c18b5aef 100644 --- a/clang/runtime/dpct-rt/include/dpct/image.hpp +++ b/clang/runtime/dpct-rt/include/dpct/image.hpp @@ -894,13 +894,6 @@ class image_accessor_ext { return detail::fetch_data()(_img_acc.read(x, _sampler)); } - /// Read data from accessor. - template - typename std::enable_if::type read_byte(float x, float y, - float z) { - return detail::fetch_data()( - _img_acc.read(sycl::float4(x / sizeof(T), y, z, 0), _sampler)); - } /// Read data from accessor. template ()( _img_acc.read(sycl::int4(x / sizeof(T), y, z, 0), _sampler)); } - /// Read data from accessor. - template - typename std::enable_if::type read_byte(float x, float y) { - return detail::fetch_data()( - _img_acc.read(sycl::float2(x / sizeof(T), y), _sampler)); - } + /// Read data from accessor. template ()( _img_acc.read(sycl::int2(x / sizeof(T), y), _sampler)); } - /// Read data from accessor. - template - typename std::enable_if::type read_byte(float x) { - return detail::fetch_data()(_img_acc.read(x / sizeof(T), _sampler)); - } + /// Read data from accessor. template ::value> From 7171b24bfcc5fda70c10f997b73bfaecaea36f04 Mon Sep 17 00:00:00 2001 From: "Chen, Sheng S" Date: Fri, 22 Nov 2024 10:23:55 +0800 Subject: [PATCH 10/21] up Signed-off-by: Chen, Sheng S --- clang/lib/DPCT/AnalysisInfo.cpp | 2 +- .../RulesLang/CallExprRewriterTexture.cpp | 43 ++++++++++--------- 2 files changed, 23 insertions(+), 22 deletions(-) diff --git a/clang/lib/DPCT/AnalysisInfo.cpp b/clang/lib/DPCT/AnalysisInfo.cpp index ab82e1bf7306..bde445cbd0cd 100644 --- a/clang/lib/DPCT/AnalysisInfo.cpp +++ b/clang/lib/DPCT/AnalysisInfo.cpp @@ -8,9 +8,9 @@ #include "AnalysisInfo.h" #include "Diagnostics/Diagnostics.h" +#include "MigrationReport/Statics.h" #include "RuleInfra/ExprAnalysis.h" #include "RuleInfra/MapNames.h" -#include "MigrationReport/Statics.h" #include "TextModification.h" #include "Utility.h" diff --git a/clang/lib/DPCT/RulesLang/CallExprRewriterTexture.cpp b/clang/lib/DPCT/RulesLang/CallExprRewriterTexture.cpp index d4892a882643..ad63c35c5b08 100644 --- a/clang/lib/DPCT/RulesLang/CallExprRewriterTexture.cpp +++ b/clang/lib/DPCT/RulesLang/CallExprRewriterTexture.cpp @@ -19,11 +19,12 @@ class TextureReadRewriterFactory : public CallExprRewriterFactoryBase { inline int getDim() const { return TexType & 0x0f; } - template + template std::shared_ptr - createRewriter(const CallExpr *C, bool RetAssign, BaseT Base) const { - const static std::string MemberName = IsSurfAPI ? "read_byte" : "read"; - + createRewriter(const CallExpr *C, bool RetAssign, bool IsSurfAPI, + BaseT Base) const { + static std::string MemberName = "read"; + MemberName = IsSurfAPI ? "read_byte" : "read"; using ReaderPrinter = decltype(makeMemberCallCreator( std::declval>(), false, MemberName, makeCallArgCreatorWithCall(Idx)...)(C)); @@ -39,14 +40,16 @@ class TextureReadRewriterFactory : public CallExprRewriterFactoryBase { std::make_pair(C, C->getArg(Idx))...); } - template + template std::shared_ptr createbindlessRewriterNormal(const CallExpr *C, bool RetAssign, - const TemplateArgumentInfo &TAI, + bool IsSurfAPI, const TemplateArgumentInfo &TAI, const std::string &VecTypeName) const { - const static std::string FuncName = [=]() -> std::string { + static std::string FuncName = ""; + FuncName = [=]() -> std::string { if (IsSurfAPI) - return MapNames::getLibraryHelperNamespace() + "experimental::sample_image"; + return MapNames::getLibraryHelperNamespace() + + "experimental::sample_image_by_byte"; return MapNames::getClNamespace() + "ext::oneapi::experimental::sample_image"; }(); @@ -69,7 +72,7 @@ class TextureReadRewriterFactory : public CallExprRewriterFactoryBase { VecType(VecTypeName, std::make_pair(C, C->getArg(Idx))...)); } - template + template std::shared_ptr createbindlessRewriterLod(const CallExpr *C, bool RetAssign, const TemplateArgumentInfo &TAI, @@ -98,7 +101,7 @@ class TextureReadRewriterFactory : public CallExprRewriterFactoryBase { std::make_pair(C, C->getArg(C->getNumArgs() - 1))); } - template + template std::shared_ptr createbindlessRewriterLayered(const CallExpr *C, bool RetAssign, const TemplateArgumentInfo &TAI, @@ -127,9 +130,8 @@ class TextureReadRewriterFactory : public CallExprRewriterFactoryBase { VecType(VecTypeName, std::make_pair(C, C->getArg(Idx))...), std::make_pair(C, C->getArg(C->getNumArgs() - 1))); } - template std::shared_ptr - createbindlessRewriter(const CallExpr *C, bool RetAssign, + createbindlessRewriter(const CallExpr *C, bool RetAssign, bool IsSurfAPI, QualType TargetType) const { TemplateArgumentInfo TAI; auto TAL = getTemplateArgsList(C); @@ -156,12 +158,11 @@ class TextureReadRewriterFactory : public CallExprRewriterFactoryBase { CallExprPrintergetArg(Idx)))...>; if ((TexType & 0xf0) == 0x10) - return createbindlessRewriterLod(C, RetAssign, TAI, - VecTypeName); + return createbindlessRewriterLod(C, RetAssign, TAI, VecTypeName); if ((TexType & 0xf0) == 0xf0) - return createbindlessRewriterLayered(C, RetAssign, - TAI, VecTypeName); - return createbindlessRewriterNormal(C, RetAssign, TAI, + return createbindlessRewriterLayered(C, RetAssign, TAI, + VecTypeName); + return createbindlessRewriterNormal(C, RetAssign, IsSurfAPI, TAI, VecTypeName); } @@ -190,7 +191,7 @@ class TextureReadRewriterFactory : public CallExprRewriterFactoryBase { } } if (DpctGlobalInfo::useExtBindlessImages()) { - return createbindlessRewriter(Call, RetAssign, TargetType); + return createbindlessRewriter(Call, RetAssign, IsSurfAPI, TargetType); } SourceExpr = SourceExpr->IgnoreImpCasts(); if (auto FD = DpctGlobalInfo::getParentFunction(Call)) { @@ -205,7 +206,7 @@ class TextureReadRewriterFactory : public CallExprRewriterFactoryBase { MemberInfo->setType( DpctGlobalInfo::getUnqualifiedTypeName(TargetType), TexType); SourceName = MemberInfo->getName(); - return createRewriter(Call, RetAssign, SourceName); + return createRewriter(Call, RetAssign, IsSurfAPI, SourceName); } } else if (auto DRE = dyn_cast(SourceExpr)) { auto CallDefRange = @@ -222,8 +223,8 @@ class TextureReadRewriterFactory : public CallExprRewriterFactoryBase { } } - return createRewriter( - Call, RetAssign, std::make_pair(Call, Call->getArg(RetAssign & 0x01))); + return createRewriter(Call, RetAssign, IsSurfAPI, + std::make_pair(Call, Call->getArg(RetAssign & 0x01))); } }; From db6ed2fe776a4002455f981b801649e17ff9749c Mon Sep 17 00:00:00 2001 From: "Chen, Sheng S" Date: Fri, 22 Nov 2024 13:23:59 +0800 Subject: [PATCH 11/21] update Signed-off-by: Chen, Sheng S --- .../lib/DPCT/RulesLang/CallExprRewriterTexture.cpp | 13 ++++++------- .../dpct/texture/surface_object_bindless_image.cu | 12 ++++++------ 2 files changed, 12 insertions(+), 13 deletions(-) diff --git a/clang/lib/DPCT/RulesLang/CallExprRewriterTexture.cpp b/clang/lib/DPCT/RulesLang/CallExprRewriterTexture.cpp index ad63c35c5b08..d5af62176167 100644 --- a/clang/lib/DPCT/RulesLang/CallExprRewriterTexture.cpp +++ b/clang/lib/DPCT/RulesLang/CallExprRewriterTexture.cpp @@ -45,8 +45,7 @@ class TextureReadRewriterFactory : public CallExprRewriterFactoryBase { createbindlessRewriterNormal(const CallExpr *C, bool RetAssign, bool IsSurfAPI, const TemplateArgumentInfo &TAI, const std::string &VecTypeName) const { - static std::string FuncName = ""; - FuncName = [=]() -> std::string { + const std::string FuncName = [=]() -> std::string { if (IsSurfAPI) return MapNames::getLibraryHelperNamespace() + "experimental::sample_image_by_byte"; @@ -54,7 +53,7 @@ class TextureReadRewriterFactory : public CallExprRewriterFactoryBase { "ext::oneapi::experimental::sample_image"; }(); using FuncNamePrinter = - TemplatedNamePrinter>; + TemplatedNamePrinter>; using ReaderPrinter = CallExprPrinter, VecType>; @@ -77,10 +76,10 @@ class TextureReadRewriterFactory : public CallExprRewriterFactoryBase { createbindlessRewriterLod(const CallExpr *C, bool RetAssign, const TemplateArgumentInfo &TAI, const std::string &VecTypeName) const { - const static std::string FuncName = + const std::string FuncName = MapNames::getClNamespace() + "ext::oneapi::experimental::sample_mipmap"; using FuncNamePrinter = - TemplatedNamePrinter>; + TemplatedNamePrinter>; using ReaderPrinter = CallExprPrinter, VecType, @@ -106,11 +105,11 @@ class TextureReadRewriterFactory : public CallExprRewriterFactoryBase { createbindlessRewriterLayered(const CallExpr *C, bool RetAssign, const TemplateArgumentInfo &TAI, const std::string &VecTypeName) const { - const static std::string FuncName = + const std::string FuncName = MapNames::getClNamespace() + "ext::oneapi::experimental::sample_image_array"; using FuncNamePrinter = - TemplatedNamePrinter>; + TemplatedNamePrinter>; using ReaderPrinter = CallExprPrinter, VecType, diff --git a/clang/test/dpct/texture/surface_object_bindless_image.cu b/clang/test/dpct/texture/surface_object_bindless_image.cu index 56ea491f31e1..059a1d10a59c 100644 --- a/clang/test/dpct/texture/surface_object_bindless_image.cu +++ b/clang/test/dpct/texture/surface_object_bindless_image.cu @@ -6,17 +6,17 @@ template __global__ void kernel(cudaSurfaceObject_t surf) { int i; float j, k, l, m; - // CHECK: dpct::experimental::sample_image(surf, float(i)); + // CHECK: dpct::experimental::sample_image_by_byte(surf, float(i)); surf1Dread(surf, i); - // CHECK: i = dpct::experimental::sample_image(surf, float(i)); + // CHECK: i = dpct::experimental::sample_image_by_byte(surf, float(i)); surf1Dread(&i, surf, i); - // CHECK: dpct::experimental::sample_image(surf, sycl::float2(j, i)); + // CHECK: dpct::experimental::sample_image_by_byte(surf, sycl::float2(j, i)); surf2Dread(surf, j, i); - // CHECK: i = dpct::experimental::sample_image(surf, sycl::float2(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(surf, sycl::float3(k, 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(surf, sycl::float3(k, j, i)); + // CHECK: i = dpct::experimental::sample_image_by_byte(surf, sycl::float3(k, j, i)); surf3Dread(&i, surf, k, j, i); } From dff0bcaf1f411abb2675700719967ba783bb719d Mon Sep 17 00:00:00 2001 From: "sheng.s.chen" Date: Sun, 24 Nov 2024 18:56:30 -0800 Subject: [PATCH 12/21] update the failed. Signed-off-by: sheng.s.chen --- clang/lib/DPCT/RulesLang/CallExprRewriterTexture.cpp | 3 +-- clang/runtime/dpct-rt/include/dpct/bindless_images.hpp | 2 +- clang/runtime/dpct-rt/include/dpct/image.hpp | 5 +++-- clang/test/dpct/surface.cu | 1 - 4 files changed, 5 insertions(+), 6 deletions(-) diff --git a/clang/lib/DPCT/RulesLang/CallExprRewriterTexture.cpp b/clang/lib/DPCT/RulesLang/CallExprRewriterTexture.cpp index d5af62176167..9e1c3fb168ce 100644 --- a/clang/lib/DPCT/RulesLang/CallExprRewriterTexture.cpp +++ b/clang/lib/DPCT/RulesLang/CallExprRewriterTexture.cpp @@ -23,8 +23,7 @@ class TextureReadRewriterFactory : public CallExprRewriterFactoryBase { std::shared_ptr createRewriter(const CallExpr *C, bool RetAssign, bool IsSurfAPI, BaseT Base) const { - static std::string MemberName = "read"; - MemberName = IsSurfAPI ? "read_byte" : "read"; + const std::string MemberName = IsSurfAPI ? "read_byte" : "read"; using ReaderPrinter = decltype(makeMemberCallCreator( std::declval>(), false, MemberName, makeCallArgCreatorWithCall(Idx)...)(C)); diff --git a/clang/runtime/dpct-rt/include/dpct/bindless_images.hpp b/clang/runtime/dpct-rt/include/dpct/bindless_images.hpp index 9d485dc82175..dff7e109c297 100644 --- a/clang/runtime/dpct-rt/include/dpct/bindless_images.hpp +++ b/clang/runtime/dpct-rt/include/dpct/bindless_images.hpp @@ -19,7 +19,7 @@ namespace dpct { namespace experimental { -#ifndef SYCL_EXT_ONEAPI_BINDLESS_IMAGES +#ifdef SYCL_EXT_ONEAPI_BINDLESS_IMAGES /// The wrapper class of bindless image memory handle. class image_mem_wrapper { diff --git a/clang/runtime/dpct-rt/include/dpct/image.hpp b/clang/runtime/dpct-rt/include/dpct/image.hpp index 1c58c18b5aef..2907423fced7 100644 --- a/clang/runtime/dpct-rt/include/dpct/image.hpp +++ b/clang/runtime/dpct-rt/include/dpct/image.hpp @@ -901,7 +901,7 @@ class image_accessor_ext { &&std::is_integral::value &&std::is_integral::value> typename std::enable_if::type read_byte(Coord0 x, Coord1 y, - Coord2 z) { + Coord2 z) { return detail::fetch_data()( _img_acc.read(sycl::int4(x / sizeof(T), y, z, 0), _sampler)); } @@ -911,7 +911,8 @@ class image_accessor_ext { bool Available = dimensions == 2 && std::is_integral::value &&std::is_integral::value> - typename std::enable_if::type read_byte(Coord0 x, Coord1 y) { + 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)); } diff --git a/clang/test/dpct/surface.cu b/clang/test/dpct/surface.cu index a868027a2418..0052461021aa 100644 --- a/clang/test/dpct/surface.cu +++ b/clang/test/dpct/surface.cu @@ -1,6 +1,5 @@ // 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 -// RUN: %if build_lit %{icpx -c -fsycl %T/surface/surface.dp.cpp -o %T/surface/surface.dp.o %} // CHECK: template void kernel(dpct::image_accessor_ext surf) { template __global__ void kernel(cudaSurfaceObject_t surf) { From 87ab98ecc9b654b3351c0eec48af4212992a387c Mon Sep 17 00:00:00 2001 From: "Chen, Sheng S" Date: Mon, 25 Nov 2024 22:37:12 -0800 Subject: [PATCH 13/21] update Signed-off-by: Chen, Sheng S --- .../RulesLang/CallExprRewriterTexture.cpp | 78 ++++++++++--------- clang/test/dpct/surface.cu | 6 ++ 2 files changed, 47 insertions(+), 37 deletions(-) diff --git a/clang/lib/DPCT/RulesLang/CallExprRewriterTexture.cpp b/clang/lib/DPCT/RulesLang/CallExprRewriterTexture.cpp index 9e1c3fb168ce..be9913e6bcea 100644 --- a/clang/lib/DPCT/RulesLang/CallExprRewriterTexture.cpp +++ b/clang/lib/DPCT/RulesLang/CallExprRewriterTexture.cpp @@ -14,6 +14,10 @@ namespace clang { namespace dpct { template class TextureReadRewriterFactory : public CallExprRewriterFactoryBase { + std::string ReadFuncName; + std::string SampleFuncName; + std::string SampleMMFuncName; + std::string SampleArrayFuncName; std::string Source; int TexType; @@ -21,38 +25,29 @@ class TextureReadRewriterFactory : public CallExprRewriterFactoryBase { template std::shared_ptr - createRewriter(const CallExpr *C, bool RetAssign, bool IsSurfAPI, - BaseT Base) const { - const std::string MemberName = IsSurfAPI ? "read_byte" : "read"; + createRewriter(const CallExpr *C, bool RetAssign, BaseT Base) const { using ReaderPrinter = decltype(makeMemberCallCreator( std::declval>(), false, - MemberName, makeCallArgCreatorWithCall(Idx)...)(C)); + ReadFuncName, 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, ReadFuncName, std::make_pair(C, C->getArg(Idx + 1))...)); } return std::make_shared>( - C, Source, Base, false, MemberName, + C, Source, Base, false, ReadFuncName, std::make_pair(C, C->getArg(Idx))...); } template std::shared_ptr createbindlessRewriterNormal(const CallExpr *C, bool RetAssign, - bool IsSurfAPI, const TemplateArgumentInfo &TAI, + const TemplateArgumentInfo &TAI, const std::string &VecTypeName) const { - const std::string FuncName = [=]() -> std::string { - if (IsSurfAPI) - return MapNames::getLibraryHelperNamespace() + - "experimental::sample_image_by_byte"; - return MapNames::getClNamespace() + - "ext::oneapi::experimental::sample_image"; - }(); using FuncNamePrinter = - TemplatedNamePrinter>; + TemplatedNamePrinter>; using ReaderPrinter = CallExprPrinter, VecType>; @@ -61,11 +56,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(SampleFuncName, {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(SampleFuncName, {TAI}), std::make_pair(C, C->getArg(0)), VecType(VecTypeName, std::make_pair(C, C->getArg(Idx))...)); } @@ -75,10 +71,8 @@ class TextureReadRewriterFactory : public CallExprRewriterFactoryBase { createbindlessRewriterLod(const CallExpr *C, bool RetAssign, const TemplateArgumentInfo &TAI, const std::string &VecTypeName) const { - const std::string FuncName = - MapNames::getClNamespace() + "ext::oneapi::experimental::sample_mipmap"; using FuncNamePrinter = - TemplatedNamePrinter>; + TemplatedNamePrinter>; using ReaderPrinter = CallExprPrinter, VecType, @@ -88,12 +82,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(SampleMMFuncName, {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(SampleMMFuncName, {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))); @@ -104,11 +99,8 @@ class TextureReadRewriterFactory : public CallExprRewriterFactoryBase { createbindlessRewriterLayered(const CallExpr *C, bool RetAssign, const TemplateArgumentInfo &TAI, const std::string &VecTypeName) const { - const std::string FuncName = - MapNames::getClNamespace() + - "ext::oneapi::experimental::sample_image_array"; using FuncNamePrinter = - TemplatedNamePrinter>; + TemplatedNamePrinter>; using ReaderPrinter = CallExprPrinter, VecType, @@ -118,18 +110,19 @@ class TextureReadRewriterFactory : public CallExprRewriterFactoryBase { BinaryOperatorPrinter>>( C, Source, DerefExpr(C->getArg(0), C), ReaderPrinter( - FuncNamePrinter(FuncName, {TAI}), std::make_pair(C, C->getArg(1)), + FuncNamePrinter(SampleArrayFuncName, {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(SampleArrayFuncName, {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, bool IsSurfAPI, + createbindlessRewriter(const CallExpr *C, bool RetAssign, QualType TargetType) const { TemplateArgumentInfo TAI; auto TAL = getTemplateArgsList(C); @@ -160,13 +153,27 @@ class TextureReadRewriterFactory : public CallExprRewriterFactoryBase { if ((TexType & 0xf0) == 0xf0) return createbindlessRewriterLayered(C, RetAssign, TAI, VecTypeName); - return createbindlessRewriterNormal(C, RetAssign, IsSurfAPI, TAI, + return createbindlessRewriterNormal(C, RetAssign, TAI, VecTypeName); } public: TextureReadRewriterFactory(std::string Name, int Tex) - : Source(std::move(Name)), TexType(Tex) {} + : Source(std::move(Name)), TexType(Tex) { + if (Source.find("surf") != std::string::npos) { + ReadFuncName = "read_byte"; + SampleFuncName = MapNames::getLibraryHelperNamespace() + + "experimental::sample_image_by_byte"; + } else { + ReadFuncName = "read"; + SampleFuncName = MapNames::getClNamespace() + + "ext::oneapi::experimental::sample_image"; + } + SampleMMFuncName = + MapNames::getClNamespace() + "ext::oneapi::experimental::sample_mipmap"; + SampleArrayFuncName = MapNames::getClNamespace() + + "ext::oneapi::experimental::sample_image_array"; + } std::shared_ptr create(const CallExpr *Call) const override { const Expr *SourceExpr = Call->getArg(0); @@ -174,9 +181,6 @@ class TextureReadRewriterFactory : public CallExprRewriterFactoryBase { QualType TargetType = Call->getType(); StringRef SourceName; bool RetAssign = false; - bool IsSurfAPI = false; - if (Source.find("surf") != std::string::npos) - IsSurfAPI = true; if (SourceExpr->getType()->isPointerType()) { TargetType = SourceExpr->getType()->getPointeeType(); SourceExpr = Call->getArg(1); @@ -189,7 +193,7 @@ class TextureReadRewriterFactory : public CallExprRewriterFactoryBase { } } if (DpctGlobalInfo::useExtBindlessImages()) { - return createbindlessRewriter(Call, RetAssign, IsSurfAPI, TargetType); + return createbindlessRewriter(Call, RetAssign, TargetType); } SourceExpr = SourceExpr->IgnoreImpCasts(); if (auto FD = DpctGlobalInfo::getParentFunction(Call)) { @@ -204,7 +208,7 @@ class TextureReadRewriterFactory : public CallExprRewriterFactoryBase { MemberInfo->setType( DpctGlobalInfo::getUnqualifiedTypeName(TargetType), TexType); SourceName = MemberInfo->getName(); - return createRewriter(Call, RetAssign, IsSurfAPI, SourceName); + return createRewriter(Call, RetAssign, SourceName); } } else if (auto DRE = dyn_cast(SourceExpr)) { auto CallDefRange = @@ -221,7 +225,7 @@ class TextureReadRewriterFactory : public CallExprRewriterFactoryBase { } } - return createRewriter(Call, RetAssign, IsSurfAPI, + return createRewriter(Call, RetAssign, std::make_pair(Call, Call->getArg(RetAssign & 0x01))); } }; diff --git a/clang/test/dpct/surface.cu b/clang/test/dpct/surface.cu index 0052461021aa..9f063e8a9321 100644 --- a/clang/test/dpct/surface.cu +++ b/clang/test/dpct/surface.cu @@ -19,6 +19,12 @@ template __global__ void kernel(cudaSurfaceObject_t surf) { surf3Dread(&i, surf, k, j, i); } +static texture tex21; + +__device__ void device01() { + // CHECK: sycl::uint2 u21 = (tex21).read(1.0f); + uint2 u21 = tex1D(tex21, 1.0f); +} int main() { // CHECK: dpct::image_wrapper_base_p surf; cudaSurfaceObject_t surf; From bb9fa92a737387d3b24ce74d691f03d40c88705a Mon Sep 17 00:00:00 2001 From: "Chen, Sheng S" Date: Tue, 26 Nov 2024 00:33:15 -0800 Subject: [PATCH 14/21] up Signed-off-by: Chen, Sheng S --- clang/lib/DPCT/RulesLang/RulesLangTexture.cpp | 3 +++ 1 file changed, 3 insertions(+) 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", From 1a0870a1914d9fc0c790c9bc05e85b8aa656224e Mon Sep 17 00:00:00 2001 From: "Chen, Sheng S" Date: Tue, 26 Nov 2024 04:27:18 -0800 Subject: [PATCH 15/21] update Signed-off-by: Chen, Sheng S --- clang/test/dpct/surface.cu | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/clang/test/dpct/surface.cu b/clang/test/dpct/surface.cu index 9f063e8a9321..4675cd181179 100644 --- a/clang/test/dpct/surface.cu +++ b/clang/test/dpct/surface.cu @@ -22,8 +22,8 @@ template __global__ void kernel(cudaSurfaceObject_t surf) { static texture tex21; __device__ void device01() { - // CHECK: sycl::uint2 u21 = (tex21).read(1.0f); - uint2 u21 = tex1D(tex21, 1.0f); + // CHECK: tex21.read(1.0f); + tex1D(tex21, 1.0f); } int main() { // CHECK: dpct::image_wrapper_base_p surf; From 938afeb6a0a4bd650e68964c67439ba2d26408e5 Mon Sep 17 00:00:00 2001 From: "Chen, Sheng S" Date: Tue, 26 Nov 2024 16:43:40 -0800 Subject: [PATCH 16/21] triage the CI Signed-off-by: Chen, Sheng S --- clang/test/dpct/texture/surface_object_bindless_image.cu | 1 - 1 file changed, 1 deletion(-) diff --git a/clang/test/dpct/texture/surface_object_bindless_image.cu b/clang/test/dpct/texture/surface_object_bindless_image.cu index 059a1d10a59c..70472a11d45f 100644 --- a/clang/test/dpct/texture/surface_object_bindless_image.cu +++ b/clang/test/dpct/texture/surface_object_bindless_image.cu @@ -19,7 +19,6 @@ template __global__ void kernel(cudaSurfaceObject_t surf) { // 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; From 7344e7e91000c1d0e8fe3a3bcec298906225ad5e Mon Sep 17 00:00:00 2001 From: "Chen, Sheng S" Date: Tue, 26 Nov 2024 22:12:21 -0800 Subject: [PATCH 17/21] up Signed-off-by: Chen, Sheng S --- clang/test/dpct/surface.cu | 1 - 1 file changed, 1 deletion(-) diff --git a/clang/test/dpct/surface.cu b/clang/test/dpct/surface.cu index 4675cd181179..5e1e5431705b 100644 --- a/clang/test/dpct/surface.cu +++ b/clang/test/dpct/surface.cu @@ -22,7 +22,6 @@ template __global__ void kernel(cudaSurfaceObject_t surf) { static texture tex21; __device__ void device01() { - // CHECK: tex21.read(1.0f); tex1D(tex21, 1.0f); } int main() { From 39a0b049ace54ded12e1ead2f833b1dd77e9e093 Mon Sep 17 00:00:00 2001 From: "Chen, Sheng S" Date: Fri, 29 Nov 2024 00:56:00 -0800 Subject: [PATCH 18/21] update the entry Signed-off-by: Chen, Sheng S --- clang/lib/DPCT/RulesLang/APINamesTexture.inc | 146 ++++++++++++++---- .../RulesLang/CallExprRewriterTexture.cpp | 73 +++++---- 2 files changed, 151 insertions(+), 68 deletions(-) diff --git a/clang/lib/DPCT/RulesLang/APINamesTexture.inc b/clang/lib/DPCT/RulesLang/APINamesTexture.inc index c6c2023120c0..099651e727cf 100644 --- a/clang/lib/DPCT/RulesLang/APINamesTexture.inc +++ b/clang/lib/DPCT/RulesLang/APINamesTexture.inc @@ -120,48 +120,105 @@ 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", 0x02, "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, + ENTRY_TEXTURE("tex1DLayered", 0xF1, + MapNames::getClNamespace() + + "ext::oneapi::experimental::sample_image_array", + 1), + FEATURE_REQUEST_FACTORY( + HelperFeatureEnum::device_ext, + ENTRY_TEXTURE("tex1DLayered", 0xF1, + MapNames::getClNamespace() + + "ext::oneapi::experimental::sample_image_array", + 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, + MapNames::getClNamespace() + + "ext::oneapi::experimental::sample_image_array", + 3, 1, 2))) FEATURE_REQUEST_FACTORY(HelperFeatureEnum::device_ext, ASSIGNABLE_FACTORY(ENTRY_BIND("cudaBindTexture", @@ -288,9 +345,36 @@ FEATURE_REQUEST_FACTORY( 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)) +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 be9913e6bcea..a99f3cd72490 100644 --- a/clang/lib/DPCT/RulesLang/CallExprRewriterTexture.cpp +++ b/clang/lib/DPCT/RulesLang/CallExprRewriterTexture.cpp @@ -14,12 +14,9 @@ namespace clang { namespace dpct { template class TextureReadRewriterFactory : public CallExprRewriterFactoryBase { - std::string ReadFuncName; - std::string SampleFuncName; - std::string SampleMMFuncName; - std::string SampleArrayFuncName; std::string Source; int TexType; + std::string TargetName; inline int getDim() const { return TexType & 0x0f; } @@ -28,16 +25,16 @@ class TextureReadRewriterFactory : public CallExprRewriterFactoryBase { createRewriter(const CallExpr *C, bool RetAssign, BaseT Base) const { using ReaderPrinter = decltype(makeMemberCallCreator( std::declval>(), false, - ReadFuncName, 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, ReadFuncName, + ReaderPrinter(std::move(Base), false, TargetName, std::make_pair(C, C->getArg(Idx + 1))...)); } return std::make_shared>( - C, Source, Base, false, ReadFuncName, + C, Source, Base, false, TargetName, std::make_pair(C, C->getArg(Idx))...); } @@ -56,12 +53,12 @@ class TextureReadRewriterFactory : public CallExprRewriterFactoryBase { BinaryOperatorPrinter>>( C, Source, DerefExpr(C->getArg(0), C), ReaderPrinter( - FuncNamePrinter(SampleFuncName, {TAI}), + 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(SampleFuncName, {TAI}), + C, Source, FuncNamePrinter(TargetName, {TAI}), std::make_pair(C, C->getArg(0)), VecType(VecTypeName, std::make_pair(C, C->getArg(Idx))...)); } @@ -82,13 +79,13 @@ class TextureReadRewriterFactory : public CallExprRewriterFactoryBase { BinaryOperatorPrinter>>( C, Source, DerefExpr(C->getArg(0), C), ReaderPrinter( - FuncNamePrinter(SampleMMFuncName, {TAI}), + 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(SampleMMFuncName, {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))); @@ -110,13 +107,13 @@ class TextureReadRewriterFactory : public CallExprRewriterFactoryBase { BinaryOperatorPrinter>>( C, Source, DerefExpr(C->getArg(0), C), ReaderPrinter( - FuncNamePrinter(SampleArrayFuncName, {TAI}), + 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(SampleArrayFuncName, {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))); @@ -158,21 +155,21 @@ class TextureReadRewriterFactory : public CallExprRewriterFactoryBase { } public: - TextureReadRewriterFactory(std::string Name, int Tex) - : Source(std::move(Name)), TexType(Tex) { - if (Source.find("surf") != std::string::npos) { - ReadFuncName = "read_byte"; - SampleFuncName = MapNames::getLibraryHelperNamespace() + - "experimental::sample_image_by_byte"; - } else { - ReadFuncName = "read"; - SampleFuncName = MapNames::getClNamespace() + - "ext::oneapi::experimental::sample_image"; - } - SampleMMFuncName = - MapNames::getClNamespace() + "ext::oneapi::experimental::sample_mipmap"; - SampleArrayFuncName = MapNames::getClNamespace() + - "ext::oneapi::experimental::sample_image_array"; + TextureReadRewriterFactory(std::string Name, int Tex, std::string TargetName) + : Source(std::move(Name)), TexType(Tex), TargetName(TargetName) { + // if (Source.find("surf") != std::string::npos) { + // ReadFuncName = "read_byte"; + // SampleFuncName = MapNames::getLibraryHelperNamespace() + + // "experimental::sample_image_by_byte"; + // } else { + // ReadFuncName = "read"; + // SampleFuncName = MapNames::getClNamespace() + + // "ext::oneapi::experimental::sample_image"; + // } + // SampleMMFuncName = + // MapNames::getClNamespace() + "ext::oneapi::experimental::sample_mipmap"; + // SampleArrayFuncName = MapNames::getClNamespace() + + // "ext::oneapi::experimental::sample_image_array"; } std::shared_ptr create(const CallExpr *Call) const override { @@ -239,14 +236,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)}, @@ -256,7 +254,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( @@ -264,8 +263,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, ...) \ From 1ca5a74396703370afcb5d519329b52f5be38efb Mon Sep 17 00:00:00 2001 From: "Chen, Sheng S" Date: Sun, 1 Dec 2024 16:41:55 -0800 Subject: [PATCH 19/21] remove comment code. Signed-off-by: Chen, Sheng S --- .../DPCT/RulesLang/CallExprRewriterTexture.cpp | 16 +--------------- 1 file changed, 1 insertion(+), 15 deletions(-) diff --git a/clang/lib/DPCT/RulesLang/CallExprRewriterTexture.cpp b/clang/lib/DPCT/RulesLang/CallExprRewriterTexture.cpp index a99f3cd72490..ea93456381ea 100644 --- a/clang/lib/DPCT/RulesLang/CallExprRewriterTexture.cpp +++ b/clang/lib/DPCT/RulesLang/CallExprRewriterTexture.cpp @@ -156,21 +156,7 @@ class TextureReadRewriterFactory : public CallExprRewriterFactoryBase { public: TextureReadRewriterFactory(std::string Name, int Tex, std::string TargetName) - : Source(std::move(Name)), TexType(Tex), TargetName(TargetName) { - // if (Source.find("surf") != std::string::npos) { - // ReadFuncName = "read_byte"; - // SampleFuncName = MapNames::getLibraryHelperNamespace() + - // "experimental::sample_image_by_byte"; - // } else { - // ReadFuncName = "read"; - // SampleFuncName = MapNames::getClNamespace() + - // "ext::oneapi::experimental::sample_image"; - // } - // SampleMMFuncName = - // MapNames::getClNamespace() + "ext::oneapi::experimental::sample_mipmap"; - // SampleArrayFuncName = MapNames::getClNamespace() + - // "ext::oneapi::experimental::sample_image_array"; - } + : Source(std::move(Name)), TexType(Tex), TargetName(TargetName) {} std::shared_ptr create(const CallExpr *Call) const override { const Expr *SourceExpr = Call->getArg(0); From 6c11e94fbd4679550e869fd56da5dfd82806d89a Mon Sep 17 00:00:00 2001 From: "Chen, Sheng S" Date: Sun, 1 Dec 2024 17:42:06 -0800 Subject: [PATCH 20/21] update INC Signed-off-by: Chen, Sheng S --- clang/lib/DPCT/RulesLang/APINamesTexture.inc | 11 +++++------ 1 file changed, 5 insertions(+), 6 deletions(-) diff --git a/clang/lib/DPCT/RulesLang/APINamesTexture.inc b/clang/lib/DPCT/RulesLang/APINamesTexture.inc index 099651e727cf..64392d481607 100644 --- a/clang/lib/DPCT/RulesLang/APINamesTexture.inc +++ b/clang/lib/DPCT/RulesLang/APINamesTexture.inc @@ -129,7 +129,7 @@ CONDITIONAL_FACTORY_ENTRY( "ext::oneapi::experimental::sample_image", 1)), FEATURE_REQUEST_FACTORY(HelperFeatureEnum::device_ext, - ENTRY_TEXTURE("tex1D", 0x02, "read", 1))) + ENTRY_TEXTURE("tex1D", 0x01, "read", 1))) CONDITIONAL_FACTORY_ENTRY( UseExtBindlessImages, @@ -197,16 +197,15 @@ CONDITIONAL_FACTORY_ENTRY( CONDITIONAL_FACTORY_ENTRY( UseExtBindlessImages, - ENTRY_TEXTURE("tex1DLayered", 0xF1, - MapNames::getClNamespace() + - "ext::oneapi::experimental::sample_image_array", - 1), FEATURE_REQUEST_FACTORY( HelperFeatureEnum::device_ext, ENTRY_TEXTURE("tex1DLayered", 0xF1, MapNames::getClNamespace() + "ext::oneapi::experimental::sample_image_array", - 2, 1))) + 1)), + FEATURE_REQUEST_FACTORY(HelperFeatureEnum::device_ext, + ENTRY_TEXTURE("tex1DLayered", 0xF1, "read", 2, 1))) + CONDITIONAL_FACTORY_ENTRY( UseExtBindlessImages, ENTRY_TEXTURE("tex2DLayered", 0xF2, From 7f83c5508aa1a501d439dd0987990191238a6204 Mon Sep 17 00:00:00 2001 From: "Chen, Sheng S" Date: Sun, 1 Dec 2024 19:13:14 -0800 Subject: [PATCH 21/21] fix the 8.0 header --- clang/lib/DPCT/RulesLang/APINamesTexture.inc | 9 +++------ 1 file changed, 3 insertions(+), 6 deletions(-) diff --git a/clang/lib/DPCT/RulesLang/APINamesTexture.inc b/clang/lib/DPCT/RulesLang/APINamesTexture.inc index 64392d481607..6ce5365b3aae 100644 --- a/clang/lib/DPCT/RulesLang/APINamesTexture.inc +++ b/clang/lib/DPCT/RulesLang/APINamesTexture.inc @@ -212,12 +212,9 @@ CONDITIONAL_FACTORY_ENTRY( MapNames::getClNamespace() + "ext::oneapi::experimental::sample_image_array", 1, 2), - FEATURE_REQUEST_FACTORY( - HelperFeatureEnum::device_ext, - ENTRY_TEXTURE("tex2DLayered", 0xF2, - MapNames::getClNamespace() + - "ext::oneapi::experimental::sample_image_array", - 3, 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",