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