274 lines · plain
1// RUN: mlir-opt -test-spirv-module-combiner -split-input-file -verify-diagnostics %s | FileCheck %s2 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 @foo8 9// CHECK-NEXT: spirv.func @use_foo10// CHECK-NEXT: spirv.mlir.addressof @foo11// CHECK-NEXT: spirv.Load12// CHECK-NEXT: spirv.ReturnValue13// CHECK-NEXT: }14 15// CHECK-NEXT: spirv.func @use_bar16// CHECK-NEXT: spirv.mlir.addressof @foo17// CHECK-NEXT: spirv.Load18// CHECK-NEXT: spirv.FAdd19// CHECK-NEXT: spirv.ReturnValue20// 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 : f3230 spirv.ReturnValue %1 : f3231 }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 : f3240 %2 = spirv.FAdd %1, %1 : f3241 spirv.ReturnValue %2 : f3242 }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_bar56// CHECK-NEXT: spirv.mlir.addressof @bar57// CHECK-NEXT: spirv.Load58// CHECK-NEXT: spirv.ReturnValue59// 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 : f3273 spirv.ReturnValue %1 : f3274 }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_bar85// CHECK-NEXT: spirv.mlir.addressof @foo86// CHECK-NEXT: spirv.Load87// CHECK-NEXT: spirv.ReturnValue88// 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 @foo116// CHECK-NEXT: spirv.ReturnValue117// CHECK-NEXT: }118 119// CHECK-NEXT: spirv.func @use_bar()120// CHECK-NEXT: %0 = spirv.mlir.referenceof @foo121// CHECK-NEXT: spirv.FAdd122// CHECK-NEXT: spirv.ReturnValue123// CHECK-NEXT: }124// CHECK-NEXT: }125// CHECK-NEXT: }126 127spirv.module Logical GLSL450 {128 spirv.SpecConstant @foo spec_id(5) = 1. : f32129 130 spirv.func @use_foo() -> (f32) "None" {131 %0 = spirv.mlir.referenceof @foo : f32132 spirv.ReturnValue %0 : f32133 }134}135 136spirv.module Logical GLSL450 {137 spirv.SpecConstant @bar spec_id(5) = 1. : f32138 139 spirv.func @use_bar() -> (f32) "None" {140 %0 = spirv.mlir.referenceof @bar : f32141 %1 = spirv.FAdd %0, %0 : f32142 spirv.ReturnValue %1 : f32143 }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]] : f32153// CHECK-NEXT: %[[MUL:.+]] = spirv.FMul %[[ADD]], %[[ARG2]] : f32154// CHECK-NEXT: spirv.ReturnValue %[[MUL]] : f32155// CHECK-NEXT: }156// CHECK-NEXT: spirv.func @foo_1(%[[ARG0:.+]]: f32, %[[ARG1:.+]]: f32, %[[ARG2:.+]]: f32)157// CHECK-NEXT: %[[ADD:.+]] = spirv.FAdd %[[ARG0]], %[[ARG2]] : f32158// CHECK-NEXT: %[[MUL:.+]] = spirv.FMul %[[ADD]], %[[ARG1]] : f32159// CHECK-NEXT: spirv.ReturnValue %[[MUL]] : f32160// 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: f32166 %mul = spirv.FMul %add, %c: f32167 spirv.ReturnValue %mul: f32168 }169}170 171spirv.module Logical GLSL450 {172 spirv.func @foo(%a: f32, %b: f32, %c: f32) -> f32 "None" {173 %add = spirv.FAdd %a, %c: f32174 %mul = spirv.FMul %add, %b: f32175 spirv.ReturnValue %mul: f32176 }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.ReturnValue189// XXXXX-NEXT: }190 191// XXXXX-NEXT: spirv.func @foo_different_body(%arg0: f32)192// XXXXX-NEXT: spirv.mlir.referenceof193// XXXXX-NEXT: spirv.ReturnValue194// XXXXX-NEXT: }195 196// XXXXX-NEXT: spirv.func @baz(%arg0: i32)197// XXXXX-NEXT: spirv.ReturnValue198// XXXXX-NEXT: }199 200// XXXXX-NEXT: spirv.func @baz_no_return(%arg0: i32)201// XXXXX-NEXT: spirv.Return202// XXXXX-NEXT: }203 204// XXXXX-NEXT: spirv.func @baz_no_return_different_control205// XXXXX-NEXT: spirv.Return206// XXXXX-NEXT: }207 208// XXXXX-NEXT: spirv.func @baz_no_return_another_control209// XXXXX-NEXT: spirv.Return210// XXXXX-NEXT: }211 212// XXXXX-NEXT: spirv.func @kernel213// XXXXX-NEXT: spirv.Return214// XXXXX-NEXT: }215 216// XXXXX-NEXT: spirv.func @kernel_different_attr217// XXXXX-NEXT: spirv.Return218// XXXXX-NEXT: }219// XXXXX-NEXT: }220// XXXXX-NEXT: }221 222module {223spirv.module Logical GLSL450 {224 spirv.SpecConstant @bar spec_id(5) = 1. : f32225 226 spirv.func @foo(%arg0: f32) -> (f32) "None" {227 spirv.ReturnValue %arg0 : f32228 }229 230 spirv.func @foo_duplicate(%arg0: f32) -> (f32) "None" {231 spirv.ReturnValue %arg0 : f32232 }233 234 spirv.func @foo_different_body(%arg0: f32) -> (f32) "None" {235 %0 = spirv.mlir.referenceof @bar : f32236 spirv.ReturnValue %arg0 : f32237 }238 239 spirv.func @baz(%arg0: i32) -> (i32) "None" {240 spirv.ReturnValue %arg0 : i32241 }242 243 spirv.func @baz_no_return(%arg0: i32) "None" {244 spirv.Return245 }246 247 spirv.func @baz_no_return_duplicate(%arg0: i32) -> () "None" {248 spirv.Return249 }250 251 spirv.func @baz_no_return_different_control(%arg0: i32) -> () "Inline" {252 spirv.Return253 }254 255 spirv.func @baz_no_return_another_control(%arg0: i32) -> () "Inline|Pure" {256 spirv.Return257 }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.Return264 }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.Return271 }272}273}274