yaxunl created this revision.
yaxunl added a reviewer: tra.
Herald added a project: All.
yaxunl requested review of this revision.
HIP is able to unbundle archive of bundled bitcode.
However currently there are two bugs:
1. archives passed by -l: are not unbundled.
2. archives passed as input files are not unbundled
Archives passed by -l: should not be prefixed with
prefix `lib` and appended with '.a', but still need to be prefixed with
paths in -L options.
Archives passed as input files should not be prefixed
or appended with anything.
https://reviews.llvm.org/D133705
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,15 @@
+// RUN: %clang_cc1 -emit-llvm %s -O3 -o - -fcuda-is-device \
+// RUN: -triple nvptx64-unknown-unknown | FileCheck %s
+
+#include "Inputs/cuda.h"
+
+// Make sure bool loaded from memory is truncated and
+// range metadata is not emitted.
+
+// CHECK: %0 = load i8, i8* %x
+// CHECK: %1 = and i8 %0, 1
+// CHECK: store i8 %1, i8* %y
+// CHECK-NOT: !range
+__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
@@ -1791,7 +1791,8 @@
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)
+ } else if (CGM.getCodeGenOpts().OptimizationLevel > 0 &&
+ !(getLangOpts().CUDA && getLangOpts().CUDAIsDevice))
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,15 @@
+// RUN: %clang_cc1 -emit-llvm %s -O3 -o - -fcuda-is-device \
+// RUN: -triple nvptx64-unknown-unknown | FileCheck %s
+
+#include "Inputs/cuda.h"
+
+// Make sure bool loaded from memory is truncated and
+// range metadata is not emitted.
+
+// CHECK: %0 = load i8, i8* %x
+// CHECK: %1 = and i8 %0, 1
+// CHECK: store i8 %1, i8* %y
+// CHECK-NOT: !range
+__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
@@ -1791,7 +1791,8 @@
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)
+ } else if (CGM.getCodeGenOpts().OptimizationLevel > 0 &&
+ !(getLangOpts().CUDA && getLangOpts().CUDAIsDevice))
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
