AuroraMiddleware/SPIRV-Tools
1
0
Fork 0
mirror of https://github.com/KhronosGroup/SPIRV-Tools synced 2024-11-26 05:10:05 +00:00
Code Issues Packages Projects Releases Wiki Activity

Fixes #1470. Don't restrict WGS storage class

Browse Source 
* Removed restriction that workgroup size can only be on Input storage
class
* added test
This commit is contained in:
Alan Baker 2018-04-12 09:22:34 -04:00
parent bc648fd76a
commit c522b697bf
2 changed files with 23 additions and 13 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

13
source/validate_builtins.cpp
View File

@ -2123,19 +2123,6 @@ spv_result_t BuiltInsValidator::ValidateWorkgroupSizeAtReference(
const Instruction& referenced_inst,
const Instruction& referenced_from_inst) {
if (spvIsVulkanEnv(_.context()->target_env)) {
const SpvStorageClass storage_class = GetStorageClass(referenced_from_inst);
if (storage_class != SpvStorageClassMax &&
storage_class != SpvStorageClassInput) {
return _.diag(SPV_ERROR_INVALID_DATA)
<< "Vulkan spec allows BuiltIn "
<< _.grammar().lookupOperandName(SPV_OPERAND_TYPE_BUILT_IN,
decoration.params()[0])
<< " to be only used for variables with Input storage class. "
<< GetReferenceDesc(decoration, built_in_inst, referenced_inst,
referenced_from_inst)
<< " " << GetStorageClassDesc(referenced_from_inst);
}
for (const SpvExecutionModel execution_model : execution_models_) {
if (execution_model != SpvExecutionModelGLCompute) {
return _.diag(SPV_ERROR_INVALID_DATA)

23
test/val/val_builtins_test.cpp
View File

@ -1656,6 +1656,29 @@ OpDecorate %workgroup_size BuiltIn WorkgroupSize
"(OpConstantComposite) has components with bit width 64."));
}
TEST_F(ValidateBuiltIns, WorkgroupSizePrivateVar) {
CodeGenerator generator = GetDefaultShaderCodeGenerator();
generator.before_types_ = R"(
OpDecorate %workgroup_size BuiltIn WorkgroupSize
)";
generator.after_types_ = R"(
%workgroup_size = OpConstantComposite %u32vec3 %u32_1 %u32_1 %u32_1
%private_ptr_u32vec3 = OpTypePointer Private %u32vec3
%var = OpVariable %private_ptr_u32vec3 Private %workgroup_size
)";
EntryPoint entry_point;
entry_point.name = "main";
entry_point.execution_model = "GLCompute";
entry_point.body = R"(
)";
generator.entry_points_.push_back(std::move(entry_point));
CompileSuccessfully(generator.Build(), SPV_ENV_VULKAN_1_0);
ASSERT_EQ(SPV_SUCCESS, ValidateInstructions(SPV_ENV_VULKAN_1_0));
}
TEST_F(ValidateBuiltIns, GeometryPositionInOutSuccess) {
CodeGenerator generator = GetDefaultShaderCodeGenerator();