Revert "[AMDGPU] Disable bool range metadata to workaround backend issue"

This reverts commit 107ee2613063183cb643cef97f0fad403508c9f0 to facilitate
investigating and fixing the root cause.

Differential Revision: https://reviews.llvm.org/D135269

diff --git a/clang/lib/CodeGen/CGExpr.cpp b/clang/lib/CodeGen/CGExpr.cpp
index b2ba3f5..afdee52 100644 (file)
--- a/clang/lib/CodeGen/CGExpr.cpp
+++ b/clang/lib/CodeGen/CGExpr.cpp
@@ -1750,10 +1750,7 @@ 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.
-    // 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())
+  } else if (CGM.getCodeGenOpts().OptimizationLevel > 0)
     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
deleted file mode 100644 (file)
index e33174a..0000000
--- a/clang/test/CodeGenCUDA/bool-range.cu
+++ /dev/null
@@ -1,23 +0,0 @@
-// 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 {{.*}}%x
-// AMD-NOT: !range
-// AMD:  %[[AND:[0-9]+]] = and i8 %[[LD]], 1
-// AMD:  store i8 %[[AND]], ptr {{.*}}%y
-__global__ void test1(bool *x, bool *y) {
-  *y = *x != false;
-}