Author: Yaxun (Sam) Liu Date: 2022-10-07T10:46:04-04:00 New Revision: 107ee2613063183cb643cef97f0fad403508c9f0
URL: https://github.com/llvm/llvm-project/commit/107ee2613063183cb643cef97f0fad403508c9f0 DIFF: https://github.com/llvm/llvm-project/commit/107ee2613063183cb643cef97f0fad403508c9f0.diff LOG: [AMDGPU] Disable bool range metadata to workaround backend issue Currently there is a middle-end or backend issue https://github.com/llvm/llvm-project/issues/58176 which causes values loaded from bool pointer incorrect when bool range metadata is emitted. Temporarily disable bool range metadata until the backend issue is fixed. Reviewed by: Artem Belevich Differential Revision: https://reviews.llvm.org/D135269 Fixes: SWDEV-344137 Added: clang/test/CodeGenCUDA/bool-range.cu Modified: clang/lib/CodeGen/CGExpr.cpp Removed: ################################################################################ diff --git a/clang/lib/CodeGen/CGExpr.cpp b/clang/lib/CodeGen/CGExpr.cpp index b58d5c29a1f9e..af753767f0328 100644 --- a/clang/lib/CodeGen/CGExpr.cpp +++ b/clang/lib/CodeGen/CGExpr.cpp @@ -1748,7 +1748,10 @@ llvm::Value *CodeGenFunction::EmitLoadOfScalar(Address Addr, bool Volatile, if (EmitScalarRangeCheck(Load, Ty, Loc)) { // In order to prevent the optimizer from throwing away the check, don't // attach range metadata to the load. - } else if (CGM.getCodeGenOpts().OptimizationLevel > 0) + // TODO: Enable range metadata for AMDGCN after issue + // https://github.com/llvm/llvm-project/issues/58176 is fixed. + } else if (CGM.getCodeGenOpts().OptimizationLevel > 0 && + !CGM.getTriple().isAMDGCN()) if (llvm::MDNode *RangeInfo = getRangeForLoadFromType(Ty)) Load->setMetadata(llvm::LLVMContext::MD_range, RangeInfo); diff --git a/clang/test/CodeGenCUDA/bool-range.cu b/clang/test/CodeGenCUDA/bool-range.cu new file mode 100644 index 0000000000000..5dc68612e7191 --- /dev/null +++ b/clang/test/CodeGenCUDA/bool-range.cu @@ -0,0 +1,23 @@ +// RUN: %clang_cc1 -emit-llvm %s -O3 -o - -fcuda-is-device \ +// RUN: -triple nvptx64-unknown-unknown | FileCheck %s -check-prefixes=NV +// RUN: %clang_cc1 -emit-llvm %s -O3 -o - -fcuda-is-device \ +// RUN: -triple amdgcn-amd-amdhsa | FileCheck %s -check-prefixes=AMD + +#include "Inputs/cuda.h" + +// NV: %[[LD:[0-9]+]] = load i8, ptr %x,{{.*}} !range ![[MD:[0-9]+]] +// NV: store i8 %[[LD]], ptr %y +// NV: ![[MD]] = !{i8 0, i8 2} + +// Make sure bool loaded from memory is truncated and +// range metadata is not emitted. +// TODO: Re-enable range metadata after issue +// https://github.com/llvm/llvm-project/issues/58176 is fixed. + +// AMD: %[[LD:[0-9]+]] = load i8, ptr addrspace(1) %x.global +// AMD-NOT: !range +// AMD: %[[AND:[0-9]+]] = and i8 %[[LD]], 1 +// AMD: store i8 %[[AND]], ptr addrspace(1) %y.global +__global__ void test1(bool *x, bool *y) { + *y = *x != false; +} _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
