[clang] 107ee26 - [AMDGPU] Disable bool range metadata to workaround backend issue

[Yaxun Liu via cfe-commits]

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;
+}