NVIDIA / NVIDIA/cccl

Implement `std::format` and `std::print`

Open
#939 10 comments 8 reactions 1 assignee Claimed by @davebayer View on GitHub
2024-2025 goal libcu++
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

Open the contributing guide

Assessment

This issue has not been assessed yet.

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.