Fixes for CUDA compiler shared library linking.

Also adds test case for shared library linking.

Closes #4912, at least on Linux. The future 0.50.0 does not yet claim to
support CUDA on systems other than Linux and backends other than Ninja.

mesonbuild/compilers/cuda.py

@@ -12,11 +12,12 @@
 # See the License for the specific language governing permissions and
 # limitations under the License.
 
-import subprocess, os.path
+import re, subprocess, os.path
 
 from .. import mlog
 from ..mesonlib import EnvironmentException, Popen_safe
-from .compilers import Compiler, cuda_buildtype_args, cuda_optimization_args, cuda_debug_args
+from .compilers import (Compiler, cuda_buildtype_args, cuda_optimization_args,
+                        cuda_debug_args, CompilerType, get_gcc_soname_args)
 
 class CudaCompiler(Compiler):
     def __init__(self, exelist, version, is_cross, exe_wrapper=None):
@@ -28,12 +29,16 @@ class CudaCompiler(Compiler):
         self.id = 'nvcc'
         default_warn_args = []
         self.warn_args = {'1': default_warn_args,
-                          '2': default_warn_args + ['-Wextra'],
+                          '2': default_warn_args + ['-Xcompiler=-Wextra'],
-                          '3': default_warn_args + ['-Wextra', '-Wpedantic']}
+                          '3': default_warn_args + ['-Xcompiler=-Wextra',
+                                                    '-Xcompiler=-Wpedantic']}
 
     def needs_static_linker(self):
         return False
 
+    def get_always_args(self):
+        return []
+
     def get_display_language(self):
         return 'Cuda'
 
@@ -136,12 +141,24 @@ __global__ void kernel (void) {
         if pe.returncode != 0:
             raise EnvironmentException('Executables created by {0} compiler {1} are not runnable.'.format(self.language, self.name_string()))
 
+    @staticmethod
+    def _cook_link_args(args):
+        """
+        Converts GNU-style arguments -Wl,-arg,-arg
+        to NVCC-style arguments -Xlinker=-arg,-arg
+        """
+        return [re.sub('^-Wl,', '-Xlinker=', arg) for arg in args]
+
     def get_output_args(self, target):
         return ['-o', target]
 
     def name_string(self):
         return ' '.join(self.exelist)
 
+    def get_soname_args(self, *args):
+        rawargs = get_gcc_soname_args(CompilerType.GCC_STANDARD, *args)
+        return self._cook_link_args(rawargs)
+
     def get_dependency_gen_args(self, outtarget, outfile):
         return []
 
@@ -177,6 +194,9 @@ __global__ void kernel (void) {
             path = '.'
         return ['-I' + path]
 
+    def get_std_shared_lib_link_args(self):
+        return ['-shared']
+
     def depfile_for_object(self, objfile):
         return objfile + '.' + self.get_depfile_suffix()
 
@@ -190,16 +210,17 @@ __global__ void kernel (void) {
         return []
 
     def build_rpath_args(self, build_dir, from_dir, rpath_paths, build_rpath, install_rpath):
-        return []
+        rawargs = self.build_unix_rpath_args(build_dir, from_dir, rpath_paths, build_rpath, install_rpath)
+        return self._cook_link_args(rawargs)
 
     def get_linker_search_args(self, dirname):
-        return ['/LIBPATH:' + dirname]
+        return ['-L' + dirname]
 
     def linker_to_compiler_args(self, args):
-        return ['/link'] + args
+        return args
 
     def get_pic_args(self):
-        return []
+        return ['-Xcompiler=-fPIC']
 
     def compute_parameters_with_absolute_paths(self, parameter_list, build_dir):
         return []

test cases/cuda/4 shared/main.cu

@@ -0,0 +1,20 @@
+#include <stdio.h>
+#include <cuda_runtime.h>
+#include "shared/kernels.h"
+
+
+int main(int argc, char **argv) {
+    int cuda_devices = 0;
+    cudaGetDeviceCount(&cuda_devices);
+    if(cuda_devices == 0) {
+        printf("No Cuda hardware found. Exiting.\n");
+        return 0;
+    }
+
+    if(run_tests() != 0){
+        printf("CUDA tests failed! Exiting.\n");
+        return 0;
+    }
+
+    return 0;
+}

test cases/cuda/4 shared/meson.build

@@ -0,0 +1,6 @@
+project('simple', 'cuda', version : '1.0.0')
+
+subdir('shared')
+
+exe = executable('prog', 'main.cu', dependencies: libkernels)
+test('cudatest', exe)

test cases/cuda/4 shared/shared/kernels.cu

@@ -0,0 +1,14 @@
+#include <stdio.h>
+#include <cuda_runtime.h>
+#include "kernels.h"
+
+
+TAG_HIDDEN __global__ void kernel (void){
+}
+
+TAG_PUBLIC int run_tests(void) {
+  kernel<<<1,1>>>();
+
+  return (int)cudaDeviceSynchronize();
+}
+

test cases/cuda/4 shared/shared/kernels.h

@@ -0,0 +1,86 @@
+/* Include Guard */
+#ifndef SHARED_KERNELS_H
+#define SHARED_KERNELS_H
+
+/**
+ * Includes
+ */
+
+#include <cuda_runtime.h>
+
+
+/**
+ * Defines
+ */
+
+/**
+ * When building a library, it is a good idea to expose as few as possible
+ * internal symbols (functions, objects, data structures). Not only does it
+ * prevent users from relying on private portions of the library that are
+ * subject to change without any notice, but it can have performance
+ * advantages:
+ *
+ *   - It can make shared libraries link faster at dynamic-load time.
+ *   - It can make internal function calls faster by bypassing the PLT.
+ *
+ * Thus, the compilation should by default hide all symbols, while the API
+ * headers will explicitly mark public the few symbols the users are permitted
+ * to use with a PUBLIC tag. We also define a HIDDEN tag, since it may be
+ * required to explicitly tag certain C++ types as visible in order for
+ * exceptions to function correctly.
+ *
+ * Additional complexity comes from non-POSIX-compliant systems, which
+ * artificially impose a requirement on knowing whether we are building or
+ * using a DLL.
+ *
+ * The above commentary and below code is inspired from
+ *                   'https://gcc.gnu.org/wiki/Visibility'
+ */
+
+#if   defined(_WIN32) || defined(__CYGWIN__)
+# define TAG_ATTRIBUTE_EXPORT __declspec(dllexport)
+# define TAG_ATTRIBUTE_IMPORT __declspec(dllimport)
+# define TAG_ATTRIBUTE_HIDDEN
+#elif __GNUC__ >= 4
+# define TAG_ATTRIBUTE_EXPORT __attribute__((visibility("default")))
+# define TAG_ATTRIBUTE_IMPORT __attribute__((visibility("default")))
+# define TAG_ATTRIBUTE_HIDDEN __attribute__((visibility("hidden")))
+#else
+# define TAG_ATTRIBUTE_EXPORT
+# define TAG_ATTRIBUTE_IMPORT
+# define TAG_ATTRIBUTE_HIDDEN
+#endif
+
+#if TAG_IS_SHARED
+# if TAG_IS_BUILDING
+#  define TAG_PUBLIC TAG_ATTRIBUTE_EXPORT
+# else
+#  define TAG_PUBLIC TAG_ATTRIBUTE_IMPORT
+# endif
+# define  TAG_HIDDEN TAG_ATTRIBUTE_HIDDEN
+#else
+# define  TAG_PUBLIC
+# define  TAG_HIDDEN
+#endif
+#define   TAG_STATIC static
+
+
+
+
+/* Extern "C" Guard */
+#ifdef __cplusplus
+extern "C" {
+#endif
+
+
+
+/* Function Prototypes */
+TAG_PUBLIC int run_tests(void);
+
+
+
+/* End Extern "C" and Include Guard */
+#ifdef __cplusplus
+}
+#endif
+#endif

test cases/cuda/4 shared/shared/meson.build

@@ -0,0 +1,5 @@
+libkernels = shared_library('kernels', 'kernels.cu',
+                            cuda_args: ['-DTAG_IS_SHARED=1', '-DTAG_IS_BUILDING=1'],
+                            gnu_symbol_visibility: 'hidden')
+libkernels = declare_dependency(compile_args: ['-DTAG_IS_SHARED=1'],
+                                link_with:    libkernels)