diff options
author | Jussi Pakkanen <jpakkane@gmail.com> | 2019-01-23 14:47:06 +0200 |
---|---|---|
committer | GitHub <noreply@github.com> | 2019-01-23 14:47:06 +0200 |
commit | 379b42c5b18cd94feded28791a5189ee253ba04e (patch) | |
tree | 32276f56c794a215bf7768b541e7fb5279e326a7 | |
parent | f72522e8890bb44283da970bf8e4545830ae81a6 (diff) | |
parent | 43147c2594f2d8d802113d1cb2bc5bb9c72d3b96 (diff) | |
download | meson-379b42c5b18cd94feded28791a5189ee253ba04e.tar.gz |
Merge pull request #3919 from mesonbuild/cudarebase
Add Cuda support
-rw-r--r-- | docs/markdown/snippets/cuda.md | 7 | ||||
-rw-r--r-- | mesonbuild/compilers/__init__.py | 2 | ||||
-rw-r--r-- | mesonbuild/compilers/compilers.py | 22 | ||||
-rw-r--r-- | mesonbuild/compilers/cuda.py | 202 | ||||
-rw-r--r-- | mesonbuild/environment.py | 44 | ||||
-rwxr-xr-x | run_project_tests.py | 1 | ||||
-rw-r--r-- | test cases/cuda/1 simple/meson.build | 5 | ||||
-rw-r--r-- | test cases/cuda/1 simple/prog.cu | 30 | ||||
-rw-r--r-- | test cases/cuda/2 split/lib.cu | 13 | ||||
-rw-r--r-- | test cases/cuda/2 split/main.cpp | 7 | ||||
-rw-r--r-- | test cases/cuda/2 split/meson.build | 7 | ||||
-rw-r--r-- | test cases/cuda/2 split/static/lib.cu | 13 | ||||
-rw-r--r-- | test cases/cuda/2 split/static/libsta.cu | 13 | ||||
-rw-r--r-- | test cases/cuda/2 split/static/main_static.cpp | 7 | ||||
-rw-r--r-- | test cases/cuda/2 split/static/meson.build | 4 |
15 files changed, 375 insertions, 2 deletions
diff --git a/docs/markdown/snippets/cuda.md b/docs/markdown/snippets/cuda.md new file mode 100644 index 000000000..a4a92cd71 --- /dev/null +++ b/docs/markdown/snippets/cuda.md @@ -0,0 +1,7 @@ +## Cuda support + +Compiling Cuda source code is now supported, though only with the +Ninja backend. This has been tested only on Linux for now. + +Because NVidia's Cuda compiler does not produce `.d` dependency files, +dependency tracking does not work. diff --git a/mesonbuild/compilers/__init__.py b/mesonbuild/compilers/__init__.py index c568a9862..60cca93c2 100644 --- a/mesonbuild/compilers/__init__.py +++ b/mesonbuild/compilers/__init__.py @@ -72,6 +72,7 @@ __all__ = [ 'JavaCompiler', 'LLVMDCompiler', 'MonoCompiler', + 'NvidiaCudaCompiler', 'VisualStudioCsCompiler', 'NAGFortranCompiler', 'ObjCCompiler', @@ -153,6 +154,7 @@ from .d import ( GnuDCompiler, LLVMDCompiler, ) +from .cuda import CudaCompiler from .fortran import ( FortranCompiler, G95FortranCompiler, diff --git a/mesonbuild/compilers/compilers.py b/mesonbuild/compilers/compilers.py index 016e704ad..b1f3cc244 100644 --- a/mesonbuild/compilers/compilers.py +++ b/mesonbuild/compilers/compilers.py @@ -37,6 +37,7 @@ lib_suffixes = ('a', 'lib', 'dll', 'dylib', 'so') lang_suffixes = { 'c': ('c',), 'cpp': ('cpp', 'cc', 'cxx', 'c++', 'hh', 'hpp', 'ipp', 'hxx'), + 'cuda': ('cu',), # f90, f95, f03, f08 are for free-form fortran ('f90' recommended) # f, for, ftn, fpp are for fixed-form fortran ('f' or 'for' recommended) 'fortran': ('f90', 'f95', 'f03', 'f08', 'f', 'for', 'ftn', 'fpp'), @@ -58,7 +59,7 @@ clib_langs = ('objcpp', 'cpp', 'objc', 'c', 'fortran',) # List of languages that can be linked with C code directly by the linker # used in build.py:process_compilers() and build.py:get_dynamic_linker() # XXX: Add Rust to this? -clink_langs = ('d',) + clib_langs +clink_langs = ('d', 'cuda') + clib_langs clink_suffixes = () for _l in clink_langs + ('vala',): clink_suffixes += lang_suffixes[_l] @@ -69,6 +70,7 @@ soregex = re.compile(r'.*\.so(\.[0-9]+)?(\.[0-9]+)?(\.[0-9]+)?$') # Environment variables that each lang uses. cflags_mapping = {'c': 'CFLAGS', 'cpp': 'CXXFLAGS', + 'cuda': 'CUFLAGS', 'objc': 'OBJCFLAGS', 'objcpp': 'OBJCXXFLAGS', 'fortran': 'FFLAGS', @@ -143,6 +145,13 @@ armclang_buildtype_args = {'plain': [], 'custom': [], } +cuda_buildtype_args = {'plain': [], + 'debug': [], + 'debugoptimized': [], + 'release': [], + 'minsize': [], + } + arm_buildtype_args = {'plain': [], 'debug': ['-O0', '--debug'], 'debugoptimized': ['-O1', '--debug'], @@ -345,6 +354,17 @@ msvc_optimization_args = {'0': [], 's': ['/O1'], # Implies /Os. } +cuda_optimization_args = {'0': [], + 'g': ['-O0'], + '1': ['-O1'], + '2': ['-O2'], + '3': ['-O3', '-Otime'], + 's': ['-O3', '-Ospace'] + } + +cuda_debug_args = {False: [], + True: ['-g']} + clike_debug_args = {False: [], True: ['-g']} diff --git a/mesonbuild/compilers/cuda.py b/mesonbuild/compilers/cuda.py new file mode 100644 index 000000000..b5f16abf8 --- /dev/null +++ b/mesonbuild/compilers/cuda.py @@ -0,0 +1,202 @@ +# Copyright 2012-2017 The Meson development team + +# Licensed under the Apache License, Version 2.0 (the "License"); +# you may not use this file except in compliance with the License. +# You may obtain a copy of the License at + +# http://www.apache.org/licenses/LICENSE-2.0 + +# Unless required by applicable law or agreed to in writing, software +# distributed under the License is distributed on an "AS IS" BASIS, +# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +# See the License for the specific language governing permissions and +# limitations under the License. + +import 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 + +class CudaCompiler(Compiler): + def __init__(self, exelist, version, is_cross, exe_wrapper=None): + if not hasattr(self, 'language'): + self.language = 'cuda' + super().__init__(exelist, version) + self.is_cross = is_cross + self.exe_wrapper = exe_wrapper + self.id = 'nvcc' + default_warn_args = [] + self.warn_args = {'1': default_warn_args, + '2': default_warn_args + ['-Wextra'], + '3': default_warn_args + ['-Wextra', '-Wpedantic']} + + def needs_static_linker(self): + return False + + def get_display_language(self): + return 'Cuda' + + def get_no_stdinc_args(self): + return [] + + def sanity_check(self, work_dir, environment): + source_name = os.path.join(work_dir, 'sanitycheckcuda.cu') + binary_name = os.path.join(work_dir, 'sanitycheckcuda') + extra_flags = self.get_cross_extra_flags(environment, link=False) + if self.is_cross: + extra_flags += self.get_compile_only_args() + + code = ''' +__global__ void kernel (void) { + +} + + int main(int argc,char** argv){ + return 0; + } + ''' + + with open(source_name, 'w') as ofile: + ofile.write(code) + pc = subprocess.Popen(self.exelist + extra_flags + [source_name, '-o', binary_name]) + pc.wait() + if pc.returncode != 0: + raise EnvironmentException('Cuda compiler %s can not compile programs.' % self.name_string()) + if self.is_cross: + # Can't check if the binaries run so we have to assume they do + return + pe = subprocess.Popen(binary_name) + pe.wait() + if pe.returncode != 0: + raise EnvironmentException('Executables created by Cuda compiler %s are not runnable.' % self.name_string()) + + def get_compiler_check_args(self): + return super().get_compiler_check_args() + [] + + def has_header_symbol(self, hname, symbol, prefix, env, extra_args=None, dependencies=None): + if super().has_header_symbol(hname, symbol, prefix, env, extra_args, dependencies): + return True + if extra_args is None: + extra_args = [] + fargs = {'prefix': prefix, 'header': hname, 'symbol': symbol} + t = '''{prefix} + #include <{header}> + using {symbol}; + int main () {{ return 0; }}''' + return self.compiles(t.format(**fargs), env, extra_args, dependencies) + + def sanity_check_impl(self, work_dir, environment, sname, code): + mlog.debug('Sanity testing ' + self.get_display_language() + ' compiler:', ' '.join(self.exelist)) + mlog.debug('Is cross compiler: %s.' % str(self.is_cross)) + + extra_flags = [] + source_name = os.path.join(work_dir, sname) + binname = sname.rsplit('.', 1)[0] + if self.is_cross: + binname += '_cross' + if self.exe_wrapper is None: + # Linking cross built apps is painful. You can't really + # tell if you should use -nostdlib or not and for example + # on OSX the compiler binary is the same but you need + # a ton of compiler flags to differentiate between + # arm and x86_64. So just compile. + extra_flags += self.get_cross_extra_flags(environment, link=False) + extra_flags += self.get_compile_only_args() + else: + extra_flags += self.get_cross_extra_flags(environment, link=True) + # Is a valid executable output for all toolchains and platforms + binname += '.exe' + # Write binary check source + binary_name = os.path.join(work_dir, binname) + with open(source_name, 'w') as ofile: + ofile.write(code) + # Compile sanity check + cmdlist = self.exelist + extra_flags + [source_name] + self.get_output_args(binary_name) + pc, stdo, stde = Popen_safe(cmdlist, cwd=work_dir) + mlog.debug('Sanity check compiler command line:', ' '.join(cmdlist)) + mlog.debug('Sanity check compile stdout:') + mlog.debug(stdo) + mlog.debug('-----\nSanity check compile stderr:') + mlog.debug(stde) + mlog.debug('-----') + if pc.returncode != 0: + raise EnvironmentException('Compiler {0} can not compile programs.'.format(self.name_string())) + # Run sanity check + if self.is_cross: + if self.exe_wrapper is None: + # Can't check if the binaries run so we have to assume they do + return + cmdlist = self.exe_wrapper + [binary_name] + else: + cmdlist = [binary_name] + mlog.debug('Running test binary command: ' + ' '.join(cmdlist)) + pe = subprocess.Popen(cmdlist) + pe.wait() + if pe.returncode != 0: + raise EnvironmentException('Executables created by {0} compiler {1} are not runnable.'.format(self.language, self.name_string())) + + def get_output_args(self, target): + return ['-o', target] + + def name_string(self): + return ' '.join(self.exelist) + + def get_dependency_gen_args(self, outtarget, outfile): + return [] + + def get_compile_only_args(self): + return ['-c'] + + def get_no_optimization_args(self): + return ['-O0'] + + def get_optimization_args(self, optimization_level): + return cuda_optimization_args[optimization_level] + + def get_debug_args(self, is_debug): + return cuda_debug_args[is_debug] + + def get_linker_exelist(self): + return self.exelist[:] + + def get_linker_output_args(self, outputname): + return ['-o', outputname] + + def get_warn_args(self, level): + return self.warn_args[level] + + def get_buildtype_args(self, buildtype): + return cuda_buildtype_args[buildtype] + + def get_include_args(self, path, is_system): + if path == '': + path = '.' + return ['-I' + path] + + def depfile_for_object(self, objfile): + return objfile + '.' + self.get_depfile_suffix() + + def get_depfile_suffix(self): + return 'd' + + def get_buildtype_linker_args(self, buildtype): + return [] + + def get_std_exe_link_args(self): + return [] + + def build_rpath_args(self, build_dir, from_dir, rpath_paths, build_rpath, install_rpath): + return [] + + def get_linker_search_args(self, dirname): + return ['/LIBPATH:' + dirname] + + def linker_to_compiler_args(self, args): + return ['/link'] + args + + def get_pic_args(self): + return [] + + def compute_parameters_with_absolute_paths(self, parameter_list, build_dir): + return [] diff --git a/mesonbuild/environment.py b/mesonbuild/environment.py index 0809ac5f7..d8530206b 100644 --- a/mesonbuild/environment.py +++ b/mesonbuild/environment.py @@ -59,6 +59,7 @@ from .compilers import ( IntelFortranCompiler, JavaCompiler, MonoCompiler, + CudaCompiler, VisualStudioCsCompiler, NAGFortranCompiler, Open64FortranCompiler, @@ -99,6 +100,17 @@ known_cpu_families = ( 'x86_64' ) +# Environment variables that each lang uses. +cflags_mapping = {'c': 'CFLAGS', + 'cpp': 'CXXFLAGS', + 'cu': 'CUFLAGS', + 'objc': 'OBJCFLAGS', + 'objcpp': 'OBJCXXFLAGS', + 'fortran': 'FFLAGS', + 'd': 'DFLAGS', + 'vala': 'VALAFLAGS'} + + def detect_gcovr(version='3.1', log=False): gcovr_exe = 'gcovr' try: @@ -410,6 +422,7 @@ class Environment: self.default_d = ['ldc2', 'ldc', 'gdc', 'dmd'] self.default_fortran = ['gfortran', 'g95', 'f95', 'f90', 'f77', 'ifort', 'pgfortran'] self.default_java = ['javac'] + self.default_cuda = ['nvcc'] self.default_rust = ['rustc'] self.default_swift = ['swiftc'] self.default_vala = ['valac'] @@ -417,6 +430,7 @@ class Environment: self.default_strip = ['strip'] self.vs_static_linker = ['lib'] self.clang_cl_static_linker = ['llvm-lib'] + self.cuda_static_linker = ['nvlink'] self.gcc_static_linker = ['gcc-ar'] self.clang_static_linker = ['llvm-ar'] self.default_pkgconfig = ['pkg-config'] @@ -737,6 +751,25 @@ class Environment: def detect_cpp_compiler(self, want_cross): return self._detect_c_or_cpp_compiler('cpp', want_cross) + def detect_cuda_compiler(self, want_cross): + popen_exceptions = {} + compilers, ccache, is_cross, exe_wrap = self._get_compilers('cuda', want_cross) + for compiler in compilers: + if isinstance(compiler, str): + compiler = [compiler] + else: + raise EnvironmentException() + arg = '--version' + try: + p, out, err = Popen_safe(compiler + [arg]) + except OSError as e: + popen_exceptions[' '.join(compiler + [arg])] = e + continue + version = search_version(out) + cls = CudaCompiler + return cls(ccache + compiler, version, is_cross, exe_wrap) + raise EnvironmentException('Could not find suitable CUDA compiler: "' + ' '.join(compilers) + '"') + def detect_fortran_compiler(self, want_cross): popen_exceptions = {} compilers, ccache, is_cross, exe_wrap = self._get_compilers('fortran', want_cross) @@ -999,6 +1032,10 @@ class Environment: comp = self.detect_objc_compiler(False) if need_cross_compiler: cross_comp = self.detect_objc_compiler(True) + elif lang == 'cuda': + comp = self.detect_cuda_compiler(False) + if need_cross_compiler: + cross_comp = self.detect_cuda_compiler(True) elif lang == 'objcpp': comp = self.detect_objcpp_compiler(False) if need_cross_compiler: @@ -1057,7 +1094,12 @@ class Environment: if linker is not None: linkers = [linker] else: - if isinstance(compiler, compilers.VisualStudioCCompiler): + evar = 'AR' + if isinstance(compiler, compilers.CudaCompiler): + linkers = [self.cuda_static_linker, self.default_static_linker] + elif evar in os.environ: + linkers = [shlex.split(os.environ[evar])] + elif isinstance(compiler, compilers.VisualStudioCCompiler): linkers = [self.vs_static_linker, self.clang_cl_static_linker] elif isinstance(compiler, compilers.GnuCompiler): # Use gcc-ar if available; needed for LTO diff --git a/run_project_tests.py b/run_project_tests.py index 02897cecd..4c6ca3b01 100755 --- a/run_project_tests.py +++ b/run_project_tests.py @@ -542,6 +542,7 @@ def detect_tests_to_run(): ('objective c++', 'objcpp', backend not in (Backend.ninja, Backend.xcode) or mesonlib.is_windows() or not have_objcpp_compiler()), ('fortran', 'fortran', backend is not Backend.ninja or not shutil.which('gfortran')), ('swift', 'swift', backend not in (Backend.ninja, Backend.xcode) or not shutil.which('swiftc')), + ('cuda', 'cuda', backend not in (Backend.ninja, Backend.xcode) or not shutil.which('nvcc')), ('python3', 'python3', backend is not Backend.ninja), ('python', 'python', backend is not Backend.ninja), ('fpga', 'fpga', shutil.which('yosys') is None), diff --git a/test cases/cuda/1 simple/meson.build b/test cases/cuda/1 simple/meson.build new file mode 100644 index 000000000..19af734af --- /dev/null +++ b/test cases/cuda/1 simple/meson.build @@ -0,0 +1,5 @@ +project('simple', 'cuda', version : '1.0.0') + +exe = executable('prog', 'prog.cu') +test('cudatest', exe) + diff --git a/test cases/cuda/1 simple/prog.cu b/test cases/cuda/1 simple/prog.cu new file mode 100644 index 000000000..7eab6738c --- /dev/null +++ b/test cases/cuda/1 simple/prog.cu @@ -0,0 +1,30 @@ +#include <iostream> + +int main(int argc, char **argv) { + int cuda_devices = 0; + std::cout << "CUDA version: " << CUDART_VERSION << "\n"; + cudaGetDeviceCount(&cuda_devices); + if(cuda_devices == 0) { + std::cout << "No Cuda hardware found. Exiting.\n"; + return 0; + } + std::cout << "This computer has " << cuda_devices << " Cuda device(s).\n"; + cudaDeviceProp props; + cudaGetDeviceProperties(&props, 0); + std::cout << "Properties of device 0.\n\n"; + + std::cout << " Name: " << props.name << "\n"; + std::cout << " Global memory: " << props.totalGlobalMem << "\n"; + std::cout << " Shared memory: " << props.sharedMemPerBlock << "\n"; + std::cout << " Constant memory: " << props.totalConstMem << "\n"; + std::cout << " Block registers: " << props.regsPerBlock << "\n"; + + std::cout << " Warp size: " << props.warpSize << "\n"; + std::cout << " Threads per block: " << props.maxThreadsPerBlock << "\n"; + std::cout << " Max block dimensions: [ " << props.maxThreadsDim[0] << ", " << props.maxThreadsDim[1] << ", " << props.maxThreadsDim[2] << " ]" << "\n"; + std::cout << " Max grid dimensions: [ " << props.maxGridSize[0] << ", " << props.maxGridSize[1] << ", " << props.maxGridSize[2] << " ]" << "\n"; + std::cout << "\n"; + + return 0; +} + diff --git a/test cases/cuda/2 split/lib.cu b/test cases/cuda/2 split/lib.cu new file mode 100644 index 000000000..c0471d048 --- /dev/null +++ b/test cases/cuda/2 split/lib.cu @@ -0,0 +1,13 @@ +#include <stdio.h> +#include <iostream> + +__global__ void kernel (void){ +} + +int do_cuda_stuff() { + kernel<<<1,1>>>(); + + printf("Hello, World!\n"); + return 0; +} + diff --git a/test cases/cuda/2 split/main.cpp b/test cases/cuda/2 split/main.cpp new file mode 100644 index 000000000..e5e6bda06 --- /dev/null +++ b/test cases/cuda/2 split/main.cpp @@ -0,0 +1,7 @@ +#include<iostream> + +int do_cuda_stuff(); + +int main(int argc, char **argv) { + return do_cuda_stuff(); +} diff --git a/test cases/cuda/2 split/meson.build b/test cases/cuda/2 split/meson.build new file mode 100644 index 000000000..51bf6ce23 --- /dev/null +++ b/test cases/cuda/2 split/meson.build @@ -0,0 +1,7 @@ +project('simple', 'cuda', 'cpp') + +exe = executable('prog', 'main.cpp', 'lib.cu') +test('cudatest', exe) + +subdir('static') + diff --git a/test cases/cuda/2 split/static/lib.cu b/test cases/cuda/2 split/static/lib.cu new file mode 100644 index 000000000..c0471d048 --- /dev/null +++ b/test cases/cuda/2 split/static/lib.cu @@ -0,0 +1,13 @@ +#include <stdio.h> +#include <iostream> + +__global__ void kernel (void){ +} + +int do_cuda_stuff() { + kernel<<<1,1>>>(); + + printf("Hello, World!\n"); + return 0; +} + diff --git a/test cases/cuda/2 split/static/libsta.cu b/test cases/cuda/2 split/static/libsta.cu new file mode 100644 index 000000000..c0471d048 --- /dev/null +++ b/test cases/cuda/2 split/static/libsta.cu @@ -0,0 +1,13 @@ +#include <stdio.h> +#include <iostream> + +__global__ void kernel (void){ +} + +int do_cuda_stuff() { + kernel<<<1,1>>>(); + + printf("Hello, World!\n"); + return 0; +} + diff --git a/test cases/cuda/2 split/static/main_static.cpp b/test cases/cuda/2 split/static/main_static.cpp new file mode 100644 index 000000000..e5e6bda06 --- /dev/null +++ b/test cases/cuda/2 split/static/main_static.cpp @@ -0,0 +1,7 @@ +#include<iostream> + +int do_cuda_stuff(); + +int main(int argc, char **argv) { + return do_cuda_stuff(); +} diff --git a/test cases/cuda/2 split/static/meson.build b/test cases/cuda/2 split/static/meson.build new file mode 100644 index 000000000..9078198d5 --- /dev/null +++ b/test cases/cuda/2 split/static/meson.build @@ -0,0 +1,4 @@ +l = static_library('clib', 'lib.cu') +exe = executable('staexe', 'main_static.cpp', + link_with : l) +test('static Cuda test', exe) |