1//===-- GPUBase.td - GPU dialect definitions ---------------*- tablegen -*-===// 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// Defines the GPU dialect 10// 11//===----------------------------------------------------------------------===// 12 13#ifndef GPU_BASE 14#define GPU_BASE 15 16include "mlir/IR/OpBase.td" 17 18//===----------------------------------------------------------------------===// 19// GPU Dialect. 20//===----------------------------------------------------------------------===// 21 22def GPU_Dialect : Dialect { 23 let name = "gpu"; 24 let cppNamespace = "::mlir::gpu"; 25 let hasOperationAttrVerify = 1; 26 27 let extraClassDeclaration = [{ 28 /// Get the name of the attribute used to annotate the modules that contain 29 /// kernel modules. 30 static StringRef getContainerModuleAttrName() { 31 return "gpu.container_module"; 32 } 33 /// Get the name of the attribute used to annotate external kernel 34 /// functions. 35 static StringRef getKernelFuncAttrName() { return "gpu.kernel"; } 36 37 /// Returns whether the given function is a kernel function, i.e., has the 38 /// 'gpu.kernel' attribute. 39 static bool isKernel(Operation *op); 40 41 /// Returns the number of workgroup (thread, block) dimensions supported in 42 /// the GPU dialect. 43 // TODO: consider generalizing this. 44 static unsigned getNumWorkgroupDimensions() { return 3; } 45 46 /// Returns the numeric value used to identify the workgroup memory address 47 /// space. 48 static unsigned getWorkgroupAddressSpace() { return 3; } 49 50 /// Returns the numeric value used to identify the private memory address 51 /// space. 52 static unsigned getPrivateAddressSpace() { return 5; } 53 }]; 54} 55 56def GPU_AsyncToken : DialectType< 57 GPU_Dialect, CPred<"$_self.isa<::mlir::gpu::AsyncTokenType>()">, "async token type">, 58 BuildableType<"mlir::gpu::AsyncTokenType::get($_builder.getContext())">; 59 60def GPU_AsyncOpInterface : OpInterface<"AsyncOpInterface"> { 61 let description = [{ 62 Interface for GPU operations that execute asynchronously on the device. 63 64 GPU operations implementing this interface take a list of dependencies 65 as `gpu.async.token` arguments and optionally return a `gpu.async.token`. 66 67 The op doesn't start executing until all depent ops producing the async 68 dependency tokens have finished executing. 69 70 If the op returns a token, the op merely schedules the execution on the 71 device and returns immediately, without waiting for the execution to 72 complete. On the hand, if the op does not return a token, the op will wait 73 for the execution to complete. 74 }]; 75 let cppNamespace = "::mlir::gpu"; 76 77 let methods = [ 78 InterfaceMethod<[{ 79 Query the operands that represent async dependency tokens. 80 }], 81 "OperandRange", "getAsyncDependencies", (ins), [{}], [{ 82 ConcreteOp op = cast<ConcreteOp>(this->getOperation()); 83 return op.asyncDependencies(); 84 }] 85 >, 86 InterfaceMethod<[{ 87 Adds a new token to the list of async dependencies. 88 }], 89 "void", "addAsyncDependency", (ins "Value":$token), 90 [{}], [{ 91 ::mlir::gpu::addAsyncDependency(this->getOperation(), token); 92 }] 93 >, 94 InterfaceMethod<[{ 95 Query the result that represents the async token to depend on. 96 }], 97 "OpResult", "getAsyncToken", (ins), [{}], [{ 98 ConcreteOp op = cast<ConcreteOp>(this->getOperation()); 99 return op.asyncToken().template dyn_cast_or_null<OpResult>(); 100 }] 101 > 102 ]; 103} 104 105#endif // GPU_BASE 106