X-Git-Url: https://git.libre-soc.org/?a=blobdiff_plain;f=src%2Fcompiler%2Fspirv%2Fspirv_to_nir.c;h=7dc6bc914d6f9374440702959f9d0f54e5758fb1;hb=44227453ec03f5462f1cff5760909a9dba95c61a;hp=fb62f3035a97b3911ead4dd4f6470048f7b8b6a7;hpb=7dfa440922f562ae41d0671219ace8ef2217804a;p=mesa.git diff --git a/src/compiler/spirv/spirv_to_nir.c b/src/compiler/spirv/spirv_to_nir.c index fb62f3035a9..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,57 +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); - for (unsigned i = 0; i < call->num_params; i++) { - unsigned arg_id = w[4 + i]; - struct vtn_value *arg = vtn_untyped_value(b, arg_id); - if (arg->value_type == vtn_value_type_pointer && - arg->pointer->ptr_type->type == NULL) { - nir_deref_var *d = vtn_pointer_to_deref_var(b, arg->pointer); - call->params[i] = nir_deref_var_clone(d, call); - } else { - struct vtn_ssa_value *arg_ssa = vtn_ssa_value(b, arg_id); - - /* Make a temporary to store the argument in */ - nir_variable *tmp = - nir_local_variable_create(b->nb.impl, arg_ssa->type, "arg_tmp"); - call->params[i] = nir_deref_var_create(call, tmp); - - vtn_local_store(b, arg_ssa, - nir_build_deref_for_chain(&b->nb, call->params[i])); - } - } - - nir_variable *out_tmp = NULL; - vtn_assert(res_type->type == callee->return_type); - if (!glsl_type_is_void(callee->return_type)) { - out_tmp = nir_local_variable_create(b->nb.impl, callee->return_type, - "out_tmp"); - call->return_deref = nir_deref_var_create(call, out_tmp); - } - - nir_builder_instr_insert(&b->nb, &call->instr); - - if (glsl_type_is_void(callee->return_type)) { - vtn_push_value(b, w[2], vtn_value_type_undef); - } else { - nir_deref_instr *return_deref = - nir_build_deref_for_chain(&b->nb, call->return_deref); - vtn_push_ssa(b, w[2], res_type, vtn_local_load(b, return_deref)); - } -} - struct vtn_ssa_value * vtn_create_ssa_value(struct vtn_builder *b, const struct glsl_type *type) { @@ -2152,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); @@ -2297,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) @@ -2410,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); } @@ -2424,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: @@ -2450,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); } @@ -2498,6 +2551,35 @@ get_ssbo_nir_atomic_op(struct vtn_builder *b, SpvOp opcode) } } +static nir_intrinsic_op +get_uniform_nir_atomic_op(struct vtn_builder *b, SpvOp opcode) +{ + switch (opcode) { +#define OP(S, N) case SpvOp##S: return nir_intrinsic_atomic_counter_ ##N; + OP(AtomicLoad, read_deref) + OP(AtomicExchange, exchange) + OP(AtomicCompareExchange, comp_swap) + OP(AtomicIIncrement, inc_deref) + OP(AtomicIDecrement, post_dec_deref) + OP(AtomicIAdd, add_deref) + OP(AtomicISub, add_deref) + OP(AtomicUMin, min_deref) + OP(AtomicUMax, max_deref) + OP(AtomicAnd, and_deref) + OP(AtomicOr, or_deref) + OP(AtomicXor, xor_deref) +#undef OP + default: + /* We left the following out: AtomicStore, AtomicSMin and + * AtomicSmax. Right now there are not nir intrinsics for them. At this + * moment Atomic Counter support is needed for ARB_spirv support, so is + * only need to support GLSL Atomic Counters that are uints and don't + * allow direct storage. + */ + unreachable("Invalid uniform atomic"); + } +} + static nir_intrinsic_op get_shared_nir_atomic_op(struct vtn_builder *b, SpvOp opcode) { @@ -2550,9 +2632,12 @@ get_deref_nir_atomic_op(struct vtn_builder *b, SpvOp opcode) } } +/* + * Handles shared atomics, ssbo atomics and atomic counters. + */ static void -vtn_handle_ssbo_or_shared_atomic(struct vtn_builder *b, SpvOp opcode, - const uint32_t *w, unsigned count) +vtn_handle_atomics(struct vtn_builder *b, SpvOp opcode, + const uint32_t *w, unsigned count) { struct vtn_pointer *ptr; nir_intrinsic_instr *atomic; @@ -2589,8 +2674,54 @@ vtn_handle_ssbo_or_shared_atomic(struct vtn_builder *b, SpvOp opcode, SpvMemorySemanticsMask semantics = w[5]; */ - if (ptr->mode == vtn_variable_mode_workgroup && - !b->options->lower_workgroup_access_to_offsets) { + /* uniform as "atomic counter uniform" */ + if (ptr->mode == vtn_variable_mode_uniform) { + nir_deref_instr *deref = vtn_pointer_to_deref(b, ptr); + const struct glsl_type *deref_type = deref->type; + nir_intrinsic_op op = get_uniform_nir_atomic_op(b, opcode); + atomic = nir_intrinsic_instr_create(b->nb.shader, op); + atomic->src[0] = nir_src_for_ssa(&deref->dest.ssa); + + /* SSBO needs to initialize index/offset. In this case we don't need to, + * as that info is already stored on the ptr->var->var nir_variable (see + * vtn_create_variable) + */ + + switch (opcode) { + case SpvOpAtomicLoad: + atomic->num_components = glsl_get_vector_elements(deref_type); + break; + + case SpvOpAtomicStore: + atomic->num_components = glsl_get_vector_elements(deref_type); + nir_intrinsic_set_write_mask(atomic, (1 << atomic->num_components) - 1); + break; + + case SpvOpAtomicExchange: + case SpvOpAtomicCompareExchange: + case SpvOpAtomicCompareExchangeWeak: + case SpvOpAtomicIIncrement: + case SpvOpAtomicIDecrement: + case SpvOpAtomicIAdd: + case SpvOpAtomicISub: + case SpvOpAtomicSMin: + case SpvOpAtomicUMin: + case SpvOpAtomicSMax: + case SpvOpAtomicUMax: + case SpvOpAtomicAnd: + case SpvOpAtomicOr: + case SpvOpAtomicXor: + /* Nothing: we don't need to call fill_common_atomic_sources here, as + * atomic counter uniforms doesn't have sources + */ + break; + + default: + unreachable("Invalid SPIR-V atomic"); + + } + } else if (ptr->mode == vtn_variable_mode_workgroup && + !b->options->lower_workgroup_access_to_offsets) { nir_deref_instr *deref = vtn_pointer_to_deref(b, ptr); const struct glsl_type *deref_type = deref->type; nir_intrinsic_op op = get_deref_nir_atomic_op(b, opcode); @@ -2631,7 +2762,7 @@ vtn_handle_ssbo_or_shared_atomic(struct vtn_builder *b, SpvOp opcode, } } else { nir_ssa_def *offset, *index; - offset = vtn_pointer_to_offset(b, ptr, &index, NULL); + offset = vtn_pointer_to_offset(b, ptr, &index); nir_intrinsic_op op; if (ptr->mode == vtn_variable_mode_ssbo) { @@ -2648,6 +2779,7 @@ vtn_handle_ssbo_or_shared_atomic(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); @@ -2656,6 +2788,7 @@ vtn_handle_ssbo_or_shared_atomic(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); @@ -2708,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; @@ -2756,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 * @@ -2973,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 = @@ -3097,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; } @@ -3293,22 +3428,25 @@ 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 SpvCapabilityAtomicStorage: 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; + case SpvCapabilityFloat64: spv_check_supported(float64, cap); break; @@ -3319,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: @@ -3405,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: @@ -3419,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"); } @@ -3458,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; @@ -3486,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]); @@ -3542,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; @@ -3589,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: @@ -3629,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; @@ -3778,7 +3949,7 @@ vtn_handle_body_instruction(struct vtn_builder *b, SpvOp opcode, vtn_handle_image(b, opcode, w, count); } else { vtn_assert(pointer->value_type == vtn_value_type_pointer); - vtn_handle_ssbo_or_shared_atomic(b, opcode, w, count); + vtn_handle_atomics(b, opcode, w, count); } break; } @@ -3789,7 +3960,7 @@ vtn_handle_body_instruction(struct vtn_builder *b, SpvOp opcode, vtn_handle_image(b, opcode, w, count); } else { vtn_assert(pointer->value_type == vtn_value_type_pointer); - vtn_handle_ssbo_or_shared_atomic(b, opcode, w, count); + vtn_handle_atomics(b, opcode, w, count); } break; }