122 lines · plain
1// RUN: mlir-opt -spirv-lower-abi-attrs -verify-diagnostics %s -o - | FileCheck %s2 3module attributes {4 spirv.target_env = #spirv.target_env<5 #spirv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>, #spirv.resource_limits<>>6} {7 8// CHECK-LABEL: spirv.module9spirv.module Logical GLSL450 {10 // CHECK-DAG: spirv.GlobalVariable [[WORKGROUPSIZE:@.*]] built_in("WorkgroupSize")11 spirv.GlobalVariable @__builtin_var_WorkgroupSize__ built_in("WorkgroupSize") : !spirv.ptr<vector<3xi32>, Input>12 // CHECK-DAG: spirv.GlobalVariable [[NUMWORKGROUPS:@.*]] built_in("NumWorkgroups")13 spirv.GlobalVariable @__builtin_var_NumWorkgroups__ built_in("NumWorkgroups") : !spirv.ptr<vector<3xi32>, Input>14 // CHECK-DAG: spirv.GlobalVariable [[LOCALINVOCATIONID:@.*]] built_in("LocalInvocationId")15 spirv.GlobalVariable @__builtin_var_LocalInvocationId__ built_in("LocalInvocationId") : !spirv.ptr<vector<3xi32>, Input>16 // CHECK-DAG: spirv.GlobalVariable [[WORKGROUPID:@.*]] built_in("WorkgroupId")17 spirv.GlobalVariable @__builtin_var_WorkgroupId__ built_in("WorkgroupId") : !spirv.ptr<vector<3xi32>, Input>18 // CHECK-DAG: spirv.GlobalVariable [[VAR0:@.*]] bind(0, 0) : !spirv.ptr<!spirv.struct<(!spirv.array<12 x !spirv.array<4 x f32, stride=4>, stride=16> [0])>, StorageBuffer>19 // CHECK-DAG: spirv.GlobalVariable [[VAR1:@.*]] bind(0, 1) : !spirv.ptr<!spirv.struct<(!spirv.array<12 x !spirv.array<4 x f32, stride=4>, stride=16> [0])>, StorageBuffer>20 // CHECK-DAG: spirv.GlobalVariable [[VAR2:@.*]] bind(0, 2) : !spirv.ptr<!spirv.struct<(!spirv.array<12 x !spirv.array<4 x f32, stride=4>, stride=16> [0])>, StorageBuffer>21 // CHECK-DAG: spirv.GlobalVariable [[VAR3:@.*]] bind(0, 3) : !spirv.ptr<!spirv.struct<(i32 [0])>, StorageBuffer>22 // CHECK-DAG: spirv.GlobalVariable [[VAR4:@.*]] bind(0, 4) : !spirv.ptr<!spirv.struct<(i32 [0])>, StorageBuffer>23 // CHECK-DAG: spirv.GlobalVariable [[VAR5:@.*]] bind(0, 5) : !spirv.ptr<!spirv.struct<(i32 [0])>, StorageBuffer>24 // CHECK-DAG: spirv.GlobalVariable [[VAR6:@.*]] bind(0, 6) : !spirv.ptr<!spirv.struct<(i32 [0])>, StorageBuffer>25 // CHECK: spirv.func [[FN:@.*]]()26 spirv.func @load_store_kernel(27 %arg0: !spirv.ptr<!spirv.struct<(!spirv.array<12 x !spirv.array<4 x f32>>)>, StorageBuffer>28 {spirv.interface_var_abi = #spirv.interface_var_abi<(0, 0)>},29 %arg1: !spirv.ptr<!spirv.struct<(!spirv.array<12 x !spirv.array<4 x f32>>)>, StorageBuffer>30 {spirv.interface_var_abi = #spirv.interface_var_abi<(0, 1)>},31 %arg2: !spirv.ptr<!spirv.struct<(!spirv.array<12 x !spirv.array<4 x f32>>)>, StorageBuffer>32 {spirv.interface_var_abi = #spirv.interface_var_abi<(0, 2)>},33 %arg3: i3234 {spirv.interface_var_abi = #spirv.interface_var_abi<(0, 3), StorageBuffer>},35 %arg4: i3236 {spirv.interface_var_abi = #spirv.interface_var_abi<(0, 4), StorageBuffer>},37 %arg5: i3238 {spirv.interface_var_abi = #spirv.interface_var_abi<(0, 5), StorageBuffer>},39 %arg6: i3240 {spirv.interface_var_abi = #spirv.interface_var_abi<(0, 6), StorageBuffer>}) "None"41 attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [32, 1, 1]>} {42 // CHECK: [[ADDRESSARG0:%.*]] = spirv.mlir.addressof [[VAR0]]43 // CHECK: [[ARG0:%.*]] = spirv.Bitcast [[ADDRESSARG0]]44 // CHECK: [[ADDRESSARG1:%.*]] = spirv.mlir.addressof [[VAR1]]45 // CHECK: [[ARG1:%.*]] = spirv.Bitcast [[ADDRESSARG1]]46 // CHECK: [[ADDRESSARG2:%.*]] = spirv.mlir.addressof [[VAR2]]47 // CHECK: [[ARG2:%.*]] = spirv.Bitcast [[ADDRESSARG2]]48 // CHECK: [[ADDRESSARG3:%.*]] = spirv.mlir.addressof [[VAR3]]49 // CHECK: [[CONST3:%.*]] = spirv.Constant 0 : i3250 // CHECK: [[ARG3PTR:%.*]] = spirv.AccessChain [[ADDRESSARG3]]{{\[}}[[CONST3]]51 // CHECK: [[ARG3:%.*]] = spirv.Load "StorageBuffer" [[ARG3PTR]]52 // CHECK: [[ADDRESSARG4:%.*]] = spirv.mlir.addressof [[VAR4]]53 // CHECK: [[CONST4:%.*]] = spirv.Constant 0 : i3254 // CHECK: [[ARG4PTR:%.*]] = spirv.AccessChain [[ADDRESSARG4]]{{\[}}[[CONST4]]55 // CHECK: [[ARG4:%.*]] = spirv.Load "StorageBuffer" [[ARG4PTR]]56 // CHECK: [[ADDRESSARG5:%.*]] = spirv.mlir.addressof [[VAR5]]57 // CHECK: [[CONST5:%.*]] = spirv.Constant 0 : i3258 // CHECK: [[ARG5PTR:%.*]] = spirv.AccessChain [[ADDRESSARG5]]{{\[}}[[CONST5]]59 // CHECK: {{%.*}} = spirv.Load "StorageBuffer" [[ARG5PTR]]60 // CHECK: [[ADDRESSARG6:%.*]] = spirv.mlir.addressof [[VAR6]]61 // CHECK: [[CONST6:%.*]] = spirv.Constant 0 : i3262 // CHECK: [[ARG6PTR:%.*]] = spirv.AccessChain [[ADDRESSARG6]]{{\[}}[[CONST6]]63 // CHECK: {{%.*}} = spirv.Load "StorageBuffer" [[ARG6PTR]] 64 %0 = spirv.mlir.addressof @__builtin_var_WorkgroupId__ : !spirv.ptr<vector<3xi32>, Input>65 %1 = spirv.Load "Input" %0 : vector<3xi32>66 %2 = spirv.CompositeExtract %1[0 : i32] : vector<3xi32>67 %3 = spirv.mlir.addressof @__builtin_var_WorkgroupId__ : !spirv.ptr<vector<3xi32>, Input>68 %4 = spirv.Load "Input" %3 : vector<3xi32>69 %5 = spirv.CompositeExtract %4[1 : i32] : vector<3xi32>70 %6 = spirv.mlir.addressof @__builtin_var_WorkgroupId__ : !spirv.ptr<vector<3xi32>, Input>71 %7 = spirv.Load "Input" %6 : vector<3xi32>72 %8 = spirv.CompositeExtract %7[2 : i32] : vector<3xi32>73 %9 = spirv.mlir.addressof @__builtin_var_LocalInvocationId__ : !spirv.ptr<vector<3xi32>, Input>74 %10 = spirv.Load "Input" %9 : vector<3xi32>75 %11 = spirv.CompositeExtract %10[0 : i32] : vector<3xi32>76 %12 = spirv.mlir.addressof @__builtin_var_LocalInvocationId__ : !spirv.ptr<vector<3xi32>, Input>77 %13 = spirv.Load "Input" %12 : vector<3xi32>78 %14 = spirv.CompositeExtract %13[1 : i32] : vector<3xi32>79 %15 = spirv.mlir.addressof @__builtin_var_LocalInvocationId__ : !spirv.ptr<vector<3xi32>, Input>80 %16 = spirv.Load "Input" %15 : vector<3xi32>81 %17 = spirv.CompositeExtract %16[2 : i32] : vector<3xi32>82 %18 = spirv.mlir.addressof @__builtin_var_NumWorkgroups__ : !spirv.ptr<vector<3xi32>, Input>83 %19 = spirv.Load "Input" %18 : vector<3xi32>84 %20 = spirv.CompositeExtract %19[0 : i32] : vector<3xi32>85 %21 = spirv.mlir.addressof @__builtin_var_NumWorkgroups__ : !spirv.ptr<vector<3xi32>, Input>86 %22 = spirv.Load "Input" %21 : vector<3xi32>87 %23 = spirv.CompositeExtract %22[1 : i32] : vector<3xi32>88 %24 = spirv.mlir.addressof @__builtin_var_NumWorkgroups__ : !spirv.ptr<vector<3xi32>, Input>89 %25 = spirv.Load "Input" %24 : vector<3xi32>90 %26 = spirv.CompositeExtract %25[2 : i32] : vector<3xi32>91 %27 = spirv.mlir.addressof @__builtin_var_WorkgroupSize__ : !spirv.ptr<vector<3xi32>, Input>92 %28 = spirv.Load "Input" %27 : vector<3xi32>93 %29 = spirv.CompositeExtract %28[0 : i32] : vector<3xi32>94 %30 = spirv.mlir.addressof @__builtin_var_WorkgroupSize__ : !spirv.ptr<vector<3xi32>, Input>95 %31 = spirv.Load "Input" %30 : vector<3xi32>96 %32 = spirv.CompositeExtract %31[1 : i32] : vector<3xi32>97 %33 = spirv.mlir.addressof @__builtin_var_WorkgroupSize__ : !spirv.ptr<vector<3xi32>, Input>98 %34 = spirv.Load "Input" %33 : vector<3xi32>99 %35 = spirv.CompositeExtract %34[2 : i32] : vector<3xi32>100 // CHECK: spirv.IAdd [[ARG3]]101 %36 = spirv.IAdd %arg3, %2 : i32102 // CHECK: spirv.IAdd [[ARG4]]103 %37 = spirv.IAdd %arg4, %11 : i32104 // CHECK: spirv.AccessChain [[ARG0]]105 %c0 = spirv.Constant 0 : i32106 %38 = spirv.AccessChain %arg0[%c0, %36, %37] : !spirv.ptr<!spirv.struct<(!spirv.array<12 x !spirv.array<4 x f32>>)>, StorageBuffer>, i32, i32, i32 -> !spirv.ptr<f32, StorageBuffer>107 %39 = spirv.Load "StorageBuffer" %38 : f32108 // CHECK: spirv.AccessChain [[ARG1]]109 %40 = spirv.AccessChain %arg1[%c0, %36, %37] : !spirv.ptr<!spirv.struct<(!spirv.array<12 x !spirv.array<4 x f32>>)>, StorageBuffer>, i32, i32, i32 -> !spirv.ptr<f32, StorageBuffer>110 %41 = spirv.Load "StorageBuffer" %40 : f32111 %42 = spirv.FAdd %39, %41 : f32112 // CHECK: spirv.AccessChain [[ARG2]]113 %43 = spirv.AccessChain %arg2[%c0, %36, %37] : !spirv.ptr<!spirv.struct<(!spirv.array<12 x !spirv.array<4 x f32>>)>, StorageBuffer>, i32, i32, i32 -> !spirv.ptr<f32, StorageBuffer>114 spirv.Store "StorageBuffer" %43, %42 : f32115 spirv.Return116 }117 // CHECK: spirv.EntryPoint "GLCompute" [[FN]], [[WORKGROUPID]], [[LOCALINVOCATIONID]], [[NUMWORKGROUPS]], [[WORKGROUPSIZE]]118 // CHECK-NEXT: spirv.ExecutionMode [[FN]] "LocalSize", 32, 1, 1119} // end spirv.module120 121} // end module122