[AMDGPU] Add dynamic stack bit info to kernel-resource-usage Rpass output

[crobeck]

Summary

In code object 5 (https://llvm.org/docs/AMDGPUUsage.html#code-object-v5-metadata) the AMDGPU backend added the .uses_dynamic_stack bit to the kernel meta data to identity kernels which have compile time indeterminable stack usage (indirect function calls and recursion mainly). This patch adds this information to the output of the kernel-resource-usage remarks.

[JonChesterfield]

What is this bit used for? Codegen and debug info emission don't need it. I'm not sure what sets CurrentProgramInfo.DynamicCallStack at present either.

[crobeck]

In D156040#4525769, @JonChesterfield wrote:

What is this bit used for? Codegen and debug info emission don't need it. I'm not sure what sets CurrentProgramInfo.DynamicCallStack at present either.

It is part of the kernel meta data that is passed to the runtime to indicate the compiler can't determine the required stack amount use of the kernel and to tell the runtime it needs to check the value set from hipDeviceSetLimit.

DynamicCallStack ends up getting set here:

https://github.com/llvm/llvm-project/blob/main/llvm/lib/Target/AMDGPU/AMDGPUAsmPrinter.cpp#L706

Which is essentially just a check for recursion or indirect function calls/runtime polymorphism (i.e. virtual functions).

But the dynamic call stack information ends up being incorporated into CodeGen here:

https://github.com/llvm/llvm-project/blob/main/llvm/lib/Target/AMDGPU/AMDGPUResourceUsageAnalysis.cpp#L198

[arsenm]

llvm/lib/Target/AMDGPU/AMDGPUAsmPrinter.cpp

1308

No reason to make this conditional on the version

llvm/test/CodeGen/AMDGPU/resource-optimization-remarks_COV5.ll

1 ↗

(On Diff #543241)

Can just add to the existing remark test. Also, -filetype=null instead of -o /dev/null

[JonChesterfield]

In D156040#4525771, @crobeck wrote:

It is part of the kernel meta data that is passed to the runtime to indicate the compiler can't determine the required stack amount use of the kernel and to tell the runtime it needs to check the value set from hipDeviceSetLimit.

I don't see how this conveys any information. The compiler writes the stack size to be allocated. If it doesn't know what is sufficient, it's going to request some maximum and hope for the best.

The runtime allocates the requested size. If it has a bit saying "but use less if you know that's safe", then it can do nothing with that bit unless it has extra information. If it has that extra information, it doesn't need this bit.

Therefore instead of adding printing stuff related to this Boolean flag, we should delete the Boolean flag.

What's the use case I'm missing which makes this flag necessary/beneficial?

[crobeck]

In D156040#4526036, @JonChesterfield wrote:

In D156040#4525771, @crobeck wrote:

It is part of the kernel meta data that is passed to the runtime to indicate the compiler can't determine the required stack amount use of the kernel and to tell the runtime it needs to check the value set from hipDeviceSetLimit.

I don't see how this conveys any information. The compiler writes the stack size to be allocated. If it doesn't know what is sufficient, it's going to request some maximum and hope for the best.

The runtime allocates the requested size. If it has a bit saying "but use less if you know that's safe", then it can do nothing with that bit unless it has extra information. If it has that extra information, it doesn't need this bit.

Therefore instead of adding printing stuff related to this Boolean flag, we should delete the Boolean flag.

What's the use case I'm missing which makes this flag necessary/beneficial?

Post Code Object 5: If the compiler knows the required stack size of a kernel, it sets it and reports that to the runtime. The runtime then uses that value.

If the compiler does not know the required amount, in the case of indirect function calls or recursion, it sets the dynamic stack use bit and reports the minimum required by the kernel to the runtime (the actual amount required by the kernel could be significantly more than the minimum).

If the value required by the kernel ends up being more than the minimum calculated by the compiler or the runtime default value, when the dynamic stack use bit is set, then the code will crash. hipDeviceSetLimit must then be used to raise the stack allocated by the compiler/runtime.

Identifying those cases where the kernels use dynamic stack, and thus developers need to consider the value set by the hipDeviceSetLimit API, can actually be somewhat difficult if the kernel is buried under many layers of templates. We're already reporting remarks about compiler identified/reported scratch use and we've been asked if we can report this flag as part of these remarks.

I suppose reporting out the minimum the compiler calculates might also be useful here but don't feel strongly one way or the other about that.

[crobeck]

Remove conditional and consolidate new test into existing remarks test

[crobeck]

Update clang frontend test

[arsenm]

In D156040#4526036, @JonChesterfield wrote:

I don't see how this conveys any information. The compiler writes the stack size to be allocated. If it doesn't know what is sufficient, it's going to request some maximum and hope for the best.

That was the old broken workaround for the old bit that was never actually implemented in the runtime. The runtime now does properly respect some field to switch to interpreting the reported size as a minimum and then allocates the max of that minimum and some API provided size value

[arsenm]

clang/test/Frontend/amdgcn-machine-analysis-remarks.cl

22

Print right after the scratch size

llvm/lib/Target/AMDGPU/AMDGPUAsmPrinter.cpp

1308

don't need std::string for simple literals? StringRef?

llvm/test/CodeGen/AMDGPU/resource-optimization-remarks.ll

177

Maybe add another that has a static component too

[b-sumner]

In D156040#4528606, @arsenm wrote:

In D156040#4526036, @JonChesterfield wrote:

I don't see how this conveys any information. The compiler writes the stack size to be allocated. If it doesn't know what is sufficient, it's going to request some maximum and hope for the best.

That was the old broken workaround for the old bit that was never actually implemented in the runtime. The runtime now does properly respect some field to switch to interpreting the reported size as a minimum and then allocates the max of that minimum and some API provided size value

"runtime" == "hip/opencl runtime". The openmp runtime needs to do the same if not being done already.

[crobeck]

Change remark order and add static stack w/ indirect function call test case

[arsenm]

lgtm with nits

llvm/lib/Target/AMDGPU/AMDGPUAsmPrinter.cpp

1298

Drop "Uses". Could also just inline the bool->string

llvm/test/CodeGen/AMDGPU/resource-optimization-remarks.ll

189

No typed pointers

[arsenm]

In D156040#4526036, @JonChesterfield wrote:

What's the use case I'm missing which makes this flag necessary/beneficial?

The metadata is also irrelevant to this patch which is just reporting optimization hint information

[crobeck]

The openmp runtime needs to do the same if not being done already.

Should be on @carlo.bertolli's TODO list

clang/test/Frontend/amdgcn-machine-analysis-remarks.cl

 // REQUIRES: amdgpu-registered-target
 // RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -target-cpu gfx908 -Rpass-analysis=kernel-resource-usage -S -O0 -verify %s -o /dev/null
 
-// expected-remark@+9 {{Function Name: foo}}
-// expected-remark@+8 {{    SGPRs: 13}}
-// expected-remark@+7 {{    VGPRs: 10}}
-// expected-remark@+6 {{    AGPRs: 12}}
-// expected-remark@+5 {{    ScratchSize [bytes/lane]: 0}}
+// expected-remark@+10 {{Function Name: foo}}
+// expected-remark@+9 {{    SGPRs: 13}}
+// expected-remark@+8 {{    VGPRs: 10}}
+// expected-remark@+7 {{    AGPRs: 12}}
+// expected-remark@+6 {{    ScratchSize [bytes/lane]: 0}}
+// expected-remark@+5 {{    Dynamic Stack: False}}
 // expected-remark@+4 {{    Occupancy [waves/SIMD]: 10}}
 // expected-remark@+3 {{    SGPRs Spill: 0}}
 // expected-remark@+2 {{    VGPRs Spill: 0}}
 // expected-remark@+1 {{    LDS Size [bytes/block]: 0}}
 __kernel void foo() {
   __asm volatile ("; clobber s8" :::"s8");
   __asm volatile ("; clobber v9" :::"v9");
   __asm volatile ("; clobber a11" :::"a11");
 }

llvm/lib/Target/AMDGPU/AMDGPUAsmPrinter.cpp

void AMDGPUAsmPrinter::emitResourceUsageRemarks(
   EmitResourceUsageRemark("FunctionName", "Function Name",
                           MF.getFunction().getName());
   EmitResourceUsageRemark("NumSGPR", "SGPRs", CurrentProgramInfo.NumSGPR);
   EmitResourceUsageRemark("NumVGPR", "VGPRs", CurrentProgramInfo.NumArchVGPR);
   if (hasMAIInsts)
     EmitResourceUsageRemark("NumAGPR", "AGPRs", CurrentProgramInfo.NumAccVGPR);
   EmitResourceUsageRemark("ScratchSize", "ScratchSize [bytes/lane]",
                           CurrentProgramInfo.ScratchSize);
+  StringRef DynamicStackStr =
+      CurrentProgramInfo.DynamicCallStack ? "True" : "False";
+  EmitResourceUsageRemark("DynamicStack", "Dynamic Stack", DynamicStackStr);
   EmitResourceUsageRemark("Occupancy", "Occupancy [waves/SIMD]",
                           CurrentProgramInfo.Occupancy);
   EmitResourceUsageRemark("SGPRSpill", "SGPRs Spill",
                           CurrentProgramInfo.SGPRSpill);
   EmitResourceUsageRemark("VGPRSpill", "VGPRs Spill",
                           CurrentProgramInfo.VGPRSpill);
   if (isModuleEntryFunction)
     EmitResourceUsageRemark("BytesLDS", "LDS Size [bytes/block]",
                             CurrentProgramInfo.LDSSize);
 }

llvm/test/CodeGen/AMDGPU/resource-optimization-remarks.ll

-; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx908 -pass-remarks-output=%t -pass-remarks-analysis=kernel-resource-usage -filetype=obj -o /dev/null %s 2>&1 | FileCheck -check-prefix=STDERR %s
+; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx908 -pass-remarks-output=%t -pass-remarks-analysis=kernel-resource-usage -filetype=null %s 2>&1 | FileCheck -check-prefix=STDERR %s
 ; RUN: FileCheck -check-prefix=REMARK %s < %t
 
 ; STDERR: remark: foo.cl:27:0: Function Name: test_kernel
 ; STDERR-NEXT: remark: foo.cl:27:0:     SGPRs: 28
 ; STDERR-NEXT: remark: foo.cl:27:0:     VGPRs: 9
 ; STDERR-NEXT: remark: foo.cl:27:0:     AGPRs: 43
 ; STDERR-NEXT: remark: foo.cl:27:0:     ScratchSize [bytes/lane]: 0
+; STDERR-NEXT: remark: foo.cl:27:0:     Dynamic Stack: False
 ; STDERR-NEXT: remark: foo.cl:27:0:     Occupancy [waves/SIMD]: 5
 ; STDERR-NEXT: remark: foo.cl:27:0:     SGPRs Spill: 0
 ; STDERR-NEXT: remark: foo.cl:27:0:     VGPRs Spill: 0
 ; STDERR-NEXT: remark: foo.cl:27:0:     LDS Size [bytes/block]: 512
 
 ; REMARK-LABEL: --- !Analysis
 ; REMARK: Pass:            kernel-resource-usage
 ; REMARK-NEXT: Name:            FunctionName

 ; REMARK-NEXT: --- !Analysis
 ; REMARK-NEXT: Pass:            kernel-resource-usage
 ; REMARK-NEXT: Name:            ScratchSize
 ; REMARK-NEXT: DebugLoc:        { File: foo.cl, Line: 27, Column: 0 }
 ; REMARK-NEXT: Function:        test_kernel
 ; REMARK-NEXT: Args:
 ; REMARK-NEXT:   - String:          '    ScratchSize [bytes/lane]: '
 ; REMARK-NEXT:   - ScratchSize:     '0'
-; REMARK-NEXT: ...
+; REMARK-NEXT: ..
+; REMARK-NEXT: --- !Analysis
+; REMARK-NEXT: Pass:            kernel-resource-usage
+; REMARK-NEXT: Name:            DynamicStack
+; REMARK-NEXT: DebugLoc:        { File: foo.cl, Line: 27, Column: 0 }
+; REMARK-NEXT: Function:        test_kernel
+; REMARK-NEXT: Args:
+; REMARK-NEXT:   - String: ' Dynamic Stack: 
+; REMARK-NEXT:   - DynamicStack: 'False' 
+; REMARK-NEXT: ..
 ; REMARK-NEXT: --- !Analysis
 ; REMARK-NEXT: Pass:            kernel-resource-usage
 ; REMARK-NEXT: Name:            Occupancy
 ; REMARK-NEXT: DebugLoc:        { File: foo.cl, Line: 27, Column: 0 }
 ; REMARK-NEXT: Function:        test_kernel
 ; REMARK-NEXT: Args:
 ; REMARK-NEXT:   - String:          '    Occupancy [waves/SIMD]: '
 ; REMARK-NEXT:   - Occupancy:       '5'
define amdgpu_kernel void @test_kernel() !dbg !3 {
   ret void
 }
 
 ; STDERR: remark: foo.cl:42:0: Function Name: test_func
 ; STDERR-NEXT: remark: foo.cl:42:0:     SGPRs: 0
 ; STDERR-NEXT: remark: foo.cl:42:0:     VGPRs: 0
 ; STDERR-NEXT: remark: foo.cl:42:0:     AGPRs: 0
 ; STDERR-NEXT: remark: foo.cl:42:0:     ScratchSize [bytes/lane]: 0
+; STDERR-NEXT: remark: foo.cl:42:0:     Dynamic Stack: False
 ; STDERR-NEXT: remark: foo.cl:42:0:     Occupancy [waves/SIMD]: 0
 ; STDERR-NEXT: remark: foo.cl:42:0:     SGPRs Spill: 0
 ; STDERR-NEXT: remark: foo.cl:42:0:     VGPRs Spill: 0
 ; STDERR-NOT: LDS Size
 define void @test_func() !dbg !6 {
   call void asm sideeffect "; clobber v17", "~{v17}"()
   call void asm sideeffect "; clobber s11", "~{s11}"()
   call void asm sideeffect "; clobber a9", "~{a9}"()
   ret void
 }
 
 ; STDERR: remark: foo.cl:8:0: Function Name: empty_kernel
 ; STDERR-NEXT: remark: foo.cl:8:0:     SGPRs: 4
 ; STDERR-NEXT: remark: foo.cl:8:0:     VGPRs: 0
 ; STDERR-NEXT: remark: foo.cl:8:0:     AGPRs: 0
 ; STDERR-NEXT: remark: foo.cl:8:0:     ScratchSize [bytes/lane]: 0
+; STDERR-NEXT: remark: foo.cl:8:0:     Dynamic Stack: False
 ; STDERR-NEXT: remark: foo.cl:8:0:     Occupancy [waves/SIMD]: 8
 ; STDERR-NEXT: remark: foo.cl:8:0:     SGPRs Spill: 0
 ; STDERR-NEXT: remark: foo.cl:8:0:     VGPRs Spill: 0
 ; STDERR-NEXT: remark: foo.cl:8:0:     LDS Size [bytes/block]: 0
 define amdgpu_kernel void @empty_kernel() !dbg !7 {
   ret void
 }
 
 ; STDERR: remark: foo.cl:52:0: Function Name: empty_func
 ; STDERR-NEXT: remark: foo.cl:52:0:     SGPRs: 0
 ; STDERR-NEXT: remark: foo.cl:52:0:     VGPRs: 0
 ; STDERR-NEXT: remark: foo.cl:52:0:     AGPRs: 0
 ; STDERR-NEXT: remark: foo.cl:52:0:     ScratchSize [bytes/lane]: 0
+; STDERR-NEXT: remark: foo.cl:52:0:     Dynamic Stack: False
 ; STDERR-NEXT: remark: foo.cl:52:0:     Occupancy [waves/SIMD]: 0
 ; STDERR-NEXT: remark: foo.cl:52:0:     SGPRs Spill: 0
 ; STDERR-NEXT: remark: foo.cl:52:0:     VGPRs Spill: 0
 define void @empty_func() !dbg !8 {
   ret void
 }
 
+; STDERR: remark: foo.cl:64:0: Function Name: test_indirect_call
+; STDERR-NEXT: remark: foo.cl:64:0:     SGPRs: 39
+; STDERR-NEXT: remark: foo.cl:64:0:     VGPRs: 32
+; STDERR-NEXT: remark: foo.cl:64:0:     AGPRs: 10
+; STDERR-NEXT: remark: foo.cl:64:0:     ScratchSize [bytes/lane]: 0
+; STDERR-NEXT: remark: foo.cl:64:0:     Dynamic Stack: True
+; STDERR-NEXT: remark: foo.cl:64:0:     Occupancy [waves/SIMD]: 8
+; STDERR-NEXT: remark: foo.cl:64:0:     SGPRs Spill: 0
+; STDERR-NEXT: remark: foo.cl:64:0:     VGPRs Spill: 0
+; STDERR-NEXT: remark: foo.cl:64:0:     LDS Size [bytes/block]: 0
[e-mail]0 = external hidden unnamed_addr addrspace(4) constant ptr, align 4
+
+define amdgpu_kernel void @test_indirect_call() !dbg !9 {
+  %fptr = load ptr, ptr addrspace(4) @gv.fptr0
+  call void %fptr()
+  ret void
+}
+
+; STDERR: remark: foo.cl:74:0: Function Name: test_indirect_w_static_stack
+; STDERR-NEXT: remark: foo.cl:74:0:     SGPRs: 39
+; STDERR-NEXT: remark: foo.cl:74:0:     VGPRs: 32
+; STDERR-NEXT: remark: foo.cl:74:0:     AGPRs: 10
+; STDERR-NEXT: remark: foo.cl:74:0:     ScratchSize [bytes/lane]: 64
+; STDERR-NEXT: remark: foo.cl:74:0:     Dynamic Stack: True
+; STDERR-NEXT: remark: foo.cl:74:0:     Occupancy [waves/SIMD]: 8
+; STDERR-NEXT: remark: foo.cl:74:0:     SGPRs Spill: 0
+; STDERR-NEXT: remark: foo.cl:74:0:     VGPRs Spill: 0
+; STDERR-NEXT: remark: foo.cl:74:0:     LDS Size [bytes/block]: 0
+
+declare void @llvm.memset.p5.i64(ptr addrspace(5) nocapture readonly, i8, i64, i1 immarg)
+ 
+define amdgpu_kernel void @test_indirect_w_static_stack() !dbg !10 {
+  %alloca = alloca <10 x i64>, align 16, addrspace(5)
+  call void @llvm.memset.p5.i64(ptr addrspace(5) %alloca, i8 0, i64 40, i1 false)
+  %fptr = load ptr, ptr addrspace(4) @gv.fptr0
+  call void %fptr()
+  ret void
+}
+
 !llvm.dbg.cu = !{!0}
 !llvm.module.flags = !{!2}
+!llvm.module.flags = !{!11}
 
 !0 = distinct !DICompileUnit(language: DW_LANG_C99, file: !1, isOptimized: true, runtimeVersion: 0, emissionKind: FullDebug)
 !1 = !DIFile(filename: "foo.cl", directory: "/tmp")
 !2 = !{i32 2, !"Debug Info Version", i32 3}
 !3 = distinct !DISubprogram(name: "test_kernel", scope: !1, file: !1, type: !4, scopeLine: 27, flags: DIFlagPrototyped, spFlags: DISPFlagDefinition, unit: !0)
 !4 = !DISubroutineType(types: !5)
 !5 = !{null}
 !6 = distinct !DISubprogram(name: "test_func", scope: !1, file: !1, type: !4, scopeLine: 42, flags: DIFlagPrototyped, spFlags: DISPFlagDefinition, unit: !0)
 !7 = distinct !DISubprogram(name: "empty_kernel", scope: !1, file: !1, type: !4, scopeLine: 8, flags: DIFlagPrototyped, spFlags: DISPFlagDefinition, unit: !0)
 !8 = distinct !DISubprogram(name: "empty_func", scope: !1, file: !1, type: !4, scopeLine: 52, flags: DIFlagPrototyped, spFlags: DISPFlagDefinition, unit: !0)
+!9 = distinct !DISubprogram(name: "test_indirect_call", scope: !1, file: !1, type: !4, scopeLine: 64, flags: DIFlagPrototyped, spFlags: DISPFlagDefinition, unit: !0)
+!10 = distinct !DISubprogram(name: "test_indirect_w_static_stack", scope: !1, file: !1, type: !4, scopeLine: 74, flags: DIFlagPrototyped, spFlags: DISPFlagDefinition, unit: !0)
+!11 = !{i32 1, !"amdgpu_code_object_version", i32 500}