; ModuleID = 'LLVMDialectModule'
source_filename = "LLVMDialectModule"
target datalayout = "e-p3:32:32-p4:32:32-p5:32:32-p6:32:32-p7:32:32-i64:64-i128:128-i256:256-v16:16-v32:32-n16:32:64"

; Function Attrs: nounwind
define ptx_kernel void @_smoke_add_one_kernel(ptr addrspace(1) %0, ptr addrspace(1) %1, i32 %2, ptr addrspace(1) readnone captures(none) %3, ptr addrspace(1) readnone captures(none) %4) local_unnamed_addr #0 !dbg !4 {
  %6 = tail call i32 @llvm.nvvm.read.ptx.sreg.ctaid.x(), !dbg !12
  %7 = shl i32 %6, 7, !dbg !13
  %8 = tail call i32 @llvm.nvvm.read.ptx.sreg.tid.x(), !dbg !14
  %9 = and i32 %8, 127, !dbg !14
  %10 = or disjoint i32 %7, %9, !dbg !13
  %11 = icmp slt i32 %10, %2, !dbg !15
  %12 = sext i32 %10 to i64, !dbg !16
  %13 = getelementptr [4 x i8], ptr addrspace(1) %0, i64 %12, !dbg !16
  %14 = tail call i32 asm sideeffect "mov.u32 $0, $1;\0A\09@$3 ld.global.b32 { $0 }, [ $2 + 0 ];", "=r,r,l,b"(i32 0, ptr addrspace(1) %13, i1 %11) #2, !dbg !17
  %15 = bitcast i32 %14 to float, !dbg !17
  %16 = getelementptr [4 x i8], ptr addrspace(1) %1, i64 %12, !dbg !18
  %17 = fadd float %15, 1.000000e+00, !dbg !19
  %18 = bitcast float %17 to i32, !dbg !20
  tail call void asm sideeffect "@$2 st.global.b32 [ $1 + 0 ], { $0 };", "r,l,b"(i32 %18, ptr addrspace(1) %16, i1 %11) #2, !dbg !20
  ret void, !dbg !21
}

; Function Attrs: mustprogress nocallback nofree nosync nounwind speculatable willreturn memory(none)
declare noundef range(i32 0, 2147483647) i32 @llvm.nvvm.read.ptx.sreg.ctaid.x() #1

; Function Attrs: mustprogress nocallback nofree nosync nounwind speculatable willreturn memory(none)
declare noundef range(i32 0, 1024) i32 @llvm.nvvm.read.ptx.sreg.tid.x() #1

attributes #0 = { nounwind "nvvm.reqntid"="128" }
attributes #1 = { mustprogress nocallback nofree nosync nounwind speculatable willreturn memory(none) }
attributes #2 = { nounwind }

!llvm.dbg.cu = !{!0}
!llvm.module.flags = !{!2, !3}

!0 = distinct !DICompileUnit(language: DW_LANG_C, file: !1, producer: "triton", isOptimized: true, runtimeVersion: 0, emissionKind: LineTablesOnly)
!1 = !DIFile(filename: "vllm_support.py", directory: "C:\\Users\\johnn\\OneDrive\\Documents\\AI\\shared\\llm_engines\\nanovllm")
!2 = !{i32 2, !"Debug Info Version", i32 3}
!3 = !{i32 4, !"nvvm-reflect-ftz", i32 1}
!4 = distinct !DISubprogram(name: "_smoke_add_one_kernel", linkageName: "_smoke_add_one_kernel", scope: !1, file: !1, line: 37, type: !5, scopeLine: 37, spFlags: DISPFlagDefinition | DISPFlagOptimized, unit: !0)
!5 = !DISubroutineType(cc: DW_CC_normal, types: !6)
!6 = !{null, !7, !7, !9, !10, !10}
!7 = !DIDerivedType(tag: DW_TAG_pointer_type, name: "pointer", baseType: !8, size: 64, dwarfAddressSpace: 1)
!8 = !DIBasicType(name: "float", size: 32, encoding: DW_ATE_float)
!9 = !DIBasicType(name: "int", size: 32, encoding: DW_ATE_signed)
!10 = !DIDerivedType(tag: DW_TAG_pointer_type, name: "pointer", baseType: !11, size: 64, dwarfAddressSpace: 1)
!11 = !DIBasicType(name: "unknown_type", encoding: DW_ATE_signed)
!12 = !DILocation(line: 38, column: 19, scope: !4)
!13 = !DILocation(line: 39, column: 20, scope: !4)
!14 = !DILocation(line: 39, column: 34, scope: !4)
!15 = !DILocation(line: 40, column: 20, scope: !4)
!16 = !DILocation(line: 41, column: 25, scope: !4)
!17 = !DILocation(line: 41, column: 17, scope: !4)
!18 = !DILocation(line: 42, column: 22, scope: !4)
!19 = !DILocation(line: 42, column: 36, scope: !4)
!20 = !DILocation(line: 42, column: 13, scope: !4)
!21 = !DILocation(line: 37, column: 9, scope: !4)