Skip to content

Commit

Permalink
Fixes #82, #85:
Browse files Browse the repository at this point in the history
* Now depending on strf version 12.2, see https://github.com/robhz786/strf/releases :
	* Some naming changes
	* Now making a `printer input` and then feeding it to a `char_printer` to actually print a value
* Made lambda-based manipulators return a reference, not a copy
* Disabled copying of `printing_ostream`s
* Fixed missing initializations in the move constructor of `printfing_ostream`
* Naming tweak in the unit tests
  • Loading branch information
Eyal Rozenberg committed Nov 18, 2020
1 parent f65e295 commit 12207a1
Show file tree
Hide file tree
Showing 4 changed files with 58 additions and 38 deletions.
2 changes: 1 addition & 1 deletion CMakeLists.txt
Original file line number Diff line number Diff line change
Expand Up @@ -34,7 +34,7 @@ project(cuda-kat
# because most of it is host-side only; and it doesn't make sense to bundle a modified
# half of the standard library just for that. Instead, we use the strf library
# (available at: https://github.com/robhz786/strf )
find_package(strf 0.10.4)
find_package(strf 0.12.2)

###############
## OPTIONS ##
Expand Down
28 changes: 21 additions & 7 deletions src/kat/on_device/streams/printfing_ostream.cuh
Original file line number Diff line number Diff line change
Expand Up @@ -37,8 +37,18 @@ public:
enum class resolution { thread, warp, block, grid };

KAT_DEV printfing_ostream(std::size_t initial_buffer_size = cout_initial_buffer_size) : main_buffer(initial_buffer_size) { }
KAT_DEV printfing_ostream(printfing_ostream&& other) : main_buffer(other.main_buffer) { }
KAT_DEV printfing_ostream(const printfing_ostream& other) : main_buffer(other.main_buffer) { }

KAT_DEV printfing_ostream(printfing_ostream&& other) :
main_buffer(other.main_buffer),
prefix(other.prefix),
flush_on_destruction(other.flush_on_destruction),
newline_on_flush(other.newline_on_flush),
use_prefix(other.use_prefix),
prefix_generator(other.prefix_generator),
printing_resolution_(other.printing_resolution_)
{ }

KAT_DEV printfing_ostream(const printfing_ostream& other) = delete;
KAT_DEV ~printfing_ostream();

// Note: You can also use strf::flush if that exists
Expand Down Expand Up @@ -86,14 +96,16 @@ public:
KAT_DEV printfing_ostream& operator<<(const T& arg)
{
if (not should_act_for_resolution(printing_resolution_)) { return *this; }
strf::print_preview<false, false> no_preview;
strf::make_printer<char>(strf::rank<5>(), strf::pack(), no_preview, arg).print_to(main_buffer);
strf::print_preview<strf::preview_size::no, strf::preview_width::no> no_preview;
// strf::make_printer<char>(strf::rank<5>(), strf::pack(), no_preview, arg).print_to(main_buffer);
auto printer_input = strf::make_printer_input<char, decltype(no_preview), decltype(strf::pack())>(no_preview, strf::pack(), arg);
strf::printer_type<char,decltype(no_preview), decltype(strf::pack()), T>(printer_input).print_to(main_buffer);
return *this;
}

// Manipulators are a clever, but confusing, idea from the C++ standard library's
// IO streams: They're functions which manipulate streams, but can also be made
// to manipulatethem by being sent to them using the << operator - which instead
// to manipulate them by being sent to them using the << operator - which instead
// of actually adding any data to the stream, invokes the manipulator function.
//
using manipulator = kat::printfing_ostream& ( kat::printfing_ostream& );
Expand Down Expand Up @@ -196,7 +208,7 @@ KAT_DEV printfing_ostream& printfing_ostream::operator<< <printfing_ostream::man

namespace manipulators {
KAT_DEV auto prefix(prefix_generator_type gen) {
return [gen](kat::printfing_ostream& os) { return os.set_prefix_generator(gen); };
return [gen](kat::printfing_ostream& os) -> kat::printfing_ostream& { return os.set_prefix_generator(gen); };
}
} // namespace manipulators

Expand All @@ -219,7 +231,9 @@ KAT_DEV printfing_ostream& operator<< (printfing_ostream& os, manipulators::pref

namespace manipulators {
KAT_DEV auto resolution(printfing_ostream::resolution new_resolution) {
return [new_resolution](kat::printfing_ostream& os) { return os.set_printing_resolution(new_resolution); };
return [new_resolution](kat::printfing_ostream& os) -> kat::printfing_ostream& {
return os.set_printing_resolution(new_resolution);
};
}
} // namespace manipulators

Expand Down
56 changes: 31 additions & 25 deletions src/kat/on_device/streams/stringstream.cuh
Original file line number Diff line number Diff line change
Expand Up @@ -51,28 +51,31 @@ KAT_DEV T* safe_malloc(std::size_t size)
* @note This class owns its buffer.
* @note nothing is dynamically allocated if the length is 0
*/
class stringstream: public ::strf::basic_outbuf<char>
class stringstream: public ::strf::basic_outbuff<char>
{
public:
using char_type = char;
// no traits_type - this is not implemented by strf
// using int_type = int; // really
// using off_type = std::off_t; // really?

using pos_type = std::size_t;
char_type* pointer_type;


protected:
// Note: initial_buffer_size + 1 bytes must be allocated
STRF_HD stringstream(char_type* initial_buffer, std::size_t initial_buffer_size) :
buffer_size(initial_buffer_size),
buffer(initial_buffer),
strf::basic_outbuf<char_type>(initial_buffer, buffer_size)
strf::basic_outbuff<char_type>(initial_buffer, buffer_size)
{
}

public:
STRF_HD stringstream(std::size_t initial_buffer_size);

STRF_HD stringstream(stringstream&& other) : strf::basic_outbuf<char_type>(other.buffer, other.buffer_size)
STRF_HD stringstream(stringstream&& other) : strf::basic_outbuff<char_type>(other.buffer, other.buffer_size)
{
if (buffer != nullptr) {
free(buffer);
Expand All @@ -99,13 +102,13 @@ public:

KAT_DEV void clear()
{
set_pos(buffer);
set_pointer(buffer);
flush();
}

KAT_DEV void flush() {
if (buffer != nullptr) {
*pos() = '\0';
*pointer() = '\0';
}
}

Expand All @@ -118,10 +121,10 @@ public:
return buffer;
}

KAT_DEV pos_type tellp() const { return pos() - buffer; }
KAT_DEV pos_type tellp() const { return pointer() - buffer; }
KAT_DEV bool empty() const { return tellp() == 0; }
// std::stringstream's don't have this
KAT_DEV stringstream& seekp(pos_type pos) { set_pos(buffer + pos); return *this; }
KAT_DEV stringstream& seekp(pos_type pos) { set_pointer(buffer + pos); return *this; }

KAT_DEV std::size_t capacity() const { return buffer_size; } // perhaps there's something else we can use instead?

Expand Down Expand Up @@ -181,7 +184,7 @@ STRF_HD stringstream::stringstream(std::size_t initial_buffer_size)

KAT_DEV void stringstream::recycle()
{
std::size_t used_size = (buffer_size == 0) ? 0 : (this->pos() - buffer);
std::size_t used_size = (buffer_size == 0) ? 0 : (this->pointer() - buffer);
// a postcondition of recycle() is that at least so much free space is available.
auto new_buffer_size = builtins::maximum(
buffer_size * 2,
Expand All @@ -191,7 +194,7 @@ KAT_DEV void stringstream::recycle()
memcpy(new_buff, buffer, sizeof(char_type) * used_size);
free(buffer);
}
this->set_pos(new_buff + used_size);
this->set_pointer(new_buff + used_size);
this->set_end(new_buff + new_buffer_size);
buffer = new_buff;
}
Expand All @@ -211,26 +214,29 @@ KAT_DEV stringstream& operator<<(stringstream& out, const T& arg)
// TODO:
// 1. Can `no_preview` be made constant?
// 2. Can't we target a specific overload rather than play with ranks?
auto no_preview = ::strf::print_preview<false, false>{};
::strf::make_printer<char>(
::strf::rank<5>(),
// strf::rank is a method for controlling matching within the overload set:
// rank objects have no members, it's only about their type. Higher rank objects can
// match lower-rank objects (i.e. match functions in the overload sets expecting lower-rank
// objects), which means they have access to more of the overload sets. If we create
// a lower-rank object here we will only be able to match a few overload set members.
::strf::pack(),
// not modifying any facets such as digit grouping or digit separator
no_preview,
// Don't know what this means actually
arg
).print_to(out);
auto no_preview = ::strf::print_preview<strf::preview_size::no, strf::preview_width::no>{};
// ::strf::make_printer<char>(
// ::strf::rank<5>(),
// // strf::rank is a method for controlling matching within the overload set:
// // rank objects have no members, it's only about their type. Higher rank objects can
// // match lower-rank objects (i.e. match functions in the overload sets expecting lower-rank
// // objects), which means they have access to more of the overload sets. If we create
// // a lower-rank object here we will only be able to match a few overload set members.
// ::strf::pack(),
// // not modifying any facets such as digit grouping or digit separator
// no_preview,
// // Don't know what this means actually
// arg
// ).print_to(out);
auto printer_input = strf::make_printer_input<char, decltype(no_preview), decltype(strf::pack())>(no_preview, strf::pack(), arg);
strf::printer_type<char,decltype(no_preview), decltype(strf::pack()), T>(printer_input).print_to(out);


// Note: This function doesn't actually rely on out being a stringstream; any
// ostream-like class would do. But for now, we don't have any ostreams other
// than the stringstream, so we'll leave it this way. Later, with could either
// have an intermediate class, or wrap basic_outbuf with an ostream class
// without a buffer, or just call basic_outbuf an ostream
// have an intermediate class, or wrap basic_outbuff with an ostream class
// without a buffer, or just call basic_outbuff an ostream

return out;
}
Expand Down
10 changes: 5 additions & 5 deletions tests/printing.cu
Original file line number Diff line number Diff line change
Expand Up @@ -216,7 +216,7 @@ __global__ void print_at_different_resolutions()
cout << "Printing at thread resolution. The printing thread is (" << blockIdx.x << "," << threadIdx.x << ")\n" << flush;
}

__device__ void sipo_for_resolution(kat::printfing_ostream& os, kat::printfing_ostream::resolution res)
__device__ void self_identifying_printfing_ostream_for_resolution(kat::printfing_ostream& os, kat::printfing_ostream::resolution res)
{
os
<< kat::manipulators::resolution(res)
Expand All @@ -232,15 +232,15 @@ __global__ void self_identifying_printfing_ostream()
using kat::flush;
namespace gi = kat::linear_grid::grid_info;

sipo_for_resolution(cout, kat::printfing_ostream::resolution::grid);
self_identifying_printfing_ostream_for_resolution(cout, kat::printfing_ostream::resolution::grid);

kat::sleep<kat::sleep_resolution::clock_cycles>(1e8);

sipo_for_resolution(cout, kat::printfing_ostream::resolution::block);
self_identifying_printfing_ostream_for_resolution(cout, kat::printfing_ostream::resolution::block);
__syncthreads();
sipo_for_resolution(cout, kat::printfing_ostream::resolution::warp);
self_identifying_printfing_ostream_for_resolution(cout, kat::printfing_ostream::resolution::warp);
__syncthreads();
sipo_for_resolution(cout, kat::printfing_ostream::resolution::thread);
self_identifying_printfing_ostream_for_resolution(cout, kat::printfing_ostream::resolution::thread);
__syncthreads();
}

Expand Down

0 comments on commit 12207a1

Please sign in to comment.