//===- TargetAndABI.td - SPIR-V Target and ABI definitions -*- tablegen -*-===//
//
// Part of the LLVM Project, 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
//
//===----------------------------------------------------------------------===//
//
// This is the base file for supporting lowering to SPIR-V dialect. This file
// defines SPIR-V attributes used for specifying the shader interface or ABI.
// This is because SPIR-V module is expected to work in an execution environment
// as specified by a client API. A SPIR-V module needs to "link" correctly with
// the execution environment regarding the resources that are used in the SPIR-V
// module and get populated with data via the client API. The shader interface
// (or ABI) is passed into SPIR-V lowering path via attributes defined in this
// file. A compilation flow targeting SPIR-V is expected to attach such
// attributes to resources and other suitable places.
//
//===----------------------------------------------------------------------===//

#ifndef MLIR_DIALECT_SPIRV_IR_TARGET_AND_ABI
#define MLIR_DIALECT_SPIRV_IR_TARGET_AND_ABI

include "mlir/Dialect/SPIRV/IR/SPIRVBase.td"

// For entry functions, this attribute specifies information related to entry
// points in the generated SPIR-V module:
// 1) WorkGroup Size.
def SPV_EntryPointABIAttr : StructAttr<"EntryPointABIAttr", SPIRV_Dialect, [
    StructFieldAttr<"local_size", I32ElementsAttr>
]>;

def SPV_ExtensionArrayAttr : TypedArrayAttrBase<
    SPV_ExtensionAttr, "SPIR-V extension array attribute">;

def SPV_CapabilityArrayAttr : TypedArrayAttrBase<
    SPV_CapabilityAttr, "SPIR-V capability array attribute">;

// Description of cooperative matrix operations supported on the
// target. Represents `VkCooperativeMatrixPropertiesNV`. See
// https://www.khronos.org/registry/vulkan/specs/1.2-extensions/man/html/VkCooperativeMatrixPropertiesNV.html
def SPV_CooperativeMatrixPropertiesNVAttr :
    StructAttr<"CooperativeMatrixPropertiesNVAttr", SPIRV_Dialect, [
    StructFieldAttr<"m_size", I32Attr>,
    StructFieldAttr<"n_size", I32Attr>,
    StructFieldAttr<"k_size", I32Attr>,
    StructFieldAttr<"a_type", TypeAttr>,
    StructFieldAttr<"b_type", TypeAttr>,
    StructFieldAttr<"c_type", TypeAttr>,
    StructFieldAttr<"result_type", TypeAttr>,
    StructFieldAttr<"scope", SPV_ScopeAttr>
]>;

def SPV_CooperativeMatrixPropertiesNVArrayAttr :
    TypedArrayAttrBase<SPV_CooperativeMatrixPropertiesNVAttr,
                       "CooperativeMatrixPropertiesNV array attribute">;

// This attribute specifies the limits for various resources on the target
// architecture.
//
// See https://www.khronos.org/registry/vulkan/specs/1.2-extensions/html/vkspec.html#limits
// for the complete list of limits and their explanation for the Vulkan API.
// The following ones are those affecting SPIR-V CodeGen. Their default value
// are the from Vulkan limit requirements:
// https://www.khronos.org/registry/vulkan/specs/1.2-extensions/html/vkspec.html#limits-minmax
def SPV_ResourceLimitsAttr : StructAttr<"ResourceLimitsAttr", SPIRV_Dialect, [
    // The maximum total storage size, in bytes, available for variables
    // declared with the Workgroup storage class.
    StructFieldAttr<"max_compute_shared_memory_size",
                    DefaultValuedAttr<I32Attr, "16384">>,

    // The maximum total number of compute shader invocations in a single local
    // workgroup.
    StructFieldAttr<"max_compute_workgroup_invocations",
                    DefaultValuedAttr<I32Attr, "128">>,
    // The maximum size of a local compute workgroup, per dimension.
    StructFieldAttr<"max_compute_workgroup_size",
                    DefaultValuedAttr<I32ElementsAttr, "{128, 128, 64}">>,

    // The default number of invocations in each subgroup.
    // 0x7FFFFFFF means unknown.
    StructFieldAttr<"subgroup_size", DefaultValuedAttr<I32Attr, "0x7FFFFFFF">>,

    // The configurations of cooperative matrix operations
    // supported. Default is an empty list.
    StructFieldAttr<
        "cooperative_matrix_properties_nv",
        DefaultValuedAttr<SPV_CooperativeMatrixPropertiesNVArrayAttr, "{}">>
]>;

#endif // MLIR_DIALECT_SPIRV_IR_TARGET_AND_ABI