mlir/test/Conversion/GPUToSPIRV/reductions.mlir

   1 // RUN: mlir-opt -split-input-file -convert-gpu-to-spirv -verify-diagnostics %s -o - | FileCheck %s
   2
   3 module attributes {
   4   gpu.container_module,
   5   spirv.target_env = #spirv.target_env<#spirv.vce<v1.3, [Kernel, Addresses, Groups, GroupNonUniformArithmetic, GroupUniformArithmeticKHR], []>, #spirv.resource_limits<>>
   6 } {
   7
   8 gpu.module @kernels {
   9   // CHECK-LABEL:  spirv.func @test
  10   //  CHECK-SAME: (%[[ARG:.*]]: f32)
  11   gpu.func @test(%arg : f32) kernel
  12     attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {
  13     // CHECK: %{{.*}} = spirv.GroupFAdd <Workgroup> <Reduce> %[[ARG]] : f32
  14     %reduced = gpu.all_reduce add %arg uniform {} : (f32) -> (f32)
  15     gpu.return
  16   }
  17 }
  18
  19 }
  20
  21 // -----
  22
  23 module attributes {
  24   gpu.container_module,
  25   spirv.target_env = #spirv.target_env<#spirv.vce<v1.3, [Kernel, Addresses, Groups, GroupNonUniformArithmetic, GroupUniformArithmeticKHR], []>, #spirv.resource_limits<>>
  26 } {
  27
  28 gpu.module @kernels {
  29   // CHECK-LABEL:  spirv.func @test
  30   //  CHECK-SAME: (%[[ARG:.*]]: f32)
  31   gpu.func @test(%arg : f32) kernel
  32     attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {
  33     // CHECK: %{{.*}} = spirv.GroupNonUniformFAdd <Workgroup> <Reduce> %[[ARG]] : f32 -> f32
  34     %reduced = gpu.all_reduce add %arg {} : (f32) -> (f32)
  35     gpu.return
  36   }
  37 }
  38
  39 }
  40
  41 // -----
  42
  43 module attributes {
  44   gpu.container_module,
  45   spirv.target_env = #spirv.target_env<#spirv.vce<v1.3, [Kernel, Addresses, Groups, GroupNonUniformArithmetic, GroupUniformArithmeticKHR], []>, #spirv.resource_limits<>>
  46 } {
  47
  48 gpu.module @kernels {
  49   // CHECK-LABEL:  spirv.func @test
  50   //  CHECK-SAME: (%[[ARG:.*]]: i32)
  51   gpu.func @test(%arg : i32) kernel
  52     attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {
  53     // CHECK: %{{.*}} = spirv.GroupIAdd <Workgroup> <Reduce> %[[ARG]] : i32
  54     %reduced = gpu.all_reduce add %arg uniform {} : (i32) -> (i32)
  55     gpu.return
  56   }
  57 }
  58
  59 }
  60
  61 // -----
  62
  63 module attributes {
  64   gpu.container_module,
  65   spirv.target_env = #spirv.target_env<#spirv.vce<v1.3, [Kernel, Addresses, Groups, GroupNonUniformArithmetic, GroupUniformArithmeticKHR], []>, #spirv.resource_limits<>>
  66 } {
  67
  68 gpu.module @kernels {
  69   // CHECK-LABEL:  spirv.func @test
  70   //  CHECK-SAME: (%[[ARG:.*]]: i32)
  71   gpu.func @test(%arg : i32) kernel
  72     attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {
  73     // CHECK: %{{.*}} = spirv.GroupNonUniformIAdd <Workgroup> <Reduce> %[[ARG]] : i32 -> i32
  74     %reduced = gpu.all_reduce add %arg {} : (i32) -> (i32)
  75     gpu.return
  76   }
  77 }
  78
  79 }
  80
  81 // -----
  82
  83 module attributes {
  84   gpu.container_module,
  85   spirv.target_env = #spirv.target_env<#spirv.vce<v1.3, [Kernel, Addresses, Groups, GroupNonUniformArithmetic, GroupUniformArithmeticKHR], []>, #spirv.resource_limits<>>
  86 } {
  87
  88 gpu.module @kernels {
  89   // CHECK-LABEL:  spirv.func @test
  90   //  CHECK-SAME: (%[[ARG:.*]]: f32)
  91   gpu.func @test(%arg : f32) kernel
  92     attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {
  93     // CHECK: %{{.*}} = spirv.GroupFAdd <Subgroup> <Reduce> %[[ARG]] : f32
  94     %reduced = gpu.subgroup_reduce add %arg uniform : (f32) -> (f32)
  95     gpu.return
  96   }
  97 }
  98
  99 }
 100
 101 // -----
 102
 103 module attributes {
 104   gpu.container_module,
 105   spirv.target_env = #spirv.target_env<#spirv.vce<v1.3, [Kernel, Addresses, Groups, GroupNonUniformArithmetic, GroupUniformArithmeticKHR], []>, #spirv.resource_limits<>>
 106 } {
 107
 108 gpu.module @kernels {
 109   // CHECK-LABEL:  spirv.func @test
 110   //  CHECK-SAME: (%[[ARG:.*]]: f32)
 111   gpu.func @test(%arg : f32) kernel
 112     attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {
 113     // CHECK: %{{.*}} = spirv.GroupNonUniformFAdd <Subgroup> <Reduce> %[[ARG]] : f32 -> f32
 114     %reduced = gpu.subgroup_reduce add %arg : (f32) -> (f32)
 115     gpu.return
 116   }
 117 }
 118
 119 }
 120
 121 // -----
 122
 123 module attributes {
 124   gpu.container_module,
 125   spirv.target_env = #spirv.target_env<#spirv.vce<v1.3, [Kernel, Addresses, Groups, GroupNonUniformArithmetic, GroupUniformArithmeticKHR], []>, #spirv.resource_limits<>>
 126 } {
 127
 128 gpu.module @kernels {
 129   // CHECK-LABEL:  spirv.func @test
 130   //  CHECK-SAME: (%[[ARG:.*]]: i32)
 131   gpu.func @test(%arg : i32) kernel
 132     attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {
 133     // CHECK: %{{.*}} = spirv.GroupIAdd <Subgroup> <Reduce> %[[ARG]] : i32
 134     %reduced = gpu.subgroup_reduce add %arg uniform : (i32) -> (i32)
 135     gpu.return
 136   }
 137 }
 138
 139 }
 140
 141 // -----
 142
 143 module attributes {
 144   gpu.container_module,
 145   spirv.target_env = #spirv.target_env<#spirv.vce<v1.3, [Kernel, Addresses, Groups, GroupNonUniformArithmetic, GroupUniformArithmeticKHR], []>, #spirv.resource_limits<>>
 146 } {
 147
 148 gpu.module @kernels {
 149   // CHECK-LABEL:  spirv.func @test
 150   //  CHECK-SAME: (%[[ARG:.*]]: i32)
 151   gpu.func @test(%arg : i32) kernel
 152     attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {
 153     // CHECK: %{{.*}} = spirv.GroupNonUniformIAdd <Subgroup> <Reduce> %[[ARG]] : i32 -> i32
 154     %reduced = gpu.subgroup_reduce add %arg : (i32) -> (i32)
 155     gpu.return
 156   }
 157 }
 158
 159 }
 160
 161 // -----
 162
 163 module attributes {
 164   gpu.container_module,
 165   spirv.target_env = #spirv.target_env<#spirv.vce<v1.3, [Kernel, Addresses, Groups, GroupNonUniformArithmetic, GroupUniformArithmeticKHR], []>, #spirv.resource_limits<>>
 166 } {
 167
 168 gpu.module @kernels {
 169   // CHECK-LABEL:  spirv.func @test
 170   //  CHECK-SAME: (%[[ARG:.*]]: f32)
 171   gpu.func @test(%arg : f32) kernel
 172     attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {
 173     // CHECK: %{{.*}} = spirv.KHR.GroupFMul <Workgroup> <Reduce> %[[ARG]] : f32
 174     %reduced = gpu.all_reduce mul %arg uniform {} : (f32) -> (f32)
 175     gpu.return
 176   }
 177 }
 178
 179 }
 180
 181 // -----
 182
 183 module attributes {
 184   gpu.container_module,
 185   spirv.target_env = #spirv.target_env<#spirv.vce<v1.3, [Kernel, Addresses, Groups, GroupNonUniformArithmetic, GroupUniformArithmeticKHR], []>, #spirv.resource_limits<>>
 186 } {
 187
 188 gpu.module @kernels {
 189   // CHECK-LABEL:  spirv.func @test
 190   //  CHECK-SAME: (%[[ARG:.*]]: f32)
 191   gpu.func @test(%arg : f32) kernel
 192     attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {
 193     // CHECK: %{{.*}} = spirv.GroupNonUniformFMul <Workgroup> <Reduce> %[[ARG]] : f32 -> f32
 194     %reduced = gpu.all_reduce mul %arg {} : (f32) -> (f32)
 195     gpu.return
 196   }
 197 }
 198
 199 }
 200
 201 // -----
 202
 203 module attributes {
 204   gpu.container_module,
 205   spirv.target_env = #spirv.target_env<#spirv.vce<v1.3, [Kernel, Addresses, Groups, GroupNonUniformArithmetic, GroupUniformArithmeticKHR], []>, #spirv.resource_limits<>>
 206 } {
 207
 208 gpu.module @kernels {
 209   // CHECK-LABEL:  spirv.func @test
 210   //  CHECK-SAME: (%[[ARG:.*]]: i32)
 211   gpu.func @test(%arg : i32) kernel
 212     attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {
 213     // CHECK: %{{.*}} = spirv.KHR.GroupIMul <Workgroup> <Reduce> %[[ARG]] : i32
 214     %reduced = gpu.all_reduce mul %arg uniform {} : (i32) -> (i32)
 215     gpu.return
 216   }
 217 }
 218
 219 }
 220
 221 // -----
 222
 223 module attributes {
 224   gpu.container_module,
 225   spirv.target_env = #spirv.target_env<#spirv.vce<v1.3, [Kernel, Addresses, Groups, GroupNonUniformArithmetic, GroupUniformArithmeticKHR], []>, #spirv.resource_limits<>>
 226 } {
 227
 228 gpu.module @kernels {
 229   // CHECK-LABEL:  spirv.func @test
 230   //  CHECK-SAME: (%[[ARG:.*]]: i32)
 231   gpu.func @test(%arg : i32) kernel
 232     attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {
 233     // CHECK: %{{.*}} = spirv.GroupNonUniformIMul <Workgroup> <Reduce> %[[ARG]] : i32 -> i32
 234     %reduced = gpu.all_reduce mul %arg {} : (i32) -> (i32)
 235     gpu.return
 236   }
 237 }
 238
 239 }
 240
 241 // -----
 242
 243 module attributes {
 244   gpu.container_module,
 245   spirv.target_env = #spirv.target_env<#spirv.vce<v1.3, [Kernel, Addresses, Groups, GroupNonUniformArithmetic, GroupUniformArithmeticKHR], []>, #spirv.resource_limits<>>
 246 } {
 247
 248 gpu.module @kernels {
 249   // CHECK-LABEL:  spirv.func @test
 250   //  CHECK-SAME: (%[[ARG:.*]]: f32)
 251   gpu.func @test(%arg : f32) kernel
 252     attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {
 253     // CHECK: %{{.*}} = spirv.KHR.GroupFMul <Subgroup> <Reduce> %[[ARG]] : f32
 254     %reduced = gpu.subgroup_reduce mul %arg uniform : (f32) -> (f32)
 255     gpu.return
 256   }
 257 }
 258
 259 }
 260
 261 // -----
 262
 263 module attributes {
 264   gpu.container_module,
 265   spirv.target_env = #spirv.target_env<#spirv.vce<v1.3, [Kernel, Addresses, Groups, GroupNonUniformArithmetic, GroupUniformArithmeticKHR], []>, #spirv.resource_limits<>>
 266 } {
 267
 268 gpu.module @kernels {
 269   // CHECK-LABEL:  spirv.func @test
 270   //  CHECK-SAME: (%[[ARG:.*]]: f32)
 271   gpu.func @test(%arg : f32) kernel
 272     attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {
 273     // CHECK: %{{.*}} = spirv.GroupNonUniformFMul <Subgroup> <Reduce> %[[ARG]] : f32 -> f32
 274     %reduced = gpu.subgroup_reduce mul %arg : (f32) -> (f32)
 275     gpu.return
 276   }
 277 }
 278
 279 }
 280
 281 // -----
 282
 283 module attributes {
 284   gpu.container_module,
 285   spirv.target_env = #spirv.target_env<#spirv.vce<v1.3, [Kernel, Addresses, Groups, GroupNonUniformArithmetic, GroupUniformArithmeticKHR], []>, #spirv.resource_limits<>>
 286 } {
 287
 288 gpu.module @kernels {
 289   // CHECK-LABEL:  spirv.func @test
 290   //  CHECK-SAME: (%[[ARG:.*]]: i32)
 291   gpu.func @test(%arg : i32) kernel
 292     attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {
 293     // CHECK: %{{.*}} = spirv.KHR.GroupIMul <Subgroup> <Reduce> %[[ARG]] : i32
 294     %reduced = gpu.subgroup_reduce mul %arg uniform : (i32) -> (i32)
 295     gpu.return
 296   }
 297 }
 298
 299 }
 300
 301 // -----
 302
 303 module attributes {
 304   gpu.container_module,
 305   spirv.target_env = #spirv.target_env<#spirv.vce<v1.3, [Kernel, Addresses, Groups, GroupNonUniformArithmetic, GroupUniformArithmeticKHR], []>, #spirv.resource_limits<>>
 306 } {
 307
 308 gpu.module @kernels {
 309   // CHECK-LABEL:  spirv.func @test
 310   //  CHECK-SAME: (%[[ARG:.*]]: i32)
 311   gpu.func @test(%arg : i32) kernel
 312     attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {
 313     // CHECK: %{{.*}} = spirv.GroupNonUniformIMul <Subgroup> <Reduce> %[[ARG]] : i32 -> i32
 314     %reduced = gpu.subgroup_reduce mul %arg : (i32) -> (i32)
 315     gpu.return
 316   }
 317 }
 318
 319 }
 320
 321 // -----
 322
 323 module attributes {
 324   gpu.container_module,
 325   spirv.target_env = #spirv.target_env<#spirv.vce<v1.3, [Kernel, Addresses, Groups, GroupNonUniformArithmetic, GroupUniformArithmeticKHR], []>, #spirv.resource_limits<>>
 326 } {
 327
 328 gpu.module @kernels {
 329   // CHECK-LABEL:  spirv.func @test
 330   //  CHECK-SAME: (%[[ARG:.*]]: f32)
 331   gpu.func @test(%arg : f32) kernel
 332     attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {
 333     // CHECK: %{{.*}} = spirv.GroupFMin <Workgroup> <Reduce> %[[ARG]] : f32
 334     %reduced = gpu.all_reduce minnumf %arg uniform {} : (f32) -> (f32)
 335     gpu.return
 336   }
 337 }
 338
 339 }
 340
 341 // -----
 342
 343 module attributes {
 344   gpu.container_module,
 345   spirv.target_env = #spirv.target_env<#spirv.vce<v1.3, [Kernel, Addresses, Groups, GroupNonUniformArithmetic, GroupUniformArithmeticKHR], []>, #spirv.resource_limits<>>
 346 } {
 347
 348 gpu.module @kernels {
 349   // CHECK-LABEL:  spirv.func @test
 350   //  CHECK-SAME: (%[[ARG:.*]]: f32)
 351   gpu.func @test(%arg : f32) kernel
 352     attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {
 353     // CHECK: %{{.*}} = spirv.GroupNonUniformFMin <Workgroup> <Reduce> %[[ARG]] : f32 -> f32
 354     %reduced = gpu.all_reduce minnumf %arg {} : (f32) -> (f32)
 355     gpu.return
 356   }
 357 }
 358
 359 }
 360
 361 // -----
 362
 363 module attributes {
 364   gpu.container_module,
 365   spirv.target_env = #spirv.target_env<#spirv.vce<v1.3, [Kernel, Addresses, Groups, GroupNonUniformArithmetic, GroupUniformArithmeticKHR], []>, #spirv.resource_limits<>>
 366 } {
 367
 368 gpu.module @kernels {
 369   // CHECK-LABEL:  spirv.func @test
 370   //  CHECK-SAME: (%[[ARG:.*]]: i32)
 371   gpu.func @test(%arg : i32) kernel
 372     attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {
 373     // CHECK: %{{.*}} = spirv.GroupSMin <Workgroup> <Reduce> %[[ARG]] : i32
 374     // CHECK: %{{.*}} = spirv.GroupUMin <Workgroup> <Reduce> %[[ARG]] : i32
 375     %r0 = gpu.all_reduce minsi %arg uniform {} : (i32) -> (i32)
 376     %r1 = gpu.all_reduce minui %arg uniform {} : (i32) -> (i32)
 377     gpu.return
 378   }
 379 }
 380
 381 }
 382
 383 // -----
 384
 385 module attributes {
 386   gpu.container_module,
 387   spirv.target_env = #spirv.target_env<#spirv.vce<v1.3, [Kernel, Addresses, Groups, GroupNonUniformArithmetic, GroupUniformArithmeticKHR], []>, #spirv.resource_limits<>>
 388 } {
 389
 390 gpu.module @kernels {
 391   // CHECK-LABEL:  spirv.func @test
 392   //  CHECK-SAME: (%[[ARG:.*]]: i32)
 393   gpu.func @test(%arg : i32) kernel
 394     attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {
 395     // CHECK: %{{.*}} = spirv.GroupNonUniformUMin <Workgroup> <Reduce> %[[ARG]] : i32 -> i32
 396     %r0 = gpu.all_reduce minsi %arg {} : (i32) -> (i32)
 397     %r1 = gpu.all_reduce minui %arg {} : (i32) -> (i32)
 398     gpu.return
 399   }
 400 }
 401
 402 }
 403
 404 // -----
 405
 406 module attributes {
 407   gpu.container_module,
 408   spirv.target_env = #spirv.target_env<#spirv.vce<v1.3, [Kernel, Addresses, Groups, GroupNonUniformArithmetic, GroupUniformArithmeticKHR], []>, #spirv.resource_limits<>>
 409 } {
 410
 411 gpu.module @kernels {
 412   // CHECK-LABEL:  spirv.func @test
 413   //  CHECK-SAME: (%[[ARG:.*]]: f32)
 414   gpu.func @test(%arg : f32) kernel
 415     attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {
 416     // CHECK: %{{.*}} = spirv.GroupFMin <Subgroup> <Reduce> %[[ARG]] : f32
 417     %reduced = gpu.subgroup_reduce minnumf %arg uniform : (f32) -> (f32)
 418     gpu.return
 419   }
 420 }
 421
 422 }
 423
 424 // -----
 425
 426 module attributes {
 427   gpu.container_module,
 428   spirv.target_env = #spirv.target_env<#spirv.vce<v1.3, [Kernel, Addresses, Groups, GroupNonUniformArithmetic, GroupUniformArithmeticKHR], []>, #spirv.resource_limits<>>
 429 } {
 430
 431 gpu.module @kernels {
 432   // CHECK-LABEL:  spirv.func @test
 433   //  CHECK-SAME: (%[[ARG:.*]]: f32)
 434   gpu.func @test(%arg : f32) kernel
 435     attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {
 436     // CHECK: %{{.*}} = spirv.GroupNonUniformFMin <Subgroup> <Reduce> %[[ARG]] : f32 -> f32
 437     %reduced = gpu.subgroup_reduce minnumf %arg : (f32) -> (f32)
 438     gpu.return
 439   }
 440 }
 441
 442 }
 443
 444 // -----
 445
 446 module attributes {
 447   gpu.container_module,
 448   spirv.target_env = #spirv.target_env<#spirv.vce<v1.3, [Kernel, Addresses, Groups, GroupNonUniformArithmetic, GroupUniformArithmeticKHR], []>, #spirv.resource_limits<>>
 449 } {
 450
 451 gpu.module @kernels {
 452   // CHECK-LABEL:  spirv.func @test
 453   //  CHECK-SAME: (%[[ARG:.*]]: i32)
 454   gpu.func @test(%arg : i32) kernel
 455     attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {
 456     // CHECK: %{{.*}} = spirv.GroupSMin <Subgroup> <Reduce> %[[ARG]] : i32
 457     // CHECK: %{{.*}} = spirv.GroupUMin <Subgroup> <Reduce> %[[ARG]] : i32
 458     %r0 = gpu.subgroup_reduce minsi %arg uniform : (i32) -> (i32)
 459     %r1 = gpu.subgroup_reduce minui %arg uniform : (i32) -> (i32)
 460     gpu.return
 461   }
 462 }
 463
 464 }
 465
 466 // -----
 467
 468 module attributes {
 469   gpu.container_module,
 470   spirv.target_env = #spirv.target_env<#spirv.vce<v1.3, [Kernel, Addresses, Groups, GroupNonUniformArithmetic, GroupUniformArithmeticKHR], []>, #spirv.resource_limits<>>
 471 } {
 472
 473 gpu.module @kernels {
 474   // CHECK-LABEL:  spirv.func @test
 475   //  CHECK-SAME: (%[[ARG:.*]]: i32)
 476   gpu.func @test(%arg : i32) kernel
 477     attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {
 478     // CHECK: %{{.*}} = spirv.GroupNonUniformSMin <Subgroup> <Reduce> %[[ARG]] : i32 -> i32
 479     // CHECK: %{{.*}} = spirv.GroupNonUniformUMin <Subgroup> <Reduce> %[[ARG]] : i32 -> i32
 480     %r0 = gpu.subgroup_reduce minsi %arg : (i32) -> (i32)
 481     %r1 = gpu.subgroup_reduce minui %arg : (i32) -> (i32)
 482     gpu.return
 483   }
 484 }
 485
 486 }
 487
 488 // -----
 489
 490 module attributes {
 491   gpu.container_module,
 492   spirv.target_env = #spirv.target_env<#spirv.vce<v1.3, [Kernel, Addresses, Groups, GroupNonUniformArithmetic, GroupUniformArithmeticKHR], []>, #spirv.resource_limits<>>
 493 } {
 494
 495 gpu.module @kernels {
 496   // CHECK-LABEL:  spirv.func @test
 497   //  CHECK-SAME: (%[[ARG:.*]]: f32)
 498   gpu.func @test(%arg : f32) kernel
 499     attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {
 500     // CHECK: %{{.*}} = spirv.GroupFMax <Workgroup> <Reduce> %[[ARG]] : f32
 501     %reduced = gpu.all_reduce maxnumf %arg uniform {} : (f32) -> (f32)
 502     gpu.return
 503   }
 504 }
 505
 506 }
 507
 508 // -----
 509
 510 module attributes {
 511   gpu.container_module,
 512   spirv.target_env = #spirv.target_env<#spirv.vce<v1.3, [Kernel, Addresses, Groups, GroupNonUniformArithmetic, GroupUniformArithmeticKHR], []>, #spirv.resource_limits<>>
 513 } {
 514
 515 gpu.module @kernels {
 516   // CHECK-LABEL:  spirv.func @test
 517   //  CHECK-SAME: (%[[ARG:.*]]: f32)
 518   gpu.func @test(%arg : f32) kernel
 519     attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {
 520     // CHECK: %{{.*}} = spirv.GroupNonUniformFMax <Workgroup> <Reduce> %[[ARG]] : f32 -> f32
 521     %reduced = gpu.all_reduce maxnumf %arg {} : (f32) -> (f32)
 522     gpu.return
 523   }
 524 }
 525
 526 }
 527
 528 // -----
 529
 530 module attributes {
 531   gpu.container_module,
 532   spirv.target_env = #spirv.target_env<#spirv.vce<v1.3, [Kernel, Addresses, Groups, GroupNonUniformArithmetic, GroupUniformArithmeticKHR], []>, #spirv.resource_limits<>>
 533 } {
 534
 535 gpu.module @kernels {
 536   // CHECK-LABEL:  spirv.func @test
 537   //  CHECK-SAME: (%[[ARG:.*]]: i32)
 538   gpu.func @test(%arg : i32) kernel
 539     attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {
 540     // CHECK: %{{.*}} = spirv.GroupSMax <Workgroup> <Reduce> %[[ARG]] : i32
 541     // CHECK: %{{.*}} = spirv.GroupUMax <Workgroup> <Reduce> %[[ARG]] : i32
 542     %r0 = gpu.all_reduce maxsi %arg uniform {} : (i32) -> (i32)
 543     %r1 = gpu.all_reduce maxui %arg uniform {} : (i32) -> (i32)
 544     gpu.return
 545   }
 546 }
 547
 548 }
 549
 550 // -----
 551
 552 module attributes {
 553   gpu.container_module,
 554   spirv.target_env = #spirv.target_env<#spirv.vce<v1.3, [Kernel, Addresses, Groups, GroupNonUniformArithmetic, GroupUniformArithmeticKHR], []>, #spirv.resource_limits<>>
 555 } {
 556
 557 gpu.module @kernels {
 558   // CHECK-LABEL:  spirv.func @test
 559   //  CHECK-SAME: (%[[ARG:.*]]: i32)
 560   gpu.func @test(%arg : i32) kernel
 561     attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {
 562     // CHECK: %{{.*}} = spirv.GroupNonUniformSMax <Workgroup> <Reduce> %[[ARG]] : i32 -> i32
 563     // CHECK: %{{.*}} = spirv.GroupNonUniformUMax <Workgroup> <Reduce> %[[ARG]] : i32 -> i32
 564     %r0 = gpu.all_reduce maxsi %arg {} : (i32) -> (i32)
 565     %r1 = gpu.all_reduce maxui %arg {} : (i32) -> (i32)
 566     gpu.return
 567   }
 568 }
 569
 570 }
 571
 572 // -----
 573
 574 module attributes {
 575   gpu.container_module,
 576   spirv.target_env = #spirv.target_env<#spirv.vce<v1.3, [Kernel, Addresses, Groups, GroupNonUniformArithmetic, GroupUniformArithmeticKHR], []>, #spirv.resource_limits<>>
 577 } {
 578
 579 gpu.module @kernels {
 580   // CHECK-LABEL:  spirv.func @test
 581   //  CHECK-SAME: (%[[ARG:.*]]: f32)
 582   gpu.func @test(%arg : f32) kernel
 583     attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {
 584     // CHECK: %{{.*}} = spirv.GroupFMax <Subgroup> <Reduce> %[[ARG]] : f32
 585     %reduced = gpu.subgroup_reduce maxnumf %arg uniform : (f32) -> (f32)
 586     gpu.return
 587   }
 588 }
 589
 590 }
 591
 592 // -----
 593
 594 module attributes {
 595   gpu.container_module,
 596   spirv.target_env = #spirv.target_env<#spirv.vce<v1.3, [Kernel, Addresses, Groups, GroupNonUniformArithmetic, GroupUniformArithmeticKHR], []>, #spirv.resource_limits<>>
 597 } {
 598
 599 gpu.module @kernels {
 600   // CHECK-LABEL:  spirv.func @test
 601   //  CHECK-SAME: (%[[ARG:.*]]: f32)
 602   gpu.func @test(%arg : f32) kernel
 603     attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {
 604     // CHECK: %{{.*}} = spirv.GroupNonUniformFMax <Subgroup> <Reduce> %[[ARG]] : f32 -> f32
 605     %reduced = gpu.subgroup_reduce maxnumf %arg : (f32) -> (f32)
 606     gpu.return
 607   }
 608 }
 609
 610 }
 611
 612 // -----
 613
 614 module attributes {
 615   gpu.container_module,
 616   spirv.target_env = #spirv.target_env<#spirv.vce<v1.3, [Kernel, Addresses, Groups, GroupNonUniformArithmetic, GroupUniformArithmeticKHR], []>, #spirv.resource_limits<>>
 617 } {
 618
 619 gpu.module @kernels {
 620   // CHECK-LABEL:  spirv.func @test
 621   //  CHECK-SAME: (%[[ARG:.*]]: i32)
 622   gpu.func @test(%arg : i32) kernel
 623     attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {
 624     // CHECK: %{{.*}} = spirv.GroupSMax <Subgroup> <Reduce> %[[ARG]] : i32
 625     // CHECK: %{{.*}} = spirv.GroupUMax <Subgroup> <Reduce> %[[ARG]] : i32
 626     %r0 = gpu.subgroup_reduce maxsi %arg uniform : (i32) -> (i32)
 627     %r1 = gpu.subgroup_reduce maxui %arg uniform : (i32) -> (i32)
 628     gpu.return
 629   }
 630 }
 631
 632 }
 633
 634 // -----
 635
 636 module attributes {
 637   gpu.container_module,
 638   spirv.target_env = #spirv.target_env<#spirv.vce<v1.3, [Kernel, Addresses, Groups, GroupNonUniformArithmetic, GroupUniformArithmeticKHR], []>, #spirv.resource_limits<>>
 639 } {
 640
 641 gpu.module @kernels {
 642   // CHECK-LABEL:  spirv.func @test
 643   //  CHECK-SAME: (%[[ARG:.*]]: i32)
 644   gpu.func @test(%arg : i32) kernel
 645     attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {
 646     // CHECK: %{{.*}} = spirv.GroupNonUniformSMax <Subgroup> <Reduce> %[[ARG]] : i32 -> i32
 647     // CHECK: %{{.*}} = spirv.GroupNonUniformUMax <Subgroup> <Reduce> %[[ARG]] : i32 -> i32
 648     %r0 = gpu.subgroup_reduce maxsi %arg : (i32) -> (i32)
 649     %r1 = gpu.subgroup_reduce maxui %arg : (i32) -> (i32)
 650     gpu.return
 651   }
 652 }
 653
 654 }
 655
 656 // -----
 657
 658 module attributes {
 659   gpu.container_module,
 660   spirv.target_env = #spirv.target_env<#spirv.vce<v1.3, [Kernel, Addresses, Groups, GroupNonUniformArithmetic, GroupUniformArithmeticKHR], []>, #spirv.resource_limits<>>
 661 } {
 662
 663 gpu.module @kernels {
 664   // CHECK-LABEL:  spirv.func @test
 665   //  CHECK-SAME: (%[[ARG:.*]]: i32)
 666   gpu.func @test(%arg : vector<1xi32>) kernel
 667     attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {
 668     // CHECK: %{{.*}} = spirv.GroupNonUniformSMax <Subgroup> <Reduce> %[[ARG]] : i32 -> i32
 669     %r0 = gpu.subgroup_reduce maxsi %arg : (vector<1xi32>) -> (vector<1xi32>)
 670     gpu.return
 671   }
 672 }
 673
 674 }
 675
 676 // -----
 677
 678 // TODO: Handle boolean reductions.
 679
 680 module attributes {
 681   gpu.container_module,
 682   spirv.target_env = #spirv.target_env<#spirv.vce<v1.3, [Kernel, Addresses, Groups, GroupNonUniformArithmetic, GroupUniformArithmeticKHR], []>, #spirv.resource_limits<>>
 683 } {
 684
 685 gpu.module @kernels {
 686   gpu.func @add(%arg : i1) kernel
 687     attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {
 688     // expected-error @+1 {{failed to legalize operation 'gpu.subgroup_reduce'}}
 689     %r0 = gpu.subgroup_reduce add %arg : (i1) -> (i1)
 690     gpu.return
 691   }
 692 }
 693 }
 694
 695 // -----
 696
 697 module attributes {
 698   gpu.container_module,
 699   spirv.target_env = #spirv.target_env<#spirv.vce<v1.3, [Kernel, Addresses, Groups, GroupNonUniformArithmetic, GroupUniformArithmeticKHR], []>, #spirv.resource_limits<>>
 700 } {
 701 gpu.module @kernels {
 702   gpu.func @mul(%arg : i1) kernel
 703     attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {
 704     // expected-error @+1 {{failed to legalize operation 'gpu.subgroup_reduce'}}
 705     %r0 = gpu.subgroup_reduce mul %arg : (i1) -> (i1)
 706     gpu.return
 707   }
 708 }
 709 }
 710
 711 // -----
 712
 713 module attributes {
 714   gpu.container_module,
 715   spirv.target_env = #spirv.target_env<#spirv.vce<v1.3, [Kernel, Addresses, Groups, GroupNonUniformArithmetic, GroupUniformArithmeticKHR], []>, #spirv.resource_limits<>>
 716 } {
 717 gpu.module @kernels {
 718   gpu.func @minsi(%arg : i1) kernel
 719     attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {
 720     // expected-error @+1 {{failed to legalize operation 'gpu.subgroup_reduce'}}
 721     %r0 = gpu.subgroup_reduce minsi %arg : (i1) -> (i1)
 722     gpu.return
 723   }
 724 }
 725 }
 726
 727 // -----
 728
 729 module attributes {
 730   gpu.container_module,
 731   spirv.target_env = #spirv.target_env<#spirv.vce<v1.3, [Kernel, Addresses, Groups, GroupNonUniformArithmetic, GroupUniformArithmeticKHR], []>, #spirv.resource_limits<>>
 732 } {
 733 gpu.module @kernels {
 734   gpu.func @minui(%arg : i1) kernel
 735     attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {
 736     // expected-error @+1 {{failed to legalize operation 'gpu.subgroup_reduce'}}
 737     %r0 = gpu.subgroup_reduce minui %arg : (i1) -> (i1)
 738     gpu.return
 739   }
 740 }
 741 }
 742
 743 // -----
 744
 745 module attributes {
 746   gpu.container_module,
 747   spirv.target_env = #spirv.target_env<#spirv.vce<v1.3, [Kernel, Addresses, Groups, GroupNonUniformArithmetic, GroupUniformArithmeticKHR], []>, #spirv.resource_limits<>>
 748 } {
 749 gpu.module @kernels {
 750   gpu.func @maxsi(%arg : i1) kernel
 751     attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {
 752     // expected-error @+1 {{failed to legalize operation 'gpu.subgroup_reduce'}}
 753     %r0 = gpu.subgroup_reduce maxsi %arg : (i1) -> (i1)
 754     gpu.return
 755   }
 756 }
 757 }
 758
 759 // -----
 760
 761 module attributes {
 762   gpu.container_module,
 763   spirv.target_env = #spirv.target_env<#spirv.vce<v1.3, [Kernel, Addresses, Groups, GroupNonUniformArithmetic, GroupUniformArithmeticKHR], []>, #spirv.resource_limits<>>
 764 } {
 765 gpu.module @kernels {
 766   gpu.func @maxui(%arg : i1) kernel
 767     attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {
 768     // expected-error @+1 {{failed to legalize operation 'gpu.subgroup_reduce'}}
 769     %r0 = gpu.subgroup_reduce maxui %arg : (i1) -> (i1)
 770     gpu.return
 771   }
 772 }
 773 }
 774
 775 // -----
 776
 777 // Vector reductions need to be lowered to scalar reductions first.
 778
 779 module attributes {
 780   gpu.container_module,
 781   spirv.target_env = #spirv.target_env<#spirv.vce<v1.3, [Kernel, Addresses, Groups, GroupNonUniformArithmetic, GroupUniformArithmeticKHR], []>, #spirv.resource_limits<>>
 782 } {
 783 gpu.module @kernels {
 784   gpu.func @maxui(%arg : vector<2xi32>) kernel
 785     attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {
 786     // expected-error @+1 {{failed to legalize operation 'gpu.subgroup_reduce'}}
 787     %r0 = gpu.subgroup_reduce maxui %arg : (vector<2xi32>) -> (vector<2xi32>)
 788     gpu.return
 789   }
 790 }
 791 }