diff --git a/CMakeModules/AFcuda_helpers.cmake b/CMakeModules/AFcuda_helpers.cmake new file mode 100644 index 0000000000..4fde494df8 --- /dev/null +++ b/CMakeModules/AFcuda_helpers.cmake @@ -0,0 +1,60 @@ +# Copyright (c) 2020, ArrayFire +# All rights reserved. +# +# This file is distributed under 3-clause BSD license. +# The complete license agreement can be obtained at: +# http://arrayfire.com/licenses/BSD-3-Clause + + +# The following macro uses a macro defined by +# FindCUDA module from cmake. +function(af_find_static_cuda_libs libname) + set(search_name + "${CMAKE_STATIC_LIBRARY_PREFIX}${libname}${CMAKE_STATIC_LIBRARY_SUFFIX}") + cuda_find_library_local_first(CUDA_${libname}_LIBRARY + ${search_name} "${libname} static library") + mark_as_advanced(CUDA_${libname}_LIBRARY) +endfunction() + +## Copied from FindCUDA.cmake +## The target_link_library needs to link with the cuda libraries using +## PRIVATE +function(cuda_add_library cuda_target) + cuda_add_cuda_include_once() + + # Separate the sources from the options + cuda_get_sources_and_options(_sources _cmake_options _options ${ARGN}) + cuda_build_shared_library(_cuda_shared_flag ${ARGN}) + # Create custom commands and targets for each file. + cuda_wrap_srcs( ${cuda_target} OBJ _generated_files ${_sources} + ${_cmake_options} ${_cuda_shared_flag} + OPTIONS ${_options} ) + + # Compute the file name of the intermedate link file used for separable + # compilation. + cuda_compute_separable_compilation_object_file_name(link_file ${cuda_target} "${${cuda_target}_SEPARABLE_COMPILATION_OBJECTS}") + + # Add the library. + add_library(${cuda_target} ${_cmake_options} + ${_generated_files} + ${_sources} + ${link_file} + ) + + # Add a link phase for the separable compilation if it has been enabled. If + # it has been enabled then the ${cuda_target}_SEPARABLE_COMPILATION_OBJECTS + # variable will have been defined. + cuda_link_separable_compilation_objects("${link_file}" ${cuda_target} "${_options}" "${${cuda_target}_SEPARABLE_COMPILATION_OBJECTS}") + + target_link_libraries(${cuda_target} + PRIVATE ${CUDA_LIBRARIES} + ) + + # We need to set the linker language based on what the expected generated file + # would be. CUDA_C_OR_CXX is computed based on CUDA_HOST_COMPILATION_CPP. + set_target_properties(${cuda_target} + PROPERTIES + LINKER_LANGUAGE ${CUDA_C_OR_CXX} + POSITION_INDEPENDENT_CODE ON + ) +endfunction() diff --git a/src/backend/cuda/CMakeLists.txt b/src/backend/cuda/CMakeLists.txt index ae29c43d7a..b6059d2166 100644 --- a/src/backend/cuda/CMakeLists.txt +++ b/src/backend/cuda/CMakeLists.txt @@ -5,13 +5,18 @@ # The complete license agreement can be obtained at: # http://arrayfire.com/licenses/BSD-3-Clause +dependency_check(CUDA_FOUND "CUDA not found.") + +include(AFcuda_helpers) +include(FileToString) include(InternalUtils) include(select_compute_arch) -dependency_check(CUDA_FOUND "CUDA not found.") - -find_cuda_helper_libs(nvrtc) -find_cuda_helper_libs(nvrtc-builtins) +# Remove cublas_device library which is no longer included with the cuda +# toolkit. Fixes issues with older CMake versions +if(DEFINED CUDA_cublas_device_LIBRARY AND NOT CUDA_cublas_device_LIBRARY) + list(REMOVE_ITEM CUDA_CUBLAS_LIBRARIES ${CUDA_cublas_device_LIBRARY}) +endif() if(NOT OPENGL_FOUND) # create a dummy gl.h header to satisfy cuda_gl_interop.h requirement @@ -24,9 +29,50 @@ if(NOT OPENGL_FOUND) file(WRITE "${dummy_gl_root}/gl.h" "// Dummy file to satisy cuda_gl_interop") endif() -get_filename_component(CUDA_LIBRARIES_PATH ${CUDA_cudart_static_LIBRARY} DIRECTORY CACHE) +# Find if CUDA Toolkit is at least 10.0 to use static +# lapack library. Otherwise, we have to use regular shared library +if(UNIX AND CUDA_VERSION_MAJOR VERSION_GREATER 10 OR CUDA_VERSION_MAJOR VERSION_EQUAL 10) + set(use_static_cuda_lapack ON) +else() + set(use_static_cuda_lapack OFF) +endif() -include(FileToString) +find_cuda_helper_libs(nvrtc) +find_cuda_helper_libs(nvrtc-builtins) +if(UNIX) + af_find_static_cuda_libs(culibos) + af_find_static_cuda_libs(cublas_static) + af_find_static_cuda_libs(cublasLt_static) + af_find_static_cuda_libs(cufft_static) + af_find_static_cuda_libs(cusparse_static) + + # FIXME When NVCC resolves this particular issue. + # NVCC doesn't like -l, hence we cannot + # use ${CMAKE_*_LIBRARY} variables in the following flags. + set(af_cuda_static_flags "-rdc=true;-dlink") + set(af_cuda_static_flags "${af_cuda_static_flags};-lculibos") + set(af_cuda_static_flags "${af_cuda_static_flags};-lcublas_static") + set(af_cuda_static_flags "${af_cuda_static_flags};-lcublasLt_static") + set(af_cuda_static_flags "${af_cuda_static_flags};-lcufft_static") + set(af_cuda_static_flags "${af_cuda_static_flags};-lcusparse_static") + + if(${use_static_cuda_lapack}) + af_find_static_cuda_libs(cusolver_static) + set(cusolver_static_lib "${CUDA_cusolver_static_LIBRARY}") + + # NVIDIA LAPACK library liblapack_static.a is a subset of LAPACK and only + # contains GPU accelerated stedc and bdsqr. The user has to link + # libcusolver_static.a with liblapack_static.a in order to build + # successfully. + af_find_static_cuda_libs(lapack_static) + + set(af_cuda_static_flags "${af_cuda_static_flags};-lcusolver_static") + else() + set(cusolver_lib "${CUDA_cusolver_LIBRARY}") + endif() +endif() + +get_filename_component(CUDA_LIBRARIES_PATH ${CUDA_cudart_static_LIBRARY} DIRECTORY CACHE) if(NOT CUDA_architecture_build_targets) cuda_detect_installed_gpus(detected_gpus) @@ -39,6 +85,7 @@ cuda_select_nvcc_arch_flags(cuda_architecture_flags ${CUDA_architecture_build_ta message(STATUS "CUDA_architecture_build_targets: ${CUDA_architecture_build_targets}") set(CUDA_NVCC_FLAGS ${CUDA_NVCC_FLAGS};${cuda_architecture_flags}) + if(${CUDA_SEPARABLE_COMPILATION}) # Enable relocatable device code generation for separable # compilation which is in turn required for any device linking done. @@ -170,54 +217,9 @@ file_to_string( NULLTERM ) -## Copied from FindCUDA.cmake -## The target_link_library needs to link with the cuda libraries using -## PRIVATE -function(cuda_add_library cuda_target) - cuda_add_cuda_include_once() - - # Separate the sources from the options - cuda_get_sources_and_options(_sources _cmake_options _options ${ARGN}) - cuda_build_shared_library(_cuda_shared_flag ${ARGN}) - # Create custom commands and targets for each file. - cuda_wrap_srcs( ${cuda_target} OBJ _generated_files ${_sources} - ${_cmake_options} ${_cuda_shared_flag} - OPTIONS ${_options} ) - - # Compute the file name of the intermedate link file used for separable - # compilation. - cuda_compute_separable_compilation_object_file_name(link_file ${cuda_target} "${${cuda_target}_SEPARABLE_COMPILATION_OBJECTS}") - - # Add the library. - add_library(${cuda_target} ${_cmake_options} - ${_generated_files} - ${_sources} - ${link_file} - ) - - # Add a link phase for the separable compilation if it has been enabled. If - # it has been enabled then the ${cuda_target}_SEPARABLE_COMPILATION_OBJECTS - # variable will have been defined. - cuda_link_separable_compilation_objects("${link_file}" ${cuda_target} "${_options}" "${${cuda_target}_SEPARABLE_COMPILATION_OBJECTS}") - - target_link_libraries(${cuda_target} - PRIVATE ${CUDA_LIBRARIES} - ) - - # We need to set the linker language based on what the expected generated file - # would be. CUDA_C_OR_CXX is computed based on CUDA_HOST_COMPILATION_CPP. - set_target_properties(${cuda_target} - PROPERTIES - LINKER_LANGUAGE ${CUDA_C_OR_CXX} - POSITION_INDEPENDENT_CODE ON - ) - -endfunction() - arrayfire_get_cuda_cxx_flags(cuda_cxx_flags) arrayfire_get_platform_definitions(platform_flags) - get_property(boost_includes TARGET Boost::boost PROPERTY INTERFACE_INCLUDE_DIRECTORIES) get_property(boost_definitions TARGET Boost::boost PROPERTY INTERFACE_COMPILE_DEFINITIONS) @@ -244,8 +246,78 @@ list(APPEND cuda_cxx_flags ${cxx_definitions}) include(kernel/scan_by_key/CMakeLists.txt) include(kernel/thrust_sort_by_key/CMakeLists.txt) +# CUDA static libraries require device linking to successfully link +# against afcuda target. Device linking requires CUDA_SEPARABLE_COMPILATION +# to be ON. Therefore, we turn on separable compilation for a subset of +# source files while compiling af_cuda_static_cuda_library target. Once +# this subset is compiled, separable compilation is reset to it's original +# value. +if(UNIX) + # Static linking cuda libs require device linking, which in turn + # requires separable compilation. + set(pior_val_CUDA_SEPARABLE_COMPILATION OFF) + if(DEFINED CUDA_SEPARABLE_COMPILATION) + set(pior_val_CUDA_SEPARABLE_COMPILATION ${CUDA_SEPARABLE_COMPILATION}) + endif() + set(CUDA_SEPARABLE_COMPILATION ON) +endif() + +cuda_add_library(af_cuda_static_cuda_library STATIC + blas.cu + blas.hpp + cufft.cu + cufft.hpp + fft.cu + sparse.cu + sparse.hpp + sparse_arith.cu + sparse_arith.hpp + sparse_blas.cu + sparse_blas.hpp + solve.cu + solve.hpp + + OPTIONS + ${platform_flags} ${cuda_cxx_flags} ${af_cuda_static_flags} + -Xcudafe \"--diag_suppress=1427\" -DAFDLL +) + +set_target_properties(af_cuda_static_cuda_library + PROPERTIES + LINKER_LANGUAGE CXX + FOLDER "Generated Targets" +) + +if(UNIX) + target_link_libraries(af_cuda_static_cuda_library + PRIVATE + Boost::boost + ${CMAKE_DL_LIBS} + ${cusolver_lib} + -Wl,--start-group + ${CUDA_culibos_LIBRARY} #also a static libary + ${CUDA_cublas_static_LIBRARY} + ${CUDA_cublasLt_static_LIBRARY} + ${CUDA_cufft_static_LIBRARY} + ${CUDA_lapack_static_LIBRARY} + ${CUDA_cusparse_static_LIBRARY} + ${cusolver_static_lib} + -Wl,--end-group + ) + set(CUDA_SEPARABLE_COMPILATION ${pior_val_CUDA_SEPARABLE_COMPILATION}) +else() + target_link_libraries(af_cuda_static_cuda_library + PRIVATE + Boost::boost + ${CUDA_CUBLAS_LIBRARIES} + ${CUDA_CUFFT_LIBRARIES} + ${CUDA_cusolver_LIBRARY} + ${CUDA_cusparse_LIBRARY} + ) +endif() + cuda_add_library(afcuda - sort.hpp + ${thrust_sort_sources} all.cu anisotropic_diffusion.cpp @@ -388,7 +460,6 @@ cuda_add_library(afcuda backend.hpp bilateral.hpp binary.hpp - blas.cpp blas.hpp canny.hpp cast.hpp @@ -405,7 +476,6 @@ cuda_add_library(afcuda cudnn.hpp cudnnModule.cpp cudnnModule.hpp - cufft.cpp cufft.hpp cusolverDn.cpp cusolverDn.hpp @@ -425,7 +495,6 @@ cuda_add_library(afcuda fast.hpp fast_pyramid.cpp fast_pyramid.hpp - fft.cpp fft.hpp fftconvolve.cpp fftconvolve.hpp @@ -507,15 +576,12 @@ cuda_add_library(afcuda shift.hpp sift.hpp sobel.hpp - solve.cpp solve.hpp + sort.hpp sort_by_key.hpp sort_index.hpp - sparse.cpp sparse.hpp - sparse_arith.cpp sparse_arith.hpp - sparse_blas.cpp sparse_blas.hpp surface.cpp surface.hpp @@ -551,16 +617,25 @@ cuda_add_library(afcuda ${scan_by_key_sources} - OPTIONS ${platform_flags} ${cuda_cxx_flags} -Xcudafe \"--diag_suppress=1427\" + OPTIONS + ${platform_flags} + ${cuda_cxx_flags} + -Xcudafe \"--diag_suppress=1427\" ) arrayfire_set_default_cxx_flags(afcuda) +# NOTE: Do not add additional CUDA specific definitions here. Add it to the +# cxx_definitions variable above. cxx_definitions is used to propigate +# definitions to the scan_by_key and thrust_sort_by_key targets as well as the +# cuda library above. target_compile_options(afcuda PRIVATE ${cxx_definitions}) add_library(ArrayFire::afcuda ALIAS afcuda) add_dependencies(afcuda ${jit_kernel_targets} ${nvrtc_kernel_targets}) +add_dependencies(af_cuda_static_cuda_library ${nvrtc_kernel_targets}) +add_dependencies(afcuda af_cuda_static_cuda_library) target_include_directories (afcuda PUBLIC @@ -577,30 +652,14 @@ target_include_directories (afcuda ${cuDNN_INCLUDE_DIRS} ) -# Remove cublas_device library which is no longer included with the cuda -# toolkit. Fixes issues with older CMake versions -if(DEFINED CUDA_cublas_device_LIBRARY AND NOT CUDA_cublas_device_LIBRARY) - list(REMOVE_ITEM CUDA_CUBLAS_LIBRARIES ${CUDA_cublas_device_LIBRARY}) -endif() - -# Remove cublas_device library which is no longer included with the cuda -# toolkit. Fixes issues with older CMake versions -if(DEFINED CUDA_cublas_device_LIBRARY AND NOT CUDA_cublas_device_LIBRARY) - list(REMOVE_ITEM CUDA_CUBLAS_LIBRARIES ${CUDA_cublas_device_LIBRARY}) -endif() - target_link_libraries(afcuda PRIVATE c_api_interface cpp_api_interface afcommon_interface - cuda_thrust_sort_by_key - ${CUDA_nvrtc_LIBRARY} - ${CUDA_CUBLAS_LIBRARIES} - ${CUDA_CUFFT_LIBRARIES} - ${CUDA_cusolver_LIBRARY} - ${CUDA_cusparse_LIBRARY} ${CMAKE_DL_LIBS} + ${CUDA_nvrtc_LIBRARY} + af_cuda_static_cuda_library ) # If the driver is not found the cuda driver api need to be linked against the @@ -695,13 +754,17 @@ function(afcu_collect_libs libname) endfunction() if(AF_INSTALL_STANDALONE) - afcu_collect_libs(cufft) afcu_collect_libs(cudnn) - afcu_collect_libs(cublas) - afcu_collect_libs(cublasLt) - afcu_collect_libs(cusolver) - afcu_collect_libs(cusparse) afcu_collect_libs(nvrtc FULL_VERSION) + if(WIN32) + afcu_collect_libs(cufft) + afcu_collect_libs(cublas) + afcu_collect_libs(cublasLt) + afcu_collect_libs(cusolver) + afcu_collect_libs(cusparse) + elseif(NOT ${use_static_cuda_lapack}) + afcu_collect_libs(cusolver) + endif() if(APPLE) afcu_collect_libs(cudart) diff --git a/src/backend/cuda/blas.cpp b/src/backend/cuda/blas.cu similarity index 99% rename from src/backend/cuda/blas.cpp rename to src/backend/cuda/blas.cu index 4d61e6439e..188a426118 100644 --- a/src/backend/cuda/blas.cpp +++ b/src/backend/cuda/blas.cu @@ -7,11 +7,7 @@ * http://arrayfire.com/licenses/BSD-3-Clause ********************************************************/ -#define NVCC #include -#include -#include -#include #include #include @@ -20,11 +16,15 @@ #include #include #include +#include +#include #include #include +#include #include #include #include +#include #include #include diff --git a/src/backend/cuda/cublas.cpp b/src/backend/cuda/cublas.cpp index 29a0023a18..4f024b8117 100644 --- a/src/backend/cuda/cublas.cpp +++ b/src/backend/cuda/cublas.cpp @@ -7,8 +7,9 @@ * http://arrayfire.com/licenses/BSD-3-Clause ********************************************************/ -#include #include + +#include #include namespace cuda { diff --git a/src/backend/cuda/cufft.cpp b/src/backend/cuda/cufft.cu similarity index 99% rename from src/backend/cuda/cufft.cpp rename to src/backend/cuda/cufft.cu index 55fcdbb415..9dd976e9fe 100644 --- a/src/backend/cuda/cufft.cpp +++ b/src/backend/cuda/cufft.cu @@ -8,6 +8,7 @@ ********************************************************/ #include + #include #include diff --git a/src/backend/cuda/fft.cpp b/src/backend/cuda/fft.cu similarity index 99% rename from src/backend/cuda/fft.cpp rename to src/backend/cuda/fft.cu index bb1219171e..634f22daeb 100644 --- a/src/backend/cuda/fft.cpp +++ b/src/backend/cuda/fft.cu @@ -7,11 +7,12 @@ * http://arrayfire.com/licenses/BSD-3-Clause ********************************************************/ +#include + #include #include #include #include -#include #include #include #include diff --git a/src/backend/cuda/kernel/thrust_sort_by_key/CMakeLists.txt b/src/backend/cuda/kernel/thrust_sort_by_key/CMakeLists.txt index 3a6f660098..6c2f7f3c49 100644 --- a/src/backend/cuda/kernel/thrust_sort_by_key/CMakeLists.txt +++ b/src/backend/cuda/kernel/thrust_sort_by_key/CMakeLists.txt @@ -1,11 +1,13 @@ -# Copyright (c) 2017, ArrayFire +# Copyright (c) 2020, ArrayFire # All rights reserved. # # This file is distributed under 3-clause BSD license. # The complete license agreement can be obtained at: # http://arrayfire.com/licenses/BSD-3-Clause -file(STRINGS "${CMAKE_CURRENT_SOURCE_DIR}/kernel/thrust_sort_by_key/thrust_sort_by_key_impl.cu" FILESTRINGS) +file(STRINGS + "${CMAKE_CURRENT_SOURCE_DIR}/kernel/thrust_sort_by_key/thrust_sort_by_key_impl.cu" + FILESTRINGS) foreach(STR ${FILESTRINGS}) if(${STR} MATCHES "// SBK_TYPES") @@ -18,35 +20,18 @@ foreach(STR ${FILESTRINGS}) endforeach() foreach(SBK_TYPE ${SBK_TYPES}) - foreach(SBK_INST ${SBK_INSTS}) - - # When using cuda_compile with older versions of FindCUDA. The generated targets - # have the same names as the source file. Since we are using the same file for - # the compilation of these targets we need to rename them before sending them - # to the cuda_compile command so that it doesn't generate multiple targets with - # the same name - file(COPY "${CMAKE_CURRENT_SOURCE_DIR}/kernel/thrust_sort_by_key/thrust_sort_by_key_impl.cu" - DESTINATION "${CMAKE_CURRENT_BINARY_DIR}/kernel/thrust_sort_by_key") - file(RENAME "${CMAKE_CURRENT_BINARY_DIR}/kernel/thrust_sort_by_key/thrust_sort_by_key_impl.cu" - "${CMAKE_CURRENT_BINARY_DIR}/kernel/thrust_sort_by_key/thrust_sort_by_key_impl_${SBK_TYPE}_${SBK_INST}.cu") - - cuda_compile(sort_by_key_gen_files - ${CMAKE_CURRENT_BINARY_DIR}/kernel/thrust_sort_by_key/thrust_sort_by_key_impl_${SBK_TYPE}_${SBK_INST}.cu - ${CMAKE_CURRENT_SOURCE_DIR}/kernel/thrust_sort_by_key_impl.hpp - OPTIONS - -DSBK_TYPE=${SBK_TYPE} - -DINSTANTIATESBK_INST=INSTANTIATE${SBK_INST} - "${platform_flags} ${cuda_cxx_flags} -DAFDLL" - ) - - list(APPEND SORT_OBJ ${sort_by_key_gen_files}) - endforeach(SBK_INST ${SBK_INSTS}) + foreach(SBK_INST ${SBK_INSTS}) + set(INSTANTIATESBK_INST "INSTANTIATE${SBK_INST}") + + configure_file( + "${CMAKE_CURRENT_SOURCE_DIR}/kernel/thrust_sort_by_key/thrust_sort_by_key_impl.cu" + "${CMAKE_CURRENT_BINARY_DIR}/kernel/thrust_sort_by_key/thrust_sort_by_key_impl_${SBK_TYPE}_${SBK_INST}.cu" + ) + + list( + APPEND + thrust_sort_sources + "${CMAKE_CURRENT_BINARY_DIR}/kernel/thrust_sort_by_key/thrust_sort_by_key_impl_${SBK_TYPE}_${SBK_INST}.cu" + ) + endforeach(SBK_INST ${SBK_INSTS}) endforeach(SBK_TYPE ${SBK_TYPES}) - -cuda_add_library(cuda_thrust_sort_by_key STATIC ${SORT_OBJ}) - -set_target_properties(cuda_thrust_sort_by_key - PROPERTIES - LINKER_LANGUAGE CXX - FOLDER "Generated Targets" - ) diff --git a/src/backend/cuda/kernel/thrust_sort_by_key/thrust_sort_by_key_impl.cu b/src/backend/cuda/kernel/thrust_sort_by_key/thrust_sort_by_key_impl.cu index cf19942149..50996bb12e 100644 --- a/src/backend/cuda/kernel/thrust_sort_by_key/thrust_sort_by_key_impl.cu +++ b/src/backend/cuda/kernel/thrust_sort_by_key/thrust_sort_by_key_impl.cu @@ -16,6 +16,8 @@ namespace cuda { namespace kernel { -INSTANTIATESBK_INST(SBK_TYPE) -} +// clang-format off +@INSTANTIATESBK_INST@ ( @SBK_TYPE@ ) +// clang-format on +} // namespace kernel } // namespace cuda diff --git a/src/backend/cuda/solve.cpp b/src/backend/cuda/solve.cu similarity index 99% rename from src/backend/cuda/solve.cpp rename to src/backend/cuda/solve.cu index 4019170d2d..d45406a77c 100644 --- a/src/backend/cuda/solve.cpp +++ b/src/backend/cuda/solve.cu @@ -7,23 +7,20 @@ * http://arrayfire.com/licenses/BSD-3-Clause ********************************************************/ -#include #include +#include +#include #include #include #include #include +#include +#include #include #include -#include - -#include -#include - -#include -#include #include +#include #include diff --git a/src/backend/cuda/sparse.cpp b/src/backend/cuda/sparse.cu similarity index 100% rename from src/backend/cuda/sparse.cpp rename to src/backend/cuda/sparse.cu diff --git a/src/backend/cuda/sparse_arith.cpp b/src/backend/cuda/sparse_arith.cu similarity index 99% rename from src/backend/cuda/sparse_arith.cpp rename to src/backend/cuda/sparse_arith.cu index a4fe734224..66fad0bac2 100644 --- a/src/backend/cuda/sparse_arith.cpp +++ b/src/backend/cuda/sparse_arith.cu @@ -7,11 +7,7 @@ * http://arrayfire.com/licenses/BSD-3-Clause ********************************************************/ -#include -#include - -#include -#include +#include #include #include @@ -20,11 +16,16 @@ #include #include #include +#include #include #include #include +#include #include +#include +#include + namespace cuda { using namespace common; diff --git a/src/backend/cuda/sparse_blas.cpp b/src/backend/cuda/sparse_blas.cu similarity index 99% rename from src/backend/cuda/sparse_blas.cpp rename to src/backend/cuda/sparse_blas.cu index 59d462780f..eb7378776c 100644 --- a/src/backend/cuda/sparse_blas.cpp +++ b/src/backend/cuda/sparse_blas.cu @@ -7,14 +7,15 @@ * http://arrayfire.com/licenses/BSD-3-Clause ********************************************************/ -#include -#include -#include #include #include #include +#include +#include #include +#include + #include #include diff --git a/src/backend/cuda/types.hpp b/src/backend/cuda/types.hpp index 93e1704ed7..d18d747db5 100644 --- a/src/backend/cuda/types.hpp +++ b/src/backend/cuda/types.hpp @@ -148,13 +148,17 @@ struct kernel_type { using data = common::half; #ifdef __CUDA_ARCH__ + // These are the types within a kernel #if __CUDA_ARCH__ >= 530 && __CUDA_ARCH__ != 610 using compute = __half; #else using compute = float; #endif -#else + using native = compute; + +#else // __CUDA_ARCH__ + // outside of a cuda kernel use float using compute = float; @@ -163,6 +167,7 @@ struct kernel_type { #else using native = common::half; #endif -#endif + +#endif // __CUDA_ARCH__ }; } // namespace common