Implement `std::cout`

[kerrmudgeon]

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:

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 <typename T>
__global__ void generic_print(T *objects) {
  cu::cout << "T" << threadIdx.x << ": " << objects[threadIdx.x] << cu::endl;
}

[donglinz]

May I expect any update on this? std::cout in device code would be a sweet feature to facilitate debugging.

[kerrmudgeon]

Is there any interest in a sweet feature like this?

Note, std::format or similar would be a perfectly viable alternative. Even a non-standard solution to enable seemingly atomic printing would be welcome.

Thanks for any consideration.

[maddyscientist]

Definitely interest from another user here for this sweet feature 😄

[wmaxey]

It's on our roadmap, we just need to prioritize it. ;) The more people that ask, the easier it would be for us to bump it into the next couple releases.

[jrhemstad]

I think std::format is going to be much more feasible than std::cout.

[wyphan]

@wmaxey I third this. Most C++ beginners prefer using std::cout and std::cerr to std::format IMHO. The latter is kinda Python-like (or is it the other way around?)