Home
last modified time | relevance | path

Searched refs:GPUFuncOp (Results 1 – 12 of 12) sorted by relevance

/external/llvm-project/mlir/lib/Dialect/GPU/IR/
DGPUDialect.cpp102 auto kernelGPUFunction = dyn_cast_or_null<gpu::GPUFuncOp>(kernelFunc); in verifyOperationAttribute()
431 GPUFuncOp kernelFunc, KernelDim3 gridSize, in build()
524 BlockArgument GPUFuncOp::addWorkgroupAttribution(Type type) { in addWorkgroupAttribution()
534 BlockArgument GPUFuncOp::addPrivateAttribution(Type type) { in addPrivateAttribution()
540 void GPUFuncOp::build(OpBuilder &builder, OperationState &result, in build()
629 result.addAttribute(GPUFuncOp::getTypeAttrName(), TypeAttr::get(type)); in parseGPUFuncOp()
632 if (failed(parseAttributions(parser, GPUFuncOp::getWorkgroupKeyword(), in parseGPUFuncOp()
639 result.addAttribute(GPUFuncOp::getNumWorkgroupAttributionsAttrName(), in parseGPUFuncOp()
643 if (failed(parseAttributions(parser, GPUFuncOp::getPrivateKeyword(), in parseGPUFuncOp()
648 if (succeeded(parser.parseOptionalKeyword(GPUFuncOp::getKernelKeyword()))) in parseGPUFuncOp()
[all …]
/external/llvm-project/mlir/include/mlir/Dialect/GPU/
DMemoryPromotion.h20 class GPUFuncOp; variable
25 void promoteToWorkgroupMemory(gpu::GPUFuncOp op, unsigned arg);
DUtils.h23 class GPUFuncOp; variable
35 gpu::GPUFuncOp outlineKernelFunc(gpu::LaunchOp launchOp, StringRef kernelFnName,
DGPUOps.td271 friend class OpTrait::FunctionLike<GPUFuncOp>;
380 OpBuilderDAG<(ins "GPUFuncOp":$kernelFunc, "KernelDim3":$gridSize,
536 def GPU_ReturnOp : GPU_Op<"return", [HasParent<"GPUFuncOp">, NoSideEffect,
/external/llvm-project/mlir/lib/Dialect/GPU/Transforms/
DKernelOutlining.cpp138 static gpu::GPUFuncOp outlineKernelFuncImpl(gpu::LaunchOp launchOp, in outlineKernelFuncImpl()
159 auto outlinedFunc = builder.create<gpu::GPUFuncOp>(loc, kernelFnName, type); in outlineKernelFuncImpl()
197 gpu::GPUFuncOp mlir::outlineKernelFunc(gpu::LaunchOp launchOp, in outlineKernelFunc()
215 gpu::GPUFuncOp kernelFunc, in convertToLaunchFuncOp()
251 gpu::GPUFuncOp outlinedFunc = in runOnOperation()
278 gpu::GPUModuleOp createKernelModule(gpu::GPUFuncOp kernelFunc, in createKernelModule()
DAllReduceLowering.cpp29 GpuAllReduceRewriter(gpu::GPUFuncOp funcOp_, gpu::AllReduceOp reduceOp_, in GpuAllReduceRewriter()
368 gpu::GPUFuncOp funcOp;
382 : RewritePattern(gpu::GPUFuncOp::getOperationName(), 1, context) {} in GpuAllReduceConversion()
386 auto funcOp = cast<gpu::GPUFuncOp>(op); in matchAndRewrite()
DMemoryPromotion.cpp159 void mlir::promoteToWorkgroupMemory(GPUFuncOp op, unsigned arg) { in promoteToWorkgroupMemory()
/external/llvm-project/mlir/test/lib/Transforms/
DTestGpuMemoryPromotion.cpp31 OperationPass<gpu::GPUFuncOp>> {
37 gpu::GPUFuncOp op = getOperation(); in runOnOperation()
/external/llvm-project/mlir/lib/Conversion/GPUToSPIRV/
DConvertGPUToSPIRV.cpp63 class GPUFuncOpConversion final : public SPIRVOpLowering<gpu::GPUFuncOp> {
65 using SPIRVOpLowering<gpu::GPUFuncOp>::SPIRVOpLowering;
68 matchAndRewrite(gpu::GPUFuncOp funcOp, ArrayRef<Value> operands,
166 lowerAsEntryFunction(gpu::GPUFuncOp funcOp, SPIRVTypeConverter &typeConverter, in lowerAsEntryFunction()
218 getDefaultABIAttrs(MLIRContext *context, gpu::GPUFuncOp funcOp, in getDefaultABIAttrs()
239 gpu::GPUFuncOp funcOp, ArrayRef<Value> operands, in matchAndRewrite()
/external/llvm-project/mlir/test/lib/Dialect/SPIRV/
DTestEntryPointAbi.cpp46 for (gpu::GPUFuncOp gpuFunc : gpuModule.getOps<gpu::GPUFuncOp>()) { in runOnOperation()
/external/llvm-project/mlir/lib/Conversion/GPUCommon/
DGPUOpsLowering.h23 : ConvertToLLVMPattern(gpu::GPUFuncOp::getOperationName(), in GPUFuncOpLowering()
31 auto gpuFuncOp = cast<gpu::GPUFuncOp>(op); in matchAndRewrite()
73 attr.first == gpu::GPUFuncOp::getNumWorkgroupAttributionsAttrName()) in matchAndRewrite()
/external/tensorflow/tensorflow/compiler/xla/service/mlir_gpu/
Dpasses.cc212 mlir::gpu::GPUFuncOp kernel = in runOnFunction()
213 module.lookupSymbol<mlir::gpu::GPUFuncOp>(launchOp.kernel()); in runOnFunction()
257 auto new_kernel = kernel_builder.create<mlir::gpu::GPUFuncOp>( in runOnFunction()