X-Git-Url: https://git.libre-soc.org/?a=blobdiff_plain;f=src%2Fcompiler%2Fspirv%2Fspirv_to_nir.c;h=7dc6bc914d6f9374440702959f9d0f54e5758fb1;hb=44227453ec03f5462f1cff5760909a9dba95c61a;hp=9d2f57cef94e803a884db4324d7d9c698cd348f0;hpb=54d7fca077d051af0b91f0684cdfa0055915b917;p=mesa.git diff --git a/src/compiler/spirv/spirv_to_nir.c b/src/compiler/spirv/spirv_to_nir.c index 9d2f57cef94..7dc6bc914d6 100644 --- a/src/compiler/spirv/spirv_to_nir.c +++ b/src/compiler/spirv/spirv_to_nir.c @@ -383,19 +383,20 @@ static void vtn_handle_extension(struct vtn_builder *b, SpvOp opcode, const uint32_t *w, unsigned count) { + const char *ext = (const char *)&w[2]; switch (opcode) { case SpvOpExtInstImport: { struct vtn_value *val = vtn_push_value(b, w[1], vtn_value_type_extension); - if (strcmp((const char *)&w[2], "GLSL.std.450") == 0) { + if (strcmp(ext, "GLSL.std.450") == 0) { val->ext_handler = vtn_handle_glsl450_instruction; - } else if ((strcmp((const char *)&w[2], "SPV_AMD_gcn_shader") == 0) + } else if ((strcmp(ext, "SPV_AMD_gcn_shader") == 0) && (b->options && b->options->caps.gcn_shader)) { val->ext_handler = vtn_handle_amd_gcn_shader_instruction; - } else if ((strcmp((const char *)&w[2], "SPV_AMD_shader_trinary_minmax") == 0) + } else if ((strcmp(ext, "SPV_AMD_shader_trinary_minmax") == 0) && (b->options && b->options->caps.trinary_minmax)) { val->ext_handler = vtn_handle_amd_shader_trinary_minmax_instruction; } else { - vtn_fail("Unsupported extension"); + vtn_fail("Unsupported extension: %s", ext); } break; } @@ -494,15 +495,19 @@ vtn_handle_decoration(struct vtn_builder *b, SpvOp opcode, case SpvOpDecorate: case SpvOpMemberDecorate: + case SpvOpDecorateStringGOOGLE: + case SpvOpMemberDecorateStringGOOGLE: case SpvOpExecutionMode: { struct vtn_value *val = vtn_untyped_value(b, target); struct vtn_decoration *dec = rzalloc(b, struct vtn_decoration); switch (opcode) { case SpvOpDecorate: + case SpvOpDecorateStringGOOGLE: dec->scope = VTN_DEC_DECORATION; break; case SpvOpMemberDecorate: + case SpvOpMemberDecorateStringGOOGLE: dec->scope = VTN_DEC_STRUCT_MEMBER0 + *(w++); vtn_fail_if(dec->scope < VTN_DEC_STRUCT_MEMBER0, /* overflow */ "Member argument of OpMemberDecorate too large"); @@ -667,6 +672,16 @@ mutable_matrix_member(struct vtn_builder *b, struct vtn_type *type, int member) return type; } +static void +vtn_handle_access_qualifier(struct vtn_builder *b, struct vtn_type *type, + int member, enum gl_access_qualifier access) +{ + type->members[member] = vtn_type_copy(b, type->members[member]); + type = type->members[member]; + + type->access |= access; +} + static void struct_member_decoration_cb(struct vtn_builder *b, struct vtn_value *val, int member, @@ -680,13 +695,21 @@ struct_member_decoration_cb(struct vtn_builder *b, assert(member < ctx->num_fields); switch (dec->decoration) { + case SpvDecorationRelaxedPrecision: + case SpvDecorationUniform: + break; /* FIXME: Do nothing with this for now. */ case SpvDecorationNonWritable: + vtn_handle_access_qualifier(b, ctx->type, member, ACCESS_NON_WRITEABLE); + break; case SpvDecorationNonReadable: - case SpvDecorationRelaxedPrecision: + vtn_handle_access_qualifier(b, ctx->type, member, ACCESS_NON_READABLE); + break; case SpvDecorationVolatile: + vtn_handle_access_qualifier(b, ctx->type, member, ACCESS_VOLATILE); + break; case SpvDecorationCoherent: - case SpvDecorationUniform: - break; /* FIXME: Do nothing with this for now. */ + vtn_handle_access_qualifier(b, ctx->type, member, ACCESS_COHERENT); + break; case SpvDecorationNoPerspective: ctx->fields[member].interpolation = INTERP_MODE_NOPERSPECTIVE; break; @@ -764,6 +787,10 @@ struct_member_decoration_cb(struct vtn_builder *b, spirv_decoration_to_string(dec->decoration)); break; + case SpvDecorationHlslSemanticGOOGLE: + /* HLSL semantic decorations can safely be ignored by the driver. */ + break; + default: vtn_fail("Unhandled decoration"); } @@ -846,16 +873,24 @@ type_decoration_cb(struct vtn_builder *b, case SpvDecorationNonWritable: case SpvDecorationNonReadable: case SpvDecorationUniform: - case SpvDecorationStream: case SpvDecorationLocation: case SpvDecorationComponent: case SpvDecorationOffset: case SpvDecorationXfbBuffer: case SpvDecorationXfbStride: + case SpvDecorationHlslSemanticGOOGLE: vtn_warn("Decoration only allowed for struct members: %s", spirv_decoration_to_string(dec->decoration)); break; + case SpvDecorationStream: + /* We don't need to do anything here, as stream is filled up when + * aplying the decoration to a variable, just check that if it is not a + * struct member, it should be a struct. + */ + vtn_assert(type->base_type == vtn_base_type_struct); + break; + case SpvDecorationRelaxedPrecision: case SpvDecorationSpecId: case SpvDecorationInvariant: @@ -1189,17 +1224,18 @@ vtn_handle_type(struct vtn_builder *b, SpvOp opcode, val->type->type = glsl_uint_type(); } - if (storage_class == SpvStorageClassWorkgroup && - b->options->lower_workgroup_access_to_offsets) { - uint32_t size, align; - val->type->deref = vtn_type_layout_std430(b, val->type->deref, - &size, &align); - val->type->length = size; - val->type->align = align; + if (storage_class == SpvStorageClassWorkgroup) { /* These can actually be stored to nir_variables and used as SSA * values so they need a real glsl_type. */ val->type->type = glsl_uint_type(); + if (b->options->lower_workgroup_access_to_offsets) { + uint32_t size, align; + val->type->deref = vtn_type_layout_std430(b, val->type->deref, + &size, &align); + val->type->length = size; + val->type->align = align; + } } break; } @@ -1432,7 +1468,7 @@ vtn_handle_constant(struct vtn_builder *b, SpvOp opcode, opcode == SpvOpSpecConstantFalse) int_val = get_specialization(b, val, int_val); - val->constant->values[0].u32[0] = int_val ? NIR_TRUE : NIR_FALSE; + val->constant->values[0].b[0] = int_val != 0; break; } @@ -1493,8 +1529,19 @@ vtn_handle_constant(struct vtn_builder *b, SpvOp opcode, spirv_op_to_string(opcode), elem_count, val->type->length); nir_constant **elems = ralloc_array(b, nir_constant *, elem_count); - for (unsigned i = 0; i < elem_count; i++) - elems[i] = vtn_value(b, w[i + 3], vtn_value_type_constant)->constant; + for (unsigned i = 0; i < elem_count; i++) { + struct vtn_value *val = vtn_untyped_value(b, w[i + 3]); + + if (val->value_type == vtn_value_type_constant) { + elems[i] = val->constant; + } else { + vtn_fail_if(val->value_type != vtn_value_type_undef, + "only constants or undefs allowed for " + "SpvOpConstantComposite"); + /* to make it easier, just insert a NULL constant for now */ + elems[i] = vtn_null_constant(b, val->type->type); + } + } switch (val->type->base_type) { case vtn_base_type_vector: { @@ -1514,6 +1561,9 @@ vtn_handle_constant(struct vtn_builder *b, SpvOp opcode, case 8: val->constant->values[0].u8[i] = elems[i]->values[0].u8[0]; break; + case 1: + val->constant->values[0].b[i] = elems[i]->values[0].b[0]; + break; default: vtn_fail("Invalid SpvOpConstantComposite bit size"); } @@ -1687,6 +1737,9 @@ vtn_handle_constant(struct vtn_builder *b, SpvOp opcode, case 8: val->constant->values[0].u8[i] = (*c)->values[col].u8[elem + i]; break; + case 1: + val->constant->values[0].b[i] = (*c)->values[col].b[elem + i]; + break; default: vtn_fail("Invalid SpvOpCompositeExtract bit size"); } @@ -1714,6 +1767,9 @@ vtn_handle_constant(struct vtn_builder *b, SpvOp opcode, case 8: (*c)->values[col].u8[elem + i] = insert->constant->values[0].u8[i]; break; + case 1: + (*c)->values[col].b[elem + i] = insert->constant->values[0].b[i]; + break; default: vtn_fail("Invalid SpvOpCompositeInsert bit size"); } @@ -1752,11 +1808,37 @@ vtn_handle_constant(struct vtn_builder *b, SpvOp opcode, nir_const_value src[4]; for (unsigned i = 0; i < count - 4; i++) { - nir_constant *c = - vtn_value(b, w[4 + i], vtn_value_type_constant)->constant; + struct vtn_value *src_val = + vtn_value(b, w[4 + i], vtn_value_type_constant); + + /* If this is an unsized source, pull the bit size from the + * source; otherwise, we'll use the bit size from the destination. + */ + if (!nir_alu_type_get_type_size(nir_op_infos[op].input_types[i])) + bit_size = glsl_get_bit_size(src_val->type->type); unsigned j = swap ? 1 - i : i; - src[j] = c->values[0]; + src[j] = src_val->constant->values[0]; + } + + /* fix up fixed size sources */ + switch (op) { + case nir_op_ishl: + case nir_op_ishr: + case nir_op_ushr: { + if (bit_size == 32) + break; + for (unsigned i = 0; i < num_components; ++i) { + switch (bit_size) { + case 64: src[1].u32[i] = src[1].u64[i]; break; + case 16: src[1].u32[i] = src[1].u16[i]; break; + case 8: src[1].u32[i] = src[1].u8[i]; break; + } + } + break; + } + default: + break; } val->constant->values[0] = @@ -1783,69 +1865,6 @@ vtn_handle_constant(struct vtn_builder *b, SpvOp opcode, vtn_foreach_decoration(b, val, handle_workgroup_size_decoration_cb, NULL); } -static void -vtn_handle_function_call(struct vtn_builder *b, SpvOp opcode, - const uint32_t *w, unsigned count) -{ - struct vtn_type *res_type = vtn_value(b, w[1], vtn_value_type_type)->type; - struct vtn_function *vtn_callee = - vtn_value(b, w[3], vtn_value_type_function)->func; - struct nir_function *callee = vtn_callee->impl->function; - - vtn_callee->referenced = true; - - nir_call_instr *call = nir_call_instr_create(b->nb.shader, callee); - - unsigned param_idx = 0; - - nir_deref_instr *ret_deref = NULL; - struct vtn_type *ret_type = vtn_callee->type->return_type; - if (ret_type->base_type != vtn_base_type_void) { - nir_variable *ret_tmp = - nir_local_variable_create(b->nb.impl, ret_type->type, "return_tmp"); - ret_deref = nir_build_deref_var(&b->nb, ret_tmp); - call->params[param_idx++] = nir_src_for_ssa(&ret_deref->dest.ssa); - } - - for (unsigned i = 0; i < vtn_callee->type->length; i++) { - struct vtn_type *arg_type = vtn_callee->type->params[i]; - unsigned arg_id = w[4 + i]; - - if (arg_type->base_type == vtn_base_type_sampled_image) { - struct vtn_sampled_image *sampled_image = - vtn_value(b, arg_id, vtn_value_type_sampled_image)->sampled_image; - - call->params[param_idx++] = - nir_src_for_ssa(&sampled_image->image->deref->dest.ssa); - call->params[param_idx++] = - nir_src_for_ssa(&sampled_image->sampler->deref->dest.ssa); - } else if (arg_type->base_type == vtn_base_type_pointer || - arg_type->base_type == vtn_base_type_image || - arg_type->base_type == vtn_base_type_sampler) { - struct vtn_pointer *pointer = - vtn_value(b, arg_id, vtn_value_type_pointer)->pointer; - call->params[param_idx++] = - nir_src_for_ssa(vtn_pointer_to_ssa(b, pointer)); - } else { - /* This is a regular SSA value and we need a temporary */ - nir_variable *tmp = - nir_local_variable_create(b->nb.impl, arg_type->type, "arg_tmp"); - nir_deref_instr *tmp_deref = nir_build_deref_var(&b->nb, tmp); - vtn_local_store(b, vtn_ssa_value(b, arg_id), tmp_deref); - call->params[param_idx++] = nir_src_for_ssa(&tmp_deref->dest.ssa); - } - } - assert(param_idx == call->num_params); - - nir_builder_instr_insert(&b->nb, &call->instr); - - if (ret_type->base_type == vtn_base_type_void) { - vtn_push_value(b, w[2], vtn_value_type_undef); - } else { - vtn_push_ssa(b, w[2], res_type, vtn_local_load(b, ret_deref)); - } -} - struct vtn_ssa_value * vtn_create_ssa_value(struct vtn_builder *b, const struct glsl_type *type) { @@ -2164,6 +2183,13 @@ vtn_handle_texture(struct vtn_builder *b, SpvOp opcode, texop = nir_texop_txf_ms; (*p++) = vtn_tex_src(b, w[idx++], nir_tex_src_ms_index); } + + if (operands & SpvImageOperandsMinLodMask) { + vtn_assert(texop == nir_texop_tex || + texop == nir_texop_txb || + texop == nir_texop_txd); + (*p++) = vtn_tex_src(b, w[idx++], nir_tex_src_min_lod); + } } /* We should have now consumed exactly all of the arguments */ vtn_assert(idx == count); @@ -2309,6 +2335,18 @@ get_image_coord(struct vtn_builder *b, uint32_t value) return nir_swizzle(&b->nb, coord->def, swizzle, 4, false); } +static nir_ssa_def * +expand_to_vec4(nir_builder *b, nir_ssa_def *value) +{ + if (value->num_components == 4) + return value; + + unsigned swiz[4]; + for (unsigned i = 0; i < 4; i++) + swiz[i] = i < value->num_components ? i : 0; + return nir_swizzle(b, value, swiz, 4, false); +} + static void vtn_handle_image(struct vtn_builder *b, SpvOp opcode, const uint32_t *w, unsigned count) @@ -2422,11 +2460,7 @@ vtn_handle_image(struct vtn_builder *b, SpvOp opcode, /* The image coordinate is always 4 components but we may not have that * many. Swizzle to compensate. */ - unsigned swiz[4]; - for (unsigned i = 0; i < 4; i++) - swiz[i] = i < image.coord->num_components ? i : 0; - intrin->src[1] = nir_src_for_ssa(nir_swizzle(&b->nb, image.coord, - swiz, 4, false)); + intrin->src[1] = nir_src_for_ssa(expand_to_vec4(&b->nb, image.coord)); intrin->src[2] = nir_src_for_ssa(image.sample); } @@ -2436,11 +2470,15 @@ vtn_handle_image(struct vtn_builder *b, SpvOp opcode, case SpvOpImageRead: break; case SpvOpAtomicStore: - intrin->src[3] = nir_src_for_ssa(vtn_ssa_value(b, w[4])->def); - break; - case SpvOpImageWrite: - intrin->src[3] = nir_src_for_ssa(vtn_ssa_value(b, w[3])->def); + case SpvOpImageWrite: { + const uint32_t value_id = opcode == SpvOpAtomicStore ? w[4] : w[3]; + nir_ssa_def *value = vtn_ssa_value(b, value_id)->def; + /* nir_intrinsic_image_deref_store always takes a vec4 value */ + assert(op == nir_intrinsic_image_deref_store); + intrin->num_components = 4; + intrin->src[3] = nir_src_for_ssa(expand_to_vec4(&b->nb, value)); break; + } case SpvOpAtomicCompareExchange: case SpvOpAtomicIIncrement: @@ -2462,23 +2500,26 @@ vtn_handle_image(struct vtn_builder *b, SpvOp opcode, vtn_fail("Invalid image opcode"); } - if (opcode != SpvOpImageWrite) { + if (opcode != SpvOpImageWrite && opcode != SpvOpAtomicStore) { struct vtn_value *val = vtn_push_value(b, w[2], vtn_value_type_ssa); struct vtn_type *type = vtn_value(b, w[1], vtn_value_type_type)->type; - unsigned dest_components = nir_intrinsic_dest_components(intrin); - if (intrin->intrinsic == nir_intrinsic_image_deref_size) { - dest_components = intrin->num_components = - glsl_get_vector_elements(type->type); - } + unsigned dest_components = glsl_get_vector_elements(type->type); + intrin->num_components = nir_intrinsic_infos[op].dest_components; + if (intrin->num_components == 0) + intrin->num_components = dest_components; nir_ssa_dest_init(&intrin->instr, &intrin->dest, - dest_components, 32, NULL); + intrin->num_components, 32, NULL); nir_builder_instr_insert(&b->nb, &intrin->instr); + nir_ssa_def *result = &intrin->dest.ssa; + if (intrin->num_components != dest_components) + result = nir_channels(&b->nb, result, (1 << dest_components) - 1); + val->ssa = vtn_create_ssa_value(b, type->type); - val->ssa->def = &intrin->dest.ssa; + val->ssa->def = result; } else { nir_builder_instr_insert(&b->nb, &intrin->instr); } @@ -2519,7 +2560,7 @@ get_uniform_nir_atomic_op(struct vtn_builder *b, SpvOp opcode) OP(AtomicExchange, exchange) OP(AtomicCompareExchange, comp_swap) OP(AtomicIIncrement, inc_deref) - OP(AtomicIDecrement, dec_deref) + OP(AtomicIDecrement, post_dec_deref) OP(AtomicIAdd, add_deref) OP(AtomicISub, add_deref) OP(AtomicUMin, min_deref) @@ -2738,6 +2779,7 @@ vtn_handle_atomics(struct vtn_builder *b, SpvOp opcode, switch (opcode) { case SpvOpAtomicLoad: atomic->num_components = glsl_get_vector_elements(ptr->type->type); + nir_intrinsic_set_align(atomic, 4, 0); if (ptr->mode == vtn_variable_mode_ssbo) atomic->src[src++] = nir_src_for_ssa(index); atomic->src[src++] = nir_src_for_ssa(offset); @@ -2746,6 +2788,7 @@ vtn_handle_atomics(struct vtn_builder *b, SpvOp opcode, case SpvOpAtomicStore: atomic->num_components = glsl_get_vector_elements(ptr->type->type); nir_intrinsic_set_write_mask(atomic, (1 << atomic->num_components) - 1); + nir_intrinsic_set_align(atomic, 4, 0); atomic->src[src++] = nir_src_for_ssa(vtn_ssa_value(b, w[4])->def); if (ptr->mode == vtn_variable_mode_ssbo) atomic->src[src++] = nir_src_for_ssa(index); @@ -2798,7 +2841,7 @@ create_vec(struct vtn_builder *b, unsigned num_components, unsigned bit_size) { nir_op op; switch (num_components) { - case 1: op = nir_op_fmov; break; + case 1: op = nir_op_imov; break; case 2: op = nir_op_vec2; break; case 3: op = nir_op_vec3; break; case 4: op = nir_op_vec4; break; @@ -2846,8 +2889,7 @@ vtn_ssa_transpose(struct vtn_builder *b, struct vtn_ssa_value *src) nir_ssa_def * vtn_vector_extract(struct vtn_builder *b, nir_ssa_def *src, unsigned index) { - unsigned swiz[4] = { index }; - return nir_swizzle(&b->nb, src, swiz, 1, true); + return nir_channel(&b->nb, src, index); } nir_ssa_def * @@ -3063,7 +3105,7 @@ vtn_handle_composite(struct vtn_builder *b, SpvOp opcode, unsigned elems = count - 3; assume(elems >= 1); if (glsl_type_is_vector_or_scalar(type)) { - nir_ssa_def *srcs[4]; + nir_ssa_def *srcs[NIR_MAX_VEC_COMPONENTS]; for (unsigned i = 0; i < elems; i++) srcs[i] = vtn_ssa_value(b, w[3 + i])->def; val->ssa->def = @@ -3187,9 +3229,12 @@ vtn_handle_barrier(struct vtn_builder *b, SpvOp opcode, switch (opcode) { case SpvOpEmitStreamVertex: - case SpvOpEndStreamPrimitive: - nir_intrinsic_set_stream_id(intrin, w[1]); + case SpvOpEndStreamPrimitive: { + unsigned stream = vtn_constant_value(b, w[1])->values[0].u32[0]; + nir_intrinsic_set_stream_id(intrin, stream); break; + } + default: break; } @@ -3383,21 +3428,21 @@ vtn_handle_preamble_instruction(struct vtn_builder *b, SpvOp opcode, case SpvCapabilityStorageImageExtendedFormats: break; - case SpvCapabilityGeometryStreams: case SpvCapabilityLinkage: case SpvCapabilityVector16: case SpvCapabilityFloat16Buffer: case SpvCapabilityFloat16: - case SpvCapabilityInt64Atomics: case SpvCapabilityStorageImageMultisample: case SpvCapabilityInt8: case SpvCapabilitySparseResidency: - case SpvCapabilityMinLod: - case SpvCapabilityTransformFeedback: vtn_warn("Unsupported SPIR-V capability: %s", spirv_capability_to_string(cap)); break; + case SpvCapabilityMinLod: + spv_check_supported(min_lod, cap); + break; + case SpvCapabilityAtomicStorage: spv_check_supported(atomic_storage, cap); break; @@ -3412,6 +3457,18 @@ vtn_handle_preamble_instruction(struct vtn_builder *b, SpvOp opcode, spv_check_supported(int16, cap); break; + case SpvCapabilityTransformFeedback: + spv_check_supported(transform_feedback, cap); + break; + + case SpvCapabilityGeometryStreams: + spv_check_supported(geometry_streams, cap); + break; + + case SpvCapabilityInt64Atomics: + spv_check_supported(int64_atomics, cap); + break; + case SpvCapabilityAddresses: case SpvCapabilityKernel: case SpvCapabilityImageBasic: @@ -3498,6 +3555,12 @@ vtn_handle_preamble_instruction(struct vtn_builder *b, SpvOp opcode, spv_check_supported(shader_viewport_index_layer, cap); break; + case SpvCapabilityStorageBuffer8BitAccess: + case SpvCapabilityUniformAndStorageBuffer8BitAccess: + case SpvCapabilityStoragePushConstant8: + spv_check_supported(storage_8bit, cap); + break; + case SpvCapabilityInputAttachmentArrayDynamicIndexingEXT: case SpvCapabilityUniformTexelBufferArrayDynamicIndexingEXT: case SpvCapabilityStorageTexelBufferArrayDynamicIndexingEXT: @@ -3512,6 +3575,10 @@ vtn_handle_preamble_instruction(struct vtn_builder *b, SpvOp opcode, spv_check_supported(stencil_export, cap); break; + case SpvCapabilitySampleMaskPostDepthCoverage: + spv_check_supported(post_depth_coverage, cap); + break; + default: vtn_fail("Unhandled capability"); } @@ -3551,6 +3618,8 @@ vtn_handle_preamble_instruction(struct vtn_builder *b, SpvOp opcode, case SpvOpMemberDecorate: case SpvOpGroupDecorate: case SpvOpGroupMemberDecorate: + case SpvOpDecorateStringGOOGLE: + case SpvOpMemberDecorateStringGOOGLE: vtn_handle_decoration(b, opcode, w, count); break; @@ -3579,6 +3648,11 @@ vtn_handle_execution_mode(struct vtn_builder *b, struct vtn_value *entry_point, b->shader->info.fs.early_fragment_tests = true; break; + case SpvExecutionModePostDepthCoverage: + vtn_assert(b->shader->info.stage == MESA_SHADER_FRAGMENT); + b->shader->info.fs.post_depth_coverage = true; + break; + case SpvExecutionModeInvocations: vtn_assert(b->shader->info.stage == MESA_SHADER_GEOMETRY); b->shader->info.gs.invocations = MAX2(1, mode->literals[0]); @@ -3635,6 +3709,8 @@ vtn_handle_execution_mode(struct vtn_builder *b, struct vtn_value *entry_point, vtn_assert(b->shader->info.stage == MESA_SHADER_GEOMETRY); b->shader->info.gs.vertices_in = vertices_in_from_spv_execution_mode(b, mode->exec_mode); + b->shader->info.gs.input_primitive = + gl_primitive_from_spv_execution_mode(b, mode->exec_mode); } break; @@ -3682,7 +3758,7 @@ vtn_handle_execution_mode(struct vtn_builder *b, struct vtn_value *entry_point, break; case SpvExecutionModeXfb: - vtn_fail("Unhandled execution mode"); + b->shader->info.has_transform_feedback_varyings = true; break; case SpvExecutionModeVecTypeHint: @@ -3722,6 +3798,8 @@ vtn_handle_variable_or_type_instruction(struct vtn_builder *b, SpvOp opcode, case SpvOpMemberDecorate: case SpvOpGroupDecorate: case SpvOpGroupMemberDecorate: + case SpvOpDecorateStringGOOGLE: + case SpvOpMemberDecorateStringGOOGLE: vtn_fail("Invalid opcode types and variables section"); break;