mdspan icon indicating copy to clipboard operation
mdspan copied to clipboard
verified publisher icon kokkos

Compilation bug regarding `extents`

Open tpadioleau opened this issue 4 years ago • 4 comments

Bug report

Steps to reproduce

Compile the following snippet.

/// main.cpp
#include <experimental/mdspan>

namespace stdex = std::experimental;

int
main( int argc, char** argv )
{
  constexpr stdex::extents< stdex::dynamic_extent, stdex::dynamic_extent > ext1( 1ul, 1ul );

  constexpr stdex::extents< stdex::dynamic_extent >  ext2( std::array< std::size_t, 1 >{1ul} );

  return 0;
}

# CMakeLists.txt
cmake_minimum_required(VERSION 3.21)

project(bug-report LANGUAGES CUDA CXX)

find_package(mdspan CONFIG REQUIRED)

add_executable(bug-report main.cpp)
target_link_libraries(bug-report PUBLIC std::mdspan)
target_compile_features(bug-report PUBLIC cxx_std_17 cuda_std_17)
set_source_files_properties(main.cpp PROPERTIES LANGUAGE CUDA)

Environment

  • mdspan/0.1-0.2
  • nvcc/11.4.48
  • g++/9.2.0
  • cmake/3.21
  • cmake options: CMAKE_CUDA_ARCHITECTURES=70
  • need to disable no_unique_adress feature to avoid problem mentionned in #97 and #65

Expected result

Successful compilation

Actual result

/gpfs/users/tpadioleau/cmake-3.21.3-linux-x86_64/bin/cmake -S/gpfs/users/tpadioleau/bug-report -B/gpfs/users/tpadioleau/bug-report/build --check-build-system CMakeFiles/Makefile.cmake 0
/gpfs/users/tpadioleau/cmake-3.21.3-linux-x86_64/bin/cmake -E cmake_progress_start /gpfs/users/tpadioleau/bug-report/build/CMakeFiles /gpfs/users/tpadioleau/bug-report/build//CMakeFiles/progress.marks
make  -f CMakeFiles/Makefile2 all
make[1] : on entre dans le répertoire « /gpfs/users/tpadioleau/bug-report/build »
make  -f CMakeFiles/bug-report.dir/build.make CMakeFiles/bug-report.dir/depend
make[2] : on entre dans le répertoire « /gpfs/users/tpadioleau/bug-report/build »
cd /gpfs/users/tpadioleau/bug-report/build && /gpfs/users/tpadioleau/cmake-3.21.3-linux-x86_64/bin/cmake -E cmake_depends "Unix Makefiles" /gpfs/users/tpadioleau/bug-report /gpfs/users/tpadioleau/bug-report /gpfs/users/tpadioleau/bug-report/build /gpfs/users/tpadioleau/bug-report/build /gpfs/users/tpadioleau/bug-report/build/CMakeFiles/bug-report.dir/DependInfo.cmake --color=
make[2] : on quitte le répertoire « /gpfs/users/tpadioleau/bug-report/build »
make  -f CMakeFiles/bug-report.dir/build.make CMakeFiles/bug-report.dir/build
make[2] : on entre dans le répertoire « /gpfs/users/tpadioleau/bug-report/build »
[ 50%] Building CUDA object CMakeFiles/bug-report.dir/main.cpp.o
/gpfs/softs/spack/opt/spack/linux-centos7-cascadelake/gcc-9.2.0/cuda-11.4.0-tqftpmxbgstj4pkv3oeejmc57us33nq2/bin/nvcc -forward-unknown-to-host-compiler  -isystem=/gpfs/users/tpadioleau/opt/include --generate-code=arch=compute_70,code=[compute_70,sm_70] -std=c++17 -MD -MT CMakeFiles/bug-report.dir/main.cpp.o -MF CMakeFiles/bug-report.dir/main.cpp.o.d -x cu -c /gpfs/users/tpadioleau/bug-report/main.cpp -o CMakeFiles/bug-report.dir/main.cpp.o
/gpfs/users/tpadioleau/opt/include/experimental/__p0009_bits/default_accessor.hpp(61): warning: __host__ annotation is ignored on a function("default_accessor") that is explicitly defaulted on its first declaration

/gpfs/users/tpadioleau/opt/include/experimental/__p0009_bits/default_accessor.hpp(61): warning: __device__ annotation is ignored on a function("default_accessor") that is explicitly defaulted on its first declaration

/gpfs/users/tpadioleau/opt/include/experimental/__p0009_bits/standard_layout_static_array.hpp(542): warning: __host__ annotation is ignored on a function("__partially_static_sizes_tagged") that is explicitly defaulted on its first declaration

/gpfs/users/tpadioleau/opt/include/experimental/__p0009_bits/standard_layout_static_array.hpp(542): warning: __device__ annotation is ignored on a function("__partially_static_sizes_tagged") that is explicitly defaulted on its first declaration

/gpfs/users/tpadioleau/bug-report/main.cpp(8): warning: variable "ext1" was declared but never referenced

/gpfs/users/tpadioleau/bug-report/main.cpp(10): warning: variable "ext2" was declared but never referenced

/gpfs/users/tpadioleau/opt/include/experimental/__p0009_bits/default_accessor.hpp(61): warning: __host__ annotation is ignored on a function("default_accessor") that is explicitly defaulted on its first declaration

/gpfs/users/tpadioleau/opt/include/experimental/__p0009_bits/default_accessor.hpp(61): warning: __device__ annotation is ignored on a function("default_accessor") that is explicitly defaulted on its first declaration

/gpfs/users/tpadioleau/opt/include/experimental/__p0009_bits/standard_layout_static_array.hpp(542): warning: __host__ annotation is ignored on a function("__partially_static_sizes_tagged") that is explicitly defaulted on its first declaration

/gpfs/users/tpadioleau/opt/include/experimental/__p0009_bits/standard_layout_static_array.hpp(542): warning: __device__ annotation is ignored on a function("__partially_static_sizes_tagged") that is explicitly defaulted on its first declaration

/gpfs/users/tpadioleau/bug-report/main.cpp: In function 'int main(int, char**)':
/gpfs/users/tpadioleau/bug-report/main.cpp:10:117: error: no matching function for call to 'std::experimental::extents<18446744073709551615>::extents(std::array<long unsigned int, 1>)'
   10 |   constexpr stdex::extents< stdex::dynamic_extent >  ext2( std::array< std::size_t, 1 >{1ul} );
      |                                                                                                                     ^
/gpfs/users/tpadioleau/opt/include/experimental/__p0009_bits/extents.hpp:273:20: note: candidate: 'constexpr std::experimental::extents<Extents>::extents(const __storage_t&) [with long unsigned int ...Extents = {18446744073709551615}; std::experimental::extents<Extents>::__storage_t = std::experimental::detail::__partially_static_sizes_tagged<std::experimental::detail::__extents_tag, 18446744073709551615>]'
  273 |   extents(__storage_t const& sto ) noexcept
      |                    ^~~~~~~
/gpfs/users/tpadioleau/opt/include/experimental/__p0009_bits/extents.hpp:273:47: note:   no known conversion for argument 1 from 'std::array<long unsigned int, 1>' to 'const __storage_t&' {aka 'const std::experimental::detail::__partially_static_sizes_tagged<std::experimental::detail::__extents_tag, 18446744073709551615>&'}
  273 |   extents(__storage_t const& sto ) noexcept
      |                            ~~~~~~~~~~~~~~~~   ^  
/gpfs/users/tpadioleau/opt/include/experimental/__p0009_bits/extents.hpp:252:1: note: candidate: 'template<class SizeType, long unsigned int N, typename std::enable_if<(((N == std::experimental::extents<18446744073709551615, 18446744073709551615>::rank()) || (N == std::experimental::extents<18446744073709551615, 18446744073709551615>::rank_dynamic())) && is_convertible_v<SizeType, long unsigned int>), int>::type <anonymous> > constexpr std::experimental::extents<Extents>::extents(const std::array<_U, _N>&)'
  252 |   extents(std::array<SizeType, N> const& exts) noexcept
      | ^ ~~~~~
/gpfs/users/tpadioleau/opt/include/experimental/__p0009_bits/extents.hpp:252:1: note:   template argument deduction/substitution failed:
/gpfs/users/tpadioleau/opt/include/experimental/__p0009_bits/extents.hpp:242:209: error: no type named 'type' in 'struct std::enable_if<false, int>'
  242 |   MDSPAN_TEMPLATE_REQUIRES(
      |                                                                                                                                                                                                                 ^
/gpfs/users/tpadioleau/opt/include/experimental/__p0009_bits/extents.hpp:224:1: note: candidate: 'template<class ... Integral, typename std::enable_if<((is_convertible_v<Integral, long unsigned int> && ...) && ((sizeof... (Integral) == std::experimental::extents<18446744073709551615, 18446744073709551615>::rank_dynamic()) || (sizeof... (Integral) == std::experimental::extents<18446744073709551615, 18446744073709551615>::rank()))), int>::type <anonymous> > constexpr std::experimental::extents<Extents>::extents(Integral ...)'
  224 |   explicit constexpr extents(Integral... exts) noexcept
      | ^ ~~~~~
/gpfs/users/tpadioleau/opt/include/experimental/__p0009_bits/extents.hpp:224:1: note:   template argument deduction/substitution failed:
/gpfs/users/tpadioleau/opt/include/experimental/__p0009_bits/extents.hpp:215:248: error: no type named 'type' in 'struct std::enable_if<false, int>'
  215 |   MDSPAN_TEMPLATE_REQUIRES(
      |                                                                                                                                                                                                                                                        ^
/gpfs/users/tpadioleau/opt/include/experimental/__p0009_bits/extents.hpp:200:1: note: candidate: 'template<long unsigned int ...OtherExtents, typename std::enable_if<decltype (std::experimental::detail::_check_compatible_extents(std::integral_constant<bool, (1 == sizeof... (OtherExtents))>{}, std::integer_sequence<long unsigned int, 18446744073709551615>{}, std::integer_sequence<long unsigned int, _Idx ...>{}))::value, int>::type <anonymous> > constexpr std::experimental::extents<Extents>::extents(const std::experimental::extents<OtherExtents ...>&)'
  200 |   constexpr extents(const extents<OtherExtents...>& __other)
      | ^ ~~~~~
/gpfs/users/tpadioleau/opt/include/experimental/__p0009_bits/extents.hpp:200:1: note:   template argument deduction/substitution failed:
/gpfs/users/tpadioleau/bug-report/main.cpp:10:117: note:   'std::array<long unsigned int, 1>' is not derived from 'const std::experimental::extents<ExtentsPack ...>'
   10 |   constexpr stdex::extents< stdex::dynamic_extent >  ext2( std::array< std::size_t, 1 >{1ul} );
      |                                                                                                                     ^
/gpfs/users/tpadioleau/opt/include/experimental/__p0009_bits/extents.hpp:182:11: note: candidate: 'constexpr std::experimental::extents<Extents>::extents(std::experimental::extents<Extents>&&) [with long unsigned int ...Extents = {18446744073709551615}]'
  182 |   MDSPAN_INLINE_FUNCTION_DEFAULTED constexpr extents(extents&&) noexcept = default;
      |           ^~~~~~~
/gpfs/users/tpadioleau/opt/include/experimental/__p0009_bits/extents.hpp:182:19: note:   no known conversion for argument 1 from 'std::array<long unsigned int, 1>' to 'std::experimental::extents<18446744073709551615>&&'
  182 |   MDSPAN_INLINE_FUNCTION_DEFAULTED constexpr extents(extents&&) noexcept = default;
      |                   ^~~~~~~~~~
/gpfs/users/tpadioleau/opt/include/experimental/__p0009_bits/extents.hpp:181:11: note: candidate: 'constexpr std::experimental::extents<Extents>::extents(const std::experimental::extents<Extents>&) [with long unsigned int ...Extents = {18446744073709551615}]'
  181 |   MDSPAN_INLINE_FUNCTION_DEFAULTED constexpr extents(extents const&) noexcept = default;
      |           ^~~~~~~
/gpfs/users/tpadioleau/opt/include/experimental/__p0009_bits/extents.hpp:181:19: note:   no known conversion for argument 1 from 'std::array<long unsigned int, 1>' to 'const std::experimental::extents<18446744073709551615>&'
  181 |   MDSPAN_INLINE_FUNCTION_DEFAULTED constexpr extents(extents const&) noexcept = default;
      |                   ^~~~~~~~~~~~~~~
/gpfs/users/tpadioleau/opt/include/experimental/__p0009_bits/extents.hpp:180:11: note: candidate: 'std::experimental::extents<Extents>::extents() [with long unsigned int ...Extents = {18446744073709551615}]'
  180 |   MDSPAN_INLINE_FUNCTION_DEFAULTED constexpr extents() noexcept = default;
      |           ^~~~~~~
/gpfs/users/tpadioleau/opt/include/experimental/__p0009_bits/extents.hpp:180:11: note:   candidate expects 0 arguments, 1 provided
make[2]: *** [CMakeFiles/bug-report.dir/main.cpp.o] Erreur 1
make[2] : on quitte le répertoire « /gpfs/users/tpadioleau/bug-report/build »
make[1]: *** [CMakeFiles/bug-report.dir/all] Erreur 2
make[1] : on quitte le répertoire « /gpfs/users/tpadioleau/bug-report/build »
make: *** [all] Erreur 2

Comments

  • The compiler fails at finding a suitable constructor. More specifically, for ext2 (of type std::experimental::extents<18446744073709551615>) the SFINAE fails at N == std::experimental::extents<18446744073709551615, 18446744073709551615>::rank(), for some reason the compiler seems to be mixing the two classes.

  • If I manually inline the calls to rank() and rank_dynamic() it compiles fine.

  • Fortunately the compiler succeeds to compile if only one of the two extents (ext1 or ext2) is defined.

  • It compiles fine using nvcc/11.4.48 and g++/10.3.0.

tpadioleau avatar Nov 23 '21 10:11 tpadioleau

I am actually working on making nvcc work in CUDA mode that should come forth in the next couple weeks. Your reported bug is something I fixed in that.

crtrott avatar Nov 30 '21 17:11 crtrott

Awesome I would be curious to know the origin of this behavior.

tpadioleau avatar Dec 01 '21 07:12 tpadioleau

This was reported to NVIDIA, its simply a bug in the compiler. My guess would be that it has to do with the host compiler and nvcc treating the attribute differently leading to a different class layout internally, and then some part in the compiler which expects consistency between the host side layout of the class and the device side layout of the class crashes.

crtrott avatar Dec 01 '21 16:12 crtrott

If I understand you correctly you seem to be making a link with the [[no_unique_address]] attribute bug ?

It seemed to me it was a different bug as I disabled the [[no_unique_address]] feature to make this bug appear.

tpadioleau avatar Dec 02 '21 09:12 tpadioleau

I think this is fixed.

crtrott avatar Nov 23 '22 00:11 crtrott