diff --git a/python/taichi/lang/kernel.py b/python/taichi/lang/kernel.py
index 604da527e91d1c..b99af5a136d730 100644
--- a/python/taichi/lang/kernel.py
+++ b/python/taichi/lang/kernel.py
@@ -458,7 +458,9 @@ def __call__(self, *args, **kwargs):
         instance_id, arg_features = self.mapper.lookup(args)
         key = (self.func, instance_id)
         self.materialize(key=key, args=args, arg_features=arg_features)
-        return self.compiled_functions[key](*args)
+        ret = self.compiled_functions[key](*args)
+        print(taichi_lang_core.pop_python_print_buffer(), end='')
+        return ret
 
 
 # For a Taichi class definition like below:
diff --git a/taichi/backends/opengl/opengl_api.cpp b/taichi/backends/opengl/opengl_api.cpp
index 23428428255f8c..23b2e4bf66edd0 100644
--- a/taichi/backends/opengl/opengl_api.cpp
+++ b/taichi/backends/opengl/opengl_api.cpp
@@ -5,6 +5,7 @@
 #include "taichi/program/kernel.h"
 #include "taichi/program/program.h"
 #include "taichi/util/environ_config.h"
+#include "taichi/python/print_buffer.h"
 
 #ifdef TI_WITH_OPENGL
 #include "glad/glad.h"
@@ -480,9 +481,9 @@ struct CompiledProgram::Impl {
             TI_WARN("[glsl] Unexpected serialization type: {}, ignoring", type);
             break;
         };
-        std::cout << str;
+        py_cout << str;
       }
-      std::cout << std::endl;
+      py_cout << '\n';
     }
     rt_buf->msg_count = 0;
     launcher->impl->runtime_ssbo->unmap();
diff --git a/taichi/python/export_misc.cpp b/taichi/python/export_misc.cpp
index 2c3c250692b6b6..e48a3b8fb9d7c3 100644
--- a/taichi/python/export_misc.cpp
+++ b/taichi/python/export_misc.cpp
@@ -7,6 +7,7 @@
 #include "taichi/common/task.h"
 #include "taichi/math/math.h"
 #include "taichi/python/exception.h"
+#include "taichi/python/print_buffer.h"
 #include "taichi/python/export.h"
 #include "taichi/system/benchmark.h"
 #include "taichi/system/profiler.h"
@@ -20,6 +21,8 @@
 
 TI_NAMESPACE_BEGIN
 
+PythonPrintBuffer py_cout;
+
 Config config_from_py_dict(py::dict &c) {
   Config config;
   for (auto item : c) {
@@ -153,6 +156,9 @@ void export_misc(py::module &m) {
     }
     printf("test was successful.\n");
   });
+  m.def("pop_python_print_buffer", []() {
+      return py_cout.pop_content();
+  });
   m.def("with_cuda", is_cuda_api_available);
   m.def("with_metal", taichi::lang::metal::is_metal_api_available);
   m.def("with_opengl", taichi::lang::opengl::is_opengl_api_available);
diff --git a/taichi/python/print_buffer.h b/taichi/python/print_buffer.h
new file mode 100644
index 00000000000000..857990cc04e505
--- /dev/null
+++ b/taichi/python/print_buffer.h
@@ -0,0 +1,25 @@
+#pragma once
+
+#include <sstream>
+
+TI_NAMESPACE_BEGIN
+
+struct PythonPrintBuffer {
+  /* holds kernel print result before switching back to python */
+  std::stringstream ss;
+
+  template <typename T>
+  PythonPrintBuffer &operator<<(const T &t) {
+    ss << t;
+    return *this;
+  }
+  std::string pop_content() {
+    auto ret = ss.str();
+    ss = std::stringstream();
+    return ret;
+  }
+};
+
+extern PythonPrintBuffer py_cout;
+
+TI_NAMESPACE_END