This commit updates the lowering of all-reduce operations to annotate
the generated barriers with `memfence [#gpu.address_space<workgroup>]`
so that these barriers do not force unrelated global memory operations
to complete. It similarly sets up the warp synchronization function in
the vectory distribuhte tests, since they also only read/write shared
memory.
In additon, this commit adds convenience builders for gpu.barrier, which
will allow it to either fence on a given address space or on the address
space of a provided memref.
This is a takeover of PR ##110527
This commit adds an optional list of memory fences to gpu.barrier,
allowing users to specify which memory scopes they wish to fence
explicitly, while leaving the default semantics (which are equivalent to
calling for a global and local fence by analogy to CUDA's __syncthreads)
unchanged. The new expanded semantics are implemented for SPIR-V and for
the AMDGPU backend.
See also
https://discourse.llvm.org/t/rfc-add-memory-scope-to-gpu-barrier/81021/2?u=fmarno,
where the default behavior of a gpu.barrier was hashed out (though note
that the examples based on VMCNT are outdated for AMDGPU in that memory
fences can now be annotated with the correct set of address spaces).
This commit also deprecates amdgpu.lds_barrier for usecases that don't
involve targeting a gfx908.
Assisted-by: Cursor/Claude code (tests and extending amdgpu.lds_barrier
pattern while copying it over)
---------
Co-authored-by: Finlay Marno <finlay.marno@codeplay.com>
Co-authored-by: Jakub Kuderski <kubakuderski@gmail.com>
Co-authored-by: Alan Li <alan.li@me.com>
This patch fixes `-Wreturn-type` warnings which happens if MLIR is built
with GCC compiler (11.5 is used for detecting)
Founded errors
```
build/llvm-llvmorg-21.1.8/mlir/lib/CAPI/Transforms/Rewrite.cpp: In function ‘MlirGreedyRewriteStrictness mlirGreedyRewriteDriverConfigGetStrictness(MlirGreedyRewriteDriverConfig)’:
build/llvm-llvmorg-21.1.8/mlir/lib/CAPI/Transforms/Rewrite.cpp:399:1: warning: control reaches end of non-void function [-Wreturn-type]
399 | }
| ^
build/llvm-llvmorg-21.1.8/mlir/lib/CAPI/Transforms/Rewrite.cpp: In function ‘MlirGreedySimplifyRegionLevel mlirGreedyRewriteDriverConfigGetRegionSimplificationLevel(MlirGreedyRewriteDriverConfig)’:
build/llvm-llvmorg-21.1.8/mlir/lib/CAPI/Transforms/Rewrite.cpp:414:1: warning: control reaches end of non-void function [-Wreturn-type]
414 | }
| ^
build/llvm-llvmorg-21.1.8/mlir/lib/Dialect/GPU/IR/GPUDialect.cpp: In member function ‘mlir::Speculation::Speculatability mlir::gpu::SubgroupBroadcastOp::getSpeculatability()’:
build/llvm-llvmorg-21.1.8/mlir/lib/Dialect/GPU/IR/GPUDialect.cpp:2522:1: warning: control reaches end of non-void function [-Wreturn-type]
2522 | }
| ^
build/llvm-llvmorg-21.1.8/mlir/lib/Dialect/GPU/IR/GPUDialect.cpp: In member function ‘llvm::LogicalResult mlir::gpu::SubgroupBroadcastOp::verify()’:
build/llvm-llvmorg-21.1.8/mlir/lib/Dialect/GPU/IR/GPUDialect.cpp:2537:1: warning: control reaches end of non-void function [-Wreturn-type]
2537 | }
| ^
build/llvm-llvmorg-21.1.8/mlir/lib/Dialect/ArmNeon/Transforms/LowerContractToNeonPatterns.cpp: In member function ‘mlir::Value {anonymous}::VectorContractRewriter::createMMLA(mlir::PatternRewriter&, mlir::Location, mlir::Value, mlir::Value, mlir::Value)’:
build/llvm-llvmorg-21.1.8/mlir/lib/Dialect/ArmNeon/Transforms/LowerContractToNeonPatterns.cpp:153:3: warning: control reaches end of non-void function [-Wreturn-type]
153 | }
| ^
build/llvm-llvmorg-21.1.8/mlir/lib/Dialect/Linalg/IR/LinalgOps.cpp: In function ‘std::pair<long int, long int> mlir::linalg::getFmrFromWinogradConv2DFmr(mlir::linalg::WinogradConv2DFmr)’:
build/llvm-llvmorg-21.1.8/mlir/lib/Dialect/Linalg/IR/LinalgOps.cpp:3776:1: warning: control reaches end of non-void function [-Wreturn-type]
3776 | }
| ^
build/llvm-llvmorg-21.1.8/mlir/test/lib/Dialect/Test/TestOpDefs.cpp: In function ‘llvm::StringLiteral getVisibilityString(mlir::SymbolTable::Visibility)’:
build/llvm-llvmorg-21.1.8/mlir/test/lib/Dialect/Test/TestOpDefs.cpp:37:1: warning: control reaches end of non-void function [-Wreturn-type]
37 | }
| ^
```
The `gpu.launch` parser attempts to add a null `asyncTokenType` to the
results list if a result is requested but the `async` keyword is
missing, leading to an assertion failure.
Explicitly verify that `asyncTokenType` is valid when
`parser.getNumResults() > 0`. Emit a diagnostic error if the `async`
keyword is missing instead of crashing.
Add a regression test to `mlir/test/Dialect/GPU/invalid.mlir`.
Fix: https://github.com/llvm/llvm-project/issues/176530
Fix this build error, which is reported by some compilers after #175815:
```
error: operands to ?: have different types ‘mlir::Operation::result_range {aka mlir::ResultRange}’ and ‘mlir::ValueRange’
return successor.isParent() ? getOperation()->getResults() : ValueRange();
```
This commit simplifies the design of the `RegionBranchOpInterface`. The
property of being a successor input is now independent of the region
branch point.
There is a new API for querying successor inputs:
`RegionBranchOpInterface::getSuccessorInputs(RegionSuccessor)`. Note
that this function does **not** take a `RegionBranchPoint` as parameter.
The `RegionSuccessor` API is now also simpler: it no longer stores
successor inputs. A region successor is simply `Region *`, wrapped
around a convenience API.
Note: This commit is mostly mechanical. Analyses / transformations that
build on top of the `RegionBranchOpInterface` (e.g.,
`visitNonControlFlowArguments` API) can likely be simplified in
follow-up commits.
Note for LLVM integration: Split
`RegionBranchOpInterface::getSuccessorRegion` implementations into two
functions: `getSuccessorRegion` and `getSuccessorInputs. (There are many
examples in this commit.)
RFC:
https://discourse.llvm.org/t/rfc-simplify-regionbranchopinterface-separate-successor-inputs-from-region-successor/89420/7
Simplify the design of `RegionSuccessor`. There is no need to store the
`Operation *` pointer when branching out of the region branch op (to the
parent). There is no API to even access the `Operation *` pointer.
Add a new helper function `RegionSuccessor::parent` to construct a
region successor that points to the parent. This aligns the
`RegionSuccessor` design and API with `RegionBranchPoint`:
* Both classes now have a `parent()` helper function.
`ClassName::parent()` can be used in documentation to precisely describe
the source/target of a region branch.
* Both classes now use `nullptr` internally to represent "parent".
This API change also protects against incorrect API usage: users can no
longer pass an incorrect parent op. If a region successor is not a
region of the region branch op, it *must* branch out of region branch op
itself ("parent"). However, the previous API allowed passing other
operations. There was one such API violation in a [test
case](https://github.com/llvm/llvm-project/pull/174945/files#diff-d5717e4a8d7344b2ff77762b8fa480bcfec0eeee97a86195c787d791a6217e13L71).
Also clean up the documentation to use the correct terminology (such as
"successor operands", "successor inputs") consistently.
Note: This PR effectively rolls back some changes from #161575. That PR
introduced `llvm::PointerUnion<Region *, Operation *>
successor{nullptr};`. It is unclear from the commit message why that
change was made.
Note for LLVM integration: You may have to slightly modify
`getSuccessorRegion` implementations: Replace
`RegionSuccessor(getOperation(), getOperation()->getResults())` with
`RegionSuccessor::parent(getResults())`.
Introduces `VerificationUtils` to consolidate common operation
verification patterns in MLIR. This initial implementation provides
`verifyDynamicDimensionCount()` to reduce code duplication across
dialect verifiers.
This is an NFC (No Functional Change) refactoring that improves code
maintainability by extracting reusable verification logic into a shared
utility.
This lets us properly annotate ranges for gpu.cluster_block_id and
gpu.cluster_dim_blocks. It also allows us to fill in the
nvvm.cluster_dim attribute for use in the NVVM backend.
This PR re-lands #165873.
This PR extends the gpu.subgroup_mma_* ops to support fp64 type.
The extension requires special handling during the lowering to nvvm due
to the return type for load ops for fragment a and b (they return a
scalar instead of a struct).
The original PR did not guard the new test based on the required
architecture (sm80) which lead to a failure on the cuda runners with T4
GPUs.
This PR extends the `gpu.subgroup_mma_*` ops to support fp64 type.
The extension requires special handling during the lowering to `nvvm`
due to the return type for load ops for fragment a and b (they return a
scalar instead of a struct).
This is still somehow a WIP, we have some issues with this interface
that are not trivial to solve. This patch tries to make the concepts of
RegionBranchPoint and RegionSuccessor more robust and aligned with their
definition:
- A `RegionBranchPoint` is either the parent (`RegionBranchOpInterface`)
op or a `RegionBranchTerminatorOpInterface` operation in a nested
region.
- A `RegionSuccessor` is either one of the nested region or the parent
`RegionBranchOpInterface`
Some new methods with reasonnable default implementation are added to
help resolving the flow of values across the RegionBranchOpInterface.
It is still not trivial in the current state to walk the def-use chain
backward with this interface. For example when you have the 3rd block
argument in the entry block of a for-loop, finding the matching operands
requires to know about the hidden loop iterator block argument and where
the iterargs start. The API is designed around forward-tracking of the
chain unfortunately.
Try to reland #161575 ; I suspect a buildbot incremental build issue.
This is still somehow a WIP, we have some issues with this interface
that are not trivial to solve. This patch tries to make the concepts of
RegionBranchPoint and RegionSuccessor more robust and aligned with their
definition:
- A `RegionBranchPoint` is either the parent (`RegionBranchOpInterface`)
op or a `RegionBranchTerminatorOpInterface` operation in a nested
region.
- A `RegionSuccessor` is either one of the nested region or the parent
`RegionBranchOpInterface`
Some new methods with reasonnable default implementation are added to
help resolving the flow of values across the RegionBranchOpInterface.
It is still not trivial in the current state to walk the def-use chain
backward with this interface. For example when you have the 3rd block
argument in the entry block of a for-loop, finding the matching operands
requires to know about the hidden loop iterator block argument and where
the iterargs start. The API is designed around forward-tracking of the
chain unfortunately.
Use wrappers around `std::accumulate` to make the code more concise and
less bug-prone: https://github.com/llvm/llvm-project/pull/162129.
With `std::accumulate`, it's the initial value that determines the
accumulator type. `llvm::sum_of` and `llvm::product_of` pick the right
accumulator type based on the range element type.
Found some funny bugs like a local accumulate helper that calculated a
sum with initial value of 1 -- we didn't hit the bug because the code
was actually dead...
This was broken and never tested.
Not only this could crash for stack-use-after-scope, but it also would
have printed something like:
```
value <block argument> of type 'memref<7x8xf64, #gpu.address_space<workgroup>>' at index: 12
```
insted of the SSA value.
It turns out the gpu.func already have a very similar helper that we can
reuse here.
Fixes#161394
`subgroup_broadcast` allow to broadcast the value from one lane to all
lanes in subgroup.
Supported modes:
* `first_active_lane` - broadcast value from the first active lane in
subgroup.
* `specific_lane` - broadcast value from the specified lane, lane index
must be within subgroup.
* `any_lane` - if `src` value is uniform across all the subgroup lanes
return it unchanged, otherwise result is poison. This variant
essentially an uniformity hint for the compiler, conveying that specific
value is uniform across all subgroup lanes. Dropping `any_lane`
broadcast should not change the code semantics.
Adds a utility getter to `warp_execute_on_lane_0` which simplifies
access to the op's terminator.
Uses are refactored to utilize the new terminator getter.
`gpu::LaunchOp` is updated the following way:
- Change the attribute type of kernel function and module from
`SymbolRefAttr` to `FlatSymbolRefAttr` to avoid nested symbol
references.
- Rename variables from camel case (kernelFunc, kernelModule) to lower
case (function, module) and update the syntax.
- `LaunchOp::build` support passing `module` and `function` attributes.
`offset` and `width` must be constants and there are constraints on
their values. Update the operation definition to use attributes instead
of operands.
This revision adds DeviceMaskingAttrInterface and extends
DeviceMappingArrayAttr to accept a union of DeviceMappingAttrInterface
and DeviceMaskingAttrInterface.
Support is added to GPUTransformOps to take advantage of this
information and lower to block/warpgroup/warp/thread specialization when
mapped to linear ids.
The revision also connects to scf::ForallOp and uses the new attribute
to implement warp specialization.
The implementation is in the form of a GPUMappingMaskAttr, which can be
additionally passed to the scf.forall.mapping attribute to specify a
mask on compute resources that should be active.
In the first implementation the masking is a bitfield that specifies for
each processing unit whether it is active or not.
In the future, we may want to implement this as a symbol to refer to
dynamically defined values.
Extending op semantics with an operand is deemed too intrusive at this
time.
---------
Co-authored-by: Oleksandr "Alex" Zinenko <git@ozinenko.com>
This revision adds a new attribute for mapping `scf.forall` to linear
lane ids.
Example:
```
// %arg2 and %arg3 map to lanes [0, 6) and are turned into epxressions
// involving threadIdx.x/y by the map_nested_forall_to_threads
// transformation. This results in a if (linear_thread_id < 6) conditional.
scf.forall (%arg2, %arg3) in (2, 3) {
...
} {mapping = [#gpu.lane<linear_dim_0>, #gpu.lane<linear_dim_1>]}
```
---------
Co-authored-by: Oleksandr "Alex" Zinenko <git@ozinenko.com>
ArrayRef has a constructor that accepts std::nullopt. This
constructor dates back to the days when we still had llvm::Optional.
Since the use of std::nullopt outside the context of std::optional is
kind of abuse and not intuitive to new comers, I would like to move
away from the constructor and eventually remove it.
This patch migrates away from TypeRagne(std::nullopt) and
ValueRange(std::nullopt).
Motivation:
Currently, the NVVMOps are not verified against the supported SM
architectures. This can manifest as an ISel failure in the NVPTX LLVM
backend during CodeGen to PTX ISA. This PR addresses this issue by
adding verifier checks for Target-SM architectures in the NVVM Dialect
itself, thereby catching the errors early on.
Summary:
* Parametric traits named `NVVMRequiresSM` and `NVVMRequiresSMa` are
added to facilitate the version checks for typical and arch-accelerated
versions respectively.
* These traits can be attached to any NVVM Op to enable the checks for
the particular Op. (example shown below)
* An attribute interface called named `TargetAttrVerifyInterface` is
added to the GPU dialect which any target attribute seeking to perform
target-verification on the module can implement.
* The checks are performed by the `NVVMTargetAttr` (implementing the
`TargetAttrVerifyInterface` interface) when called from the GPU module
verifier where it walks through the module and performs the checks for
Ops with the `NVVMRequiresSM` traits.
* A few Ops in `NVVMOps.td` have been updated to serve as examples.
Example Usage:
```
def NVVM_ReduxOp : NVVM_Op<"redux.sync"> {...}
----> def NVVM_ReduxOp : NVVM_Op<"redux.sync", [NVVMRequiresSM<80>]> {...}
def NVVM_WgmmaFenceAlignedOp : NVVM_Op<"wgmma.fence.aligned"> {...}
----> def NVVM_WgmmaFenceAlignedOp : NVVM_Op<"wgmma.fence.aligned", [NVVMRequiresSMa<[90]>]> {...}
```
---------
Co-authored-by: Guray Ozen <guray.ozen@gmail.com>
This PR adds `cmd-options` to the `gpu-lower-to-nvvm-pipeline` pipeline
and the `nvvm-attach-target` pass, allowing users to pass flags to the
downstream compiler, *ptxas*.
Example:
```
mlir-opt -gpu-lower-to-nvvm-pipeline="cubin-chip=sm_80 ptxas-cmd-options='-v --register-usage-level=8'"
```
This patch shares core interface methods dealing with argument and
result attributes from CallableOpInterface with the CallOpInterface and
makes them mandatory to gives more consistent guarantees about concrete
operations using these interfaces.
This allows adding argument attributes on call like operations, which is
sometimes required to get proper ABI, like with llvm.call (and llvm.invoke).
The patch adds optional `arg_attrs` and `res_attrs` attributes to operations using
these interfaces that did not have that already.
They can then re-use the common "rich function signature"
printing/parsing helpers if they want (for the LLVM dialect, this is
done in the next patch).
Part of RFC: https://discourse.llvm.org/t/mlir-rfc-adding-argument-and-result-attributes-to-llvm-call/84107
The GPU ID operations already implement InferIntRangeInterface, which
gives constant lower and upper bounds on those IDs when appropriate
metadata is prentent on the operations or in the surrounding context.
This commit uses that existing code to implement the
ValueBoundsOpInterface, which is used when analyzing affine operations
(unlike the integer range interface, which is used for arithmetic
optimization).
It also implements the interface for gpu.launch, where we can use it to
express the constraint that block/grid sizes are equal to their value
from outside the launch op and that the corresponding IDs are bounded
above by that size.
As a consequence, the test pass for this inference is updated to work on
a FunctionOpInterface and not a func.func, creating minor churn in other
tests.
The `parseX()` functions that are defined to support `custom<X>` in
`assemblyFormat` should return `ParseResult` rather than
`LogicalResult`. The `ParseResult` type is necessary due to tablegen
generating code that expects this type within an Op `parseX()` function.
This change allows to expose through an interface attributes wrapping
content as external resources, and the usage inside the ModuleToObject
show how we will be able to provide runtime libraries without relying on
the filesystem.
This is a follow-up of #117246.
I thought then it would be easy to edit a DictionaryAttr but it turns
out that these attributes are immutable and need to be passed during the
construction of the gpu.binary Op.
The first commit was using the NVVMTargetAttr to pass the information.
After feedback from @fabianmcg, this PR now passes the information
through a new option of the gpu-module-to-binary pass.
Please add reviewers, as you see fit.
Here is the [merged
MR](https://github.com/llvm/llvm-project/pull/116007) which caused a
failure and [was
reverted](https://github.com/llvm/llvm-project/pull/116811).
Thanks to @joker-eph for the help, I fix it (miss constructing
`ModuleObject` with callback functions in
`mlir/lib/Target/LLVM/NVVM/Target.cpp`) and split unit tests from origin
test which don't need `ptxas` to make the test runs more widely.
Updates the return type of `getNumDynamicDims` and `getNumScalableDims`
from `int64_t` to `size_t`. This is for consistency with other
helpers/methods that return "size" and to reduce the number of
`static_cast`s in various places.
