Skip to content
Merged
Show file tree
Hide file tree
Changes from all commits
Commits
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
Submodule tvm updated from f4105f to f4affc
16 changes: 16 additions & 0 deletions src/tl_templates/cuda/debug.h
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
Copy link
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,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 "
Expand Down
Loading