Skip to content

[mlir][GPU] Do not strip location info when lowering to NVVM #120432

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
merged 1 commit into from
Dec 19, 2024
Merged

[mlir][GPU] Do not strip location info when lowering to NVVM #120432

Changes from all commits
Commits
Show all changes
1 commit
Select commit
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
1 change: 0 additions & 1 deletion mlir/lib/Dialect/GPU/Pipelines/GPUToNVVMPipeline.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -73,7 +73,6 @@ void buildCommonPassPipeline(
//===----------------------------------------------------------------------===//
void buildGpuPassPipeline(OpPassManager &pm,
const mlir::gpu::GPUToNVVMPipelineOptions &options) {
pm.addNestedPass<gpu::GPUModuleOp>(createStripDebugInfoPass());
Copy link
Collaborator

@joker-eph joker-eph Dec 18, 2024

Choose a reason for hiding this comment

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

Why was this added in the first place?

Copy link
Member Author

@matthias-springer matthias-springer Dec 18, 2024

Choose a reason for hiding this comment

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

Not sure. Looking at the git history (https://reviews.llvm.org/D155463), there was a test case that had it in the pass pipeline. See also https://reviews.llvm.org/D103187. Maybe @ThomasRaoux knows.

Copy link
Contributor

@ThomasRaoux ThomasRaoux Dec 18, 2024

Choose a reason for hiding this comment

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

I remember that passing ptx with debug info used to crash nvidia driver. I had not debugged it at the time though

Copy link
Member Author

@matthias-springer matthias-springer Dec 19, 2024 •
edited
Loading

Choose a reason for hiding this comment

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

Good to know. Given that the printf integration test works on my machine, this seems to be resolved. Let's see if the build bot complains after I merge this...

ConvertGpuOpsToNVVMOpsOptions opt;
opt.useBarePtrCallConv = options.kernelUseBarePtrCallConv;
opt.indexBitwidth = options.indexBitWidth;
Expand Down
Loading