1// RUN: mlir-opt --lower-host-to-llvm %s -split-input-file | FileCheck %s 2 3module attributes {gpu.container_module, spirv.target_env = #spirv.target_env<#spirv.vce<v1.0, [Shader], [SPV_KHR_variable_pointers]>, #spirv.resource_limits<max_compute_workgroup_invocations = 128, max_compute_workgroup_size = [128, 128, 64]>>} { 4 5 // CHECK: llvm.mlir.global linkonce @__spv__foo_bar_arg_0_descriptor_set0_binding0() {addr_space = 0 : i32} : !llvm.struct<(array<6 x i32>)> 6 // CHECK: llvm.func @__spv__foo_bar() 7 8 // CHECK: spirv.module @__spv__foo 9 // CHECK: spirv.GlobalVariable @bar_arg_0 bind(0, 0) : !spirv.ptr<!spirv.struct<(!spirv.array<6 x i32, stride=4> [0])>, StorageBuffer> 10 // CHECK: spirv.func @__spv__foo_bar 11 12 // CHECK: spirv.EntryPoint "GLCompute" @__spv__foo_bar 13 // CHECK: spirv.ExecutionMode @__spv__foo_bar "LocalSize", 1, 1, 1 14 15 // CHECK-LABEL: @main 16 // CHECK: %[[SRC:.*]] = llvm.extractvalue %{{.*}}[0] : !llvm.struct<(ptr, ptr, i64, array<1 x i64>, array<1 x i64>)> 17 // CHECK-NEXT: %[[DEST:.*]] = llvm.mlir.addressof @__spv__foo_bar_arg_0_descriptor_set0_binding0 : !llvm.ptr 18 // CHECK-NEXT: "llvm.intr.memcpy"(%[[DEST]], %[[SRC]], %[[SIZE:.*]]) <{isVolatile = false}> : (!llvm.ptr, !llvm.ptr, i64) -> () 19 // CHECK-NEXT: llvm.call @__spv__foo_bar() : () -> () 20 // CHECK-NEXT: "llvm.intr.memcpy"(%[[SRC]], %[[DEST]], %[[SIZE]]) <{isVolatile = false}> : (!llvm.ptr, !llvm.ptr, i64) -> () 21 22 spirv.module @__spv__foo Logical GLSL450 requires #spirv.vce<v1.0, [Shader], [SPV_KHR_variable_pointers]> { 23 spirv.GlobalVariable @bar_arg_0 bind(0, 0) : !spirv.ptr<!spirv.struct<(!spirv.array<6 x i32, stride=4> [0])>, StorageBuffer> 24 spirv.func @bar() "None" attributes {workgroup_attributions = 0 : i64} { 25 %0 = spirv.mlir.addressof @bar_arg_0 : !spirv.ptr<!spirv.struct<(!spirv.array<6 x i32, stride=4> [0])>, StorageBuffer> 26 spirv.Return 27 } 28 spirv.EntryPoint "GLCompute" @bar 29 spirv.ExecutionMode @bar "LocalSize", 1, 1, 1 30 } 31 32 gpu.module @foo { 33 gpu.func @bar(%arg0: memref<6xi32>) kernel attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [1, 1, 1]>} { 34 gpu.return 35 } 36 } 37 38 func.func @main() { 39 %buffer = memref.alloc() : memref<6xi32> 40 %one = arith.constant 1 : index 41 gpu.launch_func @foo::@bar blocks in (%one, %one, %one) 42 threads in (%one, %one, %one) args(%buffer : memref<6xi32>) 43 return 44 } 45} 46 47// ----- 48 49// Check using a specified sym_name attribute. 50module { 51 spirv.module Logical GLSL450 attributes { sym_name = "spirv.sym" } { 52 // CHECK: spirv.func @spirv.sym_bar 53 // CHECK: spirv.EntryPoint "GLCompute" @spirv.sym_bar 54 // CHECK: spirv.ExecutionMode @spirv.sym_bar "LocalSize", 1, 1, 1 55 spirv.func @bar() "None" { 56 spirv.Return 57 } 58 spirv.EntryPoint "GLCompute" @bar 59 spirv.ExecutionMode @bar "LocalSize", 1, 1, 1 60 } 61} 62 63// ----- 64 65// Check using the default sym_name attribute. 66module { 67 spirv.module Logical GLSL450 { 68 // CHECK: spirv.func @__spv___bar 69 // CHECK: spirv.EntryPoint "GLCompute" @__spv___bar 70 // CHECK: spirv.ExecutionMode @__spv___bar "LocalSize", 1, 1, 1 71 spirv.func @bar() "None" { 72 spirv.Return 73 } 74 spirv.EntryPoint "GLCompute" @bar 75 spirv.ExecutionMode @bar "LocalSize", 1, 1, 1 76 } 77} 78