Implement `std::format` and `std::print`
- Dominant language
- C++
- Stars
- 2.5k
- Forks
- 486
- Avg merge
- 2d 6h
- Merged PRs (30d)
- 295
Description
printf() is available in CUDA, but it has several deficiencies.
- The data type specifiers in the format string must agree with the types of the variadic arguments (a DRY failure).
- printf() does not support overloads to support user-defined types.
- printing across multiple statements has undefined interleaving with other threads, limiting composability
This request is to add a similar printing facility as `std::ostream` that is accessible in device code. This should enable user-defined printing functions and provide type safety. Ideally, delimiter tokens analogous to `std::flush` and `std::endl` would enable the CUDA driver to interleave the output from CUDA threads without corruption.
Example usage:
```c++
struct Foo {
int member;
};
struct Bar {
Foo foo;
char const *name;
};
// Prints an object of type Foo.
__device__ cu::cout &operator<<(cu::cout &out, Foo const &foo) {
return out << foo.member;
}
// Prints an object of type Bar.
__device__ cu::cout &operator<<(cu::cout &out, Bar const &bar) {
return out << bar.foo << " " << name;
}
// Sample usage
__global__ void kernel() {
Bar bar;
cu::cout << "Thread " << threadIdx.x << ": " << bar << cu::endl;
__syncthreads();
cu::cout << "This is a multistatement ";
cu::cout << "output block. ";
cu::cout << "The index of the current thread is " << threadIdx.x;
if (threadIdx.x & 1) {
cu::cout << ", and it contains a bar of value " << bar;
}
cu::cout << ". The output appears contiguous despite control flow and printing over multiple statements.\n";
cu::cout << cu::flush;
}
// Generic template kernel printing all objects. Requires `operator<<(cu::ostream &, T const &rhs)` to exist in scope
template
__global__ void generic_print(T *objects) {
cu::cout << "T" << threadIdx.x << ": " << objects[threadIdx.x] << cu::endl;
}
```
Contributor guide
Assessment
This issue has not been assessed yet.