Changeset View
Changeset View
Standalone View
Standalone View
mlir/test/Conversion/SPIRVToLLVM/lower-host-to-llvm-calls.mlir
| // RUN: mlir-opt --lower-host-to-llvm %s | FileCheck %s | // RUN: mlir-opt --lower-host-to-llvm %s | FileCheck %s | ||||
| module attributes {gpu.container_module, spv.target_env = #spv.target_env<#spv.vce<v1.0, [Shader], [SPV_KHR_variable_pointers]>, #spv.resource_limits<max_compute_workgroup_invocations = 128, max_compute_workgroup_size = [128, 128, 64]>>} { | module attributes {gpu.container_module, spv.target_env = #spv.target_env<#spv.vce<v1.0, [Shader], [SPV_KHR_variable_pointers]>, #spv.resource_limits<max_compute_workgroup_invocations = 128, max_compute_workgroup_size = [128, 128, 64]>>} { | ||||
| // CHECK: llvm.mlir.global linkonce @__spv__foo_bar_arg_0_descriptor_set0_binding0() : !llvm.struct<(array<6 x i32>)> | // CHECK: llvm.mlir.global linkonce @__spv__foo_bar_arg_0_descriptor_set0_binding0() {addr_space = 0 : i32} : !llvm.struct<(array<6 x i32>)> | ||||
| // CHECK: llvm.func @__spv__foo_bar() | // CHECK: llvm.func @__spv__foo_bar() | ||||
| // CHECK: spv.module @__spv__foo | // CHECK: spv.module @__spv__foo | ||||
| // CHECK: spv.GlobalVariable @bar_arg_0 bind(0, 0) : !spv.ptr<!spv.struct<(!spv.array<6 x i32, stride=4> [0])>, StorageBuffer> | // CHECK: spv.GlobalVariable @bar_arg_0 bind(0, 0) : !spv.ptr<!spv.struct<(!spv.array<6 x i32, stride=4> [0])>, StorageBuffer> | ||||
| // CHECK: spv.func @__spv__foo_bar | // CHECK: spv.func @__spv__foo_bar | ||||
| // CHECK: spv.EntryPoint "GLCompute" @__spv__foo_bar | // CHECK: spv.EntryPoint "GLCompute" @__spv__foo_bar | ||||
| // CHECK: spv.ExecutionMode @__spv__foo_bar "LocalSize", 1, 1, 1 | // CHECK: spv.ExecutionMode @__spv__foo_bar "LocalSize", 1, 1, 1 | ||||
| Show All 34 Lines | |||||