Enable all relevant <cuda/...> headers to work with C++ Tile

[jrhemstad]

This issue can be closed by:

• First categorizing what headers are relevant for use in C++ Tile Kernels
• For all headers determined to be relevant, all tests are passing when run in tile mode

Standard Headers

• algorithm
• array
• atomic
• barrier
• bit
• bitset
• cassert
• ccomplex
• cfloat
• charconv
• chrono
• climits
• cmath
• complex
• concepts
• cstddef
• cstdint
• cstdlib
• cstring
• ctime
• execution
• expected
• functional
• initializer_list
• inplace_vector
• iterator
• latch
• limits
• linalg
• mdspan
• memory
• numbers
• numeric
• optional
• random
• ranges
• ratio
• semaphore
• source_location
• span
• string_view
• tuple
• type_traits
• utility
• variant
• version

Extension Headers <cuda/...>

• access_property
• algorithm
• anotated_ptr
• atomic
• barrier
• bit
• buffer
• cmath
• devices
• discard_memory
• functional
• hierarchy
• iterator
• latch
• launch
• mdspan
• memory
• memory_pool
• memory_resource
• numeric
• pipeline
• ptx
• random
• semaphore
• stream
• stream_ref
• tma
• type_traits
• utility
• version
• warp
• work_stealing

Unsupported Headers

(Any headers that we deem irrelevant or unsupported for Tile kernels should go here)

[miscco]

I have extensively tested the libcu++ test suite with the ToT compiler and tile mode.

The results for C++17 are as follows:

Total Discovered Tests: 3325
  Unsupported      :  135 (4.06%)
  Passed           : 2505 (75.34%)
  Expectedly Failed:  685 (20.60%)

Header that can currently not be included are:

• <cuda/access_property>
• <cuda/annotated_ptr>
• <cuda/atomic>
• <cuda/barrier"
• <cuda/buffer>
• <cuda/cmath>
• <cuda/hierarchy>
• <cuda/latch>
• <cuda/memory>
• <cuda/memory_resource>
• <cuda/pipeline>
• <cuda/ptx>
• <cuda/random>
• <cuda/semaphore>
• <cuda/discard_memory>
• <cuda/std/atomic>
• <cuda/std/cmath>
• <cuda/std/complex>
• <cuda/std/barrier>
• <cuda/std/execution>
• <cuda/std/latch>
• <cuda/std/random>
• <cuda/std/semaphore>

The main reasons we cannot include headers are:

• Use of asm or ptx is not supported inside a tile kernel
• Use of threadIdx and other CUDA compiler intrinsics
• Taking addresses of functions. This is currently not supported in tile mode
• Use of bitfields. This is rare but also breaks tile mode
• Use of unsupported compiler builtins. A lot of the compiler builtins are currently not supported in tile mode

This concerns just whether a header can actually be included, not whether a feature can be used. Features that are currently not or only partially usable in tile mode are

• cuda::std::atomic:
  • We currently implement atomic through ptx, this will require a rework to use the new compiler intrinsics
• <cuda/std/bit>:
  • A lot of the bit manipulation functions rely on compiler builtins or assembly to performance. We can work around this but its work
• <cuda/std/cmath>:
  • We are heavily relying on compiler builtins for the math functions. Those are currently not supported in tile mode
• cuda::std::complex
  • complex requires the math functions to e.g implement multiplication
• cuda::std::mdspan
  • Our implementation uses __builtin_unreachable for implementation of variadic recursion. This breaks because of the unimplemented builtin
• cuda::std::variant
  • Our implementation uses __builtin_unreachable for implementation of variadic recursion. This breaks because of the unimplemented builtin
• CPOs
  • Because of implementation challenges regarding the host STL a lot of functions are implemented as CPOs (e.g cuda::std::begin)
    We cannot access a __device__ object inside a tile kernel. This i breaking a lot of functionality