blob: 381ecea7a43b8b08b7fa0fd0a9fceb7fba373ad9 [file] [log] [blame]
# Copyright 2023 The IREE Authors
#
# Licensed under the Apache License v2.0 with LLVM Exceptions.
# See https://llvm.org/LICENSE.txt for license information.
# SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
from iree.compiler import ir
# Make sure that our dialects import.
from iree.compiler.dialects import flow, hal, stream, vm, util, iree_codegen, iree_gpu
def run(fn):
with ir.Context(), ir.Location.unknown():
module = ir.Module.create()
with ir.InsertionPoint(module.body):
print("\nTEST:", fn.__name__)
fn()
return fn
# ======================================================================
# IREE Codegen Dialect
# ======================================================================
@run
def codegen_dispatch_lowering_pass_pipeline():
pipeline_attr = iree_codegen.DispatchLoweringPassPipelineAttr.get(
iree_codegen.DispatchLoweringPassPipeline.LLVMGPUTileAndFuse
)
assert pipeline_attr is not None
assert (
pipeline_attr.value
== iree_codegen.DispatchLoweringPassPipeline.LLVMGPUTileAndFuse
)
assert pipeline_attr.raw_value == int(
iree_codegen.DispatchLoweringPassPipeline.LLVMGPUTileAndFuse
)
assert "LLVMGPUTileAndFuse" in str(pipeline_attr)
@run
def codegen_translation_info_minimal():
pipeline_attr = iree_codegen.DispatchLoweringPassPipelineAttr.get(
iree_codegen.DispatchLoweringPassPipeline.None_
)
translation_info = iree_codegen.TranslationInfoAttr.get(pipeline_attr)
assert translation_info is not None
assert str(translation_info) == "#iree_codegen.translation_info<pipeline = None>"
assert translation_info.pass_pipeline == pipeline_attr
assert translation_info.codegen_spec is None
assert translation_info.workgroup_size == []
assert translation_info.subgroup_size == 0
assert translation_info.configuration is None
@run
def codegen_translation_info_with_sizes():
pipeline_attr = iree_codegen.DispatchLoweringPassPipelineAttr.get(
iree_codegen.DispatchLoweringPassPipeline.Custom
)
translation_info = iree_codegen.TranslationInfoAttr.get(
pipeline_attr, None, [64, 4, 1], 32
)
assert translation_info is not None
assert translation_info.pass_pipeline == pipeline_attr
assert translation_info.codegen_spec is None
assert translation_info.workgroup_size == [64, 4, 1]
assert translation_info.subgroup_size == 32
assert translation_info.configuration is None
@run
def codegen_translation_info_full():
pipeline_attr = iree_codegen.DispatchLoweringPassPipelineAttr.get(
iree_codegen.DispatchLoweringPassPipeline.TransformDialectCodegen
)
foo_symbol = ir.SymbolRefAttr.get(["foo"])
configuration = ir.DictAttr.get({"A": ir.IntegerAttr.get(ir.IndexType.get(), 42)})
translation_info = iree_codegen.TranslationInfoAttr.get(
pipeline_attr, foo_symbol, [128], 32, configuration
)
assert translation_info is not None
assert translation_info.pass_pipeline == pipeline_attr
assert translation_info.codegen_spec == foo_symbol
assert translation_info.workgroup_size == [128]
assert translation_info.subgroup_size == 32
assert translation_info.configuration == configuration
# ======================================================================
# IREE GPU Dialect
# ======================================================================
@run
def gpu_pipeline_options_attr():
reorder_attr = iree_gpu.ReorderWorkgroupsStrategyAttr.get(
iree_gpu.ReorderWorkgroupsStrategy.Transpose
)
assert reorder_attr.value == iree_gpu.ReorderWorkgroupsStrategy.Transpose
gpu_attr = iree_gpu.PipelineOptionsAttr.get(
True,
False,
False,
reorder_attr,
)
assert type(gpu_attr) is iree_gpu.PipelineOptionsAttr
assert gpu_attr.prefetch_shared_memory
assert not gpu_attr.no_reduce_shared_memory_bank_conflicts
assert not gpu_attr.use_igemm_convolution
gpu_attr = iree_gpu.PipelineOptionsAttr.get(
False,
True,
True,
iree_gpu.ReorderWorkgroupsStrategyAttr.get(
iree_gpu.ReorderWorkgroupsStrategy.Transpose
),
)
assert not gpu_attr.prefetch_shared_memory
assert gpu_attr.no_reduce_shared_memory_bank_conflicts
assert gpu_attr.use_igemm_convolution
gpu_attr = iree_gpu.PipelineOptionsAttr.get()
assert (
gpu_attr.prefetch_shared_memory is None
and gpu_attr.no_reduce_shared_memory_bank_conflicts is None
and gpu_attr.use_igemm_convolution is None
and gpu_attr.reorder_workgroups_strategy is None
)
gpu_attr = iree_gpu.PipelineOptionsAttr.get(True)
assert gpu_attr.prefetch_shared_memory
assert (
gpu_attr.no_reduce_shared_memory_bank_conflicts is None
and gpu_attr.use_igemm_convolution is None
and gpu_attr.reorder_workgroups_strategy is None
)
gpu_attr = iree_gpu.PipelineOptionsAttr.get(True, False)
assert (
gpu_attr.use_igemm_convolution is None
and gpu_attr.reorder_workgroups_strategy is None
)
gpu_attr = iree_gpu.PipelineOptionsAttr.get(True, False, False)
assert gpu_attr.reorder_workgroups_strategy is None
gpu_attr = iree_gpu.PipelineOptionsAttr.get(
no_reduce_shared_memory_bank_conflicts=False
)
assert (
gpu_attr.no_reduce_shared_memory_bank_conflicts is not None
and not gpu_attr.no_reduce_shared_memory_bank_conflicts
)
assert gpu_attr.prefetch_shared_memory is None
assert gpu_attr.use_igemm_convolution is None
assert gpu_attr.reorder_workgroups_strategy is None
gpu_attr = iree_gpu.PipelineOptionsAttr.get(
reorder_workgroups_strategy=reorder_attr
)
assert gpu_attr.reorder_workgroups_strategy is not None
assert (
gpu_attr.reorder_workgroups_strategy.value
# unfortunately not `is`
== iree_gpu.ReorderWorkgroupsStrategy.Transpose
)
@run
def mma_intrinsic_attr():
mma_intrinsic_attr = iree_gpu.MMAIntrinsicAttr.get(
iree_gpu.MMAIntrinsic.MFMA_F32_32x32x8_F16
)
assert mma_intrinsic_attr is not None
assert str(mma_intrinsic_attr) == "#iree_gpu<mma_intrinsic MFMA_F32_32x32x8_F16>"
raw_value = mma_intrinsic_attr.raw_value
assert raw_value == iree_gpu.MMAIntrinsic.MFMA_F32_32x32x8_F16
value = mma_intrinsic_attr.value
assert str(value) == "MFMA_F32_32x32x8_F16"
assert int(value) == raw_value
mma_attr = iree_gpu.MMAAttr.get(raw_value)
assert mma_attr is not None
f16 = ir.F16Type.get()
f32 = ir.F32Type.get()
a_type, b_type, c_type = mma_attr.abc_element_types
assert a_type == f16
assert b_type == f16
assert c_type == f32
vec_4xf16 = ir.VectorType.get((4,), f16)
a_vec_type, b_vec_type, _c_vec_type = mma_attr.abc_vector_types
assert a_vec_type == vec_4xf16
assert b_vec_type == vec_4xf16
M, N, K = mma_attr.mnk_shape
assert M == 32
assert N == 32
assert K == 8
assert mma_intrinsic_attr.mma == mma_attr
@run
def lowering_config_attr():
attributes = ir.DictAttr.get({"reduction": ir.ArrayAttr.get([])})
lowering_config = iree_gpu.LoweringConfigAttr.get(attributes)
assert lowering_config is not None
assert lowering_config.attributes == attributes
@run
def compilation_info():
attributes = ir.DictAttr.get({"reduction": ir.ArrayAttr.get([])})
lowering_config = iree_gpu.LoweringConfigAttr.get(attributes)
pipeline_attr = iree_codegen.DispatchLoweringPassPipelineAttr.get(
iree_codegen.DispatchLoweringPassPipeline.None_
)
translation_info = iree_codegen.TranslationInfoAttr.get(pipeline_attr)
compilation_info = iree_codegen.CompilationInfoAttr.get(
lowering_config, translation_info
)
assert compilation_info is not None
assert compilation_info.lowering_config == lowering_config
assert compilation_info.translation_info == translation_info