Shuenhoy/vtk-m
1
0
Fork 0
mirror of https://gitlab.kitware.com/vtk/vtk-m synced 2024-09-16 17:22:55 +00:00
Code Issues 2 Actions Packages Projects Releases Wiki Activity

Merge topic 'turing_support'

Browse Source 
17a93921c Suppress CUDA 10 warnings about __host__ / __device__
23c5f1182 vtk-m brigand nows knows what workaround are only needed for CUDA 9.
36ca00794 Add turing to VTK-m support hardware now that CUDA 10 is out.

Acked-by: Kitware Robot <kwrobot@kitware.com>
Acked-by: Sujin Philip <sujin.philip@kitware.com>
Merge-request: !1417
This commit is contained in:
Robert Maynard 2018-09-20 19:35:43 +00:00 committed by Kitware Robot
commit 8d0f441e1b
3 changed files with 26 additions and 6 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

12
CMake/VTKmDeviceAdapters.cmake
View File

@ -188,18 +188,21 @@ if(VTKm_ENABLE_CUDA AND NOT TARGET vtkm::cuda)
# - Uses: --generate-code=arch=compute_60,code=sm_60
# 6 - volta
# - Uses: --generate-code=arch=compute_70,code=sm_70
# 7 - all
# 7 - turing
# - Uses: --generate-code=arch=compute_75code=sm_75
# 8 - all
# - Uses: --generate-code=arch=compute_30,code=sm_30
# - Uses: --generate-code=arch=compute_35,code=sm_35
# - Uses: --generate-code=arch=compute_50,code=sm_50
# - Uses: --generate-code=arch=compute_60,code=sm_60
# - Uses: --generate-code=arch=compute_70,code=sm_70
# - Uses: --generate-code=arch=compute_75,code=sm_75
# 8 - none
#
#specify the property
set(VTKm_CUDA_Architecture "native" CACHE STRING "Which GPU Architecture(s) to compile for")
set_property(CACHE VTKm_CUDA_Architecture PROPERTY STRINGS native fermi kepler maxwell pascal volta all none)
set_property(CACHE VTKm_CUDA_Architecture PROPERTY STRINGS native fermi kepler maxwell pascal volta turing all none)
#detect what the property is set too
if(VTKm_CUDA_Architecture STREQUAL "native")
@ -253,12 +256,15 @@ if(VTKm_ENABLE_CUDA AND NOT TARGET vtkm::cuda)
set(arch_flags --generate-code=arch=compute_60,code=sm_60)
elseif(VTKm_CUDA_Architecture STREQUAL "volta")
set(arch_flags --generate-code=arch=compute_70,code=sm_70)
elseif(VTKm_CUDA_Architecture STREQUAL "turing")
set(arch_flags --generate-code=arch=compute_75,code=sm_75)
elseif(VTKm_CUDA_Architecture STREQUAL "all")
set(arch_flags --generate-code=arch=compute_30,code=sm_30
--generate-code=arch=compute_35,code=sm_35
--generate-code=arch=compute_50,code=sm_50
--generate-code=arch=compute_60,code=sm_60
--generate-code=arch=compute_70,code=sm_70)
--generate-code=arch=compute_70,code=sm_70
--generate-code=arch=compute_75,code=sm_75)
endif()
string(REPLACE ";" " " arch_flags "${arch_flags}")

16
vtkm/internal/brigand.hpp
View File

@ -24,12 +24,19 @@
#define BRIGAND_COMP_CLANG
#endif
#endif
#if defined(__CUDACC__)
#if __CUDACC_VER_MAJOR__ == 9
#define BRIGAND_COMP_CUDA_9
#endif
#if __CUDACC_VER_MAJOR__ >= 9
#define BRIGAND_COMP_CUDA_9PLUS
#endif
#define BRIGAND_COMP_CUDA
#endif
#include <type_traits>
namespace brigand
{
@ -259,9 +266,12 @@ namespace brigand
};
template<std::size_t N, typename Seq> struct at_impl;
#if defined(BRIGAND_COMP_CUDA_9PLUS) || defined(BRIGAND_COMP_INTEL)
//Both CUDA 9 and the Intel 18 compiler series have a problem deducing the
//type so we are just going
#if defined(BRIGAND_COMP_CUDA_9) || defined(BRIGAND_COMP_INTEL)
//Both CUDA 9 and the Intel 18 compiler series have a problem when
//at_impl ::type typedef is produced through inheritance of a `T`
//that is deduced through an unimplemented static functions return type.
//So we don't do the inherit trick, but instead manually construct
//a ::type typedef ourself.
template <std::size_t N, template <typename...> class L, class... Ts>
struct at_impl<N, L<Ts...>>
{

4
vtkm/worklet/internal/DispatcherBase.h
View File

@ -485,6 +485,10 @@ inline void deduce(Trampoline&& trampoline, ContParams&& sig, Args&&... args)
#pragma diag_suppress 2885
#endif
#if (__CUDACC_VER_MAJOR__ >= 10)
#pragma diag_suppress 2905
#endif
#endif
//This is a separate function as the pragma guards can cause nvcc
//to have an internal compiler error (codegen #3028)