ahmed/OpenCL-CTS
1
0
Fork 0
mirror of https://github.com/KhronosGroup/OpenCL-CTS.git synced 2026-03-25 08:19:02 +00:00
Code Packages Projects Releases Wiki Activity

add SPIR-V 1.4 testing for no integer wrap decorations (#2041)

Browse Source 
This PR adds targeted testing for SPIR-V 1.4 features. Specifically,
this PR adds testing for the NoSignedWrap and NoUnsignedWrap
decorations.

We can reuse large parts of the testing for the
`cl_khr_spirv_no_integer_wrap_decoration` extension to test the SPIR-V
1.4 feature.
This commit is contained in:
Ben Ashbaugh
2024-10-22 09:47:57 -07:00
committed by GitHub
parent 94257defea
commit ec6394488a
20 changed files with 981 additions and 56 deletions
Download Patch File Download Diff File Expand all files Collapse all files

2
test_conformance/spirv_new/CMakeLists.txt
View File

@@ -4,10 +4,10 @@ set(${MODULE_NAME}_SOURCES
main.cpp main.cpp
test_basic_versions.cpp test_basic_versions.cpp
test_cl_khr_expect_assume.cpp test_cl_khr_expect_assume.cpp
test_cl_khr_spirv_no_integer_wrap_decoration.cpp
test_decorate.cpp test_decorate.cpp
test_get_program_il.cpp test_get_program_il.cpp
test_linkage.cpp test_linkage.cpp
test_no_integer_wrap_decoration.cpp
test_op_atomic.cpp test_op_atomic.cpp
test_op_branch_conditional.cpp test_op_branch_conditional.cpp
test_op_branch.cpp test_op_branch.cpp

45
test_conformance/spirv_new/spirv_asm/spv1.4/no_integer_wrap_decoration_fadd_int.spvasm32 Normal file
View File

@@ -0,0 +1,45 @@
; SPIR-V
; Version: 1.4
; Generator: Khronos SPIR-V Tools Assembler; 0
; Bound: 23
; Schema: 0
OpCapability Addresses
OpCapability Linkage
OpCapability Kernel
OpMemoryModel Physical32 OpenCL
OpEntryPoint Kernel %2 "fmath_cl" %__spirv_BuiltInGlobalInvocationId
OpSource OpenCL_C 200000
OpName %__spirv_BuiltInGlobalInvocationId "__spirv_BuiltInGlobalInvocationId"
OpName %out "out"
OpName %lhs "lhs"
OpName %rhs "rhs"
OpDecorate %7 FuncParamAttr NoWrite
%7 = OpDecorationGroup
OpDecorate %__spirv_BuiltInGlobalInvocationId BuiltIn GlobalInvocationId
OpDecorate %__spirv_BuiltInGlobalInvocationId Constant
OpDecorate %__spirv_BuiltInGlobalInvocationId LinkageAttributes "__spirv_BuiltInGlobalInvocationId" Import
OpDecorate %8 NoSignedWrap
OpGroupDecorate %7 %lhs %rhs
%uint = OpTypeInt 32 0
%v3uint = OpTypeVector %uint 3
%_ptr_Input_v3uint = OpTypePointer Input %v3uint
%void = OpTypeVoid
%_ptr_CrossWorkgroup_uint = OpTypePointer CrossWorkgroup %uint
%14 = OpTypeFunction %void %_ptr_CrossWorkgroup_uint %_ptr_CrossWorkgroup_uint %_ptr_CrossWorkgroup_uint
%__spirv_BuiltInGlobalInvocationId = OpVariable %_ptr_Input_v3uint Input
%2 = OpFunction %void None %14
%out = OpFunctionParameter %_ptr_CrossWorkgroup_uint
%lhs = OpFunctionParameter %_ptr_CrossWorkgroup_uint
%rhs = OpFunctionParameter %_ptr_CrossWorkgroup_uint
%15 = OpLabel
%16 = OpLoad %v3uint %__spirv_BuiltInGlobalInvocationId
%17 = OpCompositeExtract %uint %16 0
%18 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %lhs %17
%19 = OpLoad %uint %18 Aligned 4
%20 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %rhs %17
%21 = OpLoad %uint %20 Aligned 4
%8 = OpIAdd %uint %19 %21
%22 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %out %17
OpStore %22 %8 Aligned 4
OpReturn
OpFunctionEnd

51
test_conformance/spirv_new/spirv_asm/spv1.4/no_integer_wrap_decoration_fadd_int.spvasm64 Normal file
View File

@@ -0,0 +1,51 @@
; SPIR-V
; Version: 1.4
; Generator: Khronos SPIR-V Tools Assembler; 0
; Bound: 28
; Schema: 0
OpCapability Addresses
OpCapability Linkage
OpCapability Kernel
OpCapability Int64
OpMemoryModel Physical64 OpenCL
OpEntryPoint Kernel %2 "fmath_cl" %__spirv_BuiltInGlobalInvocationId
OpSource OpenCL_C 200000
OpName %__spirv_BuiltInGlobalInvocationId "__spirv_BuiltInGlobalInvocationId"
OpName %out "out"
OpName %lhs "lhs"
OpName %rhs "rhs"
OpDecorate %7 FuncParamAttr NoWrite
%7 = OpDecorationGroup
OpDecorate %__spirv_BuiltInGlobalInvocationId BuiltIn GlobalInvocationId
OpDecorate %__spirv_BuiltInGlobalInvocationId Constant
OpDecorate %__spirv_BuiltInGlobalInvocationId LinkageAttributes "__spirv_BuiltInGlobalInvocationId" Import
OpDecorate %8 NoSignedWrap
OpGroupDecorate %7 %lhs %rhs
%ulong = OpTypeInt 64 0
%uint = OpTypeInt 32 0
%v3ulong = OpTypeVector %ulong 3
%_ptr_Input_v3ulong = OpTypePointer Input %v3ulong
%void = OpTypeVoid
%_ptr_CrossWorkgroup_uint = OpTypePointer CrossWorkgroup %uint
%15 = OpTypeFunction %void %_ptr_CrossWorkgroup_uint %_ptr_CrossWorkgroup_uint %_ptr_CrossWorkgroup_uint
%__spirv_BuiltInGlobalInvocationId = OpVariable %_ptr_Input_v3ulong Input
%2 = OpFunction %void None %15
%out = OpFunctionParameter %_ptr_CrossWorkgroup_uint
%lhs = OpFunctionParameter %_ptr_CrossWorkgroup_uint
%rhs = OpFunctionParameter %_ptr_CrossWorkgroup_uint
%16 = OpLabel
%17 = OpLoad %v3ulong %__spirv_BuiltInGlobalInvocationId
%18 = OpCompositeExtract %ulong %17 0
%19 = OpUConvert %uint %18
%20 = OpSConvert %ulong %19
%21 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %lhs %20
%22 = OpLoad %uint %21 Aligned 4
%23 = OpSConvert %ulong %19
%24 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %rhs %23
%25 = OpLoad %uint %24 Aligned 4
%8 = OpIAdd %uint %22 %25
%26 = OpSConvert %ulong %19
%27 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %out %26
OpStore %27 %8 Aligned 4
OpReturn
OpFunctionEnd

45
test_conformance/spirv_new/spirv_asm/spv1.4/no_integer_wrap_decoration_fadd_uint.spvasm32 Normal file
View File

@@ -0,0 +1,45 @@
; SPIR-V
; Version: 1.4
; Generator: Khronos SPIR-V Tools Assembler; 0
; Bound: 23
; Schema: 0
OpCapability Addresses
OpCapability Linkage
OpCapability Kernel
OpMemoryModel Physical32 OpenCL
OpEntryPoint Kernel %2 "fmath_cl" %__spirv_BuiltInGlobalInvocationId
OpSource OpenCL_C 200000
OpName %__spirv_BuiltInGlobalInvocationId "__spirv_BuiltInGlobalInvocationId"
OpName %out "out"
OpName %lhs "lhs"
OpName %rhs "rhs"
OpDecorate %7 FuncParamAttr NoWrite
%7 = OpDecorationGroup
OpDecorate %__spirv_BuiltInGlobalInvocationId BuiltIn GlobalInvocationId
OpDecorate %__spirv_BuiltInGlobalInvocationId Constant
OpDecorate %__spirv_BuiltInGlobalInvocationId LinkageAttributes "__spirv_BuiltInGlobalInvocationId" Import
OpDecorate %8 NoUnsignedWrap
OpGroupDecorate %7 %lhs %rhs
%uint = OpTypeInt 32 0
%v3uint = OpTypeVector %uint 3
%_ptr_Input_v3uint = OpTypePointer Input %v3uint
%void = OpTypeVoid
%_ptr_CrossWorkgroup_uint = OpTypePointer CrossWorkgroup %uint
%14 = OpTypeFunction %void %_ptr_CrossWorkgroup_uint %_ptr_CrossWorkgroup_uint %_ptr_CrossWorkgroup_uint
%__spirv_BuiltInGlobalInvocationId = OpVariable %_ptr_Input_v3uint Input
%2 = OpFunction %void None %14
%out = OpFunctionParameter %_ptr_CrossWorkgroup_uint
%lhs = OpFunctionParameter %_ptr_CrossWorkgroup_uint
%rhs = OpFunctionParameter %_ptr_CrossWorkgroup_uint
%15 = OpLabel
%16 = OpLoad %v3uint %__spirv_BuiltInGlobalInvocationId
%17 = OpCompositeExtract %uint %16 0
%18 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %lhs %17
%19 = OpLoad %uint %18 Aligned 4
%20 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %rhs %17
%21 = OpLoad %uint %20 Aligned 4
%8 = OpIAdd %uint %19 %21
%22 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %out %17
OpStore %22 %8 Aligned 4
OpReturn
OpFunctionEnd

51
test_conformance/spirv_new/spirv_asm/spv1.4/no_integer_wrap_decoration_fadd_uint.spvasm64 Normal file
View File

@@ -0,0 +1,51 @@
; SPIR-V
; Version: 1.4
; Generator: Khronos SPIR-V Tools Assembler; 0
; Bound: 28
; Schema: 0
OpCapability Addresses
OpCapability Linkage
OpCapability Kernel
OpCapability Int64
OpMemoryModel Physical64 OpenCL
OpEntryPoint Kernel %2 "fmath_cl" %__spirv_BuiltInGlobalInvocationId
OpSource OpenCL_C 200000
OpName %__spirv_BuiltInGlobalInvocationId "__spirv_BuiltInGlobalInvocationId"
OpName %out "out"
OpName %lhs "lhs"
OpName %rhs "rhs"
OpDecorate %7 FuncParamAttr NoWrite
%7 = OpDecorationGroup
OpDecorate %__spirv_BuiltInGlobalInvocationId BuiltIn GlobalInvocationId
OpDecorate %__spirv_BuiltInGlobalInvocationId Constant
OpDecorate %__spirv_BuiltInGlobalInvocationId LinkageAttributes "__spirv_BuiltInGlobalInvocationId" Import
OpDecorate %8 NoUnsignedWrap
OpGroupDecorate %7 %lhs %rhs
%ulong = OpTypeInt 64 0
%uint = OpTypeInt 32 0
%v3ulong = OpTypeVector %ulong 3
%_ptr_Input_v3ulong = OpTypePointer Input %v3ulong
%void = OpTypeVoid
%_ptr_CrossWorkgroup_uint = OpTypePointer CrossWorkgroup %uint
%15 = OpTypeFunction %void %_ptr_CrossWorkgroup_uint %_ptr_CrossWorkgroup_uint %_ptr_CrossWorkgroup_uint
%__spirv_BuiltInGlobalInvocationId = OpVariable %_ptr_Input_v3ulong Input
%2 = OpFunction %void None %15
%out = OpFunctionParameter %_ptr_CrossWorkgroup_uint
%lhs = OpFunctionParameter %_ptr_CrossWorkgroup_uint
%rhs = OpFunctionParameter %_ptr_CrossWorkgroup_uint
%16 = OpLabel
%17 = OpLoad %v3ulong %__spirv_BuiltInGlobalInvocationId
%18 = OpCompositeExtract %ulong %17 0
%19 = OpUConvert %uint %18
%20 = OpSConvert %ulong %19
%21 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %lhs %20
%22 = OpLoad %uint %21 Aligned 4
%23 = OpSConvert %ulong %19
%24 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %rhs %23
%25 = OpLoad %uint %24 Aligned 4
%8 = OpIAdd %uint %22 %25
%26 = OpSConvert %ulong %19
%27 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %out %26
OpStore %27 %8 Aligned 4
OpReturn
OpFunctionEnd

45
test_conformance/spirv_new/spirv_asm/spv1.4/no_integer_wrap_decoration_fmul_int.spvasm32 Normal file
View File

@@ -0,0 +1,45 @@
; SPIR-V
; Version: 1.4
; Generator: Khronos SPIR-V Tools Assembler; 0
; Bound: 23
; Schema: 0
OpCapability Addresses
OpCapability Linkage
OpCapability Kernel
OpMemoryModel Physical32 OpenCL
OpEntryPoint Kernel %2 "fmath_cl" %__spirv_BuiltInGlobalInvocationId
OpSource OpenCL_C 200000
OpName %__spirv_BuiltInGlobalInvocationId "__spirv_BuiltInGlobalInvocationId"
OpName %out "out"
OpName %lhs "lhs"
OpName %rhs "rhs"
OpDecorate %7 FuncParamAttr NoWrite
%7 = OpDecorationGroup
OpDecorate %__spirv_BuiltInGlobalInvocationId BuiltIn GlobalInvocationId
OpDecorate %__spirv_BuiltInGlobalInvocationId Constant
OpDecorate %__spirv_BuiltInGlobalInvocationId LinkageAttributes "__spirv_BuiltInGlobalInvocationId" Import
OpDecorate %8 NoSignedWrap
OpGroupDecorate %7 %lhs %rhs
%uint = OpTypeInt 32 0
%v3uint = OpTypeVector %uint 3
%_ptr_Input_v3uint = OpTypePointer Input %v3uint
%void = OpTypeVoid
%_ptr_CrossWorkgroup_uint = OpTypePointer CrossWorkgroup %uint
%14 = OpTypeFunction %void %_ptr_CrossWorkgroup_uint %_ptr_CrossWorkgroup_uint %_ptr_CrossWorkgroup_uint
%__spirv_BuiltInGlobalInvocationId = OpVariable %_ptr_Input_v3uint Input
%2 = OpFunction %void None %14
%out = OpFunctionParameter %_ptr_CrossWorkgroup_uint
%lhs = OpFunctionParameter %_ptr_CrossWorkgroup_uint
%rhs = OpFunctionParameter %_ptr_CrossWorkgroup_uint
%15 = OpLabel
%16 = OpLoad %v3uint %__spirv_BuiltInGlobalInvocationId
%17 = OpCompositeExtract %uint %16 0
%18 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %lhs %17
%19 = OpLoad %uint %18 Aligned 4
%20 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %rhs %17
%21 = OpLoad %uint %20 Aligned 4
%8 = OpIMul %uint %19 %21
%22 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %out %17
OpStore %22 %8 Aligned 4
OpReturn
OpFunctionEnd

51
test_conformance/spirv_new/spirv_asm/spv1.4/no_integer_wrap_decoration_fmul_int.spvasm64 Normal file
View File

@@ -0,0 +1,51 @@
; SPIR-V
; Version: 1.4
; Generator: Khronos SPIR-V Tools Assembler; 0
; Bound: 28
; Schema: 0
OpCapability Addresses
OpCapability Linkage
OpCapability Kernel
OpCapability Int64
OpMemoryModel Physical64 OpenCL
OpEntryPoint Kernel %2 "fmath_cl" %__spirv_BuiltInGlobalInvocationId
OpSource OpenCL_C 200000
OpName %__spirv_BuiltInGlobalInvocationId "__spirv_BuiltInGlobalInvocationId"
OpName %out "out"
OpName %lhs "lhs"
OpName %rhs "rhs"
OpDecorate %7 FuncParamAttr NoWrite
%7 = OpDecorationGroup
OpDecorate %__spirv_BuiltInGlobalInvocationId BuiltIn GlobalInvocationId
OpDecorate %__spirv_BuiltInGlobalInvocationId Constant
OpDecorate %__spirv_BuiltInGlobalInvocationId LinkageAttributes "__spirv_BuiltInGlobalInvocationId" Import
OpDecorate %8 NoSignedWrap
OpGroupDecorate %7 %lhs %rhs
%ulong = OpTypeInt 64 0
%uint = OpTypeInt 32 0
%v3ulong = OpTypeVector %ulong 3
%_ptr_Input_v3ulong = OpTypePointer Input %v3ulong
%void = OpTypeVoid
%_ptr_CrossWorkgroup_uint = OpTypePointer CrossWorkgroup %uint
%15 = OpTypeFunction %void %_ptr_CrossWorkgroup_uint %_ptr_CrossWorkgroup_uint %_ptr_CrossWorkgroup_uint
%__spirv_BuiltInGlobalInvocationId = OpVariable %_ptr_Input_v3ulong Input
%2 = OpFunction %void None %15
%out = OpFunctionParameter %_ptr_CrossWorkgroup_uint
%lhs = OpFunctionParameter %_ptr_CrossWorkgroup_uint
%rhs = OpFunctionParameter %_ptr_CrossWorkgroup_uint
%16 = OpLabel
%17 = OpLoad %v3ulong %__spirv_BuiltInGlobalInvocationId
%18 = OpCompositeExtract %ulong %17 0
%19 = OpUConvert %uint %18
%20 = OpSConvert %ulong %19
%21 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %lhs %20
%22 = OpLoad %uint %21 Aligned 4
%23 = OpSConvert %ulong %19
%24 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %rhs %23
%25 = OpLoad %uint %24 Aligned 4
%8 = OpIMul %uint %22 %25
%26 = OpSConvert %ulong %19
%27 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %out %26
OpStore %27 %8 Aligned 4
OpReturn
OpFunctionEnd

45
test_conformance/spirv_new/spirv_asm/spv1.4/no_integer_wrap_decoration_fmul_uint.spvasm32 Normal file
View File

@@ -0,0 +1,45 @@
; SPIR-V
; Version: 1.4
; Generator: Khronos SPIR-V Tools Assembler; 0
; Bound: 23
; Schema: 0
OpCapability Addresses
OpCapability Linkage
OpCapability Kernel
OpMemoryModel Physical32 OpenCL
OpEntryPoint Kernel %2 "fmath_cl" %__spirv_BuiltInGlobalInvocationId
OpSource OpenCL_C 200000
OpName %__spirv_BuiltInGlobalInvocationId "__spirv_BuiltInGlobalInvocationId"
OpName %out "out"
OpName %lhs "lhs"
OpName %rhs "rhs"
OpDecorate %7 FuncParamAttr NoWrite
%7 = OpDecorationGroup
OpDecorate %__spirv_BuiltInGlobalInvocationId BuiltIn GlobalInvocationId
OpDecorate %__spirv_BuiltInGlobalInvocationId Constant
OpDecorate %__spirv_BuiltInGlobalInvocationId LinkageAttributes "__spirv_BuiltInGlobalInvocationId" Import
OpDecorate %8 NoUnsignedWrap
OpGroupDecorate %7 %lhs %rhs
%uint = OpTypeInt 32 0
%v3uint = OpTypeVector %uint 3
%_ptr_Input_v3uint = OpTypePointer Input %v3uint
%void = OpTypeVoid
%_ptr_CrossWorkgroup_uint = OpTypePointer CrossWorkgroup %uint
%14 = OpTypeFunction %void %_ptr_CrossWorkgroup_uint %_ptr_CrossWorkgroup_uint %_ptr_CrossWorkgroup_uint
%__spirv_BuiltInGlobalInvocationId = OpVariable %_ptr_Input_v3uint Input
%2 = OpFunction %void None %14
%out = OpFunctionParameter %_ptr_CrossWorkgroup_uint
%lhs = OpFunctionParameter %_ptr_CrossWorkgroup_uint
%rhs = OpFunctionParameter %_ptr_CrossWorkgroup_uint
%15 = OpLabel
%16 = OpLoad %v3uint %__spirv_BuiltInGlobalInvocationId
%17 = OpCompositeExtract %uint %16 0
%18 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %lhs %17
%19 = OpLoad %uint %18 Aligned 4
%20 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %rhs %17
%21 = OpLoad %uint %20 Aligned 4
%8 = OpIMul %uint %19 %21
%22 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %out %17
OpStore %22 %8 Aligned 4
OpReturn
OpFunctionEnd

51
test_conformance/spirv_new/spirv_asm/spv1.4/no_integer_wrap_decoration_fmul_uint.spvasm64 Normal file
View File

@@ -0,0 +1,51 @@
; SPIR-V
; Version: 1.3
; Generator: Khronos SPIR-V Tools Assembler; 0
; Bound: 28
; Schema: 0
OpCapability Addresses
OpCapability Linkage
OpCapability Kernel
OpCapability Int64
OpMemoryModel Physical64 OpenCL
OpEntryPoint Kernel %2 "fmath_cl" %__spirv_BuiltInGlobalInvocationId
OpSource OpenCL_C 200000
OpName %__spirv_BuiltInGlobalInvocationId "__spirv_BuiltInGlobalInvocationId"
OpName %out "out"
OpName %lhs "lhs"
OpName %rhs "rhs"
OpDecorate %7 FuncParamAttr NoWrite
%7 = OpDecorationGroup
OpDecorate %__spirv_BuiltInGlobalInvocationId BuiltIn GlobalInvocationId
OpDecorate %__spirv_BuiltInGlobalInvocationId Constant
OpDecorate %__spirv_BuiltInGlobalInvocationId LinkageAttributes "__spirv_BuiltInGlobalInvocationId" Import
OpDecorate %8 NoUnsignedWrap
OpGroupDecorate %7 %lhs %rhs
%ulong = OpTypeInt 64 0
%uint = OpTypeInt 32 0
%v3ulong = OpTypeVector %ulong 3
%_ptr_Input_v3ulong = OpTypePointer Input %v3ulong
%void = OpTypeVoid
%_ptr_CrossWorkgroup_uint = OpTypePointer CrossWorkgroup %uint
%15 = OpTypeFunction %void %_ptr_CrossWorkgroup_uint %_ptr_CrossWorkgroup_uint %_ptr_CrossWorkgroup_uint
%__spirv_BuiltInGlobalInvocationId = OpVariable %_ptr_Input_v3ulong Input
%2 = OpFunction %void None %15
%out = OpFunctionParameter %_ptr_CrossWorkgroup_uint
%lhs = OpFunctionParameter %_ptr_CrossWorkgroup_uint
%rhs = OpFunctionParameter %_ptr_CrossWorkgroup_uint
%16 = OpLabel
%17 = OpLoad %v3ulong %__spirv_BuiltInGlobalInvocationId
%18 = OpCompositeExtract %ulong %17 0
%19 = OpUConvert %uint %18
%20 = OpSConvert %ulong %19
%21 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %lhs %20
%22 = OpLoad %uint %21 Aligned 4
%23 = OpSConvert %ulong %19
%24 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %rhs %23
%25 = OpLoad %uint %24 Aligned 4
%8 = OpIMul %uint %22 %25
%26 = OpSConvert %ulong %19
%27 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %out %26
OpStore %27 %8 Aligned 4
OpReturn
OpFunctionEnd

44
test_conformance/spirv_new/spirv_asm/spv1.4/no_integer_wrap_decoration_fnegate_int.spvasm32 Normal file
View File

@@ -0,0 +1,44 @@
; SPIR-V
; Version: 1.4
; Generator: Khronos SPIR-V Tools Assembler; 0
; Bound: 22
; Schema: 0
OpCapability Addresses
OpCapability Linkage
OpCapability Kernel
OpMemoryModel Physical32 OpenCL
OpEntryPoint Kernel %2 "fmath_cl" %__spirv_BuiltInGlobalInvocationId
OpSource OpenCL_C 200000
OpName %__spirv_BuiltInGlobalInvocationId "__spirv_BuiltInGlobalInvocationId"
OpName %out "out"
OpName %lhs "lhs"
OpName %rhs "rhs"
OpDecorate %7 FuncParamAttr NoWrite
OpDecorate %8 NoSignedWrap
%7 = OpDecorationGroup
OpDecorate %__spirv_BuiltInGlobalInvocationId BuiltIn GlobalInvocationId
OpDecorate %__spirv_BuiltInGlobalInvocationId Constant
OpDecorate %__spirv_BuiltInGlobalInvocationId LinkageAttributes "__spirv_BuiltInGlobalInvocationId" Import
OpGroupDecorate %7 %lhs %rhs
%uint = OpTypeInt 32 0
%uint_0 = OpConstant %uint 0
%v3uint = OpTypeVector %uint 3
%_ptr_Input_v3uint = OpTypePointer Input %v3uint
%void = OpTypeVoid
%_ptr_CrossWorkgroup_uint = OpTypePointer CrossWorkgroup %uint
%15 = OpTypeFunction %void %_ptr_CrossWorkgroup_uint %_ptr_CrossWorkgroup_uint %_ptr_CrossWorkgroup_uint
%__spirv_BuiltInGlobalInvocationId = OpVariable %_ptr_Input_v3uint Input
%2 = OpFunction %void None %15
%out = OpFunctionParameter %_ptr_CrossWorkgroup_uint
%lhs = OpFunctionParameter %_ptr_CrossWorkgroup_uint
%rhs = OpFunctionParameter %_ptr_CrossWorkgroup_uint
%16 = OpLabel
%17 = OpLoad %v3uint %__spirv_BuiltInGlobalInvocationId
%18 = OpCompositeExtract %uint %17 0
%19 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %lhs %18
%20 = OpLoad %uint %19 Aligned 4
%8 = OpSNegate %uint %20
%21 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %out %18
OpStore %21 %8 Aligned 4
OpReturn
OpFunctionEnd

49
test_conformance/spirv_new/spirv_asm/spv1.4/no_integer_wrap_decoration_fnegate_int.spvasm64 Normal file
View File

@@ -0,0 +1,49 @@
; SPIR-V
; Version: 1.4
; Generator: Khronos SPIR-V Tools Assembler; 0
; Bound: 26
; Schema: 0
OpCapability Addresses
OpCapability Linkage
OpCapability Kernel
OpCapability Int64
OpMemoryModel Physical64 OpenCL
OpEntryPoint Kernel %2 "fmath_cl" %__spirv_BuiltInGlobalInvocationId
OpSource OpenCL_C 200000
OpName %__spirv_BuiltInGlobalInvocationId "__spirv_BuiltInGlobalInvocationId"
OpName %out "out"
OpName %lhs "lhs"
OpName %rhs "rhs"
OpDecorate %7 FuncParamAttr NoWrite
OpDecorate %8 NoSignedWrap
%7 = OpDecorationGroup
OpDecorate %__spirv_BuiltInGlobalInvocationId BuiltIn GlobalInvocationId
OpDecorate %__spirv_BuiltInGlobalInvocationId Constant
OpDecorate %__spirv_BuiltInGlobalInvocationId LinkageAttributes "__spirv_BuiltInGlobalInvocationId" Import
OpGroupDecorate %7 %lhs %rhs
%ulong = OpTypeInt 64 0
%uint = OpTypeInt 32 0
%uint_0 = OpConstant %uint 0
%v3ulong = OpTypeVector %ulong 3
%_ptr_Input_v3ulong = OpTypePointer Input %v3ulong
%void = OpTypeVoid
%_ptr_CrossWorkgroup_uint = OpTypePointer CrossWorkgroup %uint
%16 = OpTypeFunction %void %_ptr_CrossWorkgroup_uint %_ptr_CrossWorkgroup_uint %_ptr_CrossWorkgroup_uint
%__spirv_BuiltInGlobalInvocationId = OpVariable %_ptr_Input_v3ulong Input
%2 = OpFunction %void None %16
%out = OpFunctionParameter %_ptr_CrossWorkgroup_uint
%lhs = OpFunctionParameter %_ptr_CrossWorkgroup_uint
%rhs = OpFunctionParameter %_ptr_CrossWorkgroup_uint
%17 = OpLabel
%18 = OpLoad %v3ulong %__spirv_BuiltInGlobalInvocationId
%19 = OpCompositeExtract %ulong %18 0
%20 = OpUConvert %uint %19
%21 = OpSConvert %ulong %20
%22 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %lhs %21
%23 = OpLoad %uint %22 Aligned 4
%8 = OpSNegate %uint %23
%24 = OpSConvert %ulong %20
%25 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %out %24
OpStore %25 %8 Aligned 4
OpReturn
OpFunctionEnd

47
test_conformance/spirv_new/spirv_asm/spv1.4/no_integer_wrap_decoration_fshiftleft_int.spvasm32 Normal file
View File

@@ -0,0 +1,47 @@
; SPIR-V
; Version: 1.4
; Generator: Khronos SPIR-V Tools Assembler; 0
; Bound: 25
; Schema: 0
OpCapability Addresses
OpCapability Linkage
OpCapability Kernel
OpMemoryModel Physical32 OpenCL
OpEntryPoint Kernel %2 "fmath_cl" %__spirv_BuiltInGlobalInvocationId
OpSource OpenCL_C 200000
OpName %__spirv_BuiltInGlobalInvocationId "__spirv_BuiltInGlobalInvocationId"
OpName %out "out"
OpName %lhs "lhs"
OpName %rhs "rhs"
OpDecorate %7 FuncParamAttr NoWrite
%7 = OpDecorationGroup
OpDecorate %__spirv_BuiltInGlobalInvocationId BuiltIn GlobalInvocationId
OpDecorate %__spirv_BuiltInGlobalInvocationId Constant
OpDecorate %__spirv_BuiltInGlobalInvocationId LinkageAttributes "__spirv_BuiltInGlobalInvocationId" Import
OpDecorate %8 NoSignedWrap
OpGroupDecorate %7 %lhs %rhs
%uint = OpTypeInt 32 0
%uint_31 = OpConstant %uint 31
%v3uint = OpTypeVector %uint 3
%_ptr_Input_v3uint = OpTypePointer Input %v3uint
%void = OpTypeVoid
%_ptr_CrossWorkgroup_uint = OpTypePointer CrossWorkgroup %uint
%15 = OpTypeFunction %void %_ptr_CrossWorkgroup_uint %_ptr_CrossWorkgroup_uint %_ptr_CrossWorkgroup_uint
%__spirv_BuiltInGlobalInvocationId = OpVariable %_ptr_Input_v3uint Input
%2 = OpFunction %void None %15
%out = OpFunctionParameter %_ptr_CrossWorkgroup_uint
%lhs = OpFunctionParameter %_ptr_CrossWorkgroup_uint
%rhs = OpFunctionParameter %_ptr_CrossWorkgroup_uint
%16 = OpLabel
%17 = OpLoad %v3uint %__spirv_BuiltInGlobalInvocationId
%18 = OpCompositeExtract %uint %17 0
%19 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %lhs %18
%20 = OpLoad %uint %19 Aligned 4
%21 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %rhs %18
%22 = OpLoad %uint %21 Aligned 4
%23 = OpBitwiseAnd %uint %22 %uint_31
%8 = OpShiftLeftLogical %uint %20 %23
%24 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %out %18
OpStore %24 %8 Aligned 4
OpReturn
OpFunctionEnd

53
test_conformance/spirv_new/spirv_asm/spv1.4/no_integer_wrap_decoration_fshiftleft_int.spvasm64 Normal file
View File

@@ -0,0 +1,53 @@
; SPIR-V
; Version: 1.4
; Generator: Khronos SPIR-V Tools Assembler; 0
; Bound: 30
; Schema: 0
OpCapability Addresses
OpCapability Linkage
OpCapability Kernel
OpCapability Int64
OpMemoryModel Physical64 OpenCL
OpEntryPoint Kernel %2 "fmath_cl" %__spirv_BuiltInGlobalInvocationId
OpSource OpenCL_C 200000
OpName %__spirv_BuiltInGlobalInvocationId "__spirv_BuiltInGlobalInvocationId"
OpName %out "out"
OpName %lhs "lhs"
OpName %rhs "rhs"
OpDecorate %7 FuncParamAttr NoWrite
%7 = OpDecorationGroup
OpDecorate %__spirv_BuiltInGlobalInvocationId BuiltIn GlobalInvocationId
OpDecorate %__spirv_BuiltInGlobalInvocationId Constant
OpDecorate %__spirv_BuiltInGlobalInvocationId LinkageAttributes "__spirv_BuiltInGlobalInvocationId" Import
OpDecorate %8 NoSignedWrap
OpGroupDecorate %7 %lhs %rhs
%ulong = OpTypeInt 64 0
%uint = OpTypeInt 32 0
%uint_31 = OpConstant %uint 31
%v3ulong = OpTypeVector %ulong 3
%_ptr_Input_v3ulong = OpTypePointer Input %v3ulong
%void = OpTypeVoid
%_ptr_CrossWorkgroup_uint = OpTypePointer CrossWorkgroup %uint
%16 = OpTypeFunction %void %_ptr_CrossWorkgroup_uint %_ptr_CrossWorkgroup_uint %_ptr_CrossWorkgroup_uint
%__spirv_BuiltInGlobalInvocationId = OpVariable %_ptr_Input_v3ulong Input
%2 = OpFunction %void None %16
%out = OpFunctionParameter %_ptr_CrossWorkgroup_uint
%lhs = OpFunctionParameter %_ptr_CrossWorkgroup_uint
%rhs = OpFunctionParameter %_ptr_CrossWorkgroup_uint
%17 = OpLabel
%18 = OpLoad %v3ulong %__spirv_BuiltInGlobalInvocationId
%19 = OpCompositeExtract %ulong %18 0
%20 = OpUConvert %uint %19
%21 = OpSConvert %ulong %20
%22 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %lhs %21
%23 = OpLoad %uint %22 Aligned 4
%24 = OpSConvert %ulong %20
%25 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %rhs %24
%26 = OpLoad %uint %25 Aligned 4
%27 = OpBitwiseAnd %uint %26 %uint_31
%8 = OpShiftLeftLogical %uint %23 %27
%28 = OpSConvert %ulong %20
%29 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %out %28
OpStore %29 %8 Aligned 4
OpReturn
OpFunctionEnd

47
test_conformance/spirv_new/spirv_asm/spv1.4/no_integer_wrap_decoration_fshiftleft_uint.spvasm32 Normal file
View File

@@ -0,0 +1,47 @@
; SPIR-V
; Version: 1.4
; Generator: Khronos SPIR-V Tools Assembler; 0
; Bound: 25
; Schema: 0
OpCapability Addresses
OpCapability Linkage
OpCapability Kernel
OpMemoryModel Physical32 OpenCL
OpEntryPoint Kernel %2 "fmath_cl" %__spirv_BuiltInGlobalInvocationId
OpSource OpenCL_C 200000
OpName %__spirv_BuiltInGlobalInvocationId "__spirv_BuiltInGlobalInvocationId"
OpName %out "out"
OpName %lhs "lhs"
OpName %rhs "rhs"
OpDecorate %7 FuncParamAttr NoWrite
%7 = OpDecorationGroup
OpDecorate %__spirv_BuiltInGlobalInvocationId BuiltIn GlobalInvocationId
OpDecorate %__spirv_BuiltInGlobalInvocationId Constant
OpDecorate %__spirv_BuiltInGlobalInvocationId LinkageAttributes "__spirv_BuiltInGlobalInvocationId" Import
OpDecorate %8 NoUnsignedWrap
OpGroupDecorate %7 %lhs %rhs
%uint = OpTypeInt 32 0
%uint_31 = OpConstant %uint 31
%v3uint = OpTypeVector %uint 3
%_ptr_Input_v3uint = OpTypePointer Input %v3uint
%void = OpTypeVoid
%_ptr_CrossWorkgroup_uint = OpTypePointer CrossWorkgroup %uint
%15 = OpTypeFunction %void %_ptr_CrossWorkgroup_uint %_ptr_CrossWorkgroup_uint %_ptr_CrossWorkgroup_uint
%__spirv_BuiltInGlobalInvocationId = OpVariable %_ptr_Input_v3uint Input
%2 = OpFunction %void None %15
%out = OpFunctionParameter %_ptr_CrossWorkgroup_uint
%lhs = OpFunctionParameter %_ptr_CrossWorkgroup_uint
%rhs = OpFunctionParameter %_ptr_CrossWorkgroup_uint
%16 = OpLabel
%17 = OpLoad %v3uint %__spirv_BuiltInGlobalInvocationId
%18 = OpCompositeExtract %uint %17 0
%19 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %lhs %18
%20 = OpLoad %uint %19 Aligned 4
%21 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %rhs %18
%22 = OpLoad %uint %21 Aligned 4
%23 = OpBitwiseAnd %uint %22 %uint_31
%8 = OpShiftLeftLogical %uint %20 %23
%24 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %out %18
OpStore %24 %8 Aligned 4
OpReturn
OpFunctionEnd

53
test_conformance/spirv_new/spirv_asm/spv1.4/no_integer_wrap_decoration_fshiftleft_uint.spvasm64 Normal file
View File

@@ -0,0 +1,53 @@
; SPIR-V
; Version: 1.4
; Generator: Khronos SPIR-V Tools Assembler; 0
; Bound: 30
; Schema: 0
OpCapability Addresses
OpCapability Linkage
OpCapability Kernel
OpCapability Int64
OpMemoryModel Physical64 OpenCL
OpEntryPoint Kernel %2 "fmath_cl" %__spirv_BuiltInGlobalInvocationId
OpSource OpenCL_C 200000
OpName %__spirv_BuiltInGlobalInvocationId "__spirv_BuiltInGlobalInvocationId"
OpName %out "out"
OpName %lhs "lhs"
OpName %rhs "rhs"
OpDecorate %7 FuncParamAttr NoWrite
%7 = OpDecorationGroup
OpDecorate %__spirv_BuiltInGlobalInvocationId BuiltIn GlobalInvocationId
OpDecorate %__spirv_BuiltInGlobalInvocationId Constant
OpDecorate %__spirv_BuiltInGlobalInvocationId LinkageAttributes "__spirv_BuiltInGlobalInvocationId" Import
OpDecorate %8 NoUnsignedWrap
OpGroupDecorate %7 %lhs %rhs
%ulong = OpTypeInt 64 0
%uint = OpTypeInt 32 0
%uint_31 = OpConstant %uint 31
%v3ulong = OpTypeVector %ulong 3
%_ptr_Input_v3ulong = OpTypePointer Input %v3ulong
%void = OpTypeVoid
%_ptr_CrossWorkgroup_uint = OpTypePointer CrossWorkgroup %uint
%16 = OpTypeFunction %void %_ptr_CrossWorkgroup_uint %_ptr_CrossWorkgroup_uint %_ptr_CrossWorkgroup_uint
%__spirv_BuiltInGlobalInvocationId = OpVariable %_ptr_Input_v3ulong Input
%2 = OpFunction %void None %16
%out = OpFunctionParameter %_ptr_CrossWorkgroup_uint
%lhs = OpFunctionParameter %_ptr_CrossWorkgroup_uint
%rhs = OpFunctionParameter %_ptr_CrossWorkgroup_uint
%17 = OpLabel
%18 = OpLoad %v3ulong %__spirv_BuiltInGlobalInvocationId
%19 = OpCompositeExtract %ulong %18 0
%20 = OpUConvert %uint %19
%21 = OpSConvert %ulong %20
%22 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %lhs %21
%23 = OpLoad %uint %22 Aligned 4
%24 = OpSConvert %ulong %20
%25 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %rhs %24
%26 = OpLoad %uint %25 Aligned 4
%27 = OpBitwiseAnd %uint %26 %uint_31
%8 = OpShiftLeftLogical %uint %23 %27
%28 = OpSConvert %ulong %20
%29 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %out %28
OpStore %29 %8 Aligned 4
OpReturn
OpFunctionEnd

45
test_conformance/spirv_new/spirv_asm/spv1.4/no_integer_wrap_decoration_fsub_int.spvasm32 Normal file
View File

@@ -0,0 +1,45 @@
; SPIR-V
; Version: 1.4
; Generator: Khronos SPIR-V Tools Assembler; 0
; Bound: 23
; Schema: 0
OpCapability Addresses
OpCapability Linkage
OpCapability Kernel
OpMemoryModel Physical32 OpenCL
OpEntryPoint Kernel %2 "fmath_cl" %__spirv_BuiltInGlobalInvocationId
OpSource OpenCL_C 200000
OpName %__spirv_BuiltInGlobalInvocationId "__spirv_BuiltInGlobalInvocationId"
OpName %out "out"
OpName %lhs "lhs"
OpName %rhs "rhs"
OpDecorate %7 FuncParamAttr NoWrite
%7 = OpDecorationGroup
OpDecorate %__spirv_BuiltInGlobalInvocationId BuiltIn GlobalInvocationId
OpDecorate %__spirv_BuiltInGlobalInvocationId Constant
OpDecorate %__spirv_BuiltInGlobalInvocationId LinkageAttributes "__spirv_BuiltInGlobalInvocationId" Import
OpGroupDecorate %7 %lhs %rhs
OpDecorate %8 NoSignedWrap
%uint = OpTypeInt 32 0
%v3uint = OpTypeVector %uint 3
%_ptr_Input_v3uint = OpTypePointer Input %v3uint
%void = OpTypeVoid
%_ptr_CrossWorkgroup_uint = OpTypePointer CrossWorkgroup %uint
%14 = OpTypeFunction %void %_ptr_CrossWorkgroup_uint %_ptr_CrossWorkgroup_uint %_ptr_CrossWorkgroup_uint
%__spirv_BuiltInGlobalInvocationId = OpVariable %_ptr_Input_v3uint Input
%2 = OpFunction %void None %14
%out = OpFunctionParameter %_ptr_CrossWorkgroup_uint
%lhs = OpFunctionParameter %_ptr_CrossWorkgroup_uint
%rhs = OpFunctionParameter %_ptr_CrossWorkgroup_uint
%15 = OpLabel
%16 = OpLoad %v3uint %__spirv_BuiltInGlobalInvocationId
%17 = OpCompositeExtract %uint %16 0
%18 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %lhs %17
%19 = OpLoad %uint %18 Aligned 4
%20 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %rhs %17
%21 = OpLoad %uint %20 Aligned 4
%8 = OpISub %uint %19 %21
%22 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %out %17
OpStore %22 %8 Aligned 4
OpReturn
OpFunctionEnd

51
test_conformance/spirv_new/spirv_asm/spv1.4/no_integer_wrap_decoration_fsub_int.spvasm64 Normal file
View File

@@ -0,0 +1,51 @@
; SPIR-V
; Version: 1.4
; Generator: Khronos SPIR-V Tools Assembler; 0
; Bound: 28
; Schema: 0
OpCapability Addresses
OpCapability Linkage
OpCapability Kernel
OpCapability Int64
OpMemoryModel Physical64 OpenCL
OpEntryPoint Kernel %2 "fmath_cl" %__spirv_BuiltInGlobalInvocationId
OpSource OpenCL_C 200000
OpName %__spirv_BuiltInGlobalInvocationId "__spirv_BuiltInGlobalInvocationId"
OpName %out "out"
OpName %lhs "lhs"
OpName %rhs "rhs"
OpDecorate %7 FuncParamAttr NoWrite
%7 = OpDecorationGroup
OpDecorate %__spirv_BuiltInGlobalInvocationId BuiltIn GlobalInvocationId
OpDecorate %__spirv_BuiltInGlobalInvocationId Constant
OpDecorate %__spirv_BuiltInGlobalInvocationId LinkageAttributes "__spirv_BuiltInGlobalInvocationId" Import
OpGroupDecorate %7 %lhs %rhs
OpDecorate %8 NoSignedWrap
%ulong = OpTypeInt 64 0
%uint = OpTypeInt 32 0
%v3ulong = OpTypeVector %ulong 3
%_ptr_Input_v3ulong = OpTypePointer Input %v3ulong
%void = OpTypeVoid
%_ptr_CrossWorkgroup_uint = OpTypePointer CrossWorkgroup %uint
%15 = OpTypeFunction %void %_ptr_CrossWorkgroup_uint %_ptr_CrossWorkgroup_uint %_ptr_CrossWorkgroup_uint
%__spirv_BuiltInGlobalInvocationId = OpVariable %_ptr_Input_v3ulong Input
%2 = OpFunction %void None %15
%out = OpFunctionParameter %_ptr_CrossWorkgroup_uint
%lhs = OpFunctionParameter %_ptr_CrossWorkgroup_uint
%rhs = OpFunctionParameter %_ptr_CrossWorkgroup_uint
%16 = OpLabel
%17 = OpLoad %v3ulong %__spirv_BuiltInGlobalInvocationId
%18 = OpCompositeExtract %ulong %17 0
%19 = OpUConvert %uint %18
%20 = OpSConvert %ulong %19
%21 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %lhs %20
%22 = OpLoad %uint %21 Aligned 4
%23 = OpSConvert %ulong %19
%24 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %rhs %23
%25 = OpLoad %uint %24 Aligned 4
%8 = OpISub %uint %22 %25
%26 = OpSConvert %ulong %19
%27 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %out %26
OpStore %27 %8 Aligned 4
OpReturn
OpFunctionEnd

45
test_conformance/spirv_new/spirv_asm/spv1.4/no_integer_wrap_decoration_fsub_uint.spvasm32 Normal file
View File

@@ -0,0 +1,45 @@
; SPIR-V
; Version: 1.4
; Generator: Khronos SPIR-V Tools Assembler; 0
; Bound: 23
; Schema: 0
OpCapability Addresses
OpCapability Linkage
OpCapability Kernel
OpMemoryModel Physical32 OpenCL
OpEntryPoint Kernel %2 "fmath_cl" %__spirv_BuiltInGlobalInvocationId
OpSource OpenCL_C 200000
OpName %__spirv_BuiltInGlobalInvocationId "__spirv_BuiltInGlobalInvocationId"
OpName %out "out"
OpName %lhs "lhs"
OpName %rhs "rhs"
OpDecorate %7 FuncParamAttr NoWrite
%7 = OpDecorationGroup
OpDecorate %__spirv_BuiltInGlobalInvocationId BuiltIn GlobalInvocationId
OpDecorate %__spirv_BuiltInGlobalInvocationId Constant
OpDecorate %__spirv_BuiltInGlobalInvocationId LinkageAttributes "__spirv_BuiltInGlobalInvocationId" Import
OpGroupDecorate %7 %lhs %rhs
OpDecorate %8 NoUnsignedWrap
%uint = OpTypeInt 32 0
%v3uint = OpTypeVector %uint 3
%_ptr_Input_v3uint = OpTypePointer Input %v3uint
%void = OpTypeVoid
%_ptr_CrossWorkgroup_uint = OpTypePointer CrossWorkgroup %uint
%14 = OpTypeFunction %void %_ptr_CrossWorkgroup_uint %_ptr_CrossWorkgroup_uint %_ptr_CrossWorkgroup_uint
%__spirv_BuiltInGlobalInvocationId = OpVariable %_ptr_Input_v3uint Input
%2 = OpFunction %void None %14
%out = OpFunctionParameter %_ptr_CrossWorkgroup_uint
%lhs = OpFunctionParameter %_ptr_CrossWorkgroup_uint
%rhs = OpFunctionParameter %_ptr_CrossWorkgroup_uint
%15 = OpLabel
%16 = OpLoad %v3uint %__spirv_BuiltInGlobalInvocationId
%17 = OpCompositeExtract %uint %16 0
%18 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %lhs %17
%19 = OpLoad %uint %18 Aligned 4
%20 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %rhs %17
%21 = OpLoad %uint %20 Aligned 4
%8 = OpISub %uint %19 %21
%22 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %out %17
OpStore %22 %8 Aligned 4
OpReturn
OpFunctionEnd

51
test_conformance/spirv_new/spirv_asm/spv1.4/no_integer_wrap_decoration_fsub_uint.spvasm64 Normal file
View File

@@ -0,0 +1,51 @@
; SPIR-V
; Version: 1.4
; Generator: Khronos SPIR-V Tools Assembler; 0
; Bound: 28
; Schema: 0
OpCapability Addresses
OpCapability Linkage
OpCapability Kernel
OpCapability Int64
OpMemoryModel Physical64 OpenCL
OpEntryPoint Kernel %2 "fmath_cl" %__spirv_BuiltInGlobalInvocationId
OpSource OpenCL_C 200000
OpName %__spirv_BuiltInGlobalInvocationId "__spirv_BuiltInGlobalInvocationId"
OpName %out "out"
OpName %lhs "lhs"
OpName %rhs "rhs"
OpDecorate %7 FuncParamAttr NoWrite
%7 = OpDecorationGroup
OpDecorate %__spirv_BuiltInGlobalInvocationId BuiltIn GlobalInvocationId
OpDecorate %__spirv_BuiltInGlobalInvocationId Constant
OpDecorate %__spirv_BuiltInGlobalInvocationId LinkageAttributes "__spirv_BuiltInGlobalInvocationId" Import
OpGroupDecorate %7 %lhs %rhs
OpDecorate %8 NoUnsignedWrap
%ulong = OpTypeInt 64 0
%uint = OpTypeInt 32 0
%v3ulong = OpTypeVector %ulong 3
%_ptr_Input_v3ulong = OpTypePointer Input %v3ulong
%void = OpTypeVoid
%_ptr_CrossWorkgroup_uint = OpTypePointer CrossWorkgroup %uint
%15 = OpTypeFunction %void %_ptr_CrossWorkgroup_uint %_ptr_CrossWorkgroup_uint %_ptr_CrossWorkgroup_uint
%__spirv_BuiltInGlobalInvocationId = OpVariable %_ptr_Input_v3ulong Input
%2 = OpFunction %void None %15
%out = OpFunctionParameter %_ptr_CrossWorkgroup_uint
%lhs = OpFunctionParameter %_ptr_CrossWorkgroup_uint
%rhs = OpFunctionParameter %_ptr_CrossWorkgroup_uint
%16 = OpLabel
%17 = OpLoad %v3ulong %__spirv_BuiltInGlobalInvocationId
%18 = OpCompositeExtract %ulong %17 0
%19 = OpUConvert %uint %18
%20 = OpSConvert %ulong %19
%21 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %lhs %20
%22 = OpLoad %uint %21 Aligned 4
%23 = OpSConvert %ulong %19
%24 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %rhs %23
%25 = OpLoad %uint %24 Aligned 4
%8 = OpISub %uint %22 %25
%26 = OpSConvert %ulong %19
%27 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %out %26
OpStore %27 %8 Aligned 4
OpReturn
OpFunctionEnd

166
test_conformance/spirv_new/test_cl_khr_spirv_no_integer_wrap_decoration.cpp → test_conformance/spirv_new/test_no_integer_wrap_decoration.cpp
View File

@@ -15,6 +15,7 @@
// //
#include "testBase.h" #include "testBase.h"
#include "spirvInfo.hpp"
#include "types.hpp" #include "types.hpp"
#include <sstream> #include <sstream>
@@ -22,13 +23,10 @@
#include <type_traits> #include <type_traits>
template<typename T> template <typename T>
int test_ext_cl_khr_spirv_no_integer_wrap_decoration(cl_device_id deviceID, int test_no_integer_wrap_decoration(cl_device_id deviceID, cl_context context,
cl_context context, cl_command_queue queue, const char *spvName,
cl_command_queue queue, const char *funcName, const char *Tname)
const char *spvName,
const char *funcName,
const char *Tname)
{ {
cl_int err = CL_SUCCESS; cl_int err = CL_SUCCESS;
@@ -37,13 +35,10 @@ int test_ext_cl_khr_spirv_no_integer_wrap_decoration(cl_device_id deviceID,
std::vector<T> h_rhs(num); std::vector<T> h_rhs(num);
std::vector<T> expected_results(num); std::vector<T> expected_results(num);
std::vector<T> h_ref(num); std::vector<T> h_ref(num);
if (!is_extension_available(deviceID, "cl_khr_spirv_no_integer_wrap_decoration")) {
log_info("Extension cl_khr_spirv_no_integer_wrap_decoration not supported; skipping tests.\n");
return 0;
}
/*Test with some values that do not cause overflow*/ /*Test with some values that do not cause overflow*/
if (std::is_signed<T>::value == true) { if (std::is_signed<T>::value == true)
{
h_lhs.push_back((T)-25000); h_lhs.push_back((T)-25000);
h_lhs.push_back((T)-3333); h_lhs.push_back((T)-3333);
h_lhs.push_back((T)-7); h_lhs.push_back((T)-7);
@@ -54,7 +49,9 @@ int test_ext_cl_khr_spirv_no_integer_wrap_decoration(cl_device_id deviceID,
h_lhs.push_back(2048); h_lhs.push_back(2048);
h_lhs.push_back(4094); h_lhs.push_back(4094);
h_lhs.push_back(10000); h_lhs.push_back(10000);
} else { }
else
{
h_lhs.push_back(0); h_lhs.push_back(0);
h_lhs.push_back(1); h_lhs.push_back(1);
h_lhs.push_back(3); h_lhs.push_back(3);
@@ -79,16 +76,20 @@ int test_ext_cl_khr_spirv_no_integer_wrap_decoration(cl_device_id deviceID,
h_rhs.push_back(9); h_rhs.push_back(9);
size_t bytes = num * sizeof(T); size_t bytes = num * sizeof(T);
clMemWrapper lhs = clCreateBuffer(context, CL_MEM_READ_ONLY, bytes, NULL, &err); clMemWrapper lhs =
clCreateBuffer(context, CL_MEM_READ_ONLY, bytes, NULL, &err);
SPIRV_CHECK_ERROR(err, "Failed to create lhs buffer"); SPIRV_CHECK_ERROR(err, "Failed to create lhs buffer");
err = clEnqueueWriteBuffer(queue, lhs, CL_TRUE, 0, bytes, &h_lhs[0], 0, NULL, NULL); err = clEnqueueWriteBuffer(queue, lhs, CL_TRUE, 0, bytes, &h_lhs[0], 0,
NULL, NULL);
SPIRV_CHECK_ERROR(err, "Failed to copy to lhs buffer"); SPIRV_CHECK_ERROR(err, "Failed to copy to lhs buffer");
clMemWrapper rhs = clCreateBuffer(context, CL_MEM_READ_ONLY, bytes, NULL, &err); clMemWrapper rhs =
clCreateBuffer(context, CL_MEM_READ_ONLY, bytes, NULL, &err);
SPIRV_CHECK_ERROR(err, "Failed to create rhs buffer"); SPIRV_CHECK_ERROR(err, "Failed to create rhs buffer");
err = clEnqueueWriteBuffer(queue, rhs, CL_TRUE, 0, bytes, &h_rhs[0], 0, NULL, NULL); err = clEnqueueWriteBuffer(queue, rhs, CL_TRUE, 0, bytes, &h_rhs[0], 0,
NULL, NULL);
SPIRV_CHECK_ERROR(err, "Failed to copy to rhs buffer"); SPIRV_CHECK_ERROR(err, "Failed to copy to rhs buffer");
std::string kernelStr; std::string kernelStr;
@@ -101,8 +102,8 @@ int test_ext_cl_khr_spirv_no_integer_wrap_decoration(cl_device_id deviceID,
kernelStream << "#define spirv_fshiftleft(a, b) (a) << (b) \n"; kernelStream << "#define spirv_fshiftleft(a, b) (a) << (b) \n";
kernelStream << "#define spirv_fnegate(a, b) (-a) \n"; kernelStream << "#define spirv_fnegate(a, b) (-a) \n";
kernelStream << "#define T " << Tname << "\n"; kernelStream << "#define T " << Tname << "\n";
kernelStream << "#define FUNC spirv_" << funcName << "\n"; kernelStream << "#define FUNC spirv_" << funcName << "\n";
kernelStream << "__kernel void fmath_cl(__global T *out, \n"; kernelStream << "__kernel void fmath_cl(__global T *out, \n";
kernelStream << "const __global T *lhs, const __global T *rhs) \n"; kernelStream << "const __global T *lhs, const __global T *rhs) \n";
kernelStream << "{ \n"; kernelStream << "{ \n";
@@ -114,16 +115,26 @@ int test_ext_cl_khr_spirv_no_integer_wrap_decoration(cl_device_id deviceID,
const char *kernelBuf = kernelStr.c_str(); const char *kernelBuf = kernelStr.c_str();
for (int i = 0; i < num; i++) { for (int i = 0; i < num; i++)
if (std::string(funcName) == std::string("fadd")) { {
if (std::string(funcName) == std::string("fadd"))
{
expected_results[i] = h_lhs[i] + h_rhs[i]; expected_results[i] = h_lhs[i] + h_rhs[i];
} else if (std::string(funcName) == std::string("fsub")) { }
else if (std::string(funcName) == std::string("fsub"))
{
expected_results[i] = h_lhs[i] - h_rhs[i]; expected_results[i] = h_lhs[i] - h_rhs[i];
} else if (std::string(funcName) == std::string("fmul")) { }
else if (std::string(funcName) == std::string("fmul"))
{
expected_results[i] = h_lhs[i] * h_rhs[i]; expected_results[i] = h_lhs[i] * h_rhs[i];
} else if (std::string(funcName) == std::string("fshiftleft")) { }
else if (std::string(funcName) == std::string("fshiftleft"))
{
expected_results[i] = h_lhs[i] << h_rhs[i]; expected_results[i] = h_lhs[i] << h_rhs[i];
} else if (std::string(funcName) == std::string("fnegate")) { }
else if (std::string(funcName) == std::string("fnegate"))
{
expected_results[i] = 0 - h_lhs[i]; expected_results[i] = 0 - h_lhs[i];
} }
} }
@@ -136,7 +147,8 @@ int test_ext_cl_khr_spirv_no_integer_wrap_decoration(cl_device_id deviceID,
&kernelBuf, "fmath_cl"); &kernelBuf, "fmath_cl");
SPIRV_CHECK_ERROR(err, "Failed to create cl kernel"); SPIRV_CHECK_ERROR(err, "Failed to create cl kernel");
clMemWrapper ref = clCreateBuffer(context, CL_MEM_READ_WRITE, bytes, NULL, &err); clMemWrapper ref =
clCreateBuffer(context, CL_MEM_READ_WRITE, bytes, NULL, &err);
SPIRV_CHECK_ERROR(err, "Failed to create ref buffer"); SPIRV_CHECK_ERROR(err, "Failed to create ref buffer");
err = clSetKernelArg(kernel, 0, sizeof(cl_mem), &ref); err = clSetKernelArg(kernel, 0, sizeof(cl_mem), &ref);
@@ -149,16 +161,22 @@ int test_ext_cl_khr_spirv_no_integer_wrap_decoration(cl_device_id deviceID,
SPIRV_CHECK_ERROR(err, "Failed to set arg 2"); SPIRV_CHECK_ERROR(err, "Failed to set arg 2");
size_t global = num; size_t global = num;
err = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &global, NULL, 0, NULL, NULL); err = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &global, NULL, 0,
NULL, NULL);
SPIRV_CHECK_ERROR(err, "Failed to enqueue cl kernel"); SPIRV_CHECK_ERROR(err, "Failed to enqueue cl kernel");
err = clEnqueueReadBuffer(queue, ref, CL_TRUE, 0, bytes, &h_ref[0], 0, NULL, NULL); err = clEnqueueReadBuffer(queue, ref, CL_TRUE, 0, bytes, &h_ref[0], 0,
NULL, NULL);
SPIRV_CHECK_ERROR(err, "Failed to read from ref"); SPIRV_CHECK_ERROR(err, "Failed to read from ref");
} }
for (int i = 0; i < num; i++) { for (int i = 0; i < num; i++)
if (expected_results[i] != h_ref[i]) { {
log_error("Values do not match at index %d expected = %d got = %d\n", i, expected_results[i], h_ref[i]); if (expected_results[i] != h_ref[i])
{
log_error(
"Values do not match at index %d expected = %d got = %d\n", i,
expected_results[i], h_ref[i]);
return -1; return -1;
} }
} }
@@ -170,7 +188,8 @@ int test_ext_cl_khr_spirv_no_integer_wrap_decoration(cl_device_id deviceID,
clKernelWrapper kernel = clCreateKernel(prog, "fmath_cl", &err); clKernelWrapper kernel = clCreateKernel(prog, "fmath_cl", &err);
SPIRV_CHECK_ERROR(err, "Failed to create spv kernel"); SPIRV_CHECK_ERROR(err, "Failed to create spv kernel");
clMemWrapper res = clCreateBuffer(context, CL_MEM_READ_WRITE, bytes, NULL, &err); clMemWrapper res =
clCreateBuffer(context, CL_MEM_READ_WRITE, bytes, NULL, &err);
SPIRV_CHECK_ERROR(err, "Failed to create res buffer"); SPIRV_CHECK_ERROR(err, "Failed to create res buffer");
err = clSetKernelArg(kernel, 0, sizeof(cl_mem), &res); err = clSetKernelArg(kernel, 0, sizeof(cl_mem), &res);
@@ -183,39 +202,76 @@ int test_ext_cl_khr_spirv_no_integer_wrap_decoration(cl_device_id deviceID,
SPIRV_CHECK_ERROR(err, "Failed to set arg 2"); SPIRV_CHECK_ERROR(err, "Failed to set arg 2");
size_t global = num; size_t global = num;
err = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &global, NULL, 0, NULL, NULL); err = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &global, NULL, 0, NULL,
NULL);
SPIRV_CHECK_ERROR(err, "Failed to enqueue cl kernel"); SPIRV_CHECK_ERROR(err, "Failed to enqueue cl kernel");
std::vector<T> h_res(num); std::vector<T> h_res(num);
err = clEnqueueReadBuffer(queue, res, CL_TRUE, 0, bytes, &h_res[0], 0, NULL, NULL); err = clEnqueueReadBuffer(queue, res, CL_TRUE, 0, bytes, &h_res[0], 0, NULL,
NULL);
SPIRV_CHECK_ERROR(err, "Failed to read from ref"); SPIRV_CHECK_ERROR(err, "Failed to read from ref");
for (int i = 0; i < num; i++) { for (int i = 0; i < num; i++)
if (expected_results[i] != h_res[i]) { {
log_error("Values do not match at location %d expected = %d got = %d\n", i, expected_results[i], h_res[i]); if (expected_results[i] != h_res[i])
{
log_error(
"Values do not match at location %d expected = %d got = %d\n",
i, expected_results[i], h_res[i]);
return -1; return -1;
} }
} }
return 0; return TEST_PASS;
} }
#define TEST_FMATH_FUNC(TYPE, FUNC) \ #define TEST_FMATH_FUNC_KHR(TYPE, FUNC) \
TEST_SPIRV_FUNC(ext_cl_khr_spirv_no_integer_wrap_decoration_##FUNC##_##TYPE) \ TEST_SPIRV_FUNC( \
{ \ ext_cl_khr_spirv_no_integer_wrap_decoration_##FUNC##_##TYPE) \
return test_ext_cl_khr_spirv_no_integer_wrap_decoration<cl_##TYPE>(deviceID, context, queue, \ { \
"ext_cl_khr_spirv_no_integer_wrap_decoration_"#FUNC"_"#TYPE, \ if (!is_extension_available( \
#FUNC, \ deviceID, "cl_khr_spirv_no_integer_wrap_decoration")) \
#TYPE \ { \
); \ log_info("Extension cl_khr_spirv_no_integer_wrap_decoration not " \
"supported; skipping tests.\n"); \
return TEST_SKIPPED_ITSELF; \
} \
return test_no_integer_wrap_decoration<cl_##TYPE>( \
deviceID, context, queue, \
"ext_cl_khr_spirv_no_integer_wrap_decoration_" #FUNC "_" #TYPE, \
#FUNC, #TYPE); \
} }
TEST_FMATH_FUNC(int, fadd) TEST_FMATH_FUNC_KHR(int, fadd)
TEST_FMATH_FUNC(int, fsub) TEST_FMATH_FUNC_KHR(int, fsub)
TEST_FMATH_FUNC(int, fmul) TEST_FMATH_FUNC_KHR(int, fmul)
TEST_FMATH_FUNC(int, fshiftleft) TEST_FMATH_FUNC_KHR(int, fshiftleft)
TEST_FMATH_FUNC(int, fnegate) TEST_FMATH_FUNC_KHR(int, fnegate)
TEST_FMATH_FUNC(uint, fadd) TEST_FMATH_FUNC_KHR(uint, fadd)
TEST_FMATH_FUNC(uint, fsub) TEST_FMATH_FUNC_KHR(uint, fsub)
TEST_FMATH_FUNC(uint, fmul) TEST_FMATH_FUNC_KHR(uint, fmul)
TEST_FMATH_FUNC(uint, fshiftleft) TEST_FMATH_FUNC_KHR(uint, fshiftleft)
#define TEST_FMATH_FUNC_14(TYPE, FUNC) \
TEST_SPIRV_FUNC(spirv14_no_integer_wrap_decoration_##FUNC##_##TYPE) \
{ \
if (!is_spirv_version_supported(deviceID, "SPIR-V_1.4")) \
{ \
log_info("SPIR-V 1.4 not supported; skipping tests.\n"); \
return TEST_SKIPPED_ITSELF; \
} \
return test_no_integer_wrap_decoration<cl_##TYPE>( \
deviceID, context, queue, \
"spv1.4/no_integer_wrap_decoration_" #FUNC "_" #TYPE, #FUNC, \
#TYPE); \
}
TEST_FMATH_FUNC_14(int, fadd)
TEST_FMATH_FUNC_14(int, fsub)
TEST_FMATH_FUNC_14(int, fmul)
TEST_FMATH_FUNC_14(int, fshiftleft)
TEST_FMATH_FUNC_14(int, fnegate)
TEST_FMATH_FUNC_14(uint, fadd)
TEST_FMATH_FUNC_14(uint, fsub)
TEST_FMATH_FUNC_14(uint, fmul)
TEST_FMATH_FUNC_14(uint, fshiftleft)