diff --git a/mesonbuild/compilers/cuda.py b/mesonbuild/compilers/cuda.py index dd3018448..1ea67d398 100644 --- a/mesonbuild/compilers/cuda.py +++ b/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'], - '3': default_warn_args + ['-Wextra', '-Wpedantic']} + '2': default_warn_args + ['-Xcompiler=-Wextra'], + '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 [] diff --git a/test cases/cuda/4 shared/main.cu b/test cases/cuda/4 shared/main.cu new file mode 100644 index 000000000..d251167b3 --- /dev/null +++ b/test cases/cuda/4 shared/main.cu @@ -0,0 +1,20 @@ +#include +#include +#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 000000000..532aaebfb --- /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 000000000..41a95536f --- /dev/null +++ b/test cases/cuda/4 shared/shared/kernels.cu @@ -0,0 +1,14 @@ +#include +#include +#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 000000000..dbcb99d10 --- /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 + + +/** + * 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 000000000..59879166b --- /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)