From be925b253eb31ef49a9adf61218b98a52b28602b Mon Sep 17 00:00:00 2001 From: Deepak Raj H R Date: Thu, 16 Jan 2025 15:08:31 +0800 Subject: [PATCH] [SYCLomatic] Add migration support for compound OR assignment operator for CUDA texture (#2617) --- clang/lib/DPCT/RulesLang/RulesLangTexture.cpp | 3 +- clang/test/dpct/texture_object_driver.cu | 32 +++++++++++++++++++ 2 files changed, 34 insertions(+), 1 deletion(-) diff --git a/clang/lib/DPCT/RulesLang/RulesLangTexture.cpp b/clang/lib/DPCT/RulesLang/RulesLangTexture.cpp index c6bf90aa86e9..0aec74a8fa5c 100644 --- a/clang/lib/DPCT/RulesLang/RulesLangTexture.cpp +++ b/clang/lib/DPCT/RulesLang/RulesLangTexture.cpp @@ -733,7 +733,8 @@ const Expr *TextureRule::getAssignedBO(const Expr *E, ASTContext &Context) { } else if (auto BO = dyn_cast(E)) { // If E is BinaryOperator, return E only when it is assign expression, // otherwise return nullptr. - if (BO->getOpcode() == BO_Assign) + auto Opcode = BO->getOpcode(); + if (Opcode == BO_Assign || Opcode == BO_OrAssign) return BO; } else if (auto COCE = dyn_cast(E)) { if (COCE->getOperator() == OO_Equal) { diff --git a/clang/test/dpct/texture_object_driver.cu b/clang/test/dpct/texture_object_driver.cu index f5afd8b26a56..fb1a8926d9c2 100644 --- a/clang/test/dpct/texture_object_driver.cu +++ b/clang/test/dpct/texture_object_driver.cu @@ -258,4 +258,36 @@ REDEF_CUDA_RESOURCE_DESC check_typedefed_cuda_type(CUarray arr) { rdesc.res.array.hArray = arr; return rdesc; } + +enum TexSurfaceAttrType { + UnpackIntegersAsNormalizedFloats, + NormalizedAddressing, + PerformSRGBToLinearConversion +}; + +// CHECK: void checkOrAssign(TexSurfaceAttrType setting) { +// CHECK-NEXT: dpct::sampling_info texDesc; +// CHECK-NEXT: if (setting == UnpackIntegersAsNormalizedFloats) +// CHECK-NEXT: texDesc.set(sycl::coordinate_normalization_mode::unnormalized); +// CHECK-NEXT: if (setting == NormalizedAddressing) +// CHECK-NEXT: /* +// CHECK-NEXT: DPCT1074:{{[0-9]+}}: 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. +// CHECK-NEXT: */ +// CHECK-NEXT: texDesc.set(sycl::coordinate_normalization_mode::normalized); +// CHECK-NEXT: if (setting == PerformSRGBToLinearConversion) +// CHECK-NEXT: /* +// CHECK-NEXT: DPCT1074:{{[0-9]+}}: 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. +// CHECK-NEXT: */ +// CHECK-NEXT: texDesc.set(sycl::coordinate_normalization_mode::unnormalized); +// CHECK-NEXT: } +void checkOrAssign(TexSurfaceAttrType setting) { + CUDA_TEXTURE_DESC texDesc; + if (setting == UnpackIntegersAsNormalizedFloats) + texDesc.flags |= CU_TRSF_READ_AS_INTEGER; + if (setting == NormalizedAddressing) + texDesc.flags |= CU_TRSF_NORMALIZED_COORDINATES; + if (setting == PerformSRGBToLinearConversion) + texDesc.flags |= CU_TRSF_SRGB; +} + #undef REDEF_CUDA_RESOURCE_DESC