From 0ad3a475d99a21588d90805ad0fd06a68a8f56b9 Mon Sep 17 00:00:00 2001 From: "Tang, Jiajun" Date: Mon, 20 May 2024 16:27:03 +0800 Subject: [PATCH 1/3] [SYCLomatic] Refine migration of CUDA_ARRAY_DESCRIPTOR and 2 related APIs. Signed-off-by: Tang, Jiajun jiajun.tang@intel.com --- .../DPCT/Driver/cuTexRefSetAddress2D.cu | 4 +- clang/lib/DPCT/APINames.inc | 2 +- clang/lib/DPCT/APINamesTexture.inc | 23 ++++---- clang/lib/DPCT/ASTTraversal.cpp | 35 +++-------- clang/lib/DPCT/ASTTraversal.h | 2 - clang/lib/DPCT/CallExprRewriterCommon.h | 29 --------- clang/lib/DPCT/Diagnostics.inc | 4 +- clang/lib/DPCT/MapNames.cpp | 6 +- clang/lib/DPCT/TypeNames.inc | 1 - clang/runtime/dpct-rt/include/dpct/image.hpp | 17 ++++++ clang/test/dpct/know_unsupported_type.cu | 5 +- .../dpct/query_api_mapping/Driver/test.cu | 6 +- clang/test/dpct/texture_driver.cu | 59 +++++++------------ clang/test/dpct/texture_object_driver.cu | 14 ++--- 14 files changed, 75 insertions(+), 132 deletions(-) diff --git a/clang/examples/DPCT/Driver/cuTexRefSetAddress2D.cu b/clang/examples/DPCT/Driver/cuTexRefSetAddress2D.cu index 9bb01f33b8ab..f204e5040dcb 100644 --- a/clang/examples/DPCT/Driver/cuTexRefSetAddress2D.cu +++ b/clang/examples/DPCT/Driver/cuTexRefSetAddress2D.cu @@ -1,7 +1,7 @@ void test(const CUDA_ARRAY_DESCRIPTOR *pa, CUdeviceptr d, size_t s) { // Start CUtexref t; - cuTexRefSetAddress2D(t /*CUtexref*/, pa /*size_t **/, d /*CUdeviceptr*/, - s /*size_t*/); + cuTexRefSetAddress2D(t /*CUtexref*/, pa /*const CUDA_ARRAY_DESCRIPTOR **/, + d /*CUdeviceptr*/, s /*size_t*/); // End } diff --git a/clang/lib/DPCT/APINames.inc b/clang/lib/DPCT/APINames.inc index 7a185ab052b6..b47051e76d35 100644 --- a/clang/lib/DPCT/APINames.inc +++ b/clang/lib/DPCT/APINames.inc @@ -1625,7 +1625,7 @@ ENTRY(cuGraphicsUnregisterResource, cuGraphicsUnregisterResource, false, NO_FLAG ENTRY(cuInit, cuInit, true, NO_FLAG, P4, "DPCT1026/DPCT1027") ENTRY(cuArray3DCreate, cuArray3DCreate_v2, false, NO_FLAG, P4, "comment") ENTRY(cuArray3DGetDescriptor, cuArray3DGetDescriptor_v2, false, NO_FLAG, P4, "comment") -ENTRY(cuArrayCreate, cuArrayCreate_v2, true, NO_FLAG, P4, "DPCT1073") +ENTRY(cuArrayCreate, cuArrayCreate_v2, true, NO_FLAG, P4, "Successful") ENTRY(cuArrayDestroy, cuArrayDestroy, true, NO_FLAG, P4, "Successful") ENTRY(cuArrayGetDescriptor, cuArrayGetDescriptor_v2, false, NO_FLAG, P4, "comment") ENTRY(cuDeviceGetByPCIBusId, cuDeviceGetByPCIBusId, false, NO_FLAG, P4, "comment") diff --git a/clang/lib/DPCT/APINamesTexture.inc b/clang/lib/DPCT/APINamesTexture.inc index cb954ea5fd39..808ef18ee08a 100644 --- a/clang/lib/DPCT/APINamesTexture.inc +++ b/clang/lib/DPCT/APINamesTexture.inc @@ -181,17 +181,15 @@ FEATURE_REQUEST_FACTORY(HelperFeatureEnum::device_ext, ENTRY_UNSUPPORTED("cudaGetTextureObjectResourceViewDesc", Diagnostics::API_NOT_MIGRATED) -CONDITIONAL_FACTORY_ENTRY( - CheckWarning1073(1), - FEATURE_REQUEST_FACTORY( - HelperFeatureEnum::device_ext, - ASSIGNABLE_FACTORY(ASSIGN_FACTORY_ENTRY( - "cuArrayCreate_v2", DEREF(0), - NEW(MapNames::getDpctNamespace() + "image_matrix", - STRUCT_DISMANTLE(1, "channel_type_ct1", "channel_num_ct1", - "x_ct1", "y_ct1"))))), - UNSUPPORT_FACTORY_ENTRY("cuArrayCreate_v2", - Diagnostics::CANNOT_CAPUTURE_AGUMENTS, ARG(1))) +FEATURE_REQUEST_FACTORY( + HelperFeatureEnum::device_ext, + ASSIGNABLE_FACTORY(ASSIGN_FACTORY_ENTRY( + "cuArrayCreate_v2", DEREF(0), + NEW(MapNames::getDpctNamespace() + "image_matrix", + MEMBER_EXPR(ARG(1), true, LITERAL("channel_type")), + MEMBER_EXPR(ARG(1), true, LITERAL("channel_num")), + MEMBER_EXPR(ARG(1), true, LITERAL("width")), + MEMBER_EXPR(ARG(1), true, LITERAL("height")))))) ASSIGNABLE_FACTORY(DELETER_FACTORY_ENTRY("cuArrayDestroy", ARG(0))) ENTRY_UNSUPPORTED("cuTexObjectGetResourceViewDesc", Diagnostics::API_NOT_MIGRATED) FEATURE_REQUEST_FACTORY( @@ -236,8 +234,7 @@ FEATURE_REQUEST_FACTORY(HelperFeatureEnum::device_ext, FEATURE_REQUEST_FACTORY(HelperFeatureEnum::device_ext, ASSIGNABLE_FACTORY(MEMBER_CALL_FACTORY_ENTRY( "cuTexRefSetAddress2D_v3", ARG(0), true, "attach", - ARG(2), STRUCT_DISMANTLE(1, "x_ct1", "y_ct1"), - ARG(3)))) + ARG(1), ARG(2), ARG(3)))) FEATURE_REQUEST_FACTORY( HelperFeatureEnum::device_ext, diff --git a/clang/lib/DPCT/ASTTraversal.cpp b/clang/lib/DPCT/ASTTraversal.cpp index fefad266e02d..bb6b9f05808d 100644 --- a/clang/lib/DPCT/ASTTraversal.cpp +++ b/clang/lib/DPCT/ASTTraversal.cpp @@ -1753,7 +1753,7 @@ void TypeInDeclRule::registerMatcher(MatchFinder &MF) { "cudaLaunchAttributeValue", "cusparseSpSMDescr_t", "cusparseConstSpMatDescr_t", "cusparseSpSMAlg_t", "cusparseConstDnMatDescr_t", "cudaMemcpy3DParms", "CUDA_MEMCPY3D", - "CUDA_MEMCPY2D")))))) + "CUDA_MEMCPY2D", "CUDA_ARRAY_DESCRIPTOR")))))) .bind("cudaTypeDef"), this); @@ -11494,21 +11494,7 @@ TextModification *ReplaceMemberAssignAsSetMethod(const Expr *E, ME, MethodName, ReplacedArg, ExtraArg); } -void MemoryDataTypeRule::emplaceCuArrayDescDeclarations(const VarDecl *VD) { - if (DpctGlobalInfo::isCommentsEnabled()) { - emplaceTransformation(ReplaceVarDecl::getVarDeclReplacement( - VD, "// These variables are defined for info of image_matrix.")); - } - emplaceParamDecl(VD, "size_t", false, "0", "x", "y"); - emplaceParamDecl(VD, "unsigned", false, "0", "channel_num"); - emplaceParamDecl(VD, MapNames::getClNamespace() + "image_channel_type", false, - "0", "channel_type"); -} - void MemoryDataTypeRule::registerMatcher(MatchFinder &MF) { - MF.addMatcher(varDecl(hasType(namedDecl(hasAnyName("CUDA_ARRAY_DESCRIPTOR")))) - .bind("decl"), - this); MF.addMatcher(memberExpr(hasObjectExpression(declRefExpr(hasType(namedDecl( hasAnyName("CUDA_ARRAY_DESCRIPTOR")))))) .bind("arrayMember"), @@ -11530,19 +11516,12 @@ void MemoryDataTypeRule::registerMatcher(MatchFinder &MF) { } void MemoryDataTypeRule::runRule(const MatchFinder::MatchResult &Result) { - if (auto VD = getNodeAsType(Result, "decl")) { - if (isa(VD)) - return; - auto TypeName = DpctGlobalInfo::getUnqualifiedTypeName(VD->getType()); - if (TypeName == "CUDA_ARRAY_DESCRIPTOR") - emplaceCuArrayDescDeclarations(VD); - } else if (auto ME = getNodeAsType(Result, "arrayMember")) { - if (auto DRE = - dyn_cast(ME->getBase()->IgnoreImplicitAsWritten())) { - emplaceTransformation(new ReplaceStmt( - ME, getArrayDescMemberName(DRE->getDecl()->getName(), - ME->getMemberDecl()->getName().str()))); - } + if (auto ME = getNodeAsType(Result, "arrayMember")) { + const auto &Replace = MapNames::findReplacedName( + ArrayDescMemberNames, ME->getMemberDecl()->getName().str()); + if (!Replace.empty()) + emplaceTransformation(new ReplaceToken( + ME->getMemberLoc(), ME->getEndLoc(), std::string(Replace))); } else if (auto CE = getNodeAsType(Result, "makeData")) { if (auto FD = CE->getDirectCallee()) { auto Name = FD->getName(); diff --git a/clang/lib/DPCT/ASTTraversal.h b/clang/lib/DPCT/ASTTraversal.h index 2dd9fd324df4..6341c0829e3d 100644 --- a/clang/lib/DPCT/ASTTraversal.h +++ b/clang/lib/DPCT/ASTTraversal.h @@ -1527,8 +1527,6 @@ class MemoryDataTypeRule : public NamedMigrationRule { const static std::vector RemoveMember; public: - void emplaceCuArrayDescDeclarations(const VarDecl *VD); - static std::string getArrayDescMemberName(StringRef BaseName, const std::string &Member) { auto Itr = ArrayDescMemberNames.find(Member); diff --git a/clang/lib/DPCT/CallExprRewriterCommon.h b/clang/lib/DPCT/CallExprRewriterCommon.h index 45fd97746bf6..6061388e2466 100644 --- a/clang/lib/DPCT/CallExprRewriterCommon.h +++ b/clang/lib/DPCT/CallExprRewriterCommon.h @@ -324,23 +324,6 @@ makeCombinedArg(std::function Part1, }; } -inline std::function(const CallExpr *)> -makeStructDismantler(unsigned Idx, const std::vector &Suffixes) { - return [=](const CallExpr *C) -> std::vector { - std::vector Ret; - if (auto DRE = dyn_cast_or_null( - getDereferencedExpr(C->getArg(Idx)))) { - Ret.reserve(Suffixes.size()); - auto Origin = DRE->getDecl()->getName(); - std::transform(Suffixes.begin(), Suffixes.end(), std::back_inserter(Ret), - [&](StringRef Suffix) -> RenameWithSuffix { - return RenameWithSuffix(Origin, Suffix); - }); - } - return Ret; - }; -} - inline std::function makeExtendStr(unsigned Idx, const std::string Suffix) { return [=](const CallExpr *C) -> std::string { @@ -1561,17 +1544,6 @@ createDerefExprRewriterFactory( std::forward>(ArgCreator)); } -class CheckWarning1073 { - unsigned Idx; - -public: - CheckWarning1073(unsigned I) : Idx(I) {} - bool operator()(const CallExpr *C) { - auto DerefE = getDereferencedExpr(C->getArg(Idx)); - return DerefE && isa(DerefE); - } -}; - // sycl has 2 overloading of malloc_device // 1. sycl::malloc_device(Addr, Size) // 2. sycl::malloc_device(Addr, Size) @@ -2050,7 +2022,6 @@ const std::string MipmapNeedBindlessImage = #define ADDROF(x) makeAddrOfExprCreator(x) #define DEREF(x) makeDerefExprCreator(x) #define DEREF_CAST_IF_NEED(T, S) makeDerefCastIfNeedExprCreator(T, S) -#define STRUCT_DISMANTLE(idx, ...) makeStructDismantler(idx, {__VA_ARGS__}) #define ARG(x) makeCallArgCreator(x) #define ARG_WC(x) makeDerefArgCreatorWithCall(x) #define TEMPLATE_ARG(x) makeCallArgCreatorFromTemplateArg(x) diff --git a/clang/lib/DPCT/Diagnostics.inc b/clang/lib/DPCT/Diagnostics.inc index dcf41207ffc6..090b4eb426de 100644 --- a/clang/lib/DPCT/Diagnostics.inc +++ b/clang/lib/DPCT/Diagnostics.inc @@ -176,8 +176,8 @@ DEF_WARNING(OUT_OF_PLACE_FFT_EXEC, 1071, LOW_LEVEL, "deprecated") DEF_COMMENT(OUT_OF_PLACE_FFT_EXEC, 1071, LOW_LEVEL, "deprecated") DEF_WARNING(UNSUPPORT_FREE_MEMORY_SIZE, 1072, LOW_LEVEL, "SYCL currently does not support getting the available memory on the current device. You may need to adjust the code.") DEF_COMMENT(UNSUPPORT_FREE_MEMORY_SIZE, 1072, LOW_LEVEL, "SYCL currently does not support getting the available memory on the current device. You may need to adjust the code.") -DEF_WARNING(CANNOT_CAPUTURE_AGUMENTS, 1073, MEDIUM_LEVEL, "The field values of parameter '%0' could not be deduced, so the call was not migrated. You need to update this code manually.") -DEF_COMMENT(CANNOT_CAPUTURE_AGUMENTS, 1073, MEDIUM_LEVEL, "The field values of parameter '{0}' could not be deduced, so the call was not migrated. You need to update this code manually.") +DEF_WARNING(CANNOT_CAPUTURE_AGUMENTS, 1073, MEDIUM_LEVEL, "deprecated") +DEF_COMMENT(CANNOT_CAPUTURE_AGUMENTS, 1073, MEDIUM_LEVEL, "deprecated") DEF_WARNING(TEX_FLAG_UNSUPPORT, 1074, MEDIUM_LEVEL, "The SYCL Image class does not support some of the flags used in the original code. Unsupported flags were ignored. Data read from SYCL Image could not be normalized as specified in the original code.") DEF_COMMENT(TEX_FLAG_UNSUPPORT, 1074, MEDIUM_LEVEL, "The SYCL Image class does not support some of the flags used in the original code. Unsupported flags were ignored. Data read from SYCL Image could not be normalized as specified in the original code.") DEF_WARNING(CHECK_RELATED_QUEUE, 1075, LOW_LEVEL, "deprecated") diff --git a/clang/lib/DPCT/MapNames.cpp b/clang/lib/DPCT/MapNames.cpp index 94015bd001b7..ca1f220ea5cb 100644 --- a/clang/lib/DPCT/MapNames.cpp +++ b/clang/lib/DPCT/MapNames.cpp @@ -373,6 +373,8 @@ void MapNames::setExplicitNamespaceMap() { HelperFeatureEnum::device_ext)}, {"cudaMemcpyKind", std::make_shared(getDpctNamespace() + "memcpy_direction")}, + {"CUDA_ARRAY_DESCRIPTOR", + std::make_shared(getDpctNamespace() + "image_desc")}, {"cudaMemcpy3DParms", std::make_shared(getDpctNamespace() + "memcpy_parameter")}, {"CUDA_MEMCPY3D", @@ -4123,8 +4125,8 @@ const MapNames::MapTy MemoryDataTypeRule::ExtentMemberNames{ {"width", "[0]"}, {"height", "[1]"}, {"depth", "[2]"}}; const MapNames::MapTy MemoryDataTypeRule::ArrayDescMemberNames{ - {"Width", "x"}, - {"Height", "y"}, + {"Width", "width"}, + {"Height", "height"}, {"Format", "channel_type"}, {"NumChannels", "channel_num"}}; diff --git a/clang/lib/DPCT/TypeNames.inc b/clang/lib/DPCT/TypeNames.inc index 11eb36506b4b..c15d66f1a4c8 100644 --- a/clang/lib/DPCT/TypeNames.inc +++ b/clang/lib/DPCT/TypeNames.inc @@ -41,7 +41,6 @@ ENTRY_TYPE(CUgraphNode, false, NO_FLAG, P4, "comment") ENTRY_TYPE(CUgraphicsResource, false, NO_FLAG, P4, "comment") // CUDA Runtime Library -ENTRY_TYPE(CUDA_ARRAY_DESCRIPTOR*, false, NO_FLAG, P4, "comment") ENTRY_TYPE(cudaKernelNodeParams, false, NO_FLAG, P4, "comment") // cuDNN Library diff --git a/clang/runtime/dpct-rt/include/dpct/image.hpp b/clang/runtime/dpct-rt/include/dpct/image.hpp index 41e8b4f9476b..1d8d1eb4b0ac 100644 --- a/clang/runtime/dpct-rt/include/dpct/image.hpp +++ b/clang/runtime/dpct-rt/include/dpct/image.hpp @@ -139,6 +139,14 @@ static image_wrapper_base *create_image_wrapper(image_channel channel, int dims) } // namespace detail +struct image_desc { + size_t width = 0; + size_t height = 0; + sycl::image_channel_type channel_type = + sycl::image_channel_type::signed_int32; + unsigned channel_num = 0; +}; + /// Image channel info, include channel number, order, data width and type class image_channel { image_channel_data_type _type = image_channel_data_type::signed_int; @@ -626,6 +634,15 @@ class image_wrapper_base { image_wrapper_base::set_data( image_data(const_cast(data), x, y, pitch, channel)); } + /// Attach device_ptr data to this class. + void attach(const image_desc *desc, device_ptr ptr, size_t pitch) { + detach(); + image_channel channel; + channel.set_channel_num(desc->channel_num); + channel.set_channel_type(desc->channel_type); + image_wrapper_base::set_data( + image_data(ptr, desc->width, desc->height, pitch, channel)); + } /// Detach data. virtual void detach() {} diff --git a/clang/test/dpct/know_unsupported_type.cu b/clang/test/dpct/know_unsupported_type.cu index 2efd0040412d..8997de81ccdd 100644 --- a/clang/test/dpct/know_unsupported_type.cu +++ b/clang/test/dpct/know_unsupported_type.cu @@ -8,10 +8,7 @@ #include "nvml.h" #include int main(int argc, char **argv) { - // CHECK: /* - // CHECK-NEXT: DPCT1082:{{[0-9]+}}: Migration of CUDA_ARRAY_DESCRIPTOR * type is not supported. - // CHECK-NEXT: */ - // CHECK-NEXT: CUDA_ARRAY_DESCRIPTOR *pcad; + // CHECK: dpct::image_desc *pcad; CUDA_ARRAY_DESCRIPTOR *pcad; // CHECK: dpct::memcpy_parameter *p1c3d; cudaMemcpy3DParms *p1c3d; diff --git a/clang/test/dpct/query_api_mapping/Driver/test.cu b/clang/test/dpct/query_api_mapping/Driver/test.cu index b0ba3f85730e..c22eb2804d67 100644 --- a/clang/test/dpct/query_api_mapping/Driver/test.cu +++ b/clang/test/dpct/query_api_mapping/Driver/test.cu @@ -313,11 +313,11 @@ // RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=cuTexRefSetAddress2D | FileCheck %s -check-prefix=CUTEXREFSETADDRESS2D // CUTEXREFSETADDRESS2D: CUDA API: // CUTEXREFSETADDRESS2D-NEXT: CUtexref t; -// CUTEXREFSETADDRESS2D-NEXT: cuTexRefSetAddress2D(t /*CUtexref*/, pa /*size_t **/, d /*CUdeviceptr*/, -// CUTEXREFSETADDRESS2D-NEXT: s /*size_t*/); +// CUTEXREFSETADDRESS2D-NEXT: cuTexRefSetAddress2D(t /*CUtexref*/, pa /*const CUDA_ARRAY_DESCRIPTOR **/, +// CUTEXREFSETADDRESS2D-NEXT: d /*CUdeviceptr*/, s /*size_t*/); // CUTEXREFSETADDRESS2D-NEXT: Is migrated to: // CUTEXREFSETADDRESS2D-NEXT: dpct::image_wrapper_base_p t; -// CUTEXREFSETADDRESS2D-NEXT: t->attach(d, s); +// CUTEXREFSETADDRESS2D-NEXT: t->attach(pa, d, s); // RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=cuTexRefSetAddressMode | FileCheck %s -check-prefix=CUTEXREFSETADDRESSMODE // CUTEXREFSETADDRESSMODE: CUDA API: diff --git a/clang/test/dpct/texture_driver.cu b/clang/test/dpct/texture_driver.cu index 1cde81aff322..6b4959709ae1 100644 --- a/clang/test/dpct/texture_driver.cu +++ b/clang/test/dpct/texture_driver.cu @@ -1,5 +1,6 @@ // RUN: dpct --format-range=none --usm-level=none -out-root %T/texture_driver %s --cuda-include-path="%cuda-path/include" --sycl-named-lambda -- -x cuda --cuda-host-only -std=c++14 // RUN: FileCheck --input-file %T/texture_driver/texture_driver.dp.cpp --match-full-lines %s +// RUN: %if build_lit %{icpx -c -fsycl %T/texture_driver/texture_driver.dp.cpp -o %T//texture_driver/texture_driver.dp.o %} #include "cuda.h" #include @@ -19,26 +20,22 @@ void funcT(T t) {} int main() { - // CHECK: size_t halfDesc_x_ct1, halfDesc_y_ct1; - // CHECK-NEXT: unsigned halfDesc_channel_num_ct1; - // CHECK-NEXT: sycl::image_channel_type halfDesc_channel_type_ct1; - // CHECK-NEXT: halfDesc_y_ct1 = 32; - // CHECK-NEXT: halfDesc_x_ct1 = 64; - // CHECK-NEXT: halfDesc_channel_type_ct1 = sycl::image_channel_type::fp16; - // CHECK-NEXT: halfDesc_channel_num_ct1 = 1; + // CHECK: dpct::image_desc halfDesc; + // CHECK-NEXT: halfDesc.height = 32; + // CHECK-NEXT: halfDesc.width = 64; + // CHECK-NEXT: halfDesc.channel_type = sycl::image_channel_type::fp16; + // CHECK-NEXT: halfDesc.channel_num = 1; CUDA_ARRAY_DESCRIPTOR halfDesc; halfDesc.Height = 32; halfDesc.Width = 64; halfDesc.Format = CU_AD_FORMAT_HALF; halfDesc.NumChannels = 1; - // CHECK: size_t float4Desc_x_ct1, float4Desc_y_ct1; - // CHECK-NEXT: unsigned float4Desc_channel_num_ct1; - // CHECK-NEXT: sycl::image_channel_type float4Desc_channel_type_ct1; - // CHECK-NEXT: float4Desc_x_ct1 = 64; - // CHECK-NEXT: float4Desc_channel_type_ct1 = sycl::image_channel_type::fp32; - // CHECK-NEXT: float4Desc_channel_num_ct1 = 4; - // CHECK-NEXT: float4Desc_y_ct1 = 32; + // CHECK: dpct::image_desc float4Desc; + // CHECK-NEXT: float4Desc.width = 64; + // CHECK-NEXT: float4Desc.channel_type = sycl::image_channel_type::fp32; + // CHECK-NEXT: float4Desc.channel_num = 4; + // CHECK-NEXT: float4Desc.height = 32; CUDA_ARRAY_DESCRIPTOR float4Desc; float4Desc.Width = 64; float4Desc.Format = CU_AD_FORMAT_FLOAT; @@ -47,8 +44,8 @@ int main() { // CHECK: dpct::image_matrix **a_ptr = new dpct::image_matrix_p; // CHECK-NEXT: dpct::image_matrix_p a42; - // CHECK-NEXT: *a_ptr = new dpct::image_matrix(halfDesc_channel_type_ct1, halfDesc_channel_num_ct1, halfDesc_x_ct1, halfDesc_y_ct1); - // CHECK-NEXT: a42 = new dpct::image_matrix(float4Desc_channel_type_ct1, float4Desc_channel_num_ct1, float4Desc_x_ct1, float4Desc_y_ct1); + // CHECK-NEXT: *a_ptr = new dpct::image_matrix((&halfDesc)->channel_type, (&halfDesc)->channel_num, (&halfDesc)->width, (&halfDesc)->height); + // CHECK-NEXT: a42 = new dpct::image_matrix((&float4Desc)->channel_type, (&float4Desc)->channel_num, (&float4Desc)->width, (&float4Desc)->height); // CHECK-NEXT: delete (*a_ptr); // CHECK-NEXT: delete a42; // CHECK-NEXT: delete a_ptr; @@ -64,25 +61,25 @@ int main() { { int errorCode; - // CHECK: errorCode = DPCT_CHECK_ERROR(a42 = new dpct::image_matrix(float4Desc_channel_type_ct1, float4Desc_channel_num_ct1, float4Desc_x_ct1, float4Desc_y_ct1)); + // CHECK: errorCode = DPCT_CHECK_ERROR(a42 = new dpct::image_matrix((&float4Desc)->channel_type, (&float4Desc)->channel_num, (&float4Desc)->width, (&float4Desc)->height)); errorCode = cuArrayCreate(&a42, &float4Desc); // CHECK: errorCode = DPCT_CHECK_ERROR(delete a42); errorCode = cuArrayDestroy(a42); - // CHECK: cudaCheck(DPCT_CHECK_ERROR(a42 = new dpct::image_matrix(float4Desc_channel_type_ct1, float4Desc_channel_num_ct1, float4Desc_x_ct1, float4Desc_y_ct1))); + // CHECK: cudaCheck(DPCT_CHECK_ERROR(a42 = new dpct::image_matrix((&float4Desc)->channel_type, (&float4Desc)->channel_num, (&float4Desc)->width, (&float4Desc)->height))); cudaCheck(cuArrayCreate(&a42, &float4Desc)); // CHECK: cudaCheck(DPCT_CHECK_ERROR(delete a42)); cudaCheck(cuArrayDestroy(a42)); - // CHECK: func(DPCT_CHECK_ERROR(a42 = new dpct::image_matrix(float4Desc_channel_type_ct1, float4Desc_channel_num_ct1, float4Desc_x_ct1, float4Desc_y_ct1))); + // CHECK: func(DPCT_CHECK_ERROR(a42 = new dpct::image_matrix((&float4Desc)->channel_type, (&float4Desc)->channel_num, (&float4Desc)->width, (&float4Desc)->height))); func(cuArrayCreate(&a42, &float4Desc)); // CHECK: func(DPCT_CHECK_ERROR(delete a42)); func(cuArrayDestroy(a42)); - // CHECK: funcT(DPCT_CHECK_ERROR(a42 = new dpct::image_matrix(float4Desc_channel_type_ct1, float4Desc_channel_num_ct1, float4Desc_x_ct1, float4Desc_y_ct1))); + // CHECK: funcT(DPCT_CHECK_ERROR(a42 = new dpct::image_matrix((&float4Desc)->channel_type, (&float4Desc)->channel_num, (&float4Desc)->width, (&float4Desc)->height))); funcT(cuArrayCreate(&a42, &float4Desc)); // CHECK: funcT(DPCT_CHECK_ERROR(delete a42)); funcT(cuArrayDestroy(a42)); @@ -92,26 +89,14 @@ int main() { void create_array_fail() { CUarray a; unsigned i; - // CHECK: CUDA_ARRAY_DESCRIPTOR d[20], *p; + // CHECK: dpct::image_desc d[20], *p; CUDA_ARRAY_DESCRIPTOR d[20], *p; p = &d[5]; - // CHECK: /* - // CHECK-NEXT: DPCT1073:{{[0-9]+}}: The field values of parameter 'd' could not be deduced, so the call was not migrated. You need to update this code manually. - // CHECK-NEXT: */ - // CHECK-NEXT: cuArrayCreate(&a, d); - // CHECK-NEXT: /* - // CHECK-NEXT: DPCT1073:{{[0-9]+}}: The field values of parameter 'p' could not be deduced, so the call was not migrated. You need to update this code manually. - // CHECK-NEXT: */ - // CHECK-NEXT: cuArrayCreate(&a, p); - // CHECK-NEXT: /* - // CHECK-NEXT: DPCT1073:{{[0-9]+}}: The field values of parameter 'p + i' could not be deduced, so the call was not migrated. You need to update this code manually. - // CHECK-NEXT: */ - // CHECK-NEXT: cuArrayCreate(&a, p + i); - // CHECK-NEXT: /* - // CHECK-NEXT: DPCT1073:{{[0-9]+}}: The field values of parameter '&d[i]' could not be deduced, so the call was not migrated. You need to update this code manually. - // CHECK-NEXT: */ - // CHECK-NEXT: cuArrayCreate(&a, &d[i]); + // CHECK: a = new dpct::image_matrix(d->channel_type, d->channel_num, d->width, d->height); + // CHECK-NEXT: a = new dpct::image_matrix(p->channel_type, p->channel_num, p->width, p->height); + // CHECK-NEXT: a = new dpct::image_matrix((p + i)->channel_type, (p + i)->channel_num, (p + i)->width, (p + i)->height); + // CHECK-NEXT: a = new dpct::image_matrix((&d[i])->channel_type, (&d[i])->channel_num, (&d[i])->width, (&d[i])->height); cuArrayCreate(&a, d); cuArrayCreate(&a, p); cuArrayCreate(&a, p + i); diff --git a/clang/test/dpct/texture_object_driver.cu b/clang/test/dpct/texture_object_driver.cu index f5602755c184..ca942677ac76 100644 --- a/clang/test/dpct/texture_object_driver.cu +++ b/clang/test/dpct/texture_object_driver.cu @@ -45,14 +45,12 @@ int main() { // CHECK: sycl::float4 *d_data42; // CHECK-NEXT: dpct::image_matrix_p a42; - // CHECK-NEXT: size_t desc42_x_ct1, desc42_y_ct1; - // CHECK-NEXT: unsigned desc42_channel_num_ct1; - // CHECK-NEXT: sycl::image_channel_type desc42_channel_type_ct1; - // CHECK-NEXT: desc42_channel_num_ct1 = 4; - // CHECK-NEXT: desc42_channel_type_ct1 = sycl::image_channel_type::fp32; - // CHECK-NEXT: desc42_x_ct1 = 32; - // CHECK-NEXT: desc42_y_ct1 = 32; - // CHECK-NEXT: a42 = new dpct::image_matrix(desc42_channel_type_ct1, desc42_channel_num_ct1, desc42_x_ct1, desc42_y_ct1); + // CHECK-NEXT: dpct::image_desc desc42; + // CHECK-NEXT: desc42.channel_num = 4; + // CHECK-NEXT: desc42.channel_type = sycl::image_channel_type::fp32; + // CHECK-NEXT: desc42.width = 32; + // CHECK-NEXT: desc42.height = 32; + // CHECK-NEXT: a42 = new dpct::image_matrix((&desc42)->channel_type, (&desc42)->channel_num, (&desc42)->width, (&desc42)->height); // CHECK-NEXT: dpct::image_wrapper_base_p tex42; // CHECK-NEXT: dpct::image_data res42; // CHECK-NEXT: dpct::sampling_info texDesc42; From d5344d71564896b757442119a02453a5c4bd37f0 Mon Sep 17 00:00:00 2001 From: "Tang, Jiajun" Date: Mon, 27 May 2024 15:54:07 +0800 Subject: [PATCH 2/3] fix comment. --- clang/lib/DPCT/APINamesTexture.inc | 14 ++++------- clang/lib/DPCT/MapNames.cpp | 2 +- clang/runtime/dpct-rt/include/dpct/image.hpp | 7 ++++-- clang/test/dpct/know_unsupported_type.cu | 2 +- clang/test/dpct/texture_driver.cu | 26 ++++++++++---------- clang/test/dpct/texture_object_driver.cu | 4 +-- 6 files changed, 27 insertions(+), 28 deletions(-) diff --git a/clang/lib/DPCT/APINamesTexture.inc b/clang/lib/DPCT/APINamesTexture.inc index 808ef18ee08a..2f552034ddcb 100644 --- a/clang/lib/DPCT/APINamesTexture.inc +++ b/clang/lib/DPCT/APINamesTexture.inc @@ -181,15 +181,11 @@ FEATURE_REQUEST_FACTORY(HelperFeatureEnum::device_ext, ENTRY_UNSUPPORTED("cudaGetTextureObjectResourceViewDesc", Diagnostics::API_NOT_MIGRATED) -FEATURE_REQUEST_FACTORY( - HelperFeatureEnum::device_ext, - ASSIGNABLE_FACTORY(ASSIGN_FACTORY_ENTRY( - "cuArrayCreate_v2", DEREF(0), - NEW(MapNames::getDpctNamespace() + "image_matrix", - MEMBER_EXPR(ARG(1), true, LITERAL("channel_type")), - MEMBER_EXPR(ARG(1), true, LITERAL("channel_num")), - MEMBER_EXPR(ARG(1), true, LITERAL("width")), - MEMBER_EXPR(ARG(1), true, LITERAL("height")))))) +FEATURE_REQUEST_FACTORY(HelperFeatureEnum::device_ext, + ASSIGNABLE_FACTORY(ASSIGN_FACTORY_ENTRY( + "cuArrayCreate_v2", DEREF(0), + NEW(MapNames::getDpctNamespace() + "image_matrix", + ARG(1))))) ASSIGNABLE_FACTORY(DELETER_FACTORY_ENTRY("cuArrayDestroy", ARG(0))) ENTRY_UNSUPPORTED("cuTexObjectGetResourceViewDesc", Diagnostics::API_NOT_MIGRATED) FEATURE_REQUEST_FACTORY( diff --git a/clang/lib/DPCT/MapNames.cpp b/clang/lib/DPCT/MapNames.cpp index ca1f220ea5cb..71f84479ae7c 100644 --- a/clang/lib/DPCT/MapNames.cpp +++ b/clang/lib/DPCT/MapNames.cpp @@ -374,7 +374,7 @@ void MapNames::setExplicitNamespaceMap() { {"cudaMemcpyKind", std::make_shared(getDpctNamespace() + "memcpy_direction")}, {"CUDA_ARRAY_DESCRIPTOR", - std::make_shared(getDpctNamespace() + "image_desc")}, + std::make_shared(getDpctNamespace() + "matrix_desc")}, {"cudaMemcpy3DParms", std::make_shared(getDpctNamespace() + "memcpy_parameter")}, {"CUDA_MEMCPY3D", diff --git a/clang/runtime/dpct-rt/include/dpct/image.hpp b/clang/runtime/dpct-rt/include/dpct/image.hpp index 1d8d1eb4b0ac..ede252864b53 100644 --- a/clang/runtime/dpct-rt/include/dpct/image.hpp +++ b/clang/runtime/dpct-rt/include/dpct/image.hpp @@ -139,7 +139,7 @@ static image_wrapper_base *create_image_wrapper(image_channel channel, int dims) } // namespace detail -struct image_desc { +struct matrix_desc { size_t width = 0; size_t height = 0; sycl::image_channel_type channel_type = @@ -339,6 +339,9 @@ class image_matrix { } _host_data = std::malloc(_range[0] * _range[1] * _channel.get_total_size()); } + image_matrix(const matrix_desc *desc) + : image_matrix(desc->channel_type, desc->channel_num, desc->width, + desc->height) {} /// Construct a new image class with the matrix data. template sycl::image *create_image() { @@ -635,7 +638,7 @@ class image_wrapper_base { image_data(const_cast(data), x, y, pitch, channel)); } /// Attach device_ptr data to this class. - void attach(const image_desc *desc, device_ptr ptr, size_t pitch) { + void attach(const matrix_desc *desc, device_ptr ptr, size_t pitch) { detach(); image_channel channel; channel.set_channel_num(desc->channel_num); diff --git a/clang/test/dpct/know_unsupported_type.cu b/clang/test/dpct/know_unsupported_type.cu index 8997de81ccdd..c69309b1cda9 100644 --- a/clang/test/dpct/know_unsupported_type.cu +++ b/clang/test/dpct/know_unsupported_type.cu @@ -8,7 +8,7 @@ #include "nvml.h" #include int main(int argc, char **argv) { - // CHECK: dpct::image_desc *pcad; + // CHECK: dpct::matrix_desc *pcad; CUDA_ARRAY_DESCRIPTOR *pcad; // CHECK: dpct::memcpy_parameter *p1c3d; cudaMemcpy3DParms *p1c3d; diff --git a/clang/test/dpct/texture_driver.cu b/clang/test/dpct/texture_driver.cu index 6b4959709ae1..cace318f5a23 100644 --- a/clang/test/dpct/texture_driver.cu +++ b/clang/test/dpct/texture_driver.cu @@ -20,7 +20,7 @@ void funcT(T t) {} int main() { - // CHECK: dpct::image_desc halfDesc; + // CHECK: dpct::matrix_desc halfDesc; // CHECK-NEXT: halfDesc.height = 32; // CHECK-NEXT: halfDesc.width = 64; // CHECK-NEXT: halfDesc.channel_type = sycl::image_channel_type::fp16; @@ -31,7 +31,7 @@ int main() { halfDesc.Format = CU_AD_FORMAT_HALF; halfDesc.NumChannels = 1; - // CHECK: dpct::image_desc float4Desc; + // CHECK: dpct::matrix_desc float4Desc; // CHECK-NEXT: float4Desc.width = 64; // CHECK-NEXT: float4Desc.channel_type = sycl::image_channel_type::fp32; // CHECK-NEXT: float4Desc.channel_num = 4; @@ -44,8 +44,8 @@ int main() { // CHECK: dpct::image_matrix **a_ptr = new dpct::image_matrix_p; // CHECK-NEXT: dpct::image_matrix_p a42; - // CHECK-NEXT: *a_ptr = new dpct::image_matrix((&halfDesc)->channel_type, (&halfDesc)->channel_num, (&halfDesc)->width, (&halfDesc)->height); - // CHECK-NEXT: a42 = new dpct::image_matrix((&float4Desc)->channel_type, (&float4Desc)->channel_num, (&float4Desc)->width, (&float4Desc)->height); + // CHECK-NEXT: *a_ptr = new dpct::image_matrix(&halfDesc); + // CHECK-NEXT: a42 = new dpct::image_matrix(&float4Desc); // CHECK-NEXT: delete (*a_ptr); // CHECK-NEXT: delete a42; // CHECK-NEXT: delete a_ptr; @@ -61,25 +61,25 @@ int main() { { int errorCode; - // CHECK: errorCode = DPCT_CHECK_ERROR(a42 = new dpct::image_matrix((&float4Desc)->channel_type, (&float4Desc)->channel_num, (&float4Desc)->width, (&float4Desc)->height)); + // CHECK: errorCode = DPCT_CHECK_ERROR(a42 = new dpct::image_matrix(&float4Desc)); errorCode = cuArrayCreate(&a42, &float4Desc); // CHECK: errorCode = DPCT_CHECK_ERROR(delete a42); errorCode = cuArrayDestroy(a42); - // CHECK: cudaCheck(DPCT_CHECK_ERROR(a42 = new dpct::image_matrix((&float4Desc)->channel_type, (&float4Desc)->channel_num, (&float4Desc)->width, (&float4Desc)->height))); + // CHECK: cudaCheck(DPCT_CHECK_ERROR(a42 = new dpct::image_matrix(&float4Desc))); cudaCheck(cuArrayCreate(&a42, &float4Desc)); // CHECK: cudaCheck(DPCT_CHECK_ERROR(delete a42)); cudaCheck(cuArrayDestroy(a42)); - // CHECK: func(DPCT_CHECK_ERROR(a42 = new dpct::image_matrix((&float4Desc)->channel_type, (&float4Desc)->channel_num, (&float4Desc)->width, (&float4Desc)->height))); + // CHECK: func(DPCT_CHECK_ERROR(a42 = new dpct::image_matrix(&float4Desc))); func(cuArrayCreate(&a42, &float4Desc)); // CHECK: func(DPCT_CHECK_ERROR(delete a42)); func(cuArrayDestroy(a42)); - // CHECK: funcT(DPCT_CHECK_ERROR(a42 = new dpct::image_matrix((&float4Desc)->channel_type, (&float4Desc)->channel_num, (&float4Desc)->width, (&float4Desc)->height))); + // CHECK: funcT(DPCT_CHECK_ERROR(a42 = new dpct::image_matrix(&float4Desc))); funcT(cuArrayCreate(&a42, &float4Desc)); // CHECK: funcT(DPCT_CHECK_ERROR(delete a42)); funcT(cuArrayDestroy(a42)); @@ -89,14 +89,14 @@ int main() { void create_array_fail() { CUarray a; unsigned i; - // CHECK: dpct::image_desc d[20], *p; + // CHECK: dpct::matrix_desc d[20], *p; CUDA_ARRAY_DESCRIPTOR d[20], *p; p = &d[5]; - // CHECK: a = new dpct::image_matrix(d->channel_type, d->channel_num, d->width, d->height); - // CHECK-NEXT: a = new dpct::image_matrix(p->channel_type, p->channel_num, p->width, p->height); - // CHECK-NEXT: a = new dpct::image_matrix((p + i)->channel_type, (p + i)->channel_num, (p + i)->width, (p + i)->height); - // CHECK-NEXT: a = new dpct::image_matrix((&d[i])->channel_type, (&d[i])->channel_num, (&d[i])->width, (&d[i])->height); + // CHECK: a = new dpct::image_matrix(d); + // CHECK-NEXT: a = new dpct::image_matrix(p); + // CHECK-NEXT: a = new dpct::image_matrix(p + i); + // CHECK-NEXT: a = new dpct::image_matrix(&d[i]); cuArrayCreate(&a, d); cuArrayCreate(&a, p); cuArrayCreate(&a, p + i); diff --git a/clang/test/dpct/texture_object_driver.cu b/clang/test/dpct/texture_object_driver.cu index ca942677ac76..dba859dfbf54 100644 --- a/clang/test/dpct/texture_object_driver.cu +++ b/clang/test/dpct/texture_object_driver.cu @@ -45,12 +45,12 @@ int main() { // CHECK: sycl::float4 *d_data42; // CHECK-NEXT: dpct::image_matrix_p a42; - // CHECK-NEXT: dpct::image_desc desc42; + // CHECK-NEXT: dpct::matrix_desc desc42; // CHECK-NEXT: desc42.channel_num = 4; // CHECK-NEXT: desc42.channel_type = sycl::image_channel_type::fp32; // CHECK-NEXT: desc42.width = 32; // CHECK-NEXT: desc42.height = 32; - // CHECK-NEXT: a42 = new dpct::image_matrix((&desc42)->channel_type, (&desc42)->channel_num, (&desc42)->width, (&desc42)->height); + // CHECK-NEXT: a42 = new dpct::image_matrix(&desc42); // CHECK-NEXT: dpct::image_wrapper_base_p tex42; // CHECK-NEXT: dpct::image_data res42; // CHECK-NEXT: dpct::sampling_info texDesc42; From 61e1dc13fa8421a3ff21195a038504cc8e61d3af Mon Sep 17 00:00:00 2001 From: "Tang, Jiajun" Date: Fri, 31 May 2024 16:10:33 +0800 Subject: [PATCH 3/3] change matrix_desc name. --- clang/lib/DPCT/MapNames.cpp | 4 ++-- clang/runtime/dpct-rt/include/dpct/image.hpp | 6 +++--- clang/test/dpct/know_unsupported_type.cu | 2 +- clang/test/dpct/texture_driver.cu | 6 +++--- clang/test/dpct/texture_object_driver.cu | 2 +- 5 files changed, 10 insertions(+), 10 deletions(-) diff --git a/clang/lib/DPCT/MapNames.cpp b/clang/lib/DPCT/MapNames.cpp index 71f84479ae7c..f5cdcd1d3e9b 100644 --- a/clang/lib/DPCT/MapNames.cpp +++ b/clang/lib/DPCT/MapNames.cpp @@ -373,8 +373,8 @@ void MapNames::setExplicitNamespaceMap() { HelperFeatureEnum::device_ext)}, {"cudaMemcpyKind", std::make_shared(getDpctNamespace() + "memcpy_direction")}, - {"CUDA_ARRAY_DESCRIPTOR", - std::make_shared(getDpctNamespace() + "matrix_desc")}, + {"CUDA_ARRAY_DESCRIPTOR", std::make_shared( + getDpctNamespace() + "image_matrix_desc")}, {"cudaMemcpy3DParms", std::make_shared(getDpctNamespace() + "memcpy_parameter")}, {"CUDA_MEMCPY3D", diff --git a/clang/runtime/dpct-rt/include/dpct/image.hpp b/clang/runtime/dpct-rt/include/dpct/image.hpp index ede252864b53..be3fac88e695 100644 --- a/clang/runtime/dpct-rt/include/dpct/image.hpp +++ b/clang/runtime/dpct-rt/include/dpct/image.hpp @@ -139,7 +139,7 @@ static image_wrapper_base *create_image_wrapper(image_channel channel, int dims) } // namespace detail -struct matrix_desc { +struct image_matrix_desc { size_t width = 0; size_t height = 0; sycl::image_channel_type channel_type = @@ -339,7 +339,7 @@ class image_matrix { } _host_data = std::malloc(_range[0] * _range[1] * _channel.get_total_size()); } - image_matrix(const matrix_desc *desc) + image_matrix(const image_matrix_desc *desc) : image_matrix(desc->channel_type, desc->channel_num, desc->width, desc->height) {} @@ -638,7 +638,7 @@ class image_wrapper_base { image_data(const_cast(data), x, y, pitch, channel)); } /// Attach device_ptr data to this class. - void attach(const matrix_desc *desc, device_ptr ptr, size_t pitch) { + void attach(const image_matrix_desc *desc, device_ptr ptr, size_t pitch) { detach(); image_channel channel; channel.set_channel_num(desc->channel_num); diff --git a/clang/test/dpct/know_unsupported_type.cu b/clang/test/dpct/know_unsupported_type.cu index c69309b1cda9..a513bd8ffb03 100644 --- a/clang/test/dpct/know_unsupported_type.cu +++ b/clang/test/dpct/know_unsupported_type.cu @@ -8,7 +8,7 @@ #include "nvml.h" #include int main(int argc, char **argv) { - // CHECK: dpct::matrix_desc *pcad; + // CHECK: dpct::image_matrix_desc *pcad; CUDA_ARRAY_DESCRIPTOR *pcad; // CHECK: dpct::memcpy_parameter *p1c3d; cudaMemcpy3DParms *p1c3d; diff --git a/clang/test/dpct/texture_driver.cu b/clang/test/dpct/texture_driver.cu index cace318f5a23..65a29658d4ae 100644 --- a/clang/test/dpct/texture_driver.cu +++ b/clang/test/dpct/texture_driver.cu @@ -20,7 +20,7 @@ void funcT(T t) {} int main() { - // CHECK: dpct::matrix_desc halfDesc; + // CHECK: dpct::image_matrix_desc halfDesc; // CHECK-NEXT: halfDesc.height = 32; // CHECK-NEXT: halfDesc.width = 64; // CHECK-NEXT: halfDesc.channel_type = sycl::image_channel_type::fp16; @@ -31,7 +31,7 @@ int main() { halfDesc.Format = CU_AD_FORMAT_HALF; halfDesc.NumChannels = 1; - // CHECK: dpct::matrix_desc float4Desc; + // CHECK: dpct::image_matrix_desc float4Desc; // CHECK-NEXT: float4Desc.width = 64; // CHECK-NEXT: float4Desc.channel_type = sycl::image_channel_type::fp32; // CHECK-NEXT: float4Desc.channel_num = 4; @@ -89,7 +89,7 @@ int main() { void create_array_fail() { CUarray a; unsigned i; - // CHECK: dpct::matrix_desc d[20], *p; + // CHECK: dpct::image_matrix_desc d[20], *p; CUDA_ARRAY_DESCRIPTOR d[20], *p; p = &d[5]; diff --git a/clang/test/dpct/texture_object_driver.cu b/clang/test/dpct/texture_object_driver.cu index dba859dfbf54..d33b4246ae5d 100644 --- a/clang/test/dpct/texture_object_driver.cu +++ b/clang/test/dpct/texture_object_driver.cu @@ -45,7 +45,7 @@ int main() { // CHECK: sycl::float4 *d_data42; // CHECK-NEXT: dpct::image_matrix_p a42; - // CHECK-NEXT: dpct::matrix_desc desc42; + // CHECK-NEXT: dpct::image_matrix_desc desc42; // CHECK-NEXT: desc42.channel_num = 4; // CHECK-NEXT: desc42.channel_type = sycl::image_channel_type::fp32; // CHECK-NEXT: desc42.width = 32;