1//===-- CompilationAttrInterfaces.td - GPU compilation interfaces ---------===// 2// 3// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. 4// See https://llvm.org/LICENSE.txt for license information. 5// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception 6// 7//===----------------------------------------------------------------------===// 8// 9// This file defines interfaces for GPU compilation attributes. 10// 11//===----------------------------------------------------------------------===// 12 13#ifndef GPU_COMPILATIONATTRINTERFACES 14#define GPU_COMPILATIONATTRINTERFACES 15 16include "mlir/IR/AttrTypeBase.td" 17include "mlir/IR/OpBase.td" 18 19//===----------------------------------------------------------------------===// 20// GPU target attribute interface. 21//===----------------------------------------------------------------------===// 22def GPUTargetAttrInterface : AttrInterface<"TargetAttrInterface"> { 23 let description = [{ 24 Interface for GPU target attributes. Attributes implementing this interface 25 compile GPU modules into binary objects, providing an opaque interface to 26 hide implementation details. 27 }]; 28 let cppNamespace = "::mlir::gpu"; 29 let methods = [ 30 InterfaceMethod<[{ 31 Serializes a GPU module to a string containing a representation of the 32 module. 33 34 If serialization fails then the method should return `std::nullopt`. 35 36 The `module` parameter must be a GPU Module Op. The `options` parameter 37 is meant to be used for passing additional options that are not in the 38 attribute. 39 }], 40 "std::optional<::mlir::SmallVector<char, 0>>", "serializeToObject", 41 (ins "::mlir::Operation*":$module, 42 "const ::mlir::gpu::TargetOptions&":$options)>, 43 InterfaceMethod<[{ 44 Creates a GPU object attribute from a binary string. 45 46 The `module` parameter must be a `GPUModuleOp` and can be used to 47 retrieve additional information like the list of kernels in the binary. 48 The `object` parameter is a binary string. The `options` parameter is 49 meant to be used for passing additional options that are not in the 50 attribute. 51 }], "::mlir::Attribute", "createObject", 52 (ins "::mlir::Operation *":$module, 53 "const ::llvm::SmallVector<char, 0> &":$object, 54 "const ::mlir::gpu::TargetOptions &":$options)> 55 ]; 56} 57 58def GPUTargetAttr : 59 ConfinedAttr<AnyAttr, [PromisedAttrInterface<GPUTargetAttrInterface>]> { 60 let description = [{ 61 Generic GPU target attribute. These attributes must implement or promise 62 the `GPUTargetAttrInterface` interface. 63 }]; 64} 65 66def GPUTargetArrayAttr : TypedArrayAttrBase<GPUTargetAttr, 67 "array of GPU target attributes">; 68 69def GPUNonEmptyTargetArrayAttr : 70 ConfinedAttr<GPUTargetArrayAttr, [ArrayMinCount<1>]>; 71 72//===----------------------------------------------------------------------===// 73// GPU offloading translation attribute trait. 74//===----------------------------------------------------------------------===// 75 76def OffloadingTranslationAttrTrait : 77 NativeTrait<"OffloadingTranslationAttrTrait", ""> { 78 let cppNamespace = "::mlir::gpu"; 79} 80 81def HasOffloadingTranslationAttrTrait : AttrConstraint< 82 CPred<"$_self.hasTrait<::mlir::gpu::OffloadingTranslationAttrTrait>()">, 83 "with the `OffloadingTranslationAttrTrait` trait." 84>; 85 86def OffloadingTranslationAttr : 87 ConfinedAttr<AnyAttr, [HasOffloadingTranslationAttrTrait]> { 88 let description = [{ 89 Generic GPU offloading translation attribute. These attributes must 90 implement an interface for handling the translation of GPU offloading 91 operations like `gpu.binary` & `gpu.launch_func`. An example of such 92 interface is the `OffloadingLLVMTranslationAttrInterface` interface. 93 }]; 94} 95 96//===----------------------------------------------------------------------===// 97// GPU offloading LLVM translation handler attribute interface. 98//===----------------------------------------------------------------------===// 99 100def OffloadingLLVMTranslationAttrInterface : 101 AttrInterface<"OffloadingLLVMTranslationAttrInterface"> { 102 let description = [{ 103 Interface for GPU offloading LLVM translation attributes. Attributes 104 implementing this interface manage the interaction between GPU offloading 105 operations and host IR. 106 }]; 107 let cppNamespace = "::mlir::gpu"; 108 let methods = [ 109 InterfaceMethod<[{ 110 Translates a `gpu.binary` Op into a sequence of LLVM IR target-specific 111 instructions, embedding the binary into a host LLVM module. 112 113 The LLVM translation mechanism invokes this function when translating a 114 `gpu.binary`. 115 116 The first argument has to be a GPU binary operation. 117 If the function fails at any point, it must return `failure`. 118 }], 119 "::llvm::LogicalResult", "embedBinary", 120 (ins "::mlir::Operation*":$binaryOp, 121 "::llvm::IRBuilderBase&":$hostBuilder, 122 "::mlir::LLVM::ModuleTranslation&":$hostModuleTranslation) 123 >, 124 InterfaceMethod<[{ 125 Translates a `gpu.launch_func` op into a sequence of LLVM IR 126 target-specific instructions, resulting in a kernel launch on host IR. 127 128 The LLVM translation mechanism invokes this function when translating a 129 `gpu.launch_func` operation; it searches the appropriate binary and uses 130 its offloading handler. 131 132 The first two arguments must be GPU launch and binary operations, 133 respectively. If the function fails at any point, it must return 134 `failure`. 135 }], 136 "::llvm::LogicalResult", "launchKernel", 137 (ins "::mlir::Operation*":$launchFunc, "::mlir::Operation*":$binaryOp, 138 "::llvm::IRBuilderBase&":$hostBuilder, 139 "::mlir::LLVM::ModuleTranslation&":$hostModuleTranslation) 140 > 141 ]; 142} 143 144#endif // GPU_COMPILATIONATTRINTERFACES 145