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..2f552034ddcb 100644 --- a/clang/lib/DPCT/APINamesTexture.inc +++ b/clang/lib/DPCT/APINamesTexture.inc @@ -181,17 +181,11 @@ 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", + ARG(1))))) ASSIGNABLE_FACTORY(DELETER_FACTORY_ENTRY("cuArrayDestroy", ARG(0))) ENTRY_UNSUPPORTED("cuTexObjectGetResourceViewDesc", Diagnostics::API_NOT_MIGRATED) FEATURE_REQUEST_FACTORY( @@ -236,8 +230,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..f5cdcd1d3e9b 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_matrix_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..be3fac88e695 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_matrix_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; @@ -331,6 +339,9 @@ class image_matrix { } _host_data = std::malloc(_range[0] * _range[1] * _channel.get_total_size()); } + image_matrix(const image_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() { @@ -626,6 +637,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_matrix_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..a513bd8ffb03 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_matrix_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..65a29658d4ae 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_matrix_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_matrix_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); + // CHECK-NEXT: a42 = new dpct::image_matrix(&float4Desc); // 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)); 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))); 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))); 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))); 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_matrix_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); + // 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 f5602755c184..d33b4246ae5d 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_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); // CHECK-NEXT: dpct::image_wrapper_base_p tex42; // CHECK-NEXT: dpct::image_data res42; // CHECK-NEXT: dpct::sampling_info texDesc42;