Correct numerous 20054-D: dynamic initialization errors found on arm+12.2

[robertmaynard]

Description

Compile issues found by compiling libcudf with the rapidsai/devcontainers:23.10-cpp-gcc9-cuda12.2-ubuntu20.04 docker container.

Checklist

• I am familiar with the Contributing Guidelines.
• New or existing tests cover these changes.
• The documentation is up to date with these changes.

[bdice]

@robertmaynard Do you think these are "real issues" that we should fix for C++ standards compliance (like avoiding UB) or are they potentially compiler bugs? Do you see these issues with newer GCC versions or older CUDA versions on ARM?

[robertmaynard]

@robertmaynard Do you think these are "real issues" that we should fix for C++ standards compliance (like avoiding UB) or are they potentially compiler bugs? Do you see these issues with newer GCC versions or older CUDA versions on ARM?

These are all real issues that aren't exposed by the x86 compiler. I am still digging into why, but here is the background on why these failures are valid:

From the CUDA programming guide 14.5.3.1. Device Memory Space Specifiers:

A __device__, __constant__, __managed__ or __shared__ variable definition cannot have a class type with a non-empty constructor or a non-empty destructor. A constructor for a class type is considered empty at a point in the translation unit, if it is either a trivial constructor or it satisfies all of the following conditions:

• The constructor function has been defined.
• The constructor function has no parameters, the initializer list is empty and the function body is an empty compound statement.
• Its class has no virtual functions, no virtual base classes and no non-static data member initializers.
• The default constructors of all base classes of its class can be considered empty.
• For all the nonstatic data members of its class that are of class type (or array thereof), the default constructors can be considered empty.

The simpliest way to understand the above is that the __shared__ type must pass the std::is_trivially_constructible check, if it doesn't it fails to meet the requirements and could generate the warning. As an example if we look at the block_desc_s specification we can see that it fails to meet the is_trivially_constructible requirements ( https://godbolt.org/z/75rv9jd3s )

struct block_desc_s {
  block_desc_s() {};

  explicit constexpr block_desc_s(size_t offset_, uint32_t size_,
                                  uint32_t row_offset_, uint32_t first_row_,
                                  uint32_t num_rows_)
      : offset(offset_), size(size_), row_offset(row_offset_),
        first_row(first_row_), num_rows(num_rows_) {}

  size_t offset;
  uint32_t size;
  uint32_t row_offset;
  uint32_t first_row;
  uint32_t num_rows;
};

[karthikeyann]

does default here not meet 14.5.3.1. Device Memory Space Specifiers specifications?
is default not considered as trivial constructor if all non-static members have trivial constructors?

I see usage of default in this same PR.

[robertmaynard]

I need the constexpr constructor for the byte_stream_s type in page_hdr.cu. I agree that on initial read the = default constructor should be sufficient, but the compiler disagrees.

[robertmaynard]

/merge