maxjhandsome
/
llvm-project
forked from OSchip/llvm-project
1
0
Fork 0
Code Issues Pull Requests Packages Releases Wiki Activity
llvm-project/mlir/test/mlir-spirv-cpu-runner/double.mlir

Ignoring revisions in .git-blame-ignore-revs. Click here to bypass and see the normal blame view.

68 lines
2.6 KiB
MLIR
Raw Normal View History

[MLIR][mlir-spirv-cpu-runner] A SPIR-V cpu runner prototype This patch introduces a SPIR-V runner. The aim is to run a gpu kernel on a CPU via GPU -> SPIRV -> LLVM conversions. This is a first prototype, so more features will be added in due time. - Overview The runner follows similar flow as the other runners in-tree. However, having converted the kernel to SPIR-V, we encode the bind attributes of global variables that represent kernel arguments. Then SPIR-V module is converted to LLVM. On the host side, we emulate passing the data to device by creating in main module globals with the same symbolic name as in kernel module. These global variables are later linked with ones from the nested module. We copy data from kernel arguments to globals, call the kernel function from nested module and then copy the data back. - Current state At the moment, the runner is capable of running 2 modules, nested one in another. The kernel module must contain exactly one kernel function. Also, the runner supports rank 1 integer memref types as arguments (to be scaled). - Enhancement of JitRunner and ExecutionEngine To translate nested modules to LLVM IR, JitRunner and ExecutionEngine were altered to take an optional (default to `nullptr`) function reference that is a custom LLVM IR module builder. This allows to customize LLVM IR module creation from MLIR modules. Reviewed By: ftynse, mravishankar Differential Revision: https://reviews.llvm.org/D86108
2020-10-23 22:46:18 +08:00
// RUN: mlir-spirv-cpu-runner %s -e main --entry-point-result=void --shared-libs=%mlir_runner_utils_dir/libmlir_runner_utils%shlibext,%spirv_wrapper_library_dir/libmlir_test_spirv_cpu_runner_c_wrappers%shlibext
// CHECK: [8, 8, 8, 8, 8, 8]
module attributes {
gpu.container_module,
spv.target_env = #spv.target_env<
#spv.vce<v1.0, [Shader], [SPV_KHR_variable_pointers]>,
{max_compute_workgroup_invocations = 128 : i32,
max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}>
} {
gpu.module @kernels {
gpu.func @double(%arg0 : memref<6xi32>, %arg1 : memref<6xi32>)
kernel attributes { spv.entry_point_abi = {local_size = dense<[1, 1, 1]>: vector<3xi32>}} {
%factor = constant 2 : i32
%i0 = constant 0 : index
%i1 = constant 1 : index
%i2 = constant 2 : index
%i3 = constant 3 : index
%i4 = constant 4 : index
%i5 = constant 5 : index
[mlir] fix SPIR-V CPU and Vulkan runners after e2310704d890ad252aeb1ca28b4b84d29514b1d1 The commit in question changed the syntax but did not update the runner tests. This also required registering the MemRef dialect for custom parser to work correctly.
2021-03-16 01:35:40 +08:00
%x0 = memref.load %arg0[%i0] : memref<6xi32>
%x1 = memref.load %arg0[%i1] : memref<6xi32>
%x2 = memref.load %arg0[%i2] : memref<6xi32>
%x3 = memref.load %arg0[%i3] : memref<6xi32>
%x4 = memref.load %arg0[%i4] : memref<6xi32>
%x5 = memref.load %arg0[%i5] : memref<6xi32>
[MLIR][mlir-spirv-cpu-runner] A SPIR-V cpu runner prototype This patch introduces a SPIR-V runner. The aim is to run a gpu kernel on a CPU via GPU -> SPIRV -> LLVM conversions. This is a first prototype, so more features will be added in due time. - Overview The runner follows similar flow as the other runners in-tree. However, having converted the kernel to SPIR-V, we encode the bind attributes of global variables that represent kernel arguments. Then SPIR-V module is converted to LLVM. On the host side, we emulate passing the data to device by creating in main module globals with the same symbolic name as in kernel module. These global variables are later linked with ones from the nested module. We copy data from kernel arguments to globals, call the kernel function from nested module and then copy the data back. - Current state At the moment, the runner is capable of running 2 modules, nested one in another. The kernel module must contain exactly one kernel function. Also, the runner supports rank 1 integer memref types as arguments (to be scaled). - Enhancement of JitRunner and ExecutionEngine To translate nested modules to LLVM IR, JitRunner and ExecutionEngine were altered to take an optional (default to `nullptr`) function reference that is a custom LLVM IR module builder. This allows to customize LLVM IR module creation from MLIR modules. Reviewed By: ftynse, mravishankar Differential Revision: https://reviews.llvm.org/D86108
2020-10-23 22:46:18 +08:00
%y0 = muli %x0, %factor : i32
%y1 = muli %x1, %factor : i32
%y2 = muli %x2, %factor : i32
%y3 = muli %x3, %factor : i32
%y4 = muli %x4, %factor : i32
%y5 = muli %x5, %factor : i32
[mlir] fix SPIR-V CPU and Vulkan runners after e2310704d890ad252aeb1ca28b4b84d29514b1d1 The commit in question changed the syntax but did not update the runner tests. This also required registering the MemRef dialect for custom parser to work correctly.
2021-03-16 01:35:40 +08:00
memref.store %y0, %arg1[%i0] : memref<6xi32>
memref.store %y1, %arg1[%i1] : memref<6xi32>
memref.store %y2, %arg1[%i2] : memref<6xi32>
memref.store %y3, %arg1[%i3] : memref<6xi32>
memref.store %y4, %arg1[%i4] : memref<6xi32>
memref.store %y5, %arg1[%i5] : memref<6xi32>
[MLIR][mlir-spirv-cpu-runner] A SPIR-V cpu runner prototype This patch introduces a SPIR-V runner. The aim is to run a gpu kernel on a CPU via GPU -> SPIRV -> LLVM conversions. This is a first prototype, so more features will be added in due time. - Overview The runner follows similar flow as the other runners in-tree. However, having converted the kernel to SPIR-V, we encode the bind attributes of global variables that represent kernel arguments. Then SPIR-V module is converted to LLVM. On the host side, we emulate passing the data to device by creating in main module globals with the same symbolic name as in kernel module. These global variables are later linked with ones from the nested module. We copy data from kernel arguments to globals, call the kernel function from nested module and then copy the data back. - Current state At the moment, the runner is capable of running 2 modules, nested one in another. The kernel module must contain exactly one kernel function. Also, the runner supports rank 1 integer memref types as arguments (to be scaled). - Enhancement of JitRunner and ExecutionEngine To translate nested modules to LLVM IR, JitRunner and ExecutionEngine were altered to take an optional (default to `nullptr`) function reference that is a custom LLVM IR module builder. This allows to customize LLVM IR module creation from MLIR modules. Reviewed By: ftynse, mravishankar Differential Revision: https://reviews.llvm.org/D86108
2020-10-23 22:46:18 +08:00
gpu.return
}
}
func @main() {
[mlir] fix SPIR-V CPU and Vulkan runners after e2310704d890ad252aeb1ca28b4b84d29514b1d1 The commit in question changed the syntax but did not update the runner tests. This also required registering the MemRef dialect for custom parser to work correctly.
2021-03-16 01:35:40 +08:00
%input = memref.alloc() : memref<6xi32>
%output = memref.alloc() : memref<6xi32>
[MLIR][mlir-spirv-cpu-runner] A SPIR-V cpu runner prototype This patch introduces a SPIR-V runner. The aim is to run a gpu kernel on a CPU via GPU -> SPIRV -> LLVM conversions. This is a first prototype, so more features will be added in due time. - Overview The runner follows similar flow as the other runners in-tree. However, having converted the kernel to SPIR-V, we encode the bind attributes of global variables that represent kernel arguments. Then SPIR-V module is converted to LLVM. On the host side, we emulate passing the data to device by creating in main module globals with the same symbolic name as in kernel module. These global variables are later linked with ones from the nested module. We copy data from kernel arguments to globals, call the kernel function from nested module and then copy the data back. - Current state At the moment, the runner is capable of running 2 modules, nested one in another. The kernel module must contain exactly one kernel function. Also, the runner supports rank 1 integer memref types as arguments (to be scaled). - Enhancement of JitRunner and ExecutionEngine To translate nested modules to LLVM IR, JitRunner and ExecutionEngine were altered to take an optional (default to `nullptr`) function reference that is a custom LLVM IR module builder. This allows to customize LLVM IR module creation from MLIR modules. Reviewed By: ftynse, mravishankar Differential Revision: https://reviews.llvm.org/D86108
2020-10-23 22:46:18 +08:00
%four = constant 4 : i32
%zero = constant 0 : i32
[mlir] fix SPIR-V CPU and Vulkan runners after e2310704d890ad252aeb1ca28b4b84d29514b1d1 The commit in question changed the syntax but did not update the runner tests. This also required registering the MemRef dialect for custom parser to work correctly.
2021-03-16 01:35:40 +08:00
%input_casted = memref.cast %input : memref<6xi32> to memref<?xi32>
%output_casted = memref.cast %output : memref<6xi32> to memref<?xi32>
[MLIR][mlir-spirv-cpu-runner] A SPIR-V cpu runner prototype This patch introduces a SPIR-V runner. The aim is to run a gpu kernel on a CPU via GPU -> SPIRV -> LLVM conversions. This is a first prototype, so more features will be added in due time. - Overview The runner follows similar flow as the other runners in-tree. However, having converted the kernel to SPIR-V, we encode the bind attributes of global variables that represent kernel arguments. Then SPIR-V module is converted to LLVM. On the host side, we emulate passing the data to device by creating in main module globals with the same symbolic name as in kernel module. These global variables are later linked with ones from the nested module. We copy data from kernel arguments to globals, call the kernel function from nested module and then copy the data back. - Current state At the moment, the runner is capable of running 2 modules, nested one in another. The kernel module must contain exactly one kernel function. Also, the runner supports rank 1 integer memref types as arguments (to be scaled). - Enhancement of JitRunner and ExecutionEngine To translate nested modules to LLVM IR, JitRunner and ExecutionEngine were altered to take an optional (default to `nullptr`) function reference that is a custom LLVM IR module builder. This allows to customize LLVM IR module creation from MLIR modules. Reviewed By: ftynse, mravishankar Differential Revision: https://reviews.llvm.org/D86108
2020-10-23 22:46:18 +08:00
call @fillI32Buffer(%input_casted, %four) : (memref<?xi32>, i32) -> ()
call @fillI32Buffer(%output_casted, %zero) : (memref<?xi32>, i32) -> ()
%one = constant 1 : index
[mlir][gpu] NFC: Change gpu.launch_func ops to custom format. This should fix the reason for the failures after ec7780ebdab480139596c3cb08ee77d7035457b3. I will roll forward in a separate change. Reviewed By: antiagainst Differential Revision: https://reviews.llvm.org/D90410
2020-10-30 02:16:19 +08:00
gpu.launch_func @kernels::@double
blocks in (%one, %one, %one) threads in (%one, %one, %one)
args(%input : memref<6xi32>, %output : memref<6xi32>)
[mlir] fix SPIR-V CPU and Vulkan runners after e2310704d890ad252aeb1ca28b4b84d29514b1d1 The commit in question changed the syntax but did not update the runner tests. This also required registering the MemRef dialect for custom parser to work correctly.
2021-03-16 01:35:40 +08:00
%result = memref.cast %output : memref<6xi32> to memref<*xi32>
[MLIR][mlir-spirv-cpu-runner] A SPIR-V cpu runner prototype This patch introduces a SPIR-V runner. The aim is to run a gpu kernel on a CPU via GPU -> SPIRV -> LLVM conversions. This is a first prototype, so more features will be added in due time. - Overview The runner follows similar flow as the other runners in-tree. However, having converted the kernel to SPIR-V, we encode the bind attributes of global variables that represent kernel arguments. Then SPIR-V module is converted to LLVM. On the host side, we emulate passing the data to device by creating in main module globals with the same symbolic name as in kernel module. These global variables are later linked with ones from the nested module. We copy data from kernel arguments to globals, call the kernel function from nested module and then copy the data back. - Current state At the moment, the runner is capable of running 2 modules, nested one in another. The kernel module must contain exactly one kernel function. Also, the runner supports rank 1 integer memref types as arguments (to be scaled). - Enhancement of JitRunner and ExecutionEngine To translate nested modules to LLVM IR, JitRunner and ExecutionEngine were altered to take an optional (default to `nullptr`) function reference that is a custom LLVM IR module builder. This allows to customize LLVM IR module creation from MLIR modules. Reviewed By: ftynse, mravishankar Differential Revision: https://reviews.llvm.org/D86108
2020-10-23 22:46:18 +08:00
call @print_memref_i32(%result) : (memref<*xi32>) -> ()
return
}
[MLIR] Extend Symbol verification to reject public symbol declarations. - Extend the Symbol interface with `isDeclaration` to identify operations that declare a symbol as opposed to define it. - Extend verification to disallow public declarations as per the discussion in https://llvm.discourse.group/t/rfc-symbol-definition-declaration-x-visibility-checks/2140 - Adopt the new interface for `FuncOp` and fix test and code to not have/create public function declarations. Differential Revision: https://reviews.llvm.org/D91456
2020-11-14 05:04:53 +08:00
func private @fillI32Buffer(%arg0 : memref<?xi32>, %arg1 : i32)
func private @print_memref_i32(%ptr : memref<*xi32>)
[MLIR][mlir-spirv-cpu-runner] A SPIR-V cpu runner prototype This patch introduces a SPIR-V runner. The aim is to run a gpu kernel on a CPU via GPU -> SPIRV -> LLVM conversions. This is a first prototype, so more features will be added in due time. - Overview The runner follows similar flow as the other runners in-tree. However, having converted the kernel to SPIR-V, we encode the bind attributes of global variables that represent kernel arguments. Then SPIR-V module is converted to LLVM. On the host side, we emulate passing the data to device by creating in main module globals with the same symbolic name as in kernel module. These global variables are later linked with ones from the nested module. We copy data from kernel arguments to globals, call the kernel function from nested module and then copy the data back. - Current state At the moment, the runner is capable of running 2 modules, nested one in another. The kernel module must contain exactly one kernel function. Also, the runner supports rank 1 integer memref types as arguments (to be scaled). - Enhancement of JitRunner and ExecutionEngine To translate nested modules to LLVM IR, JitRunner and ExecutionEngine were altered to take an optional (default to `nullptr`) function reference that is a custom LLVM IR module builder. This allows to customize LLVM IR module creation from MLIR modules. Reviewed By: ftynse, mravishankar Differential Revision: https://reviews.llvm.org/D86108
2020-10-23 22:46:18 +08:00
}