From af8c66aedcc526b31a1d2c69f16b0b37cedd42ff Mon Sep 17 00:00:00 2001 From: Ben Ashbaugh Date: Wed, 7 Aug 2024 14:05:02 -0700 Subject: [PATCH 1/5] add testing for the SPIR-V 1.4 no integer wrap decorations --- test_conformance/spirv_new/CMakeLists.txt | 2 +- test_conformance/spirv_new/spirvInfo.hpp | 40 ++++++++++++ ..._integer_wrap_decoration_fadd_int.spvasm32 | 46 ++++++++++++++ ..._integer_wrap_decoration_fadd_int.spvasm64 | 52 ++++++++++++++++ ...integer_wrap_decoration_fadd_uint.spvasm32 | 46 ++++++++++++++ ...integer_wrap_decoration_fadd_uint.spvasm64 | 52 ++++++++++++++++ ..._integer_wrap_decoration_fmul_int.spvasm32 | 46 ++++++++++++++ ..._integer_wrap_decoration_fmul_int.spvasm64 | 52 ++++++++++++++++ ...integer_wrap_decoration_fmul_uint.spvasm32 | 46 ++++++++++++++ ...integer_wrap_decoration_fmul_uint.spvasm64 | 52 ++++++++++++++++ ...teger_wrap_decoration_fnegate_int.spvasm32 | 45 ++++++++++++++ ...teger_wrap_decoration_fnegate_int.spvasm64 | 50 +++++++++++++++ ...er_wrap_decoration_fshiftleft_int.spvasm32 | 48 +++++++++++++++ ...er_wrap_decoration_fshiftleft_int.spvasm64 | 54 ++++++++++++++++ ...r_wrap_decoration_fshiftleft_uint.spvasm32 | 48 +++++++++++++++ ...r_wrap_decoration_fshiftleft_uint.spvasm64 | 54 ++++++++++++++++ ..._integer_wrap_decoration_fsub_int.spvasm32 | 46 ++++++++++++++ ..._integer_wrap_decoration_fsub_int.spvasm64 | 52 ++++++++++++++++ ...integer_wrap_decoration_fsub_uint.spvasm32 | 46 ++++++++++++++ ...integer_wrap_decoration_fsub_uint.spvasm64 | 52 ++++++++++++++++ ...pp => test_no_integer_wrap_decoration.cpp} | 61 ++++++++++++++----- 21 files changed, 973 insertions(+), 17 deletions(-) create mode 100644 test_conformance/spirv_new/spirvInfo.hpp create mode 100644 test_conformance/spirv_new/spirv_asm/spv1.4/no_integer_wrap_decoration_fadd_int.spvasm32 create mode 100644 test_conformance/spirv_new/spirv_asm/spv1.4/no_integer_wrap_decoration_fadd_int.spvasm64 create mode 100644 test_conformance/spirv_new/spirv_asm/spv1.4/no_integer_wrap_decoration_fadd_uint.spvasm32 create mode 100644 test_conformance/spirv_new/spirv_asm/spv1.4/no_integer_wrap_decoration_fadd_uint.spvasm64 create mode 100644 test_conformance/spirv_new/spirv_asm/spv1.4/no_integer_wrap_decoration_fmul_int.spvasm32 create mode 100644 test_conformance/spirv_new/spirv_asm/spv1.4/no_integer_wrap_decoration_fmul_int.spvasm64 create mode 100644 test_conformance/spirv_new/spirv_asm/spv1.4/no_integer_wrap_decoration_fmul_uint.spvasm32 create mode 100644 test_conformance/spirv_new/spirv_asm/spv1.4/no_integer_wrap_decoration_fmul_uint.spvasm64 create mode 100644 test_conformance/spirv_new/spirv_asm/spv1.4/no_integer_wrap_decoration_fnegate_int.spvasm32 create mode 100644 test_conformance/spirv_new/spirv_asm/spv1.4/no_integer_wrap_decoration_fnegate_int.spvasm64 create mode 100644 test_conformance/spirv_new/spirv_asm/spv1.4/no_integer_wrap_decoration_fshiftleft_int.spvasm32 create mode 100644 test_conformance/spirv_new/spirv_asm/spv1.4/no_integer_wrap_decoration_fshiftleft_int.spvasm64 create mode 100644 test_conformance/spirv_new/spirv_asm/spv1.4/no_integer_wrap_decoration_fshiftleft_uint.spvasm32 create mode 100644 test_conformance/spirv_new/spirv_asm/spv1.4/no_integer_wrap_decoration_fshiftleft_uint.spvasm64 create mode 100644 test_conformance/spirv_new/spirv_asm/spv1.4/no_integer_wrap_decoration_fsub_int.spvasm32 create mode 100644 test_conformance/spirv_new/spirv_asm/spv1.4/no_integer_wrap_decoration_fsub_int.spvasm64 create mode 100644 test_conformance/spirv_new/spirv_asm/spv1.4/no_integer_wrap_decoration_fsub_uint.spvasm32 create mode 100644 test_conformance/spirv_new/spirv_asm/spv1.4/no_integer_wrap_decoration_fsub_uint.spvasm64 rename test_conformance/spirv_new/{test_cl_khr_spirv_no_integer_wrap_decoration.cpp => test_no_integer_wrap_decoration.cpp} (78%) diff --git a/test_conformance/spirv_new/CMakeLists.txt b/test_conformance/spirv_new/CMakeLists.txt index 16a61b4075..4ff98e1b15 100644 --- a/test_conformance/spirv_new/CMakeLists.txt +++ b/test_conformance/spirv_new/CMakeLists.txt @@ -4,7 +4,6 @@ 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 @@ -27,6 +26,7 @@ set(${MODULE_NAME}_SOURCES test_op_vector_extract.cpp test_op_vector_insert.cpp test_op_vector_times_scalar.cpp + test_no_integer_wrap_decoration.cpp ) set(TEST_HARNESS_SOURCES diff --git a/test_conformance/spirv_new/spirvInfo.hpp b/test_conformance/spirv_new/spirvInfo.hpp new file mode 100644 index 0000000000..412d772f11 --- /dev/null +++ b/test_conformance/spirv_new/spirvInfo.hpp @@ -0,0 +1,40 @@ +// +// Copyright (c) 2024 The Khronos Group Inc. +// +// Licensed under the Apache License, Version 2.0 (the "License"); +// you may not use this file except in compliance with the License. +// You may obtain a copy of the License at +// +// http://www.apache.org/licenses/LICENSE-2.0 +// +// Unless required by applicable law or agreed to in writing, software +// distributed under the License is distributed on an "AS IS" BASIS, +// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +// See the License for the specific language governing permissions and +// limitations under the License. +// + +#pragma once + +#include "harness/compat.h" + +#include + +extern bool gVersionSkip; + +static bool is_spirv_version_supported(cl_device_id deviceID, const char* version) +{ + std::string ilVersions = get_device_il_version_string(deviceID); + + if (gVersionSkip) + { + log_info(" Skipping version check for %s.\n", version); + return true; + } + else if (ilVersions.find(version) == std::string::npos) + { + return false; + } + + return true; +} 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 0000000000..f70e2cdc39 --- /dev/null +++ b/test_conformance/spirv_new/spirv_asm/spv1.4/no_integer_wrap_decoration_fadd_int.spvasm32 @@ -0,0 +1,46 @@ +; SPIR-V +; Version: 1.4 +; Generator: Khronos SPIR-V Tools Assembler; 0 +; Bound: 23 +; Schema: 0 + OpCapability Addresses + OpCapability Linkage + OpCapability Kernel + %1 = OpExtInstImport "OpenCL.std" + 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 0000000000..d2ef26cc16 --- /dev/null +++ b/test_conformance/spirv_new/spirv_asm/spv1.4/no_integer_wrap_decoration_fadd_int.spvasm64 @@ -0,0 +1,52 @@ +; SPIR-V +; Version: 1.4 +; Generator: Khronos SPIR-V Tools Assembler; 0 +; Bound: 28 +; Schema: 0 + OpCapability Addresses + OpCapability Linkage + OpCapability Kernel + OpCapability Int64 + %1 = OpExtInstImport "OpenCL.std" + 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 0000000000..4d5f303b8f --- /dev/null +++ b/test_conformance/spirv_new/spirv_asm/spv1.4/no_integer_wrap_decoration_fadd_uint.spvasm32 @@ -0,0 +1,46 @@ +; SPIR-V +; Version: 1.4 +; Generator: Khronos SPIR-V Tools Assembler; 0 +; Bound: 23 +; Schema: 0 + OpCapability Addresses + OpCapability Linkage + OpCapability Kernel + %1 = OpExtInstImport "OpenCL.std" + 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 0000000000..7c950cb4c5 --- /dev/null +++ b/test_conformance/spirv_new/spirv_asm/spv1.4/no_integer_wrap_decoration_fadd_uint.spvasm64 @@ -0,0 +1,52 @@ +; SPIR-V +; Version: 1.4 +; Generator: Khronos SPIR-V Tools Assembler; 0 +; Bound: 28 +; Schema: 0 + OpCapability Addresses + OpCapability Linkage + OpCapability Kernel + OpCapability Int64 + %1 = OpExtInstImport "OpenCL.std" + 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 0000000000..173ba08921 --- /dev/null +++ b/test_conformance/spirv_new/spirv_asm/spv1.4/no_integer_wrap_decoration_fmul_int.spvasm32 @@ -0,0 +1,46 @@ +; SPIR-V +; Version: 1.4 +; Generator: Khronos SPIR-V Tools Assembler; 0 +; Bound: 23 +; Schema: 0 + OpCapability Addresses + OpCapability Linkage + OpCapability Kernel + %1 = OpExtInstImport "OpenCL.std" + 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 0000000000..990b491799 --- /dev/null +++ b/test_conformance/spirv_new/spirv_asm/spv1.4/no_integer_wrap_decoration_fmul_int.spvasm64 @@ -0,0 +1,52 @@ +; SPIR-V +; Version: 1.4 +; Generator: Khronos SPIR-V Tools Assembler; 0 +; Bound: 28 +; Schema: 0 + OpCapability Addresses + OpCapability Linkage + OpCapability Kernel + OpCapability Int64 + %1 = OpExtInstImport "OpenCL.std" + 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 0000000000..df5f83aa92 --- /dev/null +++ b/test_conformance/spirv_new/spirv_asm/spv1.4/no_integer_wrap_decoration_fmul_uint.spvasm32 @@ -0,0 +1,46 @@ +; SPIR-V +; Version: 1.4 +; Generator: Khronos SPIR-V Tools Assembler; 0 +; Bound: 23 +; Schema: 0 + OpCapability Addresses + OpCapability Linkage + OpCapability Kernel + %1 = OpExtInstImport "OpenCL.std" + 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 0000000000..3a2441f522 --- /dev/null +++ b/test_conformance/spirv_new/spirv_asm/spv1.4/no_integer_wrap_decoration_fmul_uint.spvasm64 @@ -0,0 +1,52 @@ +; SPIR-V +; Version: 1.3 +; Generator: Khronos SPIR-V Tools Assembler; 0 +; Bound: 28 +; Schema: 0 + OpCapability Addresses + OpCapability Linkage + OpCapability Kernel + OpCapability Int64 + %1 = OpExtInstImport "OpenCL.std" + 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 0000000000..f4939d07e0 --- /dev/null +++ b/test_conformance/spirv_new/spirv_asm/spv1.4/no_integer_wrap_decoration_fnegate_int.spvasm32 @@ -0,0 +1,45 @@ +; SPIR-V +; Version: 1.4 +; Generator: Khronos SPIR-V Tools Assembler; 0 +; Bound: 22 +; Schema: 0 + OpCapability Addresses + OpCapability Linkage + OpCapability Kernel + %1 = OpExtInstImport "OpenCL.std" + 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 0000000000..f9c2f91bf5 --- /dev/null +++ b/test_conformance/spirv_new/spirv_asm/spv1.4/no_integer_wrap_decoration_fnegate_int.spvasm64 @@ -0,0 +1,50 @@ +; SPIR-V +; Version: 1.4 +; Generator: Khronos SPIR-V Tools Assembler; 0 +; Bound: 26 +; Schema: 0 + OpCapability Addresses + OpCapability Linkage + OpCapability Kernel + OpCapability Int64 + %1 = OpExtInstImport "OpenCL.std" + 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 0000000000..6f73fc1d41 --- /dev/null +++ b/test_conformance/spirv_new/spirv_asm/spv1.4/no_integer_wrap_decoration_fshiftleft_int.spvasm32 @@ -0,0 +1,48 @@ +; SPIR-V +; Version: 1.4 +; Generator: Khronos SPIR-V Tools Assembler; 0 +; Bound: 25 +; Schema: 0 + OpCapability Addresses + OpCapability Linkage + OpCapability Kernel + %1 = OpExtInstImport "OpenCL.std" + 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 0000000000..525d0a0a56 --- /dev/null +++ b/test_conformance/spirv_new/spirv_asm/spv1.4/no_integer_wrap_decoration_fshiftleft_int.spvasm64 @@ -0,0 +1,54 @@ +; SPIR-V +; Version: 1.4 +; Generator: Khronos SPIR-V Tools Assembler; 0 +; Bound: 30 +; Schema: 0 + OpCapability Addresses + OpCapability Linkage + OpCapability Kernel + OpCapability Int64 + %1 = OpExtInstImport "OpenCL.std" + 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 0000000000..bfa300d4f0 --- /dev/null +++ b/test_conformance/spirv_new/spirv_asm/spv1.4/no_integer_wrap_decoration_fshiftleft_uint.spvasm32 @@ -0,0 +1,48 @@ +; SPIR-V +; Version: 1.4 +; Generator: Khronos SPIR-V Tools Assembler; 0 +; Bound: 25 +; Schema: 0 + OpCapability Addresses + OpCapability Linkage + OpCapability Kernel + %1 = OpExtInstImport "OpenCL.std" + 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 0000000000..b102a2f8df --- /dev/null +++ b/test_conformance/spirv_new/spirv_asm/spv1.4/no_integer_wrap_decoration_fshiftleft_uint.spvasm64 @@ -0,0 +1,54 @@ +; SPIR-V +; Version: 1.4 +; Generator: Khronos SPIR-V Tools Assembler; 0 +; Bound: 30 +; Schema: 0 + OpCapability Addresses + OpCapability Linkage + OpCapability Kernel + OpCapability Int64 + %1 = OpExtInstImport "OpenCL.std" + 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 0000000000..2e3a8bcb37 --- /dev/null +++ b/test_conformance/spirv_new/spirv_asm/spv1.4/no_integer_wrap_decoration_fsub_int.spvasm32 @@ -0,0 +1,46 @@ +; SPIR-V +; Version: 1.4 +; Generator: Khronos SPIR-V Tools Assembler; 0 +; Bound: 23 +; Schema: 0 + OpCapability Addresses + OpCapability Linkage + OpCapability Kernel + %1 = OpExtInstImport "OpenCL.std" + 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 0000000000..be3e15758c --- /dev/null +++ b/test_conformance/spirv_new/spirv_asm/spv1.4/no_integer_wrap_decoration_fsub_int.spvasm64 @@ -0,0 +1,52 @@ +; SPIR-V +; Version: 1.4 +; Generator: Khronos SPIR-V Tools Assembler; 0 +; Bound: 28 +; Schema: 0 + OpCapability Addresses + OpCapability Linkage + OpCapability Kernel + OpCapability Int64 + %1 = OpExtInstImport "OpenCL.std" + 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 0000000000..7780593d20 --- /dev/null +++ b/test_conformance/spirv_new/spirv_asm/spv1.4/no_integer_wrap_decoration_fsub_uint.spvasm32 @@ -0,0 +1,46 @@ +; SPIR-V +; Version: 1.4 +; Generator: Khronos SPIR-V Tools Assembler; 0 +; Bound: 23 +; Schema: 0 + OpCapability Addresses + OpCapability Linkage + OpCapability Kernel + %1 = OpExtInstImport "OpenCL.std" + 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 0000000000..0082629c6a --- /dev/null +++ b/test_conformance/spirv_new/spirv_asm/spv1.4/no_integer_wrap_decoration_fsub_uint.spvasm64 @@ -0,0 +1,52 @@ +; SPIR-V +; Version: 1.4 +; Generator: Khronos SPIR-V Tools Assembler; 0 +; Bound: 28 +; Schema: 0 + OpCapability Addresses + OpCapability Linkage + OpCapability Kernel + OpCapability Int64 + %1 = OpExtInstImport "OpenCL.std" + 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 78% 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 95227b28ff..641f2993ca 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 @@ -23,9 +24,10 @@ template -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_command_queue queue, + bool test_extension, const char *spvName, const char *funcName, const char *Tname) @@ -37,9 +39,16 @@ 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; + if (test_extension) { + 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; + } + } else { + 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; + } } /*Test with some values that do not cause overflow*/ @@ -197,25 +206,45 @@ int test_ext_cl_khr_spirv_no_integer_wrap_decoration(cl_device_id deviceID, } } - 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) \ { \ - return test_ext_cl_khr_spirv_no_integer_wrap_decoration(deviceID, context, queue, \ + return test_no_integer_wrap_decoration(deviceID, context, queue, true, \ "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) \ + { \ + return test_no_integer_wrap_decoration(deviceID, context, queue, false, \ + "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) From 6e10ca9fc106293f4dd9f72bedf5567b13fd4101 Mon Sep 17 00:00:00 2001 From: Ben Ashbaugh Date: Wed, 7 Aug 2024 14:11:09 -0700 Subject: [PATCH 2/5] fix formatting --- test_conformance/spirv_new/spirvInfo.hpp | 3 +- .../test_no_integer_wrap_decoration.cpp | 144 +++++++++++------- 2 files changed, 91 insertions(+), 56 deletions(-) diff --git a/test_conformance/spirv_new/spirvInfo.hpp b/test_conformance/spirv_new/spirvInfo.hpp index 412d772f11..ed4d6c794c 100644 --- a/test_conformance/spirv_new/spirvInfo.hpp +++ b/test_conformance/spirv_new/spirvInfo.hpp @@ -22,7 +22,8 @@ extern bool gVersionSkip; -static bool is_spirv_version_supported(cl_device_id deviceID, const char* version) +static bool is_spirv_version_supported(cl_device_id deviceID, + const char* version) { std::string ilVersions = get_device_il_version_string(deviceID); diff --git a/test_conformance/spirv_new/test_no_integer_wrap_decoration.cpp b/test_conformance/spirv_new/test_no_integer_wrap_decoration.cpp index 641f2993ca..a9c9363f2a 100644 --- a/test_conformance/spirv_new/test_no_integer_wrap_decoration.cpp +++ b/test_conformance/spirv_new/test_no_integer_wrap_decoration.cpp @@ -23,14 +23,11 @@ #include -template -int test_no_integer_wrap_decoration(cl_device_id deviceID, - cl_context context, - cl_command_queue queue, - bool test_extension, - 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, bool test_extension, + const char *spvName, const char *funcName, + const char *Tname) { cl_int err = CL_SUCCESS; @@ -39,20 +36,28 @@ int test_no_integer_wrap_decoration(cl_device_id deviceID, std::vector h_rhs(num); std::vector expected_results(num); std::vector h_ref(num); - if (test_extension) { - 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"); + if (test_extension) + { + 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; } - } else { - if (!is_spirv_version_supported(deviceID, "SPIR-V_1.4")) { + } + else + { + 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; } } /*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); @@ -63,7 +68,9 @@ int test_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); @@ -88,16 +95,20 @@ int test_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; @@ -110,8 +121,8 @@ int test_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"; @@ -123,16 +134,26 @@ int test_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]; } } @@ -145,7 +166,8 @@ int test_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); @@ -158,16 +180,22 @@ int test_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; } } @@ -179,7 +207,8 @@ int test_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); @@ -192,16 +221,22 @@ int test_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; } } @@ -209,14 +244,14 @@ int test_no_integer_wrap_decoration(cl_device_id deviceID, return TEST_PASS; } -#define TEST_FMATH_FUNC_KHR(TYPE, FUNC) \ - TEST_SPIRV_FUNC(ext_cl_khr_spirv_no_integer_wrap_decoration_##FUNC##_##TYPE) \ - { \ - return test_no_integer_wrap_decoration(deviceID, context, queue, true, \ - "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) \ + { \ + return test_no_integer_wrap_decoration( \ + deviceID, context, queue, true, \ + "ext_cl_khr_spirv_no_integer_wrap_decoration_" #FUNC "_" #TYPE, \ + #FUNC, #TYPE); \ } TEST_FMATH_FUNC_KHR(int, fadd) @@ -229,14 +264,13 @@ 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) \ - { \ - return test_no_integer_wrap_decoration(deviceID, context, queue, false, \ - "spv1.4/no_integer_wrap_decoration_"#FUNC"_"#TYPE, \ - #FUNC, \ - #TYPE \ - ); \ +#define TEST_FMATH_FUNC_14(TYPE, FUNC) \ + TEST_SPIRV_FUNC(spirv14_no_integer_wrap_decoration_##FUNC##_##TYPE) \ + { \ + return test_no_integer_wrap_decoration( \ + deviceID, context, queue, false, \ + "spv1.4/no_integer_wrap_decoration_" #FUNC "_" #TYPE, #FUNC, \ + #TYPE); \ } TEST_FMATH_FUNC_14(int, fadd) From a9a03d692c5597077af53b858f6c6be1704823e2 Mon Sep 17 00:00:00 2001 From: Ben Ashbaugh Date: Fri, 16 Aug 2024 17:15:48 -0700 Subject: [PATCH 3/5] move the extension or version check out of the function --- .../test_no_integer_wrap_decoration.cpp | 39 ++++++++----------- 1 file changed, 16 insertions(+), 23 deletions(-) diff --git a/test_conformance/spirv_new/test_no_integer_wrap_decoration.cpp b/test_conformance/spirv_new/test_no_integer_wrap_decoration.cpp index a9c9363f2a..042f0757d5 100644 --- a/test_conformance/spirv_new/test_no_integer_wrap_decoration.cpp +++ b/test_conformance/spirv_new/test_no_integer_wrap_decoration.cpp @@ -25,9 +25,8 @@ template int test_no_integer_wrap_decoration(cl_device_id deviceID, cl_context context, - cl_command_queue queue, bool test_extension, - const char *spvName, const char *funcName, - const char *Tname) + cl_command_queue queue, const char *spvName, + const char *funcName, const char *Tname) { cl_int err = CL_SUCCESS; @@ -36,24 +35,6 @@ int test_no_integer_wrap_decoration(cl_device_id deviceID, cl_context context, std::vector h_rhs(num); std::vector expected_results(num); std::vector h_ref(num); - if (test_extension) - { - 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; - } - } - else - { - 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; - } - } /*Test with some values that do not cause overflow*/ if (std::is_signed::value == true) @@ -248,8 +229,15 @@ int test_no_integer_wrap_decoration(cl_device_id deviceID, cl_context context, 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, true, \ + deviceID, context, queue, \ "ext_cl_khr_spirv_no_integer_wrap_decoration_" #FUNC "_" #TYPE, \ #FUNC, #TYPE); \ } @@ -267,8 +255,13 @@ 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, false, \ + deviceID, context, queue, \ "spv1.4/no_integer_wrap_decoration_" #FUNC "_" #TYPE, #FUNC, \ #TYPE); \ } From dd07b2a3221fe16dab379ccd8cb7c444b19295a6 Mon Sep 17 00:00:00 2001 From: Ben Ashbaugh Date: Wed, 21 Aug 2024 22:26:43 -0700 Subject: [PATCH 4/5] remove unnecessary OpExtInstImport --- .../spv1.4/no_integer_wrap_decoration_fadd_int.spvasm32 | 1 - .../spv1.4/no_integer_wrap_decoration_fadd_int.spvasm64 | 1 - .../spv1.4/no_integer_wrap_decoration_fadd_uint.spvasm32 | 1 - .../spv1.4/no_integer_wrap_decoration_fadd_uint.spvasm64 | 1 - .../spv1.4/no_integer_wrap_decoration_fmul_int.spvasm32 | 1 - .../spv1.4/no_integer_wrap_decoration_fmul_int.spvasm64 | 1 - .../spv1.4/no_integer_wrap_decoration_fmul_uint.spvasm32 | 1 - .../spv1.4/no_integer_wrap_decoration_fmul_uint.spvasm64 | 1 - .../spv1.4/no_integer_wrap_decoration_fnegate_int.spvasm32 | 1 - .../spv1.4/no_integer_wrap_decoration_fnegate_int.spvasm64 | 1 - .../spv1.4/no_integer_wrap_decoration_fshiftleft_int.spvasm32 | 1 - .../spv1.4/no_integer_wrap_decoration_fshiftleft_int.spvasm64 | 1 - .../spv1.4/no_integer_wrap_decoration_fshiftleft_uint.spvasm32 | 1 - .../spv1.4/no_integer_wrap_decoration_fshiftleft_uint.spvasm64 | 1 - .../spv1.4/no_integer_wrap_decoration_fsub_int.spvasm32 | 1 - .../spv1.4/no_integer_wrap_decoration_fsub_int.spvasm64 | 1 - .../spv1.4/no_integer_wrap_decoration_fsub_uint.spvasm32 | 1 - .../spv1.4/no_integer_wrap_decoration_fsub_uint.spvasm64 | 1 - 18 files changed, 18 deletions(-) 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 index f70e2cdc39..557faae089 100644 --- 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 @@ -6,7 +6,6 @@ OpCapability Addresses OpCapability Linkage OpCapability Kernel - %1 = OpExtInstImport "OpenCL.std" OpMemoryModel Physical32 OpenCL OpEntryPoint Kernel %2 "fmath_cl" %__spirv_BuiltInGlobalInvocationId OpSource OpenCL_C 200000 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 index d2ef26cc16..4bc2c2a7cc 100644 --- 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 @@ -7,7 +7,6 @@ OpCapability Linkage OpCapability Kernel OpCapability Int64 - %1 = OpExtInstImport "OpenCL.std" OpMemoryModel Physical64 OpenCL OpEntryPoint Kernel %2 "fmath_cl" %__spirv_BuiltInGlobalInvocationId OpSource OpenCL_C 200000 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 index 4d5f303b8f..b858523c9d 100644 --- 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 @@ -6,7 +6,6 @@ OpCapability Addresses OpCapability Linkage OpCapability Kernel - %1 = OpExtInstImport "OpenCL.std" OpMemoryModel Physical32 OpenCL OpEntryPoint Kernel %2 "fmath_cl" %__spirv_BuiltInGlobalInvocationId OpSource OpenCL_C 200000 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 index 7c950cb4c5..a41f4f41e4 100644 --- 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 @@ -7,7 +7,6 @@ OpCapability Linkage OpCapability Kernel OpCapability Int64 - %1 = OpExtInstImport "OpenCL.std" OpMemoryModel Physical64 OpenCL OpEntryPoint Kernel %2 "fmath_cl" %__spirv_BuiltInGlobalInvocationId OpSource OpenCL_C 200000 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 index 173ba08921..4886c01c43 100644 --- 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 @@ -6,7 +6,6 @@ OpCapability Addresses OpCapability Linkage OpCapability Kernel - %1 = OpExtInstImport "OpenCL.std" OpMemoryModel Physical32 OpenCL OpEntryPoint Kernel %2 "fmath_cl" %__spirv_BuiltInGlobalInvocationId OpSource OpenCL_C 200000 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 index 990b491799..ff6799e71a 100644 --- 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 @@ -7,7 +7,6 @@ OpCapability Linkage OpCapability Kernel OpCapability Int64 - %1 = OpExtInstImport "OpenCL.std" OpMemoryModel Physical64 OpenCL OpEntryPoint Kernel %2 "fmath_cl" %__spirv_BuiltInGlobalInvocationId OpSource OpenCL_C 200000 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 index df5f83aa92..1993a13d17 100644 --- 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 @@ -6,7 +6,6 @@ OpCapability Addresses OpCapability Linkage OpCapability Kernel - %1 = OpExtInstImport "OpenCL.std" OpMemoryModel Physical32 OpenCL OpEntryPoint Kernel %2 "fmath_cl" %__spirv_BuiltInGlobalInvocationId OpSource OpenCL_C 200000 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 index 3a2441f522..d17b9c292e 100644 --- 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 @@ -7,7 +7,6 @@ OpCapability Linkage OpCapability Kernel OpCapability Int64 - %1 = OpExtInstImport "OpenCL.std" OpMemoryModel Physical64 OpenCL OpEntryPoint Kernel %2 "fmath_cl" %__spirv_BuiltInGlobalInvocationId OpSource OpenCL_C 200000 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 index f4939d07e0..ea9a85adbf 100644 --- 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 @@ -6,7 +6,6 @@ OpCapability Addresses OpCapability Linkage OpCapability Kernel - %1 = OpExtInstImport "OpenCL.std" OpMemoryModel Physical32 OpenCL OpEntryPoint Kernel %2 "fmath_cl" %__spirv_BuiltInGlobalInvocationId OpSource OpenCL_C 200000 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 index f9c2f91bf5..2edfb49989 100644 --- 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 @@ -7,7 +7,6 @@ OpCapability Linkage OpCapability Kernel OpCapability Int64 - %1 = OpExtInstImport "OpenCL.std" OpMemoryModel Physical64 OpenCL OpEntryPoint Kernel %2 "fmath_cl" %__spirv_BuiltInGlobalInvocationId OpSource OpenCL_C 200000 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 index 6f73fc1d41..08b5ba7152 100644 --- 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 @@ -6,7 +6,6 @@ OpCapability Addresses OpCapability Linkage OpCapability Kernel - %1 = OpExtInstImport "OpenCL.std" OpMemoryModel Physical32 OpenCL OpEntryPoint Kernel %2 "fmath_cl" %__spirv_BuiltInGlobalInvocationId OpSource OpenCL_C 200000 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 index 525d0a0a56..3487dc46ea 100644 --- 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 @@ -7,7 +7,6 @@ OpCapability Linkage OpCapability Kernel OpCapability Int64 - %1 = OpExtInstImport "OpenCL.std" OpMemoryModel Physical64 OpenCL OpEntryPoint Kernel %2 "fmath_cl" %__spirv_BuiltInGlobalInvocationId OpSource OpenCL_C 200000 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 index bfa300d4f0..ac01b182d1 100644 --- 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 @@ -6,7 +6,6 @@ OpCapability Addresses OpCapability Linkage OpCapability Kernel - %1 = OpExtInstImport "OpenCL.std" OpMemoryModel Physical32 OpenCL OpEntryPoint Kernel %2 "fmath_cl" %__spirv_BuiltInGlobalInvocationId OpSource OpenCL_C 200000 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 index b102a2f8df..a2d8ba698b 100644 --- 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 @@ -7,7 +7,6 @@ OpCapability Linkage OpCapability Kernel OpCapability Int64 - %1 = OpExtInstImport "OpenCL.std" OpMemoryModel Physical64 OpenCL OpEntryPoint Kernel %2 "fmath_cl" %__spirv_BuiltInGlobalInvocationId OpSource OpenCL_C 200000 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 index 2e3a8bcb37..946be21171 100644 --- 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 @@ -6,7 +6,6 @@ OpCapability Addresses OpCapability Linkage OpCapability Kernel - %1 = OpExtInstImport "OpenCL.std" OpMemoryModel Physical32 OpenCL OpEntryPoint Kernel %2 "fmath_cl" %__spirv_BuiltInGlobalInvocationId OpSource OpenCL_C 200000 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 index be3e15758c..26ce619ef3 100644 --- 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 @@ -7,7 +7,6 @@ OpCapability Linkage OpCapability Kernel OpCapability Int64 - %1 = OpExtInstImport "OpenCL.std" OpMemoryModel Physical64 OpenCL OpEntryPoint Kernel %2 "fmath_cl" %__spirv_BuiltInGlobalInvocationId OpSource OpenCL_C 200000 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 index 7780593d20..013614f3c1 100644 --- 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 @@ -6,7 +6,6 @@ OpCapability Addresses OpCapability Linkage OpCapability Kernel - %1 = OpExtInstImport "OpenCL.std" OpMemoryModel Physical32 OpenCL OpEntryPoint Kernel %2 "fmath_cl" %__spirv_BuiltInGlobalInvocationId OpSource OpenCL_C 200000 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 index 0082629c6a..19df8b5faf 100644 --- 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 @@ -7,7 +7,6 @@ OpCapability Linkage OpCapability Kernel OpCapability Int64 - %1 = OpExtInstImport "OpenCL.std" OpMemoryModel Physical64 OpenCL OpEntryPoint Kernel %2 "fmath_cl" %__spirv_BuiltInGlobalInvocationId OpSource OpenCL_C 200000 From 2c86330b244b45f8797ad3aece7efd7c84af6ea9 Mon Sep 17 00:00:00 2001 From: Ben Ashbaugh Date: Thu, 22 Aug 2024 09:55:54 -0700 Subject: [PATCH 5/5] keep files in CMakeLists.txt in alphabetical order --- test_conformance/spirv_new/CMakeLists.txt | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/test_conformance/spirv_new/CMakeLists.txt b/test_conformance/spirv_new/CMakeLists.txt index 4ff98e1b15..3f7d1f89a0 100644 --- a/test_conformance/spirv_new/CMakeLists.txt +++ b/test_conformance/spirv_new/CMakeLists.txt @@ -7,6 +7,7 @@ set(${MODULE_NAME}_SOURCES 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 @@ -26,7 +27,6 @@ set(${MODULE_NAME}_SOURCES test_op_vector_extract.cpp test_op_vector_insert.cpp test_op_vector_times_scalar.cpp - test_no_integer_wrap_decoration.cpp ) set(TEST_HARNESS_SOURCES