From f364013612a9a5888a9e0f90bf8d9086eadfc746 Mon Sep 17 00:00:00 2001 From: Graeme Leese Date: Wed, 1 May 2019 14:46:46 +0100 Subject: [PATCH] Add tests for OpPtrAccessChain on Workgroup storage 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 | 1 + android/cts/master/vk-master.txt | 3 + .../compute/ptr_access_chain/workgroup.amber | 111 +++++++++++++++++++++ .../ptr_access_chain/workgroup_bad_stride.amber | 111 +++++++++++++++++++++ .../ptr_access_chain/workgroup_no_stride.amber | 109 ++++++++++++++++++++ .../modules/vulkan/spirv_assembly/CMakeLists.txt | 2 + .../spirv_assembly/vktSpvAsmInstructionTests.cpp | 2 + .../vktSpvAsmPtrAccessChainTests.cpp | 69 +++++++++++++ .../vktSpvAsmPtrAccessChainTests.hpp | 37 +++++++ .../mustpass/master/vk-default-no-waivers.txt | 3 + external/vulkancts/mustpass/master/vk-default.txt | 3 + 11 files changed, 451 insertions(+) create mode 100644 external/vulkancts/data/vulkan/amber/spirv_assembly/instruction/compute/ptr_access_chain/workgroup.amber create mode 100644 external/vulkancts/data/vulkan/amber/spirv_assembly/instruction/compute/ptr_access_chain/workgroup_bad_stride.amber create mode 100644 external/vulkancts/data/vulkan/amber/spirv_assembly/instruction/compute/ptr_access_chain/workgroup_no_stride.amber create mode 100644 external/vulkancts/modules/vulkan/spirv_assembly/vktSpvAsmPtrAccessChainTests.cpp create mode 100644 external/vulkancts/modules/vulkan/spirv_assembly/vktSpvAsmPtrAccessChainTests.hpp diff --git a/AndroidGen.mk b/AndroidGen.mk index 31e8504..1403c63 100644 --- a/AndroidGen.mk +++ b/AndroidGen.mk @@ -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 \ diff --git a/android/cts/master/vk-master.txt b/android/cts/master/vk-master.txt index 954fb53..b751989 100644 --- a/android/cts/master/vk-master.txt +++ b/android/cts/master/vk-master.txt @@ -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 index 0000000..6d3855b --- /dev/null +++ b/external/vulkancts/data/vulkan/amber/spirv_assembly/instruction/compute/ptr_access_chain/workgroup.amber @@ -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 ` 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 index 0000000..f951de6 --- /dev/null +++ b/external/vulkancts/data/vulkan/amber/spirv_assembly/instruction/compute/ptr_access_chain/workgroup_bad_stride.amber @@ -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 ` 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 index 0000000..5ca6b49 --- /dev/null +++ b/external/vulkancts/data/vulkan/amber/spirv_assembly/instruction/compute/ptr_access_chain/workgroup_no_stride.amber @@ -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 ` + +[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 diff --git a/external/vulkancts/modules/vulkan/spirv_assembly/CMakeLists.txt b/external/vulkancts/modules/vulkan/spirv_assembly/CMakeLists.txt index c989bd4..acbaad7 100644 --- a/external/vulkancts/modules/vulkan/spirv_assembly/CMakeLists.txt +++ b/external/vulkancts/modules/vulkan/spirv_assembly/CMakeLists.txt @@ -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 diff --git a/external/vulkancts/modules/vulkan/spirv_assembly/vktSpvAsmInstructionTests.cpp b/external/vulkancts/modules/vulkan/spirv_assembly/vktSpvAsmInstructionTests.cpp index f93f583..aa102b3 100644 --- a/external/vulkancts/modules/vulkan/spirv_assembly/vktSpvAsmInstructionTests.cpp +++ b/external/vulkancts/modules/vulkan/spirv_assembly/vktSpvAsmInstructionTests.cpp @@ -73,6 +73,7 @@ #include "vktSpvAsmVaryingNameTests.hpp" #include "vktSpvAsmWorkgroupMemoryTests.hpp" #include "vktSpvAsmSignedIntCompareTests.hpp" +#include "vktSpvAsmPtrAccessChainTests.hpp" #include #include @@ -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 index 0000000..b2cca24 --- /dev/null +++ b/external/vulkancts/modules/vulkan/spirv_assembly/vktSpvAsmPtrAccessChainTests.cpp @@ -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 + +#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 /external/vulkancts/data/vulkan/amber//.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 index 0000000..bbbd633 --- /dev/null +++ b/external/vulkancts/modules/vulkan/spirv_assembly/vktSpvAsmPtrAccessChainTests.hpp @@ -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 diff --git a/external/vulkancts/mustpass/master/vk-default-no-waivers.txt b/external/vulkancts/mustpass/master/vk-default-no-waivers.txt index 7f1248c..b8a9d3b 100644 --- a/external/vulkancts/mustpass/master/vk-default-no-waivers.txt +++ b/external/vulkancts/mustpass/master/vk-default-no-waivers.txt @@ -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 diff --git a/external/vulkancts/mustpass/master/vk-default.txt b/external/vulkancts/mustpass/master/vk-default.txt index efe560d..9daf457 100644 --- a/external/vulkancts/mustpass/master/vk-default.txt +++ b/external/vulkancts/mustpass/master/vk-default.txt @@ -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 -- 2.7.4