diff --git a/crates/llvm-export/src/export/config.rs b/crates/llvm-export/src/export/config.rs index 9d965da301..79e93d5e44 100644 --- a/crates/llvm-export/src/export/config.rs +++ b/crates/llvm-export/src/export/config.rs @@ -15,6 +15,24 @@ pub(super) const NVPTX_DATALAYOUT_FULL: &str = "e-p:64:64:64-p3:32:32:32-i1:8:8- i16:16:16-i32:32:32-i64:64:64-i128:128:128-f32:32:32-f64:64:64-f128:128:128-\ v16:16:16-v32:32:32-v64:64:64-v128:128:128-n16:32:64-a:8:8"; +/// Device debug metadata to emit. +#[derive(Clone, Copy, Debug, Default, PartialEq, Eq)] +pub enum DebugKind { + /// Do not emit LLVM debug metadata. + #[default] + Off, + /// Emit enough metadata for source-line breakpoints and stepping. + LineTables, + /// Reserved for variable/type debug info. For now this emits line tables. + Full, +} + +impl DebugKind { + pub fn line_tables_enabled(self) -> bool { + !matches!(self, Self::Off) + } +} + /// Configuration trait for export backends (PTX, LTOIR, etc.). /// /// This trait allows different backends to customize IR generation without @@ -40,6 +58,11 @@ pub trait ExportBackendConfig { /// Whether kernel definitions should use the `ptx_kernel` calling convention. fn emit_ptx_kernel_keyword(&self) -> bool; + + /// Which device debug metadata tier to emit. + fn debug_kind(&self) -> DebugKind { + DebugKind::Off + } } /// Default PTX export configuration. diff --git a/crates/llvm-export/src/export/debug.rs b/crates/llvm-export/src/export/debug.rs new file mode 100644 index 0000000000..2a79664127 --- /dev/null +++ b/crates/llvm-export/src/export/debug.rs @@ -0,0 +1,341 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. + * SPDX-License-Identifier: Apache-2.0 + */ + +//! Line-table debug metadata emission. +//! +//! Stage 2 deliberately emits only the debug-info pieces needed to map machine +//! instructions back to source lines. Locals, argument variables, lexical +//! blocks, and Rust types need more MIR/type plumbing and belong to later stages. + +use std::{ + fmt::Write, + path::{Path, PathBuf}, +}; + +use combine::stream::position::SourcePosition; +use pliron::{ + location::{Location, Source}, + uniqued_any, +}; + +use super::state::ModuleExportState; + +impl<'a> ModuleExportState<'a> { + pub(super) fn has_debug_metadata(&self) -> bool { + self.debug_compile_unit.is_some() + } + + pub(super) fn debug_subprogram_for_function( + &mut self, + name: &str, + loc: &Location, + ) -> Option { + if !self.debug_kind.line_tables_enabled() { + return None; + } + + let (path, pos) = self.source_position_from_location(loc)?; + let cu_id = self.ensure_debug_compile_unit(&path); + let file_id = self.ensure_debug_file(&path); + let subroutine_type_id = self.ensure_debug_subroutine_type(); + let name = escape_debug_string(name); + let line = pos.line; + let id = self.alloc_metadata_id(); + + self.debug_nodes.push(( + id, + format!( + "distinct !DISubprogram(name: \"{name}\", scope: !{file_id}, file: !{file_id}, \ + line: {line}, type: !{subroutine_type_id}, scopeLine: {line}, \ + spFlags: DISPFlagDefinition, unit: !{cu_id}, retainedNodes: !{{}})" + ), + )); + self.debug_subprogram_files.insert(id, path); + self.debug_subprogram_fallbacks + .insert(id, (pos.line, pos.column)); + + Some(id) + } + + pub(super) fn attach_debug_to_last_line( + &mut self, + output: &mut String, + output_before: usize, + scope: Option, + loc: &Location, + allow_scope_fallback: bool, + ) { + if output.len() == output_before { + return; + } + + let Some(scope) = scope else { + return; + }; + let location_id = self.debug_location_for_scope(scope, loc).or_else(|| { + if allow_scope_fallback { + // LLVM rejects inlinable calls inside a debug-scoped function + // unless the call itself has a location. When rustc/pliron did + // not give the call one, point it at the function line instead + // of letting opt discard the whole debug graph. + self.debug_fallback_location_for_scope(scope) + } else { + None + } + }); + let Some(location_id) = location_id else { + return; + }; + + if output.ends_with('\n') { + output.pop(); + writeln!(output, ", !dbg !{location_id}").unwrap(); + } + } + + pub(super) fn emit_debug_metadata(&mut self, output: &mut String) { + let Some(cu_id) = self.debug_compile_unit else { + return; + }; + + let dwarf_version_id = self.alloc_metadata_id(); + let debug_info_version_id = self.alloc_metadata_id(); + + writeln!(output, "!llvm.dbg.cu = !{{!{cu_id}}}").unwrap(); + writeln!( + output, + "!llvm.module.flags = !{{!{dwarf_version_id}, !{debug_info_version_id}}}" + ) + .unwrap(); + writeln!( + output, + "!{dwarf_version_id} = !{{i32 2, !\"Dwarf Version\", i32 2}}" + ) + .unwrap(); + writeln!( + output, + "!{debug_info_version_id} = !{{i32 2, !\"Debug Info Version\", i32 3}}" + ) + .unwrap(); + + for (id, node) in &self.debug_nodes { + writeln!(output, "!{id} = {node}").unwrap(); + } + } + + fn ensure_debug_compile_unit(&mut self, path: &Path) -> usize { + if let Some(id) = self.debug_compile_unit { + return id; + } + + let file_id = self.ensure_debug_file(path); + let id = self.alloc_metadata_id(); + self.debug_nodes.push(( + id, + format!( + "distinct !DICompileUnit(language: DW_LANG_Rust, file: !{file_id}, \ + producer: \"cuda-oxide\", isOptimized: true, runtimeVersion: 0, \ + emissionKind: LineTablesOnly)" + ), + )); + self.debug_compile_unit = Some(id); + id + } + + fn ensure_debug_file(&mut self, path: &Path) -> usize { + if let Some(id) = self.debug_files.get(path).copied() { + return id; + } + + let (filename, directory) = split_file_and_directory(path); + let filename = escape_debug_string(&filename); + let directory = escape_debug_string(&directory); + let id = self.alloc_metadata_id(); + + self.debug_nodes.push(( + id, + format!("!DIFile(filename: \"{filename}\", directory: \"{directory}\")"), + )); + self.debug_files.insert(path.to_path_buf(), id); + + id + } + + fn ensure_debug_subroutine_type(&mut self) -> usize { + if let Some(id) = self.debug_subroutine_type { + return id; + } + + let id = self.alloc_metadata_id(); + self.debug_nodes + .push((id, "!DISubroutineType(types: !{null})".to_string())); + self.debug_subroutine_type = Some(id); + + id + } + + fn debug_location_for_scope(&mut self, scope: usize, loc: &Location) -> Option { + if !self.debug_kind.line_tables_enabled() { + return None; + } + + let (path, pos) = self.source_position_from_location(loc)?; + if self + .debug_subprogram_files + .get(&scope) + .is_some_and(|scope_path| scope_path.as_path() != path) + { + return None; + } + + let key = (scope, pos.line, pos.column); + if let Some(id) = self.debug_locations.get(&key).copied() { + return Some(id); + } + + let id = self.alloc_metadata_id(); + self.debug_nodes.push(( + id, + format!( + "!DILocation(line: {}, column: {}, scope: !{})", + pos.line, pos.column, scope + ), + )); + self.debug_locations.insert(key, id); + + Some(id) + } + + fn debug_fallback_location_for_scope(&mut self, scope: usize) -> Option { + let (line, column) = self.debug_subprogram_fallbacks.get(&scope).copied()?; + let key = (scope, line, column); + if let Some(id) = self.debug_locations.get(&key).copied() { + return Some(id); + } + + let id = self.alloc_metadata_id(); + self.debug_nodes.push(( + id, + format!("!DILocation(line: {line}, column: {column}, scope: !{scope})"), + )); + self.debug_locations.insert(key, id); + + Some(id) + } + + fn source_position_from_location(&self, loc: &Location) -> Option<(PathBuf, SourcePosition)> { + match loc { + Location::SrcPos { + src: Source::File(path_key), + pos, + } if pos.line > 0 && pos.column > 0 => Some(( + uniqued_any::get(self.ctx, *path_key).clone(), + SourcePosition { + line: pos.line, + column: pos.column, + }, + )), + Location::SrcPos { .. } | Location::Unknown => None, + Location::Named { child_loc, .. } => self.source_position_from_location(child_loc), + Location::Fused { locations, .. } => locations + .iter() + .find_map(|loc| self.source_position_from_location(loc)), + Location::CallSite { caller, callee } => self + .source_position_from_location(caller) + .or_else(|| self.source_position_from_location(callee)), + } + } +} + +fn split_file_and_directory(path: &Path) -> (String, String) { + let filename = path + .file_name() + .filter(|name| !name.is_empty()) + .unwrap_or(path.as_os_str()) + .to_string_lossy() + .into_owned(); + + let directory = path + .parent() + .map(|parent| { + let dir = parent.to_string_lossy(); + if dir.is_empty() { + ".".to_string() + } else { + dir.into_owned() + } + }) + .unwrap_or_else(|| ".".to_string()); + + (filename, directory) +} + +fn escape_debug_string(input: &str) -> String { + let mut out = String::with_capacity(input.len()); + for ch in input.chars() { + match ch { + '\\' => out.push_str("\\5C"), + '"' => out.push_str("\\22"), + '\n' => out.push_str("\\0A"), + '\r' => out.push_str("\\0D"), + '\t' => out.push_str("\\09"), + _ => out.push(ch), + } + } + out +} + +#[cfg(test)] +mod tests { + use super::*; + use combine::stream::position::SourcePosition; + use pliron::{context::Context, location::Source}; + + #[test] + fn debug_strings_use_llvm_metadata_escapes() { + assert_eq!(escape_debug_string("a\\b\"c\n\t"), "a\\5Cb\\22c\\0A\\09"); + } + + #[test] + fn split_file_and_directory_handles_bare_and_nested_paths() { + assert_eq!( + split_file_and_directory(Path::new("kernel.rs")), + ("kernel.rs".to_string(), ".".to_string()) + ); + assert_eq!( + split_file_and_directory(Path::new("/tmp/cuda-oxide/kernel.rs")), + ("kernel.rs".to_string(), "/tmp/cuda-oxide".to_string()) + ); + } + + #[test] + fn source_position_from_location_unwraps_named_locations() { + let mut ctx = Context::new(); + let loc = Location::Named { + name: "lowered".to_string(), + child_loc: Box::new(Location::SrcPos { + src: Source::new_from_file(&mut ctx, PathBuf::from("/tmp/kernel.rs")), + pos: SourcePosition { + line: 12, + column: 4, + }, + }), + }; + let state = ModuleExportState::new( + &ctx, + false, + true, + super::super::config::DebugKind::LineTables, + ); + + let (path, pos) = state + .source_position_from_location(&loc) + .expect("location should unwrap"); + + assert_eq!(path, PathBuf::from("/tmp/kernel.rs")); + assert_eq!(pos.line, 12); + assert_eq!(pos.column, 4); + } +} diff --git a/crates/llvm-export/src/export/function.rs b/crates/llvm-export/src/export/function.rs index 2cc596ce61..ec9f24bacf 100644 --- a/crates/llvm-export/src/export/function.rs +++ b/crates/llvm-export/src/export/function.rs @@ -21,6 +21,7 @@ use pliron::{ }, context::Ptr, linked_list::ContainsLinkedList, + location::Located, op::Op, operation::Operation, r#type::Typed, @@ -227,6 +228,9 @@ impl<'a> ModuleExportState<'a> { .next(); if let Some(entry_block) = entry_block_opt { + let func_loc = func.get_operation().deref(self.ctx).loc(); + let debug_scope = self.debug_subprogram_for_function(&fixed_func_name, &func_loc); + write!(output, "define ").unwrap(); if is_kernel && self.emit_ptx_kernel_keyword { write!(output, "ptx_kernel ").unwrap(); @@ -274,7 +278,11 @@ impl<'a> ModuleExportState<'a> { // gets its `bar.sync.aligned` pushed into a `tid`-dependent branch // and deadlocks. opt's FunctionAttrs strips `convergent` from // functions it proves never reach a convergent op. - writeln!(output, ") #0 {{").unwrap(); + if let Some(scope_id) = debug_scope { + writeln!(output, ") #0 !dbg !{scope_id} {{").unwrap(); + } else { + writeln!(output, ") #0 {{").unwrap(); + } self.convergent_used = true; // Assign labels to all blocks @@ -447,6 +455,7 @@ impl<'a> ModuleExportState<'a> { &block_labels, &pred_map, i == 0, + debug_scope, output, )?; } @@ -482,6 +491,7 @@ impl<'a> ModuleExportState<'a> { block_labels: &HashMap, String>, pred_map: &PredecessorMap, is_entry: bool, + debug_scope: Option, output: &mut String, ) -> Result<(), String> { // Always print label to ensure it can be referenced by PHI nodes @@ -535,7 +545,14 @@ impl<'a> ModuleExportState<'a> { } for op in block.deref(self.ctx).iter(self.ctx) { - self.export_op(op, value_names, next_value_id, block_labels, output)?; + self.export_op( + op, + value_names, + next_value_id, + block_labels, + debug_scope, + output, + )?; } Ok(()) } diff --git a/crates/llvm-export/src/export/metadata.rs b/crates/llvm-export/src/export/metadata.rs index d511dee608..342d5035e0 100644 --- a/crates/llvm-export/src/export/metadata.rs +++ b/crates/llvm-export/src/export/metadata.rs @@ -3,78 +3,110 @@ * SPDX-License-Identifier: Apache-2.0 */ -//! NVVM annotations, version metadata, and llvm.used emission. +//! NVVM annotations and version metadata. +use std::collections::HashSet; use std::fmt::Write; use super::state::ModuleExportState; +pub(super) fn needs_nvvm_annotations( + state: &ModuleExportState, + emit_all_annotations: bool, +) -> bool { + let has_special_kernels = + !state.cluster_kernels.is_empty() || !state.launch_bounds_kernels.is_empty(); + has_special_kernels || (emit_all_annotations && !state.all_kernels.is_empty()) +} + /// Emit `!nvvm.annotations` metadata nodes for kernels. /// -/// Used by the `export_module_with_externs` path. +/// Metadata IDs come from `ModuleExportState` because `!nvvm.annotations`, +/// `!nvvmir.version`, and future debug-info nodes all share LLVM's flat +/// module metadata namespace. pub(super) fn emit_nvvm_annotations( output: &mut String, - state: &ModuleExportState, + state: &mut ModuleExportState, emit_all_annotations: bool, ) { let mut metadata_refs = Vec::new(); - let mut md_id = 0; // Collect names of kernels that have special configs - let special_kernel_names: std::collections::HashSet<&str> = state + let special_kernel_names: HashSet = state .cluster_kernels .iter() - .map(|k| k.name.as_str()) - .chain(state.launch_bounds_kernels.iter().map(|k| k.name.as_str())) + .map(|k| k.name.clone()) + .chain(state.launch_bounds_kernels.iter().map(|k| k.name.clone())) .collect(); // Emit basic annotation for kernels WITHOUT special configs if emit_all_annotations { - for kernel in state.all_kernels.iter() { - if !special_kernel_names.contains(kernel.name.as_str()) { - writeln!( - output, - "!{} = !{{ptr @{}, !\"kernel\", i32 1}}", - md_id, kernel.name - ) - .unwrap(); - metadata_refs.push(format!("!{}", md_id)); - md_id += 1; - } + let basic_kernels: Vec = state + .all_kernels + .iter() + .filter(|kernel| !special_kernel_names.contains(&kernel.name)) + .map(|kernel| kernel.name.clone()) + .collect(); + + for kernel_name in basic_kernels { + let md_id = state.alloc_metadata_id(); + writeln!( + output, + "!{} = !{{ptr @{}, !\"kernel\", i32 1}}", + md_id, kernel_name + ) + .unwrap(); + metadata_refs.push(format!("!{}", md_id)); } } // Emit cluster config annotations - for cfg in state.cluster_kernels.iter() { + let cluster_kernels: Vec<_> = state + .cluster_kernels + .iter() + .map(|cfg| (cfg.name.clone(), cfg.dim_x, cfg.dim_y, cfg.dim_z)) + .collect(); + + for (name, dim_x, dim_y, dim_z) in cluster_kernels { + let md_id = state.alloc_metadata_id(); writeln!( output, "!{} = !{{ptr @{}, !\"kernel\", i32 1, !\"cluster_dim_x\", i32 {}, !\"cluster_dim_y\", i32 {}, !\"cluster_dim_z\", i32 {}}}", - md_id, cfg.name, cfg.dim_x, cfg.dim_y, cfg.dim_z + md_id, name, dim_x, dim_y, dim_z ) .unwrap(); metadata_refs.push(format!("!{}", md_id)); - md_id += 1; } // Emit launch bounds annotations - for bounds in state.launch_bounds_kernels.iter() { - if let Some(min_blocks) = bounds.min_blocks { + let launch_bounds_kernels: Vec<_> = state + .launch_bounds_kernels + .iter() + .map(|bounds| (bounds.name.clone(), bounds.max_threads, bounds.min_blocks)) + .collect(); + + for (name, max_threads, min_blocks) in launch_bounds_kernels { + for (key, value) in [("maxntidx", max_threads), ("maxntidy", 1), ("maxntidz", 1)] { + let md_id = state.alloc_metadata_id(); writeln!( output, - "!{} = !{{ptr @{}, !\"kernel\", i32 1, !\"maxntidx\", i32 {}, !\"minctasm\", i32 {}}}", - md_id, bounds.name, bounds.max_threads, min_blocks + "!{} = !{{ptr @{}, !\"{}\", i32 {}}}", + md_id, name, key, value ) .unwrap(); - } else { + metadata_refs.push(format!("!{}", md_id)); + } + + if let Some(min_blocks) = min_blocks { + let md_id = state.alloc_metadata_id(); writeln!( output, - "!{} = !{{ptr @{}, !\"kernel\", i32 1, !\"maxntidx\", i32 {}}}", - md_id, bounds.name, bounds.max_threads + "!{} = !{{ptr @{}, !\"minctasm\", i32 {}}}", + md_id, name, min_blocks ) .unwrap(); + metadata_refs.push(format!("!{}", md_id)); } - metadata_refs.push(format!("!{}", md_id)); - md_id += 1; } // Emit named metadata referencing all annotation nodes @@ -88,34 +120,103 @@ pub(super) fn emit_nvvm_annotations( } } -/// Calculate the next metadata ID after annotations (for `!nvvmir.version`). -pub(super) fn md_id_after_annotations(state: &ModuleExportState) -> usize { - let mut count = state.all_kernels.len(); +pub(super) fn emit_nvvmir_version( + output: &mut String, + state: &mut ModuleExportState, + version: [i32; 4], +) { + let md_id = state.alloc_metadata_id(); + writeln!(output, "!nvvmir.version = !{{!{}}}", md_id).unwrap(); + writeln!( + output, + "!{} = !{{i32 {}, i32 {}, i32 {}, i32 {}}}", + md_id, version[0], version[1], version[2], version[3] + ) + .unwrap(); +} - // Subtract kernels that have special configs (they're not double-counted) - let special_kernel_names: std::collections::HashSet<&str> = state - .cluster_kernels - .iter() - .map(|k| k.name.as_str()) - .chain(state.launch_bounds_kernels.iter().map(|k| k.name.as_str())) - .collect(); +#[cfg(test)] +mod tests { + use super::*; + use crate::export::config::DebugKind; + use crate::export::state::{ + KernelClusterConfig, KernelInfo, KernelLaunchBounds, ModuleExportState, + }; + use pliron::context::Context; - for kernel in &state.all_kernels { - if special_kernel_names.contains(kernel.name.as_str()) { - count -= 1; - } + fn test_state<'a>(ctx: &'a Context) -> ModuleExportState<'a> { + ModuleExportState::new(ctx, true, false, DebugKind::Off) } - // Add cluster kernels - count += state.cluster_kernels.len(); + #[test] + fn allocator_returns_contiguous_module_metadata_ids() { + let ctx = Context::new(); + let mut state = test_state(&ctx); - // Add launch bounds kernels (each has multiple metadata entries) - for cfg in &state.launch_bounds_kernels { - count += 3; // maxntidx, maxntidy, maxntidz - if cfg.min_blocks.is_some() { - count += 1; // minctasm - } + assert_eq!(state.alloc_metadata_id(), 0); + assert_eq!(state.alloc_metadata_id(), 1); + assert_eq!(state.alloc_metadata_id(), 2); + assert_eq!(state.next_metadata_id(), 3); } - count + #[test] + fn nvvm_metadata_uses_one_shared_allocator() { + let ctx = Context::new(); + let mut state = test_state(&ctx); + state.all_kernels.push(KernelInfo { + name: "plain".into(), + }); + state.all_kernels.push(KernelInfo { + name: "clustered".into(), + }); + state.all_kernels.push(KernelInfo { + name: "bounded".into(), + }); + state.cluster_kernels.push(KernelClusterConfig { + name: "clustered".into(), + dim_x: 2, + dim_y: 3, + dim_z: 4, + }); + state.launch_bounds_kernels.push(KernelLaunchBounds { + name: "bounded".into(), + max_threads: 256, + min_blocks: Some(2), + }); + + let mut output = String::new(); + emit_nvvm_annotations(&mut output, &mut state, true); + emit_nvvmir_version(&mut output, &mut state, [2, 0, 3, 2]); + + assert_eq!( + output, + concat!( + "!0 = !{ptr @plain, !\"kernel\", i32 1}\n", + "!1 = !{ptr @clustered, !\"kernel\", i32 1, !\"cluster_dim_x\", i32 2, !\"cluster_dim_y\", i32 3, !\"cluster_dim_z\", i32 4}\n", + "!2 = !{ptr @bounded, !\"maxntidx\", i32 256}\n", + "!3 = !{ptr @bounded, !\"maxntidy\", i32 1}\n", + "!4 = !{ptr @bounded, !\"maxntidz\", i32 1}\n", + "!5 = !{ptr @bounded, !\"minctasm\", i32 2}\n", + "!nvvm.annotations = !{!0, !1, !2, !3, !4, !5}\n", + "!nvvmir.version = !{!6}\n", + "!6 = !{i32 2, i32 0, i32 3, i32 2}\n", + ) + ); + assert_eq!(state.next_metadata_id(), 7); + } + + #[test] + fn basic_kernel_annotations_are_skipped_when_backend_does_not_need_them() { + let ctx = Context::new(); + let mut state = test_state(&ctx); + state.all_kernels.push(KernelInfo { + name: "plain".into(), + }); + + let mut output = String::new(); + emit_nvvm_annotations(&mut output, &mut state, false); + + assert!(output.is_empty()); + assert_eq!(state.next_metadata_id(), 0); + } } diff --git a/crates/llvm-export/src/export/mod.rs b/crates/llvm-export/src/export/mod.rs index dd08d7cebc..38985cca92 100644 --- a/crates/llvm-export/src/export/mod.rs +++ b/crates/llvm-export/src/export/mod.rs @@ -37,6 +37,7 @@ //! - `metadata`: nvvm annotations and version, llvm.used mod config; +mod debug; mod externs; mod function; mod literals; @@ -47,7 +48,7 @@ mod ops; mod state; mod types; -pub use config::{ExportBackendConfig, NvvmExportConfig, PtxExportConfig}; +pub use config::{DebugKind, ExportBackendConfig, NvvmExportConfig, PtxExportConfig}; pub use externs::{AsDeviceExtern, DeviceExternAttrs, DeviceExternDecl}; use pliron::{builtin::ops::ModuleOp, context::Context}; diff --git a/crates/llvm-export/src/export/module.rs b/crates/llvm-export/src/export/module.rs index 358b951798..7eff71d5b4 100644 --- a/crates/llvm-export/src/export/module.rs +++ b/crates/llvm-export/src/export/module.rs @@ -23,7 +23,7 @@ use crate::ops; use super::{ config::ExportBackendConfig, externs::DeviceExternDecl, - metadata::{emit_nvvm_annotations, md_id_after_annotations}, + metadata::{emit_nvvm_annotations, emit_nvvmir_version, needs_nvvm_annotations}, state::ModuleExportState, }; @@ -37,7 +37,12 @@ pub(super) fn export_module_with_externs_impl( let mut output = String::new(); let emit_all_annotations = config.emit_all_kernel_annotations(); let emit_ptx_kernel_keyword = config.emit_ptx_kernel_keyword(); - let mut state = ModuleExportState::new(ctx, emit_all_annotations, emit_ptx_kernel_keyword); + let mut state = ModuleExportState::new( + ctx, + emit_all_annotations, + emit_ptx_kernel_keyword, + config.debug_kind(), + ); // 1. Header writeln!( @@ -168,28 +173,21 @@ pub(super) fn export_module_with_externs_impl( } // 6. nvvm.annotations metadata - let has_special_kernels = - !state.cluster_kernels.is_empty() || !state.launch_bounds_kernels.is_empty(); - let needs_annotations = - has_special_kernels || (emit_all_annotations && !state.all_kernels.is_empty()); - - if needs_annotations { + if needs_nvvm_annotations(&state, emit_all_annotations) { writeln!(&mut output).unwrap(); - emit_nvvm_annotations(&mut output, &state, emit_all_annotations); + emit_nvvm_annotations(&mut output, &mut state, emit_all_annotations); } // 7. nvvmir.version metadata (if backend requires) - // Must use a unique metadata ID that doesn't conflict with nvvm.annotations if config.emit_nvvmir_version() { writeln!(&mut output).unwrap(); - let ver = config.nvvmir_version(); - let md_id = md_id_after_annotations(&state); - writeln!( - &mut output, - "!nvvmir.version = !{{!{}}}\n!{} = !{{i32 {}, i32 {}, i32 {}, i32 {}}}", - md_id, md_id, ver[0], ver[1], ver[2], ver[3] - ) - .unwrap(); + emit_nvvmir_version(&mut output, &mut state, config.nvvmir_version()); + } + + // 8. DWARF line-table metadata (if requested and source locations exist) + if state.has_debug_metadata() { + writeln!(&mut output).unwrap(); + state.emit_debug_metadata(&mut output); } Ok(output) @@ -207,7 +205,12 @@ pub(super) fn export_module_to_string_with_config( let mut output = String::new(); let emit_all_annotations = config.emit_all_kernel_annotations(); let emit_ptx_kernel_keyword = config.emit_ptx_kernel_keyword(); - let mut state = ModuleExportState::new(ctx, emit_all_annotations, emit_ptx_kernel_keyword); + let mut state = ModuleExportState::new( + ctx, + emit_all_annotations, + emit_ptx_kernel_keyword, + config.debug_kind(), + ); // 1. Header writeln!( @@ -308,133 +311,22 @@ pub(super) fn export_module_to_string_with_config( // Emit nvvm.annotations metadata // - Default: Only for kernels with cluster configuration or launch bounds // - Alternate backends: May require annotations for ALL kernels - let has_special_kernels = - !state.cluster_kernels.is_empty() || !state.launch_bounds_kernels.is_empty(); - let needs_annotations = - has_special_kernels || (emit_all_annotations && !state.all_kernels.is_empty()); - - if needs_annotations { + if needs_nvvm_annotations(&state, emit_all_annotations) { writeln!(&mut output).unwrap(); - - let mut metadata_refs = Vec::new(); - let mut md_id = 0; - - // If backend requires annotations for all kernels, emit basic annotations first - // (unless they have cluster/launch_bounds which will be emitted below with more detail) - if emit_all_annotations { - // Collect names of kernels that have special configs (they'll get detailed annotations) - let special_kernel_names: std::collections::HashSet<&str> = state - .cluster_kernels - .iter() - .map(|k| k.name.as_str()) - .chain(state.launch_bounds_kernels.iter().map(|k| k.name.as_str())) - .collect(); - - // Emit basic annotation for kernels WITHOUT special configs - for kernel in state.all_kernels.iter() { - if !special_kernel_names.contains(kernel.name.as_str()) { - // Basic kernel annotation: !{ptr @kernel_name, !"kernel", i32 1} - writeln!( - &mut output, - "!{} = !{{ptr @{}, !\"kernel\", i32 1}}", - md_id, kernel.name - ) - .unwrap(); - metadata_refs.push(format!("!{}", md_id)); - md_id += 1; - } - } - } - - // Each kernel with cluster config gets its own metadata node - // Format: !{ptr @kernel_name, !"kernel", i32 1, !"cluster_dim_x", i32 X, ...} - for cfg in state.cluster_kernels.iter() { - writeln!( - &mut output, - "!{} = !{{ptr @{}, !\"kernel\", i32 1, !\"cluster_dim_x\", i32 {}, !\"cluster_dim_y\", i32 {}, !\"cluster_dim_z\", i32 {}}}", - md_id, cfg.name, cfg.dim_x, cfg.dim_y, cfg.dim_z - ) - .unwrap(); - metadata_refs.push(format!("!{}", md_id)); - md_id += 1; - } - - // Each kernel with launch bounds gets its own metadata node - // LLVM NVPTX expects separate annotations: !"maxntidx", !"maxntidy", !"maxntidz", !"minctapersm" - // See: https://llvm.org/docs/NVPTXUsage.html - for cfg in state.launch_bounds_kernels.iter() { - // Emit maxntidx (we use the single max_threads value for 1D block size) - writeln!( - &mut output, - "!{} = !{{ptr @{}, !\"maxntidx\", i32 {}}}", - md_id, cfg.name, cfg.max_threads - ) - .unwrap(); - metadata_refs.push(format!("!{}", md_id)); - md_id += 1; - - // Emit maxntidy = 1 (for complete 3D specification) - writeln!( - &mut output, - "!{} = !{{ptr @{}, !\"maxntidy\", i32 1}}", - md_id, cfg.name - ) - .unwrap(); - metadata_refs.push(format!("!{}", md_id)); - md_id += 1; - - // Emit maxntidz = 1 (for complete 3D specification) - writeln!( - &mut output, - "!{} = !{{ptr @{}, !\"maxntidz\", i32 1}}", - md_id, cfg.name - ) - .unwrap(); - metadata_refs.push(format!("!{}", md_id)); - md_id += 1; - - // Emit minctasm as separate metadata node if specified (generates .minnctapersm in PTX) - if let Some(min_blocks) = cfg.min_blocks { - writeln!( - &mut output, - "!{} = !{{ptr @{}, !\"minctasm\", i32 {}}}", - md_id, cfg.name, min_blocks - ) - .unwrap(); - metadata_refs.push(format!("!{}", md_id)); - md_id += 1; - } - } - - // The nvvm.annotations named metadata references all kernel metadata - writeln!( - &mut output, - "!nvvm.annotations = !{{{}}}", - metadata_refs.join(", ") - ) - .unwrap(); + emit_nvvm_annotations(&mut output, &mut state, emit_all_annotations); } // Emit !nvvmir.version metadata if backend requests it if config.emit_nvvmir_version() { writeln!(&mut output).unwrap(); - let version = config.nvvmir_version(); - writeln!( - &mut output, - "!nvvmir.version = !{{!{}}}", - md_id_after_annotations(&state) - ) - .unwrap(); - writeln!( - &mut output, - "!{} = !{{i32 {}, i32 {}, i32 {}, i32 {}}}", - md_id_after_annotations(&state), - version[0], - version[1], - version[2], - version[3] - ) - .unwrap(); + emit_nvvmir_version(&mut output, &mut state, config.nvvmir_version()); + } + + // Emit DWARF line-table metadata if debug export requested it and at least + // one function had a real source location. + if state.has_debug_metadata() { + writeln!(&mut output).unwrap(); + state.emit_debug_metadata(&mut output); } Ok(output) diff --git a/crates/llvm-export/src/export/ops.rs b/crates/llvm-export/src/export/ops.rs index 5c399652bc..79b9e6de27 100644 --- a/crates/llvm-export/src/export/ops.rs +++ b/crates/llvm-export/src/export/ops.rs @@ -17,6 +17,7 @@ use pliron::{ op_interfaces::{CallOpCallable, CallOpInterface}, }, context::Ptr, + location::Located, op::Op, operation::Operation, value::Value, @@ -216,6 +217,19 @@ impl<'op> TryFrom<&'op dyn Op> for LlvmOp<'op> { } } +impl LlvmOp<'_> { + fn emits_real_instruction(&self) -> bool { + !matches!( + self, + Self::Undef(_) | Self::Constant(_) | Self::AddressOf(_) + ) + } + + fn needs_scoped_debug_location(&self) -> bool { + matches!(self, Self::Call(_)) + } +} + impl<'a> ModuleExportState<'a> { pub(super) fn export_op( &mut self, @@ -223,10 +237,18 @@ impl<'a> ModuleExportState<'a> { value_names: &mut HashMap, next_value_id: &mut usize, block_labels: &HashMap, String>, + debug_scope: Option, output: &mut String, ) -> Result<(), String> { let op_ref = op.deref(self.ctx); + let op_loc = op_ref.loc(); let op_obj = Operation::get_op_dyn(op, self.ctx); + let llvm_op = LlvmOp::try_from(op_obj.as_ref()).ok(); + let should_attach_debug = llvm_op.as_ref().is_some_and(LlvmOp::emits_real_instruction); + let allow_scope_debug_fallback = llvm_op + .as_ref() + .is_some_and(LlvmOp::needs_scoped_debug_location); + let output_before = output.len(); // Register result names (skip if already named in pre-pass) for res in op_ref.results() { @@ -237,7 +259,7 @@ impl<'a> ModuleExportState<'a> { }); } - match LlvmOp::try_from(op_obj.as_ref()).ok() { + match llvm_op { // Terminators Some(LlvmOp::Return(op)) => self.emit_return(op, value_names, output)?, Some(LlvmOp::Unreachable(_)) => writeln!(output, " unreachable").unwrap(), @@ -372,6 +394,16 @@ impl<'a> ModuleExportState<'a> { .unwrap(), } + if should_attach_debug { + self.attach_debug_to_last_line( + output, + output_before, + debug_scope, + &op_loc, + allow_scope_debug_fallback, + ); + } + Ok(()) } diff --git a/crates/llvm-export/src/export/state.rs b/crates/llvm-export/src/export/state.rs index 7a2c37cde2..5e8da45ead 100644 --- a/crates/llvm-export/src/export/state.rs +++ b/crates/llvm-export/src/export/state.rs @@ -6,7 +6,9 @@ //! Exporter state and kernel bookkeeping. use pliron::{basic_block::BasicBlock, context::Ptr, value::Value}; -use std::collections::HashMap; +use std::{collections::HashMap, path::PathBuf}; + +use super::config::DebugKind; /// Map from block to its predecessors with the values passed to each predecessor. /// Used for PHI node generation when exporting to LLVM IR. @@ -48,6 +50,28 @@ pub(super) struct ModuleExportState<'a> { pub(super) emit_ptx_kernel_keyword: bool, /// Track device function names for @llvm.used (standalone device fn compilation) pub(super) device_functions: Vec, + /// Next `!N` metadata ID in this module. + /// + /// LLVM has one flat numbered metadata namespace per module. Today this is + /// used for NVVM annotations/version nodes; debug-info nodes will use the + /// same counter so the exporter never has to guess which IDs are free. + next_metadata_id: usize, + /// Which debug metadata tier this export should emit. + pub(super) debug_kind: DebugKind, + /// The single compile unit used for Stage 2 line-table debug info. + pub(super) debug_compile_unit: Option, + /// `DIFile` nodes keyed by the source path they describe. + pub(super) debug_files: HashMap, + /// Shared empty function type used by all line-table-only subprograms. + pub(super) debug_subroutine_type: Option, + /// `DISubprogram` file paths, used to avoid attaching locations to the wrong scope. + pub(super) debug_subprogram_files: HashMap, + /// Fallback line/column for calls that LLVM requires to have a location. + pub(super) debug_subprogram_fallbacks: HashMap, + /// `DILocation` nodes keyed by `(scope, line, column)`. + pub(super) debug_locations: HashMap<(usize, i32, i32), usize>, + /// Numbered debug metadata definitions, in allocation order. + pub(super) debug_nodes: Vec<(usize, String)>, } impl<'a> ModuleExportState<'a> { @@ -55,6 +79,7 @@ impl<'a> ModuleExportState<'a> { ctx: &'a pliron::context::Context, track_all_kernels: bool, emit_ptx_kernel_keyword: bool, + debug_kind: DebugKind, ) -> Self { Self { ctx, @@ -65,9 +90,29 @@ impl<'a> ModuleExportState<'a> { track_all_kernels, emit_ptx_kernel_keyword, device_functions: Vec::new(), + next_metadata_id: 0, + debug_kind, + debug_compile_unit: None, + debug_files: HashMap::new(), + debug_subroutine_type: None, + debug_subprogram_files: HashMap::new(), + debug_subprogram_fallbacks: HashMap::new(), + debug_locations: HashMap::new(), + debug_nodes: Vec::new(), } } + pub(super) fn alloc_metadata_id(&mut self) -> usize { + let id = self.next_metadata_id; + self.next_metadata_id += 1; + id + } + + #[cfg(test)] + pub(super) fn next_metadata_id(&self) -> usize { + self.next_metadata_id + } + /// Check if a function name is a known convergent intrinsic. /// /// These intrinsics require warp-synchronous execution semantics and must diff --git a/crates/llvm-export/tests/export_test.rs b/crates/llvm-export/tests/export_test.rs index 1b0f0b5a0c..ef528d987e 100644 --- a/crates/llvm-export/tests/export_test.rs +++ b/crates/llvm-export/tests/export_test.rs @@ -3,21 +3,76 @@ * SPDX-License-Identifier: Apache-2.0 */ +use combine::stream::position::SourcePosition; use llvm_export::{ - export::export_module_to_string, - ops::{AddressOfOp, BrOp, FuncOp, GepIndex, GetElementPtrOp, GlobalOp, InlineAsmOp, ReturnOp}, + export::{ + DebugKind, ExportBackendConfig, NvvmExportConfig, PtxExportConfig, export_module_to_string, + export_module_to_string_with_config, + }, + ops::{ + AddressOfOp, BrOp, CallOp, FuncOp, GepIndex, GetElementPtrOp, GlobalOp, InlineAsmOp, + ReturnOp, + }, types::{FuncType, VoidType}, }; use pliron::{ basic_block::BasicBlock, builtin::{ + attributes::{IntegerAttr, StringAttr}, + op_interfaces::CallOpCallable, ops::ModuleOp, types::{IntegerType, Signedness}, }, context::Context, + identifier::Identifier, linked_list::ContainsLinkedList, + location::{Located, Location, Source}, op::Op, + utils::apint::APInt, }; +use std::{num::NonZero, path::PathBuf}; + +struct DebugConfig { + inner: C, + debug_kind: DebugKind, +} + +impl ExportBackendConfig for DebugConfig { + fn datalayout(&self) -> &str { + self.inner.datalayout() + } + + fn emit_llvm_used(&self) -> bool { + self.inner.emit_llvm_used() + } + + fn emit_nvvmir_version(&self) -> bool { + self.inner.emit_nvvmir_version() + } + + fn nvvmir_version(&self) -> [i32; 4] { + self.inner.nvvmir_version() + } + + fn emit_all_kernel_annotations(&self) -> bool { + self.inner.emit_all_kernel_annotations() + } + + fn emit_ptx_kernel_keyword(&self) -> bool { + self.inner.emit_ptx_kernel_keyword() + } + + fn debug_kind(&self) -> DebugKind { + self.debug_kind + } +} + +fn src_location(ctx: &mut Context, file: &str, line: i32, column: i32) -> Location { + Location::SrcPos { + src: Source::new_from_file(ctx, PathBuf::from(file)), + pos: SourcePosition { line, column }, + } +} #[test] fn export_addressof_uses_symbol_when_definition_block_prints_later() { @@ -192,6 +247,254 @@ fn export_inline_asm_escapes_llvm_string_literals() { ); } +#[test] +fn nvvm_metadata_version_uses_next_allocated_metadata_id() { + let mut ctx = Context::new(); + + let module = ModuleOp::new(&mut ctx, "test_module".try_into().unwrap()); + let module_region = module.get_operation().deref(&ctx).get_region(0); + let module_block = { + let region = module_region.deref(&ctx); + region.iter(&ctx).next().unwrap() + }; + + let void_ty = VoidType::get(&ctx); + let func_ty = FuncType::get(&mut ctx, void_ty.to_ptr(), vec![], false); + let func = FuncOp::new(&mut ctx, "bounded_kernel".try_into().unwrap(), func_ty); + let entry = func.get_or_create_entry_block(&mut ctx); + ReturnOp::new(&mut ctx, None) + .get_operation() + .insert_at_back(entry, &ctx); + + let u32_ty = IntegerType::get(&mut ctx, 32, Signedness::Unsigned); + let width = NonZero::new(32).unwrap(); + let max_threads = IntegerAttr::new(u32_ty, APInt::from_u32(256, width)); + let min_blocks = IntegerAttr::new(u32_ty, APInt::from_u32(2, width)); + + { + let attrs = &mut func.get_operation().deref_mut(&ctx).attributes; + attrs.set( + Identifier::try_from("gpu_kernel").unwrap(), + StringAttr::new("true".into()), + ); + attrs.set(Identifier::try_from("maxntid").unwrap(), max_threads); + attrs.set(Identifier::try_from("minctasm").unwrap(), min_blocks); + } + + func.get_operation().insert_at_back(module_block, &ctx); + + let ir = export_module_to_string_with_config(&ctx, &module, &NvvmExportConfig) + .expect("NVVM export succeeds"); + + assert!( + ir.contains("!nvvm.annotations = !{!0, !1, !2, !3}"), + "launch-bounds annotations should occupy !0..!3:\n{ir}" + ); + assert!( + ir.contains("!nvvmir.version = !{!4}\n!4 = !{i32 2, i32 0, i32 3, i32 2}"), + "version metadata should use the next allocated ID:\n{ir}" + ); +} + +#[test] +fn line_table_debug_metadata_emits_function_scope_and_instruction_locations() { + let mut ctx = Context::new(); + + let module = ModuleOp::new(&mut ctx, "test_module".try_into().unwrap()); + let module_region = module.get_operation().deref(&ctx).get_region(0); + let module_block = { + let region = module_region.deref(&ctx); + region.iter(&ctx).next().unwrap() + }; + + let void_ty = VoidType::get(&ctx); + let func_ty = FuncType::get(&mut ctx, void_ty.to_ptr(), vec![], false); + let func = FuncOp::new(&mut ctx, "debug_kernel".try_into().unwrap(), func_ty); + let func_loc = src_location(&mut ctx, "/tmp/cuda-oxide/tests/kernel.rs", 7, 1); + func.get_operation().deref_mut(&ctx).set_loc(func_loc); + + let entry = func.get_or_create_entry_block(&mut ctx); + let ret = ReturnOp::new(&mut ctx, None); + let ret_loc = src_location(&mut ctx, "/tmp/cuda-oxide/tests/kernel.rs", 8, 5); + ret.get_operation().deref_mut(&ctx).set_loc(ret_loc); + ret.get_operation().insert_at_back(entry, &ctx); + + func.get_operation().insert_at_back(module_block, &ctx); + + let config = DebugConfig { + inner: PtxExportConfig, + debug_kind: DebugKind::LineTables, + }; + let ir = + export_module_to_string_with_config(&ctx, &module, &config).expect("debug export succeeds"); + + let define_line = ir + .lines() + .find(|line| line.starts_with("define ")) + .expect("function definition"); + assert!( + define_line.contains("!dbg !"), + "function definition should reference its DISubprogram:\n{ir}" + ); + + let ret_line = ir + .lines() + .find(|line| line.trim_start().starts_with("ret void")) + .expect("return instruction"); + assert!( + ret_line.contains(", !dbg !"), + "real instructions should carry DILocation attachments:\n{ir}" + ); + + assert!( + ir.contains("!llvm.dbg.cu = !{!"), + "module should reference a compile unit:\n{ir}" + ); + assert!( + ir.contains("!llvm.module.flags = !{!"), + "module should declare debug-info flags:\n{ir}" + ); + assert!( + ir.contains("!DIFile(filename: \"kernel.rs\", directory: \"/tmp/cuda-oxide/tests\")"), + "source path should be split into DIFile filename and directory:\n{ir}" + ); + assert!( + ir.contains("distinct !DICompileUnit(language: DW_LANG_Rust"), + "debug export should describe the Rust compile unit:\n{ir}" + ); + assert!( + ir.contains("distinct !DISubprogram(name: \"debug_kernel\""), + "function definition should get a DISubprogram:\n{ir}" + ); + assert!( + ir.contains("!DILocation(line: 8, column: 5, scope: !"), + "instruction location should preserve the source line and column:\n{ir}" + ); +} + +#[test] +fn debug_metadata_shares_allocator_with_nvvm_metadata() { + let mut ctx = Context::new(); + + let module = ModuleOp::new(&mut ctx, "test_module".try_into().unwrap()); + let module_region = module.get_operation().deref(&ctx).get_region(0); + let module_block = { + let region = module_region.deref(&ctx); + region.iter(&ctx).next().unwrap() + }; + + let void_ty = VoidType::get(&ctx); + let func_ty = FuncType::get(&mut ctx, void_ty.to_ptr(), vec![], false); + let func = FuncOp::new(&mut ctx, "debug_kernel".try_into().unwrap(), func_ty); + let func_loc = src_location(&mut ctx, "/tmp/cuda-oxide/tests/kernel.rs", 10, 1); + func.get_operation().deref_mut(&ctx).set_loc(func_loc); + + { + let attrs = &mut func.get_operation().deref_mut(&ctx).attributes; + attrs.set( + Identifier::try_from("gpu_kernel").unwrap(), + StringAttr::new("true".into()), + ); + } + + let entry = func.get_or_create_entry_block(&mut ctx); + let ret = ReturnOp::new(&mut ctx, None); + let ret_loc = src_location(&mut ctx, "/tmp/cuda-oxide/tests/kernel.rs", 11, 5); + ret.get_operation().deref_mut(&ctx).set_loc(ret_loc); + ret.get_operation().insert_at_back(entry, &ctx); + + func.get_operation().insert_at_back(module_block, &ctx); + + let config = DebugConfig { + inner: NvvmExportConfig, + debug_kind: DebugKind::LineTables, + }; + let ir = export_module_to_string_with_config(&ctx, &module, &config) + .expect("debug NVVM export succeeds"); + + assert!( + ir.contains("!0 = !DIFile(filename: \"kernel.rs\", directory: \"/tmp/cuda-oxide/tests\")"), + "debug file node should take the first metadata ID:\n{ir}" + ); + assert!( + ir.contains("!4 = !DILocation(line: 11, column: 5, scope: !3)"), + "instruction location should be allocated before NVVM metadata:\n{ir}" + ); + assert!( + ir.contains("!5 = !{ptr @debug_kernel, !\"kernel\", i32 1}"), + "NVVM annotations should continue after debug metadata:\n{ir}" + ); + assert!( + ir.contains("!nvvm.annotations = !{!5}"), + "named NVVM metadata should reference its allocated node:\n{ir}" + ); + assert!( + ir.contains("!nvvmir.version = !{!6}\n!6 = !{i32 2, i32 0, i32 3, i32 2}"), + "NVVM version should use the next free metadata ID:\n{ir}" + ); + assert!( + ir.contains("!llvm.module.flags = !{!7, !8}"), + "debug module flags should also use the shared allocator:\n{ir}" + ); +} + +#[test] +fn line_table_debug_metadata_adds_fallback_locations_to_calls() { + let mut ctx = Context::new(); + + let module = ModuleOp::new(&mut ctx, "test_module".try_into().unwrap()); + let module_region = module.get_operation().deref(&ctx).get_region(0); + let module_block = { + let region = module_region.deref(&ctx); + region.iter(&ctx).next().unwrap() + }; + + let i32_ty = IntegerType::get(&mut ctx, 32, Signedness::Signless); + let void_ty = VoidType::get(&ctx); + let helper_ty = FuncType::get(&mut ctx, i32_ty.to_ptr(), vec![], false); + let helper = FuncOp::new(&mut ctx, "helper".try_into().unwrap(), helper_ty); + helper.get_operation().insert_at_back(module_block, &ctx); + + let caller_ty = FuncType::get(&mut ctx, void_ty.to_ptr(), vec![], false); + let caller = FuncOp::new(&mut ctx, "debug_kernel".try_into().unwrap(), caller_ty); + let caller_loc = src_location(&mut ctx, "/tmp/cuda-oxide/tests/kernel.rs", 20, 3); + caller.get_operation().deref_mut(&ctx).set_loc(caller_loc); + + let entry = caller.get_or_create_entry_block(&mut ctx); + let call = CallOp::new( + &mut ctx, + CallOpCallable::Direct("helper".try_into().unwrap()), + helper_ty, + vec![], + ); + call.get_operation().insert_at_back(entry, &ctx); + ReturnOp::new(&mut ctx, None) + .get_operation() + .insert_at_back(entry, &ctx); + + caller.get_operation().insert_at_back(module_block, &ctx); + + let config = DebugConfig { + inner: PtxExportConfig, + debug_kind: DebugKind::LineTables, + }; + let ir = + export_module_to_string_with_config(&ctx, &module, &config).expect("debug export succeeds"); + + let call_line = ir + .lines() + .find(|line| line.contains("call i32 @helper()")) + .expect("call instruction"); + assert!( + call_line.contains(", !dbg !"), + "calls without their own source span should use the function fallback location:\n{ir}" + ); + assert!( + ir.contains("!DILocation(line: 20, column: 3, scope: !"), + "fallback call location should point at the caller's function line:\n{ir}" + ); +} + /// Scans the textual LLVM IR and asserts that every `%vN` token appearing in /// an operand position has a corresponding `%vN = ...` definition somewhere /// in the module. Operates on `%v` temporaries only because that's the diff --git a/crates/mir-importer/src/pipeline.rs b/crates/mir-importer/src/pipeline.rs index 9e557c795d..0a4f5e55cc 100644 --- a/crates/mir-importer/src/pipeline.rs +++ b/crates/mir-importer/src/pipeline.rs @@ -33,6 +33,7 @@ //! //! Override with `CUDA_OXIDE_TARGET=` environment variable. +use llvm_export::export::{DebugKind, ExportBackendConfig}; use pliron::common_traits::Verify; use rustc_public::mir::mono::Instance; @@ -173,6 +174,9 @@ pub struct PipelineConfig { /// Currently supports NVVM 20 dialect (Blackwell+). Architecture is /// controlled by `--arch` flag. pub emit_nvvm_ir: bool, + /// Device debug metadata tier. Full variable/type debug is not implemented + /// yet, so `Full` currently emits line tables. + pub debug_kind: DebugKind, } impl Default for PipelineConfig { @@ -184,6 +188,7 @@ impl Default for PipelineConfig { show_mir_dialect: false, show_llvm_dialect: false, emit_nvvm_ir: false, + debug_kind: DebugKind::Off, } } } @@ -387,7 +392,14 @@ pub fn run_pipeline( eprintln!("\n=== Exporting to LLVM IR ({} mode) ===", mode); } let ll_path = config.output_dir.join(format!("{}.ll", config.output_name)); - let _llvm_ir = export_llvm_ir(&ctx, module_op_ptr, device_externs, &ll_path, emit_nvvm_ir)?; + let _llvm_ir = export_llvm_ir( + &ctx, + module_op_ptr, + device_externs, + &ll_path, + emit_nvvm_ir, + config.debug_kind, + )?; if config.verbose { eprintln!("LLVM IR written to {}", ll_path.display()); } @@ -650,16 +662,23 @@ fn export_llvm_ir( device_externs: &[DeviceExternDecl], path: &Path, emit_nvvm_ir: bool, + debug_kind: DebugKind, ) -> Result { let module_op = Operation::get_op::(module_op_ptr, ctx) .ok_or_else(|| PipelineError::Export("Not a module op".to_string()))?; let llvm_ir = if emit_nvvm_ir { - let config = llvm_export::export::NvvmExportConfig; + let config = PipelineExportConfig { + inner: llvm_export::export::NvvmExportConfig, + debug_kind, + }; llvm_export::export::export_module_with_externs(ctx, &module_op, device_externs, &config) .map_err(PipelineError::Export)? } else { - let config = llvm_export::export::PtxExportConfig; + let config = PipelineExportConfig { + inner: llvm_export::export::PtxExportConfig, + debug_kind, + }; llvm_export::export::export_module_with_externs(ctx, &module_op, device_externs, &config) .map_err(PipelineError::Export)? }; @@ -669,6 +688,41 @@ fn export_llvm_ir( Ok(llvm_ir) } +struct PipelineExportConfig { + inner: C, + debug_kind: DebugKind, +} + +impl ExportBackendConfig for PipelineExportConfig { + fn datalayout(&self) -> &str { + self.inner.datalayout() + } + + fn emit_llvm_used(&self) -> bool { + self.inner.emit_llvm_used() + } + + fn emit_nvvmir_version(&self) -> bool { + self.inner.emit_nvvmir_version() + } + + fn nvvmir_version(&self) -> [i32; 4] { + self.inner.nvvmir_version() + } + + fn emit_all_kernel_annotations(&self) -> bool { + self.inner.emit_all_kernel_annotations() + } + + fn emit_ptx_kernel_keyword(&self) -> bool { + self.inner.emit_ptx_kernel_keyword() + } + + fn debug_kind(&self) -> DebugKind { + self.debug_kind + } +} + /// Checks for WGMMA instructions (Hopper sm_90a only, NOT forward-compatible). /// /// WGMMA (Warpgroup Matrix Multiply-Accumulate) requires sm_90a specifically. @@ -1141,6 +1195,7 @@ mod tests { assert!(!config.show_mir_dialect); assert!(!config.show_llvm_dialect); assert!(!config.emit_nvvm_ir); + assert_eq!(config.debug_kind, DebugKind::Off); } #[test] diff --git a/crates/mir-importer/src/translator/body.rs b/crates/mir-importer/src/translator/body.rs index 8a9b8908aa..70e26bc31c 100644 --- a/crates/mir-importer/src/translator/body.rs +++ b/crates/mir-importer/src/translator/body.rs @@ -22,6 +22,7 @@ use super::block; use super::types; use crate::error::{TranslationErr, TranslationResult}; +use crate::translator::location::span_to_location; use crate::translator::values::{self, SlotAddrSpaceMap, ValueMap}; use dialect_mir::ops::MirFuncOp; use dialect_mir::types::address_space; @@ -30,7 +31,7 @@ use pliron::builtin::op_interfaces::SymbolOpInterface; use pliron::context::{Context, Ptr}; use pliron::identifier::{Identifier, Legaliser}; use pliron::input_err_noloc; -use pliron::location::{Located, Location}; +use pliron::location::Located; use pliron::op::Op; use pliron::operation::Operation; @@ -420,12 +421,9 @@ pub fn translate_body( 1, // 1 region for function body ); - // Set function location - // Use body span for location - let loc = Location::Named { - name: format!("{:?}", body.span), - child_loc: Box::new(Location::Unknown), - }; + // Set the function location from rustc's body span. This becomes the + // default scope line for line-table debug info once LLVM export is enabled. + let loc = span_to_location(ctx, body.span); op_ptr.deref_mut(ctx).set_loc(loc); // Create MirFuncOp and set the function type attribute and symbol name diff --git a/crates/mir-importer/src/translator/location.rs b/crates/mir-importer/src/translator/location.rs new file mode 100644 index 0000000000..5c6bcee5c0 --- /dev/null +++ b/crates/mir-importer/src/translator/location.rs @@ -0,0 +1,39 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. + * SPDX-License-Identifier: Apache-2.0 + */ + +//! Source-location helpers for MIR translation. +//! +//! rustc gives us source spans while we still know which MIR statement or +//! terminator produced an op. Store those spans as structured pliron locations +//! here, before later lowering stages only see generic operations. + +use std::path::PathBuf; + +use combine::stream::position::SourcePosition; +use pliron::context::Context; +use pliron::location::{Location, Source}; +use rustc_public::ty::Span; + +/// Convert a rustc source span into a pliron source position. +/// +/// A span with no real file or no positive line number cannot produce useful +/// line-table debug info, so it becomes `Unknown` and the LLVM exporter skips +/// `!dbg` for that operation. +pub(crate) fn span_to_location(ctx: &mut Context, span: Span) -> Location { + let file = span.get_filename(); + let lines = span.get_lines(); + + if file.is_empty() || lines.start_line == 0 || lines.start_col == 0 { + return Location::Unknown; + } + + Location::SrcPos { + src: Source::new_from_file(ctx, PathBuf::from(file)), + pos: SourcePosition { + line: lines.start_line as i32, + column: lines.start_col as i32, + }, + } +} diff --git a/crates/mir-importer/src/translator/mod.rs b/crates/mir-importer/src/translator/mod.rs index 8a68c41c0c..f32a2b3533 100644 --- a/crates/mir-importer/src/translator/mod.rs +++ b/crates/mir-importer/src/translator/mod.rs @@ -46,6 +46,7 @@ pub mod block; pub mod body; pub(crate) mod layout; +pub(crate) mod location; pub mod rvalue; pub mod statement; pub mod terminator; diff --git a/crates/mir-importer/src/translator/statement.rs b/crates/mir-importer/src/translator/statement.rs index 97a43745d2..c4388ad940 100644 --- a/crates/mir-importer/src/translator/statement.rs +++ b/crates/mir-importer/src/translator/statement.rs @@ -35,6 +35,7 @@ use super::types; use crate::error::{TranslationErr, TranslationResult}; +use crate::translator::location::span_to_location; use crate::translator::rvalue; use crate::translator::values::ValueMap; use dialect_mir::ops::{MirStorageDeadOp, MirStorageLiveOp, MirStoreOp}; @@ -65,11 +66,7 @@ pub fn translate_statement( block_ptr: Ptr, prev_op: Option>, ) -> TranslationResult>> { - // Use Debug representation of the span as location - let loc = Location::Named { - name: format!("{:?}", stmt.span), - child_loc: Box::new(Location::Unknown), - }; + let loc = span_to_location(ctx, stmt.span); match &stmt.kind { mir::StatementKind::Assign(place, rvalue) => { diff --git a/crates/mir-importer/src/translator/terminator/mod.rs b/crates/mir-importer/src/translator/terminator/mod.rs index 0892594005..6e37cce77f 100644 --- a/crates/mir-importer/src/translator/terminator/mod.rs +++ b/crates/mir-importer/src/translator/terminator/mod.rs @@ -60,6 +60,7 @@ pub mod intrinsics; use super::types; use crate::error::{TranslationErr, TranslationResult}; +use crate::translator::location::span_to_location; use crate::translator::rvalue; use crate::translator::values::ValueMap; use dialect_mir::ops::{ @@ -110,11 +111,7 @@ pub fn translate_terminator( block_map: &[Ptr], legaliser: &mut Legaliser, ) -> TranslationResult> { - // Use Debug representation of the span as location - let loc = Location::Named { - name: format!("{:?}", term.span), - child_loc: Box::new(Location::Unknown), - }; + let loc = span_to_location(ctx, term.span); match &term.kind { mir::TerminatorKind::Return => translate_return(ctx, value_map, block_ptr, prev_op, loc), diff --git a/crates/rustc-codegen-cuda/src/device_codegen.rs b/crates/rustc-codegen-cuda/src/device_codegen.rs index 85daeae3e8..0f4364673d 100644 --- a/crates/rustc-codegen-cuda/src/device_codegen.rs +++ b/crates/rustc-codegen-cuda/src/device_codegen.rs @@ -94,6 +94,7 @@ use crate::collector::{CollectedFunction, DeviceExternDecl}; use rustc_middle::ty::{Ty, TyCtxt, TyKind}; +use rustc_session::config::DebugInfo; use std::path::PathBuf; /// Convert a rustc type to an LLVM type string for device extern declarations. @@ -459,6 +460,8 @@ pub fn generate_device_code<'tcx>( } } + let debug_kind = device_debug_kind(tcx.sess.opts.debuginfo); + // Create pipeline config let pipeline_config = mir_importer::PipelineConfig { output_dir: output_dir.clone(), @@ -467,6 +470,7 @@ pub fn generate_device_code<'tcx>( show_mir_dialect: show_mir, show_llvm_dialect: show_llvm, emit_nvvm_ir, + debug_kind, }; // Run the cuda-oxide pipeline! @@ -527,6 +531,37 @@ pub fn generate_device_code<'tcx>( } } +fn device_debug_kind(rustc_debug: DebugInfo) -> llvm_export::export::DebugKind { + device_debug_kind_with_override( + rustc_debug, + std::env::var("CUDA_OXIDE_DEBUG").ok().as_deref(), + ) +} + +fn device_debug_kind_with_override( + rustc_debug: DebugInfo, + override_value: Option<&str>, +) -> llvm_export::export::DebugKind { + if let Some(value) = override_value { + match value.trim().to_ascii_lowercase().as_str() { + "0" | "off" | "none" => return llvm_export::export::DebugKind::Off, + "1" | "line" | "lines" | "line-tables" | "line-tables-only" => { + return llvm_export::export::DebugKind::LineTables; + } + "2" | "full" => return llvm_export::export::DebugKind::Full, + _ => {} + } + } + + match rustc_debug { + DebugInfo::None => llvm_export::export::DebugKind::Off, + DebugInfo::LineDirectivesOnly + | DebugInfo::LineTablesOnly + | DebugInfo::Limited + | DebugInfo::Full => llvm_export::export::DebugKind::LineTables, + } +} + fn read_compilation_artifact( result: &mir_importer::CompilationResult, ) -> Result, DeviceCodegenError> { @@ -564,6 +599,38 @@ mod tests { assert_eq!(config.output_name, "kernel"); } + #[test] + fn device_debug_kind_follows_rustc_debuginfo() { + assert_eq!( + device_debug_kind_with_override(DebugInfo::None, None), + llvm_export::export::DebugKind::Off + ); + assert_eq!( + device_debug_kind_with_override(DebugInfo::LineTablesOnly, None), + llvm_export::export::DebugKind::LineTables + ); + assert_eq!( + device_debug_kind_with_override(DebugInfo::Full, None), + llvm_export::export::DebugKind::LineTables + ); + } + + #[test] + fn device_debug_kind_env_override_wins() { + assert_eq!( + device_debug_kind_with_override(DebugInfo::Full, Some("off")), + llvm_export::export::DebugKind::Off + ); + assert_eq!( + device_debug_kind_with_override(DebugInfo::None, Some(" Line-Tables ")), + llvm_export::export::DebugKind::LineTables + ); + assert_eq!( + device_debug_kind_with_override(DebugInfo::None, Some("full")), + llvm_export::export::DebugKind::Full + ); + } + #[test] fn read_compilation_artifact_uses_declared_nvvm_ir_path() { let temp_dir = unique_temp_dir("cuda-codegen-artifact");