diff options
Diffstat (limited to 'llvm/test/CodeGen/SPIRV')
30 files changed, 310 insertions, 139 deletions
diff --git a/llvm/test/CodeGen/SPIRV/AtomicCompareExchange.ll b/llvm/test/CodeGen/SPIRV/AtomicCompareExchange.ll index f8207c56a565..5ce4a1954c5f 100644 --- a/llvm/test/CodeGen/SPIRV/AtomicCompareExchange.ll +++ b/llvm/test/CodeGen/SPIRV/AtomicCompareExchange.ll @@ -1,6 +1,6 @@ ; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV -; CHECK-SPIRV: %[[#Int:]] = OpTypeInt 32 0 +; CHECK-SPIRV-DAG: %[[#Int:]] = OpTypeInt 32 0 ; CHECK-SPIRV-DAG: %[[#MemScope_CrossDevice:]] = OpConstant %[[#Int]] 0 ; CHECK-SPIRV-DAG: %[[#MemSemEqual_SeqCst:]] = OpConstant %[[#Int]] 16 ; CHECK-SPIRV-DAG: %[[#MemSemUnequal_Acquire:]] = OpConstant %[[#Int]] 2 diff --git a/llvm/test/CodeGen/SPIRV/event-zero-const.ll b/llvm/test/CodeGen/SPIRV/event-zero-const.ll index f3f20a0fb2f4..523d2ad9825f 100644 --- a/llvm/test/CodeGen/SPIRV/event-zero-const.ll +++ b/llvm/test/CodeGen/SPIRV/event-zero-const.ll @@ -4,10 +4,10 @@ ; RUN: llc -verify-machineinstrs -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s ; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv32-unknown-unknown %s -o - -filetype=obj | spirv-val %} -; CHECK: %[[#LongTy:]] = OpTypeInt 64 0 -; CHECK: %[[#EventTy:]] = OpTypeEvent -; CHECK: %[[#LongNull:]] = OpConstantNull %[[#LongTy]] -; CHECK: %[[#EventNull:]] = OpConstantNull %[[#EventTy]] +; CHECK-DAG: %[[#LongTy:]] = OpTypeInt 64 0 +; CHECK-DAG: %[[#EventTy:]] = OpTypeEvent +; CHECK-DAG: %[[#LongNull:]] = OpConstantNull %[[#LongTy]] +; CHECK-DAG: %[[#EventNull:]] = OpConstantNull %[[#EventTy]] ; CHECK: OpFunction ; CHECK: OpINotEqual %[[#]] %[[#]] %[[#LongNull]] ; CHECK: OpGroupAsyncCopy %[[#EventTy]] %[[#]] %[[#]] %[[#]] %[[#]] %[[#]] %[[#EventNull]] diff --git a/llvm/test/CodeGen/SPIRV/extensions/SPV_INTEL_function_pointers/fp-simple-hierarchy.ll b/llvm/test/CodeGen/SPIRV/extensions/SPV_INTEL_function_pointers/fp-simple-hierarchy.ll index 368c5d4a3298..80309e96db00 100644 --- a/llvm/test/CodeGen/SPIRV/extensions/SPV_INTEL_function_pointers/fp-simple-hierarchy.ll +++ b/llvm/test/CodeGen/SPIRV/extensions/SPV_INTEL_function_pointers/fp-simple-hierarchy.ll @@ -1,16 +1,17 @@ -; RUN: llc -O0 -mtriple=spirv32-unknown-unknown --spirv-ext=+SPV_INTEL_function_pointers %s -o - | FileCheck %s +; RUN: llc -verify-machineinstrs -O0 -mtriple=spirv32-unknown-unknown --spirv-ext=+SPV_INTEL_function_pointers %s -o - | FileCheck %s ; TODO: %if spirv-tools %{ llc -O0 -mtriple=spirv64-unknown-unknown %s -o - -filetype=obj | spirv-val %} -; TODO: This test currently fails with LLVM_ENABLE_EXPENSIVE_CHECKS enabled -; XFAIL: expensive_checks - ; CHECK-DAG: OpName %[[I9:.*]] "_ZN13BaseIncrement9incrementEPi" ; CHECK-DAG: OpName %[[I29:.*]] "_ZN12IncrementBy29incrementEPi" ; CHECK-DAG: OpName %[[I49:.*]] "_ZN12IncrementBy49incrementEPi" ; CHECK-DAG: OpName %[[I89:.*]] "_ZN12IncrementBy89incrementEPi" +; CHECK-DAG: OpName %[[Foo:.*]] "foo" ; CHECK-DAG: %[[TyVoid:.*]] = OpTypeVoid -; CHECK-DAG: %[[TyArr:.*]] = OpTypeArray +; CHECK-DAG: %[[TyInt32:.*]] = OpTypeInt 32 0 +; CHECK-DAG: %[[TyInt8:.*]] = OpTypeInt 8 0 +; CHECK-DAG: %[[Const8:.*]] = OpConstant %[[TyInt32]] 8 +; CHECK-DAG: %[[TyArr:.*]] = OpTypeArray %[[TyInt8]] %[[Const8]] ; CHECK-DAG: %[[TyStruct1:.*]] = OpTypeStruct %[[TyArr]] ; CHECK-DAG: %[[TyStruct2:.*]] = OpTypeStruct %[[TyStruct1]] ; CHECK-DAG: %[[TyPtrStruct2:.*]] = OpTypePointer Generic %[[TyStruct2]] @@ -18,16 +19,21 @@ ; CHECK-DAG: %[[TyPtrFun:.*]] = OpTypePointer Generic %[[TyFun]] ; CHECK-DAG: %[[TyPtrPtrFun:.*]] = OpTypePointer Generic %[[TyPtrFun]] -; CHECK: %[[I9]] = OpFunction -; CHECK: %[[I29]] = OpFunction -; CHECK: %[[I49]] = OpFunction -; CHECK: %[[I89]] = OpFunction +; CHECK-DAG: %[[I9]] = OpFunction +; CHECK-DAG: %[[I29]] = OpFunction +; CHECK-DAG: %[[I49]] = OpFunction +; CHECK-DAG: %[[I89]] = OpFunction + +; CHECK: %[[Foo]] = OpFunction +; CHECK-4: OpFunctionParameter ; CHECK: %[[Arg1:.*]] = OpPhi %[[TyPtrStruct2]] ; CHECK: %[[VTbl:.*]] = OpBitcast %[[TyPtrPtrFun]] %[[#]] ; CHECK: %[[FP:.*]] = OpLoad %[[TyPtrFun]] %[[VTbl]] ; CHECK: %[[#]] = OpFunctionPointerCallINTEL %[[TyVoid]] %[[FP]] %[[Arg1]] %[[#]] +; CHECK-NO: OpFunction + %"cls::id" = type { %"cls::detail::array" } %"cls::detail::array" = type { [1 x i64] } %struct.obj_storage_t = type { %"struct.aligned_storage<BaseIncrement, IncrementBy2, IncrementBy4, IncrementBy8>::type" } diff --git a/llvm/test/CodeGen/SPIRV/extensions/SPV_INTEL_function_pointers/fp_const.ll b/llvm/test/CodeGen/SPIRV/extensions/SPV_INTEL_function_pointers/fp_const.ll index 75ad382f05ff..b96da631c0a8 100644 --- a/llvm/test/CodeGen/SPIRV/extensions/SPV_INTEL_function_pointers/fp_const.ll +++ b/llvm/test/CodeGen/SPIRV/extensions/SPV_INTEL_function_pointers/fp_const.ll @@ -1,9 +1,6 @@ -; RUN: llc -O0 -mtriple=spirv32-unknown-unknown --spirv-ext=+SPV_INTEL_function_pointers %s -o - | FileCheck %s +; RUN: llc -verify-machineinstrs -O0 -mtriple=spirv32-unknown-unknown --spirv-ext=+SPV_INTEL_function_pointers %s -o - | FileCheck %s ; TODO: %if spirv-tools %{ llc -O0 -mtriple=spirv64-unknown-unknown %s -o - -filetype=obj | spirv-val %} -; TODO: This test currently fails with LLVM_ENABLE_EXPENSIVE_CHECKS enabled -; XFAIL: expensive_checks - ; CHECK-DAG: OpCapability FunctionPointersINTEL ; CHECK-DAG: OpCapability Int64 ; CHECK: OpExtension "SPV_INTEL_function_pointers" diff --git a/llvm/test/CodeGen/SPIRV/extensions/SPV_INTEL_function_pointers/fun-ptr-addrcast.ll b/llvm/test/CodeGen/SPIRV/extensions/SPV_INTEL_function_pointers/fun-ptr-addrcast.ll index d38de216c223..8edecc1329d0 100644 --- a/llvm/test/CodeGen/SPIRV/extensions/SPV_INTEL_function_pointers/fun-ptr-addrcast.ll +++ b/llvm/test/CodeGen/SPIRV/extensions/SPV_INTEL_function_pointers/fun-ptr-addrcast.ll @@ -2,15 +2,9 @@ ; pointers/PtrCast-null-in-OpSpecConstantOp.ll (that is OpSpecConstantOp with ptr-cast operation) correctly ; work also for function pointers. -; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - --spirv-ext=+SPV_INTEL_function_pointers | FileCheck %s +; RUN: llc -verify-machineinstrs -O0 -mtriple=spirv32-unknown-unknown %s -o - --spirv-ext=+SPV_INTEL_function_pointers | FileCheck %s ; TODO: %if spirv-tools %{ llc -O0 -mtriple=spirv64-unknown-unknown %s -o - -filetype=obj | spirv-val %} -; TODO: This test currently fails with LLVM_ENABLE_EXPENSIVE_CHECKS enabled -; XFAIL: expensive_checks - -; Running with -verify-machineinstrs would lead to "Reading virtual register without a def" -; error, because OpConstantFunctionPointerINTEL forward-refers to a function definition. - ; CHECK-COUNT-3: %[[#]] = OpSpecConstantOp %[[#]] 121 %[[#]] ; CHECK-COUNT-3: OpPtrCastToGeneric diff --git a/llvm/test/CodeGen/SPIRV/extensions/SPV_INTEL_inline_assembly/inline_asm.ll b/llvm/test/CodeGen/SPIRV/extensions/SPV_INTEL_inline_assembly/inline_asm.ll index e006651d49e4..91286d5bf32e 100644 --- a/llvm/test/CodeGen/SPIRV/extensions/SPV_INTEL_inline_assembly/inline_asm.ll +++ b/llvm/test/CodeGen/SPIRV/extensions/SPV_INTEL_inline_assembly/inline_asm.ll @@ -31,20 +31,20 @@ ; CHECK-DAG: %[[#Const123:]] = OpConstant %[[#Int32Ty]] 123 ; CHECK-DAG: %[[#Const42:]] = OpConstant %[[#DoubleTy:]] 42 -; CHECK: %[[#Dialect:]] = OpAsmTargetINTEL "spirv64-unknown-unknown" +; CHECK-DAG: %[[#Dialect:]] = OpAsmTargetINTEL "spirv64-unknown-unknown" ; CHECK-NO: OpAsmTargetINTEL -; CHECK: %[[#Asm1:]] = OpAsmINTEL %[[#VoidTy]] %[[#Fun1Ty]] %[[#Dialect]] "" "" -; CHECK: %[[#Asm2:]] = OpAsmINTEL %[[#VoidTy]] %[[#Fun1Ty]] %[[#Dialect]] "nop" "" -; CHECK: %[[#Asm3:]] = OpAsmINTEL %[[#VoidTy]] %[[#Fun1Ty]] %[[#Dialect]] "" "~{cc},~{memory}" -; CHECK: %[[#Asm4:]] = OpAsmINTEL %[[#Int32Ty]] %[[#Fun2Ty:]] %[[#Dialect]] "clobber_out $0" "=&r" -; CHECK: %[[#Asm5:]] = OpAsmINTEL %[[#Int32Ty]] %[[#Fun3Ty]] %[[#Dialect]] "icmd $0 $1" "=r,r" -; CHECK: %[[#Asm6:]] = OpAsmINTEL %[[#FloatTy]] %[[#Fun4Ty]] %[[#Dialect]] "fcmd $0 $1" "=r,r" -; CHECK: %[[#Asm7:]] = OpAsmINTEL %[[#HalfTy]] %[[#Fun5Ty]] %[[#Dialect]] "fcmdext $0 $1 $2" "=r,r,r" -; CHECK: %[[#Asm8:]] = OpAsmINTEL %[[#Int8Ty]] %[[#Fun6Ty]] %[[#Dialect]] "cmdext $0 $3 $1 $2" "=r,r,r,r" -; CHECK: %[[#Asm9:]] = OpAsmINTEL %[[#Int64Ty]] %[[#Fun7Ty]] %[[#Dialect]] "icmdext $0 $3 $1 $2" "=r,r,r,r" -; CHECK: %[[#Asm10:]] = OpAsmINTEL %[[#VoidTy]] %[[#Fun8Ty]] %[[#Dialect]] "constcmd $0 $1" "r,r" -; CHECK: %[[#Asm11:]] = OpAsmINTEL %[[#VoidTy]] %[[#Fun8Ty]] %[[#Dialect]] "constcmd $0 $1" "i,i" +; CHECK-DAG: %[[#Asm1:]] = OpAsmINTEL %[[#VoidTy]] %[[#Fun1Ty]] %[[#Dialect]] "" "" +; CHECK-DAG: %[[#Asm2:]] = OpAsmINTEL %[[#VoidTy]] %[[#Fun1Ty]] %[[#Dialect]] "nop" "" +; CHECK-DAG: %[[#Asm3:]] = OpAsmINTEL %[[#VoidTy]] %[[#Fun1Ty]] %[[#Dialect]] "" "~{cc},~{memory}" +; CHECK-DAG: %[[#Asm4:]] = OpAsmINTEL %[[#Int32Ty]] %[[#Fun2Ty:]] %[[#Dialect]] "clobber_out $0" "=&r" +; CHECK-DAG: %[[#Asm5:]] = OpAsmINTEL %[[#Int32Ty]] %[[#Fun3Ty]] %[[#Dialect]] "icmd $0 $1" "=r,r" +; CHECK-DAG: %[[#Asm6:]] = OpAsmINTEL %[[#FloatTy]] %[[#Fun4Ty]] %[[#Dialect]] "fcmd $0 $1" "=r,r" +; CHECK-DAG: %[[#Asm7:]] = OpAsmINTEL %[[#HalfTy]] %[[#Fun5Ty]] %[[#Dialect]] "fcmdext $0 $1 $2" "=r,r,r" +; CHECK-DAG: %[[#Asm8:]] = OpAsmINTEL %[[#Int8Ty]] %[[#Fun6Ty]] %[[#Dialect]] "cmdext $0 $3 $1 $2" "=r,r,r,r" +; CHECK-DAG: %[[#Asm9:]] = OpAsmINTEL %[[#Int64Ty]] %[[#Fun7Ty]] %[[#Dialect]] "icmdext $0 $3 $1 $2" "=r,r,r,r" +; CHECK-DAG: %[[#Asm10:]] = OpAsmINTEL %[[#VoidTy]] %[[#Fun8Ty]] %[[#Dialect]] "constcmd $0 $1" "r,r" +; CHECK-DAG: %[[#Asm11:]] = OpAsmINTEL %[[#VoidTy]] %[[#Fun8Ty]] %[[#Dialect]] "constcmd $0 $1" "i,i" ; CHECK-NO: OpAsmINTEL ; CHECK: OpFunction diff --git a/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_shader_clock/shader_clock.ll b/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_shader_clock/shader_clock.ll index 8ecd0a2b25eb..bd07ba1316ec 100644 --- a/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_shader_clock/shader_clock.ll +++ b/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_shader_clock/shader_clock.ll @@ -6,9 +6,11 @@ ; CHECK: OpCapability ShaderClockKHR ; CHECK: OpExtension "SPV_KHR_shader_clock" -; CHECK-DAG: [[uint:%[a-z0-9_]+]] = OpTypeInt 32 +; CHECK-DAG: [[uint:%[a-z0-9_]+]] = OpTypeInt 32 0 ; CHECK-DAG: [[ulong:%[a-z0-9_]+]] = OpTypeInt 64 ; CHECK-DAG: [[v2uint:%[a-z0-9_]+]] = OpTypeVector [[uint]] 2 +; CHECK-DAG: OpConstant [[uint]] 8 +; CHECK-DAG: OpConstant [[uint]] 16 ; CHECK-DAG: [[uint_1:%[a-z0-9_]+]] = OpConstant [[uint]] 1 ; CHECK-DAG: [[uint_2:%[a-z0-9_]+]] = OpConstant [[uint]] 2 ; CHECK-DAG: [[uint_3:%[a-z0-9_]+]] = OpConstant [[uint]] 3 diff --git a/llvm/test/CodeGen/SPIRV/global-var-name-linkage.ll b/llvm/test/CodeGen/SPIRV/global-var-name-linkage.ll new file mode 100644 index 000000000000..4501819ce494 --- /dev/null +++ b/llvm/test/CodeGen/SPIRV/global-var-name-linkage.ll @@ -0,0 +1,59 @@ +; Check names and decoration of global variables. + +; RUN: llc -verify-machineinstrs -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s +; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv64-unknown-unknown %s -o - -filetype=obj | spirv-val %} + +; RUN: llc -verify-machineinstrs -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s +; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv32-unknown-unknown %s -o - -filetype=obj | spirv-val %} + +; CHECK-DAG: OpName %[[#id18:]] "G1" +; CHECK-DAG: OpName %[[#id22:]] "g1" +; CHECK-DAG: OpName %[[#id23:]] "g2" +; CHECK-DAG: OpName %[[#id27:]] "g4" +; CHECK-DAG: OpName %[[#id30:]] "c1" +; CHECK-DAG: OpName %[[#id31:]] "n_t" +; CHECK-DAG: OpName %[[#id32:]] "w" +; CHECK-DAG: OpName %[[#id34:]] "a.b" +; CHECK-DAG: OpName %[[#id35:]] "e" +; CHECK-DAG: OpName %[[#id36:]] "y.z" +; CHECK-DAG: OpName %[[#id38:]] "x" + +; CHECK-NOT: OpDecorate %[[#id18]] LinkageAttributes +; CHECK-DAG: OpDecorate %[[#id18]] Constant +; CHECK-DAG: OpDecorate %[[#id22]] Alignment 4 +; CHECK-DAG: OpDecorate %[[#id22]] LinkageAttributes "g1" Export +; CHECK-DAG: OpDecorate %[[#id23]] Alignment 4 +; CHECK-DAG: OpDecorate %[[#id27]] Alignment 4 +; CHECK-DAG: OpDecorate %[[#id27]] LinkageAttributes "g4" Export +; CHECK-DAG: OpDecorate %[[#id30]] Constant +; CHECK-DAG: OpDecorate %[[#id30]] Alignment 4 +; CHECK-DAG: OpDecorate %[[#id30]] LinkageAttributes "c1" Export +; CHECK-DAG: OpDecorate %[[#id31]] Constant +; CHECK-DAG: OpDecorate %[[#id31]] LinkageAttributes "n_t" Import +; CHECK-DAG: OpDecorate %[[#id32]] Constant +; CHECK-DAG: OpDecorate %[[#id32]] Alignment 4 +; CHECK-DAG: OpDecorate %[[#id32]] LinkageAttributes "w" Export +; CHECK-DAG: OpDecorate %[[#id34]] Constant +; CHECK-DAG: OpDecorate %[[#id34]] Alignment 4 +; CHECK-DAG: OpDecorate %[[#id35]] LinkageAttributes "e" Import +; CHECK-DAG: OpDecorate %[[#id36]] Alignment 4 +; CHECK-DAG: OpDecorate %[[#id38]] Constant +; CHECK-DAG: OpDecorate %[[#id38]] Alignment 4 + +%"class.sycl::_V1::nd_item" = type { i8 } + +@G1 = private unnamed_addr addrspace(1) constant %"class.sycl::_V1::nd_item" poison, align 1 +@g1 = addrspace(1) global i32 1, align 4 +@g2 = internal addrspace(1) global i32 2, align 4 +@g4 = common addrspace(1) global i32 0, align 4 +@c1 = addrspace(2) constant [2 x i32] [i32 0, i32 1], align 4 +@n_t = external addrspace(2) constant [256 x i32] +@w = addrspace(1) constant i32 0, align 4 +@a.b = internal addrspace(2) constant [2 x i32] [i32 2, i32 3], align 4 +@e = external addrspace(1) global i32 +@y.z = internal addrspace(1) global i32 0, align 4 +@x = internal addrspace(2) constant float 1.000000e+00, align 4 + +define internal spir_func void @foo(ptr addrspace(4) align 1 %arg) { + ret void +} diff --git a/llvm/test/CodeGen/SPIRV/hlsl-intrinsics/SV_GroupID.ll b/llvm/test/CodeGen/SPIRV/hlsl-intrinsics/SV_GroupID.ll new file mode 100644 index 000000000000..92947f7865ce --- /dev/null +++ b/llvm/test/CodeGen/SPIRV/hlsl-intrinsics/SV_GroupID.ll @@ -0,0 +1,52 @@ +; RUN: llc -O0 -verify-machineinstrs -mtriple=spirv-vulkan-unknown %s -o - | FileCheck %s +; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv-vulkan-unknown %s -o - -filetype=obj | spirv-val %} + +; CHECK-DAG: %[[#int:]] = OpTypeInt 32 0 +; CHECK-DAG: %[[#v3int:]] = OpTypeVector %[[#int]] 3 +; CHECK-DAG: %[[#ptr_Input_v3int:]] = OpTypePointer Input %[[#v3int]] +; CHECK-DAG: %[[#tempvar:]] = OpUndef %[[#v3int]] +; CHECK-DAG: %[[#WorkgroupId:]] = OpVariable %[[#ptr_Input_v3int]] Input + +; CHECK-DAG: OpEntryPoint GLCompute {{.*}} %[[#WorkgroupId]] +; CHECK-DAG: OpName %[[#WorkgroupId]] "__spirv_BuiltInWorkgroupId" +; CHECK-DAG: OpDecorate %[[#WorkgroupId]] LinkageAttributes "__spirv_BuiltInWorkgroupId" Import +; CHECK-DAG: OpDecorate %[[#WorkgroupId]] BuiltIn WorkgroupId + +target triple = "spirv-unknown-vulkan-library" + +declare void @group_id_user(<3 x i32>) + +; Function Attrs: convergent noinline norecurse +define void @main() #1 { +entry: + +; CHECK: %[[#load:]] = OpLoad %[[#v3int]] %[[#WorkgroupId]] +; CHECK: %[[#load0:]] = OpCompositeExtract %[[#int]] %[[#load]] 0 + %1 = call i32 @llvm.spv.group.id(i32 0) + +; CHECK: %[[#tempvar:]] = OpCompositeInsert %[[#v3int]] %[[#load0]] %[[#tempvar]] + %2 = insertelement <3 x i32> poison, i32 %1, i64 0 + +; CHECK: %[[#load:]] = OpLoad %[[#v3int]] %[[#WorkgroupId]] +; CHECK: %[[#load1:]] = OpCompositeExtract %[[#int]] %[[#load]] 1 + %3 = call i32 @llvm.spv.group.id(i32 1) + +; CHECK: %[[#tempvar:]] = OpCompositeInsert %[[#v3int]] %[[#load1]] %[[#tempvar]] 1 + %4 = insertelement <3 x i32> %2, i32 %3, i64 1 + +; CHECK: %[[#load:]] = OpLoad %[[#v3int]] %[[#WorkgroupId]] +; CHECK: %[[#load2:]] = OpCompositeExtract %[[#int]] %[[#load]] 2 + %5 = call i32 @llvm.spv.group.id(i32 2) + +; CHECK: %[[#tempvar:]] = OpCompositeInsert %[[#v3int]] %[[#load2]] %[[#tempvar]] 2 + %6 = insertelement <3 x i32> %4, i32 %5, i64 2 + + call spir_func void @group_id_user(<3 x i32> %6) + ret void +} + +; Function Attrs: nounwind willreturn memory(none) +declare i32 @llvm.spv.group.id(i32) #3 + +attributes #1 = { convergent noinline norecurse "hlsl.numthreads"="1,1,1" "hlsl.shader"="compute" "no-trapping-math"="true" "stack-protector-buffer-size"="8" } +attributes #3 = { nounwind willreturn memory(none) } diff --git a/llvm/test/CodeGen/SPIRV/hlsl-intrinsics/cross.ll b/llvm/test/CodeGen/SPIRV/hlsl-intrinsics/cross.ll index 2e0eb8c429ac..b1625c07111e 100644 --- a/llvm/test/CodeGen/SPIRV/hlsl-intrinsics/cross.ll +++ b/llvm/test/CodeGen/SPIRV/hlsl-intrinsics/cross.ll @@ -15,7 +15,7 @@ entry: ; CHECK: %[[#arg0:]] = OpFunctionParameter %[[#vec3_float_16]]
; CHECK: %[[#arg1:]] = OpFunctionParameter %[[#vec3_float_16]]
; CHECK: %[[#]] = OpExtInst %[[#vec3_float_16]] %[[#op_ext_glsl]] Cross %[[#arg0]] %[[#arg1]]
- %hlsl.cross = call <3 x half> @llvm.spv.cross.v4f16(<3 x half> %a, <3 x half> %b)
+ %hlsl.cross = call <3 x half> @llvm.spv.cross.v3f16(<3 x half> %a, <3 x half> %b)
ret <3 x half> %hlsl.cross
}
@@ -25,9 +25,9 @@ entry: ; CHECK: %[[#arg0:]] = OpFunctionParameter %[[#vec3_float_32]]
; CHECK: %[[#arg1:]] = OpFunctionParameter %[[#vec3_float_32]]
; CHECK: %[[#]] = OpExtInst %[[#vec3_float_32]] %[[#op_ext_glsl]] Cross %[[#arg0]] %[[#arg1]]
- %hlsl.cross = call <3 x float> @llvm.spv.cross.v4f32(<3 x float> %a, <3 x float> %b)
+ %hlsl.cross = call <3 x float> @llvm.spv.cross.v3f32(<3 x float> %a, <3 x float> %b)
ret <3 x float> %hlsl.cross
}
-declare <3 x half> @llvm.spv.cross.v4f16(<3 x half>, <3 x half>)
-declare <3 x float> @llvm.spv.cross.v4f32(<3 x float>, <3 x float>)
+declare <3 x half> @llvm.spv.cross.v3f16(<3 x half>, <3 x half>)
+declare <3 x float> @llvm.spv.cross.v3f32(<3 x float>, <3 x float>)
diff --git a/llvm/test/CodeGen/SPIRV/hlsl-intrinsics/distance.ll b/llvm/test/CodeGen/SPIRV/hlsl-intrinsics/distance.ll new file mode 100644 index 000000000000..85a24a0127ae --- /dev/null +++ b/llvm/test/CodeGen/SPIRV/hlsl-intrinsics/distance.ll @@ -0,0 +1,33 @@ +; RUN: llc -O0 -mtriple=spirv-unknown-unknown %s -o - | FileCheck %s
+; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv-unknown-unknown %s -o - -filetype=obj | spirv-val %}
+
+; Make sure SPIRV operation function calls for distance are lowered correctly.
+
+; CHECK-DAG: %[[#op_ext_glsl:]] = OpExtInstImport "GLSL.std.450"
+; CHECK-DAG: %[[#float_16:]] = OpTypeFloat 16
+; CHECK-DAG: %[[#vec4_float_16:]] = OpTypeVector %[[#float_16]] 4
+; CHECK-DAG: %[[#float_32:]] = OpTypeFloat 32
+; CHECK-DAG: %[[#vec4_float_32:]] = OpTypeVector %[[#float_32]] 4
+
+define noundef half @distance_half4(<4 x half> noundef %a, <4 x half> noundef %b) {
+entry:
+ ; CHECK: %[[#]] = OpFunction %[[#float_16]] None %[[#]]
+ ; CHECK: %[[#arg0:]] = OpFunctionParameter %[[#vec4_float_16]]
+ ; CHECK: %[[#arg1:]] = OpFunctionParameter %[[#vec4_float_16]]
+ ; CHECK: %[[#]] = OpExtInst %[[#float_16]] %[[#op_ext_glsl]] Distance %[[#arg0]] %[[#arg1]]
+ %spv.distance = call half @llvm.spv.distance.f16(<4 x half> %a, <4 x half> %b)
+ ret half %spv.distance
+}
+
+define noundef float @distance_float4(<4 x float> noundef %a, <4 x float> noundef %b) {
+entry:
+ ; CHECK: %[[#]] = OpFunction %[[#float_32]] None %[[#]]
+ ; CHECK: %[[#arg0:]] = OpFunctionParameter %[[#vec4_float_32]]
+ ; CHECK: %[[#arg1:]] = OpFunctionParameter %[[#vec4_float_32]]
+ ; CHECK: %[[#]] = OpExtInst %[[#float_32]] %[[#op_ext_glsl]] Distance %[[#arg0]] %[[#arg1]]
+ %spv.distance = call float @llvm.spv.distance.f32(<4 x float> %a, <4 x float> %b)
+ ret float %spv.distance
+}
+
+declare half @llvm.spv.distance.f16(<4 x half>, <4 x half>)
+declare float @llvm.spv.distance.f32(<4 x float>, <4 x float>)
diff --git a/llvm/test/CodeGen/SPIRV/hlsl-intrinsics/length.ll b/llvm/test/CodeGen/SPIRV/hlsl-intrinsics/length.ll index b4a9d8e0664b..1ac862b79a3f 100644 --- a/llvm/test/CodeGen/SPIRV/hlsl-intrinsics/length.ll +++ b/llvm/test/CodeGen/SPIRV/hlsl-intrinsics/length.ll @@ -11,19 +11,21 @@ define noundef half @length_half4(<4 x half> noundef %a) {
entry:
- ; CHECK: %[[#arg0:]] = OpFunctionParameter %[[#]]
+ ; CHECK: %[[#]] = OpFunction %[[#float_16]] None %[[#]]
+ ; CHECK: %[[#arg0:]] = OpFunctionParameter %[[#vec4_float_16]]
; CHECK: %[[#]] = OpExtInst %[[#float_16]] %[[#op_ext_glsl]] Length %[[#arg0]]
- %hlsl.length = call half @llvm.spv.length.v4f16(<4 x half> %a)
+ %hlsl.length = call half @llvm.spv.length.f16(<4 x half> %a)
ret half %hlsl.length
}
define noundef float @length_float4(<4 x float> noundef %a) {
entry:
- ; CHECK: %[[#arg0:]] = OpFunctionParameter %[[#]]
+ ; CHECK: %[[#]] = OpFunction %[[#float_32]] None %[[#]]
+ ; CHECK: %[[#arg0:]] = OpFunctionParameter %[[#vec4_float_32]]
; CHECK: %[[#]] = OpExtInst %[[#float_32]] %[[#op_ext_glsl]] Length %[[#arg0]]
- %hlsl.length = call float @llvm.spv.length.v4f32(<4 x float> %a)
+ %hlsl.length = call float @llvm.spv.length.f32(<4 x float> %a)
ret float %hlsl.length
}
-declare half @llvm.spv.length.v4f16(<4 x half>)
-declare float @llvm.spv.length.v4f32(<4 x float>)
+declare half @llvm.spv.length.f16(<4 x half>)
+declare float @llvm.spv.length.f32(<4 x float>)
diff --git a/llvm/test/CodeGen/SPIRV/iaddcarry-builtin.ll b/llvm/test/CodeGen/SPIRV/iaddcarry-builtin.ll index 8f14eba21b63..49aaa45afed1 100644 --- a/llvm/test/CodeGen/SPIRV/iaddcarry-builtin.ll +++ b/llvm/test/CodeGen/SPIRV/iaddcarry-builtin.ll @@ -25,9 +25,7 @@ ; CHECK-SPIRV-DAG: [[v4uint:%[a-z0-9_]+]] = OpTypeVector [[uint]] 4 ; CHECK-SPIRV-DAG: [[vecstruct:%[a-z0-9_]+]] = OpTypeStruct [[v4uint]] [[v4uint]] ; CHECK-SPIRV-DAG: [[_ptr_Function_vecstruct:%[a-z0-9_]+]] = OpTypePointer Function [[vecstruct]] -; CHECK-SPIRV-DAG: [[struct_anon:%[a-z0-9_.]+]] = OpTypeStruct [[uint]] [[uint]] -; CHECK-SPIRV-DAG: [[_ptr_Function_struct_anon:%[a-z0-9_]+]] = OpTypePointer Function [[struct_anon]] -; CHECK-SPIRV-DAG: [[_ptr_Generic_struct_anon:%[a-z0-9_]+]] = OpTypePointer Generic [[struct_anon]] +; CHECK-SPIRV-DAG: [[_ptr_Generic_i32struct:%[a-z0-9_]+]] = OpTypePointer Generic [[i32struct]] define spir_func void @test_builtin_iaddcarrycc(i8 %a, i8 %b) { entry: @@ -116,9 +114,9 @@ define spir_func void @test_builtin_iaddcarry_anon(i32 %a, i32 %b) { ; CHECK-SPIRV: [[a_4:%[a-z0-9_]+]] = OpFunctionParameter [[uint]] ; CHECK-SPIRV: [[b_4:%[a-z0-9_]+]] = OpFunctionParameter [[uint]] ; CHECK-SPIRV: [[entry_4:%[a-z0-9_]+]] = OpLabel -; CHECK-SPIRV: [[var_59:%[a-z0-9_]+]] = OpVariable [[_ptr_Function_struct_anon]] Function -; CHECK-SPIRV: [[var_61:%[a-z0-9_]+]] = OpPtrCastToGeneric [[_ptr_Generic_struct_anon]] [[var_59]] -; CHECK-SPIRV: [[var_62:%[a-z0-9_]+]] = OpIAddCarry [[struct_anon]] [[a_4]] [[b_4]] +; CHECK-SPIRV: [[var_59:%[a-z0-9_]+]] = OpVariable [[_ptr_Function_i32struct]] Function +; CHECK-SPIRV: [[var_61:%[a-z0-9_]+]] = OpPtrCastToGeneric [[_ptr_Generic_i32struct]] [[var_59]] +; CHECK-SPIRV: [[var_62:%[a-z0-9_]+]] = OpIAddCarry [[i32struct]] [[a_4]] [[b_4]] ; CHECK-SPIRV: OpStore [[var_61]] [[var_62]] declare void @_Z17__spirv_IAddCarryIiiE4anonIT_T0_ES1_S2_(ptr addrspace(4) sret(%struct.anon) align 4, i32, i32) diff --git a/llvm/test/CodeGen/SPIRV/image-unoptimized.ll b/llvm/test/CodeGen/SPIRV/image-unoptimized.ll index 0ce9c73ab103..d7d5b1d6b756 100644 --- a/llvm/test/CodeGen/SPIRV/image-unoptimized.ll +++ b/llvm/test/CodeGen/SPIRV/image-unoptimized.ll @@ -1,7 +1,7 @@ ; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s -; CHECK: %[[#TypeImage:]] = OpTypeImage -; CHECK: %[[#TypeSampler:]] = OpTypeSampler +; CHECK-DAG: %[[#TypeImage:]] = OpTypeImage +; CHECK-DAG: %[[#TypeSampler:]] = OpTypeSampler ; CHECK-DAG: %[[#TypeImagePtr:]] = OpTypePointer {{.*}} %[[#TypeImage]] ; CHECK-DAG: %[[#TypeSamplerPtr:]] = OpTypePointer {{.*}} %[[#TypeSampler]] diff --git a/llvm/test/CodeGen/SPIRV/isubborrow-builtin.ll b/llvm/test/CodeGen/SPIRV/isubborrow-builtin.ll index 08b4d2a1fa8e..ca842d2f9557 100644 --- a/llvm/test/CodeGen/SPIRV/isubborrow-builtin.ll +++ b/llvm/test/CodeGen/SPIRV/isubborrow-builtin.ll @@ -23,9 +23,7 @@ ; CHECK-SPIRV-DAG: [[v4uint:%[a-z0-9_]+]] = OpTypeVector [[uint]] 4 ; CHECK-SPIRV-DAG: [[vecstruct:%[a-z0-9_]+]] = OpTypeStruct [[v4uint]] [[v4uint]] ; CHECK-SPIRV-DAG: [[_ptr_Function_vecstruct:%[a-z0-9_]+]] = OpTypePointer Function [[vecstruct]] -; CHECK-SPIRV-DAG: [[struct_anon:%[a-z0-9_.]+]] = OpTypeStruct [[uint]] [[uint]] -; CHECK-SPIRV-DAG: [[_ptr_Function_struct_anon:%[a-z0-9_]+]] = OpTypePointer Function [[struct_anon]] -; CHECK-SPIRV-DAG: [[_ptr_Generic_struct_anon:%[a-z0-9_]+]] = OpTypePointer Generic [[struct_anon]] +; CHECK-SPIRV-DAG: [[_ptr_Generic_i32struct:%[a-z0-9_]+]] = OpTypePointer Generic [[i32struct]] define spir_func void @test_builtin_isubborrowcc(i8 %a, i8 %b) { entry: @@ -114,9 +112,9 @@ define spir_func void @test_builtin_isubborrow_anon(i32 %a, i32 %b) { ; CHECK-SPIRV: [[a_4:%[a-z0-9_]+]] = OpFunctionParameter [[uint]] ; CHECK-SPIRV: [[b_4:%[a-z0-9_]+]] = OpFunctionParameter [[uint]] ; CHECK-SPIRV: [[entry_4:%[a-z0-9_]+]] = OpLabel -; CHECK-SPIRV: [[var_59:%[a-z0-9_]+]] = OpVariable [[_ptr_Function_struct_anon]] Function -; CHECK-SPIRV: [[var_61:%[a-z0-9_]+]] = OpPtrCastToGeneric [[_ptr_Generic_struct_anon]] [[var_59]] -; CHECK-SPIRV: [[var_62:%[a-z0-9_]+]] = OpISubBorrow [[struct_anon]] [[a_4]] [[b_4]] +; CHECK-SPIRV: [[var_59:%[a-z0-9_]+]] = OpVariable [[_ptr_Function_i32struct]] Function +; CHECK-SPIRV: [[var_61:%[a-z0-9_]+]] = OpPtrCastToGeneric [[_ptr_Generic_i32struct]] [[var_59]] +; CHECK-SPIRV: [[var_62:%[a-z0-9_]+]] = OpISubBorrow [[i32struct]] [[a_4]] [[b_4]] ; CHECK-SPIRV: OpStore [[var_61]] [[var_62]] declare void @_Z18__spirv_ISubBorrowIiiE4anonIT_T0_ES1_S2_(ptr addrspace(4) sret(%struct.anon) align 4, i32, i32) diff --git a/llvm/test/CodeGen/SPIRV/keep-tracked-const.ll b/llvm/test/CodeGen/SPIRV/keep-tracked-const.ll index 0dc86233f8e1..61d06fe2752e 100644 --- a/llvm/test/CodeGen/SPIRV/keep-tracked-const.ll +++ b/llvm/test/CodeGen/SPIRV/keep-tracked-const.ll @@ -3,9 +3,9 @@ ; RUN: llc -verify-machineinstrs -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV ; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv64-unknown-unknown %s -o - -filetype=obj | spirv-val %} -; CHECK-SPIRV: %[[#Int:]] = OpTypeInt 32 0 -; CHECK-SPIRV: %[[#C0:]] = OpConstant %[[#Int]] 0 -; CHECK-SPIRV: %[[#C1:]] = OpConstant %[[#Int]] 1 +; CHECK-SPIRV-DAG: %[[#Int:]] = OpTypeInt 32 0 +; CHECK-SPIRV-DAG: %[[#C0:]] = OpConstant %[[#Int]] 0 +; CHECK-SPIRV-DAG: %[[#C1:]] = OpConstant %[[#Int]] 1 ; CHECK-SPIRV: OpSelect %[[#Int]] %[[#]] %[[#C1]] %[[#C0]] diff --git a/llvm/test/CodeGen/SPIRV/llvm-intrinsics/fshl.ll b/llvm/test/CodeGen/SPIRV/llvm-intrinsics/fshl.ll index 2d5b30978aa2..25b530461012 100644 --- a/llvm/test/CodeGen/SPIRV/llvm-intrinsics/fshl.ll +++ b/llvm/test/CodeGen/SPIRV/llvm-intrinsics/fshl.ll @@ -1,21 +1,21 @@ ; RUN: llc -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV -; CHECK-SPIRV: OpName %[[#NAME_FSHL_FUNC_32:]] "spirv.llvm_fshl_i32" -; CHECK-SPIRV: OpName %[[#NAME_FSHL_FUNC_16:]] "spirv.llvm_fshl_i16" -; CHECK-SPIRV: OpName %[[#NAME_FSHL_FUNC_VEC_INT_16:]] "spirv.llvm_fshl_v2i16" -; CHECK-SPIRV: %[[#TYPE_INT_32:]] = OpTypeInt 32 0 -; CHECK-SPIRV: %[[#TYPE_ORIG_FUNC_32:]] = OpTypeFunction %[[#TYPE_INT_32]] %[[#TYPE_INT_32]] %[[#TYPE_INT_32]] -; CHECK-SPIRV: %[[#TYPE_INT_16:]] = OpTypeInt 16 0 -; CHECK-SPIRV: %[[#TYPE_ORIG_FUNC_16:]] = OpTypeFunction %[[#TYPE_INT_16]] %[[#TYPE_INT_16]] %[[#TYPE_INT_16]] -; CHECK-SPIRV: %[[#TYPE_VEC_INT_16:]] = OpTypeVector %[[#TYPE_INT_16]] 2 -; CHECK-SPIRV: %[[#TYPE_ORIG_FUNC_VEC_INT_16:]] = OpTypeFunction %[[#TYPE_VEC_INT_16]] %[[#TYPE_VEC_INT_16]] %[[#TYPE_VEC_INT_16]] -; CHECK-SPIRV: %[[#TYPE_FSHL_FUNC_32:]] = OpTypeFunction %[[#TYPE_INT_32]] %[[#TYPE_INT_32]] %[[#TYPE_INT_32]] %[[#TYPE_INT_32]] -; CHECK-SPIRV: %[[#TYPE_FSHL_FUNC_16:]] = OpTypeFunction %[[#TYPE_INT_16]] %[[#TYPE_INT_16]] %[[#TYPE_INT_16]] %[[#TYPE_INT_16]] -; CHECK-SPIRV: %[[#TYPE_FSHL_FUNC_VEC_INT_16:]] = OpTypeFunction %[[#TYPE_VEC_INT_16]] %[[#TYPE_VEC_INT_16]] %[[#TYPE_VEC_INT_16]] %[[#TYPE_VEC_INT_16]] -; CHECK-SPIRV-DAG: %[[#CONST_ROTATE_32:]] = OpConstant %[[#TYPE_INT_32]] 8 -; CHECK-SPIRV-DAG: %[[#CONST_ROTATE_16:]] = OpConstant %[[#TYPE_INT_16]] 8 -; CHECK-SPIRV: %[[#CONST_ROTATE_VEC_INT_16:]] = OpConstantComposite %[[#TYPE_VEC_INT_16]] %[[#CONST_ROTATE_16]] %[[#CONST_ROTATE_16]] -; CHECK-SPIRV-DAG: %[[#CONST_TYPE_SIZE_32:]] = OpConstant %[[#TYPE_INT_32]] 32 +; CHECK-SPIRV-DAG: OpName %[[#NAME_FSHL_FUNC_32:]] "spirv.llvm_fshl_i32" +; CHECK-SPIRV-DAG: OpName %[[#NAME_FSHL_FUNC_16:]] "spirv.llvm_fshl_i16" +; CHECK-SPIRV-DAG: OpName %[[#NAME_FSHL_FUNC_VEC_INT_16:]] "spirv.llvm_fshl_v2i16" +; CHECK-SPIRV-DAG: %[[#TYPE_INT_32:]] = OpTypeInt 32 0 +; CHECK-SPIRV-DAG: %[[#TYPE_ORIG_FUNC_32:]] = OpTypeFunction %[[#TYPE_INT_32]] %[[#TYPE_INT_32]] %[[#TYPE_INT_32]] +; CHECK-SPIRV-DAG: %[[#TYPE_INT_16:]] = OpTypeInt 16 0 +; CHECK-SPIRV-DAG: %[[#TYPE_ORIG_FUNC_16:]] = OpTypeFunction %[[#TYPE_INT_16]] %[[#TYPE_INT_16]] %[[#TYPE_INT_16]] +; CHECK-SPIRV-DAG: %[[#TYPE_VEC_INT_16:]] = OpTypeVector %[[#TYPE_INT_16]] 2 +; CHECK-SPIRV-DAG: %[[#TYPE_ORIG_FUNC_VEC_INT_16:]] = OpTypeFunction %[[#TYPE_VEC_INT_16]] %[[#TYPE_VEC_INT_16]] %[[#TYPE_VEC_INT_16]] +; CHECK-SPIRV-DAG: %[[#TYPE_FSHL_FUNC_32:]] = OpTypeFunction %[[#TYPE_INT_32]] %[[#TYPE_INT_32]] %[[#TYPE_INT_32]] %[[#TYPE_INT_32]] +; CHECK-SPIRV-DAG: %[[#TYPE_FSHL_FUNC_16:]] = OpTypeFunction %[[#TYPE_INT_16]] %[[#TYPE_INT_16]] %[[#TYPE_INT_16]] %[[#TYPE_INT_16]] +; CHECK-SPIRV-DAG: %[[#TYPE_FSHL_FUNC_VEC_INT_16:]] = OpTypeFunction %[[#TYPE_VEC_INT_16]] %[[#TYPE_VEC_INT_16]] %[[#TYPE_VEC_INT_16]] %[[#TYPE_VEC_INT_16]] +; CHECK-SPIRV-DAG: %[[#CONST_ROTATE_32:]] = OpConstant %[[#TYPE_INT_32]] 8 +; CHECK-SPIRV-DAG: %[[#CONST_ROTATE_16:]] = OpConstant %[[#TYPE_INT_16]] 8 +; CHECK-SPIRV-DAG: %[[#CONST_ROTATE_VEC_INT_16:]] = OpConstantComposite %[[#TYPE_VEC_INT_16]] %[[#CONST_ROTATE_16]] %[[#CONST_ROTATE_16]] +; CHECK-SPIRV-DAG: %[[#CONST_TYPE_SIZE_32:]] = OpConstant %[[#TYPE_INT_32]] 32 ; CHECK-SPIRV: %[[#]] = OpFunction %[[#TYPE_INT_32]] {{.*}} %[[#TYPE_ORIG_FUNC_32]] ; CHECK-SPIRV: %[[#X:]] = OpFunctionParameter %[[#TYPE_INT_32]] diff --git a/llvm/test/CodeGen/SPIRV/llvm-intrinsics/fshr.ll b/llvm/test/CodeGen/SPIRV/llvm-intrinsics/fshr.ll index 4cf5ca53a411..55fb2d9079e1 100644 --- a/llvm/test/CodeGen/SPIRV/llvm-intrinsics/fshr.ll +++ b/llvm/test/CodeGen/SPIRV/llvm-intrinsics/fshr.ll @@ -1,20 +1,20 @@ ; RUN: llc -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV -; CHECK-SPIRV: OpName %[[#NAME_FSHR_FUNC_32:]] "spirv.llvm_fshr_i32" -; CHECK-SPIRV: OpName %[[#NAME_FSHR_FUNC_16:]] "spirv.llvm_fshr_i16" -; CHECK-SPIRV: OpName %[[#NAME_FSHR_FUNC_VEC_INT_16:]] "spirv.llvm_fshr_v2i16" -; CHECK-SPIRV: %[[#TYPE_INT_32:]] = OpTypeInt 32 0 -; CHECK-SPIRV: %[[#TYPE_ORIG_FUNC_32:]] = OpTypeFunction %[[#TYPE_INT_32]] %[[#TYPE_INT_32]] %[[#TYPE_INT_32]] -; CHECK-SPIRV: %[[#TYPE_INT_16:]] = OpTypeInt 16 0 -; CHECK-SPIRV: %[[#TYPE_ORIG_FUNC_16:]] = OpTypeFunction %[[#TYPE_INT_16]] %[[#TYPE_INT_16]] %[[#TYPE_INT_16]] -; CHECK-SPIRV: %[[#TYPE_VEC_INT_16:]] = OpTypeVector %[[#TYPE_INT_16]] 2 -; CHECK-SPIRV: %[[#TYPE_ORIG_FUNC_VEC_INT_16:]] = OpTypeFunction %[[#TYPE_VEC_INT_16]] %[[#TYPE_VEC_INT_16]] %[[#TYPE_VEC_INT_16]] -; CHECK-SPIRV: %[[#TYPE_FSHR_FUNC_32:]] = OpTypeFunction %[[#TYPE_INT_32]] %[[#TYPE_INT_32]] %[[#TYPE_INT_32]] %[[#TYPE_INT_32]] -; CHECK-SPIRV: %[[#TYPE_FSHR_FUNC_16:]] = OpTypeFunction %[[#TYPE_INT_16]] %[[#TYPE_INT_16]] %[[#TYPE_INT_16]] %[[#TYPE_INT_16]] -; CHECK-SPIRV: %[[#TYPE_FSHR_FUNC_VEC_INT_16:]] = OpTypeFunction %[[#TYPE_VEC_INT_16]] %[[#TYPE_VEC_INT_16]] %[[#TYPE_VEC_INT_16]] %[[#TYPE_VEC_INT_16]] +; CHECK-SPIRV-DAG: OpName %[[#NAME_FSHR_FUNC_32:]] "spirv.llvm_fshr_i32" +; CHECK-SPIRV-DAG: OpName %[[#NAME_FSHR_FUNC_16:]] "spirv.llvm_fshr_i16" +; CHECK-SPIRV-DAG: OpName %[[#NAME_FSHR_FUNC_VEC_INT_16:]] "spirv.llvm_fshr_v2i16" +; CHECK-SPIRV-DAG: %[[#TYPE_INT_32:]] = OpTypeInt 32 0 +; CHECK-SPIRV-DAG: %[[#TYPE_ORIG_FUNC_32:]] = OpTypeFunction %[[#TYPE_INT_32]] %[[#TYPE_INT_32]] %[[#TYPE_INT_32]] +; CHECK-SPIRV-DAG: %[[#TYPE_INT_16:]] = OpTypeInt 16 0 +; CHECK-SPIRV-DAG: %[[#TYPE_ORIG_FUNC_16:]] = OpTypeFunction %[[#TYPE_INT_16]] %[[#TYPE_INT_16]] %[[#TYPE_INT_16]] +; CHECK-SPIRV-DAG: %[[#TYPE_VEC_INT_16:]] = OpTypeVector %[[#TYPE_INT_16]] 2 +; CHECK-SPIRV-DAG: %[[#TYPE_ORIG_FUNC_VEC_INT_16:]] = OpTypeFunction %[[#TYPE_VEC_INT_16]] %[[#TYPE_VEC_INT_16]] %[[#TYPE_VEC_INT_16]] +; CHECK-SPIRV-DAG: %[[#TYPE_FSHR_FUNC_32:]] = OpTypeFunction %[[#TYPE_INT_32]] %[[#TYPE_INT_32]] %[[#TYPE_INT_32]] %[[#TYPE_INT_32]] +; CHECK-SPIRV-DAG: %[[#TYPE_FSHR_FUNC_16:]] = OpTypeFunction %[[#TYPE_INT_16]] %[[#TYPE_INT_16]] %[[#TYPE_INT_16]] %[[#TYPE_INT_16]] +; CHECK-SPIRV-DAG: %[[#TYPE_FSHR_FUNC_VEC_INT_16:]] = OpTypeFunction %[[#TYPE_VEC_INT_16]] %[[#TYPE_VEC_INT_16]] %[[#TYPE_VEC_INT_16]] %[[#TYPE_VEC_INT_16]] ; CHECK-SPIRV-DAG: %[[#CONST_ROTATE_32:]] = OpConstant %[[#TYPE_INT_32]] 8 ; CHECK-SPIRV-DAG: %[[#CONST_ROTATE_16:]] = OpConstant %[[#TYPE_INT_16]] 8 -; CHECK-SPIRV: %[[#CONST_ROTATE_VEC_INT_16:]] = OpConstantComposite %[[#TYPE_VEC_INT_16]] %[[#CONST_ROTATE_16]] %[[#CONST_ROTATE_16]] +; CHECK-SPIRV-DAG: %[[#CONST_ROTATE_VEC_INT_16:]] = OpConstantComposite %[[#TYPE_VEC_INT_16]] %[[#CONST_ROTATE_16]] %[[#CONST_ROTATE_16]] ; CHECK-SPIRV-DAG: %[[#CONST_TYPE_SIZE_32:]] = OpConstant %[[#TYPE_INT_32]] 32 ; CHECK-SPIRV: %[[#]] = OpFunction %[[#TYPE_INT_32]] {{.*}} %[[#TYPE_ORIG_FUNC_32]] diff --git a/llvm/test/CodeGen/SPIRV/llvm-intrinsics/memset.ll b/llvm/test/CodeGen/SPIRV/llvm-intrinsics/memset.ll index e7a986980f25..d5e70ae9e7aa 100644 --- a/llvm/test/CodeGen/SPIRV/llvm-intrinsics/memset.ll +++ b/llvm/test/CodeGen/SPIRV/llvm-intrinsics/memset.ll @@ -12,17 +12,17 @@ ; CHECK-DAG: %[[#Int8Ptr:]] = OpTypePointer Generic %[[#Int8]] ; CHECK-DAG: %[[#Const4:]] = OpConstant %[[#Int32]] 4 -; CHECK: %[[#Int8x4:]] = OpTypeArray %[[#Int8]] %[[#Const4]] +; CHECK-DAG: %[[#Int8x4:]] = OpTypeArray %[[#Int8]] %[[#Const4]] ; CHECK-DAG: %[[#Const12:]] = OpConstant %[[#Int32]] 12 -; CHECK: %[[#Int8x12:]] = OpTypeArray %[[#Int8]] %[[#Const12]] +; CHECK-DAG: %[[#Int8x12:]] = OpTypeArray %[[#Int8]] %[[#Const12]] ; CHECK-DAG: %[[#Const21:]] = OpConstant %[[#Int8]] 21 ; CHECK-DAG: %[[#False:]] = OpConstantFalse %[[#]] ; CHECK-DAG: %[[#ConstComp:]] = OpConstantComposite %[[#Int8x4]] %[[#Const21]] %[[#Const21]] %[[#Const21]] %[[#Const21]] ; CHECK-DAG: %[[#ConstNull:]] = OpConstantNull %[[#Int8x12]] -; CHECK: %[[#VarComp:]] = OpVariable %[[#]] UniformConstant %[[#ConstComp]] -; CHECK: %[[#VarNull:]] = OpVariable %[[#]] UniformConstant %[[#ConstNull]] +; CHECK-DAG: %[[#VarComp:]] = OpVariable %[[#]] UniformConstant %[[#ConstComp]] +; CHECK-DAG: %[[#VarNull:]] = OpVariable %[[#]] UniformConstant %[[#ConstNull]] ; CHECK-DAG: %[[#Int8PtrConst:]] = OpTypePointer UniformConstant %[[#Int8]] ; CHECK: OpCopyMemorySized %[[#Target:]] %[[#Source:]] %[[#Const12]] Aligned 4 diff --git a/llvm/test/CodeGen/SPIRV/logical-access-chain.ll b/llvm/test/CodeGen/SPIRV/logical-access-chain.ll index 39f6d33712ef..d56678ecfc2c 100644 --- a/llvm/test/CodeGen/SPIRV/logical-access-chain.ll +++ b/llvm/test/CodeGen/SPIRV/logical-access-chain.ll @@ -1,10 +1,10 @@ ; RUN: llc -O0 -mtriple=spirv-unknown-unknown %s -o - | FileCheck %s -; CHECK: [[uint:%[0-9]+]] = OpTypeInt 32 0 -; CHECK: [[uint2:%[0-9]+]] = OpTypeVector [[uint]] 2 -; CHECK: [[uint_1:%[0-9]+]] = OpConstant [[uint]] 1 -; CHECK: [[ptr_uint:%[0-9]+]] = OpTypePointer Function [[uint]] -; CHECK: [[ptr_uint2:%[0-9]+]] = OpTypePointer Function [[uint2]] +; CHECK-DAG: [[uint:%[0-9]+]] = OpTypeInt 32 0 +; CHECK-DAG: [[uint2:%[0-9]+]] = OpTypeVector [[uint]] 2 +; CHECK-DAG: [[uint_1:%[0-9]+]] = OpConstant [[uint]] 1 +; CHECK-DAG: [[ptr_uint:%[0-9]+]] = OpTypePointer Function [[uint]] +; CHECK-DAG: [[ptr_uint2:%[0-9]+]] = OpTypePointer Function [[uint2]] define void @main() #1 { entry: diff --git a/llvm/test/CodeGen/SPIRV/opencl/degrees.ll b/llvm/test/CodeGen/SPIRV/opencl/degrees.ll index 88f97835fe71..b8d4f52a2879 100644 --- a/llvm/test/CodeGen/SPIRV/opencl/degrees.ll +++ b/llvm/test/CodeGen/SPIRV/opencl/degrees.ll @@ -3,7 +3,7 @@ ; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv64-unknown-unknown %s -o - -filetype=obj | spirv-val %} ; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv32-unknown-unknown %s -o - -filetype=obj | spirv-val %} -; CHECK-DAG: %[[#op_ext_glsl:]] = OpExtInstImport "OpenCL.std" +; CHECK-DAG: %[[#op_ext_ocl:]] = OpExtInstImport "OpenCL.std" ; CHECK-DAG: %[[#float_32:]] = OpTypeFloat 32 ; CHECK-DAG: %[[#float_16:]] = OpTypeFloat 16 @@ -20,7 +20,7 @@ declare <4 x half> @llvm.spv.degrees.v4f16(<4 x half>) define noundef float @degrees_float(float noundef %a) { entry: ; CHECK: %[[#float_32_arg:]] = OpFunctionParameter %[[#float_32]] -; CHECK: %[[#]] = OpExtInst %[[#float_32]] %[[#op_ext_glsl]] degrees %[[#float_32_arg]] +; CHECK: %[[#]] = OpExtInst %[[#float_32]] %[[#op_ext_ocl]] degrees %[[#float_32_arg]] %elt.degrees = call float @llvm.spv.degrees.f32(float %a) ret float %elt.degrees } @@ -28,7 +28,7 @@ entry: define noundef half @degrees_half(half noundef %a) { entry: ; CHECK: %[[#float_16_arg:]] = OpFunctionParameter %[[#float_16]] -; CHECK: %[[#]] = OpExtInst %[[#float_16]] %[[#op_ext_glsl]] degrees %[[#float_16_arg]] +; CHECK: %[[#]] = OpExtInst %[[#float_16]] %[[#op_ext_ocl]] degrees %[[#float_16_arg]] %elt.degrees = call half @llvm.spv.degrees.f16(half %a) ret half %elt.degrees } @@ -36,7 +36,7 @@ entry: define noundef <4 x float> @degrees_float_vector(<4 x float> noundef %a) { entry: ; CHECK: %[[#vec4_float_32_arg:]] = OpFunctionParameter %[[#vec4_float_32]] -; CHECK: %[[#]] = OpExtInst %[[#vec4_float_32]] %[[#op_ext_glsl]] degrees %[[#vec4_float_32_arg]] +; CHECK: %[[#]] = OpExtInst %[[#vec4_float_32]] %[[#op_ext_ocl]] degrees %[[#vec4_float_32_arg]] %elt.degrees = call <4 x float> @llvm.spv.degrees.v4f32(<4 x float> %a) ret <4 x float> %elt.degrees } @@ -44,7 +44,7 @@ entry: define noundef <4 x half> @degrees_half_vector(<4 x half> noundef %a) { entry: ; CHECK: %[[#vec4_float_16_arg:]] = OpFunctionParameter %[[#vec4_float_16]] -; CHECK: %[[#]] = OpExtInst %[[#vec4_float_16]] %[[#op_ext_glsl]] degrees %[[#vec4_float_16_arg]] +; CHECK: %[[#]] = OpExtInst %[[#vec4_float_16]] %[[#op_ext_ocl]] degrees %[[#vec4_float_16_arg]] %elt.degrees = call <4 x half> @llvm.spv.degrees.v4f16(<4 x half> %a) ret <4 x half> %elt.degrees } diff --git a/llvm/test/CodeGen/SPIRV/opencl/distance.ll b/llvm/test/CodeGen/SPIRV/opencl/distance.ll new file mode 100644 index 000000000000..ac18804c00c9 --- /dev/null +++ b/llvm/test/CodeGen/SPIRV/opencl/distance.ll @@ -0,0 +1,34 @@ +; RUN: llc -verify-machineinstrs -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s +; RUN: llc -verify-machineinstrs -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s +; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv64-unknown-unknown %s -o - -filetype=obj | spirv-val %} +; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv32-unknown-unknown %s -o - -filetype=obj | spirv-val %} + +; CHECK-DAG: %[[#op_ext_cl:]] = OpExtInstImport "OpenCL.std" + +; CHECK-DAG: %[[#float_16:]] = OpTypeFloat 16 +; CHECK-DAG: %[[#vec4_float_16:]] = OpTypeVector %[[#float_16]] 4 +; CHECK-DAG: %[[#float_32:]] = OpTypeFloat 32 +; CHECK-DAG: %[[#vec4_float_32:]] = OpTypeVector %[[#float_32]] 4 + +define noundef half @distance_half4(<4 x half> noundef %a, <4 x half> noundef %b) { +entry: + ; CHECK: %[[#]] = OpFunction %[[#float_16]] None %[[#]] + ; CHECK: %[[#arg0:]] = OpFunctionParameter %[[#vec4_float_16]] + ; CHECK: %[[#arg1:]] = OpFunctionParameter %[[#vec4_float_16]] + ; CHECK: %[[#]] = OpExtInst %[[#float_16]] %[[#op_ext_cl]] distance %[[#arg0]] %[[#arg1]] + %spv.distance = call half @llvm.spv.distance.f16(<4 x half> %a, <4 x half> %b) + ret half %spv.distance +} + +define noundef float @distance_float4(<4 x float> noundef %a, <4 x float> noundef %b) { +entry: + ; CHECK: %[[#]] = OpFunction %[[#float_32]] None %[[#]] + ; CHECK: %[[#arg0:]] = OpFunctionParameter %[[#vec4_float_32]] + ; CHECK: %[[#arg1:]] = OpFunctionParameter %[[#vec4_float_32]] + ; CHECK: %[[#]] = OpExtInst %[[#float_32]] %[[#op_ext_cl]] distance %[[#arg0]] %[[#arg1]] + %spv.distance = call float @llvm.spv.distance.f32(<4 x float> %a, <4 x float> %b) + ret float %spv.distance +} + +declare half @llvm.spv.distance.f16(<4 x half>, <4 x half>) +declare float @llvm.spv.distance.f32(<4 x float>, <4 x float>) diff --git a/llvm/test/CodeGen/SPIRV/opencl/radians.ll b/llvm/test/CodeGen/SPIRV/opencl/radians.ll index f7bb8d5226cd..5b4f26a13a4c 100644 --- a/llvm/test/CodeGen/SPIRV/opencl/radians.ll +++ b/llvm/test/CodeGen/SPIRV/opencl/radians.ll @@ -3,7 +3,7 @@ ; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv64-unknown-unknown %s -o - -filetype=obj | spirv-val %} ; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv32-unknown-unknown %s -o - -filetype=obj | spirv-val %} -; CHECK-DAG: %[[#op_ext_glsl:]] = OpExtInstImport "OpenCL.std" +; CHECK-DAG: %[[#op_ext_ocl:]] = OpExtInstImport "OpenCL.std" ; CHECK-DAG: %[[#float_32:]] = OpTypeFloat 32 ; CHECK-DAG: %[[#float_16:]] = OpTypeFloat 16 @@ -20,7 +20,7 @@ declare <4 x half> @llvm.spv.radians.v4f16(<4 x half>) define noundef float @radians_float(float noundef %a) { entry: ; CHECK: %[[#float_32_arg:]] = OpFunctionParameter %[[#float_32]] -; CHECK: %[[#]] = OpExtInst %[[#float_32]] %[[#op_ext_glsl]] radians %[[#float_32_arg]] +; CHECK: %[[#]] = OpExtInst %[[#float_32]] %[[#op_ext_ocl]] radians %[[#float_32_arg]] %elt.radians = call float @llvm.spv.radians.f32(float %a) ret float %elt.radians } @@ -28,7 +28,7 @@ entry: define noundef half @radians_half(half noundef %a) { entry: ; CHECK: %[[#float_16_arg:]] = OpFunctionParameter %[[#float_16]] -; CHECK: %[[#]] = OpExtInst %[[#float_16]] %[[#op_ext_glsl]] radians %[[#float_16_arg]] +; CHECK: %[[#]] = OpExtInst %[[#float_16]] %[[#op_ext_ocl]] radians %[[#float_16_arg]] %elt.radians = call half @llvm.spv.radians.f16(half %a) ret half %elt.radians } @@ -36,7 +36,7 @@ entry: define noundef <4 x float> @radians_float_vector(<4 x float> noundef %a) { entry: ; CHECK: %[[#vec4_float_32_arg:]] = OpFunctionParameter %[[#vec4_float_32]] -; CHECK: %[[#]] = OpExtInst %[[#vec4_float_32]] %[[#op_ext_glsl]] radians %[[#vec4_float_32_arg]] +; CHECK: %[[#]] = OpExtInst %[[#vec4_float_32]] %[[#op_ext_ocl]] radians %[[#vec4_float_32_arg]] %elt.radians = call <4 x float> @llvm.spv.radians.v4f32(<4 x float> %a) ret <4 x float> %elt.radians } @@ -44,7 +44,7 @@ entry: define noundef <4 x half> @radians_half_vector(<4 x half> noundef %a) { entry: ; CHECK: %[[#vec4_float_16_arg:]] = OpFunctionParameter %[[#vec4_float_16]] -; CHECK: %[[#]] = OpExtInst %[[#vec4_float_16]] %[[#op_ext_glsl]] radians %[[#vec4_float_16_arg]] +; CHECK: %[[#]] = OpExtInst %[[#vec4_float_16]] %[[#op_ext_ocl]] radians %[[#vec4_float_16_arg]] %elt.radians = call <4 x half> @llvm.spv.radians.v4f16(<4 x half> %a) ret <4 x half> %elt.radians } diff --git a/llvm/test/CodeGen/SPIRV/pointers/PtrCast-null-in-OpSpecConstantOp.ll b/llvm/test/CodeGen/SPIRV/pointers/PtrCast-null-in-OpSpecConstantOp.ll index 99e2c3e6d396..dee16da7e099 100644 --- a/llvm/test/CodeGen/SPIRV/pointers/PtrCast-null-in-OpSpecConstantOp.ll +++ b/llvm/test/CodeGen/SPIRV/pointers/PtrCast-null-in-OpSpecConstantOp.ll @@ -5,10 +5,8 @@ ; CHECK-DAG: %[[Struct:.*]] = OpTypeStruct %[[Array]] ; CHECK-DAG: %[[Zero:.*]] = OpTypeInt 64 0 ; CHECK-DAG: %[[Null:.*]] = OpConstantNull %[[Zero]] -; CHECK-DAG: %[[R1:.*]] = OpConstantComposite %[[Array]] %[[Null]] -; CHECK-DAG: %[[#]] = OpConstantComposite %[[Struct]] %[[R1]] -; CHECK-DAG: %[[R2:.*]] = OpConstantComposite %[[Array]] %[[Null]] -; CHECK-DAG: %[[#]] = OpConstantComposite %[[Struct]] %[[R2]] +; CHECK-DAG: %[[R:.*]] = OpConstantComposite %[[Array]] %[[Null]] +; CHECK-DAG: %[[#]] = OpConstantComposite %[[Struct]] %[[R]] @G1 = addrspace(1) constant { [1 x ptr addrspace(4)] } { [1 x ptr addrspace(4)] [ptr addrspace(4) addrspacecast (ptr null to ptr addrspace(4))] } @G2 = addrspace(1) constant { [1 x ptr addrspace(4)] } { [1 x ptr addrspace(4)] [ptr addrspace(4) addrspacecast (ptr addrspace(1) null to ptr addrspace(4))] } diff --git a/llvm/test/CodeGen/SPIRV/pointers/struct-opaque-pointers.ll b/llvm/test/CodeGen/SPIRV/pointers/struct-opaque-pointers.ll index 03ecf5e8d839..59a24231769c 100644 --- a/llvm/test/CodeGen/SPIRV/pointers/struct-opaque-pointers.ll +++ b/llvm/test/CodeGen/SPIRV/pointers/struct-opaque-pointers.ll @@ -1,12 +1,12 @@ ; RUN: llc -verify-machineinstrs -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s ; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv32-unknown-unknown %s -o - -filetype=obj | spirv-val %} -; CHECK: %[[TyInt64:.*]] = OpTypeInt 64 0 -; CHECK: %[[TyInt64Ptr:.*]] = OpTypePointer {{[a-zA-Z]+}} %[[TyInt64]] -; CHECK: %[[TyStruct:.*]] = OpTypeStruct %[[TyInt64Ptr]] %[[TyInt64Ptr]] -; CHECK: %[[ConstStruct:.*]] = OpConstantComposite %[[TyStruct]] %[[ConstField:.*]] %[[ConstField]] -; CHECK: %[[TyStructPtr:.*]] = OpTypePointer {{[a-zA-Z]+}} %[[TyStruct]] -; CHECK: OpVariable %[[TyStructPtr]] {{[a-zA-Z]+}} %[[ConstStruct]] +; CHECK-DAG: %[[TyInt64:.*]] = OpTypeInt 64 0 +; CHECK-DAG: %[[TyInt64Ptr:.*]] = OpTypePointer {{[a-zA-Z]+}} %[[TyInt64]] +; CHECK-DAG: %[[TyStruct:.*]] = OpTypeStruct %[[TyInt64Ptr]] %[[TyInt64Ptr]] +; CHECK-DAG: %[[ConstStruct:.*]] = OpConstantComposite %[[TyStruct]] %[[ConstField:.*]] %[[ConstField]] +; CHECK-DAG: %[[TyStructPtr:.*]] = OpTypePointer {{[a-zA-Z]+}} %[[TyStruct]] +; CHECK-DAG: OpVariable %[[TyStructPtr]] {{[a-zA-Z]+}} %[[ConstStruct]] @a = addrspace(1) constant i64 42 @struct = addrspace(1) global {ptr addrspace(1), ptr addrspace(1)} { ptr addrspace(1) @a, ptr addrspace(1) @a } diff --git a/llvm/test/CodeGen/SPIRV/transcoding/SampledImage.ll b/llvm/test/CodeGen/SPIRV/transcoding/SampledImage.ll index e4c7bdb9e9c8..8a90e40e8881 100644 --- a/llvm/test/CodeGen/SPIRV/transcoding/SampledImage.ll +++ b/llvm/test/CodeGen/SPIRV/transcoding/SampledImage.ll @@ -24,12 +24,10 @@ ; CHECK-SPIRV: OpName %[[#sample_kernel_float:]] "sample_kernel_float" ; CHECK-SPIRV: OpName %[[#sample_kernel_int:]] "sample_kernel_int" -; CHECK-SPIRV: %[[#TypeSampler:]] = OpTypeSampler +; CHECK-SPIRV-DAG: %[[#TypeSampler:]] = OpTypeSampler ; CHECK-SPIRV-DAG: %[[#SampledImageTy:]] = OpTypeSampledImage ; CHECK-SPIRV-DAG: %[[#ConstSampler1:]] = OpConstantSampler %[[#TypeSampler]] None 0 Linear ; CHECK-SPIRV-DAG: %[[#ConstSampler2:]] = OpConstantSampler %[[#TypeSampler]] Repeat 0 Nearest -; CHECK-SPIRV-DAG: %[[#ConstSampler3:]] = OpConstantSampler %[[#TypeSampler]] None 0 Linear -; CHECK-SPIRV-DAG: %[[#ConstSampler4:]] = OpConstantSampler %[[#TypeSampler]] Repeat 0 Nearest ; CHECK-SPIRV: %[[#sample_kernel_float]] = OpFunction %{{.*}} ; CHECK-SPIRV: %[[#InputImage:]] = OpFunctionParameter %{{.*}} @@ -65,13 +63,13 @@ declare spir_func target("spirv.Sampler") @__translate_sampler_initializer(i32) ; CHECK-SPIRV: %[[#InputImage:]] = OpFunctionParameter %{{.*}} ; CHECK-SPIRV: %[[#argSampl:]] = OpFunctionParameter %[[#TypeSampler]] -; CHECK-SPIRV: %[[#SampledImage4:]] = OpSampledImage %[[#SampledImageTy]] %[[#InputImage]] %[[#ConstSampler3]] +; CHECK-SPIRV: %[[#SampledImage4:]] = OpSampledImage %[[#SampledImageTy]] %[[#InputImage]] %[[#ConstSampler1]] ; CHECK-SPIRV: %[[#]] = OpImageSampleExplicitLod %[[#]] %[[#SampledImage4]] ; CHECK-SPIRV: %[[#SampledImage5:]] = OpSampledImage %[[#SampledImageTy]] %[[#InputImage]] %[[#argSampl]] ; CHECK-SPIRV: %[[#]] = OpImageSampleExplicitLod %[[#]] %[[#SampledImage5]] -; CHECK-SPIRV: %[[#SampledImage6:]] = OpSampledImage %[[#SampledImageTy]] %[[#InputImage]] %[[#ConstSampler4]] +; CHECK-SPIRV: %[[#SampledImage6:]] = OpSampledImage %[[#SampledImageTy]] %[[#InputImage]] %[[#ConstSampler2]] ; CHECK-SPIRV: %[[#]] = OpImageSampleExplicitLod %[[#]] %[[#SampledImage6]] define dso_local spir_kernel void @sample_kernel_int(target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 0) %input, <2 x float> noundef %coords, <4 x i32> addrspace(1)* nocapture noundef writeonly %results, target("spirv.Sampler") %argSampl) local_unnamed_addr { diff --git a/llvm/test/CodeGen/SPIRV/transcoding/cl-types.ll b/llvm/test/CodeGen/SPIRV/transcoding/cl-types.ll index 8b326e265502..55f1125706f6 100644 --- a/llvm/test/CodeGen/SPIRV/transcoding/cl-types.ll +++ b/llvm/test/CodeGen/SPIRV/transcoding/cl-types.ll @@ -39,7 +39,7 @@ ; CHECK-SPIRV-DAG: %[[#SAMP:]] = OpTypeSampler ; CHECK-SPIRV-DAG: %[[#SAMPIMG:]] = OpTypeSampledImage %[[#IMG2D_RD]] -; CHECK-SPIRV: %[[#SAMP_CONST:]] = OpConstantSampler %[[#SAMP]] None 0 Linear +; CHECK-SPIRV-DAG: %[[#SAMP_CONST:]] = OpConstantSampler %[[#SAMP]] None 0 Linear ; CHECK-SPIRV: %[[#]] = OpFunctionParameter %[[#PIPE_RD]] ; CHECK-SPIRV: %[[#]] = OpFunctionParameter %[[#PIPE_WR]] diff --git a/llvm/test/CodeGen/SPIRV/transcoding/spirv-private-array-initialization.ll b/llvm/test/CodeGen/SPIRV/transcoding/spirv-private-array-initialization.ll index 74dbaab63e03..5810d9c59ee3 100644 --- a/llvm/test/CodeGen/SPIRV/transcoding/spirv-private-array-initialization.ll +++ b/llvm/test/CodeGen/SPIRV/transcoding/spirv-private-array-initialization.ll @@ -14,19 +14,19 @@ ; CHECK-SPIRV-DAG: %[[#twelve:]] = OpConstant %[[#i32]] 12 ; CHECK-SPIRV-DAG: %[[#const_i32x3_ptr:]] = OpTypePointer UniformConstant %[[#i32x3]] -; CHECK-SPIRV: %[[#test_arr2:]] = OpVariable %[[#const_i32x3_ptr]] UniformConstant %[[#test_arr_init]] -; CHECK-SPIRV: %[[#test_arr:]] = OpVariable %[[#const_i32x3_ptr]] UniformConstant %[[#test_arr_init]] +; CHECK-SPIRV-DAG: %[[#test_arr1:]] = OpVariable %[[#const_i32x3_ptr]] UniformConstant %[[#test_arr_init]] +; CHECK-SPIRV-DAG: %[[#test_arr2:]] = OpVariable %[[#const_i32x3_ptr]] UniformConstant %[[#test_arr_init]] ; CHECK-SPIRV-DAG: %[[#i32x3_ptr:]] = OpTypePointer Function %[[#i32x3]] -; CHECK-SPIRV: %[[#arr:]] = OpVariable %[[#i32x3_ptr]] Function +; CHECK-SPIRV: %[[#arr1:]] = OpVariable %[[#i32x3_ptr]] Function ; CHECK-SPIRV: %[[#arr2:]] = OpVariable %[[#i32x3_ptr]] Function -; CHECK-SPIRV-32: OpCopyMemorySized %[[#arr]] %[[#test_arr]] %[[#twelve]] Aligned 4 +; CHECK-SPIRV-32: OpCopyMemorySized %[[#arr1]] %[[#test_arr1]] %[[#twelve]] Aligned 4 ; CHECK-SPIRV-32: OpCopyMemorySized %[[#arr2]] %[[#test_arr2]] %[[#twelve]] Aligned 4 ; CHECK-SPIRV-64: %[[#twelvezext1:]] = OpUConvert %[[#i64:]] %[[#twelve:]] -; CHECK-SPIRV-64: OpCopyMemorySized %[[#arr]] %[[#test_arr]] %[[#twelvezext1]] Aligned 4 +; CHECK-SPIRV-64: OpCopyMemorySized %[[#arr1]] %[[#test_arr1]] %[[#twelvezext1]] Aligned 4 ; CHECK-SPIRV-64: %[[#twelvezext2:]] = OpUConvert %[[#i64:]] %[[#twelve:]] ; CHECK-SPIRV-64: OpCopyMemorySized %[[#arr2]] %[[#test_arr2]] %[[#twelvezext2]] Aligned 4 diff --git a/llvm/test/CodeGen/SPIRV/transcoding/sub_group_non_uniform_arithmetic.ll b/llvm/test/CodeGen/SPIRV/transcoding/sub_group_non_uniform_arithmetic.ll index adf73fe153de..62b09f6fe685 100644 --- a/llvm/test/CodeGen/SPIRV/transcoding/sub_group_non_uniform_arithmetic.ll +++ b/llvm/test/CodeGen/SPIRV/transcoding/sub_group_non_uniform_arithmetic.ll @@ -329,12 +329,12 @@ ; CHECK-SPIRV-DAG: %[[#double:]] = OpTypeFloat 64 ; CHECK-SPIRV-DAG: %[[#false:]] = OpConstantFalse %[[#bool]] +; CHECK-SPIRV-DAG: %[[#int_32:]] = OpConstant %[[#int]] 32 ; CHECK-SPIRV-DAG: %[[#ScopeSubgroup:]] = OpConstant %[[#int]] 3 ; CHECK-SPIRV-DAG: %[[#char_0:]] = OpConstant %[[#char]] 0 ; CHECK-SPIRV-DAG: %[[#char_10:]] = OpConstant %[[#char]] 10 ; CHECK-SPIRV-DAG: %[[#short_0:]] = OpConstant %[[#short]] 0 ; CHECK-SPIRV-DAG: %[[#int_0:]] = OpConstant %[[#int]] 0 -; CHECK-SPIRV-DAG: %[[#int_32:]] = OpConstant %[[#int]] 32 ; CHECK-SPIRV-DAG: %[[#long_0:]] = OpConstantNull %[[#long]] ; CHECK-SPIRV-DAG: %[[#half_0:]] = OpConstant %[[#half]] 0 ; CHECK-SPIRV-DAG: %[[#float_0:]] = OpConstant %[[#float]] 0 diff --git a/llvm/test/CodeGen/SPIRV/unnamed-global.ll b/llvm/test/CodeGen/SPIRV/unnamed-global.ll index f72334bd7752..90bac50507c0 100644 --- a/llvm/test/CodeGen/SPIRV/unnamed-global.ll +++ b/llvm/test/CodeGen/SPIRV/unnamed-global.ll @@ -4,10 +4,10 @@ ; RUN: llc -verify-machineinstrs -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s ; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv32-unknown-unknown %s -o - -filetype=obj | spirv-val %} -; CHECK: %[[TyInt:.*]] = OpTypeInt 8 0 -; CHECK: %[[ConstInt:.*]] = OpConstant %[[TyInt]] 123 -; CHECK: %[[TyPtr:.*]] = OpTypePointer {{[a-zA-Z]+}} %[[TyInt]] -; CHECK: %[[VarId:.*]] = OpVariable %[[TyPtr]] {{[a-zA-Z]+}} %[[ConstInt]] +; CHECK-DAG: %[[TyInt:.*]] = OpTypeInt 8 0 +; CHECK-DAG: %[[ConstInt:.*]] = OpConstant %[[TyInt]] 123 +; CHECK-DAG: %[[TyPtr:.*]] = OpTypePointer {{[a-zA-Z]+}} %[[TyInt]] +; CHECK-DAG: %[[VarId:.*]] = OpVariable %[[TyPtr]] {{[a-zA-Z]+}} %[[ConstInt]] @0 = addrspace(1) global i8 123 |