Merge pull request #25 from robertmaynard/workaround_cuda_11.0_centos7_bug

Workaround CUDA 11.0 ICE on centos7

nvbench/CMakeLists.txt

@@ -10,7 +10,6 @@ set(srcs
   device_manager.cu
   float64_axis.cu
   int64_axis.cu
-  json_printer.cu
   markdown_printer.cu
   named_values.cu
   option_parser.cu
@@ -27,6 +26,16 @@ set(srcs
   detail/state_generator.cu
 )
 
+# CUDA 11.0 can't compile json_printer without crashing
+# So for that version fall back to C++ with degraded
+# output ( no PTX version info )
+if(CMAKE_CUDA_COMPILER_ID STREQUAL NVIDIA AND
+   CMAKE_CUDA_COMPILER_VERSION VERSION_LESS 11.1)
+  list(APPEND srcs json_printer.cxx)
+else()
+  list(APPEND srcs json_printer.cu)
+endif()
+
 # Generate doc strings from md files:
 include("../cmake/FileToString.cmake")
 file_to_string("../docs/cli_help.md"
@@ -40,13 +49,16 @@ file_to_string("../docs/cli_help_axis.md"
   cli_help_axis_text
 )
 
+find_package(CUDAToolkit)
+
 add_library(nvbench STATIC ${srcs})
 add_library(nvbench::nvbench ALIAS nvbench)
 # TODO generator expressions for installed paths
 target_include_directories(nvbench PUBLIC "${NVBench_SOURCE_DIR}")
 target_include_directories(nvbench PRIVATE "${CMAKE_CURRENT_BINARY_DIR}")
+target_include_directories(nvbench PRIVATE "${CUDAToolkit_INCLUDE_DIRS}")
 target_link_libraries(nvbench PRIVATE fmt::fmt nvbench_json)
-target_compile_features(nvbench PUBLIC cuda_std_17)
+target_compile_features(nvbench PUBLIC cuda_std_17 PRIVATE cxx_std_17)
 set_target_properties(nvbench PROPERTIES
   ARCHIVE_OUTPUT_DIRECTORY "${NVBench_LIBRARY_OUTPUT_DIR}"
   LIBRARY_OUTPUT_DIRECTORY "${NVBench_LIBRARY_OUTPUT_DIR}"

nvbench/axes_metadata.cuh

@@ -116,7 +116,7 @@ axes_metadata::axes_metadata(nvbench::type_list<TypeAxes...>)
       typedef typename decltype(wrapped_type)::type type_list;
       auto axis = std::make_unique<nvbench::type_axis>(std::move(*names_iter++),
                                                        type_axis_index);
-      axis->set_inputs<type_list>();
+      axis->template set_inputs<type_list>();
       axes.push_back(std::move(axis));
     });
 }

nvbench/detail/l2flush.cuh

@@ -35,7 +35,8 @@ struct l2flush
       cudaDeviceGetAttribute(&m_l2_size, cudaDevAttrL2CacheSize, dev_id));
     if (m_l2_size > 0)
     {
-      NVBENCH_CUDA_CALL(cudaMalloc(&m_l2_buffer, m_l2_size));
+      void* buffer = m_l2_buffer;
+      NVBENCH_CUDA_CALL(cudaMalloc(&buffer, m_l2_size));
     }
   }
 

nvbench/detail/statistics.cuh

@@ -20,6 +20,7 @@
 
 #include <nvbench/detail/transform_reduce.cuh>
 
+#include <cmath>
 #include <functional>
 #include <limits>
 #include <numeric>

nvbench/device_info.cuh

@@ -215,7 +215,7 @@ try
   nvbench::detail::device_scope _{dev_id};
   cudaFuncAttributes attr{};
   NVBENCH_CUDA_CALL(
-    cudaFuncGetAttributes(&attr, nvbench::detail::noop_kernel_ptr));
+    cudaFuncGetAttributes(&attr, ((const void*)nvbench::detail::noop_kernel_ptr) ));
   return attr.ptxVersion * 10;
 }
 catch(...)

nvbench/json_printer.cxx

@@ -0,0 +1,19 @@
+/*
+ *  Copyright 2021 NVIDIA Corporation
+ *
+ *  Licensed under the Apache License, Version 2.0 with the LLVM exception
+ *  (the "License"); you may not use this file except in compliance with
+ *  the License.
+ *
+ *  You may obtain a copy of the License at
+ *
+ *      http://llvm.org/foundation/relicensing/LICENSE.txt
+ *
+ *  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.
+ */
+
+#include <nvbench/json_printer.cu>