fix printf when reaching the end of the printf buffer (#585)

stop printing if the args size overflow the printf buffer size
add a new test

src/printf.cpp

@@ -156,7 +156,8 @@ std::string print_part(const std::string& fmt, const char* data, size_t size) {
     return std::string(out.data());
 }
 
-void process_printf(char*& data, const printf_descriptor_map_t& descs) {
+void process_printf(char*& data, const printf_descriptor_map_t& descs,
+                    char* data_end) {
 
     uint32_t printf_id = read_inc_buff<uint32_t>(data);
     auto& format_string = descs.at(printf_id).format_string;
@@ -192,6 +193,11 @@ void process_printf(char*& data, const printf_descriptor_map_t& descs) {
         // The size of the argument that this format part will consume
         auto& size = descs.at(printf_id).arg_sizes[arg_idx];
 
+        if (data + size > data_end) {
+            data += size;
+            return;
+        }
+
         // Check to see if we have a vector format specifier
         int vec_len = 0;
         int el_size = 0;
@@ -239,13 +245,13 @@ cl_int cvk_printf(cvk_mem* printf_buffer,
     char* data = static_cast<char*>(printf_buffer->host_va());
     auto buffer_size = printf_buffer->size();
     const auto bytes_written_size = sizeof(uint32_t);
-    const auto data_size = buffer_size - bytes_written_size;
-    auto bytes_written = read_inc_buff<uint32_t>(data) * 4;
-    auto* data_start = data;
+    const size_t data_size = buffer_size - bytes_written_size;
+    const size_t bytes_written = read_inc_buff<uint32_t>(data) * 4;
+    const size_t limit = std::min(bytes_written, data_size);
+    auto* data_end = data + limit;
 
-    while (static_cast<size_t>(data - data_start) < bytes_written &&
-           static_cast<size_t>(data - data_start) < data_size) {
-        process_printf(data, descriptors);
+    while (data < data_end) {
+        process_printf(data, descriptors, data_end);
     }
 
     printf_buffer->unmap();

tests/api/printf.cpp

@@ -129,7 +129,9 @@ TEST_F(WithCommandQueue, SimplePrintf) {
 }
 
 TEST_F(WithCommandQueue, TooLongPrintf) {
-    clvk_override_printf_buffer_size(24);
+    // each print takes 12 bytes (4 for the printf_id, and 2*4 for the 2 integer
+    // to print) + 4 for the byte written counter
+    clvk_override_printf_buffer_size(28);
 
     temp_folder_deletion temp;
     stdoutFileName = getStdoutFileName(temp);
@@ -140,7 +142,43 @@ TEST_F(WithCommandQueue, TooLongPrintf) {
     const char* source = R"(
     kernel void test_printf() {
       for (unsigned i = 0; i < 3; i++){
-        printf("get_global_id(%u) = %u\n", i, get_global_id(i));
+        printf("get_global_id(%u) = %u\n", i, (unsigned)get_global_id(i));
+      }
+    }
+    )";
+    auto kernel = CreateKernel(source, "test_printf");
+
+    size_t gws = 1;
+    size_t lws = 1;
+    EnqueueNDRangeKernel(kernel, 1, nullptr, &gws, &lws, 0, nullptr, nullptr);
+    Finish();
+
+    releaseStdout(fd);
+    auto printf_buffer = getStdoutContent();
+    ASSERT_NE(printf_buffer, nullptr);
+
+    // We only get the first 2 prints because the buffer is too small to get the
+    // last one.
+    const char* message = "get_global_id(0) = 0\nget_global_id(1) = 0\n";
+    ASSERT_STREQ(printf_buffer, message);
+}
+
+TEST_F(WithCommandQueue, TooLongPrintf2) {
+    // each print takes 12 bytes (4 for the printf_id, and 2*4 for the 2 integer
+    // to print) + 4 for the byte written counter + 8 which are not enough for
+    // the third print, but should not cause any issue in clvk
+    clvk_override_printf_buffer_size(36);
+
+    temp_folder_deletion temp;
+    stdoutFileName = getStdoutFileName(temp);
+
+    int fd;
+    ASSERT_TRUE(getStdout(fd));
+
+    const char* source = R"(
+    kernel void test_printf() {
+      for (unsigned i = 0; i < 3; i++){
+        printf("get_global_id(%u) = %u\n", i, (unsigned)get_global_id(i));
       }
     }
     )";