From 10c3b884eb79a8c2d1cb0d7a37c77227f04f39b3 Mon Sep 17 00:00:00 2001 From: Valentin Churavy Date: Tue, 16 Jan 2024 11:36:21 -0500 Subject: [PATCH 1/5] include the right dialects --- src/Dialects.jl | 19 +------------------ 1 file changed, 1 insertion(+), 18 deletions(-) diff --git a/src/Dialects.jl b/src/Dialects.jl index e3c1f179..b5a9f124 100644 --- a/src/Dialects.jl +++ b/src/Dialects.jl @@ -27,26 +27,9 @@ let You might need a newer version of MLIR.jl for this version of Julia.""") end - for path in readdir(joinpath(@__DIR__, "Dialects"); join=true) + for path in readdir(dir; join=true) include(path) end end -# module arith - -# module Predicates -# const eq = 0 -# const ne = 1 -# const slt = 2 -# const sle = 3 -# const sgt = 4 -# const sge = 5 -# const ult = 6 -# const ule = 7 -# const ugt = 8 -# const uge = 9 -# end - -# end # module arith - end # module Dialects From 3335ecf4f574f0812648f232689c4d7dc0967e2a Mon Sep 17 00:00:00 2001 From: Valentin Churavy Date: Tue, 16 Jan 2024 15:49:08 -0500 Subject: [PATCH 2/5] update bindings --- src/Dialects/14/GPU.jl | 2 +- src/Dialects/14/LLVMIR.jl | 164 +++++++++++++++---------------- src/Dialects/14/OpenACC.jl | 18 ++-- src/Dialects/14/OpenMP.jl | 4 +- src/Dialects/14/PDL.jl | 4 +- src/Dialects/14/X86Vector.jl | 52 +++++----- src/Dialects/15/GPU.jl | 2 +- src/Dialects/15/OpenACC.jl | 18 ++-- src/Dialects/15/OpenMP.jl | 8 +- src/Dialects/15/PDL.jl | 4 +- src/Dialects/15/X86Vector.jl | 52 +++++----- src/Dialects/16/Bufferization.jl | 2 +- src/Dialects/16/GPU.jl | 2 +- src/Dialects/16/NVGPU.jl | 4 +- src/Dialects/16/OpenACC.jl | 18 ++-- src/Dialects/16/OpenMP.jl | 16 +-- src/Dialects/16/PDL.jl | 4 +- src/Dialects/16/Tensor.jl | 2 +- src/Dialects/16/X86Vector.jl | 52 +++++----- 19 files changed, 214 insertions(+), 214 deletions(-) diff --git a/src/Dialects/14/GPU.jl b/src/Dialects/14/GPU.jl index dbaeaab2..0cc84420 100644 --- a/src/Dialects/14/GPU.jl +++ b/src/Dialects/14/GPU.jl @@ -460,7 +460,7 @@ module attributes {gpu.container_module} { } ``` """ -function launch_func(asyncDependencies::Vector{Value}, gridSizeX::Value, gridSizeY::Value, gridSizeZ::Value, blockSizeX::Value, blockSizeY::Value, blockSizeZ::Value, dynamicSharedMemorySize=nothing::Union{Nothing, Value}, operands::Vector{Value}; asyncToken=nothing::Union{Nothing, MLIRType}, kernel, location=Location()) +function launch_func(asyncDependencies::Vector{Value}, gridSizeX::Value, gridSizeY::Value, gridSizeZ::Value, blockSizeX::Value, blockSizeY::Value, blockSizeZ::Value, dynamicSharedMemorySize=nothing::Union{Nothing, Value}; operands::Vector{Value}, asyncToken=nothing::Union{Nothing, MLIRType}, kernel, location=Location()) results = MLIRType[] operands = Value[asyncDependencies..., gridSizeX, gridSizeY, gridSizeZ, blockSizeX, blockSizeY, blockSizeZ, operands..., ] owned_regions = Region[] diff --git a/src/Dialects/14/LLVMIR.jl b/src/Dialects/14/LLVMIR.jl index cbe21a08..dfae218a 100644 --- a/src/Dialects/14/LLVMIR.jl +++ b/src/Dialects/14/LLVMIR.jl @@ -497,10 +497,10 @@ function intr_copysign(a::Value, b::Value; res=nothing::Union{Nothing, MLIRType} end """ -`intr_coro.align` +`intr_coro_align` """ -function intr_coro.align(; res::MLIRType, location=Location()) +function intr_coro_align(; res::MLIRType, location=Location()) results = MLIRType[res, ] operands = Value[] owned_regions = Region[] @@ -516,10 +516,10 @@ function intr_coro.align(; res::MLIRType, location=Location()) end """ -`intr_coro.begin` +`intr_coro_begin` """ -function intr_coro.begin(token::Value, mem::Value; res::MLIRType, location=Location()) +function intr_coro_begin(token::Value, mem::Value; res::MLIRType, location=Location()) results = MLIRType[res, ] operands = Value[token, mem, ] owned_regions = Region[] @@ -535,10 +535,10 @@ function intr_coro.begin(token::Value, mem::Value; res::MLIRType, location=Locat end """ -`intr_coro.end` +`intr_coro_end` """ -function intr_coro.end(handle::Value, unwind::Value; res::MLIRType, location=Location()) +function intr_coro_end(handle::Value, unwind::Value; res::MLIRType, location=Location()) results = MLIRType[res, ] operands = Value[handle, unwind, ] owned_regions = Region[] @@ -554,10 +554,10 @@ function intr_coro.end(handle::Value, unwind::Value; res::MLIRType, location=Loc end """ -`intr_coro.free` +`intr_coro_free` """ -function intr_coro.free(id::Value, handle::Value; res::MLIRType, location=Location()) +function intr_coro_free(id::Value, handle::Value; res::MLIRType, location=Location()) results = MLIRType[res, ] operands = Value[id, handle, ] owned_regions = Region[] @@ -573,10 +573,10 @@ function intr_coro.free(id::Value, handle::Value; res::MLIRType, location=Locati end """ -`intr_coro.id` +`intr_coro_id` """ -function intr_coro.id(align::Value, promise::Value, coroaddr::Value, fnaddrs::Value; res::MLIRType, location=Location()) +function intr_coro_id(align::Value, promise::Value, coroaddr::Value, fnaddrs::Value; res::MLIRType, location=Location()) results = MLIRType[res, ] operands = Value[align, promise, coroaddr, fnaddrs, ] owned_regions = Region[] @@ -592,10 +592,10 @@ function intr_coro.id(align::Value, promise::Value, coroaddr::Value, fnaddrs::Va end """ -`intr_coro.resume` +`intr_coro_resume` """ -function intr_coro.resume(handle::Value; location=Location()) +function intr_coro_resume(handle::Value; location=Location()) results = MLIRType[] operands = Value[handle, ] owned_regions = Region[] @@ -611,10 +611,10 @@ function intr_coro.resume(handle::Value; location=Location()) end """ -`intr_coro.save` +`intr_coro_save` """ -function intr_coro.save(handle::Value; res::MLIRType, location=Location()) +function intr_coro_save(handle::Value; res::MLIRType, location=Location()) results = MLIRType[res, ] operands = Value[handle, ] owned_regions = Region[] @@ -630,10 +630,10 @@ function intr_coro.save(handle::Value; res::MLIRType, location=Location()) end """ -`intr_coro.size` +`intr_coro_size` """ -function intr_coro.size(; res::MLIRType, location=Location()) +function intr_coro_size(; res::MLIRType, location=Location()) results = MLIRType[res, ] operands = Value[] owned_regions = Region[] @@ -649,10 +649,10 @@ function intr_coro.size(; res::MLIRType, location=Location()) end """ -`intr_coro.suspend` +`intr_coro_suspend` """ -function intr_coro.suspend(save::Value, final::Value; res::MLIRType, location=Location()) +function intr_coro_suspend(save::Value, final::Value; res::MLIRType, location=Location()) results = MLIRType[res, ] operands = Value[save, final, ] owned_regions = Region[] @@ -746,10 +746,10 @@ function intr_ctpop(in::Value; res=nothing::Union{Nothing, MLIRType}, location=L end """ -`intr_eh.typeid.for` +`intr_eh_typeid_for` """ -function intr_eh.typeid.for(type_info::Value; res::MLIRType, location=Location()) +function intr_eh_typeid_for(type_info::Value; res::MLIRType, location=Location()) results = MLIRType[res, ] operands = Value[type_info, ] owned_regions = Region[] @@ -1223,10 +1223,10 @@ function getelementptr(base::Value, indices::Vector{Value}; res::MLIRType, struc end """ -`intr_get.active.lane.mask` +`intr_get_active_lane_mask` """ -function intr_get.active.lane.mask(base::Value, n::Value; res::MLIRType, location=Location()) +function intr_get_active_lane_mask(base::Value, n::Value; res::MLIRType, location=Location()) results = MLIRType[res, ] operands = Value[base, n, ] owned_regions = Region[] @@ -1732,10 +1732,10 @@ function intr_log(in::Value; res=nothing::Union{Nothing, MLIRType}, location=Loc end """ -`intr_masked.load` +`intr_masked_load` """ -function intr_masked.load(data::Value, mask::Value, pass_thru::Vector{Value}; res::MLIRType, alignment, location=Location()) +function intr_masked_load(data::Value, mask::Value, pass_thru::Vector{Value}; res::MLIRType, alignment, location=Location()) results = MLIRType[res, ] operands = Value[data, mask, pass_thru..., ] owned_regions = Region[] @@ -1751,10 +1751,10 @@ function intr_masked.load(data::Value, mask::Value, pass_thru::Vector{Value}; re end """ -`intr_masked.store` +`intr_masked_store` """ -function intr_masked.store(value::Value, data::Value, mask::Value; alignment, location=Location()) +function intr_masked_store(value::Value, data::Value, mask::Value; alignment, location=Location()) results = MLIRType[] operands = Value[value, data, mask, ] owned_regions = Region[] @@ -1770,10 +1770,10 @@ function intr_masked.store(value::Value, data::Value, mask::Value; alignment, lo end """ -`intr_matrix.column.major.load` +`intr_matrix_column_major_load` """ -function intr_matrix.column.major.load(data::Value, stride::Value; res::MLIRType, isVolatile, rows, columns, location=Location()) +function intr_matrix_column_major_load(data::Value, stride::Value; res::MLIRType, isVolatile, rows, columns, location=Location()) results = MLIRType[res, ] operands = Value[data, stride, ] owned_regions = Region[] @@ -1789,10 +1789,10 @@ function intr_matrix.column.major.load(data::Value, stride::Value; res::MLIRType end """ -`intr_matrix.column.major.store` +`intr_matrix_column_major_store` """ -function intr_matrix.column.major.store(matrix::Value, data::Value, stride::Value; isVolatile, rows, columns, location=Location()) +function intr_matrix_column_major_store(matrix::Value, data::Value, stride::Value; isVolatile, rows, columns, location=Location()) results = MLIRType[] operands = Value[matrix, data, stride, ] owned_regions = Region[] @@ -1808,10 +1808,10 @@ function intr_matrix.column.major.store(matrix::Value, data::Value, stride::Valu end """ -`intr_matrix.multiply` +`intr_matrix_multiply` """ -function intr_matrix.multiply(lhs::Value, rhs::Value; res::MLIRType, lhs_rows, lhs_columns, rhs_columns, location=Location()) +function intr_matrix_multiply(lhs::Value, rhs::Value; res::MLIRType, lhs_rows, lhs_columns, rhs_columns, location=Location()) results = MLIRType[res, ] operands = Value[lhs, rhs, ] owned_regions = Region[] @@ -1827,10 +1827,10 @@ function intr_matrix.multiply(lhs::Value, rhs::Value; res::MLIRType, lhs_rows, l end """ -`intr_matrix.transpose` +`intr_matrix_transpose` """ -function intr_matrix.transpose(matrix::Value; res::MLIRType, rows, columns, location=Location()) +function intr_matrix_transpose(matrix::Value; res::MLIRType, rows, columns, location=Location()) results = MLIRType[res, ] operands = Value[matrix, ] owned_regions = Region[] @@ -1886,10 +1886,10 @@ function intr_maximum(a::Value, b::Value; res=nothing::Union{Nothing, MLIRType}, end """ -`intr_memcpy.inline` +`intr_memcpy_inline` """ -function intr_memcpy.inline(dst::Value, src::Value, len::Value, isVolatile::Value; location=Location()) +function intr_memcpy_inline(dst::Value, src::Value, len::Value, isVolatile::Value; location=Location()) results = MLIRType[] operands = Value[dst, src, len, isVolatile, ] owned_regions = Region[] @@ -2217,10 +2217,10 @@ function return_(args::Vector{Value}; location=Location()) end """ -`intr_sadd.with.overflow` +`intr_sadd_with_overflow` """ -function intr_sadd.with.overflow(operand_0::Value, operand_1::Value; res::MLIRType, location=Location()) +function intr_sadd_with_overflow(operand_0::Value, operand_1::Value; res::MLIRType, location=Location()) results = MLIRType[res, ] operands = Value[operand_0, operand_1, ] owned_regions = Region[] @@ -2334,10 +2334,10 @@ function intr_smin(a::Value, b::Value; res=nothing::Union{Nothing, MLIRType}, lo end """ -`intr_smul.with.overflow` +`intr_smul_with_overflow` """ -function intr_smul.with.overflow(operand_0::Value, operand_1::Value; res::MLIRType, location=Location()) +function intr_smul_with_overflow(operand_0::Value, operand_1::Value; res::MLIRType, location=Location()) results = MLIRType[res, ] operands = Value[operand_0, operand_1, ] owned_regions = Region[] @@ -2373,10 +2373,10 @@ function srem(lhs::Value, rhs::Value; res=nothing::Union{Nothing, MLIRType}, loc end """ -`intr_ssub.with.overflow` +`intr_ssub_with_overflow` """ -function intr_ssub.with.overflow(operand_0::Value, operand_1::Value; res::MLIRType, location=Location()) +function intr_ssub_with_overflow(operand_0::Value, operand_1::Value; res::MLIRType, location=Location()) results = MLIRType[res, ] operands = Value[operand_0, operand_1, ] owned_regions = Region[] @@ -2615,10 +2615,10 @@ function trunc(arg::Value; res::MLIRType, location=Location()) end """ -`intr_uadd.with.overflow` +`intr_uadd_with_overflow` """ -function intr_uadd.with.overflow(operand_0::Value, operand_1::Value; res::MLIRType, location=Location()) +function intr_uadd_with_overflow(operand_0::Value, operand_1::Value; res::MLIRType, location=Location()) results = MLIRType[res, ] operands = Value[operand_0, operand_1, ] owned_regions = Region[] @@ -2713,10 +2713,10 @@ function intr_umin(a::Value, b::Value; res=nothing::Union{Nothing, MLIRType}, lo end """ -`intr_umul.with.overflow` +`intr_umul_with_overflow` """ -function intr_umul.with.overflow(operand_0::Value, operand_1::Value; res::MLIRType, location=Location()) +function intr_umul_with_overflow(operand_0::Value, operand_1::Value; res::MLIRType, location=Location()) results = MLIRType[res, ] operands = Value[operand_0, operand_1, ] owned_regions = Region[] @@ -2752,10 +2752,10 @@ function urem(lhs::Value, rhs::Value; res=nothing::Union{Nothing, MLIRType}, loc end """ -`intr_usub.with.overflow` +`intr_usub_with_overflow` """ -function intr_usub.with.overflow(operand_0::Value, operand_1::Value; res::MLIRType, location=Location()) +function intr_usub_with_overflow(operand_0::Value, operand_1::Value; res::MLIRType, location=Location()) results = MLIRType[res, ] operands = Value[operand_0, operand_1, ] owned_regions = Region[] @@ -2859,10 +2859,10 @@ function zext(arg::Value; res::MLIRType, location=Location()) end """ -`intr_masked.compressstore` +`intr_masked_compressstore` """ -function intr_masked.compressstore(operand_0::Value, operand_1::Value, operand_2::Value; location=Location()) +function intr_masked_compressstore(operand_0::Value, operand_1::Value, operand_2::Value; location=Location()) results = MLIRType[] operands = Value[operand_0, operand_1, operand_2, ] owned_regions = Region[] @@ -2878,10 +2878,10 @@ function intr_masked.compressstore(operand_0::Value, operand_1::Value, operand_2 end """ -`intr_masked.expandload` +`intr_masked_expandload` """ -function intr_masked.expandload(operand_0::Value, operand_1::Value, operand_2::Value; res::MLIRType, location=Location()) +function intr_masked_expandload(operand_0::Value, operand_1::Value, operand_2::Value; res::MLIRType, location=Location()) results = MLIRType[res, ] operands = Value[operand_0, operand_1, operand_2, ] owned_regions = Region[] @@ -2897,10 +2897,10 @@ function intr_masked.expandload(operand_0::Value, operand_1::Value, operand_2::V end """ -`intr_masked.gather` +`intr_masked_gather` """ -function intr_masked.gather(ptrs::Value, mask::Value, pass_thru::Vector{Value}; res::MLIRType, alignment, location=Location()) +function intr_masked_gather(ptrs::Value, mask::Value, pass_thru::Vector{Value}; res::MLIRType, alignment, location=Location()) results = MLIRType[res, ] operands = Value[ptrs, mask, pass_thru..., ] owned_regions = Region[] @@ -2916,10 +2916,10 @@ function intr_masked.gather(ptrs::Value, mask::Value, pass_thru::Vector{Value}; end """ -`intr_masked.scatter` +`intr_masked_scatter` """ -function intr_masked.scatter(value::Value, ptrs::Value, mask::Value; alignment, location=Location()) +function intr_masked_scatter(value::Value, ptrs::Value, mask::Value; alignment, location=Location()) results = MLIRType[] operands = Value[value, ptrs, mask, ] owned_regions = Region[] @@ -2935,10 +2935,10 @@ function intr_masked.scatter(value::Value, ptrs::Value, mask::Value; alignment, end """ -`intr_vector.reduce.add` +`intr_vector_reduce_add` """ -function intr_vector.reduce.add(operand_0::Value; res::MLIRType, location=Location()) +function intr_vector_reduce_add(operand_0::Value; res::MLIRType, location=Location()) results = MLIRType[res, ] operands = Value[operand_0, ] owned_regions = Region[] @@ -2954,10 +2954,10 @@ function intr_vector.reduce.add(operand_0::Value; res::MLIRType, location=Locati end """ -`intr_vector.reduce.and` +`intr_vector_reduce_and` """ -function intr_vector.reduce.and(operand_0::Value; res::MLIRType, location=Location()) +function intr_vector_reduce_and(operand_0::Value; res::MLIRType, location=Location()) results = MLIRType[res, ] operands = Value[operand_0, ] owned_regions = Region[] @@ -2973,10 +2973,10 @@ function intr_vector.reduce.and(operand_0::Value; res::MLIRType, location=Locati end """ -`intr_vector.reduce.fadd` +`intr_vector_reduce_fadd` """ -function intr_vector.reduce.fadd(operand_0::Value, operand_1::Value; res::MLIRType, reassoc=nothing, location=Location()) +function intr_vector_reduce_fadd(operand_0::Value, operand_1::Value; res::MLIRType, reassoc=nothing, location=Location()) results = MLIRType[res, ] operands = Value[operand_0, operand_1, ] owned_regions = Region[] @@ -2993,10 +2993,10 @@ function intr_vector.reduce.fadd(operand_0::Value, operand_1::Value; res::MLIRTy end """ -`intr_vector.reduce.fmax` +`intr_vector_reduce_fmax` """ -function intr_vector.reduce.fmax(operand_0::Value; res::MLIRType, location=Location()) +function intr_vector_reduce_fmax(operand_0::Value; res::MLIRType, location=Location()) results = MLIRType[res, ] operands = Value[operand_0, ] owned_regions = Region[] @@ -3012,10 +3012,10 @@ function intr_vector.reduce.fmax(operand_0::Value; res::MLIRType, location=Locat end """ -`intr_vector.reduce.fmin` +`intr_vector_reduce_fmin` """ -function intr_vector.reduce.fmin(operand_0::Value; res::MLIRType, location=Location()) +function intr_vector_reduce_fmin(operand_0::Value; res::MLIRType, location=Location()) results = MLIRType[res, ] operands = Value[operand_0, ] owned_regions = Region[] @@ -3031,10 +3031,10 @@ function intr_vector.reduce.fmin(operand_0::Value; res::MLIRType, location=Locat end """ -`intr_vector.reduce.fmul` +`intr_vector_reduce_fmul` """ -function intr_vector.reduce.fmul(operand_0::Value, operand_1::Value; res::MLIRType, reassoc=nothing, location=Location()) +function intr_vector_reduce_fmul(operand_0::Value, operand_1::Value; res::MLIRType, reassoc=nothing, location=Location()) results = MLIRType[res, ] operands = Value[operand_0, operand_1, ] owned_regions = Region[] @@ -3051,10 +3051,10 @@ function intr_vector.reduce.fmul(operand_0::Value, operand_1::Value; res::MLIRTy end """ -`intr_vector.reduce.mul` +`intr_vector_reduce_mul` """ -function intr_vector.reduce.mul(operand_0::Value; res::MLIRType, location=Location()) +function intr_vector_reduce_mul(operand_0::Value; res::MLIRType, location=Location()) results = MLIRType[res, ] operands = Value[operand_0, ] owned_regions = Region[] @@ -3070,10 +3070,10 @@ function intr_vector.reduce.mul(operand_0::Value; res::MLIRType, location=Locati end """ -`intr_vector.reduce.or` +`intr_vector_reduce_or` """ -function intr_vector.reduce.or(operand_0::Value; res::MLIRType, location=Location()) +function intr_vector_reduce_or(operand_0::Value; res::MLIRType, location=Location()) results = MLIRType[res, ] operands = Value[operand_0, ] owned_regions = Region[] @@ -3089,10 +3089,10 @@ function intr_vector.reduce.or(operand_0::Value; res::MLIRType, location=Locatio end """ -`intr_vector.reduce.smax` +`intr_vector_reduce_smax` """ -function intr_vector.reduce.smax(operand_0::Value; res::MLIRType, location=Location()) +function intr_vector_reduce_smax(operand_0::Value; res::MLIRType, location=Location()) results = MLIRType[res, ] operands = Value[operand_0, ] owned_regions = Region[] @@ -3108,10 +3108,10 @@ function intr_vector.reduce.smax(operand_0::Value; res::MLIRType, location=Locat end """ -`intr_vector.reduce.smin` +`intr_vector_reduce_smin` """ -function intr_vector.reduce.smin(operand_0::Value; res::MLIRType, location=Location()) +function intr_vector_reduce_smin(operand_0::Value; res::MLIRType, location=Location()) results = MLIRType[res, ] operands = Value[operand_0, ] owned_regions = Region[] @@ -3127,10 +3127,10 @@ function intr_vector.reduce.smin(operand_0::Value; res::MLIRType, location=Locat end """ -`intr_vector.reduce.umax` +`intr_vector_reduce_umax` """ -function intr_vector.reduce.umax(operand_0::Value; res::MLIRType, location=Location()) +function intr_vector_reduce_umax(operand_0::Value; res::MLIRType, location=Location()) results = MLIRType[res, ] operands = Value[operand_0, ] owned_regions = Region[] @@ -3146,10 +3146,10 @@ function intr_vector.reduce.umax(operand_0::Value; res::MLIRType, location=Locat end """ -`intr_vector.reduce.umin` +`intr_vector_reduce_umin` """ -function intr_vector.reduce.umin(operand_0::Value; res::MLIRType, location=Location()) +function intr_vector_reduce_umin(operand_0::Value; res::MLIRType, location=Location()) results = MLIRType[res, ] operands = Value[operand_0, ] owned_regions = Region[] @@ -3165,10 +3165,10 @@ function intr_vector.reduce.umin(operand_0::Value; res::MLIRType, location=Locat end """ -`intr_vector.reduce.xor` +`intr_vector_reduce_xor` """ -function intr_vector.reduce.xor(operand_0::Value; res::MLIRType, location=Location()) +function intr_vector_reduce_xor(operand_0::Value; res::MLIRType, location=Location()) results = MLIRType[res, ] operands = Value[operand_0, ] owned_regions = Region[] diff --git a/src/Dialects/14/OpenACC.jl b/src/Dialects/14/OpenACC.jl index e3fcffd6..099a545a 100644 --- a/src/Dialects/14/OpenACC.jl +++ b/src/Dialects/14/OpenACC.jl @@ -23,7 +23,7 @@ acc.data present(%a: memref<10x10xf32>, %b: memref<10x10xf32>, } ``` """ -function data(ifCond=nothing::Union{Nothing, Value}, copyOperands::Vector{Value}, copyinOperands::Vector{Value}, copyinReadonlyOperands::Vector{Value}, copyoutOperands::Vector{Value}, copyoutZeroOperands::Vector{Value}, createOperands::Vector{Value}, createZeroOperands::Vector{Value}, noCreateOperands::Vector{Value}, presentOperands::Vector{Value}, deviceptrOperands::Vector{Value}, attachOperands::Vector{Value}; defaultAttr=nothing, region::Region, location=Location()) +function data(ifCond=nothing::Union{Nothing, Value}; copyOperands::Vector{Value}, copyinOperands::Vector{Value}, copyinReadonlyOperands::Vector{Value}, copyoutOperands::Vector{Value}, copyoutZeroOperands::Vector{Value}, createOperands::Vector{Value}, createZeroOperands::Vector{Value}, noCreateOperands::Vector{Value}, presentOperands::Vector{Value}, deviceptrOperands::Vector{Value}, attachOperands::Vector{Value}, defaultAttr=nothing, region::Region, location=Location()) results = MLIRType[] operands = Value[copyOperands..., copyinOperands..., copyinReadonlyOperands..., copyoutOperands..., copyoutZeroOperands..., createOperands..., createZeroOperands..., noCreateOperands..., presentOperands..., deviceptrOperands..., attachOperands..., ] owned_regions = Region[region, ] @@ -52,7 +52,7 @@ The \"acc.enter_data\" operation represents the OpenACC enter data directive. acc.enter_data create(%d1 : memref<10xf32>) attributes {async} ``` """ -function enter_data(ifCond=nothing::Union{Nothing, Value}, asyncOperand=nothing::Union{Nothing, Value}, waitDevnum=nothing::Union{Nothing, Value}, waitOperands::Vector{Value}, copyinOperands::Vector{Value}, createOperands::Vector{Value}, createZeroOperands::Vector{Value}, attachOperands::Vector{Value}; async=nothing, wait=nothing, location=Location()) +function enter_data(ifCond=nothing::Union{Nothing, Value}; asyncOperand=nothing::Union{Nothing, Value}, waitDevnum=nothing::Union{Nothing, Value}, waitOperands::Vector{Value}, copyinOperands::Vector{Value}, createOperands::Vector{Value}, createZeroOperands::Vector{Value}, attachOperands::Vector{Value}, async=nothing, wait=nothing, location=Location()) results = MLIRType[] operands = Value[waitOperands..., copyinOperands..., createOperands..., createZeroOperands..., attachOperands..., ] owned_regions = Region[] @@ -84,7 +84,7 @@ The \"acc.exit_data\" operation represents the OpenACC exit data directive. acc.exit_data delete(%d1 : memref<10xf32>) attributes {async} ``` """ -function exit_data(ifCond=nothing::Union{Nothing, Value}, asyncOperand=nothing::Union{Nothing, Value}, waitDevnum=nothing::Union{Nothing, Value}, waitOperands::Vector{Value}, copyoutOperands::Vector{Value}, deleteOperands::Vector{Value}, detachOperands::Vector{Value}; async=nothing, wait=nothing, finalize=nothing, location=Location()) +function exit_data(ifCond=nothing::Union{Nothing, Value}; asyncOperand=nothing::Union{Nothing, Value}, waitDevnum=nothing::Union{Nothing, Value}, waitOperands::Vector{Value}, copyoutOperands::Vector{Value}, deleteOperands::Vector{Value}, detachOperands::Vector{Value}, async=nothing, wait=nothing, finalize=nothing, location=Location()) results = MLIRType[] operands = Value[waitOperands..., copyoutOperands..., deleteOperands..., detachOperands..., ] owned_regions = Region[] @@ -119,7 +119,7 @@ acc.init acc.init device_num(%dev1 : i32) ``` """ -function init(deviceTypeOperands::Vector{Value}, deviceNumOperand=nothing::Union{Nothing, Value}, ifCond=nothing::Union{Nothing, Value}; location=Location()) +function init(deviceTypeOperands::Vector{Value}, deviceNumOperand=nothing::Union{Nothing, Value}; ifCond=nothing::Union{Nothing, Value}, location=Location()) results = MLIRType[] operands = Value[deviceTypeOperands..., ] owned_regions = Region[] @@ -157,7 +157,7 @@ acc.loop gang vector { } attributes { collapse = 3 } ``` """ -function loop(gangNum=nothing::Union{Nothing, Value}, gangStatic=nothing::Union{Nothing, Value}, workerNum=nothing::Union{Nothing, Value}, vectorLength=nothing::Union{Nothing, Value}, tileOperands::Vector{Value}, privateOperands::Vector{Value}, reductionOperands::Vector{Value}; results::Vector{MLIRType}, collapse=nothing, seq=nothing, independent=nothing, auto_=nothing, reductionOp=nothing, exec_mapping=nothing, region::Region, location=Location()) +function loop(gangNum=nothing::Union{Nothing, Value}; gangStatic=nothing::Union{Nothing, Value}, workerNum=nothing::Union{Nothing, Value}, vectorLength=nothing::Union{Nothing, Value}, tileOperands::Vector{Value}, privateOperands::Vector{Value}, reductionOperands::Vector{Value}, results::Vector{MLIRType}, collapse=nothing, seq=nothing, independent=nothing, auto_=nothing, reductionOp=nothing, exec_mapping=nothing, region::Region, location=Location()) results = MLIRType[results..., ] operands = Value[tileOperands..., privateOperands..., reductionOperands..., ] owned_regions = Region[region, ] @@ -198,7 +198,7 @@ acc.parallel num_gangs(%c10) num_workers(%c10) } ``` """ -function parallel(async=nothing::Union{Nothing, Value}, waitOperands::Vector{Value}, numGangs=nothing::Union{Nothing, Value}, numWorkers=nothing::Union{Nothing, Value}, vectorLength=nothing::Union{Nothing, Value}, ifCond=nothing::Union{Nothing, Value}, selfCond=nothing::Union{Nothing, Value}, reductionOperands::Vector{Value}, copyOperands::Vector{Value}, copyinOperands::Vector{Value}, copyinReadonlyOperands::Vector{Value}, copyoutOperands::Vector{Value}, copyoutZeroOperands::Vector{Value}, createOperands::Vector{Value}, createZeroOperands::Vector{Value}, noCreateOperands::Vector{Value}, presentOperands::Vector{Value}, devicePtrOperands::Vector{Value}, attachOperands::Vector{Value}, gangPrivateOperands::Vector{Value}, gangFirstPrivateOperands::Vector{Value}; asyncAttr=nothing, waitAttr=nothing, selfAttr=nothing, reductionOp=nothing, defaultAttr=nothing, region::Region, location=Location()) +function parallel(async=nothing::Union{Nothing, Value}; waitOperands::Vector{Value}, numGangs=nothing::Union{Nothing, Value}, numWorkers=nothing::Union{Nothing, Value}, vectorLength=nothing::Union{Nothing, Value}, ifCond=nothing::Union{Nothing, Value}, selfCond=nothing::Union{Nothing, Value}, reductionOperands::Vector{Value}, copyOperands::Vector{Value}, copyinOperands::Vector{Value}, copyinReadonlyOperands::Vector{Value}, copyoutOperands::Vector{Value}, copyoutZeroOperands::Vector{Value}, createOperands::Vector{Value}, createZeroOperands::Vector{Value}, noCreateOperands::Vector{Value}, presentOperands::Vector{Value}, devicePtrOperands::Vector{Value}, attachOperands::Vector{Value}, gangPrivateOperands::Vector{Value}, gangFirstPrivateOperands::Vector{Value}, asyncAttr=nothing, waitAttr=nothing, selfAttr=nothing, reductionOp=nothing, defaultAttr=nothing, region::Region, location=Location()) results = MLIRType[] operands = Value[waitOperands..., reductionOperands..., copyOperands..., copyinOperands..., copyinReadonlyOperands..., copyoutOperands..., copyoutZeroOperands..., createOperands..., createZeroOperands..., noCreateOperands..., presentOperands..., devicePtrOperands..., attachOperands..., gangPrivateOperands..., gangFirstPrivateOperands..., ] owned_regions = Region[region, ] @@ -238,7 +238,7 @@ acc.shutdown acc.shutdown device_num(%dev1 : i32) ``` """ -function shutdown(deviceTypeOperands::Vector{Value}, deviceNumOperand=nothing::Union{Nothing, Value}, ifCond=nothing::Union{Nothing, Value}; location=Location()) +function shutdown(deviceTypeOperands::Vector{Value}, deviceNumOperand=nothing::Union{Nothing, Value}; ifCond=nothing::Union{Nothing, Value}, location=Location()) results = MLIRType[] operands = Value[deviceTypeOperands..., ] owned_regions = Region[] @@ -293,7 +293,7 @@ add to \$hostOperands. acc.update device(%d1 : memref<10xf32>) attributes {async} ``` """ -function update(ifCond=nothing::Union{Nothing, Value}, asyncOperand=nothing::Union{Nothing, Value}, waitDevnum=nothing::Union{Nothing, Value}, waitOperands::Vector{Value}, deviceTypeOperands::Vector{Value}, hostOperands::Vector{Value}, deviceOperands::Vector{Value}; async=nothing, wait=nothing, ifPresent=nothing, location=Location()) +function update(ifCond=nothing::Union{Nothing, Value}; asyncOperand=nothing::Union{Nothing, Value}, waitDevnum=nothing::Union{Nothing, Value}, waitOperands::Vector{Value}, deviceTypeOperands::Vector{Value}, hostOperands::Vector{Value}, deviceOperands::Vector{Value}, async=nothing, wait=nothing, ifPresent=nothing, location=Location()) results = MLIRType[] operands = Value[waitOperands..., deviceTypeOperands..., hostOperands..., deviceOperands..., ] owned_regions = Region[] @@ -328,7 +328,7 @@ acc.wait(%value1: index) acc.wait() async(%async1: i32) ``` """ -function wait(waitOperands::Vector{Value}, asyncOperand=nothing::Union{Nothing, Value}, waitDevnum=nothing::Union{Nothing, Value}, ifCond=nothing::Union{Nothing, Value}; async=nothing, location=Location()) +function wait(waitOperands::Vector{Value}, asyncOperand=nothing::Union{Nothing, Value}; waitDevnum=nothing::Union{Nothing, Value}, ifCond=nothing::Union{Nothing, Value}, async=nothing, location=Location()) results = MLIRType[] operands = Value[waitOperands..., ] owned_regions = Region[] diff --git a/src/Dialects/14/OpenMP.jl b/src/Dialects/14/OpenMP.jl index 163b53fb..a419db61 100644 --- a/src/Dialects/14/OpenMP.jl +++ b/src/Dialects/14/OpenMP.jl @@ -363,7 +363,7 @@ that specify the memory allocator to be used to obtain storage for private value The optional \$proc_bind_val attribute controls the thread affinity for the execution of the parallel region. """ -function parallel(if_expr_var=nothing::Union{Nothing, Value}, num_threads_var=nothing::Union{Nothing, Value}, private_vars::Vector{Value}, firstprivate_vars::Vector{Value}, shared_vars::Vector{Value}, copyin_vars::Vector{Value}, allocate_vars::Vector{Value}, allocators_vars::Vector{Value}; default_val=nothing, proc_bind_val=nothing, region::Region, location=Location()) +function parallel(if_expr_var=nothing::Union{Nothing, Value}; num_threads_var=nothing::Union{Nothing, Value}, private_vars::Vector{Value}, firstprivate_vars::Vector{Value}, shared_vars::Vector{Value}, copyin_vars::Vector{Value}, allocate_vars::Vector{Value}, allocators_vars::Vector{Value}, default_val=nothing, proc_bind_val=nothing, region::Region, location=Location()) results = MLIRType[] operands = Value[private_vars..., firstprivate_vars..., shared_vars..., copyin_vars..., allocate_vars..., allocators_vars..., ] owned_regions = Region[region, ] @@ -534,7 +534,7 @@ even if the target task is not yet completed. TODO: private, map, is_device_ptr, firstprivate, depend, defaultmap, in_reduction """ -function target(if_expr=nothing::Union{Nothing, Value}, device=nothing::Union{Nothing, Value}, thread_limit=nothing::Union{Nothing, Value}; nowait=nothing, region::Region, location=Location()) +function target(if_expr=nothing::Union{Nothing, Value}; device=nothing::Union{Nothing, Value}, thread_limit=nothing::Union{Nothing, Value}, nowait=nothing, region::Region, location=Location()) results = MLIRType[] operands = Value[] owned_regions = Region[region, ] diff --git a/src/Dialects/14/PDL.jl b/src/Dialects/14/PDL.jl index 20a1fc7d..71e56649 100644 --- a/src/Dialects/14/PDL.jl +++ b/src/Dialects/14/PDL.jl @@ -413,7 +413,7 @@ pdl.replace %root with (%vals : !pdl.range) pdl.replace %root with %otherOp ``` """ -function replace(operation::Value, replOperation=nothing::Union{Nothing, Value}, replValues::Vector{Value}; location=Location()) +function replace(operation::Value, replOperation=nothing::Union{Nothing, Value}; replValues::Vector{Value}, location=Location()) results = MLIRType[] operands = Value[operation, replValues..., ] owned_regions = Region[] @@ -552,7 +552,7 @@ pdl.rewrite { } ``` """ -function rewrite(root=nothing::Union{Nothing, Value}, externalArgs::Vector{Value}; name=nothing, externalConstParams=nothing, body::Region, location=Location()) +function rewrite(root=nothing::Union{Nothing, Value}; externalArgs::Vector{Value}, name=nothing, externalConstParams=nothing, body::Region, location=Location()) results = MLIRType[] operands = Value[externalArgs..., ] owned_regions = Region[body, ] diff --git a/src/Dialects/14/X86Vector.jl b/src/Dialects/14/X86Vector.jl index 35d1d80f..4f7e8fee 100644 --- a/src/Dialects/14/X86Vector.jl +++ b/src/Dialects/14/X86Vector.jl @@ -6,10 +6,10 @@ import ...API """ -`avx_intr.dp.ps.256` +`avx_intr_dp_ps_256` """ -function avx_intr.dp.ps.256(a::Value, b::Value, c::Value; res=nothing::Union{Nothing, MLIRType}, location=Location()) +function avx_intr_dp_ps_256(a::Value, b::Value, c::Value; res=nothing::Union{Nothing, MLIRType}, location=Location()) results = MLIRType[] operands = Value[a, b, c, ] owned_regions = Region[] @@ -26,7 +26,7 @@ function avx_intr.dp.ps.256(a::Value, b::Value, c::Value; res=nothing::Union{Not end """ -`avx_intr.dot` +`avx_intr_dot` Computes the 4-way dot products of the lower and higher parts of the source vectors and broadcasts the two results to the lower and higher elements of @@ -43,7 +43,7 @@ dot product of the two source vectors. %d = arith.addf %1, %2 : f32 ``` """ -function avx_intr.dot(a::Value, b::Value; res=nothing::Union{Nothing, MLIRType}, location=Location()) +function avx_intr_dot(a::Value, b::Value; res=nothing::Union{Nothing, MLIRType}, location=Location()) results = MLIRType[] operands = Value[a, b, ] owned_regions = Region[] @@ -60,10 +60,10 @@ function avx_intr.dot(a::Value, b::Value; res=nothing::Union{Nothing, MLIRType}, end """ -`avx512_intr.mask.compress` +`avx512_intr_mask_compress` """ -function avx512_intr.mask.compress(a::Value, src::Value, k::Value; res=nothing::Union{Nothing, MLIRType}, location=Location()) +function avx512_intr_mask_compress(a::Value, src::Value, k::Value; res=nothing::Union{Nothing, MLIRType}, location=Location()) results = MLIRType[] operands = Value[a, src, k, ] owned_regions = Region[] @@ -80,7 +80,7 @@ function avx512_intr.mask.compress(a::Value, src::Value, k::Value; res=nothing:: end """ -`avx512_mask.compress` +`avx512_mask_compress` The mask.compress op is an AVX512 specific op that can lower to the `llvm.mask.compress` instruction. Instead of `src`, a constant vector @@ -94,7 +94,7 @@ Contiguously store the active integer/floating-point elements in `a` (those with their respective bit set in writemask `k`) to `dst`, and pass through the remaining elements from `src`. """ -function avx512_mask.compress(k::Value, a::Value, src=nothing::Union{Nothing, Value}; dst=nothing::Union{Nothing, MLIRType}, constant_src=nothing, location=Location()) +function avx512_mask_compress(k::Value, a::Value, src=nothing::Union{Nothing, Value}; dst=nothing::Union{Nothing, MLIRType}, constant_src=nothing, location=Location()) results = MLIRType[] operands = Value[k, a, ] owned_regions = Region[] @@ -113,7 +113,7 @@ function avx512_mask.compress(k::Value, a::Value, src=nothing::Union{Nothing, Va end """ -`avx512_mask.rndscale` +`avx512_mask_rndscale` The mask.rndscale op is an AVX512 specific op that can lower to the proper LLVMAVX512 operation: `llvm.mask.rndscale.ps.512` or @@ -126,7 +126,7 @@ Round packed floating-point elements in `a` to the number of fraction bits specified by `imm`, and store the results in `dst` using writemask `k` (elements are copied from src when the corresponding mask bit is not set). """ -function avx512_mask.rndscale(src::Value, k::Value, a::Value, imm::Value, rounding::Value; dst=nothing::Union{Nothing, MLIRType}, location=Location()) +function avx512_mask_rndscale(src::Value, k::Value, a::Value, imm::Value, rounding::Value; dst=nothing::Union{Nothing, MLIRType}, location=Location()) results = MLIRType[] operands = Value[src, k, a, imm, rounding, ] owned_regions = Region[] @@ -143,10 +143,10 @@ function avx512_mask.rndscale(src::Value, k::Value, a::Value, imm::Value, roundi end """ -`avx512_intr.mask.rndscale.pd.512` +`avx512_intr_mask_rndscale_pd_512` """ -function avx512_intr.mask.rndscale.pd.512(src::Value, k::Value, a::Value, imm::Value, rounding::Value; res=nothing::Union{Nothing, MLIRType}, location=Location()) +function avx512_intr_mask_rndscale_pd_512(src::Value, k::Value, a::Value, imm::Value, rounding::Value; res=nothing::Union{Nothing, MLIRType}, location=Location()) results = MLIRType[] operands = Value[src, k, a, imm, rounding, ] owned_regions = Region[] @@ -163,10 +163,10 @@ function avx512_intr.mask.rndscale.pd.512(src::Value, k::Value, a::Value, imm::V end """ -`avx512_intr.mask.rndscale.ps.512` +`avx512_intr_mask_rndscale_ps_512` """ -function avx512_intr.mask.rndscale.ps.512(src::Value, k::Value, a::Value, imm::Value, rounding::Value; res=nothing::Union{Nothing, MLIRType}, location=Location()) +function avx512_intr_mask_rndscale_ps_512(src::Value, k::Value, a::Value, imm::Value, rounding::Value; res=nothing::Union{Nothing, MLIRType}, location=Location()) results = MLIRType[] operands = Value[src, k, a, imm, rounding, ] owned_regions = Region[] @@ -183,7 +183,7 @@ function avx512_intr.mask.rndscale.ps.512(src::Value, k::Value, a::Value, imm::V end """ -`avx512_mask.scalef` +`avx512_mask_scalef` The `mask.scalef` op is an AVX512 specific op that can lower to the proper LLVMAVX512 operation: `llvm.mask.scalef.ps.512` or @@ -196,7 +196,7 @@ Scale the packed floating-point elements in `a` using values from `b`, and store the results in `dst` using writemask `k` (elements are copied from src when the corresponding mask bit is not set). """ -function avx512_mask.scalef(src::Value, a::Value, b::Value, k::Value, rounding::Value; dst=nothing::Union{Nothing, MLIRType}, location=Location()) +function avx512_mask_scalef(src::Value, a::Value, b::Value, k::Value, rounding::Value; dst=nothing::Union{Nothing, MLIRType}, location=Location()) results = MLIRType[] operands = Value[src, a, b, k, rounding, ] owned_regions = Region[] @@ -213,10 +213,10 @@ function avx512_mask.scalef(src::Value, a::Value, b::Value, k::Value, rounding:: end """ -`avx512_intr.mask.scalef.pd.512` +`avx512_intr_mask_scalef_pd_512` """ -function avx512_intr.mask.scalef.pd.512(src::Value, a::Value, b::Value, k::Value, rounding::Value; res=nothing::Union{Nothing, MLIRType}, location=Location()) +function avx512_intr_mask_scalef_pd_512(src::Value, a::Value, b::Value, k::Value, rounding::Value; res=nothing::Union{Nothing, MLIRType}, location=Location()) results = MLIRType[] operands = Value[src, a, b, k, rounding, ] owned_regions = Region[] @@ -233,10 +233,10 @@ function avx512_intr.mask.scalef.pd.512(src::Value, a::Value, b::Value, k::Value end """ -`avx512_intr.mask.scalef.ps.512` +`avx512_intr_mask_scalef_ps_512` """ -function avx512_intr.mask.scalef.ps.512(src::Value, a::Value, b::Value, k::Value, rounding::Value; res=nothing::Union{Nothing, MLIRType}, location=Location()) +function avx512_intr_mask_scalef_ps_512(src::Value, a::Value, b::Value, k::Value, rounding::Value; res=nothing::Union{Nothing, MLIRType}, location=Location()) results = MLIRType[] operands = Value[src, a, b, k, rounding, ] owned_regions = Region[] @@ -253,10 +253,10 @@ function avx512_intr.mask.scalef.ps.512(src::Value, a::Value, b::Value, k::Value end """ -`avx_intr.rsqrt.ps.256` +`avx_intr_rsqrt_ps_256` """ -function avx_intr.rsqrt.ps.256(a::Value; res=nothing::Union{Nothing, MLIRType}, location=Location()) +function avx_intr_rsqrt_ps_256(a::Value; res=nothing::Union{Nothing, MLIRType}, location=Location()) results = MLIRType[] operands = Value[a, ] owned_regions = Region[] @@ -293,10 +293,10 @@ function avx_rsqrt(a::Value; b=nothing::Union{Nothing, MLIRType}, location=Locat end """ -`avx512_intr.vp2intersect.d.512` +`avx512_intr_vp2intersect_d_512` """ -function avx512_intr.vp2intersect.d.512(a::Value, b::Value; res::MLIRType, location=Location()) +function avx512_intr_vp2intersect_d_512(a::Value, b::Value; res::MLIRType, location=Location()) results = MLIRType[res, ] operands = Value[a, b, ] owned_regions = Region[] @@ -343,10 +343,10 @@ function avx512_vp2intersect(a::Value, b::Value; k1::MLIRType, k2::MLIRType, loc end """ -`avx512_intr.vp2intersect.q.512` +`avx512_intr_vp2intersect_q_512` """ -function avx512_intr.vp2intersect.q.512(a::Value, b::Value; res::MLIRType, location=Location()) +function avx512_intr_vp2intersect_q_512(a::Value, b::Value; res::MLIRType, location=Location()) results = MLIRType[res, ] operands = Value[a, b, ] owned_regions = Region[] diff --git a/src/Dialects/15/GPU.jl b/src/Dialects/15/GPU.jl index 947e3e64..ba13de41 100644 --- a/src/Dialects/15/GPU.jl +++ b/src/Dialects/15/GPU.jl @@ -515,7 +515,7 @@ module attributes {gpu.container_module} { } ``` """ -function launch_func(asyncDependencies::Vector{Value}, gridSizeX::Value, gridSizeY::Value, gridSizeZ::Value, blockSizeX::Value, blockSizeY::Value, blockSizeZ::Value, dynamicSharedMemorySize=nothing::Union{Nothing, Value}, operands::Vector{Value}; asyncToken=nothing::Union{Nothing, MLIRType}, kernel, location=Location()) +function launch_func(asyncDependencies::Vector{Value}, gridSizeX::Value, gridSizeY::Value, gridSizeZ::Value, blockSizeX::Value, blockSizeY::Value, blockSizeZ::Value, dynamicSharedMemorySize=nothing::Union{Nothing, Value}; operands::Vector{Value}, asyncToken=nothing::Union{Nothing, MLIRType}, kernel, location=Location()) results = MLIRType[] operands = Value[asyncDependencies..., gridSizeX, gridSizeY, gridSizeZ, blockSizeX, blockSizeY, blockSizeZ, operands..., ] owned_regions = Region[] diff --git a/src/Dialects/15/OpenACC.jl b/src/Dialects/15/OpenACC.jl index e3fcffd6..099a545a 100644 --- a/src/Dialects/15/OpenACC.jl +++ b/src/Dialects/15/OpenACC.jl @@ -23,7 +23,7 @@ acc.data present(%a: memref<10x10xf32>, %b: memref<10x10xf32>, } ``` """ -function data(ifCond=nothing::Union{Nothing, Value}, copyOperands::Vector{Value}, copyinOperands::Vector{Value}, copyinReadonlyOperands::Vector{Value}, copyoutOperands::Vector{Value}, copyoutZeroOperands::Vector{Value}, createOperands::Vector{Value}, createZeroOperands::Vector{Value}, noCreateOperands::Vector{Value}, presentOperands::Vector{Value}, deviceptrOperands::Vector{Value}, attachOperands::Vector{Value}; defaultAttr=nothing, region::Region, location=Location()) +function data(ifCond=nothing::Union{Nothing, Value}; copyOperands::Vector{Value}, copyinOperands::Vector{Value}, copyinReadonlyOperands::Vector{Value}, copyoutOperands::Vector{Value}, copyoutZeroOperands::Vector{Value}, createOperands::Vector{Value}, createZeroOperands::Vector{Value}, noCreateOperands::Vector{Value}, presentOperands::Vector{Value}, deviceptrOperands::Vector{Value}, attachOperands::Vector{Value}, defaultAttr=nothing, region::Region, location=Location()) results = MLIRType[] operands = Value[copyOperands..., copyinOperands..., copyinReadonlyOperands..., copyoutOperands..., copyoutZeroOperands..., createOperands..., createZeroOperands..., noCreateOperands..., presentOperands..., deviceptrOperands..., attachOperands..., ] owned_regions = Region[region, ] @@ -52,7 +52,7 @@ The \"acc.enter_data\" operation represents the OpenACC enter data directive. acc.enter_data create(%d1 : memref<10xf32>) attributes {async} ``` """ -function enter_data(ifCond=nothing::Union{Nothing, Value}, asyncOperand=nothing::Union{Nothing, Value}, waitDevnum=nothing::Union{Nothing, Value}, waitOperands::Vector{Value}, copyinOperands::Vector{Value}, createOperands::Vector{Value}, createZeroOperands::Vector{Value}, attachOperands::Vector{Value}; async=nothing, wait=nothing, location=Location()) +function enter_data(ifCond=nothing::Union{Nothing, Value}; asyncOperand=nothing::Union{Nothing, Value}, waitDevnum=nothing::Union{Nothing, Value}, waitOperands::Vector{Value}, copyinOperands::Vector{Value}, createOperands::Vector{Value}, createZeroOperands::Vector{Value}, attachOperands::Vector{Value}, async=nothing, wait=nothing, location=Location()) results = MLIRType[] operands = Value[waitOperands..., copyinOperands..., createOperands..., createZeroOperands..., attachOperands..., ] owned_regions = Region[] @@ -84,7 +84,7 @@ The \"acc.exit_data\" operation represents the OpenACC exit data directive. acc.exit_data delete(%d1 : memref<10xf32>) attributes {async} ``` """ -function exit_data(ifCond=nothing::Union{Nothing, Value}, asyncOperand=nothing::Union{Nothing, Value}, waitDevnum=nothing::Union{Nothing, Value}, waitOperands::Vector{Value}, copyoutOperands::Vector{Value}, deleteOperands::Vector{Value}, detachOperands::Vector{Value}; async=nothing, wait=nothing, finalize=nothing, location=Location()) +function exit_data(ifCond=nothing::Union{Nothing, Value}; asyncOperand=nothing::Union{Nothing, Value}, waitDevnum=nothing::Union{Nothing, Value}, waitOperands::Vector{Value}, copyoutOperands::Vector{Value}, deleteOperands::Vector{Value}, detachOperands::Vector{Value}, async=nothing, wait=nothing, finalize=nothing, location=Location()) results = MLIRType[] operands = Value[waitOperands..., copyoutOperands..., deleteOperands..., detachOperands..., ] owned_regions = Region[] @@ -119,7 +119,7 @@ acc.init acc.init device_num(%dev1 : i32) ``` """ -function init(deviceTypeOperands::Vector{Value}, deviceNumOperand=nothing::Union{Nothing, Value}, ifCond=nothing::Union{Nothing, Value}; location=Location()) +function init(deviceTypeOperands::Vector{Value}, deviceNumOperand=nothing::Union{Nothing, Value}; ifCond=nothing::Union{Nothing, Value}, location=Location()) results = MLIRType[] operands = Value[deviceTypeOperands..., ] owned_regions = Region[] @@ -157,7 +157,7 @@ acc.loop gang vector { } attributes { collapse = 3 } ``` """ -function loop(gangNum=nothing::Union{Nothing, Value}, gangStatic=nothing::Union{Nothing, Value}, workerNum=nothing::Union{Nothing, Value}, vectorLength=nothing::Union{Nothing, Value}, tileOperands::Vector{Value}, privateOperands::Vector{Value}, reductionOperands::Vector{Value}; results::Vector{MLIRType}, collapse=nothing, seq=nothing, independent=nothing, auto_=nothing, reductionOp=nothing, exec_mapping=nothing, region::Region, location=Location()) +function loop(gangNum=nothing::Union{Nothing, Value}; gangStatic=nothing::Union{Nothing, Value}, workerNum=nothing::Union{Nothing, Value}, vectorLength=nothing::Union{Nothing, Value}, tileOperands::Vector{Value}, privateOperands::Vector{Value}, reductionOperands::Vector{Value}, results::Vector{MLIRType}, collapse=nothing, seq=nothing, independent=nothing, auto_=nothing, reductionOp=nothing, exec_mapping=nothing, region::Region, location=Location()) results = MLIRType[results..., ] operands = Value[tileOperands..., privateOperands..., reductionOperands..., ] owned_regions = Region[region, ] @@ -198,7 +198,7 @@ acc.parallel num_gangs(%c10) num_workers(%c10) } ``` """ -function parallel(async=nothing::Union{Nothing, Value}, waitOperands::Vector{Value}, numGangs=nothing::Union{Nothing, Value}, numWorkers=nothing::Union{Nothing, Value}, vectorLength=nothing::Union{Nothing, Value}, ifCond=nothing::Union{Nothing, Value}, selfCond=nothing::Union{Nothing, Value}, reductionOperands::Vector{Value}, copyOperands::Vector{Value}, copyinOperands::Vector{Value}, copyinReadonlyOperands::Vector{Value}, copyoutOperands::Vector{Value}, copyoutZeroOperands::Vector{Value}, createOperands::Vector{Value}, createZeroOperands::Vector{Value}, noCreateOperands::Vector{Value}, presentOperands::Vector{Value}, devicePtrOperands::Vector{Value}, attachOperands::Vector{Value}, gangPrivateOperands::Vector{Value}, gangFirstPrivateOperands::Vector{Value}; asyncAttr=nothing, waitAttr=nothing, selfAttr=nothing, reductionOp=nothing, defaultAttr=nothing, region::Region, location=Location()) +function parallel(async=nothing::Union{Nothing, Value}; waitOperands::Vector{Value}, numGangs=nothing::Union{Nothing, Value}, numWorkers=nothing::Union{Nothing, Value}, vectorLength=nothing::Union{Nothing, Value}, ifCond=nothing::Union{Nothing, Value}, selfCond=nothing::Union{Nothing, Value}, reductionOperands::Vector{Value}, copyOperands::Vector{Value}, copyinOperands::Vector{Value}, copyinReadonlyOperands::Vector{Value}, copyoutOperands::Vector{Value}, copyoutZeroOperands::Vector{Value}, createOperands::Vector{Value}, createZeroOperands::Vector{Value}, noCreateOperands::Vector{Value}, presentOperands::Vector{Value}, devicePtrOperands::Vector{Value}, attachOperands::Vector{Value}, gangPrivateOperands::Vector{Value}, gangFirstPrivateOperands::Vector{Value}, asyncAttr=nothing, waitAttr=nothing, selfAttr=nothing, reductionOp=nothing, defaultAttr=nothing, region::Region, location=Location()) results = MLIRType[] operands = Value[waitOperands..., reductionOperands..., copyOperands..., copyinOperands..., copyinReadonlyOperands..., copyoutOperands..., copyoutZeroOperands..., createOperands..., createZeroOperands..., noCreateOperands..., presentOperands..., devicePtrOperands..., attachOperands..., gangPrivateOperands..., gangFirstPrivateOperands..., ] owned_regions = Region[region, ] @@ -238,7 +238,7 @@ acc.shutdown acc.shutdown device_num(%dev1 : i32) ``` """ -function shutdown(deviceTypeOperands::Vector{Value}, deviceNumOperand=nothing::Union{Nothing, Value}, ifCond=nothing::Union{Nothing, Value}; location=Location()) +function shutdown(deviceTypeOperands::Vector{Value}, deviceNumOperand=nothing::Union{Nothing, Value}; ifCond=nothing::Union{Nothing, Value}, location=Location()) results = MLIRType[] operands = Value[deviceTypeOperands..., ] owned_regions = Region[] @@ -293,7 +293,7 @@ add to \$hostOperands. acc.update device(%d1 : memref<10xf32>) attributes {async} ``` """ -function update(ifCond=nothing::Union{Nothing, Value}, asyncOperand=nothing::Union{Nothing, Value}, waitDevnum=nothing::Union{Nothing, Value}, waitOperands::Vector{Value}, deviceTypeOperands::Vector{Value}, hostOperands::Vector{Value}, deviceOperands::Vector{Value}; async=nothing, wait=nothing, ifPresent=nothing, location=Location()) +function update(ifCond=nothing::Union{Nothing, Value}; asyncOperand=nothing::Union{Nothing, Value}, waitDevnum=nothing::Union{Nothing, Value}, waitOperands::Vector{Value}, deviceTypeOperands::Vector{Value}, hostOperands::Vector{Value}, deviceOperands::Vector{Value}, async=nothing, wait=nothing, ifPresent=nothing, location=Location()) results = MLIRType[] operands = Value[waitOperands..., deviceTypeOperands..., hostOperands..., deviceOperands..., ] owned_regions = Region[] @@ -328,7 +328,7 @@ acc.wait(%value1: index) acc.wait() async(%async1: i32) ``` """ -function wait(waitOperands::Vector{Value}, asyncOperand=nothing::Union{Nothing, Value}, waitDevnum=nothing::Union{Nothing, Value}, ifCond=nothing::Union{Nothing, Value}; async=nothing, location=Location()) +function wait(waitOperands::Vector{Value}, asyncOperand=nothing::Union{Nothing, Value}; waitDevnum=nothing::Union{Nothing, Value}, ifCond=nothing::Union{Nothing, Value}, async=nothing, location=Location()) results = MLIRType[] operands = Value[waitOperands..., ] owned_regions = Region[] diff --git a/src/Dialects/15/OpenMP.jl b/src/Dialects/15/OpenMP.jl index 1139835c..4a6e3b04 100644 --- a/src/Dialects/15/OpenMP.jl +++ b/src/Dialects/15/OpenMP.jl @@ -416,7 +416,7 @@ threads complete. The optional \$proc_bind_val attribute controls the thread affinity for the execution of the parallel region. """ -function parallel(if_expr_var=nothing::Union{Nothing, Value}, num_threads_var=nothing::Union{Nothing, Value}, allocate_vars::Vector{Value}, allocators_vars::Vector{Value}, reduction_vars::Vector{Value}; reductions=nothing, proc_bind_val=nothing, region::Region, location=Location()) +function parallel(if_expr_var=nothing::Union{Nothing, Value}; num_threads_var=nothing::Union{Nothing, Value}, allocate_vars::Vector{Value}, allocators_vars::Vector{Value}, reduction_vars::Vector{Value}, reductions=nothing, proc_bind_val=nothing, region::Region, location=Location()) results = MLIRType[] operands = Value[allocate_vars..., allocators_vars..., reduction_vars..., ] owned_regions = Region[region, ] @@ -654,7 +654,7 @@ even if the target task is not yet completed. TODO: map, is_device_ptr, depend, defaultmap, in_reduction """ -function target(if_expr=nothing::Union{Nothing, Value}, device=nothing::Union{Nothing, Value}, thread_limit=nothing::Union{Nothing, Value}; nowait=nothing, region::Region, location=Location()) +function target(if_expr=nothing::Union{Nothing, Value}; device=nothing::Union{Nothing, Value}, thread_limit=nothing::Union{Nothing, Value}, nowait=nothing, region::Region, location=Location()) results = MLIRType[] operands = Value[] owned_regions = Region[region, ] @@ -813,7 +813,7 @@ construct. Thus, the taskloop construct creates an implicit taskgroup region. If the `nogroup` clause is present, no implicit taskgroup region is created. """ -function taskloop(lowerBound::Vector{Value}, upperBound::Vector{Value}, step::Vector{Value}, if_expr=nothing::Union{Nothing, Value}, final_expr=nothing::Union{Nothing, Value}, in_reduction_vars::Vector{Value}, reduction_vars::Vector{Value}, priority=nothing::Union{Nothing, Value}, allocate_vars::Vector{Value}, allocators_vars::Vector{Value}, grain_size=nothing::Union{Nothing, Value}, num_tasks=nothing::Union{Nothing, Value}; inclusive=nothing, untied=nothing, mergeable=nothing, in_reductions=nothing, reductions=nothing, nogroup=nothing, region::Region, location=Location()) +function taskloop(lowerBound::Vector{Value}, upperBound::Vector{Value}, step::Vector{Value}, if_expr=nothing::Union{Nothing, Value}; final_expr=nothing::Union{Nothing, Value}, in_reduction_vars::Vector{Value}, reduction_vars::Vector{Value}, priority=nothing::Union{Nothing, Value}, allocate_vars::Vector{Value}, allocators_vars::Vector{Value}, grain_size=nothing::Union{Nothing, Value}, num_tasks=nothing::Union{Nothing, Value}, inclusive=nothing, untied=nothing, mergeable=nothing, in_reductions=nothing, reductions=nothing, nogroup=nothing, region::Region, location=Location()) results = MLIRType[] operands = Value[lowerBound..., upperBound..., step..., in_reduction_vars..., reduction_vars..., allocate_vars..., allocators_vars..., ] owned_regions = Region[region, ] @@ -882,7 +882,7 @@ The `allocators_vars` and `allocate_vars` arguments are a variadic list of values that specify the memory allocator to be used to obtain storage for private values. """ -function task(if_expr=nothing::Union{Nothing, Value}, final_expr=nothing::Union{Nothing, Value}, in_reduction_vars::Vector{Value}, priority=nothing::Union{Nothing, Value}, allocate_vars::Vector{Value}, allocators_vars::Vector{Value}; untied=nothing, mergeable=nothing, in_reductions=nothing, region::Region, location=Location()) +function task(if_expr=nothing::Union{Nothing, Value}; final_expr=nothing::Union{Nothing, Value}, in_reduction_vars::Vector{Value}, priority=nothing::Union{Nothing, Value}, allocate_vars::Vector{Value}, allocators_vars::Vector{Value}, untied=nothing, mergeable=nothing, in_reductions=nothing, region::Region, location=Location()) results = MLIRType[] operands = Value[in_reduction_vars..., allocate_vars..., allocators_vars..., ] owned_regions = Region[region, ] diff --git a/src/Dialects/15/PDL.jl b/src/Dialects/15/PDL.jl index 736b61e6..51471728 100644 --- a/src/Dialects/15/PDL.jl +++ b/src/Dialects/15/PDL.jl @@ -400,7 +400,7 @@ pdl.replace %root with (%vals : !pdl.range) pdl.replace %root with %otherOp ``` """ -function replace(operation::Value, replOperation=nothing::Union{Nothing, Value}, replValues::Vector{Value}; location=Location()) +function replace(operation::Value, replOperation=nothing::Union{Nothing, Value}; replValues::Vector{Value}, location=Location()) results = MLIRType[] operands = Value[operation, replValues..., ] owned_regions = Region[] @@ -538,7 +538,7 @@ pdl.rewrite { } ``` """ -function rewrite(root=nothing::Union{Nothing, Value}, externalArgs::Vector{Value}; name=nothing, body::Region, location=Location()) +function rewrite(root=nothing::Union{Nothing, Value}; externalArgs::Vector{Value}, name=nothing, body::Region, location=Location()) results = MLIRType[] operands = Value[externalArgs..., ] owned_regions = Region[body, ] diff --git a/src/Dialects/15/X86Vector.jl b/src/Dialects/15/X86Vector.jl index 35d1d80f..4f7e8fee 100644 --- a/src/Dialects/15/X86Vector.jl +++ b/src/Dialects/15/X86Vector.jl @@ -6,10 +6,10 @@ import ...API """ -`avx_intr.dp.ps.256` +`avx_intr_dp_ps_256` """ -function avx_intr.dp.ps.256(a::Value, b::Value, c::Value; res=nothing::Union{Nothing, MLIRType}, location=Location()) +function avx_intr_dp_ps_256(a::Value, b::Value, c::Value; res=nothing::Union{Nothing, MLIRType}, location=Location()) results = MLIRType[] operands = Value[a, b, c, ] owned_regions = Region[] @@ -26,7 +26,7 @@ function avx_intr.dp.ps.256(a::Value, b::Value, c::Value; res=nothing::Union{Not end """ -`avx_intr.dot` +`avx_intr_dot` Computes the 4-way dot products of the lower and higher parts of the source vectors and broadcasts the two results to the lower and higher elements of @@ -43,7 +43,7 @@ dot product of the two source vectors. %d = arith.addf %1, %2 : f32 ``` """ -function avx_intr.dot(a::Value, b::Value; res=nothing::Union{Nothing, MLIRType}, location=Location()) +function avx_intr_dot(a::Value, b::Value; res=nothing::Union{Nothing, MLIRType}, location=Location()) results = MLIRType[] operands = Value[a, b, ] owned_regions = Region[] @@ -60,10 +60,10 @@ function avx_intr.dot(a::Value, b::Value; res=nothing::Union{Nothing, MLIRType}, end """ -`avx512_intr.mask.compress` +`avx512_intr_mask_compress` """ -function avx512_intr.mask.compress(a::Value, src::Value, k::Value; res=nothing::Union{Nothing, MLIRType}, location=Location()) +function avx512_intr_mask_compress(a::Value, src::Value, k::Value; res=nothing::Union{Nothing, MLIRType}, location=Location()) results = MLIRType[] operands = Value[a, src, k, ] owned_regions = Region[] @@ -80,7 +80,7 @@ function avx512_intr.mask.compress(a::Value, src::Value, k::Value; res=nothing:: end """ -`avx512_mask.compress` +`avx512_mask_compress` The mask.compress op is an AVX512 specific op that can lower to the `llvm.mask.compress` instruction. Instead of `src`, a constant vector @@ -94,7 +94,7 @@ Contiguously store the active integer/floating-point elements in `a` (those with their respective bit set in writemask `k`) to `dst`, and pass through the remaining elements from `src`. """ -function avx512_mask.compress(k::Value, a::Value, src=nothing::Union{Nothing, Value}; dst=nothing::Union{Nothing, MLIRType}, constant_src=nothing, location=Location()) +function avx512_mask_compress(k::Value, a::Value, src=nothing::Union{Nothing, Value}; dst=nothing::Union{Nothing, MLIRType}, constant_src=nothing, location=Location()) results = MLIRType[] operands = Value[k, a, ] owned_regions = Region[] @@ -113,7 +113,7 @@ function avx512_mask.compress(k::Value, a::Value, src=nothing::Union{Nothing, Va end """ -`avx512_mask.rndscale` +`avx512_mask_rndscale` The mask.rndscale op is an AVX512 specific op that can lower to the proper LLVMAVX512 operation: `llvm.mask.rndscale.ps.512` or @@ -126,7 +126,7 @@ Round packed floating-point elements in `a` to the number of fraction bits specified by `imm`, and store the results in `dst` using writemask `k` (elements are copied from src when the corresponding mask bit is not set). """ -function avx512_mask.rndscale(src::Value, k::Value, a::Value, imm::Value, rounding::Value; dst=nothing::Union{Nothing, MLIRType}, location=Location()) +function avx512_mask_rndscale(src::Value, k::Value, a::Value, imm::Value, rounding::Value; dst=nothing::Union{Nothing, MLIRType}, location=Location()) results = MLIRType[] operands = Value[src, k, a, imm, rounding, ] owned_regions = Region[] @@ -143,10 +143,10 @@ function avx512_mask.rndscale(src::Value, k::Value, a::Value, imm::Value, roundi end """ -`avx512_intr.mask.rndscale.pd.512` +`avx512_intr_mask_rndscale_pd_512` """ -function avx512_intr.mask.rndscale.pd.512(src::Value, k::Value, a::Value, imm::Value, rounding::Value; res=nothing::Union{Nothing, MLIRType}, location=Location()) +function avx512_intr_mask_rndscale_pd_512(src::Value, k::Value, a::Value, imm::Value, rounding::Value; res=nothing::Union{Nothing, MLIRType}, location=Location()) results = MLIRType[] operands = Value[src, k, a, imm, rounding, ] owned_regions = Region[] @@ -163,10 +163,10 @@ function avx512_intr.mask.rndscale.pd.512(src::Value, k::Value, a::Value, imm::V end """ -`avx512_intr.mask.rndscale.ps.512` +`avx512_intr_mask_rndscale_ps_512` """ -function avx512_intr.mask.rndscale.ps.512(src::Value, k::Value, a::Value, imm::Value, rounding::Value; res=nothing::Union{Nothing, MLIRType}, location=Location()) +function avx512_intr_mask_rndscale_ps_512(src::Value, k::Value, a::Value, imm::Value, rounding::Value; res=nothing::Union{Nothing, MLIRType}, location=Location()) results = MLIRType[] operands = Value[src, k, a, imm, rounding, ] owned_regions = Region[] @@ -183,7 +183,7 @@ function avx512_intr.mask.rndscale.ps.512(src::Value, k::Value, a::Value, imm::V end """ -`avx512_mask.scalef` +`avx512_mask_scalef` The `mask.scalef` op is an AVX512 specific op that can lower to the proper LLVMAVX512 operation: `llvm.mask.scalef.ps.512` or @@ -196,7 +196,7 @@ Scale the packed floating-point elements in `a` using values from `b`, and store the results in `dst` using writemask `k` (elements are copied from src when the corresponding mask bit is not set). """ -function avx512_mask.scalef(src::Value, a::Value, b::Value, k::Value, rounding::Value; dst=nothing::Union{Nothing, MLIRType}, location=Location()) +function avx512_mask_scalef(src::Value, a::Value, b::Value, k::Value, rounding::Value; dst=nothing::Union{Nothing, MLIRType}, location=Location()) results = MLIRType[] operands = Value[src, a, b, k, rounding, ] owned_regions = Region[] @@ -213,10 +213,10 @@ function avx512_mask.scalef(src::Value, a::Value, b::Value, k::Value, rounding:: end """ -`avx512_intr.mask.scalef.pd.512` +`avx512_intr_mask_scalef_pd_512` """ -function avx512_intr.mask.scalef.pd.512(src::Value, a::Value, b::Value, k::Value, rounding::Value; res=nothing::Union{Nothing, MLIRType}, location=Location()) +function avx512_intr_mask_scalef_pd_512(src::Value, a::Value, b::Value, k::Value, rounding::Value; res=nothing::Union{Nothing, MLIRType}, location=Location()) results = MLIRType[] operands = Value[src, a, b, k, rounding, ] owned_regions = Region[] @@ -233,10 +233,10 @@ function avx512_intr.mask.scalef.pd.512(src::Value, a::Value, b::Value, k::Value end """ -`avx512_intr.mask.scalef.ps.512` +`avx512_intr_mask_scalef_ps_512` """ -function avx512_intr.mask.scalef.ps.512(src::Value, a::Value, b::Value, k::Value, rounding::Value; res=nothing::Union{Nothing, MLIRType}, location=Location()) +function avx512_intr_mask_scalef_ps_512(src::Value, a::Value, b::Value, k::Value, rounding::Value; res=nothing::Union{Nothing, MLIRType}, location=Location()) results = MLIRType[] operands = Value[src, a, b, k, rounding, ] owned_regions = Region[] @@ -253,10 +253,10 @@ function avx512_intr.mask.scalef.ps.512(src::Value, a::Value, b::Value, k::Value end """ -`avx_intr.rsqrt.ps.256` +`avx_intr_rsqrt_ps_256` """ -function avx_intr.rsqrt.ps.256(a::Value; res=nothing::Union{Nothing, MLIRType}, location=Location()) +function avx_intr_rsqrt_ps_256(a::Value; res=nothing::Union{Nothing, MLIRType}, location=Location()) results = MLIRType[] operands = Value[a, ] owned_regions = Region[] @@ -293,10 +293,10 @@ function avx_rsqrt(a::Value; b=nothing::Union{Nothing, MLIRType}, location=Locat end """ -`avx512_intr.vp2intersect.d.512` +`avx512_intr_vp2intersect_d_512` """ -function avx512_intr.vp2intersect.d.512(a::Value, b::Value; res::MLIRType, location=Location()) +function avx512_intr_vp2intersect_d_512(a::Value, b::Value; res::MLIRType, location=Location()) results = MLIRType[res, ] operands = Value[a, b, ] owned_regions = Region[] @@ -343,10 +343,10 @@ function avx512_vp2intersect(a::Value, b::Value; k1::MLIRType, k2::MLIRType, loc end """ -`avx512_intr.vp2intersect.q.512` +`avx512_intr_vp2intersect_q_512` """ -function avx512_intr.vp2intersect.q.512(a::Value, b::Value; res::MLIRType, location=Location()) +function avx512_intr_vp2intersect_q_512(a::Value, b::Value; res::MLIRType, location=Location()) results = MLIRType[res, ] operands = Value[a, b, ] owned_regions = Region[] diff --git a/src/Dialects/16/Bufferization.jl b/src/Dialects/16/Bufferization.jl index 65cc0709..61236116 100644 --- a/src/Dialects/16/Bufferization.jl +++ b/src/Dialects/16/Bufferization.jl @@ -56,7 +56,7 @@ return %0 : tensor : tensor ``` """ -function alloc_tensor(dynamic_sizes::Vector{Value}, copy=nothing::Union{Nothing, Value}, size_hint=nothing::Union{Nothing, Value}; result::MLIRType, memory_space=nothing, location=Location()) +function alloc_tensor(dynamic_sizes::Vector{Value}, copy=nothing::Union{Nothing, Value}; size_hint=nothing::Union{Nothing, Value}, result::MLIRType, memory_space=nothing, location=Location()) results = MLIRType[result, ] operands = Value[dynamic_sizes..., ] owned_regions = Region[] diff --git a/src/Dialects/16/GPU.jl b/src/Dialects/16/GPU.jl index a893162e..87719577 100644 --- a/src/Dialects/16/GPU.jl +++ b/src/Dialects/16/GPU.jl @@ -530,7 +530,7 @@ module attributes {gpu.container_module} { } ``` """ -function launch_func(asyncDependencies::Vector{Value}, gridSizeX::Value, gridSizeY::Value, gridSizeZ::Value, blockSizeX::Value, blockSizeY::Value, blockSizeZ::Value, dynamicSharedMemorySize=nothing::Union{Nothing, Value}, kernelOperands::Vector{Value}; asyncToken=nothing::Union{Nothing, MLIRType}, kernel, location=Location()) +function launch_func(asyncDependencies::Vector{Value}, gridSizeX::Value, gridSizeY::Value, gridSizeZ::Value, blockSizeX::Value, blockSizeY::Value, blockSizeZ::Value, dynamicSharedMemorySize=nothing::Union{Nothing, Value}; kernelOperands::Vector{Value}, asyncToken=nothing::Union{Nothing, MLIRType}, kernel, location=Location()) results = MLIRType[] operands = Value[asyncDependencies..., gridSizeX, gridSizeY, gridSizeZ, blockSizeX, blockSizeY, blockSizeZ, kernelOperands..., ] owned_regions = Region[] diff --git a/src/Dialects/16/NVGPU.jl b/src/Dialects/16/NVGPU.jl index 9904c291..ab2ffaed 100644 --- a/src/Dialects/16/NVGPU.jl +++ b/src/Dialects/16/NVGPU.jl @@ -173,7 +173,7 @@ function ldmatrix(srcMemref::Value, indices::Vector{Value}; res::MLIRType, trans end """ -`mma_sp.sync` +`mma_sp_sync` The `nvgu.mma.sp.sync` operation performs a warp-distributed MMA operation where operand A is \"structured sparse\". In this case, the `matrixA` operand @@ -196,7 +196,7 @@ nvgpu.mma.sp.sync (%a, %b, %c) metadata (%meta) {mmaShape = [16, 8, 32]} : (vector<4x2xf16>, vector<2x2xf16>, vector<2x2xf16>) -> vector<2x2xf16> ``` """ -function mma_sp.sync(matrixA::Value, matrixB::Value, matrixC::Value, sparseMetadata::Value; res::MLIRType, mmaShape, sparsitySelector=nothing, tf32Enabled=nothing, location=Location()) +function mma_sp_sync(matrixA::Value, matrixB::Value, matrixC::Value, sparseMetadata::Value; res::MLIRType, mmaShape, sparsitySelector=nothing, tf32Enabled=nothing, location=Location()) results = MLIRType[res, ] operands = Value[matrixA, matrixB, matrixC, sparseMetadata, ] owned_regions = Region[] diff --git a/src/Dialects/16/OpenACC.jl b/src/Dialects/16/OpenACC.jl index e3fcffd6..099a545a 100644 --- a/src/Dialects/16/OpenACC.jl +++ b/src/Dialects/16/OpenACC.jl @@ -23,7 +23,7 @@ acc.data present(%a: memref<10x10xf32>, %b: memref<10x10xf32>, } ``` """ -function data(ifCond=nothing::Union{Nothing, Value}, copyOperands::Vector{Value}, copyinOperands::Vector{Value}, copyinReadonlyOperands::Vector{Value}, copyoutOperands::Vector{Value}, copyoutZeroOperands::Vector{Value}, createOperands::Vector{Value}, createZeroOperands::Vector{Value}, noCreateOperands::Vector{Value}, presentOperands::Vector{Value}, deviceptrOperands::Vector{Value}, attachOperands::Vector{Value}; defaultAttr=nothing, region::Region, location=Location()) +function data(ifCond=nothing::Union{Nothing, Value}; copyOperands::Vector{Value}, copyinOperands::Vector{Value}, copyinReadonlyOperands::Vector{Value}, copyoutOperands::Vector{Value}, copyoutZeroOperands::Vector{Value}, createOperands::Vector{Value}, createZeroOperands::Vector{Value}, noCreateOperands::Vector{Value}, presentOperands::Vector{Value}, deviceptrOperands::Vector{Value}, attachOperands::Vector{Value}, defaultAttr=nothing, region::Region, location=Location()) results = MLIRType[] operands = Value[copyOperands..., copyinOperands..., copyinReadonlyOperands..., copyoutOperands..., copyoutZeroOperands..., createOperands..., createZeroOperands..., noCreateOperands..., presentOperands..., deviceptrOperands..., attachOperands..., ] owned_regions = Region[region, ] @@ -52,7 +52,7 @@ The \"acc.enter_data\" operation represents the OpenACC enter data directive. acc.enter_data create(%d1 : memref<10xf32>) attributes {async} ``` """ -function enter_data(ifCond=nothing::Union{Nothing, Value}, asyncOperand=nothing::Union{Nothing, Value}, waitDevnum=nothing::Union{Nothing, Value}, waitOperands::Vector{Value}, copyinOperands::Vector{Value}, createOperands::Vector{Value}, createZeroOperands::Vector{Value}, attachOperands::Vector{Value}; async=nothing, wait=nothing, location=Location()) +function enter_data(ifCond=nothing::Union{Nothing, Value}; asyncOperand=nothing::Union{Nothing, Value}, waitDevnum=nothing::Union{Nothing, Value}, waitOperands::Vector{Value}, copyinOperands::Vector{Value}, createOperands::Vector{Value}, createZeroOperands::Vector{Value}, attachOperands::Vector{Value}, async=nothing, wait=nothing, location=Location()) results = MLIRType[] operands = Value[waitOperands..., copyinOperands..., createOperands..., createZeroOperands..., attachOperands..., ] owned_regions = Region[] @@ -84,7 +84,7 @@ The \"acc.exit_data\" operation represents the OpenACC exit data directive. acc.exit_data delete(%d1 : memref<10xf32>) attributes {async} ``` """ -function exit_data(ifCond=nothing::Union{Nothing, Value}, asyncOperand=nothing::Union{Nothing, Value}, waitDevnum=nothing::Union{Nothing, Value}, waitOperands::Vector{Value}, copyoutOperands::Vector{Value}, deleteOperands::Vector{Value}, detachOperands::Vector{Value}; async=nothing, wait=nothing, finalize=nothing, location=Location()) +function exit_data(ifCond=nothing::Union{Nothing, Value}; asyncOperand=nothing::Union{Nothing, Value}, waitDevnum=nothing::Union{Nothing, Value}, waitOperands::Vector{Value}, copyoutOperands::Vector{Value}, deleteOperands::Vector{Value}, detachOperands::Vector{Value}, async=nothing, wait=nothing, finalize=nothing, location=Location()) results = MLIRType[] operands = Value[waitOperands..., copyoutOperands..., deleteOperands..., detachOperands..., ] owned_regions = Region[] @@ -119,7 +119,7 @@ acc.init acc.init device_num(%dev1 : i32) ``` """ -function init(deviceTypeOperands::Vector{Value}, deviceNumOperand=nothing::Union{Nothing, Value}, ifCond=nothing::Union{Nothing, Value}; location=Location()) +function init(deviceTypeOperands::Vector{Value}, deviceNumOperand=nothing::Union{Nothing, Value}; ifCond=nothing::Union{Nothing, Value}, location=Location()) results = MLIRType[] operands = Value[deviceTypeOperands..., ] owned_regions = Region[] @@ -157,7 +157,7 @@ acc.loop gang vector { } attributes { collapse = 3 } ``` """ -function loop(gangNum=nothing::Union{Nothing, Value}, gangStatic=nothing::Union{Nothing, Value}, workerNum=nothing::Union{Nothing, Value}, vectorLength=nothing::Union{Nothing, Value}, tileOperands::Vector{Value}, privateOperands::Vector{Value}, reductionOperands::Vector{Value}; results::Vector{MLIRType}, collapse=nothing, seq=nothing, independent=nothing, auto_=nothing, reductionOp=nothing, exec_mapping=nothing, region::Region, location=Location()) +function loop(gangNum=nothing::Union{Nothing, Value}; gangStatic=nothing::Union{Nothing, Value}, workerNum=nothing::Union{Nothing, Value}, vectorLength=nothing::Union{Nothing, Value}, tileOperands::Vector{Value}, privateOperands::Vector{Value}, reductionOperands::Vector{Value}, results::Vector{MLIRType}, collapse=nothing, seq=nothing, independent=nothing, auto_=nothing, reductionOp=nothing, exec_mapping=nothing, region::Region, location=Location()) results = MLIRType[results..., ] operands = Value[tileOperands..., privateOperands..., reductionOperands..., ] owned_regions = Region[region, ] @@ -198,7 +198,7 @@ acc.parallel num_gangs(%c10) num_workers(%c10) } ``` """ -function parallel(async=nothing::Union{Nothing, Value}, waitOperands::Vector{Value}, numGangs=nothing::Union{Nothing, Value}, numWorkers=nothing::Union{Nothing, Value}, vectorLength=nothing::Union{Nothing, Value}, ifCond=nothing::Union{Nothing, Value}, selfCond=nothing::Union{Nothing, Value}, reductionOperands::Vector{Value}, copyOperands::Vector{Value}, copyinOperands::Vector{Value}, copyinReadonlyOperands::Vector{Value}, copyoutOperands::Vector{Value}, copyoutZeroOperands::Vector{Value}, createOperands::Vector{Value}, createZeroOperands::Vector{Value}, noCreateOperands::Vector{Value}, presentOperands::Vector{Value}, devicePtrOperands::Vector{Value}, attachOperands::Vector{Value}, gangPrivateOperands::Vector{Value}, gangFirstPrivateOperands::Vector{Value}; asyncAttr=nothing, waitAttr=nothing, selfAttr=nothing, reductionOp=nothing, defaultAttr=nothing, region::Region, location=Location()) +function parallel(async=nothing::Union{Nothing, Value}; waitOperands::Vector{Value}, numGangs=nothing::Union{Nothing, Value}, numWorkers=nothing::Union{Nothing, Value}, vectorLength=nothing::Union{Nothing, Value}, ifCond=nothing::Union{Nothing, Value}, selfCond=nothing::Union{Nothing, Value}, reductionOperands::Vector{Value}, copyOperands::Vector{Value}, copyinOperands::Vector{Value}, copyinReadonlyOperands::Vector{Value}, copyoutOperands::Vector{Value}, copyoutZeroOperands::Vector{Value}, createOperands::Vector{Value}, createZeroOperands::Vector{Value}, noCreateOperands::Vector{Value}, presentOperands::Vector{Value}, devicePtrOperands::Vector{Value}, attachOperands::Vector{Value}, gangPrivateOperands::Vector{Value}, gangFirstPrivateOperands::Vector{Value}, asyncAttr=nothing, waitAttr=nothing, selfAttr=nothing, reductionOp=nothing, defaultAttr=nothing, region::Region, location=Location()) results = MLIRType[] operands = Value[waitOperands..., reductionOperands..., copyOperands..., copyinOperands..., copyinReadonlyOperands..., copyoutOperands..., copyoutZeroOperands..., createOperands..., createZeroOperands..., noCreateOperands..., presentOperands..., devicePtrOperands..., attachOperands..., gangPrivateOperands..., gangFirstPrivateOperands..., ] owned_regions = Region[region, ] @@ -238,7 +238,7 @@ acc.shutdown acc.shutdown device_num(%dev1 : i32) ``` """ -function shutdown(deviceTypeOperands::Vector{Value}, deviceNumOperand=nothing::Union{Nothing, Value}, ifCond=nothing::Union{Nothing, Value}; location=Location()) +function shutdown(deviceTypeOperands::Vector{Value}, deviceNumOperand=nothing::Union{Nothing, Value}; ifCond=nothing::Union{Nothing, Value}, location=Location()) results = MLIRType[] operands = Value[deviceTypeOperands..., ] owned_regions = Region[] @@ -293,7 +293,7 @@ add to \$hostOperands. acc.update device(%d1 : memref<10xf32>) attributes {async} ``` """ -function update(ifCond=nothing::Union{Nothing, Value}, asyncOperand=nothing::Union{Nothing, Value}, waitDevnum=nothing::Union{Nothing, Value}, waitOperands::Vector{Value}, deviceTypeOperands::Vector{Value}, hostOperands::Vector{Value}, deviceOperands::Vector{Value}; async=nothing, wait=nothing, ifPresent=nothing, location=Location()) +function update(ifCond=nothing::Union{Nothing, Value}; asyncOperand=nothing::Union{Nothing, Value}, waitDevnum=nothing::Union{Nothing, Value}, waitOperands::Vector{Value}, deviceTypeOperands::Vector{Value}, hostOperands::Vector{Value}, deviceOperands::Vector{Value}, async=nothing, wait=nothing, ifPresent=nothing, location=Location()) results = MLIRType[] operands = Value[waitOperands..., deviceTypeOperands..., hostOperands..., deviceOperands..., ] owned_regions = Region[] @@ -328,7 +328,7 @@ acc.wait(%value1: index) acc.wait() async(%async1: i32) ``` """ -function wait(waitOperands::Vector{Value}, asyncOperand=nothing::Union{Nothing, Value}, waitDevnum=nothing::Union{Nothing, Value}, ifCond=nothing::Union{Nothing, Value}; async=nothing, location=Location()) +function wait(waitOperands::Vector{Value}, asyncOperand=nothing::Union{Nothing, Value}; waitDevnum=nothing::Union{Nothing, Value}, ifCond=nothing::Union{Nothing, Value}, async=nothing, location=Location()) results = MLIRType[] operands = Value[waitOperands..., ] owned_regions = Region[] diff --git a/src/Dialects/16/OpenMP.jl b/src/Dialects/16/OpenMP.jl index 60fed369..2551b4cd 100644 --- a/src/Dialects/16/OpenMP.jl +++ b/src/Dialects/16/OpenMP.jl @@ -416,7 +416,7 @@ threads complete. The optional \$proc_bind_val attribute controls the thread affinity for the execution of the parallel region. """ -function parallel(if_expr_var=nothing::Union{Nothing, Value}, num_threads_var=nothing::Union{Nothing, Value}, allocate_vars::Vector{Value}, allocators_vars::Vector{Value}, reduction_vars::Vector{Value}; reductions=nothing, proc_bind_val=nothing, region::Region, location=Location()) +function parallel(if_expr_var=nothing::Union{Nothing, Value}; num_threads_var=nothing::Union{Nothing, Value}, allocate_vars::Vector{Value}, allocators_vars::Vector{Value}, reduction_vars::Vector{Value}, reductions=nothing, proc_bind_val=nothing, region::Region, location=Location()) results = MLIRType[] operands = Value[allocate_vars..., allocators_vars..., reduction_vars..., ] owned_regions = Region[region, ] @@ -611,7 +611,7 @@ for (%i1, %i2) : index = (%c0, %c0) to (%c10, %c10) step (%c1, %c1) { } ``` """ -function simdloop(lowerBound::Vector{Value}, upperBound::Vector{Value}, step::Vector{Value}, aligned_vars::Vector{Value}, if_expr=nothing::Union{Nothing, Value}, nontemporal_vars::Vector{Value}; alignment_values=nothing, order_val=nothing, simdlen=nothing, safelen=nothing, inclusive=nothing, region::Region, location=Location()) +function simdloop(lowerBound::Vector{Value}, upperBound::Vector{Value}, step::Vector{Value}, aligned_vars::Vector{Value}, if_expr=nothing::Union{Nothing, Value}; nontemporal_vars::Vector{Value}, alignment_values=nothing, order_val=nothing, simdlen=nothing, safelen=nothing, inclusive=nothing, region::Region, location=Location()) results = MLIRType[] operands = Value[lowerBound..., upperBound..., step..., aligned_vars..., nontemporal_vars..., ] owned_regions = Region[region, ] @@ -679,7 +679,7 @@ even if the target task is not yet completed. TODO: map, is_device_ptr, depend, defaultmap, in_reduction """ -function target(if_expr=nothing::Union{Nothing, Value}, device=nothing::Union{Nothing, Value}, thread_limit=nothing::Union{Nothing, Value}; nowait=nothing, region::Region, location=Location()) +function target(if_expr=nothing::Union{Nothing, Value}; device=nothing::Union{Nothing, Value}, thread_limit=nothing::Union{Nothing, Value}, nowait=nothing, region::Region, location=Location()) results = MLIRType[] operands = Value[] owned_regions = Region[region, ] @@ -730,7 +730,7 @@ The \$map_types specifies the types and modifiers for the map clause. TODO: depend clause and map_type_modifier values iterator and mapper. """ -function target_data(if_expr=nothing::Union{Nothing, Value}, device=nothing::Union{Nothing, Value}, use_device_ptr::Vector{Value}, use_device_addr::Vector{Value}, map_operands::Vector{Value}; map_types, region::Region, location=Location()) +function target_data(if_expr=nothing::Union{Nothing, Value}; device=nothing::Union{Nothing, Value}, use_device_ptr::Vector{Value}, use_device_addr::Vector{Value}, map_operands::Vector{Value}, map_types, region::Region, location=Location()) results = MLIRType[] operands = Value[use_device_ptr..., use_device_addr..., map_operands..., ] owned_regions = Region[region, ] @@ -772,7 +772,7 @@ The \$map_types specifies the types and modifiers for the map clause. TODO: depend clause and map_type_modifier values iterator and mapper. """ -function target_enter_data(if_expr=nothing::Union{Nothing, Value}, device=nothing::Union{Nothing, Value}, map_operands::Vector{Value}; nowait=nothing, map_types, location=Location()) +function target_enter_data(if_expr=nothing::Union{Nothing, Value}; device=nothing::Union{Nothing, Value}, map_operands::Vector{Value}, nowait=nothing, map_types, location=Location()) results = MLIRType[] operands = Value[map_operands..., ] owned_regions = Region[] @@ -815,7 +815,7 @@ The \$map_types specifies the types and modifiers for the map clause. TODO: depend clause and map_type_modifier values iterator and mapper. """ -function target_exit_data(if_expr=nothing::Union{Nothing, Value}, device=nothing::Union{Nothing, Value}, map_operands::Vector{Value}; nowait=nothing, map_types, location=Location()) +function target_exit_data(if_expr=nothing::Union{Nothing, Value}; device=nothing::Union{Nothing, Value}, map_operands::Vector{Value}, nowait=nothing, map_types, location=Location()) results = MLIRType[] operands = Value[map_operands..., ] owned_regions = Region[] @@ -973,7 +973,7 @@ construct. Thus, the taskloop construct creates an implicit taskgroup region. If the `nogroup` clause is present, no implicit taskgroup region is created. """ -function taskloop(lowerBound::Vector{Value}, upperBound::Vector{Value}, step::Vector{Value}, if_expr=nothing::Union{Nothing, Value}, final_expr=nothing::Union{Nothing, Value}, in_reduction_vars::Vector{Value}, reduction_vars::Vector{Value}, priority=nothing::Union{Nothing, Value}, allocate_vars::Vector{Value}, allocators_vars::Vector{Value}, grain_size=nothing::Union{Nothing, Value}, num_tasks=nothing::Union{Nothing, Value}; inclusive=nothing, untied=nothing, mergeable=nothing, in_reductions=nothing, reductions=nothing, nogroup=nothing, region::Region, location=Location()) +function taskloop(lowerBound::Vector{Value}, upperBound::Vector{Value}, step::Vector{Value}, if_expr=nothing::Union{Nothing, Value}; final_expr=nothing::Union{Nothing, Value}, in_reduction_vars::Vector{Value}, reduction_vars::Vector{Value}, priority=nothing::Union{Nothing, Value}, allocate_vars::Vector{Value}, allocators_vars::Vector{Value}, grain_size=nothing::Union{Nothing, Value}, num_tasks=nothing::Union{Nothing, Value}, inclusive=nothing, untied=nothing, mergeable=nothing, in_reductions=nothing, reductions=nothing, nogroup=nothing, region::Region, location=Location()) results = MLIRType[] operands = Value[lowerBound..., upperBound..., step..., in_reduction_vars..., reduction_vars..., allocate_vars..., allocators_vars..., ] owned_regions = Region[region, ] @@ -1042,7 +1042,7 @@ The `allocators_vars` and `allocate_vars` arguments are a variadic list of values that specify the memory allocator to be used to obtain storage for private values. """ -function task(if_expr=nothing::Union{Nothing, Value}, final_expr=nothing::Union{Nothing, Value}, in_reduction_vars::Vector{Value}, priority=nothing::Union{Nothing, Value}, allocate_vars::Vector{Value}, allocators_vars::Vector{Value}; untied=nothing, mergeable=nothing, in_reductions=nothing, region::Region, location=Location()) +function task(if_expr=nothing::Union{Nothing, Value}; final_expr=nothing::Union{Nothing, Value}, in_reduction_vars::Vector{Value}, priority=nothing::Union{Nothing, Value}, allocate_vars::Vector{Value}, allocators_vars::Vector{Value}, untied=nothing, mergeable=nothing, in_reductions=nothing, region::Region, location=Location()) results = MLIRType[] operands = Value[in_reduction_vars..., allocate_vars..., allocators_vars..., ] owned_regions = Region[region, ] diff --git a/src/Dialects/16/PDL.jl b/src/Dialects/16/PDL.jl index d5bdffad..f20df0ec 100644 --- a/src/Dialects/16/PDL.jl +++ b/src/Dialects/16/PDL.jl @@ -442,7 +442,7 @@ pdl.replace %root with (%vals : !pdl.range) pdl.replace %root with %otherOp ``` """ -function replace(opValue::Value, replOperation=nothing::Union{Nothing, Value}, replValues::Vector{Value}; location=Location()) +function replace(opValue::Value, replOperation=nothing::Union{Nothing, Value}; replValues::Vector{Value}, location=Location()) results = MLIRType[] operands = Value[opValue, replValues..., ] owned_regions = Region[] @@ -580,7 +580,7 @@ pdl.rewrite { } ``` """ -function rewrite(root=nothing::Union{Nothing, Value}, externalArgs::Vector{Value}; name=nothing, bodyRegion::Region, location=Location()) +function rewrite(root=nothing::Union{Nothing, Value}; externalArgs::Vector{Value}, name=nothing, bodyRegion::Region, location=Location()) results = MLIRType[] operands = Value[externalArgs..., ] owned_regions = Region[bodyRegion, ] diff --git a/src/Dialects/16/Tensor.jl b/src/Dialects/16/Tensor.jl index 72bce95b..327e0802 100644 --- a/src/Dialects/16/Tensor.jl +++ b/src/Dialects/16/Tensor.jl @@ -632,7 +632,7 @@ Example NC_to_NCnc with padding: inner_tiles = [8, 2] into %arg1 : tensor<13x15xf32> -> tensor<2x8x8x2xf32> ``` """ -function pack(source::Value, dest::Value, padding_value=nothing::Union{Nothing, Value}, inner_tiles::Vector{Value}; result=nothing::Union{Nothing, MLIRType}, outer_dims_perm=nothing, inner_dims_pos, static_inner_tiles, location=Location()) +function pack(source::Value, dest::Value, padding_value=nothing::Union{Nothing, Value}; inner_tiles::Vector{Value}, result=nothing::Union{Nothing, MLIRType}, outer_dims_perm=nothing, inner_dims_pos, static_inner_tiles, location=Location()) results = MLIRType[] operands = Value[source, dest, inner_tiles..., ] owned_regions = Region[] diff --git a/src/Dialects/16/X86Vector.jl b/src/Dialects/16/X86Vector.jl index 30799bb9..43f5bd89 100644 --- a/src/Dialects/16/X86Vector.jl +++ b/src/Dialects/16/X86Vector.jl @@ -6,10 +6,10 @@ import ...API """ -`avx_intr.dp.ps.256` +`avx_intr_dp_ps_256` """ -function avx_intr.dp.ps.256(a::Value, b::Value, c::Value; res=nothing::Union{Nothing, MLIRType}, location=Location()) +function avx_intr_dp_ps_256(a::Value, b::Value, c::Value; res=nothing::Union{Nothing, MLIRType}, location=Location()) results = MLIRType[] operands = Value[a, b, c, ] owned_regions = Region[] @@ -26,7 +26,7 @@ function avx_intr.dp.ps.256(a::Value, b::Value, c::Value; res=nothing::Union{Not end """ -`avx_intr.dot` +`avx_intr_dot` Computes the 4-way dot products of the lower and higher parts of the source vectors and broadcasts the two results to the lower and higher elements of @@ -43,7 +43,7 @@ dot product of the two source vectors. %d = arith.addf %1, %2 : f32 ``` """ -function avx_intr.dot(a::Value, b::Value; res=nothing::Union{Nothing, MLIRType}, location=Location()) +function avx_intr_dot(a::Value, b::Value; res=nothing::Union{Nothing, MLIRType}, location=Location()) results = MLIRType[] operands = Value[a, b, ] owned_regions = Region[] @@ -60,10 +60,10 @@ function avx_intr.dot(a::Value, b::Value; res=nothing::Union{Nothing, MLIRType}, end """ -`avx512_intr.mask.compress` +`avx512_intr_mask_compress` """ -function avx512_intr.mask.compress(a::Value, src::Value, k::Value; res=nothing::Union{Nothing, MLIRType}, location=Location()) +function avx512_intr_mask_compress(a::Value, src::Value, k::Value; res=nothing::Union{Nothing, MLIRType}, location=Location()) results = MLIRType[] operands = Value[a, src, k, ] owned_regions = Region[] @@ -80,7 +80,7 @@ function avx512_intr.mask.compress(a::Value, src::Value, k::Value; res=nothing:: end """ -`avx512_mask.compress` +`avx512_mask_compress` The mask.compress op is an AVX512 specific op that can lower to the `llvm.mask.compress` instruction. Instead of `src`, a constant vector @@ -94,7 +94,7 @@ Contiguously store the active integer/floating-point elements in `a` (those with their respective bit set in writemask `k`) to `dst`, and pass through the remaining elements from `src`. """ -function avx512_mask.compress(k::Value, a::Value, src=nothing::Union{Nothing, Value}; dst=nothing::Union{Nothing, MLIRType}, constant_src=nothing, location=Location()) +function avx512_mask_compress(k::Value, a::Value, src=nothing::Union{Nothing, Value}; dst=nothing::Union{Nothing, MLIRType}, constant_src=nothing, location=Location()) results = MLIRType[] operands = Value[k, a, ] owned_regions = Region[] @@ -113,7 +113,7 @@ function avx512_mask.compress(k::Value, a::Value, src=nothing::Union{Nothing, Va end """ -`avx512_mask.rndscale` +`avx512_mask_rndscale` The mask.rndscale op is an AVX512 specific op that can lower to the proper LLVMAVX512 operation: `llvm.mask.rndscale.ps.512` or @@ -126,7 +126,7 @@ Round packed floating-point elements in `a` to the number of fraction bits specified by `imm`, and store the results in `dst` using writemask `k` (elements are copied from src when the corresponding mask bit is not set). """ -function avx512_mask.rndscale(src::Value, k::Value, a::Value, imm::Value, rounding::Value; dst=nothing::Union{Nothing, MLIRType}, location=Location()) +function avx512_mask_rndscale(src::Value, k::Value, a::Value, imm::Value, rounding::Value; dst=nothing::Union{Nothing, MLIRType}, location=Location()) results = MLIRType[] operands = Value[src, k, a, imm, rounding, ] owned_regions = Region[] @@ -143,10 +143,10 @@ function avx512_mask.rndscale(src::Value, k::Value, a::Value, imm::Value, roundi end """ -`avx512_intr.mask.rndscale.pd.512` +`avx512_intr_mask_rndscale_pd_512` """ -function avx512_intr.mask.rndscale.pd.512(src::Value, k::Value, a::Value, imm::Value, rounding::Value; res=nothing::Union{Nothing, MLIRType}, location=Location()) +function avx512_intr_mask_rndscale_pd_512(src::Value, k::Value, a::Value, imm::Value, rounding::Value; res=nothing::Union{Nothing, MLIRType}, location=Location()) results = MLIRType[] operands = Value[src, k, a, imm, rounding, ] owned_regions = Region[] @@ -163,10 +163,10 @@ function avx512_intr.mask.rndscale.pd.512(src::Value, k::Value, a::Value, imm::V end """ -`avx512_intr.mask.rndscale.ps.512` +`avx512_intr_mask_rndscale_ps_512` """ -function avx512_intr.mask.rndscale.ps.512(src::Value, k::Value, a::Value, imm::Value, rounding::Value; res=nothing::Union{Nothing, MLIRType}, location=Location()) +function avx512_intr_mask_rndscale_ps_512(src::Value, k::Value, a::Value, imm::Value, rounding::Value; res=nothing::Union{Nothing, MLIRType}, location=Location()) results = MLIRType[] operands = Value[src, k, a, imm, rounding, ] owned_regions = Region[] @@ -183,7 +183,7 @@ function avx512_intr.mask.rndscale.ps.512(src::Value, k::Value, a::Value, imm::V end """ -`avx512_mask.scalef` +`avx512_mask_scalef` The `mask.scalef` op is an AVX512 specific op that can lower to the proper LLVMAVX512 operation: `llvm.mask.scalef.ps.512` or @@ -196,7 +196,7 @@ Scale the packed floating-point elements in `a` using values from `b`, and store the results in `dst` using writemask `k` (elements are copied from src when the corresponding mask bit is not set). """ -function avx512_mask.scalef(src::Value, a::Value, b::Value, k::Value, rounding::Value; dst=nothing::Union{Nothing, MLIRType}, location=Location()) +function avx512_mask_scalef(src::Value, a::Value, b::Value, k::Value, rounding::Value; dst=nothing::Union{Nothing, MLIRType}, location=Location()) results = MLIRType[] operands = Value[src, a, b, k, rounding, ] owned_regions = Region[] @@ -213,10 +213,10 @@ function avx512_mask.scalef(src::Value, a::Value, b::Value, k::Value, rounding:: end """ -`avx512_intr.mask.scalef.pd.512` +`avx512_intr_mask_scalef_pd_512` """ -function avx512_intr.mask.scalef.pd.512(src::Value, a::Value, b::Value, k::Value, rounding::Value; res=nothing::Union{Nothing, MLIRType}, location=Location()) +function avx512_intr_mask_scalef_pd_512(src::Value, a::Value, b::Value, k::Value, rounding::Value; res=nothing::Union{Nothing, MLIRType}, location=Location()) results = MLIRType[] operands = Value[src, a, b, k, rounding, ] owned_regions = Region[] @@ -233,10 +233,10 @@ function avx512_intr.mask.scalef.pd.512(src::Value, a::Value, b::Value, k::Value end """ -`avx512_intr.mask.scalef.ps.512` +`avx512_intr_mask_scalef_ps_512` """ -function avx512_intr.mask.scalef.ps.512(src::Value, a::Value, b::Value, k::Value, rounding::Value; res=nothing::Union{Nothing, MLIRType}, location=Location()) +function avx512_intr_mask_scalef_ps_512(src::Value, a::Value, b::Value, k::Value, rounding::Value; res=nothing::Union{Nothing, MLIRType}, location=Location()) results = MLIRType[] operands = Value[src, a, b, k, rounding, ] owned_regions = Region[] @@ -253,10 +253,10 @@ function avx512_intr.mask.scalef.ps.512(src::Value, a::Value, b::Value, k::Value end """ -`avx_intr.rsqrt.ps.256` +`avx_intr_rsqrt_ps_256` """ -function avx_intr.rsqrt.ps.256(a::Value; res=nothing::Union{Nothing, MLIRType}, location=Location()) +function avx_intr_rsqrt_ps_256(a::Value; res=nothing::Union{Nothing, MLIRType}, location=Location()) results = MLIRType[] operands = Value[a, ] owned_regions = Region[] @@ -293,10 +293,10 @@ function avx_rsqrt(a::Value; b=nothing::Union{Nothing, MLIRType}, location=Locat end """ -`avx512_intr.vp2intersect.d.512` +`avx512_intr_vp2intersect_d_512` """ -function avx512_intr.vp2intersect.d.512(a::Value, b::Value; res::MLIRType, location=Location()) +function avx512_intr_vp2intersect_d_512(a::Value, b::Value; res::MLIRType, location=Location()) results = MLIRType[res, ] operands = Value[a, b, ] owned_regions = Region[] @@ -345,10 +345,10 @@ function avx512_vp2intersect(a::Value, b::Value; k1=nothing::Union{Nothing, MLIR end """ -`avx512_intr.vp2intersect.q.512` +`avx512_intr_vp2intersect_q_512` """ -function avx512_intr.vp2intersect.q.512(a::Value, b::Value; res::MLIRType, location=Location()) +function avx512_intr_vp2intersect_q_512(a::Value, b::Value; res::MLIRType, location=Location()) results = MLIRType[res, ] operands = Value[a, b, ] owned_regions = Region[] From 6fdb226842523ffe347c28d46fbb2234f711d001 Mon Sep 17 00:00:00 2001 From: Valentin Churavy Date: Tue, 16 Jan 2024 15:46:34 -0500 Subject: [PATCH 3/5] Manual fixes --- src/Dialects/14/ArmNeon.jl | 26 +++---- src/Dialects/14/SPIRV.jl | 2 +- src/Dialects/15/ArmNeon.jl | 26 +++---- src/Dialects/15/MLProgram.jl | 130 +++++++++++++++++------------------ src/Dialects/15/SPIRV.jl | 2 +- src/Dialects/16/ArmNeon.jl | 26 +++---- src/Dialects/16/MLProgram.jl | 130 +++++++++++++++++------------------ src/Dialects/16/SPIRV.jl | 2 +- 8 files changed, 172 insertions(+), 172 deletions(-) diff --git a/src/Dialects/14/ArmNeon.jl b/src/Dialects/14/ArmNeon.jl index e531fce6..7fafa0f7 100644 --- a/src/Dialects/14/ArmNeon.jl +++ b/src/Dialects/14/ArmNeon.jl @@ -43,20 +43,20 @@ corresponding entry of `a`: res[i] := a[i] + dot_product(b[i, ...], c[i, ...]) ``` """ -function 2d_sdot(a::Value, b::Value, c::Value; res::MLIRType, location=Location()) - results = MLIRType[res, ] - operands = Value[a, b, c, ] - owned_regions = Region[] - successors = Block[] - attributes = NamedAttribute[] +# function 2d_sdot(a::Value, b::Value, c::Value; res::MLIRType, location=Location()) +# results = MLIRType[res, ] +# operands = Value[a, b, c, ] +# owned_regions = Region[] +# successors = Block[] +# attributes = NamedAttribute[] - create_operation( - "arm_neon.2d.sdot", location; - operands, owned_regions, successors, attributes, - results=results, - result_inference=false - ) -end +# create_operation( +# "arm_neon.2d.sdot", location; +# operands, owned_regions, successors, attributes, +# results=results, +# result_inference=false +# ) +# end """ `intr_sdot` diff --git a/src/Dialects/14/SPIRV.jl b/src/Dialects/14/SPIRV.jl index f3720670..c980f7dc 100644 --- a/src/Dialects/14/SPIRV.jl +++ b/src/Dialects/14/SPIRV.jl @@ -4502,7 +4502,7 @@ spv.GlobalVariable @var2 bind(1, 2) : !spv.ptr spv.GlobalVariable @var3 built_in(\"GlobalInvocationId\") : !spv.ptr, Input> ``` """ -function GlobalVariable(; type, sym_name, initializer=nothing, location=nothing, binding=nothing, descriptorSet=nothing, builtin=nothing, location=Location()) +function GlobalVariable(; type, sym_name, initializer=nothing, binding=nothing, descriptorSet=nothing, builtin=nothing, location=Location()) results = MLIRType[] operands = Value[] owned_regions = Region[] diff --git a/src/Dialects/15/ArmNeon.jl b/src/Dialects/15/ArmNeon.jl index e531fce6..7fafa0f7 100644 --- a/src/Dialects/15/ArmNeon.jl +++ b/src/Dialects/15/ArmNeon.jl @@ -43,20 +43,20 @@ corresponding entry of `a`: res[i] := a[i] + dot_product(b[i, ...], c[i, ...]) ``` """ -function 2d_sdot(a::Value, b::Value, c::Value; res::MLIRType, location=Location()) - results = MLIRType[res, ] - operands = Value[a, b, c, ] - owned_regions = Region[] - successors = Block[] - attributes = NamedAttribute[] +# function 2d_sdot(a::Value, b::Value, c::Value; res::MLIRType, location=Location()) +# results = MLIRType[res, ] +# operands = Value[a, b, c, ] +# owned_regions = Region[] +# successors = Block[] +# attributes = NamedAttribute[] - create_operation( - "arm_neon.2d.sdot", location; - operands, owned_regions, successors, attributes, - results=results, - result_inference=false - ) -end +# create_operation( +# "arm_neon.2d.sdot", location; +# operands, owned_regions, successors, attributes, +# results=results, +# result_inference=false +# ) +# end """ `intr_sdot` diff --git a/src/Dialects/15/MLProgram.jl b/src/Dialects/15/MLProgram.jl index 99f6fd45..59f540e2 100644 --- a/src/Dialects/15/MLProgram.jl +++ b/src/Dialects/15/MLProgram.jl @@ -55,20 +55,20 @@ such a load can be considered to have no side effects. %0 = ml_program.global_load_const @foobar : tensor ``` """ -function global_load_const(; result::MLIRType, global, location=Location()) - results = MLIRType[result, ] - operands = Value[] - owned_regions = Region[] - successors = Block[] - attributes = NamedAttribute[namedattribute("global", global), ] +# function global_load_const(; result::MLIRType, global, location=Location()) +# results = MLIRType[result, ] +# operands = Value[] +# owned_regions = Region[] +# successors = Block[] +# attributes = NamedAttribute[namedattribute("global", global), ] - create_operation( - "ml_program.global_load_const", location; - operands, owned_regions, successors, attributes, - results=results, - result_inference=false - ) -end +# create_operation( +# "ml_program.global_load_const", location; +# operands, owned_regions, successors, attributes, +# results=results, +# result_inference=false +# ) +# end """ `global_load_graph` @@ -90,20 +90,20 @@ without additional consideration to evaluation order constraints. ordering (%token -> !ml_program.token) : tensor ``` """ -function global_load_graph(consumeTokens::Vector{Value}; result::MLIRType, produceToken::MLIRType, global, location=Location()) - results = MLIRType[result, produceToken, ] - operands = Value[consumeTokens..., ] - owned_regions = Region[] - successors = Block[] - attributes = NamedAttribute[namedattribute("global", global), ] +# function global_load_graph(consumeTokens::Vector{Value}; result::MLIRType, produceToken::MLIRType, global, location=Location()) +# results = MLIRType[result, produceToken, ] +# operands = Value[consumeTokens..., ] +# owned_regions = Region[] +# successors = Block[] +# attributes = NamedAttribute[namedattribute("global", global), ] - create_operation( - "ml_program.global_load_graph", location; - operands, owned_regions, successors, attributes, - results=results, - result_inference=false - ) -end +# create_operation( +# "ml_program.global_load_graph", location; +# operands, owned_regions, successors, attributes, +# results=results, +# result_inference=false +# ) +# end """ `global_load` @@ -125,20 +125,20 @@ without additional consideration to evaluation order constraints. See %0 = ml_program.global_load @foobar : tensor ``` """ -function global_load(; result::MLIRType, global, location=Location()) - results = MLIRType[result, ] - operands = Value[] - owned_regions = Region[] - successors = Block[] - attributes = NamedAttribute[namedattribute("global", global), ] +# function global_load(; result::MLIRType, global, location=Location()) +# results = MLIRType[result, ] +# operands = Value[] +# owned_regions = Region[] +# successors = Block[] +# attributes = NamedAttribute[namedattribute("global", global), ] - create_operation( - "ml_program.global_load", location; - operands, owned_regions, successors, attributes, - results=results, - result_inference=false - ) -end +# create_operation( +# "ml_program.global_load", location; +# operands, owned_regions, successors, attributes, +# results=results, +# result_inference=false +# ) +# end """ `global_` @@ -208,20 +208,20 @@ without additional consideration to evaluation order constraints. ordering (%in_token -> !ml_program.token) : tensor ``` """ -function global_store_graph(value::Value, consumeTokens::Vector{Value}; produceToken::MLIRType, global, location=Location()) - results = MLIRType[produceToken, ] - operands = Value[value, consumeTokens..., ] - owned_regions = Region[] - successors = Block[] - attributes = NamedAttribute[namedattribute("global", global), ] +# function global_store_graph(value::Value, consumeTokens::Vector{Value}; produceToken::MLIRType, global, location=Location()) +# results = MLIRType[produceToken, ] +# operands = Value[value, consumeTokens..., ] +# owned_regions = Region[] +# successors = Block[] +# attributes = NamedAttribute[namedattribute("global", global), ] - create_operation( - "ml_program.global_store_graph", location; - operands, owned_regions, successors, attributes, - results=results, - result_inference=false - ) -end +# create_operation( +# "ml_program.global_store_graph", location; +# operands, owned_regions, successors, attributes, +# results=results, +# result_inference=false +# ) +# end """ `global_store` @@ -243,20 +243,20 @@ without additional consideration to evaluation order constraints. See ml_program.global_store @foobar = %0 : tensor ``` """ -function global_store(value::Value; global, location=Location()) - results = MLIRType[] - operands = Value[value, ] - owned_regions = Region[] - successors = Block[] - attributes = NamedAttribute[namedattribute("global", global), ] +# function global_store(value::Value; global, location=Location()) +# results = MLIRType[] +# operands = Value[value, ] +# owned_regions = Region[] +# successors = Block[] +# attributes = NamedAttribute[namedattribute("global", global), ] - create_operation( - "ml_program.global_store", location; - operands, owned_regions, successors, attributes, - results=results, - result_inference=false - ) -end +# create_operation( +# "ml_program.global_store", location; +# operands, owned_regions, successors, attributes, +# results=results, +# result_inference=false +# ) +# end """ `output` diff --git a/src/Dialects/15/SPIRV.jl b/src/Dialects/15/SPIRV.jl index c6c7b5a2..eec360dc 100644 --- a/src/Dialects/15/SPIRV.jl +++ b/src/Dialects/15/SPIRV.jl @@ -5201,7 +5201,7 @@ spv.GlobalVariable @var2 bind(1, 2) : !spv.ptr spv.GlobalVariable @var3 built_in(\"GlobalInvocationId\") : !spv.ptr, Input> ``` """ -function GlobalVariable(; type, sym_name, initializer=nothing, location=nothing, binding=nothing, descriptor_set=nothing, builtin=nothing, location=Location()) +function GlobalVariable(; type, sym_name, initializer=nothing, binding=nothing, descriptor_set=nothing, builtin=nothing, location=Location()) results = MLIRType[] operands = Value[] owned_regions = Region[] diff --git a/src/Dialects/16/ArmNeon.jl b/src/Dialects/16/ArmNeon.jl index e531fce6..7fafa0f7 100644 --- a/src/Dialects/16/ArmNeon.jl +++ b/src/Dialects/16/ArmNeon.jl @@ -43,20 +43,20 @@ corresponding entry of `a`: res[i] := a[i] + dot_product(b[i, ...], c[i, ...]) ``` """ -function 2d_sdot(a::Value, b::Value, c::Value; res::MLIRType, location=Location()) - results = MLIRType[res, ] - operands = Value[a, b, c, ] - owned_regions = Region[] - successors = Block[] - attributes = NamedAttribute[] +# function 2d_sdot(a::Value, b::Value, c::Value; res::MLIRType, location=Location()) +# results = MLIRType[res, ] +# operands = Value[a, b, c, ] +# owned_regions = Region[] +# successors = Block[] +# attributes = NamedAttribute[] - create_operation( - "arm_neon.2d.sdot", location; - operands, owned_regions, successors, attributes, - results=results, - result_inference=false - ) -end +# create_operation( +# "arm_neon.2d.sdot", location; +# operands, owned_regions, successors, attributes, +# results=results, +# result_inference=false +# ) +# end """ `intr_sdot` diff --git a/src/Dialects/16/MLProgram.jl b/src/Dialects/16/MLProgram.jl index 9d99c0ea..105b9b7c 100644 --- a/src/Dialects/16/MLProgram.jl +++ b/src/Dialects/16/MLProgram.jl @@ -57,20 +57,20 @@ such a load can be considered to have no side effects. %0 = ml_program.global_load_const @foobar : tensor ``` """ -function global_load_const(; result::MLIRType, global, location=Location()) - results = MLIRType[result, ] - operands = Value[] - owned_regions = Region[] - successors = Block[] - attributes = NamedAttribute[namedattribute("global", global), ] +# function global_load_const(; result::MLIRType, global, location=Location()) +# results = MLIRType[result, ] +# operands = Value[] +# owned_regions = Region[] +# successors = Block[] +# attributes = NamedAttribute[namedattribute("global", global), ] - create_operation( - "ml_program.global_load_const", location; - operands, owned_regions, successors, attributes, - results=results, - result_inference=false - ) -end +# create_operation( +# "ml_program.global_load_const", location; +# operands, owned_regions, successors, attributes, +# results=results, +# result_inference=false +# ) +# end """ `global_load_graph` @@ -92,20 +92,20 @@ without additional consideration to evaluation order constraints. ordering (%token -> !ml_program.token) : tensor ``` """ -function global_load_graph(consumeTokens::Vector{Value}; result::MLIRType, produceToken::MLIRType, global, location=Location()) - results = MLIRType[result, produceToken, ] - operands = Value[consumeTokens..., ] - owned_regions = Region[] - successors = Block[] - attributes = NamedAttribute[namedattribute("global", global), ] +# function global_load_graph(consumeTokens::Vector{Value}; result::MLIRType, produceToken::MLIRType, global, location=Location()) +# results = MLIRType[result, produceToken, ] +# operands = Value[consumeTokens..., ] +# owned_regions = Region[] +# successors = Block[] +# attributes = NamedAttribute[namedattribute("global", global), ] - create_operation( - "ml_program.global_load_graph", location; - operands, owned_regions, successors, attributes, - results=results, - result_inference=false - ) -end +# create_operation( +# "ml_program.global_load_graph", location; +# operands, owned_regions, successors, attributes, +# results=results, +# result_inference=false +# ) +# end """ `global_load` @@ -127,20 +127,20 @@ without additional consideration to evaluation order constraints. See %0 = ml_program.global_load @foobar : tensor ``` """ -function global_load(; result::MLIRType, global, location=Location()) - results = MLIRType[result, ] - operands = Value[] - owned_regions = Region[] - successors = Block[] - attributes = NamedAttribute[namedattribute("global", global), ] +# function global_load(; result::MLIRType, global, location=Location()) +# results = MLIRType[result, ] +# operands = Value[] +# owned_regions = Region[] +# successors = Block[] +# attributes = NamedAttribute[namedattribute("global", global), ] - create_operation( - "ml_program.global_load", location; - operands, owned_regions, successors, attributes, - results=results, - result_inference=false - ) -end +# create_operation( +# "ml_program.global_load", location; +# operands, owned_regions, successors, attributes, +# results=results, +# result_inference=false +# ) +# end """ `global_` @@ -210,20 +210,20 @@ without additional consideration to evaluation order constraints. ordering (%in_token -> !ml_program.token) : tensor ``` """ -function global_store_graph(value::Value, consumeTokens::Vector{Value}; produceToken::MLIRType, global, location=Location()) - results = MLIRType[produceToken, ] - operands = Value[value, consumeTokens..., ] - owned_regions = Region[] - successors = Block[] - attributes = NamedAttribute[namedattribute("global", global), ] +# function global_store_graph(value::Value, consumeTokens::Vector{Value}; produceToken::MLIRType, global, location=Location()) +# results = MLIRType[produceToken, ] +# operands = Value[value, consumeTokens..., ] +# owned_regions = Region[] +# successors = Block[] +# attributes = NamedAttribute[namedattribute("global", global), ] - create_operation( - "ml_program.global_store_graph", location; - operands, owned_regions, successors, attributes, - results=results, - result_inference=false - ) -end +# create_operation( +# "ml_program.global_store_graph", location; +# operands, owned_regions, successors, attributes, +# results=results, +# result_inference=false +# ) +# end """ `global_store` @@ -245,20 +245,20 @@ without additional consideration to evaluation order constraints. See ml_program.global_store @foobar = %0 : tensor ``` """ -function global_store(value::Value; global, location=Location()) - results = MLIRType[] - operands = Value[value, ] - owned_regions = Region[] - successors = Block[] - attributes = NamedAttribute[namedattribute("global", global), ] +# function global_store(value::Value; global, location=Location()) +# results = MLIRType[] +# operands = Value[value, ] +# owned_regions = Region[] +# successors = Block[] +# attributes = NamedAttribute[namedattribute("global", global), ] - create_operation( - "ml_program.global_store", location; - operands, owned_regions, successors, attributes, - results=results, - result_inference=false - ) -end +# create_operation( +# "ml_program.global_store", location; +# operands, owned_regions, successors, attributes, +# results=results, +# result_inference=false +# ) +# end """ `output` diff --git a/src/Dialects/16/SPIRV.jl b/src/Dialects/16/SPIRV.jl index a8b3f388..bddaf5ac 100644 --- a/src/Dialects/16/SPIRV.jl +++ b/src/Dialects/16/SPIRV.jl @@ -5354,7 +5354,7 @@ spirv.GlobalVariable @var2 bind(1, 2) : !spirv.ptr spirv.GlobalVariable @var3 built_in(\"GlobalInvocationId\") : !spirv.ptr, Input> ``` """ -function GlobalVariable(; type, sym_name, initializer=nothing, location=nothing, binding=nothing, descriptor_set=nothing, builtin=nothing, location=Location()) +function GlobalVariable(; type, sym_name, initializer=nothing, binding=nothing, descriptor_set=nothing, builtin=nothing, location=Location()) results = MLIRType[] operands = Value[] owned_regions = Region[] From 8a63b5010fe95741bf13191d8e0b61918ac9da74 Mon Sep 17 00:00:00 2001 From: Paul Berg Date: Wed, 10 Jan 2024 16:10:23 +0100 Subject: [PATCH 4/5] update toy brutus drop support for 1.9 --- examples/brutus.jl | 118 ++++++++++++++++++++++++++------------------- src/Dialects.jl | 8 +-- src/IR/IR.jl | 31 ------------ test/examples.jl | 4 +- 4 files changed, 73 insertions(+), 88 deletions(-) diff --git a/examples/brutus.jl b/examples/brutus.jl index bba5835b..5d29a4fa 100644 --- a/examples/brutus.jl +++ b/examples/brutus.jl @@ -1,33 +1,55 @@ +""" +Brutus is a toy implementation of a Julia typed IR to MLIR conversion, the name +is a reference to the [brutus](https://github.com/JuliaLabs/brutus) project from +the MIT JuliaLabs which performs a similar conversion (with a lot more language constructs supported) +but from C++. +""" module Brutus import LLVM using MLIR.IR -using MLIR.Dialects: arith, func, cf, std +using MLIR.Dialects: arith, func, cf using Core: PhiNode, GotoNode, GotoIfNot, SSAValue, Argument, ReturnNode, PiNode const BrutusScalar = Union{Bool,Int64,Int32,Float32,Float64} +module Predicates +const eq = 0 +const ne = 1 +const slt = 2 +const sle = 3 +const sgt = 4 +const sge = 5 +const ult = 6 +const ule = 7 +const ugt = 8 +const uge = 9 +end + function cmpi_pred(predicate) - function(ops; loc=Location()) - arith.cmpi(predicate, ops; loc) + function (ops...; location = Location()) + arith.cmpi(ops...; result=IR.MLIRType(Bool), predicate, location) end end function single_op_wrapper(fop) - (block::Block, args::Vector{Value}; loc=Location()) -> push!(block, fop(args; loc)) + (block::Block, args::Vector{Value}; location=Location()) -> push!(block, fop(args...; location)) end const intrinsics_to_mlir = Dict([ Base.add_int => single_op_wrapper(arith.addi), - Base.sle_int => single_op_wrapper(cmpi_pred(arith.Predicates.sle)), - Base.slt_int => single_op_wrapper(cmpi_pred(arith.Predicates.slt)), - Base.:(===) => single_op_wrapper(cmpi_pred(arith.Predicates.eq)), + Base.sle_int => single_op_wrapper(cmpi_pred(Predicates.sle)), + Base.slt_int => single_op_wrapper(cmpi_pred(Predicates.slt)), + Base.:(===) => single_op_wrapper(cmpi_pred(Predicates.eq)), Base.mul_int => single_op_wrapper(arith.muli), Base.mul_float => single_op_wrapper(arith.mulf), - Base.not_int => function(block, args; loc=Location()) + Base.not_int => function (block, args; location=Location()) arg = only(args) - ones = push!(block, arith.constant(-1, IR.get_type(arg); loc)) |> IR.get_result - push!(block, arith.xori(Value[arg, ones]; loc)) + mT = IR.get_type(arg) + T = IR.julia_type(mT) + ones = push!(block, arith.constant(value=typemax(UInt64) % T; + result=mT, location)) |> IR.get_result + push!(block, arith.xori(arg, ones; location)) end, ]) @@ -85,7 +107,7 @@ function code_mlir(f, types) values = Vector{Value}(undef, length(ir.stmts)) - for dialect in (LLVM.version() >= v"15" ? ("func", "cf") : ("std",)) + for dialect in ("func", "cf") IR.get_or_load_dialect!(dialect) end @@ -107,7 +129,7 @@ function code_mlir(f, types) elseif x isa Core.Argument IR.get_argument(entry_block, x.n - 1) elseif x isa BrutusScalar - IR.get_result(push!(current_block, arith.constant(x))) + IR.get_result(push!(current_block, arith.constant(;value=x))) else error("could not use value $x inside MLIR") end @@ -137,8 +159,8 @@ function code_mlir(f, types) fop! = intrinsics_to_mlir[called_func] args = get_value.(@view inst.args[begin+1:end]) - loc = Location(string(line.file), line.line, 0) - res = IR.get_result(fop!(current_block, args; loc)) + location = Location(string(line.file), line.line, 0) + res = IR.get_result(fop!(current_block, args; location)) values[sidx] = res elseif inst isa PhiNode @@ -148,9 +170,8 @@ function code_mlir(f, types) elseif inst isa GotoNode args = get_value.(collect_value_arguments(ir, block_id, inst.label)) dest = blocks[inst.label] - loc = Location(string(line.file), line.line, 0) - brop = LLVM.version() >= v"15" ? cf.br : std.br - push!(current_block, brop(dest, args; loc)) + location = Location(string(line.file), line.line, 0) + push!(current_block, cf.br(args; dest, location)) elseif inst isa GotoIfNot false_args = get_value.(collect_value_arguments(ir, block_id, inst.dest)) cond = get_value(inst.cond) @@ -160,15 +181,13 @@ function code_mlir(f, types) other_dest = blocks[other_dest] dest = blocks[inst.dest] - loc = Location(string(line.file), line.line, 0) - cond_brop = LLVM.version() >= v"15" ? cf.cond_br : std.cond_br - cond_br = cond_brop(cond, other_dest, dest, true_args, false_args; loc) + location = Location(string(line.file), line.line, 0) + cond_br = cf.cond_br(cond, true_args, false_args; trueDest=other_dest, falseDest=dest, location) push!(current_block, cond_br) elseif inst isa ReturnNode line = ir.linetable[stmt[:line]] - retop = LLVM.version() >= v"15" ? func.return_ : std.return_ - loc = Location(string(line.file), line.line, 0) - push!(current_block, retop([get_value(inst.val)]; loc)) + location = Location(string(line.file), line.line, 0) + push!(current_block, func.return_([get_value(inst.val)]; location)) elseif Meta.isexpr(inst, :code_coverage_effect) # Skip else @@ -194,13 +213,13 @@ function code_mlir(f, types) ftype = MLIRType(input_types => result_types) op = IR.create_operation( - LLVM15 ? "func.func" : "builtin.func", + "func.func", Location(); - attributes = [ + attributes=[ NamedAttribute("sym_name", IR.Attribute(string(func_name))), - NamedAttribute(LLVM15 ? "function_type" : "type", IR.Attribute(ftype)), + NamedAttribute("function_type", IR.Attribute(ftype)), ], - owned_regions = Region[region], + owned_regions=Region[region], result_inference=false, ) @@ -227,16 +246,18 @@ macro code_mlir(call) end end +export code_mlir, @code_mlir + end # module Brutus # --- function pow(x::F, n) where {F} - p = one(F) - for _ in 1:n - p *= x - end - p + p = one(F) + for _ in 1:n + p *= x + end + p end function f(x) @@ -252,29 +273,28 @@ end using Test using MLIR.IR, MLIR -ctx = Context() -# IR.enable_multithreading!(ctx, false) +fptr = IR.context!(IR.Context()) do + op = Brutus.code_mlir(pow, Tuple{Int,Int}) -op = Brutus.code_mlir(pow, Tuple{Int, Int}) + mod = MModule(Location()) + body = IR.get_body(mod) + push!(body, op) -mod = MModule(Location()) -body = IR.get_body(mod) -push!(body, op) + pm = IR.PassManager() + opm = IR.OpPassManager(pm) -pm = IR.PassManager() -opm = IR.OpPassManager(pm) + # IR.enable_ir_printing!(pm) + IR.enable_verifier!(pm, true) -# IR.enable_ir_printing!(pm) -IR.enable_verifier!(pm, true) + MLIR.API.mlirRegisterAllPasses() + MLIR.API.mlirRegisterAllLLVMTranslations(IR.context()) + IR.add_pipeline!(opm, "convert-arith-to-llvm,convert-func-to-llvm") -MLIR.API.mlirRegisterAllPasses() -MLIR.API.mlirRegisterAllLLVMTranslations(ctx) -IR.add_pipeline!(opm, Brutus.LLVM.version() >= v"15" ? "convert-arith-to-llvm,convert-func-to-llvm" : "convert-std-to-llvm") + IR.run!(pm, mod) -IR.run!(pm, mod) - -jit = MLIR.API.mlirExecutionEngineCreate(mod, 0, 0, C_NULL) -fptr = MLIR.API.mlirExecutionEngineLookup(jit, "pow") + jit = MLIR.API.mlirExecutionEngineCreate(mod, 0, 0, C_NULL) + MLIR.API.mlirExecutionEngineLookup(jit, "pow") +end x, y = 3, 4 diff --git a/src/Dialects.jl b/src/Dialects.jl index b5a9f124..591ac8c4 100644 --- a/src/Dialects.jl +++ b/src/Dialects.jl @@ -11,13 +11,7 @@ function namedattribute(name, val::NamedAttribute) return val end -operandsegmentsizes(segments) = namedattribute( - "operand_segment_sizes", - Attribute(API.mlirDenseI32ArrayGet( - context().context, - length(segments), - Int32.(segments) - ))) +operandsegmentsizes(segments) = namedattribute("operand_segment_sizes", Attribute(Int32.(segments))) let ver = string(LLVM.version().major) diff --git a/src/IR/IR.jl b/src/IR/IR.jl index 35ea6133..eff2e367 100644 --- a/src/IR/IR.jl +++ b/src/IR/IR.jl @@ -217,37 +217,6 @@ function Base.eltype(type::MLIRType) end end -function show_inner(io::IO, type::MLIRType) - if API.mlirTypeIsAInteger(type) - is_signless = API.mlirIntegerTypeIsSignless(type) - is_signed = API.mlirIntegerTypeIsSigned(type) - - width = API.mlirIntegerTypeGetWidth(type) - t = if is_signed - "si" - elseif is_signless - "i" - else - "u" - end - print(io, t, width) - elseif API.mlirTypeIsAF64(type) - print(io, "f64") - elseif API.mlirTypeIsAF32(type) - print(io, "f32") - elseif API.mlirTypeIsARankedTensor(type) - print(io, "tensor<") - s = size(type) - print(io, join(s, "x"), "x") - show_inner(io, eltype(type)) - print(io, ">") - elseif API.mlirTypeIsAIndex(type) - print(io, "index") - else - print(io, "unknown") - end -end - function Base.show(io::IO, type::MLIRType) print(io, "MLIRType(#= ") c_print_callback = @cfunction(print_callback, Cvoid, (MlirStringRef, Any)) diff --git a/test/examples.jl b/test/examples.jl index a2e3da49..4864eab6 100644 --- a/test/examples.jl +++ b/test/examples.jl @@ -13,8 +13,10 @@ end examples_dir = joinpath(@__DIR__, "..", "examples") examples = find_sources(examples_dir) + filter!(file -> readline(file) != "# EXCLUDE FROM TESTING", examples) filter!(file -> !occursin("Kaleidoscope", file), examples) +filter!(file -> VERSION >= v"1.10" || !contains(file, "brutus.jl"), examples) cd(examples_dir) do examples = relpath.(examples, Ref(examples_dir)) @@ -28,4 +30,4 @@ cd(examples_dir) do end end -end \ No newline at end of file +end From b3fd104da75fe637c90300f360f679d7fc52ad25 Mon Sep 17 00:00:00 2001 From: Paul Berg Date: Thu, 11 Jan 2024 18:11:26 +0100 Subject: [PATCH 5/5] use sizeof(::String) instead of length(::String) for conversion to MlirStringRef MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit otherwise, we can't use 🍕 as a string attribute. --- src/MLIR.jl | 2 +- test/runtests.jl | 9 +++++++++ 2 files changed, 10 insertions(+), 1 deletion(-) diff --git a/src/MLIR.jl b/src/MLIR.jl index 5b1f7fea..d13ff806 100644 --- a/src/MLIR.jl +++ b/src/MLIR.jl @@ -33,7 +33,7 @@ Base.cconvert(::Type{API.MlirStringRef}, s::AbstractString) = # Directly create `MlirStringRef` instead of adding an extra ccall. function Base.unsafe_convert(::Type{API.MlirStringRef}, s::Union{Symbol, String, AbstractVector{UInt8}}) p = Base.unsafe_convert(Ptr{Cchar}, s) - return API.MlirStringRef(p, length(s)) + return API.MlirStringRef(p, sizeof(s)) end module IR diff --git a/test/runtests.jl b/test/runtests.jl index e0d9aaba..9b482217 100644 --- a/test/runtests.jl +++ b/test/runtests.jl @@ -3,3 +3,12 @@ using Test include("examples.jl") include("executionengine.jl") + +@testset "MlirStringRef conversion" begin + s = "mlir 😄 α γ 🍕" + + ms = Base.unsafe_convert(MLIR.API.MlirStringRef, s) + reconstructed = unsafe_string(Ptr{Cchar}(ms.data), ms.length) + + @test s == reconstructed +end