diff --git a/test_conformance/spirv_new/CMakeLists.txt b/test_conformance/spirv_new/CMakeLists.txt index 66219208..67faecf8 100644 --- a/test_conformance/spirv_new/CMakeLists.txt +++ b/test_conformance/spirv_new/CMakeLists.txt @@ -4,10 +4,10 @@ set(${MODULE_NAME}_SOURCES main.cpp test_basic_versions.cpp test_cl_khr_expect_assume.cpp - test_cl_khr_spirv_no_integer_wrap_decoration.cpp test_decorate.cpp test_get_program_il.cpp test_linkage.cpp + test_no_integer_wrap_decoration.cpp test_op_atomic.cpp test_op_branch_conditional.cpp test_op_branch.cpp diff --git a/test_conformance/spirv_new/spirv_asm/spv1.4/no_integer_wrap_decoration_fadd_int.spvasm32 b/test_conformance/spirv_new/spirv_asm/spv1.4/no_integer_wrap_decoration_fadd_int.spvasm32 new file mode 100644 index 00000000..557faae0 --- /dev/null +++ b/test_conformance/spirv_new/spirv_asm/spv1.4/no_integer_wrap_decoration_fadd_int.spvasm32 @@ -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 diff --git a/test_conformance/spirv_new/spirv_asm/spv1.4/no_integer_wrap_decoration_fadd_int.spvasm64 b/test_conformance/spirv_new/spirv_asm/spv1.4/no_integer_wrap_decoration_fadd_int.spvasm64 new file mode 100644 index 00000000..4bc2c2a7 --- /dev/null +++ b/test_conformance/spirv_new/spirv_asm/spv1.4/no_integer_wrap_decoration_fadd_int.spvasm64 @@ -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 diff --git a/test_conformance/spirv_new/spirv_asm/spv1.4/no_integer_wrap_decoration_fadd_uint.spvasm32 b/test_conformance/spirv_new/spirv_asm/spv1.4/no_integer_wrap_decoration_fadd_uint.spvasm32 new file mode 100644 index 00000000..b858523c --- /dev/null +++ b/test_conformance/spirv_new/spirv_asm/spv1.4/no_integer_wrap_decoration_fadd_uint.spvasm32 @@ -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 diff --git a/test_conformance/spirv_new/spirv_asm/spv1.4/no_integer_wrap_decoration_fadd_uint.spvasm64 b/test_conformance/spirv_new/spirv_asm/spv1.4/no_integer_wrap_decoration_fadd_uint.spvasm64 new file mode 100644 index 00000000..a41f4f41 --- /dev/null +++ b/test_conformance/spirv_new/spirv_asm/spv1.4/no_integer_wrap_decoration_fadd_uint.spvasm64 @@ -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 diff --git a/test_conformance/spirv_new/spirv_asm/spv1.4/no_integer_wrap_decoration_fmul_int.spvasm32 b/test_conformance/spirv_new/spirv_asm/spv1.4/no_integer_wrap_decoration_fmul_int.spvasm32 new file mode 100644 index 00000000..4886c01c --- /dev/null +++ b/test_conformance/spirv_new/spirv_asm/spv1.4/no_integer_wrap_decoration_fmul_int.spvasm32 @@ -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 diff --git a/test_conformance/spirv_new/spirv_asm/spv1.4/no_integer_wrap_decoration_fmul_int.spvasm64 b/test_conformance/spirv_new/spirv_asm/spv1.4/no_integer_wrap_decoration_fmul_int.spvasm64 new file mode 100644 index 00000000..ff6799e7 --- /dev/null +++ b/test_conformance/spirv_new/spirv_asm/spv1.4/no_integer_wrap_decoration_fmul_int.spvasm64 @@ -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 diff --git a/test_conformance/spirv_new/spirv_asm/spv1.4/no_integer_wrap_decoration_fmul_uint.spvasm32 b/test_conformance/spirv_new/spirv_asm/spv1.4/no_integer_wrap_decoration_fmul_uint.spvasm32 new file mode 100644 index 00000000..1993a13d --- /dev/null +++ b/test_conformance/spirv_new/spirv_asm/spv1.4/no_integer_wrap_decoration_fmul_uint.spvasm32 @@ -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 diff --git a/test_conformance/spirv_new/spirv_asm/spv1.4/no_integer_wrap_decoration_fmul_uint.spvasm64 b/test_conformance/spirv_new/spirv_asm/spv1.4/no_integer_wrap_decoration_fmul_uint.spvasm64 new file mode 100644 index 00000000..d17b9c29 --- /dev/null +++ b/test_conformance/spirv_new/spirv_asm/spv1.4/no_integer_wrap_decoration_fmul_uint.spvasm64 @@ -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 diff --git a/test_conformance/spirv_new/spirv_asm/spv1.4/no_integer_wrap_decoration_fnegate_int.spvasm32 b/test_conformance/spirv_new/spirv_asm/spv1.4/no_integer_wrap_decoration_fnegate_int.spvasm32 new file mode 100644 index 00000000..ea9a85ad --- /dev/null +++ b/test_conformance/spirv_new/spirv_asm/spv1.4/no_integer_wrap_decoration_fnegate_int.spvasm32 @@ -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 diff --git a/test_conformance/spirv_new/spirv_asm/spv1.4/no_integer_wrap_decoration_fnegate_int.spvasm64 b/test_conformance/spirv_new/spirv_asm/spv1.4/no_integer_wrap_decoration_fnegate_int.spvasm64 new file mode 100644 index 00000000..2edfb499 --- /dev/null +++ b/test_conformance/spirv_new/spirv_asm/spv1.4/no_integer_wrap_decoration_fnegate_int.spvasm64 @@ -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 diff --git a/test_conformance/spirv_new/spirv_asm/spv1.4/no_integer_wrap_decoration_fshiftleft_int.spvasm32 b/test_conformance/spirv_new/spirv_asm/spv1.4/no_integer_wrap_decoration_fshiftleft_int.spvasm32 new file mode 100644 index 00000000..08b5ba71 --- /dev/null +++ b/test_conformance/spirv_new/spirv_asm/spv1.4/no_integer_wrap_decoration_fshiftleft_int.spvasm32 @@ -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 diff --git a/test_conformance/spirv_new/spirv_asm/spv1.4/no_integer_wrap_decoration_fshiftleft_int.spvasm64 b/test_conformance/spirv_new/spirv_asm/spv1.4/no_integer_wrap_decoration_fshiftleft_int.spvasm64 new file mode 100644 index 00000000..3487dc46 --- /dev/null +++ b/test_conformance/spirv_new/spirv_asm/spv1.4/no_integer_wrap_decoration_fshiftleft_int.spvasm64 @@ -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 diff --git a/test_conformance/spirv_new/spirv_asm/spv1.4/no_integer_wrap_decoration_fshiftleft_uint.spvasm32 b/test_conformance/spirv_new/spirv_asm/spv1.4/no_integer_wrap_decoration_fshiftleft_uint.spvasm32 new file mode 100644 index 00000000..ac01b182 --- /dev/null +++ b/test_conformance/spirv_new/spirv_asm/spv1.4/no_integer_wrap_decoration_fshiftleft_uint.spvasm32 @@ -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 diff --git a/test_conformance/spirv_new/spirv_asm/spv1.4/no_integer_wrap_decoration_fshiftleft_uint.spvasm64 b/test_conformance/spirv_new/spirv_asm/spv1.4/no_integer_wrap_decoration_fshiftleft_uint.spvasm64 new file mode 100644 index 00000000..a2d8ba69 --- /dev/null +++ b/test_conformance/spirv_new/spirv_asm/spv1.4/no_integer_wrap_decoration_fshiftleft_uint.spvasm64 @@ -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 diff --git a/test_conformance/spirv_new/spirv_asm/spv1.4/no_integer_wrap_decoration_fsub_int.spvasm32 b/test_conformance/spirv_new/spirv_asm/spv1.4/no_integer_wrap_decoration_fsub_int.spvasm32 new file mode 100644 index 00000000..946be211 --- /dev/null +++ b/test_conformance/spirv_new/spirv_asm/spv1.4/no_integer_wrap_decoration_fsub_int.spvasm32 @@ -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 diff --git a/test_conformance/spirv_new/spirv_asm/spv1.4/no_integer_wrap_decoration_fsub_int.spvasm64 b/test_conformance/spirv_new/spirv_asm/spv1.4/no_integer_wrap_decoration_fsub_int.spvasm64 new file mode 100644 index 00000000..26ce619e --- /dev/null +++ b/test_conformance/spirv_new/spirv_asm/spv1.4/no_integer_wrap_decoration_fsub_int.spvasm64 @@ -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 diff --git a/test_conformance/spirv_new/spirv_asm/spv1.4/no_integer_wrap_decoration_fsub_uint.spvasm32 b/test_conformance/spirv_new/spirv_asm/spv1.4/no_integer_wrap_decoration_fsub_uint.spvasm32 new file mode 100644 index 00000000..013614f3 --- /dev/null +++ b/test_conformance/spirv_new/spirv_asm/spv1.4/no_integer_wrap_decoration_fsub_uint.spvasm32 @@ -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 diff --git a/test_conformance/spirv_new/spirv_asm/spv1.4/no_integer_wrap_decoration_fsub_uint.spvasm64 b/test_conformance/spirv_new/spirv_asm/spv1.4/no_integer_wrap_decoration_fsub_uint.spvasm64 new file mode 100644 index 00000000..19df8b5f --- /dev/null +++ b/test_conformance/spirv_new/spirv_asm/spv1.4/no_integer_wrap_decoration_fsub_uint.spvasm64 @@ -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 diff --git a/test_conformance/spirv_new/test_cl_khr_spirv_no_integer_wrap_decoration.cpp b/test_conformance/spirv_new/test_no_integer_wrap_decoration.cpp similarity index 53% rename from test_conformance/spirv_new/test_cl_khr_spirv_no_integer_wrap_decoration.cpp rename to test_conformance/spirv_new/test_no_integer_wrap_decoration.cpp index 95227b28..042f0757 100644 --- a/test_conformance/spirv_new/test_cl_khr_spirv_no_integer_wrap_decoration.cpp +++ b/test_conformance/spirv_new/test_no_integer_wrap_decoration.cpp @@ -15,6 +15,7 @@ // #include "testBase.h" +#include "spirvInfo.hpp" #include "types.hpp" #include @@ -22,13 +23,10 @@ #include -template -int test_ext_cl_khr_spirv_no_integer_wrap_decoration(cl_device_id deviceID, - cl_context context, - cl_command_queue queue, - const char *spvName, - const char *funcName, - const char *Tname) +template +int test_no_integer_wrap_decoration(cl_device_id deviceID, cl_context context, + cl_command_queue queue, const char *spvName, + const char *funcName, const char *Tname) { 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 h_rhs(num); std::vector expected_results(num); std::vector 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*/ - if (std::is_signed::value == true) { + if (std::is_signed::value == true) + { h_lhs.push_back((T)-25000); h_lhs.push_back((T)-3333); 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(4094); h_lhs.push_back(10000); - } else { + } + else + { h_lhs.push_back(0); h_lhs.push_back(1); 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); 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"); - 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"); - 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"); - 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"); 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_fnegate(a, b) (-a) \n"; - kernelStream << "#define T " << Tname << "\n"; - kernelStream << "#define FUNC spirv_" << funcName << "\n"; + kernelStream << "#define T " << Tname << "\n"; + kernelStream << "#define FUNC spirv_" << funcName << "\n"; kernelStream << "__kernel void fmath_cl(__global T *out, \n"; kernelStream << "const __global T *lhs, const __global T *rhs) \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(); - for (int i = 0; i < num; i++) { - if (std::string(funcName) == std::string("fadd")) { + for (int i = 0; i < num; i++) + { + if (std::string(funcName) == std::string("fadd")) + { 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]; - } 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]; - } 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]; - } else if (std::string(funcName) == std::string("fnegate")) { + } + else if (std::string(funcName) == std::string("fnegate")) + { 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"); 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"); 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"); 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"); - 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"); } - 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]); + 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]); 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); 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"); 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"); 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"); std::vector 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"); - 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]); + 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]); return -1; } } - return 0; + return TEST_PASS; } -#define TEST_FMATH_FUNC(TYPE, FUNC) \ - TEST_SPIRV_FUNC(ext_cl_khr_spirv_no_integer_wrap_decoration_##FUNC##_##TYPE) \ - { \ - return test_ext_cl_khr_spirv_no_integer_wrap_decoration(deviceID, context, queue, \ - "ext_cl_khr_spirv_no_integer_wrap_decoration_"#FUNC"_"#TYPE, \ - #FUNC, \ - #TYPE \ - ); \ +#define TEST_FMATH_FUNC_KHR(TYPE, FUNC) \ + TEST_SPIRV_FUNC( \ + ext_cl_khr_spirv_no_integer_wrap_decoration_##FUNC##_##TYPE) \ + { \ + 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 TEST_SKIPPED_ITSELF; \ + } \ + return test_no_integer_wrap_decoration( \ + deviceID, context, queue, \ + "ext_cl_khr_spirv_no_integer_wrap_decoration_" #FUNC "_" #TYPE, \ + #FUNC, #TYPE); \ } -TEST_FMATH_FUNC(int, fadd) -TEST_FMATH_FUNC(int, fsub) -TEST_FMATH_FUNC(int, fmul) -TEST_FMATH_FUNC(int, fshiftleft) -TEST_FMATH_FUNC(int, fnegate) -TEST_FMATH_FUNC(uint, fadd) -TEST_FMATH_FUNC(uint, fsub) -TEST_FMATH_FUNC(uint, fmul) -TEST_FMATH_FUNC(uint, fshiftleft) +TEST_FMATH_FUNC_KHR(int, fadd) +TEST_FMATH_FUNC_KHR(int, fsub) +TEST_FMATH_FUNC_KHR(int, fmul) +TEST_FMATH_FUNC_KHR(int, fshiftleft) +TEST_FMATH_FUNC_KHR(int, fnegate) +TEST_FMATH_FUNC_KHR(uint, fadd) +TEST_FMATH_FUNC_KHR(uint, fsub) +TEST_FMATH_FUNC_KHR(uint, fmul) +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( \ + 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)