Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Add critical convergent attributes to functions for OpenCL #805

Closed
seven-mile opened this issue Aug 24, 2024 · 1 comment · Fixed by #840
Closed

Add critical convergent attributes to functions for OpenCL #805

seven-mile opened this issue Aug 24, 2024 · 1 comment · Fixed by #840
Labels
invalid This doesn't seem right

Comments

@seven-mile
Copy link
Collaborator

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.

@seven-mile
Copy link
Collaborator Author

Note that call also requires these attributes.

Generating of attributes of Func or Call in ClangIR should be unified with mechanism like ConstructAttributeList. While for LLVM lowering, amending of LLVM attributes for Call op is not yet implemented.

Hugobros3 pushed a commit to shady-gang/clangir that referenced this issue Oct 2, 2024
…ages (llvm#840)

Fix llvm#805. This PR includes end-to-end implementation.

The `convergent` attribute is set depending on languages, which is
wrapped as `langOpts.assumeFunctionsAreConvergent()`.

Therefore, in ClangIR, every `cir.func` under `#cir.lang<opencl_c>` is
set to be convergent.

After lowering to LLVM IR, `PostOrderFunctionAttrs` pass will remove
unnecessary `convergent` then. In other words, we will still see
`convergent` on every function with `-O0`, but not with default
optimization level.

The test taken from `CodeGenOpenCL/convergent.cl` is a bit complicated.
However, the core of it is that `convergent` is set properly for
`convfun()` `non_convfun()` `f()` and `g()`. Merge of two `if` is more
or less a result of generating the same LLVM IR as OG.
smeenai pushed a commit to smeenai/clangir that referenced this issue Oct 9, 2024
…ages (llvm#840)

Fix llvm#805. This PR includes end-to-end implementation.

The `convergent` attribute is set depending on languages, which is
wrapped as `langOpts.assumeFunctionsAreConvergent()`.

Therefore, in ClangIR, every `cir.func` under `#cir.lang<opencl_c>` is
set to be convergent.

After lowering to LLVM IR, `PostOrderFunctionAttrs` pass will remove
unnecessary `convergent` then. In other words, we will still see
`convergent` on every function with `-O0`, but not with default
optimization level.

The test taken from `CodeGenOpenCL/convergent.cl` is a bit complicated.
However, the core of it is that `convergent` is set properly for
`convfun()` `non_convfun()` `f()` and `g()`. Merge of two `if` is more
or less a result of generating the same LLVM IR as OG.
smeenai pushed a commit to smeenai/clangir that referenced this issue Oct 9, 2024
…ages (llvm#840)

Fix llvm#805. This PR includes end-to-end implementation.

The `convergent` attribute is set depending on languages, which is
wrapped as `langOpts.assumeFunctionsAreConvergent()`.

Therefore, in ClangIR, every `cir.func` under `#cir.lang<opencl_c>` is
set to be convergent.

After lowering to LLVM IR, `PostOrderFunctionAttrs` pass will remove
unnecessary `convergent` then. In other words, we will still see
`convergent` on every function with `-O0`, but not with default
optimization level.

The test taken from `CodeGenOpenCL/convergent.cl` is a bit complicated.
However, the core of it is that `convergent` is set properly for
`convfun()` `non_convfun()` `f()` and `g()`. Merge of two `if` is more
or less a result of generating the same LLVM IR as OG.
keryell pushed a commit to keryell/clangir that referenced this issue Oct 19, 2024
…ages (llvm#840)

Fix llvm#805. This PR includes end-to-end implementation.

The `convergent` attribute is set depending on languages, which is
wrapped as `langOpts.assumeFunctionsAreConvergent()`.

Therefore, in ClangIR, every `cir.func` under `#cir.lang<opencl_c>` is
set to be convergent.

After lowering to LLVM IR, `PostOrderFunctionAttrs` pass will remove
unnecessary `convergent` then. In other words, we will still see
`convergent` on every function with `-O0`, but not with default
optimization level.

The test taken from `CodeGenOpenCL/convergent.cl` is a bit complicated.
However, the core of it is that `convergent` is set properly for
`convfun()` `non_convfun()` `f()` and `g()`. Merge of two `if` is more
or less a result of generating the same LLVM IR as OG.
lanza pushed a commit that referenced this issue Nov 5, 2024
…ages (#840)

Fix #805. This PR includes end-to-end implementation.

The `convergent` attribute is set depending on languages, which is
wrapped as `langOpts.assumeFunctionsAreConvergent()`.

Therefore, in ClangIR, every `cir.func` under `#cir.lang<opencl_c>` is
set to be convergent.

After lowering to LLVM IR, `PostOrderFunctionAttrs` pass will remove
unnecessary `convergent` then. In other words, we will still see
`convergent` on every function with `-O0`, but not with default
optimization level.

The test taken from `CodeGenOpenCL/convergent.cl` is a bit complicated.
However, the core of it is that `convergent` is set properly for
`convfun()` `non_convfun()` `f()` and `g()`. Merge of two `if` is more
or less a result of generating the same LLVM IR as OG.
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
invalid This doesn't seem right
Projects
None yet
Development

Successfully merging a pull request may close this issue.

1 participant