Setting a new include path at runtime in the PTX compiler now works

lib/Interpreter/IncrementalCUDADeviceCompiler.cpp

@@ -234,7 +234,6 @@ namespace cling {
   }
 
   bool cling::IncrementalCUDADeviceCompiler::generatePTX() {
-
     std::error_code EC;
     llvm::raw_fd_ostream os(m_PTXFilePath, EC, llvm::sys::fs::F_None);
     if (EC) {

lib/Interpreter/IncrementalCUDADeviceCompiler.h

@@ -156,6 +156,12 @@ namespace cling {
         const cling::InvocationOptions& invocationOptions,
         const clang::CompilerInstance& CI);
 
+    ///\brief Returns a reference to the PTX interpreter
+    ///
+    ///\return std::unique_ptr< cling::Interpreter >&
+    ///
+    std::unique_ptr<Interpreter>& getPTXInterpreter() { return m_PTX_interp; }
+
     ///\brief Generate an new fatbin file with the path in
     /// CudaGpuBinaryFileNames.
     /// It will add the content of input, to the existing source code, which was

lib/Interpreter/Interpreter.cpp

@@ -579,6 +579,9 @@ namespace cling {
                               E.Group == frontend::Angled);
       }
     }
+
+    if (m_CUDACompiler)
+      m_CUDACompiler->getPTXInterpreter()->AddIncludePaths(PathStr, Delm);
   }
 
   void Interpreter::AddIncludePath(llvm::StringRef PathsStr) {

test/CUDADeviceCode/CUDADefineArg.C

@@ -36,5 +36,12 @@ __global__ void g(){
 
 .rawInput 0
 
+// Runing the kernel is neccessary because FileCheck has problems whith the
+// error output of the PTX compiler. Therefore I need an error message from
+// the host interpreter.
+g<<<1,1>>>();
+cudaGetLastError()
+// CHECK: (cudaError_t) (cudaError::cudaSuccess) : (unsigned int) 0
+
 // expected-no-diagnostics
 .q

test/CUDADeviceCode/CUDAInclude.C

@@ -0,0 +1,45 @@
+//------------------------------------------------------------------------------
+// CLING - the C++ LLVM-based InterpreterG :)
+//
+// This file is dual-licensed: you can choose to license it under the University
+// of Illinois Open Source License or the GNU Lesser General Public License. See
+// LICENSE.TXT for details.
+//------------------------------------------------------------------------------
+
+// The test checks whether setting a new include path at runtime also works for
+// the PTX compiler.
+// RUN: cat %s | %cling -DTEST_PATH="\"%/p/\"" -x cuda -Xclang -verify 2>&1 | FileCheck %s
+// REQUIRES: cuda-runtime
+
+#include <iostream>
+
+// Check if cuda driver is available
+int version;
+cudaDriverGetVersion(&version)
+// CHECK: (cudaError_t) (cudaError::cudaSuccess) : (unsigned int) 0
+
+// Check if a CUDA compatible device (GPU) is available.
+int device_count = 0;
+cudaGetDeviceCount(&device_count)
+// CHECK: (cudaError_t) (cudaError::cudaSuccess) : (unsigned int) 0
+device_count > 0
+// CHECK: (bool) true
+
+#include "cling/Interpreter/Interpreter.h"
+gCling->AddIncludePaths(TEST_PATH "include");
+
+#include "foo.h"
+
+foo()
+// CHECK: (int) 3
+
+// Runing the kernel is neccessary because FileCheck has problems whith the
+// error output of the PTX compiler. Therefore I need an error message from
+// the host interpreter.
+bar<<<1,1>>>();
+cudaGetLastError()
+// CHECK: (cudaError_t) (cudaError::cudaSuccess) : (unsigned int) 0
+
+
+// expected-no-diagnostics
+.q

test/CUDADeviceCode/include/foo.h

@@ -0,0 +1,8 @@
+#ifndef FOO_H
+#define FOO_H
+
+int foo() { return 3; }
+
+__global__ void bar() { int i = 4; }
+
+#endif