| /external/llvm-project/mlir/lib/Conversion/GPUToVulkan/ |
| D | ConvertGPULaunchFuncToVulkanLaunchFunc.cpp | 53 void convertGpuLaunchFunc(gpu::LaunchFuncOp launchOp);
69 gpu::LaunchFuncOp launchOp);
101 Location loc, gpu::LaunchFuncOp launchOp) { in declareVulkanLaunchFunc() argument
106 SmallVector<Type, 8> gpuLaunchTypes(launchOp.getOperandTypes()); in declareVulkanLaunchFunc()
121 return launchOp.emitError() << type << " is unsupported to run on Vulkan"; in declareVulkanLaunchFunc()
150 gpu::LaunchFuncOp launchOp) { in convertGpuLaunchFunc() argument
152 OpBuilder builder(launchOp); in convertGpuLaunchFunc()
153 Location loc = launchOp.getLoc(); in convertGpuLaunchFunc()
161 if (failed(declareVulkanLaunchFunc(loc, launchOp))) in convertGpuLaunchFunc()
164 SmallVector<Value, 8> gpuLaunchOperands(launchOp.getOperands()); in convertGpuLaunchFunc()
[all …]
|
| /external/llvm-project/mlir/lib/Dialect/GPU/Transforms/ |
| D | KernelOutlining.cpp | 106 LogicalResult mlir::sinkOperationsIntoLaunchOp(gpu::LaunchOp launchOp) { in sinkOperationsIntoLaunchOp() argument
107 Region &launchOpBody = launchOp.body(); in sinkOperationsIntoLaunchOp()
131 launchOp.body()); in sinkOperationsIntoLaunchOp()
138 static gpu::GPUFuncOp outlineKernelFuncImpl(gpu::LaunchOp launchOp, in outlineKernelFuncImpl() argument
141 Location loc = launchOp.getLoc(); in outlineKernelFuncImpl()
144 OpBuilder builder(launchOp.getContext()); in outlineKernelFuncImpl()
145 Region &launchOpBody = launchOp.body(); in outlineKernelFuncImpl()
158 FunctionType::get(kernelOperandTypes, {}, launchOp.getContext()); in outlineKernelFuncImpl()
197 gpu::GPUFuncOp mlir::outlineKernelFunc(gpu::LaunchOp launchOp, in outlineKernelFunc() argument
203 auto funcOp = outlineKernelFuncImpl(launchOp, kernelFnName, operandSet); in outlineKernelFunc()
[all …]
|
| /external/llvm-project/mlir/lib/Conversion/GPUCommon/ |
| D | ConvertLaunchFuncToRuntimeCalls.cpp | 248 Value generateParamsArray(gpu::LaunchFuncOp launchOp,
254 matchAndRewrite(gpu::LaunchFuncOp launchOp, ArrayRef<Value> operands,
485 gpu::LaunchFuncOp launchOp, ArrayRef<Value> operands, in generateParamsArray() argument
487 auto loc = launchOp.getLoc(); in generateParamsArray()
488 auto numKernelOperands = launchOp.getNumKernelOperands(); in generateParamsArray()
490 loc, launchOp.getOperands().take_back(numKernelOperands), in generateParamsArray()
566 gpu::LaunchFuncOp launchOp, ArrayRef<Value> operands, in matchAndRewrite() argument
568 if (failed(areAllLLVMTypes(launchOp, operands, rewriter))) in matchAndRewrite()
571 if (launchOp.asyncDependencies().size() > 1) in matchAndRewrite()
573 launchOp, "Cannot convert with more than one async dependency."); in matchAndRewrite()
[all …]
|
| /external/llvm-project/mlir/lib/Conversion/SCFToGPU/ |
| D | SCFToGPU.cpp | 218 auto launchOp = builder.create<gpu::LaunchOp>( in createLaunch() local
231 launchOp.body().front().getOperations().splice( in createLaunch()
232 launchOp.body().front().begin(), in createLaunch()
239 builder.setInsertionPointToStart(&launchOp.body().front()); in createLaunch()
245 ? getDim3Value(launchOp.getBlockIds(), en.index()) in createLaunch()
246 : getDim3Value(launchOp.getThreadIds(), en.index() - numBlockDims); in createLaunch()
378 ParallelOp parallelOp, gpu::LaunchOp launchOp, in processParallelLoop() argument
392 auto launchIndependent = [&launchOp](Value val) { in processParallelLoop()
393 return val.getParentRegion()->isAncestor(launchOp->getParentRegion()); in processParallelLoop()
421 launchOp.body().getArgument(getLaunchOpArgumentNum(processor)); in processParallelLoop()
[all …]
|
| /external/llvm-project/mlir/lib/Conversion/SPIRVToLLVM/ |
| D | ConvertLaunchFuncToLLVMCalls.cpp | 154 matchAndRewrite(gpu::LaunchFuncOp launchOp, ArrayRef<Value> operands, in matchAndRewrite() argument
156 auto *op = launchOp.getOperation(); in matchAndRewrite()
158 auto module = launchOp->getParentOfType<ModuleOp>(); in matchAndRewrite()
164 StringRef kernelModuleName = launchOp.getKernelModuleName(); in matchAndRewrite()
168 return launchOp.emitOpError("SPIR-V kernel module '") in matchAndRewrite()
178 StringRef kernelFuncName = launchOp.getKernelName(); in matchAndRewrite()
189 rewriter.setInsertionPoint(launchOp); in matchAndRewrite()
199 Location loc = launchOp.getLoc(); in matchAndRewrite()
201 auto numKernelOperands = launchOp.getNumKernelOperands(); in matchAndRewrite()
205 auto memRefType = launchOp.getKernelOperand(operand.index()) in matchAndRewrite()
[all …]
|
| /external/tensorflow/tensorflow/compiler/xla/service/mlir_gpu/ |
| D | passes.cc | 211 getFunction().walk([&](mlir::gpu::LaunchFuncOp launchOp) { in runOnFunction() argument
213 module.lookupSymbol<mlir::gpu::GPUFuncOp>(launchOp.kernel()); in runOnFunction()
232 for (int i = 0, e = launchOp.getNumKernelOperands(); i < e; ++i) { in runOnFunction()
233 if (launchOp.getKernelOperand(i) == arg) { in runOnFunction()
242 for (int i = 0, e = launchOp.getNumKernelOperands(); i < e; ++i) { in runOnFunction()
243 if (launchOp.getKernelOperand(i) == result) { in runOnFunction()
303 mlir::OpBuilder launch_builder(launchOp); in runOnFunction()
305 launchOp.getLoc(), new_kernel, launchOp.getGridSizeOperandValues(), in runOnFunction()
306 launchOp.getBlockSizeOperandValues(), new_operands); in runOnFunction()
309 launchOp.erase(); in runOnFunction()
|
| /external/llvm-project/mlir/include/mlir/Dialect/GPU/ |
| D | Utils.h | 35 gpu::GPUFuncOp outlineKernelFunc(gpu::LaunchOp launchOp, StringRef kernelFnName,
41 LogicalResult sinkOperationsIntoLaunchOp(gpu::LaunchOp launchOp);
|
| /external/llvm-project/mlir/lib/Dialect/GPU/IR/ |
| D | GPUDialect.cpp | 80 auto walkResult = module.walk([&module](LaunchFuncOp launchOp) -> WalkResult { in verifyOperationAttribute() argument
83 if (!launchOp->getParentOp() || in verifyOperationAttribute()
84 launchOp->getParentOp()->getParentOp() != module) in verifyOperationAttribute()
89 if (!launchOp->getAttrOfType<SymbolRefAttr>( in verifyOperationAttribute()
94 StringRef kernelModuleName = launchOp.getKernelModuleName(); in verifyOperationAttribute()
97 return launchOp.emitOpError() in verifyOperationAttribute()
101 Operation *kernelFunc = module.lookupSymbol(launchOp.kernel()); in verifyOperationAttribute()
105 return launchOp.emitOpError("kernel function '") in verifyOperationAttribute()
106 << launchOp.kernel() << "' is undefined"; in verifyOperationAttribute()
109 return launchOp.emitOpError("kernel function is missing the '") in verifyOperationAttribute()
[all …]
|