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.

5 files changed, 131 insertions, 0 deletions

diff --git a/test cases/cuda/4 shared/main.cu b/test cases/cuda/4 shared/main.cu
new file mode 100644
index 0000000..d251167
--- /dev/null
+++ b/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;
+}
diff --git a/test cases/cuda/4 shared/meson.build b/test cases/cuda/4 shared/meson.build
new file mode 100644
index 0000000..532aaeb
--- /dev/null
+++ b/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)
diff --git a/test cases/cuda/4 shared/shared/kernels.cu b/test cases/cuda/4 shared/shared/kernels.cu
new file mode 100644
index 0000000..41a9553
--- /dev/null
+++ b/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();
+}
+
diff --git a/test cases/cuda/4 shared/shared/kernels.h b/test cases/cuda/4 shared/shared/kernels.h
new file mode 100644
index 0000000..dbcb99d
--- /dev/null
+++ b/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
diff --git a/test cases/cuda/4 shared/shared/meson.build b/test cases/cuda/4 shared/shared/meson.build
new file mode 100644
index 0000000..5987916
--- /dev/null
+++ b/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)