Add tests for OpPtrAccessChain on Workgroup storage
authorGraeme Leese <gleese@broadcom.com>
Wed, 1 May 2019 13:46:46 +0000 (14:46 +0100)
committerAlexander Galazin <Alexander.Galazin@arm.com>
Wed, 3 Jul 2019 13:42:13 +0000 (09:42 -0400)
OpPtrAccessChain is supposed to work differently on workgroup storage
than on any other storage class but there were no tests for it. Add some
very basic testing.

Components: Vulkan
New Tests: dEQP-VK.spirv_assembly.instruction.compute.ptr_access_chain.*

Change-Id: I91d95fb239d3be6bfb86db88f842e19101f5976d

AndroidGen.mk
android/cts/master/vk-master.txt
external/vulkancts/data/vulkan/amber/spirv_assembly/instruction/compute/ptr_access_chain/workgroup.amber [new file with mode: 0644]
external/vulkancts/data/vulkan/amber/spirv_assembly/instruction/compute/ptr_access_chain/workgroup_bad_stride.amber [new file with mode: 0644]
external/vulkancts/data/vulkan/amber/spirv_assembly/instruction/compute/ptr_access_chain/workgroup_no_stride.amber [new file with mode: 0644]
external/vulkancts/modules/vulkan/spirv_assembly/CMakeLists.txt
external/vulkancts/modules/vulkan/spirv_assembly/vktSpvAsmInstructionTests.cpp
external/vulkancts/modules/vulkan/spirv_assembly/vktSpvAsmPtrAccessChainTests.cpp [new file with mode: 0644]
external/vulkancts/modules/vulkan/spirv_assembly/vktSpvAsmPtrAccessChainTests.hpp [new file with mode: 0644]
external/vulkancts/mustpass/master/vk-default-no-waivers.txt
external/vulkancts/mustpass/master/vk-default.txt

index 31e8504..1403c63 100644 (file)
@@ -310,6 +310,7 @@ LOCAL_SRC_FILES := \
        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/vktSpvAsmPtrAccessChainTests.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 \
index 954fb53..b751989 100644 (file)
@@ -264445,6 +264445,9 @@ dEQP-VK.spirv_assembly.instruction.compute.unused_variables.variable_set_0_bindi
 dEQP-VK.spirv_assembly.instruction.compute.unused_variables.function_set_0_binding_5
 dEQP-VK.spirv_assembly.instruction.compute.unused_variables.variable_set_5_binding_5
 dEQP-VK.spirv_assembly.instruction.compute.unused_variables.function_set_5_binding_5
+dEQP-VK.spirv_assembly.instruction.compute.ptr_access_chain.workgroup
+dEQP-VK.spirv_assembly.instruction.compute.ptr_access_chain.workgroup_no_stride
+dEQP-VK.spirv_assembly.instruction.compute.ptr_access_chain.workgroup_bad_stride
 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/ptr_access_chain/workgroup.amber b/external/vulkancts/data/vulkan/amber/spirv_assembly/instruction/compute/ptr_access_chain/workgroup.amber
new file mode 100644 (file)
index 0000000..6d3855b
--- /dev/null
@@ -0,0 +1,111 @@
+# Test OpPtrAccessChain applied to workgroup memory.
+#
+# Derived from the following OpenCL C:
+#
+#int get_data(local int *d);
+#
+#int get_data(local int *d) {
+#   return d[1];
+#}
+#
+#kernel void foo(global int *A, global int*B, global int*C) __attribute__((reqd_work_group_size(16,1,1))) {
+#   local int data[17];
+#   uint i = get_local_id(0);
+#   data[i] = A[i] * B[i];
+#
+#   if (i == 0) data[16] = 0;
+#
+#   C[i] = get_data(&data[i]);
+#}
+#
+# Compiled with `clspv -no-inline-single -cl-opt-disable <X.clc>` with a (correct) ArrayStride decoration
+# added. This is allowed but not needed so provides a baseline for comparison.
+
+[compute shader spirv]
+               OpCapability Shader
+               OpCapability VariablePointers
+               OpExtension "SPV_KHR_storage_buffer_storage_class"
+               OpExtension "SPV_KHR_variable_pointers"
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %30 "main" %gl_LocalInvocationID
+               OpExecutionMode %30 LocalSize 16 1 1
+               OpSource OpenCL_C 120
+               OpDecorate %_runtimearr_uint ArrayStride 4
+               OpMemberDecorate %_struct_3 0 Offset 0
+               OpDecorate %_struct_3 Block
+               OpDecorate %gl_LocalInvocationID BuiltIn LocalInvocationId
+               OpDecorate %22 DescriptorSet 0
+               OpDecorate %22 Binding 0
+               OpDecorate %23 DescriptorSet 0
+               OpDecorate %23 Binding 1
+               OpDecorate %24 DescriptorSet 0
+               OpDecorate %24 Binding 2
+               OpDecorate %_arr_uint_uint_17 ArrayStride 4
+               OpDecorate %_ptr_Workgroup_uint ArrayStride 4
+       %uint = OpTypeInt 32 0
+%_runtimearr_uint = OpTypeRuntimeArray %uint
+  %_struct_3 = OpTypeStruct %_runtimearr_uint
+%_ptr_StorageBuffer__struct_3 = OpTypePointer StorageBuffer %_struct_3
+%_ptr_Workgroup_uint = OpTypePointer Workgroup %uint
+          %6 = OpTypeFunction %uint %_ptr_Workgroup_uint
+       %void = OpTypeVoid
+          %8 = OpTypeFunction %void
+     %v3uint = OpTypeVector %uint 3
+%_ptr_Input_v3uint = OpTypePointer Input %v3uint
+%_ptr_Input_uint = OpTypePointer Input %uint
+%_ptr_StorageBuffer_uint = OpTypePointer StorageBuffer %uint
+    %uint_17 = OpConstant %uint 17
+%_arr_uint_uint_17 = OpTypeArray %uint %uint_17
+%_ptr_Workgroup__arr_uint_uint_17 = OpTypePointer Workgroup %_arr_uint_uint_17
+       %bool = OpTypeBool
+     %uint_1 = OpConstant %uint 1
+     %uint_0 = OpConstant %uint 0
+    %uint_16 = OpConstant %uint 16
+         %20 = OpVariable %_ptr_Workgroup__arr_uint_uint_17 Workgroup
+%gl_LocalInvocationID = OpVariable %_ptr_Input_v3uint Input
+         %22 = OpVariable %_ptr_StorageBuffer__struct_3 StorageBuffer
+         %23 = OpVariable %_ptr_StorageBuffer__struct_3 StorageBuffer
+         %24 = OpVariable %_ptr_StorageBuffer__struct_3 StorageBuffer
+         %25 = OpFunction %uint Pure %6
+         %26 = OpFunctionParameter %_ptr_Workgroup_uint
+         %27 = OpLabel
+         %28 = OpPtrAccessChain %_ptr_Workgroup_uint %26 %uint_1
+         %29 = OpLoad %uint %28
+               OpReturnValue %29
+               OpFunctionEnd
+         %30 = OpFunction %void None %8
+         %31 = OpLabel
+         %32 = OpAccessChain %_ptr_Input_uint %gl_LocalInvocationID %uint_0
+         %33 = OpLoad %uint %32
+         %34 = OpAccessChain %_ptr_StorageBuffer_uint %22 %uint_0 %33
+         %35 = OpLoad %uint %34
+         %36 = OpAccessChain %_ptr_StorageBuffer_uint %23 %uint_0 %33
+         %37 = OpLoad %uint %36
+         %38 = OpIMul %uint %37 %35
+         %39 = OpAccessChain %_ptr_Workgroup_uint %20 %33
+               OpStore %39 %38
+         %40 = OpIEqual %bool %33 %uint_0
+               OpSelectionMerge %43 None
+               OpBranchConditional %40 %41 %43
+         %41 = OpLabel
+         %42 = OpAccessChain %_ptr_Workgroup_uint %20 %uint_16
+               OpStore %42 %uint_0
+               OpBranch %43
+         %43 = OpLabel
+         %44 = OpFunctionCall %uint %25 %39
+         %45 = OpAccessChain %_ptr_StorageBuffer_uint %24 %uint_0 %33
+               OpStore %45 %44
+               OpReturn
+               OpFunctionEnd
+
+[test]
+# A[]
+ssbo 0:0 subdata int 0  0 1 2 3 4 5 6 7 8 9 10 11 12 13 14 15
+# B[]
+ssbo 0:1 subdata int 0  1 1 1 1 1 1 1 1 1 1 1 1 1 1 1 1
+# The answer array C[]
+ssbo 0:2 subdata int 0  -1 -1 -1 -1 -1 -1 -1 -1 -1 -1 -1 -1 -1 -1 -1 -1
+
+compute 1 1 1
+
+probe ssbo int 0:2 0 ==  1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 0
diff --git a/external/vulkancts/data/vulkan/amber/spirv_assembly/instruction/compute/ptr_access_chain/workgroup_bad_stride.amber b/external/vulkancts/data/vulkan/amber/spirv_assembly/instruction/compute/ptr_access_chain/workgroup_bad_stride.amber
new file mode 100644 (file)
index 0000000..f951de6
--- /dev/null
@@ -0,0 +1,111 @@
+# Test OpPtrAccessChain applied to workgroup memory.
+#
+# Derived from the following OpenCL C:
+#
+#int get_data(local int *d);
+#
+#int get_data(local int *d) {
+#   return d[1];
+#}
+#
+#kernel void foo(global int *A, global int*B, global int*C) __attribute__((reqd_work_group_size(16,1,1))) {
+#   local int data[17];
+#   uint i = get_local_id(0);
+#   data[i] = A[i] * B[i];
+#
+#   if (i == 0) data[16] = 0;
+#
+#   C[i] = get_data(&data[i]);
+#}
+#
+# Compiled with `clspv -no-inline-single -cl-opt-disable <X.clc>` with an incorrect ArrayStride decoration
+# added. This decoration should be ignored, so it should give the same results as ArrayStride == 4.
+
+[compute shader spirv]
+               OpCapability Shader
+               OpCapability VariablePointers
+               OpExtension "SPV_KHR_storage_buffer_storage_class"
+               OpExtension "SPV_KHR_variable_pointers"
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %30 "main" %gl_LocalInvocationID
+               OpExecutionMode %30 LocalSize 16 1 1
+               OpSource OpenCL_C 120
+               OpDecorate %_runtimearr_uint ArrayStride 4
+               OpMemberDecorate %_struct_3 0 Offset 0
+               OpDecorate %_struct_3 Block
+               OpDecorate %gl_LocalInvocationID BuiltIn LocalInvocationId
+               OpDecorate %22 DescriptorSet 0
+               OpDecorate %22 Binding 0
+               OpDecorate %23 DescriptorSet 0
+               OpDecorate %23 Binding 1
+               OpDecorate %24 DescriptorSet 0
+               OpDecorate %24 Binding 2
+               OpDecorate %_arr_uint_uint_17 ArrayStride 4
+               OpDecorate %_ptr_Workgroup_uint ArrayStride 8
+       %uint = OpTypeInt 32 0
+%_runtimearr_uint = OpTypeRuntimeArray %uint
+  %_struct_3 = OpTypeStruct %_runtimearr_uint
+%_ptr_StorageBuffer__struct_3 = OpTypePointer StorageBuffer %_struct_3
+%_ptr_Workgroup_uint = OpTypePointer Workgroup %uint
+          %6 = OpTypeFunction %uint %_ptr_Workgroup_uint
+       %void = OpTypeVoid
+          %8 = OpTypeFunction %void
+     %v3uint = OpTypeVector %uint 3
+%_ptr_Input_v3uint = OpTypePointer Input %v3uint
+%_ptr_Input_uint = OpTypePointer Input %uint
+%_ptr_StorageBuffer_uint = OpTypePointer StorageBuffer %uint
+    %uint_17 = OpConstant %uint 17
+%_arr_uint_uint_17 = OpTypeArray %uint %uint_17
+%_ptr_Workgroup__arr_uint_uint_17 = OpTypePointer Workgroup %_arr_uint_uint_17
+       %bool = OpTypeBool
+     %uint_1 = OpConstant %uint 1
+     %uint_0 = OpConstant %uint 0
+    %uint_16 = OpConstant %uint 16
+         %20 = OpVariable %_ptr_Workgroup__arr_uint_uint_17 Workgroup
+%gl_LocalInvocationID = OpVariable %_ptr_Input_v3uint Input
+         %22 = OpVariable %_ptr_StorageBuffer__struct_3 StorageBuffer
+         %23 = OpVariable %_ptr_StorageBuffer__struct_3 StorageBuffer
+         %24 = OpVariable %_ptr_StorageBuffer__struct_3 StorageBuffer
+         %25 = OpFunction %uint Pure %6
+         %26 = OpFunctionParameter %_ptr_Workgroup_uint
+         %27 = OpLabel
+         %28 = OpPtrAccessChain %_ptr_Workgroup_uint %26 %uint_1
+         %29 = OpLoad %uint %28
+               OpReturnValue %29
+               OpFunctionEnd
+         %30 = OpFunction %void None %8
+         %31 = OpLabel
+         %32 = OpAccessChain %_ptr_Input_uint %gl_LocalInvocationID %uint_0
+         %33 = OpLoad %uint %32
+         %34 = OpAccessChain %_ptr_StorageBuffer_uint %22 %uint_0 %33
+         %35 = OpLoad %uint %34
+         %36 = OpAccessChain %_ptr_StorageBuffer_uint %23 %uint_0 %33
+         %37 = OpLoad %uint %36
+         %38 = OpIMul %uint %37 %35
+         %39 = OpAccessChain %_ptr_Workgroup_uint %20 %33
+               OpStore %39 %38
+         %40 = OpIEqual %bool %33 %uint_0
+               OpSelectionMerge %43 None
+               OpBranchConditional %40 %41 %43
+         %41 = OpLabel
+         %42 = OpAccessChain %_ptr_Workgroup_uint %20 %uint_16
+               OpStore %42 %uint_0
+               OpBranch %43
+         %43 = OpLabel
+         %44 = OpFunctionCall %uint %25 %39
+         %45 = OpAccessChain %_ptr_StorageBuffer_uint %24 %uint_0 %33
+               OpStore %45 %44
+               OpReturn
+               OpFunctionEnd
+
+[test]
+# A[]
+ssbo 0:0 subdata int 0  0 1 2 3 4 5 6 7 8 9 10 11 12 13 14 15
+# B[]
+ssbo 0:1 subdata int 0  1 1 1 1 1 1 1 1 1 1 1 1 1 1 1 1
+# The answer array C[]
+ssbo 0:2 subdata int 0  -1 -1 -1 -1 -1 -1 -1 -1 -1 -1 -1 -1 -1 -1 -1 -1
+
+compute 1 1 1
+
+probe ssbo int 0:2 0 ==  1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 0
diff --git a/external/vulkancts/data/vulkan/amber/spirv_assembly/instruction/compute/ptr_access_chain/workgroup_no_stride.amber b/external/vulkancts/data/vulkan/amber/spirv_assembly/instruction/compute/ptr_access_chain/workgroup_no_stride.amber
new file mode 100644 (file)
index 0000000..5ca6b49
--- /dev/null
@@ -0,0 +1,109 @@
+# Test OpPtrAccessChain applied to workgroup memory.
+#
+# Derived from the following OpenCL C:
+#
+#int get_data(local int *d);
+#
+#int get_data(local int *d) {
+#   return d[1];
+#}
+#
+#kernel void foo(global int *A, global int*B, global int*C) __attribute__((reqd_work_group_size(16,1,1))) {
+#   local int data[17];
+#   uint i = get_local_id(0);
+#   data[i] = A[i] * B[i];
+#
+#   if (i == 0) data[16] = 0;
+#
+#   C[i] = get_data(&data[i]);
+#}
+#
+# Compiled with `clspv -no-inline-single -cl-opt-disable <X.clc>`
+
+[compute shader spirv]
+               OpCapability Shader
+               OpCapability VariablePointers
+               OpExtension "SPV_KHR_storage_buffer_storage_class"
+               OpExtension "SPV_KHR_variable_pointers"
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %30 "main" %gl_LocalInvocationID
+               OpExecutionMode %30 LocalSize 16 1 1
+               OpSource OpenCL_C 120
+               OpDecorate %_runtimearr_uint ArrayStride 4
+               OpMemberDecorate %_struct_3 0 Offset 0
+               OpDecorate %_struct_3 Block
+               OpDecorate %gl_LocalInvocationID BuiltIn LocalInvocationId
+               OpDecorate %22 DescriptorSet 0
+               OpDecorate %22 Binding 0
+               OpDecorate %23 DescriptorSet 0
+               OpDecorate %23 Binding 1
+               OpDecorate %24 DescriptorSet 0
+               OpDecorate %24 Binding 2
+               OpDecorate %_arr_uint_uint_17 ArrayStride 4
+       %uint = OpTypeInt 32 0
+%_runtimearr_uint = OpTypeRuntimeArray %uint
+  %_struct_3 = OpTypeStruct %_runtimearr_uint
+%_ptr_StorageBuffer__struct_3 = OpTypePointer StorageBuffer %_struct_3
+%_ptr_Workgroup_uint = OpTypePointer Workgroup %uint
+          %6 = OpTypeFunction %uint %_ptr_Workgroup_uint
+       %void = OpTypeVoid
+          %8 = OpTypeFunction %void
+     %v3uint = OpTypeVector %uint 3
+%_ptr_Input_v3uint = OpTypePointer Input %v3uint
+%_ptr_Input_uint = OpTypePointer Input %uint
+%_ptr_StorageBuffer_uint = OpTypePointer StorageBuffer %uint
+    %uint_17 = OpConstant %uint 17
+%_arr_uint_uint_17 = OpTypeArray %uint %uint_17
+%_ptr_Workgroup__arr_uint_uint_17 = OpTypePointer Workgroup %_arr_uint_uint_17
+       %bool = OpTypeBool
+     %uint_1 = OpConstant %uint 1
+     %uint_0 = OpConstant %uint 0
+    %uint_16 = OpConstant %uint 16
+         %20 = OpVariable %_ptr_Workgroup__arr_uint_uint_17 Workgroup
+%gl_LocalInvocationID = OpVariable %_ptr_Input_v3uint Input
+         %22 = OpVariable %_ptr_StorageBuffer__struct_3 StorageBuffer
+         %23 = OpVariable %_ptr_StorageBuffer__struct_3 StorageBuffer
+         %24 = OpVariable %_ptr_StorageBuffer__struct_3 StorageBuffer
+         %25 = OpFunction %uint Pure %6
+         %26 = OpFunctionParameter %_ptr_Workgroup_uint
+         %27 = OpLabel
+         %28 = OpPtrAccessChain %_ptr_Workgroup_uint %26 %uint_1
+         %29 = OpLoad %uint %28
+               OpReturnValue %29
+               OpFunctionEnd
+         %30 = OpFunction %void None %8
+         %31 = OpLabel
+         %32 = OpAccessChain %_ptr_Input_uint %gl_LocalInvocationID %uint_0
+         %33 = OpLoad %uint %32
+         %34 = OpAccessChain %_ptr_StorageBuffer_uint %22 %uint_0 %33
+         %35 = OpLoad %uint %34
+         %36 = OpAccessChain %_ptr_StorageBuffer_uint %23 %uint_0 %33
+         %37 = OpLoad %uint %36
+         %38 = OpIMul %uint %37 %35
+         %39 = OpAccessChain %_ptr_Workgroup_uint %20 %33
+               OpStore %39 %38
+         %40 = OpIEqual %bool %33 %uint_0
+               OpSelectionMerge %43 None
+               OpBranchConditional %40 %41 %43
+         %41 = OpLabel
+         %42 = OpAccessChain %_ptr_Workgroup_uint %20 %uint_16
+               OpStore %42 %uint_0
+               OpBranch %43
+         %43 = OpLabel
+         %44 = OpFunctionCall %uint %25 %39
+         %45 = OpAccessChain %_ptr_StorageBuffer_uint %24 %uint_0 %33
+               OpStore %45 %44
+               OpReturn
+               OpFunctionEnd
+
+[test]
+# A[]
+ssbo 0:0 subdata int 0  0 1 2 3 4 5 6 7 8 9 10 11 12 13 14 15
+# B[]
+ssbo 0:1 subdata int 0  1 1 1 1 1 1 1 1 1 1 1 1 1 1 1 1
+# The answer array C[]
+ssbo 0:2 subdata int 0  -1 -1 -1 -1 -1 -1 -1 -1 -1 -1 -1 -1 -1 -1 -1 -1
+
+compute 1 1 1
+
+probe ssbo int 0:2 0 ==  1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 0
index c989bd4..acbaad7 100644 (file)
@@ -53,6 +53,8 @@ set(DEQP_VK_SPIRV_ASSEMBLY_SRCS
        vktSpvAsmVaryingNameTests.hpp
        vktSpvAsmWorkgroupMemoryTests.cpp
        vktSpvAsmWorkgroupMemoryTests.hpp
+       vktSpvAsmPtrAccessChainTests.cpp
+       vktSpvAsmPtrAccessChainTests.hpp
        )
 
 set(DEQP_VK_SPIRV_ASSEMBLY_LIBS
index f93f583..aa102b3 100644 (file)
@@ -73,6 +73,7 @@
 #include "vktSpvAsmVaryingNameTests.hpp"
 #include "vktSpvAsmWorkgroupMemoryTests.hpp"
 #include "vktSpvAsmSignedIntCompareTests.hpp"
+#include "vktSpvAsmPtrAccessChainTests.hpp"
 
 #include <cmath>
 #include <limits>
@@ -18659,6 +18660,7 @@ tcu::TestCaseGroup* createInstructionTests (tcu::TestContext& testCtx)
        computeTests->addChild(createSpirvIdsAbuseGroup(testCtx));
        computeTests->addChild(createSignedIntCompareGroup(testCtx));
        computeTests->addChild(createUnusedVariableComputeTests(testCtx));
+       computeTests->addChild(createPtrAccessChainGroup(testCtx));
 
        graphicsTests->addChild(createCrossStageInterfaceTests(testCtx));
        graphicsTests->addChild(createSpivVersionCheckTests(testCtx, !testComputePipeline));
diff --git a/external/vulkancts/modules/vulkan/spirv_assembly/vktSpvAsmPtrAccessChainTests.cpp b/external/vulkancts/modules/vulkan/spirv_assembly/vktSpvAsmPtrAccessChainTests.cpp
new file mode 100644 (file)
index 0000000..b2cca24
--- /dev/null
@@ -0,0 +1,69 @@
+/*------------------------------------------------------------------------
+ * 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.
+ *
+ *--------------------------------------------------------------------*/
+
+#include <string>
+
+#include "vktAmberTestCaseUtil.hpp"
+#include "vktSpvAsmPtrAccessChainTests.hpp"
+
+namespace vkt
+{
+namespace SpirVAssembly
+{
+namespace
+{
+
+void createTests (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[] =
+       {
+               { "workgroup", "OpPtrAccessChain with correct ArrayStride decoration" },
+               { "workgroup_no_stride", "OpPtrAccessChain with no ArrayStride decoration" },
+               { "workgroup_bad_stride", "OpPtrAccessChain with incorrect ArrayStride decoration" },
+       };
+
+       for (unsigned i = 0; i < sizeof(cases)/sizeof(cases[0]) ; ++i)
+       {
+               std::string                                     file            = std::string(cases[i].basename) + ".amber";
+               cts_amber::AmberTestCase        *testCase       = cts_amber::createAmberTestCase(testCtx, cases[i].basename, cases[i].description, data_dir, file);
+
+               tests->addChild(testCase);
+       }
+}
+
+} // anonymous
+
+tcu::TestCaseGroup* createPtrAccessChainGroup (tcu::TestContext& testCtx)
+{
+       // Location of the Amber script files under the data/vulkan/amber source tree.
+       const char* data_dir = "spirv_assembly/instruction/compute/ptr_access_chain";
+       return createTestGroup(testCtx, "ptr_access_chain", "OpPtrAccessChain edge cases", createTests, data_dir);
+}
+
+} // SpirVAssembly
+} // vkt
diff --git a/external/vulkancts/modules/vulkan/spirv_assembly/vktSpvAsmPtrAccessChainTests.hpp b/external/vulkancts/modules/vulkan/spirv_assembly/vktSpvAsmPtrAccessChainTests.hpp
new file mode 100644 (file)
index 0000000..bbbd633
--- /dev/null
@@ -0,0 +1,37 @@
+#ifndef _VKTSPVASMPTRACCESSCHAINTESTS_HPP
+#define _VKTSPVASMPTRACCESSCHAINTESTS_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.
+ *
+ *--------------------------------------------------------------------*/
+
+#include "tcuDefs.hpp"
+#include "tcuTestCase.hpp"
+
+namespace vkt
+{
+namespace SpirVAssembly
+{
+
+tcu::TestCaseGroup*    createPtrAccessChainGroup (tcu::TestContext& testCtx);
+
+} // SpirVAssembly
+} // vkt
+
+#endif // _VKTSPVASMPTRACCESSCHAINTESTS_HPP
index 7f1248c..b8a9d3b 100644 (file)
@@ -264459,6 +264459,9 @@ dEQP-VK.spirv_assembly.instruction.compute.unused_variables.variable_set_0_bindi
 dEQP-VK.spirv_assembly.instruction.compute.unused_variables.function_set_0_binding_5
 dEQP-VK.spirv_assembly.instruction.compute.unused_variables.variable_set_5_binding_5
 dEQP-VK.spirv_assembly.instruction.compute.unused_variables.function_set_5_binding_5
+dEQP-VK.spirv_assembly.instruction.compute.ptr_access_chain.workgroup
+dEQP-VK.spirv_assembly.instruction.compute.ptr_access_chain.workgroup_no_stride
+dEQP-VK.spirv_assembly.instruction.compute.ptr_access_chain.workgroup_bad_stride
 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
index efe560d..9daf457 100644 (file)
@@ -264459,6 +264459,9 @@ dEQP-VK.spirv_assembly.instruction.compute.unused_variables.variable_set_0_bindi
 dEQP-VK.spirv_assembly.instruction.compute.unused_variables.function_set_0_binding_5
 dEQP-VK.spirv_assembly.instruction.compute.unused_variables.variable_set_5_binding_5
 dEQP-VK.spirv_assembly.instruction.compute.unused_variables.function_set_5_binding_5
+dEQP-VK.spirv_assembly.instruction.compute.ptr_access_chain.workgroup
+dEQP-VK.spirv_assembly.instruction.compute.ptr_access_chain.workgroup_no_stride
+dEQP-VK.spirv_assembly.instruction.compute.ptr_access_chain.workgroup_bad_stride
 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