yaxunl created this revision. yaxunl added reviewers: tra, ronl. Herald added subscribers: kosarev, t-tye, tpr, dstuttard, kzhuravl. Herald added a project: All. yaxunl requested review of this revision. Herald added a subscriber: wdng.
Currently there is some backend issue 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. https://reviews.llvm.org/D135269 Files: clang/lib/CodeGen/CGExpr.cpp clang/test/CodeGenCUDA/bool-range.cu Index: clang/test/CodeGenCUDA/bool-range.cu =================================================================== --- /dev/null +++ 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" + +// Make sure bool loaded from memory is truncated and +// range metadata is not emitted. + +// NV: %[[LD:[0-9]+]] = load i8, ptr %x,{{.*}} !range ![[MD:[0-9]+]] +// NV: store i8 %[[LD]], ptr %y +// NV: ![[MD]] = !{i8 0, i8 2} + +// TODO: Re-enable range metadata after backend issue 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; +} Index: clang/lib/CodeGen/CGExpr.cpp =================================================================== --- clang/lib/CodeGen/CGExpr.cpp +++ clang/lib/CodeGen/CGExpr.cpp @@ -1789,7 +1789,9 @@ 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 backend issue is fixed. + } else if (CGM.getCodeGenOpts().OptimizationLevel > 0 && + !CGM.getTriple().isAMDGCN()) if (llvm::MDNode *RangeInfo = getRangeForLoadFromType(Ty)) Load->setMetadata(llvm::LLVMContext::MD_range, RangeInfo);
Index: clang/test/CodeGenCUDA/bool-range.cu =================================================================== --- /dev/null +++ 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" + +// Make sure bool loaded from memory is truncated and +// range metadata is not emitted. + +// NV: %[[LD:[0-9]+]] = load i8, ptr %x,{{.*}} !range ![[MD:[0-9]+]] +// NV: store i8 %[[LD]], ptr %y +// NV: ![[MD]] = !{i8 0, i8 2} + +// TODO: Re-enable range metadata after backend issue 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; +} Index: clang/lib/CodeGen/CGExpr.cpp =================================================================== --- clang/lib/CodeGen/CGExpr.cpp +++ clang/lib/CodeGen/CGExpr.cpp @@ -1789,7 +1789,9 @@ 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 backend issue is fixed. + } else if (CGM.getCodeGenOpts().OptimizationLevel > 0 && + !CGM.getTriple().isAMDGCN()) if (llvm::MDNode *RangeInfo = getRangeForLoadFromType(Ty)) Load->setMetadata(llvm::LLVMContext::MD_range, RangeInfo);
_______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits