1// RUN: mlir-opt -mlir-disable-threading -test-spirv-target-env %s | FileCheck %s 2 3// Note: The following tests check that a spv.target_env can properly control 4// the conversion target and filter unavailable ops during the conversion. 5// We don't care about the op argument consistency too much; so certain enum 6// values for enum attributes may not make much sense for the test op. 7 8// spv.AtomicCompareExchangeWeak is available from SPIR-V 1.0 to 1.3 under 9// Kernel capability. 10// spv.AtomicCompareExchangeWeak has two memory semantics enum attribute, 11// whose value, if containing AtomicCounterMemory bit, additionally requires 12// AtomicStorage capability. 13 14// spv.BitReverse is available in all SPIR-V versions under Shader capability. 15 16// spv.GroupNonUniformBallot is available starting from SPIR-V 1.3 under 17// GroupNonUniform capability. 18 19// spv.SubgroupBallotKHR is available under in all SPIR-V versions under 20// SubgroupBallotKHR capability and SPV_KHR_shader_ballot extension. 21 22// The GeometryPointSize capability implies the Geometry capability, which 23// implies the Shader capability. 24 25// PhysicalStorageBuffer64 addressing model is available via extension 26// SPV_EXT_physical_storage_buffer or SPV_KHR_physical_storage_buffer; 27// both extensions are incorporated into SPIR-V 1.5. 28 29// Vulkan memory model is available via extension SPV_KHR_vulkan_memory_model, 30// which extensions are incorporated into SPIR-V 1.5. 31 32//===----------------------------------------------------------------------===// 33// MaxVersion 34//===----------------------------------------------------------------------===// 35 36// CHECK-LABEL: @cmp_exchange_weak_suitable_version_capabilities 37func @cmp_exchange_weak_suitable_version_capabilities(%ptr: !spv.ptr<i32, Workgroup>, %value: i32, %comparator: i32) -> i32 attributes { 38 spv.target_env = #spv.target_env<#spv.vce<v1.1, [Kernel, AtomicStorage], []>, {max_compute_workgroup_invocations = 128 : i32, max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}> 39} { 40 // CHECK: spv.AtomicCompareExchangeWeak "Workgroup" "AcquireRelease|AtomicCounterMemory" "Acquire" 41 %0 = "test.convert_to_atomic_compare_exchange_weak_op"(%ptr, %value, %comparator): (!spv.ptr<i32, Workgroup>, i32, i32) -> (i32) 42 return %0: i32 43} 44 45// CHECK-LABEL: @cmp_exchange_weak_unsupported_version 46func @cmp_exchange_weak_unsupported_version(%ptr: !spv.ptr<i32, Workgroup>, %value: i32, %comparator: i32) -> i32 attributes { 47 spv.target_env = #spv.target_env<#spv.vce<v1.4, [Kernel, AtomicStorage], []>, {max_compute_workgroup_invocations = 128 : i32, max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}> 48} { 49 // CHECK: test.convert_to_atomic_compare_exchange_weak_op 50 %0 = "test.convert_to_atomic_compare_exchange_weak_op"(%ptr, %value, %comparator): (!spv.ptr<i32, Workgroup>, i32, i32) -> (i32) 51 return %0: i32 52} 53 54//===----------------------------------------------------------------------===// 55// MinVersion 56//===----------------------------------------------------------------------===// 57 58// CHECK-LABEL: @group_non_uniform_ballot_suitable_version 59func @group_non_uniform_ballot_suitable_version(%predicate: i1) -> vector<4xi32> attributes { 60 spv.target_env = #spv.target_env<#spv.vce<v1.4, [GroupNonUniformBallot], []>, {max_compute_workgroup_invocations = 128 : i32, max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}> 61} { 62 // CHECK: spv.GroupNonUniformBallot "Workgroup" 63 %0 = "test.convert_to_group_non_uniform_ballot_op"(%predicate): (i1) -> (vector<4xi32>) 64 return %0: vector<4xi32> 65} 66 67// CHECK-LABEL: @group_non_uniform_ballot_unsupported_version 68func @group_non_uniform_ballot_unsupported_version(%predicate: i1) -> vector<4xi32> attributes { 69 spv.target_env = #spv.target_env<#spv.vce<v1.1, [GroupNonUniformBallot], []>, {max_compute_workgroup_invocations = 128 : i32, max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}> 70} { 71 // CHECK: test.convert_to_group_non_uniform_ballot_op 72 %0 = "test.convert_to_group_non_uniform_ballot_op"(%predicate): (i1) -> (vector<4xi32>) 73 return %0: vector<4xi32> 74} 75 76//===----------------------------------------------------------------------===// 77// Capability 78//===----------------------------------------------------------------------===// 79 80// CHECK-LABEL: @cmp_exchange_weak_missing_capability_kernel 81func @cmp_exchange_weak_missing_capability_kernel(%ptr: !spv.ptr<i32, Workgroup>, %value: i32, %comparator: i32) -> i32 attributes { 82 spv.target_env = #spv.target_env<#spv.vce<v1.3, [AtomicStorage], []>, {max_compute_workgroup_invocations = 128 : i32, max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}> 83} { 84 // CHECK: test.convert_to_atomic_compare_exchange_weak_op 85 %0 = "test.convert_to_atomic_compare_exchange_weak_op"(%ptr, %value, %comparator): (!spv.ptr<i32, Workgroup>, i32, i32) -> (i32) 86 return %0: i32 87} 88 89// CHECK-LABEL: @cmp_exchange_weak_missing_capability_atomic_storage 90func @cmp_exchange_weak_missing_capability_atomic_storage(%ptr: !spv.ptr<i32, Workgroup>, %value: i32, %comparator: i32) -> i32 attributes { 91 spv.target_env = #spv.target_env<#spv.vce<v1.3, [Kernel], []>, {max_compute_workgroup_invocations = 128 : i32, max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}> 92} { 93 // CHECK: test.convert_to_atomic_compare_exchange_weak_op 94 %0 = "test.convert_to_atomic_compare_exchange_weak_op"(%ptr, %value, %comparator): (!spv.ptr<i32, Workgroup>, i32, i32) -> (i32) 95 return %0: i32 96} 97 98// CHECK-LABEL: @subgroup_ballot_missing_capability 99func @subgroup_ballot_missing_capability(%predicate: i1) -> vector<4xi32> attributes { 100 spv.target_env = #spv.target_env<#spv.vce<v1.4, [], [SPV_KHR_shader_ballot]>, {max_compute_workgroup_invocations = 128 : i32, max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}> 101} { 102 // CHECK: test.convert_to_subgroup_ballot_op 103 %0 = "test.convert_to_subgroup_ballot_op"(%predicate): (i1) -> (vector<4xi32>) 104 return %0: vector<4xi32> 105} 106 107// CHECK-LABEL: @bit_reverse_directly_implied_capability 108func @bit_reverse_directly_implied_capability(%operand: i32) -> i32 attributes { 109 spv.target_env = #spv.target_env<#spv.vce<v1.0, [Geometry], []>, {max_compute_workgroup_invocations = 128 : i32, max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}> 110} { 111 // CHECK: spv.BitReverse 112 %0 = "test.convert_to_bit_reverse_op"(%operand): (i32) -> (i32) 113 return %0: i32 114} 115 116// CHECK-LABEL: @bit_reverse_recursively_implied_capability 117func @bit_reverse_recursively_implied_capability(%operand: i32) -> i32 attributes { 118 spv.target_env = #spv.target_env<#spv.vce<v1.0, [GeometryPointSize], []>, {max_compute_workgroup_invocations = 128 : i32, max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}> 119} { 120 // CHECK: spv.BitReverse 121 %0 = "test.convert_to_bit_reverse_op"(%operand): (i32) -> (i32) 122 return %0: i32 123} 124 125//===----------------------------------------------------------------------===// 126// Extension 127//===----------------------------------------------------------------------===// 128 129// CHECK-LABEL: @subgroup_ballot_suitable_extension 130func @subgroup_ballot_suitable_extension(%predicate: i1) -> vector<4xi32> attributes { 131 spv.target_env = #spv.target_env<#spv.vce<v1.4, [SubgroupBallotKHR], [SPV_KHR_shader_ballot]>, {max_compute_workgroup_invocations = 128 : i32, max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}> 132} { 133 // CHECK: spv.SubgroupBallotKHR 134 %0 = "test.convert_to_subgroup_ballot_op"(%predicate): (i1) -> (vector<4xi32>) 135 return %0: vector<4xi32> 136} 137 138// CHECK-LABEL: @subgroup_ballot_missing_extension 139func @subgroup_ballot_missing_extension(%predicate: i1) -> vector<4xi32> attributes { 140 spv.target_env = #spv.target_env<#spv.vce<v1.4, [SubgroupBallotKHR], []>, {max_compute_workgroup_invocations = 128 : i32, max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}> 141} { 142 // CHECK: test.convert_to_subgroup_ballot_op 143 %0 = "test.convert_to_subgroup_ballot_op"(%predicate): (i1) -> (vector<4xi32>) 144 return %0: vector<4xi32> 145} 146 147// CHECK-LABEL: @module_suitable_extension1 148func @module_suitable_extension1() attributes { 149 spv.target_env = #spv.target_env<#spv.vce<v1.0, [VulkanMemoryModel, PhysicalStorageBufferAddresses], [SPV_KHR_vulkan_memory_model, SPV_EXT_physical_storage_buffer]>, {max_compute_workgroup_invocations = 128 : i32, max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}> 150} { 151 // CHECK: spv.module PhysicalStorageBuffer64 Vulkan 152 "test.convert_to_module_op"() : () ->() 153 return 154} 155 156// CHECK-LABEL: @module_suitable_extension2 157func @module_suitable_extension2() attributes { 158 spv.target_env = #spv.target_env<#spv.vce<v1.0, [VulkanMemoryModel, PhysicalStorageBufferAddresses], [SPV_KHR_vulkan_memory_model, SPV_KHR_physical_storage_buffer]>, {max_compute_workgroup_invocations = 128 : i32, max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}> 159} { 160 // CHECK: spv.module PhysicalStorageBuffer64 Vulkan 161 "test.convert_to_module_op"() : () -> () 162 return 163} 164 165// CHECK-LABEL: @module_missing_extension_mm 166func @module_missing_extension_mm() attributes { 167 spv.target_env = #spv.target_env<#spv.vce<v1.0, [VulkanMemoryModel, PhysicalStorageBufferAddresses], [SPV_KHR_physical_storage_buffer]>, {max_compute_workgroup_invocations = 128 : i32, max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}> 168} { 169 // CHECK: test.convert_to_module_op 170 "test.convert_to_module_op"() : () -> () 171 return 172} 173 174// CHECK-LABEL: @module_missing_extension_am 175func @module_missing_extension_am() attributes { 176 spv.target_env = #spv.target_env<#spv.vce<v1.0, [VulkanMemoryModel, PhysicalStorageBufferAddresses], [SPV_KHR_vulkan_memory_model]>, {max_compute_workgroup_invocations = 128 : i32, max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}> 177} { 178 // CHECK: test.convert_to_module_op 179 "test.convert_to_module_op"() : () -> () 180 return 181} 182 183// CHECK-LABEL: @module_implied_extension 184func @module_implied_extension() attributes { 185 // Version 1.5 implies SPV_KHR_vulkan_memory_model and SPV_KHR_physical_storage_buffer. 186 spv.target_env = #spv.target_env<#spv.vce<v1.5, [VulkanMemoryModel, PhysicalStorageBufferAddresses], []>, {max_compute_workgroup_invocations = 128 : i32, max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}> 187} { 188 // CHECK: spv.module PhysicalStorageBuffer64 Vulkan 189 "test.convert_to_module_op"() : () -> () 190 return 191} 192