833 lines · plain
1// RUN: mlir-opt -split-input-file -convert-gpu-to-spirv -verify-diagnostics %s -o - | FileCheck %s2 3module 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 8gpu.module @kernels {9 // CHECK-LABEL: spirv.func @test10 // CHECK-SAME: (%[[ARG:.*]]: f32)11 gpu.func @test(%arg : f32) kernel12 attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {13 // CHECK: %{{.*}} = spirv.GroupFAdd <Workgroup> <Reduce> %[[ARG]] : f3214 %reduced = gpu.all_reduce add %arg uniform {} : (f32) -> (f32)15 gpu.return16 }17}18 19}20 21// -----22 23module 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 28gpu.module @kernels {29 // CHECK-LABEL: spirv.func @test30 // CHECK-SAME: (%[[ARG:.*]]: f32)31 gpu.func @test(%arg : f32) kernel32 attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {33 // CHECK: %{{.*}} = spirv.GroupNonUniformFAdd <Workgroup> <Reduce> %[[ARG]] : f32 -> f3234 %reduced = gpu.all_reduce add %arg {} : (f32) -> (f32)35 gpu.return36 }37}38 39}40 41// -----42 43module 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 48gpu.module @kernels {49 // CHECK-LABEL: spirv.func @test50 // CHECK-SAME: (%[[ARG:.*]]: i32)51 gpu.func @test(%arg : i32) kernel52 attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {53 // CHECK: %{{.*}} = spirv.GroupIAdd <Workgroup> <Reduce> %[[ARG]] : i3254 %reduced = gpu.all_reduce add %arg uniform {} : (i32) -> (i32)55 gpu.return56 }57}58 59}60 61// -----62 63module 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 68gpu.module @kernels {69 // CHECK-LABEL: spirv.func @test70 // CHECK-SAME: (%[[ARG:.*]]: i32)71 gpu.func @test(%arg : i32) kernel72 attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {73 // CHECK: %{{.*}} = spirv.GroupNonUniformIAdd <Workgroup> <Reduce> %[[ARG]] : i32 -> i3274 %reduced = gpu.all_reduce add %arg {} : (i32) -> (i32)75 gpu.return76 }77}78 79}80 81// -----82 83module 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 88gpu.module @kernels {89 // CHECK-LABEL: spirv.func @test90 // CHECK-SAME: (%[[ARG:.*]]: f32)91 gpu.func @test(%arg : f32) kernel92 attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {93 // CHECK: %{{.*}} = spirv.GroupFAdd <Subgroup> <Reduce> %[[ARG]] : f3294 %reduced = gpu.subgroup_reduce add %arg uniform : (f32) -> (f32)95 gpu.return96 }97}98 99}100 101// -----102 103module 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 108gpu.module @kernels {109 // CHECK-LABEL: spirv.func @test110 // CHECK-SAME: (%[[ARG:.*]]: f32)111 gpu.func @test(%arg : f32) kernel112 attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {113 // CHECK: %{{.*}} = spirv.GroupNonUniformFAdd <Subgroup> <Reduce> %[[ARG]] : f32 -> f32114 %reduced = gpu.subgroup_reduce add %arg : (f32) -> (f32)115 gpu.return116 }117}118 119}120 121// -----122 123module 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 128gpu.module @kernels {129 // CHECK-LABEL: spirv.func @test130 // CHECK-SAME: (%[[ARG:.*]]: i32)131 gpu.func @test(%arg : i32) kernel132 attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {133 // CHECK: %{{.*}} = spirv.GroupIAdd <Subgroup> <Reduce> %[[ARG]] : i32134 %reduced = gpu.subgroup_reduce add %arg uniform : (i32) -> (i32)135 gpu.return136 }137}138 139}140 141// -----142 143module 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 148gpu.module @kernels {149 // CHECK-LABEL: spirv.func @test150 // CHECK-SAME: (%[[ARG:.*]]: i32)151 gpu.func @test(%arg : i32) kernel152 attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {153 // CHECK: %{{.*}} = spirv.GroupNonUniformIAdd <Subgroup> <Reduce> %[[ARG]] : i32 -> i32154 %reduced = gpu.subgroup_reduce add %arg : (i32) -> (i32)155 gpu.return156 }157}158 159}160 161// -----162 163module 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 168gpu.module @kernels {169 // CHECK-LABEL: spirv.func @test170 // CHECK-SAME: (%[[ARG:.*]]: f32)171 gpu.func @test(%arg : f32) kernel172 attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {173 // CHECK: %{{.*}} = spirv.KHR.GroupFMul <Workgroup> <Reduce> %[[ARG]] : f32174 %reduced = gpu.all_reduce mul %arg uniform {} : (f32) -> (f32)175 gpu.return176 }177}178 179}180 181// -----182 183module 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 188gpu.module @kernels {189 // CHECK-LABEL: spirv.func @test190 // CHECK-SAME: (%[[ARG:.*]]: f32)191 gpu.func @test(%arg : f32) kernel192 attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {193 // CHECK: %{{.*}} = spirv.GroupNonUniformFMul <Workgroup> <Reduce> %[[ARG]] : f32 -> f32194 %reduced = gpu.all_reduce mul %arg {} : (f32) -> (f32)195 gpu.return196 }197}198 199}200 201// -----202 203module 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 208gpu.module @kernels {209 // CHECK-LABEL: spirv.func @test210 // CHECK-SAME: (%[[ARG:.*]]: i32)211 gpu.func @test(%arg : i32) kernel212 attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {213 // CHECK: %{{.*}} = spirv.KHR.GroupIMul <Workgroup> <Reduce> %[[ARG]] : i32214 %reduced = gpu.all_reduce mul %arg uniform {} : (i32) -> (i32)215 gpu.return216 }217}218 219}220 221// -----222 223module 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 228gpu.module @kernels {229 // CHECK-LABEL: spirv.func @test230 // CHECK-SAME: (%[[ARG:.*]]: i32)231 gpu.func @test(%arg : i32) kernel232 attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {233 // CHECK: %{{.*}} = spirv.GroupNonUniformIMul <Workgroup> <Reduce> %[[ARG]] : i32 -> i32234 %reduced = gpu.all_reduce mul %arg {} : (i32) -> (i32)235 gpu.return236 }237}238 239}240 241// -----242 243module 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 248gpu.module @kernels {249 // CHECK-LABEL: spirv.func @test250 // CHECK-SAME: (%[[ARG:.*]]: f32)251 gpu.func @test(%arg : f32) kernel252 attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {253 // CHECK: %{{.*}} = spirv.KHR.GroupFMul <Subgroup> <Reduce> %[[ARG]] : f32254 %reduced = gpu.subgroup_reduce mul %arg uniform : (f32) -> (f32)255 gpu.return256 }257}258 259}260 261// -----262 263module 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 268gpu.module @kernels {269 // CHECK-LABEL: spirv.func @test270 // CHECK-SAME: (%[[ARG:.*]]: f32)271 gpu.func @test(%arg : f32) kernel272 attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {273 // CHECK: %{{.*}} = spirv.GroupNonUniformFMul <Subgroup> <Reduce> %[[ARG]] : f32 -> f32274 %reduced = gpu.subgroup_reduce mul %arg : (f32) -> (f32)275 gpu.return276 }277}278 279}280 281// -----282 283module 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 288gpu.module @kernels {289 // CHECK-LABEL: spirv.func @test290 // CHECK-SAME: (%[[ARG:.*]]: i32)291 gpu.func @test(%arg : i32) kernel292 attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {293 // CHECK: %{{.*}} = spirv.KHR.GroupIMul <Subgroup> <Reduce> %[[ARG]] : i32294 %reduced = gpu.subgroup_reduce mul %arg uniform : (i32) -> (i32)295 gpu.return296 }297}298 299}300 301// -----302 303module 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 308gpu.module @kernels {309 // CHECK-LABEL: spirv.func @test310 // CHECK-SAME: (%[[ARG:.*]]: i32)311 gpu.func @test(%arg : i32) kernel312 attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {313 // CHECK: %{{.*}} = spirv.GroupNonUniformIMul <Subgroup> <Reduce> %[[ARG]] : i32 -> i32314 %reduced = gpu.subgroup_reduce mul %arg : (i32) -> (i32)315 gpu.return316 }317}318 319}320 321// -----322 323module 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 328gpu.module @kernels {329 // CHECK-LABEL: spirv.func @test330 // CHECK-SAME: (%[[ARG:.*]]: f32)331 gpu.func @test(%arg : f32) kernel332 attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {333 // CHECK: %{{.*}} = spirv.GroupFMin <Workgroup> <Reduce> %[[ARG]] : f32334 %reduced = gpu.all_reduce minnumf %arg uniform {} : (f32) -> (f32)335 gpu.return336 }337}338 339}340 341// -----342 343module 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 348gpu.module @kernels {349 // CHECK-LABEL: spirv.func @test350 // CHECK-SAME: (%[[ARG:.*]]: f32)351 gpu.func @test(%arg : f32) kernel352 attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {353 // CHECK: %{{.*}} = spirv.GroupNonUniformFMin <Workgroup> <Reduce> %[[ARG]] : f32 -> f32354 %reduced = gpu.all_reduce minnumf %arg {} : (f32) -> (f32)355 gpu.return356 }357}358 359}360 361// -----362 363module 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 368gpu.module @kernels {369 // CHECK-LABEL: spirv.func @test370 // CHECK-SAME: (%[[ARG:.*]]: i32)371 gpu.func @test(%arg : i32) kernel372 attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {373 // CHECK: %{{.*}} = spirv.GroupSMin <Workgroup> <Reduce> %[[ARG]] : i32374 // CHECK: %{{.*}} = spirv.GroupUMin <Workgroup> <Reduce> %[[ARG]] : i32375 %r0 = gpu.all_reduce minsi %arg uniform {} : (i32) -> (i32)376 %r1 = gpu.all_reduce minui %arg uniform {} : (i32) -> (i32)377 gpu.return378 }379}380 381}382 383// -----384 385module 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 390gpu.module @kernels {391 // CHECK-LABEL: spirv.func @test392 // CHECK-SAME: (%[[ARG:.*]]: i32)393 gpu.func @test(%arg : i32) kernel394 attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {395 // CHECK: %{{.*}} = spirv.GroupNonUniformUMin <Workgroup> <Reduce> %[[ARG]] : i32 -> i32396 %r0 = gpu.all_reduce minsi %arg {} : (i32) -> (i32)397 %r1 = gpu.all_reduce minui %arg {} : (i32) -> (i32)398 gpu.return399 }400}401 402}403 404// -----405 406module 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 411gpu.module @kernels {412 // CHECK-LABEL: spirv.func @test413 // CHECK-SAME: (%[[ARG:.*]]: f32)414 gpu.func @test(%arg : f32) kernel415 attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {416 // CHECK: %{{.*}} = spirv.GroupFMin <Subgroup> <Reduce> %[[ARG]] : f32417 %reduced = gpu.subgroup_reduce minnumf %arg uniform : (f32) -> (f32)418 gpu.return419 }420}421 422}423 424// -----425 426module 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 431gpu.module @kernels {432 // CHECK-LABEL: spirv.func @test433 // CHECK-SAME: (%[[ARG:.*]]: f32)434 gpu.func @test(%arg : f32) kernel435 attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {436 // CHECK: %{{.*}} = spirv.GroupNonUniformFMin <Subgroup> <Reduce> %[[ARG]] : f32 -> f32437 %reduced = gpu.subgroup_reduce minnumf %arg : (f32) -> (f32)438 gpu.return439 }440}441 442}443 444// -----445 446module 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 451gpu.module @kernels {452 // CHECK-LABEL: spirv.func @test453 // CHECK-SAME: (%[[ARG:.*]]: i32)454 gpu.func @test(%arg : i32) kernel455 attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {456 // CHECK: %{{.*}} = spirv.GroupSMin <Subgroup> <Reduce> %[[ARG]] : i32457 // CHECK: %{{.*}} = spirv.GroupUMin <Subgroup> <Reduce> %[[ARG]] : i32458 %r0 = gpu.subgroup_reduce minsi %arg uniform : (i32) -> (i32)459 %r1 = gpu.subgroup_reduce minui %arg uniform : (i32) -> (i32)460 gpu.return461 }462}463 464}465 466// -----467 468module 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 473gpu.module @kernels {474 // CHECK-LABEL: spirv.func @test475 // CHECK-SAME: (%[[ARG:.*]]: i32)476 gpu.func @test(%arg : i32) kernel477 attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {478 // CHECK: %{{.*}} = spirv.GroupNonUniformSMin <Subgroup> <Reduce> %[[ARG]] : i32 -> i32479 // CHECK: %{{.*}} = spirv.GroupNonUniformUMin <Subgroup> <Reduce> %[[ARG]] : i32 -> i32480 %r0 = gpu.subgroup_reduce minsi %arg : (i32) -> (i32)481 %r1 = gpu.subgroup_reduce minui %arg : (i32) -> (i32)482 gpu.return483 }484}485 486}487 488// -----489 490module 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 495gpu.module @kernels {496 // CHECK-LABEL: spirv.func @test497 // CHECK-SAME: (%[[ARG:.*]]: f32)498 gpu.func @test(%arg : f32) kernel499 attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {500 // CHECK: %{{.*}} = spirv.GroupFMax <Workgroup> <Reduce> %[[ARG]] : f32501 %reduced = gpu.all_reduce maxnumf %arg uniform {} : (f32) -> (f32)502 gpu.return503 }504}505 506}507 508// -----509 510module 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 515gpu.module @kernels {516 // CHECK-LABEL: spirv.func @test517 // CHECK-SAME: (%[[ARG:.*]]: f32)518 gpu.func @test(%arg : f32) kernel519 attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {520 // CHECK: %{{.*}} = spirv.GroupNonUniformFMax <Workgroup> <Reduce> %[[ARG]] : f32 -> f32521 %reduced = gpu.all_reduce maxnumf %arg {} : (f32) -> (f32)522 gpu.return523 }524}525 526}527 528// -----529 530module 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 535gpu.module @kernels {536 // CHECK-LABEL: spirv.func @test537 // CHECK-SAME: (%[[ARG:.*]]: i32)538 gpu.func @test(%arg : i32) kernel539 attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {540 // CHECK: %{{.*}} = spirv.GroupSMax <Workgroup> <Reduce> %[[ARG]] : i32541 // CHECK: %{{.*}} = spirv.GroupUMax <Workgroup> <Reduce> %[[ARG]] : i32542 %r0 = gpu.all_reduce maxsi %arg uniform {} : (i32) -> (i32)543 %r1 = gpu.all_reduce maxui %arg uniform {} : (i32) -> (i32)544 gpu.return545 }546}547 548}549 550// -----551 552module 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 557gpu.module @kernels {558 // CHECK-LABEL: spirv.func @test559 // CHECK-SAME: (%[[ARG:.*]]: i32)560 gpu.func @test(%arg : i32) kernel561 attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {562 // CHECK: %{{.*}} = spirv.GroupNonUniformSMax <Workgroup> <Reduce> %[[ARG]] : i32 -> i32563 // CHECK: %{{.*}} = spirv.GroupNonUniformUMax <Workgroup> <Reduce> %[[ARG]] : i32 -> i32564 %r0 = gpu.all_reduce maxsi %arg {} : (i32) -> (i32)565 %r1 = gpu.all_reduce maxui %arg {} : (i32) -> (i32)566 gpu.return567 }568}569 570}571 572// -----573 574module 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 579gpu.module @kernels {580 // CHECK-LABEL: spirv.func @test581 // CHECK-SAME: (%[[ARG:.*]]: f32)582 gpu.func @test(%arg : f32) kernel583 attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {584 // CHECK: %{{.*}} = spirv.GroupFMax <Subgroup> <Reduce> %[[ARG]] : f32585 %reduced = gpu.subgroup_reduce maxnumf %arg uniform : (f32) -> (f32)586 gpu.return587 }588}589 590}591 592// -----593 594module 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 599gpu.module @kernels {600 // CHECK-LABEL: spirv.func @test601 // CHECK-SAME: (%[[ARG:.*]]: f32)602 gpu.func @test(%arg : f32) kernel603 attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {604 // CHECK: %{{.*}} = spirv.GroupNonUniformFMax <Subgroup> <Reduce> %[[ARG]] : f32 -> f32605 %reduced = gpu.subgroup_reduce maxnumf %arg : (f32) -> (f32)606 gpu.return607 }608}609 610}611 612// -----613 614module 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 619gpu.module @kernels {620 // CHECK-LABEL: spirv.func @test621 // CHECK-SAME: (%[[ARG:.*]]: i32)622 gpu.func @test(%arg : i32) kernel623 attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {624 // CHECK: %{{.*}} = spirv.GroupSMax <Subgroup> <Reduce> %[[ARG]] : i32625 // CHECK: %{{.*}} = spirv.GroupUMax <Subgroup> <Reduce> %[[ARG]] : i32626 %r0 = gpu.subgroup_reduce maxsi %arg uniform : (i32) -> (i32)627 %r1 = gpu.subgroup_reduce maxui %arg uniform : (i32) -> (i32)628 gpu.return629 }630}631 632}633 634// -----635 636module 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 641gpu.module @kernels {642 // CHECK-LABEL: spirv.func @test643 // CHECK-SAME: (%[[ARG:.*]]: i32)644 gpu.func @test(%arg : i32) kernel645 attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {646 // CHECK: %{{.*}} = spirv.GroupNonUniformSMax <Subgroup> <Reduce> %[[ARG]] : i32 -> i32647 // CHECK: %{{.*}} = spirv.GroupNonUniformUMax <Subgroup> <Reduce> %[[ARG]] : i32 -> i32648 %r0 = gpu.subgroup_reduce maxsi %arg : (i32) -> (i32)649 %r1 = gpu.subgroup_reduce maxui %arg : (i32) -> (i32)650 gpu.return651 }652}653 654}655 656// -----657 658module 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 663gpu.module @kernels {664 // CHECK-LABEL: spirv.func @test665 // CHECK-SAME: (%[[ARG:.*]]: i32)666 gpu.func @test(%arg : vector<1xi32>) kernel667 attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {668 // CHECK: %{{.*}} = spirv.GroupNonUniformSMax <Subgroup> <Reduce> %[[ARG]] : i32 -> i32669 %r0 = gpu.subgroup_reduce maxsi %arg : (vector<1xi32>) -> (vector<1xi32>)670 gpu.return671 }672}673 674}675 676// -----677 678// TODO: Handle boolean reductions.679 680module 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 685gpu.module @kernels {686 gpu.func @add(%arg : i1) kernel687 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.return691 }692}693}694 695// -----696 697module 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} {701gpu.module @kernels {702 gpu.func @mul(%arg : i1) kernel703 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.return707 }708}709}710 711// -----712 713module 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} {717gpu.module @kernels {718 gpu.func @minsi(%arg : i1) kernel719 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.return723 }724}725}726 727// -----728 729module 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} {733gpu.module @kernels {734 gpu.func @minui(%arg : i1) kernel735 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.return739 }740}741}742 743// -----744 745module 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} {749gpu.module @kernels {750 gpu.func @maxsi(%arg : i1) kernel751 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.return755 }756}757}758 759// -----760 761module 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} {765gpu.module @kernels {766 gpu.func @maxui(%arg : i1) kernel767 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.return771 }772}773}774 775// -----776 777// Vector reductions need to be lowered to scalar reductions first.778 779module 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} {783gpu.module @kernels {784 gpu.func @maxui(%arg : vector<2xi32>) kernel785 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.return789 }790}791}792 793// -----794 795module attributes {796 gpu.container_module,797 spirv.target_env = #spirv.target_env<#spirv.vce<v1.3, [Kernel, Addresses, Groups, GroupUniformArithmeticKHR, GroupNonUniformClustered], []>, #spirv.resource_limits<>>798} {799 800gpu.module @kernels {801 // CHECK-LABEL: spirv.func @test_subgroup_reduce_clustered802 // CHECK-SAME: (%[[ARG:.*]]: f32)803 // CHECK: %[[CLUSTER_SIZE:.*]] = spirv.Constant 8 : i32804 gpu.func @test_subgroup_reduce_clustered(%arg : f32) kernel805 attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {806 // CHECK: %{{.*}} = spirv.GroupNonUniformFAdd <Subgroup> <ClusteredReduce> %[[ARG]] cluster_size(%[[CLUSTER_SIZE]]) : f32, i32 -> f32807 %reduced = gpu.subgroup_reduce add %arg cluster(size = 8) : (f32) -> (f32)808 gpu.return809 }810}811 812}813 814// -----815 816// Subgrop reduce with cluster stride > 1 is not yet supported.817 818module attributes {819 gpu.container_module,820 spirv.target_env = #spirv.target_env<#spirv.vce<v1.3, [Kernel, Addresses, Groups, GroupUniformArithmeticKHR, GroupNonUniformClustered], []>, #spirv.resource_limits<>>821} {822 823gpu.module @kernels {824 gpu.func @test_invalid_subgroup_reduce_clustered_stride(%arg : f32) kernel825 attributes {spirv.entry_point_abi = #spirv.entry_point_abi<workgroup_size = [16, 1, 1]>} {826 // expected-error @+1 {{failed to legalize operation 'gpu.subgroup_reduce'}}827 %reduced = gpu.subgroup_reduce add %arg cluster(size = 8, stride = 2) : (f32) -> (f32)828 gpu.return829 }830}831 832}833