Skip to content
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.

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

[Feat] Debug improvements #474

Merged
merged 7 commits into from
Feb 18, 2025

Conversation

wingertge
Copy link
Contributor

@wingertge wingertge commented Feb 10, 2025

  • Moves debug info from separate instructions that need to be kept in order, to an optional field on each instruction. The scope keeps track of the current debug location if debug symbols are enabled, and attaches that location to each instruction.
  • The scope also keeps track of the unique set of kernel function sources so consumers don't need to scan over all instructions to find any NonSemantic::Sources and correlate them to specific functions.
  • Adds #line directives to C++ so we have at least some debug info. Example:
#line 108 "crates/cubecl-linalg/src/matmul/components/batch/one_to_one.rs"
  const uint l_3 = l_2 - uint(1);
  const uint l_4 = info[info[uint(9)] + l_3];
  const uint l_13 = info[uint(8)];
#line 22 "crates/cubecl-linalg/src/matmul/components/batch/shared.rs"
  const uint l_14 = l_13 - uint(2);
  const uint l_15 = info[info[uint(11)] + l_14];
#line 22 "crates/cubecl-linalg/src/matmul/components/batch/shared.rs"
  const uint l_16 = blockIdx.z * l_15;
  • Variables that represent actual bindings in Rust now have their names tracked so they can be properly labeled in the debug info.

Finally, the actual Rust source can now be attached in one of two ways:

  1. Setting the src_file argument on at least one #[cube] macro per file. This uses the fact that include_str! is relative by just including a file in itself. But to do that we need to know the actual name of the file.
  2. Using a nightly compiler and setting the CUBECL_DEBUG_NIGHTLY environment variable during compilation. This will use the unstable Span::source_file API to automatically fetch the source and include it in the source_expand invocation.

Copy link
Member

@nathanielsimard nathanielsimard left a comment

Choose a reason for hiding this comment

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

I would also add an example output in the PR!

@@ -614,6 +618,7 @@ for ({i_ty} {i} = {start}; {i} {cmp} {end}; {increment}) {{
}
}
Instruction::Pipeline(pipeline_ops) => write!(f, "{pipeline_ops}"),
Instruction::Line { file, line } => writeln!(f, "#line {line} \"{file}\""),
Copy link
Member

Choose a reason for hiding this comment

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

I'm curious what it looks like!

Copy link
Contributor Author

Choose a reason for hiding this comment

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

I added an example to the PR text. A debugger/profiler should use these to keep track of actual source locations, but I don't know how well they actually keep track of things especially with the JIT compiler.

Comment on lines +51 to +52
pub source_loc: Option<SourceLoc>,
pub entry_loc: Option<SourceLoc>,
Copy link
Member

Choose a reason for hiding this comment

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

What's the difference between the two?

Copy link
Contributor Author

Choose a reason for hiding this comment

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

The entry_loc is basically just the first source_loc ever set for a kernel, so we can set the source loc at the start of a kernel without having to look ahead to the first update. Otherwise you'd need to always scan ahead which is harder than just keeping track of the first value.

@nathanielsimard nathanielsimard merged commit 67e6d26 into tracel-ai:main Feb 18, 2025
5 checks passed
@wingertge wingertge deleted the feat/debug-improvements branch February 22, 2025 11:31
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
None yet
Projects
None yet
Development

Successfully merging this pull request may close these issues.

2 participants