Skip to content

Add critical convergent attributes to functions for OpenCL #805

@seven-mile

Description

@seven-mile

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 convergents.

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.

Metadata

Metadata

Assignees

No one assigned

    Labels

    invalidThis doesn't seem right

    Type

    No type

    Projects

    No projects

    Milestone

    No milestone

    Relationships

    None yet

    Development

    No branches or pull requests

    Issue actions