Skip to content

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

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
LeiWang1999 merged 4 commits into from
Nov 19, 2025
Merged

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

Show file tree
Hide file tree
Changes from 3 commits
Commits
Show all changes
4 commits
Select commit Hold shift + click to select a range
0100fa5
fix for bool dtype
LeiWang1999 Nov 19, 2025
a0cdb6a
lint fix
LeiWang1999 Nov 19, 2025
1e6a00a
fix
LeiWang1999 Nov 19, 2025
709f3ec
ci fix
LeiWang1999 Nov 19, 2025
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
2 changes: 1 addition & 1 deletion 3rdparty/tvm
Open in desktop
Submodule tvm updated from f4105f to f4affc
138 changes: 138 additions & 0 deletions src/tl_templates/cuda/debug.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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);
}
Comment on lines +32 to +38

@coderabbitai coderabbitai Bot Nov 19, 2025

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

⚠️ Potential issue | 🟠 Major

Approve char specialization, but buffer variant is missing.

The implementation is correct. Casting to (int) is appropriate for printing plain char since its signedness is implementation-defined.

However, notice that all other types in this file have both debug_print_var and debug_print_buffer_value specializations. The debug_print_buffer_value<char> specialization is missing (should be added around line 155-165 following the pattern of other types).

Add the missing specialization:

// Specialization for plain char type
template <>
__device__ void
debug_print_buffer_value<char>(const char *msg, const char *buf_name,
                               int index, char var) {
  printf("msg='%s' BlockIdx=(%d, %d, %d), ThreadIdx=(%d, %d, %d): buffer=%s, "
         "index=%d, dtype=char value=%d\n",
         msg, blockIdx.x, blockIdx.y, blockIdx.z, threadIdx.x, threadIdx.y,
         threadIdx.z, buf_name, index, (int)var);
}
🤖 Prompt for AI Agents 
In src/tl_templates/cuda/debug.h (existing char specialization at ~lines 32-38;
add buffer specialization around lines 155-165), add the missing __device__
template<> debug_print_buffer_value<char> implementation that mirrors other
types: accept (const char *msg, const char *buf_name, int index, char var), and
call printf with the same header (msg, BlockIdx, ThreadIdx) then print buffer
name, index and the char value cast to (int) for safe numeric output; follow the
exact formatting used by other buffer specializations to keep consistency.


// Specialization for unsigned char type
template <>
__device__ void debug_print_var<unsigned char>(const char *msg,
Expand Down Expand Up @@ -58,6 +66,92 @@ __device__ void debug_print_var<unsigned int>(const char *msg,
threadIdx.z, var);
}

// Specialization for short type
template <> __device__ void debug_print_var<short>(const char *msg, short var) {
printf("msg='%s' BlockIdx=(%d, %d, %d), ThreadIdx=(%d, %d, %d): dtype=short "
"value=%d\n",
msg, blockIdx.x, blockIdx.y, blockIdx.z, threadIdx.x, threadIdx.y,
threadIdx.z, (int)var);
}

// Specialization for unsigned short type
template <>
__device__ void debug_print_var<unsigned short>(const char *msg,
unsigned short var) {
printf(
"msg='%s' BlockIdx=(%d, %d, %d), ThreadIdx=(%d, %d, %d): dtype=unsigned "
"short value=%u\n",
msg, blockIdx.x, blockIdx.y, blockIdx.z, threadIdx.x, threadIdx.y,
threadIdx.z, (unsigned int)var);
}

// Specialization for long type
template <> __device__ void debug_print_var<long>(const char *msg, long var) {
printf("msg='%s' BlockIdx=(%d, %d, %d), ThreadIdx=(%d, %d, %d): dtype=long "
"value=%ld\n",
msg, blockIdx.x, blockIdx.y, blockIdx.z, threadIdx.x, threadIdx.y,
threadIdx.z, var);
}

// Specialization for unsigned long type
template <>
__device__ void debug_print_var<unsigned long>(const char *msg,
unsigned long var) {
printf(
"msg='%s' BlockIdx=(%d, %d, %d), ThreadIdx=(%d, %d, %d): dtype=unsigned "
"long value=%lu\n",
msg, blockIdx.x, blockIdx.y, blockIdx.z, threadIdx.x, threadIdx.y,
threadIdx.z, var);
}

// Specialization for long long type
template <>
__device__ void debug_print_var<long long>(const char *msg, long long var) {
printf("msg='%s' BlockIdx=(%d, %d, %d), ThreadIdx=(%d, %d, %d): dtype=long "
"long value=%lld\n",
msg, blockIdx.x, blockIdx.y, blockIdx.z, threadIdx.x, threadIdx.y,
threadIdx.z, var);
}

// Specialization for unsigned long long type
template <>
__device__ void debug_print_var<unsigned long long>(const char *msg,
unsigned long long var) {
printf(
"msg='%s' BlockIdx=(%d, %d, %d), ThreadIdx=(%d, %d, %d): dtype=unsigned "
"long long value=%llu\n",
msg, blockIdx.x, blockIdx.y, blockIdx.z, threadIdx.x, threadIdx.y,
threadIdx.z, var);
}

// Specialization for int64_t type
template <>
__device__ void debug_print_var<int64_t>(const char *msg, int64_t var) {
printf(
"msg='%s' BlockIdx=(%d, %d, %d), ThreadIdx=(%d, %d, %d): dtype=int64_t "
"value=%lld\n",
msg, blockIdx.x, blockIdx.y, blockIdx.z, threadIdx.x, threadIdx.y,
threadIdx.z, (long long)var);
}

// Specialization for uint64_t type
template <>
__device__ void debug_print_var<uint64_t>(const char *msg, uint64_t var) {
printf(
"msg='%s' BlockIdx=(%d, %d, %d), ThreadIdx=(%d, %d, %d): dtype=uint64_t "
"value=%llu\n",
msg, blockIdx.x, blockIdx.y, blockIdx.z, threadIdx.x, threadIdx.y,
threadIdx.z, (unsigned long long)var);
}
Comment thread
coderabbitai[bot] marked this conversation as resolved.
Outdated

// 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");
}
Comment thread
coderabbitai[bot] marked this conversation as resolved.

// 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 "
Expand Down Expand Up @@ -170,6 +264,17 @@ debug_print_buffer_value<unsigned int>(const char *msg, const char *buf_name,
threadIdx.z, buf_name, index, var);
}

// Specialization for bool type
template <>
__device__ void debug_print_buffer_value<bool>(const char *msg,
const char *buf_name, int index,
bool var) {
printf("msg='%s' BlockIdx=(%d, %d, %d), ThreadIdx=(%d, %d, %d): buffer=%s, "
"index=%d, dtype=bool value=%s\n",
msg, blockIdx.x, blockIdx.y, blockIdx.z, threadIdx.x, threadIdx.y,
threadIdx.z, buf_name, index, var ? "true" : "false");
}

// Specialization for float type
template <>
__device__ void debug_print_buffer_value<float>(const char *msg,
Expand Down Expand Up @@ -258,6 +363,39 @@ __device__ void debug_print_buffer_value<int16_t>(const char *msg,
threadIdx.z, buf_name, index, (int32_t)var);
}

// Specialization for uint16 type
template <>
__device__ void debug_print_buffer_value<uint16_t>(const char *msg,
const char *buf_name,
int index, uint16_t var) {
printf("msg='%s' BlockIdx=(%d, %d, %d), ThreadIdx=(%d, %d, %d): buffer=%s, "
"index=%d, dtype=uint16_t value=%u\n",
msg, blockIdx.x, blockIdx.y, blockIdx.z, threadIdx.x, threadIdx.y,
threadIdx.z, buf_name, index, (unsigned int)var);
}

// Specialization for int64 type
template <>
__device__ void debug_print_buffer_value<int64_t>(const char *msg,
const char *buf_name,
int index, int64_t var) {
printf("msg='%s' BlockIdx=(%d, %d, %d), ThreadIdx=(%d, %d, %d): buffer=%s, "
"index=%d, dtype=int64_t value=%lld\n",
msg, blockIdx.x, blockIdx.y, blockIdx.z, threadIdx.x, threadIdx.y,
threadIdx.z, buf_name, index, (long long)var);
}

// Specialization for uint64 type
template <>
__device__ void debug_print_buffer_value<uint64_t>(const char *msg,
const char *buf_name,
int index, uint64_t var) {
printf("msg='%s' BlockIdx=(%d, %d, %d), ThreadIdx=(%d, %d, %d): buffer=%s, "
"index=%d, dtype=uint64_t value=%llu\n",
msg, blockIdx.x, blockIdx.y, blockIdx.z, threadIdx.x, threadIdx.y,
threadIdx.z, buf_name, index, (unsigned long long)var);
}

TL_DEVICE void device_assert(bool cond) { assert(cond); }

TL_DEVICE void device_assert_with_msg(bool cond, const char *msg) {
Expand Down
Loading