-
Notifications
You must be signed in to change notification settings - Fork 160
Description
ClangIR is still adding more function attributes sequentially. Most of them are not that important at this early stage. This is already tracked by in-code MissingFeatures::setFunctionAttributes()
.
But for OpenCL and other SIMT programs, the unit attribute convergent
for functions are especially critical. Missing of it would potentially cause some misoptimizations.
The good news is that the frontend does not need to do heavy analysis to determine the presence of convergent
. In OG CodeGen, the logic is basically to add convergent
to all functions for all SIMT languages. Then LLVM will utilize DivergenceAnalysis to remove unnecessary convergent
s.
clangir/clang/lib/CodeGen/CGCall.cpp
Lines 2012 to 2019 in 826abe4
if (LangOpts.assumeFunctionsAreConvergent()) { | |
// Conservatively, mark all functions and calls in CUDA and OpenCL as | |
// convergent (meaning, they may call an intrinsically convergent op, such | |
// as __syncthreads() / barrier(), and so can't have certain optimizations | |
// applied around them). LLVM will remove this attribute where it safely | |
// can. | |
FuncAttrs.addAttribute(llvm::Attribute::Convergent); | |
} |
And for non-SIMT languages, there should not be any functional changes.
The only challenge here is that the test case for OG OpenCL CodeGen: clang/test/CodeGenOpenCL/convergent.cl
is a bit beyond the current capabilities of ClangIR. (It may relies on #803 to avoid strange poison propagation.) I would recommend writing a trivial test by ourselves to exercise the code paths for ConvergentAttr.