[Bugfix] Supply missing `T.print` for bool type (#1279)

* fix for bool dtype

* lint fix

* fix

* ci fix

tvm @ f4affc7f

-Subproject commit f4105f89a646622acc9818584d1d91e2ca3f533d
+Subproject commit f4affc7f31e36e7f88c0fe1c715b03215c6a0c62

src/tl_templates/cuda/debug.h

@@ -29,6 +29,14 @@ __device__ void debug_print_var<signed char>(const char *msg, signed char var) {
          threadIdx.z, var);
 }
 
+// Specialization for plain char type
+template <> __device__ void debug_print_var<char>(const char *msg, char var) {
+  printf("msg='%s' BlockIdx=(%d, %d, %d), ThreadIdx=(%d, %d, %d): dtype=char "
+         "value=%d\n",
+         msg, blockIdx.x, blockIdx.y, blockIdx.z, threadIdx.x, threadIdx.y,
+         threadIdx.z, (int)var);
+}
+
 // Specialization for unsigned char type
 template <>
 __device__ void debug_print_var<unsigned char>(const char *msg,
@@ -58,6 +66,14 @@ __device__ void debug_print_var<unsigned int>(const char *msg,
          threadIdx.z, var);
 }
 
+// Specialization for bool type
+template <> __device__ void debug_print_var<bool>(const char *msg, bool var) {
+  printf("msg='%s' BlockIdx=(%d, %d, %d), ThreadIdx=(%d, %d, %d): dtype=bool "
+         "value=%s\n",
+         msg, blockIdx.x, blockIdx.y, blockIdx.z, threadIdx.x, threadIdx.y,
+         threadIdx.z, var ? "true" : "false");
+}
+
 // Specialization for float type
 template <> __device__ void debug_print_var<float>(const char *msg, float var) {
   printf("msg='%s' BlockIdx=(%d, %d, %d), ThreadIdx=(%d, %d, %d): dtype=float "