longjmp(b->fail_jump, 1);
}
-struct spec_constant_value {
- bool is_double;
- union {
- uint32_t data32;
- uint64_t data64;
- };
-};
-
static struct vtn_ssa_value *
vtn_undef_ssa_value(struct vtn_builder *b, const struct glsl_type *type)
{
if (dec->decoration != SpvDecorationSpecId)
return;
- struct spec_constant_value *const_value = data;
-
+ nir_const_value *value = data;
for (unsigned i = 0; i < b->num_specializations; i++) {
if (b->specializations[i].id == dec->operands[0]) {
- if (const_value->is_double)
- const_value->data64 = b->specializations[i].data64;
- else
- const_value->data32 = b->specializations[i].data32;
+ *value = b->specializations[i].value;
return;
}
}
}
-static uint32_t
-get_specialization(struct vtn_builder *b, struct vtn_value *val,
- uint32_t const_value)
-{
- struct spec_constant_value data;
- data.is_double = false;
- data.data32 = const_value;
- vtn_foreach_decoration(b, val, spec_constant_decoration_cb, &data);
- return data.data32;
-}
-
-static uint64_t
-get_specialization64(struct vtn_builder *b, struct vtn_value *val,
- uint64_t const_value)
-{
- struct spec_constant_value data;
- data.is_double = true;
- data.data64 = const_value;
- vtn_foreach_decoration(b, val, spec_constant_decoration_cb, &data);
- return data.data64;
-}
-
static void
handle_workgroup_size_decoration_cb(struct vtn_builder *b,
struct vtn_value *val,
"Result type of %s must be OpTypeBool",
spirv_op_to_string(opcode));
- uint32_t int_val = (opcode == SpvOpConstantTrue ||
- opcode == SpvOpSpecConstantTrue);
+ bool bval = (opcode == SpvOpConstantTrue ||
+ opcode == SpvOpSpecConstantTrue);
+
+ nir_const_value u32val = nir_const_value_for_uint(bval, 32);
if (opcode == SpvOpSpecConstantTrue ||
opcode == SpvOpSpecConstantFalse)
- int_val = get_specialization(b, val, int_val);
+ vtn_foreach_decoration(b, val, spec_constant_decoration_cb, &u32val);
- val->constant->values[0].b = int_val != 0;
+ val->constant->values[0].b = u32val.u32 != 0;
break;
}
- case SpvOpConstant: {
+ case SpvOpConstant:
+ case SpvOpSpecConstant: {
vtn_fail_if(val->type->base_type != vtn_base_type_scalar,
"Result type of %s must be a scalar",
spirv_op_to_string(opcode));
default:
vtn_fail("Unsupported SpvOpConstant bit size: %u", bit_size);
}
- break;
- }
- case SpvOpSpecConstant: {
- vtn_fail_if(val->type->base_type != vtn_base_type_scalar,
- "Result type of %s must be a scalar",
- spirv_op_to_string(opcode));
- int bit_size = glsl_get_bit_size(val->type->type);
- switch (bit_size) {
- case 64:
- val->constant->values[0].u64 =
- get_specialization64(b, val, vtn_u64_literal(&w[3]));
- break;
- case 32:
- val->constant->values[0].u32 = get_specialization(b, val, w[3]);
- break;
- case 16:
- val->constant->values[0].u16 = get_specialization(b, val, w[3]);
- break;
- case 8:
- val->constant->values[0].u8 = get_specialization(b, val, w[3]);
- break;
- default:
- vtn_fail("Unsupported SpvOpSpecConstant bit size");
- }
+ if (opcode == SpvOpSpecConstant)
+ vtn_foreach_decoration(b, val, spec_constant_decoration_cb,
+ &val->constant->values[0]);
break;
}
}
case SpvOpSpecConstantOp: {
- SpvOp opcode = get_specialization(b, val, w[3]);
+ nir_const_value u32op = nir_const_value_for_uint(w[3], 32);
+ vtn_foreach_decoration(b, val, spec_constant_decoration_cb, &u32op);
+ SpvOp opcode = u32op.u32;
switch (opcode) {
case SpvOpVectorShuffle: {
struct vtn_value *v0 = &b->values[w[4]];
return dest;
}
-nir_ssa_def *
-vtn_vector_insert(struct vtn_builder *b, nir_ssa_def *src, nir_ssa_def *insert,
- unsigned index)
-{
- nir_alu_instr *vec = create_vec(b, src->num_components,
- src->bit_size);
-
- for (unsigned i = 0; i < src->num_components; i++) {
- if (i == index) {
- vec->src[i].src = nir_src_for_ssa(insert);
- } else {
- vec->src[i].src = nir_src_for_ssa(src);
- vec->src[i].swizzle[0] = i;
- }
- }
-
- nir_builder_instr_insert(&b->nb, &vec->instr);
-
- return &vec->dest.dest.ssa;
-}
-
-nir_ssa_def *
-vtn_vector_insert_dynamic(struct vtn_builder *b, nir_ssa_def *src,
- nir_ssa_def *insert, nir_ssa_def *index)
-{
- nir_const_value per_comp_idx_const[NIR_MAX_VEC_COMPONENTS];
- for (unsigned i = 0; i < NIR_MAX_VEC_COMPONENTS; i++)
- per_comp_idx_const[i] = nir_const_value_for_int(i, index->bit_size);
- nir_ssa_def *per_comp_idx =
- nir_build_imm(&b->nb, src->num_components,
- index->bit_size, per_comp_idx_const);
-
- /* nir_builder will automatically splat out scalars to vectors so an insert
- * is as simple as "if I'm the channel, replace me with the scalar."
- */
- return nir_bcsel(&b->nb, nir_ieq(&b->nb, index, per_comp_idx), insert, src);
-}
-
static nir_ssa_def *
vtn_vector_shuffle(struct vtn_builder *b, unsigned num_components,
nir_ssa_def *src0, nir_ssa_def *src1,
* the index to insert the scalar into the vector.
*/
- cur->def = vtn_vector_insert(b, cur->def, insert->def, indices[i]);
+ cur->def = nir_vector_insert_imm(&b->nb, cur->def, insert->def, indices[i]);
} else {
vtn_fail_if(indices[i] >= glsl_get_length(cur->type),
"All indices in an OpCompositeInsert must be in-bounds");
break;
case SpvOpVectorInsertDynamic:
- ssa->def = vtn_vector_insert_dynamic(b, vtn_ssa_value(b, w[3])->def,
- vtn_ssa_value(b, w[4])->def,
- vtn_ssa_value(b, w[5])->def);
+ ssa->def = nir_vector_insert(&b->nb, vtn_ssa_value(b, w[3])->def,
+ vtn_ssa_value(b, w[4])->def,
+ vtn_ssa_value(b, w[5])->def);
break;
case SpvOpVectorShuffle:
break;
case SpvOpCopyLogical:
- case SpvOpCopyObject:
ssa = vtn_composite_copy(b, vtn_ssa_value(b, w[3]));
break;
+ case SpvOpCopyObject:
+ vtn_copy_value(b, w[3], w[2]);
+ return;
default:
vtn_fail_with_opcode("unknown composite operation", opcode);
/* GLSLang, prior to commit 8297936dd6eb3, emitted OpControlBarrier with
* memory semantics of None for GLSL barrier().
+ * And before that, prior to c3f1cdfa, emitted the OpControlBarrier with
+ * Device instead of Workgroup for execution scope.
*/
if (b->wa_glslang_cs_barrier &&
b->nb.shader->info.stage == MESA_SHADER_COMPUTE &&
- execution_scope == SpvScopeWorkgroup &&
+ (execution_scope == SpvScopeWorkgroup ||
+ execution_scope == SpvScopeDevice) &&
memory_semantics == SpvMemorySemanticsMaskNone) {
+ execution_scope = SpvScopeWorkgroup;
memory_scope = SpvScopeWorkgroup;
memory_semantics = SpvMemorySemanticsAcquireReleaseMask |
SpvMemorySemanticsWorkgroupMemoryMask;
}
case SpvOpReadClockKHR: {
- assert(vtn_constant_uint(b, w[3]) == SpvScopeSubgroup);
+ SpvScope scope = vtn_constant_uint(b, w[3]);
+ nir_scope nir_scope;
+
+ switch (scope) {
+ case SpvScopeDevice:
+ nir_scope = NIR_SCOPE_DEVICE;
+ break;
+ case SpvScopeSubgroup:
+ nir_scope = NIR_SCOPE_SUBGROUP;
+ break;
+ default:
+ vtn_fail("invalid read clock scope");
+ }
/* Operation supports two result types: uvec2 and uint64_t. The NIR
* intrinsic gives uvec2, so pack the result for the other case.
nir_intrinsic_instr *intrin =
nir_intrinsic_instr_create(b->nb.shader, nir_intrinsic_shader_clock);
nir_ssa_dest_init(&intrin->instr, &intrin->dest, 2, 32, NULL);
+ nir_intrinsic_set_memory_scope(intrin, nir_scope);
nir_builder_instr_insert(&b->nb, &intrin->instr);
struct vtn_type *type = vtn_value(b, w[1], vtn_value_type_type)->type;
projects / mesa.git / blobdiff