[maleadt]

The AMDGPU back-end assumes that a call to a function without norecurse may recurse. It then gives the calling kernel a dynamic stack (.amdhsa_uses_dynamic_stack 1) and enables scratch memory, even when the call is never made. Julia marks @noinline functions only noinline, and our pipeline infers no norecurse, so every kernel with a non-inlined call pays for this. #1065 measured a 2x slowdown on an MI210 for a call on a branch that is never taken. The check is the same in every LLVM from 15 to main, and the AMDGPU attributor doesn't infer norecurse either.

Device code cannot call a kernel, so this PR marks kernel entry points norecurse, like clang does for CUDA and HIP __global__ functions. It then runs ReversePostOrderFunctionAttrsPass after the runtime is linked and the final inlining, to propagate the attribute to the internal functions the kernel calls. That is the top-down pass LLVM's default pipeline runs at the start of its optimization pipeline. The bottom-up function-attrs pass would only cover leaf functions, and doesn't infer norecurse on LLVM 15 when a callee calls intrinsics.

The native target is excluded: a kernel there is an ordinary function, which code outside the module could call back through a function pointer.

Effect on the other back-ends:

• NVPTX: no change. PTX can't express the attribute, and ptxas works out the stack from the call graph itself; ptxas resource usage and SASS were identical.
• SPIR-V: binaries are identical with both back-ends.
• Metal: the attribute reaches AIR, but on an M1 the kernel from Manually adding norecurse can significantly improve the performance on an AMD MI210 #1065 ran at the same speed with or without it, and a non-inlined call doesn't cost anything to begin with.

Fixes #1065, closes #262.

[codecov]

Codecov Report

✅ All modified and coverable lines are covered by tests.
✅ Project coverage is 92.42%. Comparing base (272a9bd) to head (930cbe9).
⚠️ Report is 1 commits behind head on main.

Additional details and impacted files

@@           Coverage Diff           @@
##             main    #1066   +/-   ##
=======================================
  Coverage   92.42%   92.42%           
=======================================
  Files          31       31           
  Lines        7323     7326    +3     
=======================================
+ Hits         6768     6771    +3     
  Misses        555      555           

☔ View full report in Codecov by Harness.
📢 Have feedback on the report? Share it here.

🚀 New features to boost your workflow:

• ❄️ Test Analytics: Detect flaky tests, report on failures, and find test suite problems.