6 Coro gpu_work() { 7 cudaStream_t stream; 8 cudaStreamCreate(stream); 9 gpu_matmul<<<8, 256, 0, stream>>>(matA, matB); 10 while(cudaStreamQuery(stream) != cudaSuccess) { 11 co_await std::suspend_always{}; std::suspend_always{}; 12 } 13 cudaStreamDestory(stream); 14 } 2 struct Coro { 3 4 struct promise_type { 5 std::suspend_always initial_suspend() noexcept { return {}; } 6 std::suspend_always final_suspend() Polling D C 7 auto task_a = taro.emplace([&]() { 8 cuda.wait([&](cudaStream_t stream) { 9 kernel_a1<<<32, 256, 0, stream>>>(); 10 }); // synchronize 11 }); 31Taro’s Programming Model – Example 1 Taro:

取了，从而浪费了 2 倍带宽。 绕过缓存，直接写入： _mm_stream_si32 • 因此需要把 16 次 float 用 SIMD 指令合并成一次 写入，且写入的地址要对齐到 64 字节，才能避免 浪费读取的带宽。这样的条件实在有点苛刻，毕 竟小彭老师的电脑还不支持 AVX512 。 • 可以用 _mm_stream_si32 指令代替直接赋值的 写入，他能够绕开缓存，将一个 4 bitcast 一下参数。 stream 的特点：不会读到缓存里 • 因为 _mm_stream_si32 会绕开缓存，直 接把数据写到内存，之后读取的话，反而 需要等待 stream 写回执行完成，然后重 新读取到缓存，反而更低效。 • 因此，仅当这些情况： 1. 该数组只有写入，之前完全没有读取过 。 2. 之后没有再读取该数组的地方。 • 才应该用 stream 指令。 4 倍矢量化的版本： 倍矢量化的版本： _mm_stream_ps • _mm_stream_si32 可以一次性写入 4 字 节到挂起队列。而 _mm_stream_ps 可以 一次性写入 16 字节到挂起队列，更加高 效了。 • 他的第二参数是一个 __m128 类型，可以 配合其他手写的 SIMD 指令使用。 • 不过， _mm_stream_ps 写入的地址必须 对齐到 16 字节，否则会产生段错误等异

typename DynamicBufferV2, typename CompletionToken> decltype(auto) async_read(AsyncReadStream& stream, DynamicBufferV2 buffer, CompletionToken&& token); 9 Dynamic Buffer Requirements struct /* . typename ConstBufferSequence, typename CompletionToken> decltype(auto) async_write(AsyncWriteStream& stream, message_type type, async_write_state& state, ConstBufferSequence payload, CompletionToken&& WaitableTimer, typename Events, typename CompletionToken> decltype(auto) async_heartbeat(AsyncWriteStream& stream, async_write_state& state, WaitableTimer& timer, Events events, CompletionToken&& token); 16

_srvs::srv::GetPath::Response> response) { if (map_.get_data().size() == 0) { RCLCPP_ERROR_STREAM(this->get_logger(), "MAP IS EMPTY!!"); response->code.code = example_srvs::msg::GetPathCodes _srvs::srv::GetPath::Response> response) { if (map_.get_data().size() == 0) { RCLCPP_ERROR_STREAM(this->get_logger(), "MAP IS EMPTY!!"); response->code.code = example_srvs::msg::GetPathCodes _srvs::srv::GetPath::Response> response) { if (map_.get_data().size() == 0) { RCLCPP_ERROR_STREAM(this->get_logger(), "MAP IS EMPTY!!"); response->code.code = example_srvs::msg::GetPathCodes

DISK CAPACITY int fscanf(FILE* stream, char const* format, …); int fprintf(FILE* stream, char const* format, …); int fgetc(FILE* stream); int fputc(int ch, FILE* stream); size_t fread(void* buffer, size_t size, size_t count, FILE* stream); size_t fwrite(void* buffer, size_t size, size_t count, FILE* stream);MORE RAM, MORE DISK CAPACITY • std::filesystem • Directories, links, block files, character mode); • BUFSIZ • int setvbuf(FILE* stream, char* buffer, int mode, size_t size); • size_t fread(void* buffer, size_t size, size_t count, std::FILE* stream);MOVING DATA INTO AND OUT OF RAM • fstream

enum Key to Compressed String Hu�man Coding for compression Output struct: Mapping of Key -> bit stream location/length Compressed bitstream Hu�man tree encoded into array Lookup -> iterator over compressed char -> bitstream cache Build encoded bit stream Build index -> bit position / string length lookup Package Hu�man tree to array Combine Hu�man tree, bit stream and index lookup into result 10 . 7/ Constrain tree = BuildHuffmanTree(makeStringsLambda); ResultType stream; // fill output stream... return stream } 10 . 11/ constexpr auto charLookup = MakeCharacterLookupTable();

rpc ListFeatures(Rectangle) returns (stream Feature) {} // Client streaming RPC. rpc RecordRoute(stream Point) returns (RouteSummary) {} // Bidirectional streaming RPC. rpc RouteChat(stream RouteNote) returns (stream RouteNote) {} }service RouteGuide { // Unary RPC. rpc ListFeatures(Rectangle) returns (stream Feature) {} // Client streaming RPC. rpc RecordRoute(stream Point) returns (RouteSummary) {} //

>>= pipes::transform(numberDashColor) >>= pipes::intersperse("\n") >>= pipes::to_out_stream(std::cout); 46 auto const numbers = std::vector{ 1, 2, 3 }; auto const colors = std::vector{ >>= pipes::transform(numberDashColor) >>= pipes::intersperse("\n") >>= pipes::to_out_stream(std::cout); 47 auto const numbers = std::vector{ 1, 2, 3 }; auto const colors = std::vector{ >>= pipes::transform(numberDashNumber) >>= pipes::intersperse("\n") >>= pipes::to_out_stream(std::cout);49 auto const inputs = std::vector{ "one", "two", "three", "four" }; p

Back to Basics: Classic STL Copyright © 2021 Bob Steagall Reading Elements (Bytes) From a FILE Stream 31 Action Operation Read element and advance b = fgetc(p) Compare for end-of-file equality b == EOF, feof(p) Make a copy (assign) q = p O(1) - constant time • Consider a pointer to a FILE stream opened for input • What can you do with it? bn-1 bn-2 bn-4 bn-5 bn-6 … … … … … bn-3 bn FILE* p Classic STL Copyright © 2021 Bob Steagall Writing Elements (Bytes) To a FILE Stream • Consider a pointer to a FILE stream opened for output 32 • What can you do with it? Action Operation Write element

level either a stream or packetsStream Depending on your low level either a stream or packetsStream Message 4 Message 3 Message 2 Message 1 Depending on your low level either a stream or packetsPotential