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

[matthias-springer]

This is needed for a subsequent commit that reads location information when lowering gpu.assert.

[llvmbot]

@llvm/pr-subscribers-mlir

@llvm/pr-subscribers-mlir-gpu

Author: Matthias Springer (matthias-springer)

Changes

This is needed for a subsequent commit that reads location information when lowering gpu.assert.

---

Full diff: https://github.com/llvm/llvm-project/pull/120432.diff

1 Files Affected:

• (modified) mlir/lib/Dialect/GPU/Pipelines/GPUToNVVMPipeline.cpp (-1)

diff --git a/mlir/lib/Dialect/GPU/Pipelines/GPUToNVVMPipeline.cpp b/mlir/lib/Dialect/GPU/Pipelines/GPUToNVVMPipeline.cpp
index fb440756e0c1d5..20d7372eef85d5 100644
--- a/mlir/lib/Dialect/GPU/Pipelines/GPUToNVVMPipeline.cpp
+++ b/mlir/lib/Dialect/GPU/Pipelines/GPUToNVVMPipeline.cpp
@@ -73,7 +73,6 @@ void buildCommonPassPipeline(
 //===----------------------------------------------------------------------===//
 void buildGpuPassPipeline(OpPassManager &pm,
                           const mlir::gpu::GPUToNVVMPipelineOptions &options) {
-  pm.addNestedPass<gpu::GPUModuleOp>(createStripDebugInfoPass());
   ConvertGpuOpsToNVVMOpsOptions opt;
   opt.useBarePtrCallConv = options.kernelUseBarePtrCallConv;
   opt.indexBitwidth = options.indexBitWidth;

[joker-eph]

Why was this added in the first place?

[matthias-springer]

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.

[ThomasRaoux]

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

[matthias-springer]

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...