[mlir] replace 'emit_c_wrappers' func->llvm conversion option with a passThe 'emit_c_wrappers' option in the FuncToLLVM conversion requests C interfacewrappers to be emitted for every builtin func
[mlir] replace 'emit_c_wrappers' func->llvm conversion option with a passThe 'emit_c_wrappers' option in the FuncToLLVM conversion requests C interfacewrappers to be emitted for every builtin function in the module. While this hasbeen useful to bootstrap the interface, it is problematic in the longer term asit may unintentionally affect the functions that should retain their existinginterface, e.g., libm functions obtained by lowering math operations (seeD126964 for an example). Since D77314, we have a finer-grain control overinterface generation via an attribute that avoids the problem entirely. Removethe 'emit_c_wrappers' option. Introduce the '-llvm-request-c-wrappers' passthat can be run in any pipeline that needs blanket emission of functions toannotate all builtin functions with the attribute before performing the usuallowering that accounts for the attribute.Reviewed By: cheliniDifferential Revision: https://reviews.llvm.org/D127952
show more ...
[mlir][GPUToNVVM] Fix bug in mma elementwise loweringThe maxf implementation of wmma elementwise op was incorrect as theoperands of the select to check for Nan were swapped.Differential Revision
[mlir][GPUToNVVM] Fix bug in mma elementwise loweringThe maxf implementation of wmma elementwise op was incorrect as theoperands of the select to check for Nan were swapped.Differential Revision: https://reviews.llvm.org/D127879
[mlir][gpu] Move async copy ops to NVGPU and add caching hintsMove async copy operations to NVGPU as they only exist on NV target and aredesigned to match ptx semantic. This allows us to also add
[mlir][gpu] Move async copy ops to NVGPU and add caching hintsMove async copy operations to NVGPU as they only exist on NV target and aredesigned to match ptx semantic. This allows us to also add more fine graincaching hint attribute to the op.Add hint to bypass L1 and hook it up to NVVM op.Differential Revision: https://reviews.llvm.org/D125244
[mlir][nvgpu] Move mma.sync and ldmatrix in nvgpu dialectMove gpu operation mma.sync and ldmatrix in nvgpu as they are specificto nvidia target.Differential Revision: https://reviews.llvm.org/D1
[mlir][nvgpu] Move mma.sync and ldmatrix in nvgpu dialectMove gpu operation mma.sync and ldmatrix in nvgpu as they are specificto nvidia target.Differential Revision: https://reviews.llvm.org/D123824
[MLIR][GPU] Add GPU ops nvvm.mma.sync, nvvm.mma.ldmatrix, lane_idThis change adds three new operations to the GPU dialect: gpu.mma.sync,gpu.mma.ldmatrix, and gpu.lane_id. The former two are meant
[MLIR][GPU] Add GPU ops nvvm.mma.sync, nvvm.mma.ldmatrix, lane_idThis change adds three new operations to the GPU dialect: gpu.mma.sync,gpu.mma.ldmatrix, and gpu.lane_id. The former two are meant to targetthe lower level nvvm.mma.sync and nvvm.ldmatrix instructions, respectively.Lowerings are added for the new GPU operations for conversion toNVVM.Reviewed By: ThomasRaouxDifferential Revision: https://reviews.llvm.org/D123647
[mlir] Move the Builtin FuncOp to the Func dialectThis commit moves FuncOp out of the builtin dialect, and into the Funcdialect. This move has been planned in some capacity from the momentwe made
[mlir] Move the Builtin FuncOp to the Func dialectThis commit moves FuncOp out of the builtin dialect, and into the Funcdialect. This move has been planned in some capacity from the momentwe made FuncOp an operation (years ago). This commit handles thefunctional aspects of the move, but various aspects are left untouchedto ease migration: func::FuncOp is re-exported into mlir to reducethe actual API churn, the assembly format still accepts the unqualified`func`. These temporary measures will remain for a little while tosimplify migration before being removed.Differential Revision: https://reviews.llvm.org/D121266
[mlir] Rename the Standard dialect to the Func dialectThe last remaining operations in the standard dialect all revolve aroundFuncOp/function related constructs. This patch simply handles the init
[mlir] Rename the Standard dialect to the Func dialectThe last remaining operations in the standard dialect all revolve aroundFuncOp/function related constructs. This patch simply handles the initialrenaming (which by itself is already huge), but there are a large numberof cleanups unlocked/necessary afterwards:* Removing a bunch of unnecessary dependencies on Func* Cleaning up the From/ToStandard conversion passes* Preparing for the move of FuncOp to the Func dialectSee the discussion at https://discourse.llvm.org/t/standard-dialect-the-final-chapter/6061Differential Revision: https://reviews.llvm.org/D120624
[mlir][gpu] Add device side async copy operationsAdd new operations to the gpu dialect to represent device sideasynchronous copies. This also add the lowering of those operations tonvvm dialect.
[mlir][gpu] Add device side async copy operationsAdd new operations to the gpu dialect to represent device sideasynchronous copies. This also add the lowering of those operations tonvvm dialect.Those ops are meant to be low level and map directly to llvm dialectslike nvvm or rocdl.We can further add higher level of abstraction by building on top ofthose operations.This has been discuss here:https://discourse.llvm.org/t/modeling-gpu-async-copy-ampere-feature/4924Differential Revision: https://reviews.llvm.org/D119191
[mlir] Split out a new ControlFlow dialect from StandardThis dialect is intended to model lower level/branch based control-flow constructs. The initial setof operations are: AssertOp, BranchOp, Co
[mlir] Split out a new ControlFlow dialect from StandardThis dialect is intended to model lower level/branch based control-flow constructs. The initial setof operations are: AssertOp, BranchOp, CondBranchOp, SwitchOp; all split out from the currentstandard dialect.See https://discourse.llvm.org/t/standard-dialect-the-final-chapter/6061Differential Revision: https://reviews.llvm.org/D118966
Fix bugs in GPUToNVVM loweringThe current lowering from GPU to NVVM doesnot correctly handle the following cases whenlowering the gpu shuffle op.1. When the active width is set to 32 (all lanes
Fix bugs in GPUToNVVM loweringThe current lowering from GPU to NVVM doesnot correctly handle the following cases whenlowering the gpu shuffle op.1. When the active width is set to 32 (all lanes),then the current approach computes (1 << 32) -1 whichresults in poison values in the LLVM IR. We fix this bydefining the active mask as (-1) >> (32 - width).2. In the case of shuffle up, the computation of the thirdoperand c has to be different from the other 3 modes due tothe op definition in the ISA reference.(https://docs.nvidia.com/cuda/parallel-thread-execution/index.html)Specifically, the predicate value is computed as j >= maxLanefor up and j <= maxLane for all other modes. We fix this bycomputing maskAndClamp as 32 - width for this mode.TEST: We modify the existing test and add more checks for the up mode.Reviewed By: ThomasRaouxDifferential Revision: https://reviews.llvm.org/D118086
[mlir] Replace StrEnumAttr -> EnumAttr in core dialectsRemoves uses of `StrEnumAttr` in core dialectsReviewed By: mehdi_amini, rriddleDifferential Revision: https://reviews.llvm.org/D117514
[mlir][gpu] Extend shuffle op modes and add nvvm loweringAdd up, down and idx modes to gpu shuffle ops, also change the mode fromstring to enumDifferential Revision: https://reviews.llvm.org/D11
[mlir][gpu] Extend shuffle op modes and add nvvm loweringAdd up, down and idx modes to gpu shuffle ops, also change the mode fromstring to enumDifferential Revision: https://reviews.llvm.org/D114188
[mlir][nvvm] Remove special case ptr arithmetic lowering in gpu to nvvmUse existing helper instead of handling only a subset of indices loweringarithmetic. Also relax the restriction on the memref
[mlir][nvvm] Remove special case ptr arithmetic lowering in gpu to nvvmUse existing helper instead of handling only a subset of indices loweringarithmetic. Also relax the restriction on the memref rank for the GPU mma opsas we can now support any rank.Differential Revision: https://reviews.llvm.org/D113383
[mlir][gpu] Add basic support to do elementwise ops on mma matrix typeIn order to support fusion with mma matrix type we need to be able toexecute elementwise operations on them. This add an op to
[mlir][gpu] Add basic support to do elementwise ops on mma matrix typeIn order to support fusion with mma matrix type we need to be able toexecute elementwise operations on them. This add an op to be able tosupport some basic elementwise operations. This is a is not a fullsolution as it only supports a limited scope or operations. Ideally we wouldwant to be able to fuse with more kind of operations.Differential Revision: https://reviews.llvm.org/D112857
[mlir][nvvm] Generalize wmma ops to handle more types and shapeswmma intrinsics have a large number of combinations, ideally we want to be ableto target all the different variants. To avoid a comb
[mlir][nvvm] Generalize wmma ops to handle more types and shapeswmma intrinsics have a large number of combinations, ideally we want to be ableto target all the different variants. To avoid a combinatorial explosion in thenumber of mlir op we use attributes to represent the different variation ofload/store/mma ops. We also can generate with tablegen helpers to know whichcombinations are available. Using this we can avoid having too hardcode a pathfor specific shapes and can support more types.This patch also adds boiler plates for tf32 op support.Differential Revision: https://reviews.llvm.org/D112689
[mlir][GPUtoNVVM] Relax restriction on wmma op loweringAllow lowering of wmma ops with 64bits indexes. Change the defaultversion of the test to use default layout.Differential Revision: https://
[mlir][GPUtoNVVM] Relax restriction on wmma op loweringAllow lowering of wmma ops with 64bits indexes. Change the defaultversion of the test to use default layout.Differential Revision: https://reviews.llvm.org/D112479
[MLIR] Replace std ops with arith dialect opsPrecursor: https://reviews.llvm.org/D110200Removed redundant ops from the standard dialect that were moved to the`arith` or `math` dialects.Renamed
[MLIR] Replace std ops with arith dialect opsPrecursor: https://reviews.llvm.org/D110200Removed redundant ops from the standard dialect that were moved to the`arith` or `math` dialects.Renamed all instances of operations in the codebase and in tests.Reviewed By: rriddle, jpienaarDifferential Revision: https://reviews.llvm.org/D110797
Add a new interface allowing to set a default dialect to be used for printing/parsing regionsCurrently the builtin dialect is the default namespace used for parsingand printing. As such module and
Add a new interface allowing to set a default dialect to be used for printing/parsing regionsCurrently the builtin dialect is the default namespace used for parsingand printing. As such module and func don't need to be prefixed.In the case of some dialects that defines new regions for their ownpurpose (like SpirV modules for example), it can be beneficial tochange the default dialect in order to improve readability.Differential Revision: https://reviews.llvm.org/D107236
[mlir][nvvm]: Add math::Exp2Op lowering to NVVM.Differential Revision: https://reviews.llvm.org/D106050
[mlir] Fix wrong type in WmmaConstantOpToNVVMLoweringInsertElement takes a scalar integer attribute not an array of integer.Differential Revision: https://reviews.llvm.org/D105174
[mlir][gpu] Add op to create MMA constant matrixThis allow creating a matrix with all elements set to a given value. This isneeded to be able to implement a simple dot op.Differential Revision:
[mlir][gpu] Add op to create MMA constant matrixThis allow creating a matrix with all elements set to a given value. This isneeded to be able to implement a simple dot op.Differential Revision: https://reviews.llvm.org/D103870
[mlir][gpu] Relax restriction on MMA store op to allow chain of mma ops.In order to allow large matmul operations using the MMA ops we need to chainoperations this is not possible unless "DOp" and
[mlir][gpu] Relax restriction on MMA store op to allow chain of mma ops.In order to allow large matmul operations using the MMA ops we need to chainoperations this is not possible unless "DOp" and "COp" type have matchinglayout so remove the "DOp" layout and force accumulator and result type tomatch.Added a test for the case where the MMA value is accumulated.Differential Revision: https://reviews.llvm.org/D103023
[MLIR][GPU][NVVM] Add conversion of warp synchronous matrix-multiply accumulate GPU opsAdd conversion of warp synchronous matrix-multiplyaccumulate GPU opsAdd conversion of warp synchronous matri
[MLIR][GPU][NVVM] Add conversion of warp synchronous matrix-multiply accumulate GPU opsAdd conversion of warp synchronous matrix-multiplyaccumulate GPU opsAdd conversion of warp synchronous matrix-multiply accumulate GPU ops toNVVM ops. The following conversions are added :- 1.) subgroup_mma_load_matrix -> wmma.m16n16k16.load.[a,b,c]..row.stride 2.) subgroup_mma_store_matrix -> wmma.m16n16k16.store.d.[f16,f32].row.stride 3.) subgroup_mma_compute -> wmma.m16n16k16.mma.row.row.[f16,f32].[f16,f32]Reviewed By: bondhugula, ftynseDifferential Revision: https://reviews.llvm.org/D95331
Lower math.expm1 to intrinsics in the GPUToNVVM and GPUToROCDL conversions.This adds the lowering for expm1 for GPU backends.Differential Revision: https://reviews.llvm.org/D96756
[mlir] Use target-specific GPU kernel attributes in lowering pipelinesUntil now, the GPU translation to NVVM or ROCDL intrinsics relied on thepresence of the generic `gpu.kernel` attribute to atta
[mlir] Use target-specific GPU kernel attributes in lowering pipelinesUntil now, the GPU translation to NVVM or ROCDL intrinsics relied on thepresence of the generic `gpu.kernel` attribute to attach additional LLVM IRmetadata to the relevant functions. This would be problematic if each dialectwere to handle the conversion of its own options, which is the intendeddirection for the translation infrastructure. Introduce `nvvm.kernel` and`rocdl.kernel` in addition to `gpu.kernel` and base translation on these newattributes instead.Reviewed By: herhutDifferential Revision: https://reviews.llvm.org/D96591
123