1// RUN: mlir-opt -test-spirv-module-combiner -split-input-file -verify-diagnostics %s | FileCheck %s 2 3// Deduplicate 2 global variables with the same descriptor set and binding. 4 5// CHECK: module { 6// CHECK-NEXT: spirv.module Logical GLSL450 { 7// CHECK-NEXT: spirv.GlobalVariable @foo 8 9// CHECK-NEXT: spirv.func @use_foo 10// CHECK-NEXT: spirv.mlir.addressof @foo 11// CHECK-NEXT: spirv.Load 12// CHECK-NEXT: spirv.ReturnValue 13// CHECK-NEXT: } 14 15// CHECK-NEXT: spirv.func @use_bar 16// CHECK-NEXT: spirv.mlir.addressof @foo 17// CHECK-NEXT: spirv.Load 18// CHECK-NEXT: spirv.FAdd 19// CHECK-NEXT: spirv.ReturnValue 20// CHECK-NEXT: } 21// CHECK-NEXT: } 22// CHECK-NEXT: } 23 24spirv.module Logical GLSL450 { 25 spirv.GlobalVariable @foo bind(1, 0) : !spirv.ptr<f32, Input> 26 27 spirv.func @use_foo() -> f32 "None" { 28 %0 = spirv.mlir.addressof @foo : !spirv.ptr<f32, Input> 29 %1 = spirv.Load "Input" %0 : f32 30 spirv.ReturnValue %1 : f32 31 } 32} 33 34spirv.module Logical GLSL450 { 35 spirv.GlobalVariable @bar bind(1, 0) : !spirv.ptr<f32, Input> 36 37 spirv.func @use_bar() -> f32 "None" { 38 %0 = spirv.mlir.addressof @bar : !spirv.ptr<f32, Input> 39 %1 = spirv.Load "Input" %0 : f32 40 %2 = spirv.FAdd %1, %1 : f32 41 spirv.ReturnValue %2 : f32 42 } 43} 44 45// ----- 46 47// Deduplicate 2 global variables with the same descriptor set and binding but different types. 48 49// CHECK: module { 50// CHECK-NEXT: spirv.module Logical GLSL450 { 51// CHECK-NEXT: spirv.GlobalVariable @foo bind(1, 0) 52 53// CHECK-NEXT: spirv.GlobalVariable @bar bind(1, 0) 54 55// CHECK-NEXT: spirv.func @use_bar 56// CHECK-NEXT: spirv.mlir.addressof @bar 57// CHECK-NEXT: spirv.Load 58// CHECK-NEXT: spirv.ReturnValue 59// CHECK-NEXT: } 60// CHECK-NEXT: } 61// CHECK-NEXT: } 62 63spirv.module Logical GLSL450 { 64 spirv.GlobalVariable @foo bind(1, 0) : !spirv.ptr<i32, Input> 65} 66 67spirv.module Logical GLSL450 { 68 spirv.GlobalVariable @bar bind(1, 0) : !spirv.ptr<f32, Input> 69 70 spirv.func @use_bar() -> f32 "None" { 71 %0 = spirv.mlir.addressof @bar : !spirv.ptr<f32, Input> 72 %1 = spirv.Load "Input" %0 : f32 73 spirv.ReturnValue %1 : f32 74 } 75} 76 77// ----- 78 79// Deduplicate 2 global variables with the same built-in attribute. 80 81// CHECK: module { 82// CHECK-NEXT: spirv.module Logical GLSL450 { 83// CHECK-NEXT: spirv.GlobalVariable @foo built_in("GlobalInvocationId") 84// CHECK-NEXT: spirv.func @use_bar 85// CHECK-NEXT: spirv.mlir.addressof @foo 86// CHECK-NEXT: spirv.Load 87// CHECK-NEXT: spirv.ReturnValue 88// CHECK-NEXT: } 89// CHECK-NEXT: } 90// CHECK-NEXT: } 91 92spirv.module Logical GLSL450 { 93 spirv.GlobalVariable @foo built_in("GlobalInvocationId") : !spirv.ptr<vector<3xi32>, Input> 94} 95 96spirv.module Logical GLSL450 { 97 spirv.GlobalVariable @bar built_in("GlobalInvocationId") : !spirv.ptr<vector<3xi32>, Input> 98 99 spirv.func @use_bar() -> vector<3xi32> "None" { 100 %0 = spirv.mlir.addressof @bar : !spirv.ptr<vector<3xi32>, Input> 101 %1 = spirv.Load "Input" %0 : vector<3xi32> 102 spirv.ReturnValue %1 : vector<3xi32> 103 } 104} 105 106// ----- 107 108// Deduplicate 2 spec constants with the same spec ID. 109 110// CHECK: module { 111// CHECK-NEXT: spirv.module Logical GLSL450 { 112// CHECK-NEXT: spirv.SpecConstant @foo spec_id(5) 113 114// CHECK-NEXT: spirv.func @use_foo() 115// CHECK-NEXT: %0 = spirv.mlir.referenceof @foo 116// CHECK-NEXT: spirv.ReturnValue 117// CHECK-NEXT: } 118 119// CHECK-NEXT: spirv.func @use_bar() 120// CHECK-NEXT: %0 = spirv.mlir.referenceof @foo 121// CHECK-NEXT: spirv.FAdd 122// CHECK-NEXT: spirv.ReturnValue 123// CHECK-NEXT: } 124// CHECK-NEXT: } 125// CHECK-NEXT: } 126 127spirv.module Logical GLSL450 { 128 spirv.SpecConstant @foo spec_id(5) = 1. : f32 129 130 spirv.func @use_foo() -> (f32) "None" { 131 %0 = spirv.mlir.referenceof @foo : f32 132 spirv.ReturnValue %0 : f32 133 } 134} 135 136spirv.module Logical GLSL450 { 137 spirv.SpecConstant @bar spec_id(5) = 1. : f32 138 139 spirv.func @use_bar() -> (f32) "None" { 140 %0 = spirv.mlir.referenceof @bar : f32 141 %1 = spirv.FAdd %0, %0 : f32 142 spirv.ReturnValue %1 : f32 143 } 144} 145 146// ----- 147 148// Don't deduplicate functions with similar ops but different operands. 149 150// CHECK: spirv.module Logical GLSL450 { 151// CHECK-NEXT: spirv.func @foo(%[[ARG0:.+]]: f32, %[[ARG1:.+]]: f32, %[[ARG2:.+]]: f32) 152// CHECK-NEXT: %[[ADD:.+]] = spirv.FAdd %[[ARG0]], %[[ARG1]] : f32 153// CHECK-NEXT: %[[MUL:.+]] = spirv.FMul %[[ADD]], %[[ARG2]] : f32 154// CHECK-NEXT: spirv.ReturnValue %[[MUL]] : f32 155// CHECK-NEXT: } 156// CHECK-NEXT: spirv.func @foo_1(%[[ARG0:.+]]: f32, %[[ARG1:.+]]: f32, %[[ARG2:.+]]: f32) 157// CHECK-NEXT: %[[ADD:.+]] = spirv.FAdd %[[ARG0]], %[[ARG2]] : f32 158// CHECK-NEXT: %[[MUL:.+]] = spirv.FMul %[[ADD]], %[[ARG1]] : f32 159// CHECK-NEXT: spirv.ReturnValue %[[MUL]] : f32 160// CHECK-NEXT: } 161// CHECK-NEXT: } 162 163spirv.module Logical GLSL450 { 164 spirv.func @foo(%a: f32, %b: f32, %c: f32) -> f32 "None" { 165 %add = spirv.FAdd %a, %b: f32 166 %mul = spirv.FMul %add, %c: f32 167 spirv.ReturnValue %mul: f32 168 } 169} 170 171spirv.module Logical GLSL450 { 172 spirv.func @foo(%a: f32, %b: f32, %c: f32) -> f32 "None" { 173 %add = spirv.FAdd %a, %c: f32 174 %mul = spirv.FMul %add, %b: f32 175 spirv.ReturnValue %mul: f32 176 } 177} 178 179// ----- 180 181// TODO: re-enable this test once we have better function deduplication. 182 183// XXXXX: module { 184// XXXXX-NEXT: spirv.module Logical GLSL450 { 185// XXXXX-NEXT: spirv.SpecConstant @bar spec_id(5) 186 187// XXXXX-NEXT: spirv.func @foo(%arg0: f32) 188// XXXXX-NEXT: spirv.ReturnValue 189// XXXXX-NEXT: } 190 191// XXXXX-NEXT: spirv.func @foo_different_body(%arg0: f32) 192// XXXXX-NEXT: spirv.mlir.referenceof 193// XXXXX-NEXT: spirv.ReturnValue 194// XXXXX-NEXT: } 195 196// XXXXX-NEXT: spirv.func @baz(%arg0: i32) 197// XXXXX-NEXT: spirv.ReturnValue 198// XXXXX-NEXT: } 199 200// XXXXX-NEXT: spirv.func @baz_no_return(%arg0: i32) 201// XXXXX-NEXT: spirv.Return 202// XXXXX-NEXT: } 203 204// XXXXX-NEXT: spirv.func @baz_no_return_different_control 205// XXXXX-NEXT: spirv.Return 206// XXXXX-NEXT: } 207 208// XXXXX-NEXT: spirv.func @baz_no_return_another_control 209// XXXXX-NEXT: spirv.Return 210// XXXXX-NEXT: } 211 212// XXXXX-NEXT: spirv.func @kernel 213// XXXXX-NEXT: spirv.Return 214// XXXXX-NEXT: } 215 216// XXXXX-NEXT: spirv.func @kernel_different_attr 217// XXXXX-NEXT: spirv.Return 218// XXXXX-NEXT: } 219// XXXXX-NEXT: } 220// XXXXX-NEXT: } 221 222module { 223spirv.module Logical GLSL450 { 224 spirv.SpecConstant @bar spec_id(5) = 1. : f32 225 226 spirv.func @foo(%arg0: f32) -> (f32) "None" { 227 spirv.ReturnValue %arg0 : f32 228 } 229 230 spirv.func @foo_duplicate(%arg0: f32) -> (f32) "None" { 231 spirv.ReturnValue %arg0 : f32 232 } 233 234 spirv.func @foo_different_body(%arg0: f32) -> (f32) "None" { 235 %0 = spirv.mlir.referenceof @bar : f32 236 spirv.ReturnValue %arg0 : f32 237 } 238 239 spirv.func @baz(%arg0: i32) -> (i32) "None" { 240 spirv.ReturnValue %arg0 : i32 241 } 242 243 spirv.func @baz_no_return(%arg0: i32) "None" { 244 spirv.Return 245 } 246 247 spirv.func @baz_no_return_duplicate(%arg0: i32) -> () "None" { 248 spirv.Return 249 } 250 251 spirv.func @baz_no_return_different_control(%arg0: i32) -> () "Inline" { 252 spirv.Return 253 } 254 255 spirv.func @baz_no_return_another_control(%arg0: i32) -> () "Inline|Pure" { 256 spirv.Return 257 } 258 259 spirv.func @kernel( 260 %arg0: f32, 261 %arg1: !spirv.ptr<!spirv.struct<(!spirv.array<12 x f32>)>, CrossWorkgroup>) "None" 262 attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [32, 1, 1]>} { 263 spirv.Return 264 } 265 266 spirv.func @kernel_different_attr( 267 %arg0: f32, 268 %arg1: !spirv.ptr<!spirv.struct<(!spirv.array<12 x f32>)>, CrossWorkgroup>) "None" 269 attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [64, 1, 1]>} { 270 spirv.Return 271 } 272} 273} 274