|
|
|
source_filename = "LLVMDialectModule" |
|
|
|
define void @triton__0d1de(ptr addrspace(1) %0, i32 %1) local_unnamed_addr !dbg !5 { |
|
%3 = tail call i32 @llvm.nvvm.read.ptx.sreg.tid.x(), !dbg !8 |
|
%4 = shl i32 %3, 1, !dbg !8 |
|
%5 = and i32 %4, 510, !dbg !8 |
|
%6 = tail call i32 asm "mov.u32 $0, %ctaid.x;", "=r"() #1, !dbg !9 |
|
%7 = shl i32 %6, 9, !dbg !10 |
|
%8 = or i32 %7, %5, !dbg !11 |
|
%9 = icmp slt i32 %8, 12865792, !dbg !12 |
|
%10 = sext i32 %8 to i64, !dbg !13 |
|
%11 = getelementptr float, ptr addrspace(1) %0, i64 %10, !dbg !13 |
|
tail call void asm sideeffect "@$3 st.global.v2.b32 [ $2 + 0 ], { $0, $1 };", "r,r,l,b"(i32 0, i32 0, ptr addrspace(1) %11, i1 %9) #1, !dbg !14 |
|
ret void, !dbg !15 |
|
} |
|
|
|
|
|
declare noundef i32 @llvm.nvvm.read.ptx.sreg.tid.x() #0 |
|
|
|
attributes #0 = { mustprogress nocallback nofree nosync nounwind speculatable willreturn memory(none) } |
|
attributes #1 = { nounwind } |
|
|
|
!llvm.module.flags = !{!0} |
|
!llvm.dbg.cu = !{!1} |
|
!nvvm.annotations = !{!3, !4, !4, !3} |
|
|
|
!0 = !{i32 2, !"Debug Info Version", i32 3} |
|
!1 = distinct !DICompileUnit(language: DW_LANG_C, file: !2, producer: "triton", isOptimized: true, runtimeVersion: 0, emissionKind: FullDebug) |
|
!2 = !DIFile(filename: "c4yseldwmu3to52pbh2md2oeufrq3fcdmapkt4nxdzmyqtgd2ysp.py", directory: "/tmp/torchinductor_root/4y") |
|
!3 = !{ptr @triton__0d1de, !"kernel", i32 1} |
|
!4 = !{ptr @triton__0d1de, !"maxntidx", i32 256} |
|
!5 = distinct !DISubprogram(name: "triton__0d1de", linkageName: "triton__0d1de", scope: !2, file: !2, line: 18, type: !6, scopeLine: 18, spFlags: DISPFlagDefinition | DISPFlagOptimized, unit: !1) |
|
!6 = !DISubroutineType(cc: DW_CC_normal, types: !7) |
|
!7 = !{} |
|
!8 = !DILocation(line: 21, column: 36, scope: !5) |
|
!9 = !DILocation(line: 20, column: 28, scope: !5) |
|
!10 = !DILocation(line: 20, column: 33, scope: !5) |
|
!11 = !DILocation(line: 21, column: 23, scope: !5) |
|
!12 = !DILocation(line: 22, column: 21, scope: !5) |
|
!13 = !DILocation(line: 25, column: 25, scope: !5) |
|
!14 = !DILocation(line: 25, column: 36, scope: !5) |
|
!15 = !DILocation(line: 25, column: 4, scope: !5) |
|
|
