Fixed cling argument --cuda-path

interpreter/cling/lib/Interpreter/IncrementalCUDADeviceCompiler.h → interpreter/cling/include/cling/Interpreter/IncrementalCUDADeviceCompiler.h

@@ -87,6 +87,9 @@ namespace cling {
     ///\brief Contains the PTX code of the current input
     llvm::SmallString<1024> m_PTX_code;
 
+    ///\brief Keep the ptx compiler args for reflection during runtime.
+    std::vector<std::string> argv;
+
     ///\brief Add the include paths from the interpreter runtime to a argument
     /// list.
     ///

interpreter/cling/include/cling/Interpreter/Interpreter.h

@@ -693,6 +693,10 @@ namespace cling {
     clang::Sema& getSema() const;
     clang::DiagnosticsEngine& getDiagnostics() const;
 
+    IncrementalCUDADeviceCompiler* getCUDACompiler() const {
+      return m_CUDACompiler.get();
+    }
+
     ///\brief Create suitable default compilation options.
     CompilationOptions makeDefaultCompilationOpts() const;
 

interpreter/cling/lib/Interpreter/CIFactory.cpp

@@ -873,14 +873,16 @@ static void stringifyPreprocSetting(PreprocessorOptions& PPOpts,
   }
 
   static llvm::IntrusiveRefCntPtr<DiagnosticsEngine>
-  SetupDiagnostics(DiagnosticOptions& DiagOpts) {
+  SetupDiagnostics(DiagnosticOptions& DiagOpts, const std::string& ExeName) {
     // The compiler invocation is the owner of the diagnostic options.
     // Everything else points to them.
     llvm::IntrusiveRefCntPtr<clang::DiagnosticIDs> DiagIDs(new DiagnosticIDs());
 
     std::unique_ptr<TextDiagnosticPrinter>
       DiagnosticPrinter(new TextDiagnosticPrinter(cling::errs(), &DiagOpts));
 
+    DiagnosticPrinter->setPrefix(ExeName);
+
     llvm::IntrusiveRefCntPtr<DiagnosticsEngine>
       Diags(new DiagnosticsEngine(DiagIDs, &DiagOpts,
                                   DiagnosticPrinter.get(), /*Owns it*/ true));
@@ -1332,8 +1334,13 @@ static void stringifyPreprocSetting(PreprocessorOptions& PPOpts,
     // The compiler invocation is the owner of the diagnostic options.
     // Everything else points to them.
     DiagnosticOptions& DiagOpts = InvocationPtr->getDiagnosticOpts();
+    // add prefix to diagnostic messages if second compiler instance is existing
+    // e.g. in CUDA mode
+    std::string ExeName = "";
+    if (COpts.CUDAHost)
+      ExeName = COpts.CUDADevice ? "cling-ptx" : "cling";
     llvm::IntrusiveRefCntPtr<DiagnosticsEngine> Diags =
-        SetupDiagnostics(DiagOpts);
+        SetupDiagnostics(DiagOpts, ExeName);
     if (!Diags) {
       cling::errs() << "Could not setup diagnostic engine.\n";
       return nullptr;

interpreter/cling/lib/Interpreter/IncrementalCUDADeviceCompiler.cpp

@@ -7,8 +7,7 @@
 // LICENSE.TXT for details.
 //------------------------------------------------------------------------------
 
-#include "IncrementalCUDADeviceCompiler.h"
-
+#include "cling/Interpreter/IncrementalCUDADeviceCompiler.h"
 #include "cling/Interpreter/Interpreter.h"
 #include "cling/Interpreter/InvocationOptions.h"
 #include "cling/Interpreter/Transaction.h"
@@ -52,16 +51,15 @@ namespace cling {
 
     // cling -std=c++xx -Ox -x cuda -S --cuda-gpu-arch=sm_xx --cuda-device-only
     // ${include headers} ${-I/paths} [-v] [-g] ${m_CuArgs->additionalPtxOpt}
-    std::vector<std::string> argv = {
-        "cling",
-        m_CuArgs->cppStdVersion.c_str(),
-        "-O" + std::to_string(optLevel),
-        "-x",
-        "cuda",
-        "-S",
-        std::string("--cuda-gpu-arch=sm_")
-            .append(std::to_string(m_CuArgs->smVersion)),
-        "--cuda-device-only"};
+    argv = {"cling",
+            m_CuArgs->cppStdVersion.c_str(),
+            "-O" + std::to_string(optLevel),
+            "-x",
+            "cuda",
+            "-S",
+            std::string("--cuda-gpu-arch=sm_")
+                .append(std::to_string(m_CuArgs->smVersion)),
+            "--cuda-device-only"};
 
     addHeaderSearchPathFlags(argv, CI.getHeaderSearchOptsPtr());
 
@@ -159,6 +157,11 @@ namespace cling {
         additionalPtxOpt.push_back(s);
     }
 
+    // use custom CUDA SDK path
+    if(!invocationOptions.CompilerOpts.CUDAPath.empty()){
+      additionalPtxOpt.push_back("--cuda-path=" + invocationOptions.CompilerOpts.CUDAPath);
+    }
+
     enum FatBinFlags {
       AddressSize64 = 0x01,
       HasDebugInfo = 0x02,

interpreter/cling/lib/Interpreter/Interpreter.cpp

@@ -18,7 +18,6 @@
 #include "EnterUserCodeRAII.h"
 #include "ExternalInterpreterSource.h"
 #include "ForwardDeclPrinter.h"
-#include "IncrementalCUDADeviceCompiler.h"
 #include "IncrementalExecutor.h"
 #include "IncrementalParser.h"
 #include "MultiplexInterpreterCallbacks.h"
@@ -32,6 +31,7 @@
 #include "cling/Interpreter/DynamicExprInfo.h"
 #include "cling/Interpreter/DynamicLibraryManager.h"
 #include "cling/Interpreter/Exception.h"
+#include "cling/Interpreter/IncrementalCUDADeviceCompiler.h"
 #include "cling/Interpreter/LookupHelper.h"
 #include "cling/Interpreter/Transaction.h"
 #include "cling/Interpreter/Value.h"

interpreter/cling/test/CUDADeviceCode/CUDADefineArg.C

@@ -9,7 +9,7 @@
 
 // The Test checks whether a define argument (-DTEST=3) is passed to the PTX
 // compiler. If it works, it should not throw an error.
-// RUN: cat %s | %cling -DTEST=3 -x cuda -Xclang -verify 2>&1 | FileCheck %s
+// RUN: cat %s | %cling -DTEST=3 -x cuda --cuda-path=%cudapath -Xclang -verify 2>&1 | FileCheck %s
 // REQUIRES: cuda-runtime
 
 #include <iostream>

interpreter/cling/test/CUDADeviceCode/CUDAHostPrefix.C

@@ -9,7 +9,7 @@
 
 // The Test checks if a function with __host__ and __device__ prefix available
 // on host and device side.
-// RUN: cat %s | %cling -x cuda -Xclang -verify 2>&1 | FileCheck %s
+// RUN: cat %s | %cling -x cuda --cuda-path=%cudapath -Xclang -verify 2>&1 | FileCheck %s
 // REQUIRES: cuda-runtime
 
 .rawInput 1

interpreter/cling/test/CUDADeviceCode/CUDAInclude.C

@@ -9,7 +9,7 @@
 
 // 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
+// RUN: cat %s | %cling -DTEST_PATH="\"%/p/\"" -x cuda --cuda-path=%cudapath -Xclang -verify 2>&1 | FileCheck %s
 // REQUIRES: cuda-runtime
 
 #include <iostream>

interpreter/cling/test/CUDADeviceCode/CUDAKernelArgument.C

@@ -9,7 +9,7 @@
 
 // The Test checks if a CUDA kernel works with a arguments and built-in
 // functions.
-// RUN: cat %s | %cling -x cuda -Xclang -verify 2>&1 | FileCheck %s
+// RUN: cat %s | %cling -x cuda --cuda-path=%cudapath -Xclang -verify 2>&1 | FileCheck %s
 // REQUIRES: cuda-runtime
 
 // Test, if a simple kernel with arguments works.

interpreter/cling/test/CUDADeviceCode/CUDAKernelTemplateComplex.C

@@ -8,7 +8,7 @@
 //------------------------------------------------------------------------------
 
 // The Test checks if templated CUDA kernel in some special cases works.
-// RUN: cat %s | %cling -x cuda -Xclang -verify 2>&1 | FileCheck %s
+// RUN: cat %s | %cling -x cuda --cuda-path=%cudapath -Xclang -verify 2>&1 | FileCheck %s
 // REQUIRES: cuda-runtime
 
 // Check if templated CUDA kernel works, without explicit template type declaration.

interpreter/cling/test/CUDADeviceCode/CUDAKernelTemplateSimple.C

@@ -8,7 +8,7 @@
 //------------------------------------------------------------------------------
 
 // The Test checks if templated CUDA kernel works.
-// RUN: cat %s | %cling -x cuda -Xclang -verify 2>&1 | FileCheck %s
+// RUN: cat %s | %cling -x cuda --cuda-path=%cudapath -Xclang -verify 2>&1 | FileCheck %s
 // REQUIRES: cuda-runtime
 
 // Check if template device side resoultion works.

interpreter/cling/test/CUDADeviceCode/CUDARegression.C

@@ -9,7 +9,7 @@
 
 // The test checks if the interface functions process(), declare() and parse()
 // of cling::Interpreter also work in the cuda mode.
-// RUN: cat %s | %cling -DTEST_PATH="\"%/p/\"" -x cuda -Xclang -verify 2>&1 | FileCheck %s
+// RUN: cat %s | %cling -DTEST_PATH="\"%/p/\"" -x cuda --cuda-path=%cudapath -Xclang -verify 2>&1 | FileCheck %s
 // REQUIRES: cuda-runtime
 
 #include "cling/Interpreter/Interpreter.h"

interpreter/cling/test/CUDADeviceCode/CUDASharedMemory.C

@@ -8,7 +8,7 @@
 //------------------------------------------------------------------------------
 
 // The Test checks if runtime shared memory works.
-// RUN: cat %s | %cling -x cuda -Xclang -verify 2>&1 | FileCheck %s
+// RUN: cat %s | %cling -x cuda --cuda-path=%cudapath -Xclang -verify 2>&1 | FileCheck %s
 // REQUIRES: cuda-runtime
 
 const unsigned int numberOfThreads = 4;

interpreter/cling/test/CUDADeviceCode/CUDASimpleKernel.C

@@ -9,7 +9,7 @@
 
 // The Test checks if a CUDA compatible device is available and checks, if simple
 // __global__ and __device__ kernels are running.
-// RUN: cat %s | %cling -x cuda -Xclang -verify 2>&1 | FileCheck %s
+// RUN: cat %s | %cling -x cuda --cuda-path=%cudapath -Xclang -verify 2>&1 | FileCheck %s
 // REQUIRES: cuda-runtime
 
 // Check if cuda driver is available

interpreter/cling/test/CUDADeviceCode/CUDAStreams.C

@@ -8,7 +8,7 @@
 //------------------------------------------------------------------------------
 
 // The Test checks if cuda streams works.
-// RUN: cat %s | %cling -x cuda -Xclang -verify 2>&1 | FileCheck %s
+// RUN: cat %s | %cling -x cuda --cuda-path=%cudapath -Xclang -verify 2>&1 | FileCheck %s
 // REQUIRES: cuda-runtime
 
 const unsigned int numberOfThreads = 4;

interpreter/cling/test/CodeGeneration/CUDACtorDtor.C

@@ -9,7 +9,7 @@
 // The Test checks, if the symbols __cuda_module_ctor and __cuda_module_dtor are
 // unique for every module. Attention, for a working test case, a cuda
 // fatbinary is necessary.
-// RUN: cat %s | %cling -x cuda -Xclang -verify 2>&1 | FileCheck %s
+// RUN: cat %s | %cling -x cuda --cuda-path=%cudapath -Xclang -verify 2>&1 | FileCheck %s
 // REQUIRES: cuda-runtime
 
 #include "cling/Interpreter/Interpreter.h"

interpreter/cling/test/Driver/CUDAMode.C

@@ -9,8 +9,8 @@
 
 // The Test starts the cling with the arg -xcuda and checks if the cuda mode
 // is enabled.
-// RUN: cat %s | %cling -x cuda -Xclang -verify 2>&1 | FileCheck %s
-// RUN: cat %s | %cling -x cuda -fsyntax-only -Xclang -verify 2>&1
+// RUN: cat %s | %cling -x cuda --cuda-path=%cudapath -Xclang -verify 2>&1 | FileCheck %s
+// RUN: cat %s | %cling -x cuda --cuda-path=%cudapath -fsyntax-only -Xclang -verify 2>&1
 // REQUIRES: cuda-runtime
 
 

interpreter/cling/test/lit.cfg

@@ -334,8 +334,16 @@ config.substitutions.append(('%testexecdir', config.test_exec_root))
 config.substitutions.append(('%llvmsrcdir', config.llvm_src_root))
 config.available_features.add('vanilla-cling')
 
-if lit.util.which('libcudart.so', config.environment.get('LD_LIBRARY_PATH','')) is not None:
+libcudart_path = lit.util.which('libcudart.so', config.environment.get('LD_LIBRARY_PATH',''))
+if libcudart_path is not None:
   config.available_features.add('cuda-runtime')
+  # set the CUDA SDK root path
+  # necessary if CUDA is not installed under /usr/local/
+  config.substitutions.append(('%cudapath', libcudart_path[0:-len('/lib64/libcudart.so')]))
+  # limit the number of usable GPUs in a system
+  # https://developer.nvidia.com/blog/cuda-pro-tip-control-gpu-visibility-cuda_visible_devices/
+  if 'CUDA_VISIBLE_DEVICES' in os.environ:
+      config.environment['CUDA_VISIBLE_DEVICES'] = os.environ['CUDA_VISIBLE_DEVICES']
 
 # Loadable module
 # FIXME: This should be supplied by Makefile or autoconf.