mlir/test/Dialect/SPIRV/Linking/ModuleCombiner/deduplication.mlir

   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
  24 spirv.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
  34 spirv.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
  63 spirv.module Logical GLSL450 {
  64   spirv.GlobalVariable @foo bind(1, 0) : !spirv.ptr<i32, Input>
  65 }
  66
  67 spirv.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
  92 spirv.module Logical GLSL450 {
  93   spirv.GlobalVariable @foo built_in("GlobalInvocationId") : !spirv.ptr<vector<3xi32>, Input>
  94 }
  95
  96 spirv.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
 127 spirv.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
 136 spirv.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
 163 spirv.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
 171 spirv.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
 222 module {
 223 spirv.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 }