Add scalar uint with signed int compares as Amber cases
Logs an error and throws an exception if the script fails to parse.
Components: Vulkan
Affects:
dEQP-VK.spirv_assembly.instruction.compute.signed_int_compare.uint_sgreaterthan
dEQP-VK.spirv_assembly.instruction.compute.signed_int_compare.uint_sgreaterthanequal
dEQP-VK.spirv_assembly.instruction.compute.signed_int_compare.uint_slessthan
dEQP-VK.spirv_assembly.instruction.compute.signed_int_compare.uint_slessthanequal
VK-GL-CTS issue: 1147
Change-Id: Ic8ce35bff0e05360dc9e674cce313d441e2f4953
diff --git a/AndroidGen.mk b/AndroidGen.mk
index 319a65d..60ee1ae 100644
--- a/AndroidGen.mk
+++ b/AndroidGen.mk
@@ -297,6 +297,7 @@
external/vulkancts/modules/vulkan/spirv_assembly/vktSpvAsmLoopDepInfTests.cpp \
external/vulkancts/modules/vulkan/spirv_assembly/vktSpvAsmLoopDepLenTests.cpp \
external/vulkancts/modules/vulkan/spirv_assembly/vktSpvAsmPointerParameterTests.cpp \
+ external/vulkancts/modules/vulkan/spirv_assembly/vktSpvAsmSignedIntCompareTests.cpp \
external/vulkancts/modules/vulkan/spirv_assembly/vktSpvAsmSpirvVersionTests.cpp \
external/vulkancts/modules/vulkan/spirv_assembly/vktSpvAsmTests.cpp \
external/vulkancts/modules/vulkan/spirv_assembly/vktSpvAsmTypeTests.cpp \
diff --git a/android/cts/master/vk-master.txt b/android/cts/master/vk-master.txt
index 8757d60..8b02480 100755
--- a/android/cts/master/vk-master.txt
+++ b/android/cts/master/vk-master.txt
@@ -228095,6 +228095,10 @@
dEQP-VK.spirv_assembly.instruction.compute.workgroup_memory.uint8
dEQP-VK.spirv_assembly.instruction.compute.spirv_ids_abuse.sparse_ids
dEQP-VK.spirv_assembly.instruction.compute.spirv_ids_abuse.lots_ids
+dEQP-VK.spirv_assembly.instruction.compute.signed_int_compare.uint_sgreaterthanequal
+dEQP-VK.spirv_assembly.instruction.compute.signed_int_compare.uint_sgreaterthan
+dEQP-VK.spirv_assembly.instruction.compute.signed_int_compare.uint_slessthan
+dEQP-VK.spirv_assembly.instruction.compute.signed_int_compare.uint_slessthanequal
dEQP-VK.spirv_assembly.instruction.graphics.cross_stage.basic_type.flat
dEQP-VK.spirv_assembly.instruction.graphics.cross_stage.basic_type.no_perspective
dEQP-VK.spirv_assembly.instruction.graphics.cross_stage.basic_type.relaxedprecision
diff --git a/external/vulkancts/data/vulkan/amber/spirv_assembly/instruction/compute/signed_int_compare/uint_sgreaterthan.amber b/external/vulkancts/data/vulkan/amber/spirv_assembly/instruction/compute/signed_int_compare/uint_sgreaterthan.amber
new file mode 100644
index 0000000..3c3d788
--- /dev/null
+++ b/external/vulkancts/data/vulkan/amber/spirv_assembly/instruction/compute/signed_int_compare/uint_sgreaterthan.amber
@@ -0,0 +1,69 @@
+# Test SGreaterThan with unsigned int params
+# Google bug b/73133282
+#
+# Derived from the following OpenCL C, but cleaned up to be more generic.
+#
+# kernel void foo(global int *A, global int*B, global int*C) __attribute__((reqd_work_group_size(1,1,1))) {
+# uint i = get_global_id(0);
+# C[i] = A[i] > B[i];
+# }
+
+[compute shader spirv]
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %18 "main" %gl_GlobalInvocationID
+ OpExecutionMode %18 LocalSize 1 1 1
+ OpSource OpenCL_C 120
+ OpDecorate %_runtimearr_uint ArrayStride 4
+ OpMemberDecorate %_struct_3 0 Offset 0
+ OpDecorate %_struct_3 BufferBlock
+ OpDecorate %gl_GlobalInvocationID BuiltIn GlobalInvocationId
+ OpDecorate %15 DescriptorSet 0
+ OpDecorate %15 Binding 0
+ OpDecorate %16 DescriptorSet 0
+ OpDecorate %16 Binding 1
+ OpDecorate %17 DescriptorSet 0
+ OpDecorate %17 Binding 2
+ %uint = OpTypeInt 32 0
+%_runtimearr_uint = OpTypeRuntimeArray %uint
+ %_struct_3 = OpTypeStruct %_runtimearr_uint
+%_ptr_Uniform__struct_3 = OpTypePointer Uniform %_struct_3
+ %void = OpTypeVoid
+ %6 = OpTypeFunction %void
+ %v3uint = OpTypeVector %uint 3
+%_ptr_Input_v3uint = OpTypePointer Input %v3uint
+%_ptr_Input_uint = OpTypePointer Input %uint
+%_ptr_Uniform_uint = OpTypePointer Uniform %uint
+ %bool = OpTypeBool
+ %uint_0 = OpConstant %uint 0
+ %uint_1 = OpConstant %uint 1
+%gl_GlobalInvocationID = OpVariable %_ptr_Input_v3uint Input
+ %15 = OpVariable %_ptr_Uniform__struct_3 Uniform
+ %16 = OpVariable %_ptr_Uniform__struct_3 Uniform
+ %17 = OpVariable %_ptr_Uniform__struct_3 Uniform
+ %18 = OpFunction %void None %6
+ %19 = OpLabel
+ %20 = OpAccessChain %_ptr_Input_uint %gl_GlobalInvocationID %uint_0
+ %21 = OpLoad %uint %20
+ %22 = OpAccessChain %_ptr_Uniform_uint %15 %uint_0 %21
+ %23 = OpLoad %uint %22
+ %24 = OpAccessChain %_ptr_Uniform_uint %16 %uint_0 %21
+ %25 = OpLoad %uint %24
+ %26 = OpSGreaterThan %bool %23 %25
+ %27 = OpSelect %uint %26 %uint_1 %uint_0
+ %28 = OpAccessChain %_ptr_Uniform_uint %17 %uint_0 %21
+ OpStore %28 %27
+ OpReturn
+ OpFunctionEnd
+
+[test]
+# A[]
+ssbo 0:0 subdata int 0 -8 -7 -6 -5 -4 -3 -2 0 0 1 2 3 4 5 6 7
+# B[]
+ssbo 0:1 subdata int 0 -9 -7 -5 2 -1 1 0 0 1 0 2 -2 4 8 4 -4
+# The answer array C[]
+ssbo 0:2 subdata int 0 8 8 8 8 8 8 8 8 8 8 8 8 8 8 8 8
+
+compute 16 1 1
+
+probe ssbo int 0:2 0 == 1 0 0 0 0 0 0 0 0 1 0 1 0 0 1 1
diff --git a/external/vulkancts/data/vulkan/amber/spirv_assembly/instruction/compute/signed_int_compare/uint_sgreaterthanequal.amber b/external/vulkancts/data/vulkan/amber/spirv_assembly/instruction/compute/signed_int_compare/uint_sgreaterthanequal.amber
new file mode 100644
index 0000000..607a53b
--- /dev/null
+++ b/external/vulkancts/data/vulkan/amber/spirv_assembly/instruction/compute/signed_int_compare/uint_sgreaterthanequal.amber
@@ -0,0 +1,69 @@
+# Test SGreaterThanEqual with unsigned int params
+# Google bug b/73133282
+#
+# Derived from the following OpenCL C, but cleaned up to be more generic.
+#
+# kernel void foo(global int *A, global int*B, global int*C) __attribute__((reqd_work_group_size(1,1,1))) {
+# uint i = get_global_id(0);
+# C[i] = A[i] >= B[i];
+# }
+
+[compute shader spirv]
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %18 "main" %gl_GlobalInvocationID
+ OpExecutionMode %18 LocalSize 1 1 1
+ OpSource OpenCL_C 120
+ OpDecorate %_runtimearr_uint ArrayStride 4
+ OpMemberDecorate %_struct_3 0 Offset 0
+ OpDecorate %_struct_3 BufferBlock
+ OpDecorate %gl_GlobalInvocationID BuiltIn GlobalInvocationId
+ OpDecorate %15 DescriptorSet 0
+ OpDecorate %15 Binding 0
+ OpDecorate %16 DescriptorSet 0
+ OpDecorate %16 Binding 1
+ OpDecorate %17 DescriptorSet 0
+ OpDecorate %17 Binding 2
+ %uint = OpTypeInt 32 0
+%_runtimearr_uint = OpTypeRuntimeArray %uint
+ %_struct_3 = OpTypeStruct %_runtimearr_uint
+%_ptr_Uniform__struct_3 = OpTypePointer Uniform %_struct_3
+ %void = OpTypeVoid
+ %6 = OpTypeFunction %void
+ %v3uint = OpTypeVector %uint 3
+%_ptr_Input_v3uint = OpTypePointer Input %v3uint
+%_ptr_Input_uint = OpTypePointer Input %uint
+%_ptr_Uniform_uint = OpTypePointer Uniform %uint
+ %bool = OpTypeBool
+ %uint_0 = OpConstant %uint 0
+ %uint_1 = OpConstant %uint 1
+%gl_GlobalInvocationID = OpVariable %_ptr_Input_v3uint Input
+ %15 = OpVariable %_ptr_Uniform__struct_3 Uniform
+ %16 = OpVariable %_ptr_Uniform__struct_3 Uniform
+ %17 = OpVariable %_ptr_Uniform__struct_3 Uniform
+ %18 = OpFunction %void None %6
+ %19 = OpLabel
+ %20 = OpAccessChain %_ptr_Input_uint %gl_GlobalInvocationID %uint_0
+ %21 = OpLoad %uint %20
+ %22 = OpAccessChain %_ptr_Uniform_uint %15 %uint_0 %21
+ %23 = OpLoad %uint %22
+ %24 = OpAccessChain %_ptr_Uniform_uint %16 %uint_0 %21
+ %25 = OpLoad %uint %24
+ %26 = OpSGreaterThanEqual %bool %23 %25
+ %27 = OpSelect %uint %26 %uint_1 %uint_0
+ %28 = OpAccessChain %_ptr_Uniform_uint %17 %uint_0 %21
+ OpStore %28 %27
+ OpReturn
+ OpFunctionEnd
+
+[test]
+# A[]
+ssbo 0:0 subdata int 0 -8 -7 -6 -5 -4 -3 -2 0 0 1 2 3 4 5 6 7
+# B[]
+ssbo 0:1 subdata int 0 -9 -7 -5 2 -1 1 0 0 1 0 2 -2 4 8 4 -4
+# The answer array C[]
+ssbo 0:2 subdata int 0 8 8 8 8 8 8 8 8 8 8 8 8 8 8 8 8
+
+compute 16 1 1
+
+probe ssbo int 0:2 0 == 1 1 0 0 0 0 0 1 0 1 1 1 1 0 1 1
diff --git a/external/vulkancts/data/vulkan/amber/spirv_assembly/instruction/compute/signed_int_compare/uint_slessthan.amber b/external/vulkancts/data/vulkan/amber/spirv_assembly/instruction/compute/signed_int_compare/uint_slessthan.amber
new file mode 100644
index 0000000..2a75eed
--- /dev/null
+++ b/external/vulkancts/data/vulkan/amber/spirv_assembly/instruction/compute/signed_int_compare/uint_slessthan.amber
@@ -0,0 +1,69 @@
+# Test SLessThan with unsigned int params
+# Google bug b/73133282
+#
+# Derived from the following OpenCL C, but cleaned up to be more generic.
+#
+# kernel void foo(global int *A, global int*B, global int*C) __attribute__((reqd_work_group_size(1,1,1))) {
+# uint i = get_global_id(0);
+# C[i] = A[i] < B[i];
+# }
+
+[compute shader spirv]
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %18 "main" %gl_GlobalInvocationID
+ OpExecutionMode %18 LocalSize 1 1 1
+ OpSource OpenCL_C 120
+ OpDecorate %_runtimearr_uint ArrayStride 4
+ OpMemberDecorate %_struct_3 0 Offset 0
+ OpDecorate %_struct_3 BufferBlock
+ OpDecorate %gl_GlobalInvocationID BuiltIn GlobalInvocationId
+ OpDecorate %15 DescriptorSet 0
+ OpDecorate %15 Binding 0
+ OpDecorate %16 DescriptorSet 0
+ OpDecorate %16 Binding 1
+ OpDecorate %17 DescriptorSet 0
+ OpDecorate %17 Binding 2
+ %uint = OpTypeInt 32 0
+%_runtimearr_uint = OpTypeRuntimeArray %uint
+ %_struct_3 = OpTypeStruct %_runtimearr_uint
+%_ptr_Uniform__struct_3 = OpTypePointer Uniform %_struct_3
+ %void = OpTypeVoid
+ %6 = OpTypeFunction %void
+ %v3uint = OpTypeVector %uint 3
+%_ptr_Input_v3uint = OpTypePointer Input %v3uint
+%_ptr_Input_uint = OpTypePointer Input %uint
+%_ptr_Uniform_uint = OpTypePointer Uniform %uint
+ %bool = OpTypeBool
+ %uint_0 = OpConstant %uint 0
+ %uint_1 = OpConstant %uint 1
+%gl_GlobalInvocationID = OpVariable %_ptr_Input_v3uint Input
+ %15 = OpVariable %_ptr_Uniform__struct_3 Uniform
+ %16 = OpVariable %_ptr_Uniform__struct_3 Uniform
+ %17 = OpVariable %_ptr_Uniform__struct_3 Uniform
+ %18 = OpFunction %void None %6
+ %19 = OpLabel
+ %20 = OpAccessChain %_ptr_Input_uint %gl_GlobalInvocationID %uint_0
+ %21 = OpLoad %uint %20
+ %22 = OpAccessChain %_ptr_Uniform_uint %15 %uint_0 %21
+ %23 = OpLoad %uint %22
+ %24 = OpAccessChain %_ptr_Uniform_uint %16 %uint_0 %21
+ %25 = OpLoad %uint %24
+ %26 = OpSLessThan %bool %23 %25
+ %27 = OpSelect %uint %26 %uint_1 %uint_0
+ %28 = OpAccessChain %_ptr_Uniform_uint %17 %uint_0 %21
+ OpStore %28 %27
+ OpReturn
+ OpFunctionEnd
+
+[test]
+# A[]
+ssbo 0:0 subdata int 0 -8 -7 -6 -5 -4 -3 -2 0 0 1 2 3 4 5 6 7
+# B[]
+ssbo 0:1 subdata int 0 -9 -7 -5 2 -1 1 0 0 1 0 2 -2 4 8 4 -4
+# The answer array C[]
+ssbo 0:2 subdata int 0 8 8 8 8 8 8 8 8 8 8 8 8 8 8 8 8
+
+compute 16 1 1
+
+probe ssbo int 0:2 0 == 0 0 1 1 1 1 1 0 1 0 0 0 0 1 0 0
diff --git a/external/vulkancts/data/vulkan/amber/spirv_assembly/instruction/compute/signed_int_compare/uint_slessthanequal.amber b/external/vulkancts/data/vulkan/amber/spirv_assembly/instruction/compute/signed_int_compare/uint_slessthanequal.amber
new file mode 100644
index 0000000..736d42c
--- /dev/null
+++ b/external/vulkancts/data/vulkan/amber/spirv_assembly/instruction/compute/signed_int_compare/uint_slessthanequal.amber
@@ -0,0 +1,69 @@
+# Test SLessThanEqual with unsigned int params
+# Google bug b/73133282
+#
+# Derived from the following OpenCL C, but cleaned up to be more generic.
+#
+# kernel void foo(global int *A, global int*B, global int*C) __attribute__((reqd_work_group_size(1,1,1))) {
+# uint i = get_global_id(0);
+# C[i] = A[i] < B[i];
+# }
+
+[compute shader spirv]
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %18 "main" %gl_GlobalInvocationID
+ OpExecutionMode %18 LocalSize 1 1 1
+ OpSource OpenCL_C 120
+ OpDecorate %_runtimearr_uint ArrayStride 4
+ OpMemberDecorate %_struct_3 0 Offset 0
+ OpDecorate %_struct_3 BufferBlock
+ OpDecorate %gl_GlobalInvocationID BuiltIn GlobalInvocationId
+ OpDecorate %15 DescriptorSet 0
+ OpDecorate %15 Binding 0
+ OpDecorate %16 DescriptorSet 0
+ OpDecorate %16 Binding 1
+ OpDecorate %17 DescriptorSet 0
+ OpDecorate %17 Binding 2
+ %uint = OpTypeInt 32 0
+%_runtimearr_uint = OpTypeRuntimeArray %uint
+ %_struct_3 = OpTypeStruct %_runtimearr_uint
+%_ptr_Uniform__struct_3 = OpTypePointer Uniform %_struct_3
+ %void = OpTypeVoid
+ %6 = OpTypeFunction %void
+ %v3uint = OpTypeVector %uint 3
+%_ptr_Input_v3uint = OpTypePointer Input %v3uint
+%_ptr_Input_uint = OpTypePointer Input %uint
+%_ptr_Uniform_uint = OpTypePointer Uniform %uint
+ %bool = OpTypeBool
+ %uint_0 = OpConstant %uint 0
+ %uint_1 = OpConstant %uint 1
+%gl_GlobalInvocationID = OpVariable %_ptr_Input_v3uint Input
+ %15 = OpVariable %_ptr_Uniform__struct_3 Uniform
+ %16 = OpVariable %_ptr_Uniform__struct_3 Uniform
+ %17 = OpVariable %_ptr_Uniform__struct_3 Uniform
+ %18 = OpFunction %void None %6
+ %19 = OpLabel
+ %20 = OpAccessChain %_ptr_Input_uint %gl_GlobalInvocationID %uint_0
+ %21 = OpLoad %uint %20
+ %22 = OpAccessChain %_ptr_Uniform_uint %15 %uint_0 %21
+ %23 = OpLoad %uint %22
+ %24 = OpAccessChain %_ptr_Uniform_uint %16 %uint_0 %21
+ %25 = OpLoad %uint %24
+ %26 = OpSLessThanEqual %bool %23 %25
+ %27 = OpSelect %uint %26 %uint_1 %uint_0
+ %28 = OpAccessChain %_ptr_Uniform_uint %17 %uint_0 %21
+ OpStore %28 %27
+ OpReturn
+ OpFunctionEnd
+
+[test]
+# A[]
+ssbo 0:0 subdata int 0 -8 -7 -6 -5 -4 -3 -2 0 0 1 2 3 4 5 6 7
+# B[]
+ssbo 0:1 subdata int 0 -9 -7 -5 2 -1 1 0 0 1 0 2 -2 4 8 4 -4
+# The answer array C[]
+ssbo 0:2 subdata int 0 8 8 8 8 8 8 8 8 8 8 8 8 8 8 8 8
+
+compute 16 1 1
+
+probe ssbo int 0:2 0 == 0 1 1 1 1 1 1 1 1 0 1 0 1 1 0 0
diff --git a/external/vulkancts/modules/vulkan/amber/vktAmberTestCase.cpp b/external/vulkancts/modules/vulkan/amber/vktAmberTestCase.cpp
index d2c7497..d0710e8 100644
--- a/external/vulkancts/modules/vulkan/amber/vktAmberTestCase.cpp
+++ b/external/vulkancts/modules/vulkan/amber/vktAmberTestCase.cpp
@@ -76,9 +76,14 @@
{
getTestContext().getLog()
<< tcu::TestLog::Message
+ << "Failed to parse Amber test "
+ << readFilename
+ << ": "
<< r.Error()
<< "\n"
<< tcu::TestLog::EndMessage;
+ // TODO(dneto): Enhance Amber to not require this.
+ m_recipe->SetImpl(DE_NULL);
return false;
}
return true;
diff --git a/external/vulkancts/modules/vulkan/spirv_assembly/CMakeLists.txt b/external/vulkancts/modules/vulkan/spirv_assembly/CMakeLists.txt
index f092099..4c6f9ea 100644
--- a/external/vulkancts/modules/vulkan/spirv_assembly/CMakeLists.txt
+++ b/external/vulkancts/modules/vulkan/spirv_assembly/CMakeLists.txt
@@ -31,6 +31,8 @@
vktSpvAsmInstructionTests.hpp
vktSpvAsmPointerParameterTests.cpp
vktSpvAsmPointerParameterTests.hpp
+ vktSpvAsmSignedIntCompareTests.cpp
+ vktSpvAsmSignedIntCompareTests.hpp
vktSpvAsmTypeTests.cpp
vktSpvAsmTypeTests.hpp
vktSpvAsmTests.cpp
@@ -54,11 +56,21 @@
)
set(DEQP_VK_SPIRV_ASSEMBLY_LIBS
+ libamber
tcutil
vkutil
+ deqp-vk-amber
)
PCH(DEQP_VK_SPIRV_ASSEMBLY_SRCS ../pch.cpp)
+if (DE_COMPILER_IS_GCC OR DE_COMPILER_IS_CLANG)
+ set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -Wno-unknown-pragmas")
+elseif (DE_COMPILER_IS_MSC)
+ set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} /wd4068")
+endif()
+include_directories("../../../../amber/src/include")
+include_directories("../amber")
+
add_library(deqp-vk-spirv-assembly STATIC ${DEQP_VK_SPIRV_ASSEMBLY_SRCS})
target_link_libraries(deqp-vk-spirv-assembly ${DEQP_VK_SPIRV_ASSEMBLY_LIBS})
diff --git a/external/vulkancts/modules/vulkan/spirv_assembly/vktSpvAsmInstructionTests.cpp b/external/vulkancts/modules/vulkan/spirv_assembly/vktSpvAsmInstructionTests.cpp
index 3023827..096243b 100644
--- a/external/vulkancts/modules/vulkan/spirv_assembly/vktSpvAsmInstructionTests.cpp
+++ b/external/vulkancts/modules/vulkan/spirv_assembly/vktSpvAsmInstructionTests.cpp
@@ -71,6 +71,7 @@
#include "vktSpvAsmCompositeInsertTests.hpp"
#include "vktSpvAsmVaryingNameTests.hpp"
#include "vktSpvAsmWorkgroupMemoryTests.hpp"
+#include "vktSpvAsmSignedIntCompareTests.hpp"
#include <cmath>
#include <limits>
@@ -18412,6 +18413,7 @@
computeTests->addChild(createBoolGroup(testCtx));
computeTests->addChild(createWorkgroupMemoryComputeGroup(testCtx));
computeTests->addChild(createSpirvIdsAbuseGroup(testCtx));
+ computeTests->addChild(createSignedIntCompareGroup(testCtx));
graphicsTests->addChild(createCrossStageInterfaceTests(testCtx));
graphicsTests->addChild(createSpivVersionCheckTests(testCtx, !testComputePipeline));
diff --git a/external/vulkancts/modules/vulkan/spirv_assembly/vktSpvAsmSignedIntCompareTests.cpp b/external/vulkancts/modules/vulkan/spirv_assembly/vktSpvAsmSignedIntCompareTests.cpp
new file mode 100644
index 0000000..3a0509b
--- /dev/null
+++ b/external/vulkancts/modules/vulkan/spirv_assembly/vktSpvAsmSignedIntCompareTests.cpp
@@ -0,0 +1,89 @@
+/*------------------------------------------------------------------------
+ * Vulkan Conformance Tests
+ * ------------------------
+ *
+ * Copyright (c) 2019 Google LLC
+ * Copyright (c) 2019 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.
+ *
+ *//*!
+ * \file
+ * \brief SPIR-V signed int compare on unsigned scalars values. Google bug b/73133282
+ *//*--------------------------------------------------------------------*/
+
+#include <string>
+#include <amber/amber.h>
+
+#include "tcuDefs.hpp"
+
+#include "vktAmberTestCase.hpp"
+#include "vktSpvAsmSignedIntCompareTests.hpp"
+#include "vktTestGroupUtil.hpp"
+
+namespace vkt
+{
+namespace SpirVAssembly
+{
+namespace
+{
+
+void createSignedIntCompareTests (tcu::TestCaseGroup* tests, const char* data_dir)
+{
+ tcu::TestContext& testCtx = tests->getTestContext();
+
+ // Shader test files are saved in <path>/external/vulkancts/data/vulkan/amber/<data_dir>/<basename>.amber
+ struct Case {
+ const char *basename;
+ const char *description;
+ };
+ const Case cases[] =
+ {
+ { "uint_sgreaterthanequal", "32bit unsigned int with OpSGreaterThanEqual" },
+ { "uint_sgreaterthan", "32bit unsigned int with OpSGreaterThan" },
+ { "uint_slessthan", "32bit unsigned int with OpSLessThan" },
+ { "uint_slessthanequal", "32bit unsigned int with OpSLessThanEqual" }
+ // For testing fail-to-parse case:
+ //, { "foo", "Amber syntax error" }
+ };
+
+ for (unsigned i = 0; i < sizeof(cases)/sizeof(cases[0]) ; ++i)
+ {
+
+ cts_amber::AmberTestCase *testCase = new cts_amber::AmberTestCase(testCtx, cases[i].basename, cases[i].description);
+ // Make sure the input can be parsed before we use it.
+ std::string file = std::string(cases[i].basename) + ".amber";
+ if (testCase->parse(data_dir, file.c_str()))
+ {
+ tests->addChild(testCase);
+ }
+ else
+ {
+ delete testCase;
+ std::string message = "Failed to parse Amber test " + file + ". Check log for details";
+ TCU_THROW(Exception, message.c_str());
+ }
+ }
+}
+
+} // anonymous
+
+tcu::TestCaseGroup* createSignedIntCompareGroup (tcu::TestContext& testCtx)
+{
+ // Location of the Amber script files under the data/vulkan/amber source tree.
+ const char* data_dir = "spirv_assembly/instruction/compute/signed_int_compare";
+ return createTestGroup(testCtx, "signed_int_compare", "Signed int compare over uint values", createSignedIntCompareTests, data_dir);
+}
+
+} // SpirVAssembly
+} // vkt
diff --git a/external/vulkancts/modules/vulkan/spirv_assembly/vktSpvAsmSignedIntCompareTests.hpp b/external/vulkancts/modules/vulkan/spirv_assembly/vktSpvAsmSignedIntCompareTests.hpp
new file mode 100644
index 0000000..e78d41e
--- /dev/null
+++ b/external/vulkancts/modules/vulkan/spirv_assembly/vktSpvAsmSignedIntCompareTests.hpp
@@ -0,0 +1,40 @@
+#ifndef _VKTSPVASMSIGNEDINTCOMPARETESTS_HPP
+#define _VKTSPVASMSIGNEDINTCOMPARETESTS_HPP
+/*------------------------------------------------------------------------
+ * Vulkan Conformance Tests
+ * ------------------------
+ *
+ * Copyright (c) 2019 Google LLC
+ * Copyright (c) 2019 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.
+ *
+ *//*!
+ * \file
+ * \brief Functional signed integer compare tests using Amber
+ *//*--------------------------------------------------------------------*/
+
+#include "tcuDefs.hpp"
+#include "tcuTestCase.hpp"
+
+namespace vkt
+{
+namespace SpirVAssembly
+{
+
+tcu::TestCaseGroup* createSignedIntCompareGroup (tcu::TestContext& testCtx);
+
+} // SpirVAssembly
+} // vkt
+
+#endif // _VKTSPVASMSIGNEDINTCOMPARETESTS_HPP
diff --git a/external/vulkancts/mustpass/1.1.4/vk-default-no-waivers.txt b/external/vulkancts/mustpass/1.1.4/vk-default-no-waivers.txt
index 7d07bb8..e018d8a 100644
--- a/external/vulkancts/mustpass/1.1.4/vk-default-no-waivers.txt
+++ b/external/vulkancts/mustpass/1.1.4/vk-default-no-waivers.txt
@@ -228109,6 +228109,10 @@
dEQP-VK.spirv_assembly.instruction.compute.workgroup_memory.uint8
dEQP-VK.spirv_assembly.instruction.compute.spirv_ids_abuse.sparse_ids
dEQP-VK.spirv_assembly.instruction.compute.spirv_ids_abuse.lots_ids
+dEQP-VK.spirv_assembly.instruction.compute.signed_int_compare.uint_sgreaterthanequal
+dEQP-VK.spirv_assembly.instruction.compute.signed_int_compare.uint_sgreaterthan
+dEQP-VK.spirv_assembly.instruction.compute.signed_int_compare.uint_slessthan
+dEQP-VK.spirv_assembly.instruction.compute.signed_int_compare.uint_slessthanequal
dEQP-VK.spirv_assembly.instruction.graphics.cross_stage.basic_type.flat
dEQP-VK.spirv_assembly.instruction.graphics.cross_stage.basic_type.no_perspective
dEQP-VK.spirv_assembly.instruction.graphics.cross_stage.basic_type.relaxedprecision
diff --git a/external/vulkancts/mustpass/1.1.4/vk-default.txt b/external/vulkancts/mustpass/1.1.4/vk-default.txt
index 2686c7e..5553b51 100644
--- a/external/vulkancts/mustpass/1.1.4/vk-default.txt
+++ b/external/vulkancts/mustpass/1.1.4/vk-default.txt
@@ -228109,6 +228109,10 @@
dEQP-VK.spirv_assembly.instruction.compute.workgroup_memory.uint8
dEQP-VK.spirv_assembly.instruction.compute.spirv_ids_abuse.sparse_ids
dEQP-VK.spirv_assembly.instruction.compute.spirv_ids_abuse.lots_ids
+dEQP-VK.spirv_assembly.instruction.compute.signed_int_compare.uint_sgreaterthanequal
+dEQP-VK.spirv_assembly.instruction.compute.signed_int_compare.uint_sgreaterthan
+dEQP-VK.spirv_assembly.instruction.compute.signed_int_compare.uint_slessthan
+dEQP-VK.spirv_assembly.instruction.compute.signed_int_compare.uint_slessthanequal
dEQP-VK.spirv_assembly.instruction.graphics.cross_stage.basic_type.flat
dEQP-VK.spirv_assembly.instruction.graphics.cross_stage.basic_type.no_perspective
dEQP-VK.spirv_assembly.instruction.graphics.cross_stage.basic_type.relaxedprecision
