Index: lib/Target/NVPTX/NVPTXAsmPrinter.cpp
===================================================================
--- lib/Target/NVPTX/NVPTXAsmPrinter.cpp
+++ lib/Target/NVPTX/NVPTXAsmPrinter.cpp
@@ -93,6 +93,11 @@
 
 #define DEPOTNAME "__local_depot"
 
+static cl::opt<bool>
+    NoCudaDebug("no-cuda-debug",
+                cl::desc("Do not mark ptx file as having debug info"),
+                cl::init(false));
+
 /// DiscoverDependentGlobals - Return a set of GlobalVariables on which \p V
 /// depends.
 static void
@@ -876,7 +881,7 @@
     O << ", texmode_independent";
 
   // FIXME: remove comment once debug info is properly supported.
-  if (MMI && MMI->hasDebugInfo())
+  if (MMI && MMI->hasDebugInfo() && !NoCudaDebug)
     O << "//, debug";
 
   O << "\n";
Index: test/DebugInfo/NVPTX/debug-info.ll
===================================================================
--- test/DebugInfo/NVPTX/debug-info.ll
+++ test/DebugInfo/NVPTX/debug-info.ll
@@ -1,4 +1,5 @@
-; RUN: llc < %s -mtriple=nvptx64-nvidia-cuda | FileCheck %s
+; RUN: llc < %s -mtriple=nvptx64-nvidia-cuda | FileCheck %s --check-prefixes=CHECK,DEBUG
+; RUN: llc < %s -mtriple=nvptx64-nvidia-cuda -no-cuda-debug | FileCheck %s --check-prefixes=CHECK,NODEBUG
 
 ; // Bitcode int this test case is reduced version of compiled code below:
 ;__device__ inline void res(float x, float y, float *res) { *res = x + y; }
@@ -9,7 +10,8 @@
 ;    res(a * x[i], y[i], &y[i]);
 ;}
 
-; CHECK: .target sm_{{[0-9]+}}//, debug
+; DEBUG: .target sm_{{[0-9]+}}//, debug
+; NODEBUG: .target sm_{{[0-9]+$}}
 
 ; CHECK: .visible .entry _Z5saxpyifPfS_(
 ; CHECK: .param .u32 {{.+}},