godotengine/godot-nir-static
1
0
Fork 0
mirror of https://github.com/godotengine/godot-nir-static.git synced 2025-12-31 13:48:20 +03:00
Code Issues Packages Projects Releases Wiki Activity
Files
809e2710956e8d83cdc8126d1cf45731786cc35c
godot-nir-static/godot-patches/01_godot_nir_goodies.patch

1126 lines
44 KiB
Diff
Raw Blame History

diff --git a/godot-mesa/src/compiler/nir/nir_dominance_lca.c b/godot-mesa/src/compiler/nir/nir_dominance_lca.c
index 84a7f1d..40c0811 100644
--- a/godot-mesa/src/compiler/nir/nir_dominance_lca.c
+++ b/godot-mesa/src/compiler/nir/nir_dominance_lca.c
@@ -167,7 +167,7 @@ nir_dominance_lca(nir_block *b1, nir_block *b2)
uint32_t i1 = dom_lca_representative(b1);
uint32_t i2 = dom_lca_representative(b2);
if (i1 > i2)
- SWAP(i1, i2);
+ MESA_SWAP(i1, i2);
uint32_t index = range_minimum_query(&impl->dom_lca_info.table, i1, i2 + 1);
nir_block *result = impl->dom_lca_info.block_from_idx[index];
diff --git a/godot-mesa/src/compiler/nir/nir_intrinsics.py b/godot-mesa/src/compiler/nir/nir_intrinsics.py
index 41371db..dd3454f 100644
--- a/godot-mesa/src/compiler/nir/nir_intrinsics.py
+++ b/godot-mesa/src/compiler/nir/nir_intrinsics.py
@@ -1262,6 +1262,9 @@ load("push_constant", [1], [BASE, RANGE, ALIGN_MUL, ALIGN_OFFSET], [CAN_ELIMINAT
# src[] = { offset }.
load("constant", [1], [BASE, RANGE, ACCESS, ALIGN_MUL, ALIGN_OFFSET],
[CAN_ELIMINATE, CAN_REORDER])
+# src[] = { offset }.
+load("constant_non_opt", [1], [BASE, RANGE, ACCESS, ALIGN_MUL, ALIGN_OFFSET],
+ [CAN_ELIMINATE, CAN_REORDER])
# src[] = { address }.
load("global", [1], [ACCESS, ALIGN_MUL, ALIGN_OFFSET], [CAN_ELIMINATE])
# src[] = { base_address, offset, bound }.
diff --git a/godot-mesa/src/compiler/nir/nir_loop_analyze.c b/godot-mesa/src/compiler/nir/nir_loop_analyze.c
index aa7369c..34b8a99 100644
--- a/godot-mesa/src/compiler/nir/nir_loop_analyze.c
+++ b/godot-mesa/src/compiler/nir/nir_loop_analyze.c
@@ -1020,7 +1020,7 @@ try_find_trip_count_vars_in_logical_op(nir_scalar *cond,
if (!nir_scalar_is_alu(logical_op) || !nir_scalar_is_const(zero)) {
/* Maybe we had it the wrong way, flip things around */
- SWAP(zero, logical_op);
+ MESA_SWAP(zero, logical_op);
/* If we still didn't find what we need then return */
if (!nir_scalar_is_const(zero))
diff --git a/godot-mesa/src/compiler/nir/nir_lower_cooperative_matrix.c b/godot-mesa/src/compiler/nir/nir_lower_cooperative_matrix.c
index 830746b..74c8df6 100644
--- a/godot-mesa/src/compiler/nir/nir_lower_cooperative_matrix.c
+++ b/godot-mesa/src/compiler/nir/nir_lower_cooperative_matrix.c
@@ -530,7 +530,7 @@ split_cmat_load_store(nir_builder *b,
col_offset = (i / split->num_col_splits) * desc.rows;
if (layout == GLSL_MATRIX_LAYOUT_ROW_MAJOR)
- SWAP(row_offset, col_offset);
+ MESA_SWAP(row_offset, col_offset);
ptr_deref = nir_build_deref_cast(b, &addr_deref->def, addr_deref->modes, scalar_type, elem_size);
stride = nir_udiv_imm(b, nir_imul_imm(b, stride, deref_bytes_size), elem_size);
diff --git a/godot-mesa/src/compiler/nir/nir_lower_input_attachments.c b/godot-mesa/src/compiler/nir/nir_lower_input_attachments.c
index 323b188..09460ff 100644
--- a/godot-mesa/src/compiler/nir/nir_lower_input_attachments.c
+++ b/godot-mesa/src/compiler/nir/nir_lower_input_attachments.c
@@ -106,12 +106,53 @@ load_coord(nir_builder *b, nir_deref_instr *deref,
}
}
+static const struct glsl_type *
+get_texture_type_for_image(const struct glsl_type *type)
+{
+ if (glsl_type_is_array(type)) {
+ const struct glsl_type *elem_type =
+ get_texture_type_for_image(glsl_get_array_element(type));
+ return glsl_array_type(elem_type, glsl_get_length(type), 0 /*explicit size*/);
+ }
+
+ assert((glsl_type_is_image(type)));
+ return glsl_texture_type(glsl_get_sampler_dim(type),
+ glsl_sampler_type_is_array(type),
+ glsl_get_sampler_result_type(type));
+}
+
+static bool
+replace_image_type_with_texture(nir_deref_instr *deref)
+{
+ const struct glsl_type *type = deref->type;
+
+ /* If we've already chased up the deref chain this far from a different intrinsic, we're done */
+ if (!glsl_type_is_image(glsl_without_array(type)))
+ return false;
+
+ deref->type = get_texture_type_for_image(type);
+ deref->modes = nir_var_uniform;
+ if (deref->deref_type == nir_deref_type_var) {
+ type = deref->var->type;
+ if (glsl_type_is_image(glsl_without_array(type))) {
+ deref->var->type = get_texture_type_for_image(type);
+ deref->var->data.mode = nir_var_uniform;
+ memset(&deref->var->data.sampler, 0, sizeof(deref->var->data.sampler));
+ }
+ } else {
+ nir_deref_instr *parent = nir_deref_instr_parent(deref);
+ if (parent)
+ replace_image_type_with_texture(parent);
+ }
+
+ return true;
+}
+
static bool
try_lower_input_load(nir_builder *b, nir_intrinsic_instr *load,
const nir_input_attachment_options *options)
{
nir_deref_instr *deref = nir_src_as_deref(load->src[0]);
- assert(glsl_type_is_image(deref->type));
enum glsl_sampler_dim image_dim = glsl_get_sampler_dim(deref->type);
if (image_dim != GLSL_SAMPLER_DIM_SUBPASS &&
@@ -172,6 +213,8 @@ try_lower_input_load(nir_builder *b, nir_intrinsic_instr *load,
&tex->def);
}
+ replace_image_type_with_texture(deref);
+
return true;
}
@@ -204,6 +247,8 @@ try_lower_input_texop(nir_builder *b, nir_tex_instr *tex,
nir_src_rewrite(&tex->src[coord_src_idx].src, coord);
+ replace_image_type_with_texture(deref);
+
return true;
}
diff --git a/godot-mesa/src/compiler/nir/nir_opt_reassociate.c b/godot-mesa/src/compiler/nir/nir_opt_reassociate.c
index 09566a5..15c5ad1 100644
--- a/godot-mesa/src/compiler/nir/nir_opt_reassociate.c
+++ b/godot-mesa/src/compiler/nir/nir_opt_reassociate.c
@@ -118,7 +118,7 @@ get_pair_key(nir_op op, nir_scalar a, nir_scalar b)
if ((a.def->index > b.def->index) ||
((a.def->index == b.def->index) && (a.comp > b.comp))) {
- SWAP(a, b);
+ MESA_SWAP(a, b);
}
return (struct pair_key){
@@ -460,8 +460,8 @@ reassociate_chain(struct chain *c, void *pair_freq)
if (best_pair.i != best_pair.j) {
/* Pin the best pair at the front. The rest is sorted by rank. */
- SWAP(c->srcs[0], c->srcs[best_pair.i]);
- SWAP(c->srcs[1], c->srcs[best_pair.j]);
+ MESA_SWAP(c->srcs[0], c->srcs[best_pair.i]);
+ MESA_SWAP(c->srcs[1], c->srcs[best_pair.j]);
pinned = 2;
}
}
diff --git a/godot-mesa/src/compiler/spirv/spirv_to_nir.c b/godot-mesa/src/compiler/spirv/spirv_to_nir.c
index 7150419..61b0b7c 100644
--- a/godot-mesa/src/compiler/spirv/spirv_to_nir.c
+++ b/godot-mesa/src/compiler/spirv/spirv_to_nir.c
@@ -45,6 +45,8 @@
#include <stdio.h>
+#include "drivers/d3d12/d3d12_godot_nir_bridge.h"
+
#ifndef PATH_MAX
#define PATH_MAX 4096
#endif
@@ -2481,7 +2483,7 @@ vtn_null_constant(struct vtn_builder *b, struct vtn_type *type)
}
static void
-spec_constant_decoration_cb(struct vtn_builder *b, UNUSED struct vtn_value *val,
+spec_constant_decoration_cb(struct vtn_builder *b, struct vtn_value *val,
ASSERTED int member,
const struct vtn_decoration *dec, void *data)
{
@@ -2489,13 +2491,8 @@ spec_constant_decoration_cb(struct vtn_builder *b, UNUSED struct vtn_value *val,
if (dec->decoration != SpvDecorationSpecId)
return;
- nir_const_value *value = data;
- for (unsigned i = 0; i < b->num_specializations; i++) {
- if (b->specializations[i].id == dec->operands[0]) {
- *value = b->specializations[i].value;
- return;
- }
- }
+ val->is_sc = true;
+ val->sc_id = dec->operands[0];
}
static void
@@ -2519,6 +2516,12 @@ vtn_handle_constant(struct vtn_builder *b, SpvOp opcode,
const uint32_t *w, unsigned count)
{
struct vtn_value *val = vtn_push_value(b, w[2], vtn_value_type_constant);
+ if (opcode == SpvOpSpecConstantComposite || opcode == SpvOpSpecConstantOp) {
+ val->value_type = vtn_value_type_ssa;
+ val->ssa = NULL;
+ return;
+ }
+
val->constant = rzalloc(b, nir_constant);
switch (opcode) {
case SpvOpConstantTrue:
@@ -2536,7 +2539,7 @@ vtn_handle_constant(struct vtn_builder *b, SpvOp opcode,
if (opcode == SpvOpSpecConstantTrue ||
opcode == SpvOpSpecConstantFalse)
- vtn_foreach_decoration(b, val, spec_constant_decoration_cb, &u32val);
+ vtn_foreach_decoration(b, val, spec_constant_decoration_cb, NULL);
val->constant->values[0].b = u32val.u32 != 0;
break;
@@ -2567,14 +2570,12 @@ vtn_handle_constant(struct vtn_builder *b, SpvOp opcode,
if (opcode == SpvOpSpecConstant)
vtn_foreach_decoration(b, val, spec_constant_decoration_cb,
- &val->constant->values[0]);
+ NULL);
break;
}
- case SpvOpSpecConstantComposite:
case SpvOpConstantComposite:
- case SpvOpConstantCompositeReplicateEXT:
- case SpvOpSpecConstantCompositeReplicateEXT: {
+ case SpvOpConstantCompositeReplicateEXT: {
const unsigned elem_count =
val->type->base_type == vtn_base_type_cooperative_matrix ?
1 : val->type->length;
@@ -2648,315 +2649,6 @@ vtn_handle_constant(struct vtn_builder *b, SpvOp opcode,
break;
}
- case SpvOpSpecConstantOp: {
- 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 SpvOpBitcast: {
- struct vtn_value *src = &b->values[w[4]];
-
- vtn_assert(src->value_type == vtn_value_type_constant ||
- src->value_type == vtn_value_type_undef);
-
- unsigned src_len = glsl_get_vector_elements(src->type->type);
- unsigned dst_len = glsl_get_vector_elements(val->type->type);
-
- unsigned src_bit_size = glsl_get_bit_size(src->type->type);
- unsigned dst_bit_size = glsl_get_bit_size(val->type->type);
-
- vtn_assert(src_len * src_bit_size == dst_len * dst_bit_size);
-
- /* This will end up being zero */
- if (src->value_type == vtn_value_type_undef)
- break;
-
- if (src_bit_size == dst_bit_size) {
- /* This is just a copy */
- for (unsigned i = 0; i < src_len; i++)
- val->constant->values[i] = src->constant->values[i];
- } else {
- /* You can't non-trivially bitcast booleans */
- vtn_assert(src_bit_size >= 8 && dst_bit_size >= 8);
- const unsigned src_byte_size = src_bit_size / 8;
- const unsigned dst_byte_size = dst_bit_size / 8;
-
- vtn_assert(src_len <= NIR_MAX_VEC_COMPONENTS &&
- dst_len <= NIR_MAX_VEC_COMPONENTS);
-
- uint8_t bits[NIR_MAX_VEC_COMPONENTS * sizeof(nir_const_value)];
-
- for (unsigned i = 0; i < src_len; i++) {
- uint64_t v = nir_const_value_as_int(src->constant->values[i],
- src_bit_size);
- memcpy(bits + i * src_byte_size, &v, src_byte_size);
- }
-
- for (unsigned i = 0; i < dst_len; i++) {
- uint64_t v = 0;
- memcpy(&v, bits + i * dst_byte_size, dst_byte_size);
- val->constant->values[i] =
- nir_const_value_for_uint(v, dst_bit_size);
- }
- }
- break;
- }
- case SpvOpVectorShuffle: {
- struct vtn_value *v0 = &b->values[w[4]];
- struct vtn_value *v1 = &b->values[w[5]];
-
- vtn_assert(v0->value_type == vtn_value_type_constant ||
- v0->value_type == vtn_value_type_undef);
- vtn_assert(v1->value_type == vtn_value_type_constant ||
- v1->value_type == vtn_value_type_undef);
-
- unsigned len0 = glsl_get_vector_elements(v0->type->type);
- unsigned len1 = glsl_get_vector_elements(v1->type->type);
-
- vtn_assert(len0 + len1 < 16);
-
- unsigned bit_size = glsl_get_bit_size(val->type->type);
- unsigned bit_size0 = glsl_get_bit_size(v0->type->type);
- unsigned bit_size1 = glsl_get_bit_size(v1->type->type);
-
- vtn_assert(bit_size == bit_size0 && bit_size == bit_size1);
- (void)bit_size0; (void)bit_size1;
-
- nir_const_value undef = { .u64 = 0xdeadbeefdeadbeef };
- nir_const_value combined[NIR_MAX_VEC_COMPONENTS * 2];
-
- if (v0->value_type == vtn_value_type_constant) {
- for (unsigned i = 0; i < len0; i++)
- combined[i] = v0->constant->values[i];
- }
- if (v1->value_type == vtn_value_type_constant) {
- for (unsigned i = 0; i < len1; i++)
- combined[len0 + i] = v1->constant->values[i];
- }
-
- for (unsigned i = 0, j = 0; i < count - 6; i++, j++) {
- uint32_t comp = w[i + 6];
- if (comp == (uint32_t)-1) {
- /* If component is not used, set the value to a known constant
- * to detect if it is wrongly used.
- */
- val->constant->values[j] = undef;
- } else {
- vtn_fail_if(comp >= len0 + len1,
- "All Component literals must either be FFFFFFFF "
- "or in [0, N - 1] (inclusive).");
- val->constant->values[j] = combined[comp];
- }
- }
- break;
- }
-
- case SpvOpCompositeExtract:
- case SpvOpCompositeInsert: {
- struct vtn_value *comp;
- unsigned deref_start;
- struct nir_constant **c;
- if (opcode == SpvOpCompositeExtract) {
- comp = vtn_value(b, w[4], vtn_value_type_constant);
- deref_start = 5;
- c = &comp->constant;
- } else {
- comp = vtn_value(b, w[5], vtn_value_type_constant);
- deref_start = 6;
- val->constant = nir_constant_clone(comp->constant, b->shader);
- c = &val->constant;
- }
-
- int elem = -1;
- const struct vtn_type *type = comp->type;
- for (unsigned i = deref_start; i < count; i++) {
- if (type->base_type == vtn_base_type_cooperative_matrix) {
- /* Cooperative matrices are always scalar constants. We don't
- * care about the index w[i] because it's always replicated.
- */
- type = type->component_type;
- } else {
- vtn_fail_if(w[i] > type->length,
- "%uth index of %s is %u but the type has only "
- "%u elements", i - deref_start,
- spirv_op_to_string(opcode), w[i], type->length);
-
- switch (type->base_type) {
- case vtn_base_type_vector:
- elem = w[i];
- type = type->array_element;
- break;
-
- case vtn_base_type_matrix:
- case vtn_base_type_array:
- c = &(*c)->elements[w[i]];
- type = type->array_element;
- break;
-
- case vtn_base_type_struct:
- c = &(*c)->elements[w[i]];
- type = type->members[w[i]];
- break;
-
- default:
- vtn_fail("%s must only index into composite types",
- spirv_op_to_string(opcode));
- }
- }
- }
-
- if (opcode == SpvOpCompositeExtract) {
- if (elem == -1) {
- val->constant = *c;
- } else {
- unsigned num_components = type->length;
- for (unsigned i = 0; i < num_components; i++)
- val->constant->values[i] = (*c)->values[elem + i];
- }
- } else {
- struct vtn_value *insert =
- vtn_value(b, w[4], vtn_value_type_constant);
- vtn_assert(insert->type == type);
- if (elem == -1) {
- *c = insert->constant;
- } else {
- unsigned num_components = type->length;
- for (unsigned i = 0; i < num_components; i++)
- (*c)->values[elem + i] = insert->constant->values[i];
- }
- }
- break;
- }
-
- default: {
- bool swap;
-
- const glsl_type *org_dst_type = val->type->type;
- const glsl_type *org_src_type = org_dst_type;
-
- const bool saturate = vtn_has_decoration(b, val, SpvDecorationSaturatedToLargestFloat8NormalConversionEXT);
- unsigned num_components = glsl_get_vector_elements(val->type->type);
-
- vtn_assert(count <= 7);
-
- switch (opcode) {
- case SpvOpSConvert:
- case SpvOpFConvert:
- case SpvOpUConvert:
- /* We have a different source type in a conversion. */
- org_src_type = vtn_get_value_type(b, w[4])->type;
- break;
- default:
- break;
- };
-
- const glsl_type *dst_type = org_dst_type;
- if (glsl_type_is_bfloat_16(dst_type) || glsl_type_is_e4m3fn(dst_type) || glsl_type_is_e5m2(dst_type))
- dst_type = glsl_float_type();
-
- const glsl_type *src_type = org_src_type;
- if (glsl_type_is_bfloat_16(src_type) || glsl_type_is_e4m3fn(src_type) || glsl_type_is_e5m2(src_type))
- src_type = glsl_float_type();
-
- bool exact;
- nir_op op = vtn_nir_alu_op_for_spirv_opcode(b, opcode, &swap, &exact,
- src_type, dst_type);
-
- /* No SPIR-V opcodes handled through this path should set exact.
- * Since it is ignored, assert on it.
- */
- assert(!exact);
-
- unsigned bit_size = glsl_get_bit_size(dst_type);
- nir_const_value src[3][NIR_MAX_VEC_COMPONENTS];
-
- for (unsigned i = 0; i < count - 4; i++) {
- 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])) {
- if (org_src_type != src_type) {
- /* Small float conversion. */
- assert(i == 0);
- bit_size = glsl_get_bit_size(src_type);
- } else {
- bit_size = glsl_get_bit_size(src_val->type->type);
- }
- }
-
- unsigned src_comps = nir_op_infos[op].input_sizes[i] ?
- nir_op_infos[op].input_sizes[i] :
- num_components;
-
- unsigned j = swap ? 1 - i : i;
- for (unsigned c = 0; c < src_comps; c++) {
- src[j][c] = src_val->constant->values[c];
- if (glsl_type_is_bfloat_16(org_src_type))
- src[j][c].f32 = _mesa_bfloat16_bits_to_float(src[j][c].u16);
- else if (glsl_type_is_e4m3fn(org_src_type))
- src[j][c].f32 = _mesa_e4m3fn_to_float(src[j][c].u8);
- else if (glsl_type_is_e5m2(org_src_type))
- src[j][c].f32 = _mesa_e5m2_to_float(src[j][c].u8);
- }
- }
-
- /* 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][i].u32 = src[1][i].u64; break;
- case 16: src[1][i].u32 = src[1][i].u16; break;
- case 8: src[1][i].u32 = src[1][i].u8; break;
- }
- }
- break;
- }
- default:
- break;
- }
-
- nir_const_value *srcs[3] = {
- src[0], src[1], src[2],
- };
- nir_eval_const_opcode(op, val->constant->values,
- num_components, bit_size, srcs,
- b->shader->info.float_controls_execution_mode);
-
- for (int i = 0; i < num_components; i++) {
- uint16_t conv;
- if (glsl_type_is_bfloat_16(org_dst_type)) {
- conv = _mesa_float_to_bfloat16_bits_rte(val->constant->values[i].f32);
- } else if (glsl_type_is_e4m3fn(org_dst_type)) {
- if (saturate)
- conv = _mesa_float_to_e4m3fn_sat(val->constant->values[i].f32);
- else
- conv = _mesa_float_to_e4m3fn(val->constant->values[i].f32);
- } else if (glsl_type_is_e5m2(org_dst_type)) {
- if (saturate)
- conv = _mesa_float_to_e5m2_sat(val->constant->values[i].f32);
- else
- conv = _mesa_float_to_e5m2(val->constant->values[i].f32);
- } else {
- continue;
- }
-
- val->constant->values[i] = nir_const_value_for_raw_uint(conv, glsl_get_bit_size(org_dst_type));
- }
-
- break;
- } /* default */
- }
- break;
- }
-
case SpvOpConstantNull:
val->constant = vtn_null_constant(b, val->type);
val->is_null_constant = true;
@@ -7018,6 +6710,93 @@ vtn_handle_body_instruction(struct vtn_builder *b, SpvOp opcode,
return true;
}
+static bool
+vtn_handle_spec_constant_instructions(struct vtn_builder* b, SpvOp opcode,
+ const uint32_t* w, unsigned count)
+{
+ switch (opcode) {
+ case SpvOpSpecConstantTrue:
+ case SpvOpSpecConstantFalse:
+ case SpvOpSpecConstant:
+ case SpvOpSpecConstantComposite:
+ case SpvOpSpecConstantOp:
+ break;
+ default:
+ return true;
+ }
+
+ struct vtn_value* val = vtn_untyped_value(b, w[2]);
+
+ switch (opcode) {
+ case SpvOpSpecConstantTrue:
+ case SpvOpSpecConstantFalse:
+ case SpvOpSpecConstant: {
+ vtn_assert(val->is_sc);
+ vtn_assert(val->value_type == vtn_value_type_constant || val->value_type == vtn_value_type_ssa);
+
+ val->value_type = vtn_value_type_ssa;
+ val->ssa = vtn_create_ssa_value(b, val->type->type);
+
+ nir_def *sc_imm = nir_imm_int(&b->nb, GODOT_NIR_SC_SENTINEL_MAGIC | val->sc_id);
+ nir_def *non_opt_const = nir_load_constant_non_opt(&b->nb, 1, 32, sc_imm);
+
+ vtn_assert(b->nb.cursor.option == nir_cursor_after_instr);
+ vtn_assert(b->nb.cursor.instr->type == nir_instr_type_intrinsic);
+
+ val->ssa = vtn_create_ssa_value(b, val->type->type);
+ if (val->type->type == glsl_uint_type()) {
+ val->ssa->def = non_opt_const;
+ } else if (val->type->type == glsl_bool_type()) {
+ val->ssa->def = nir_build_alu(
+ &b->nb,
+ nir_op_ine,
+ non_opt_const,
+ nir_imm_int(&b->nb, 0),
+ NULL,
+ NULL);
+ } else if (val->type->type == glsl_float_type()) {
+ val->ssa->def = non_opt_const;
+ } else {
+ vtn_assert(false);
+ }
+ } break;
+
+ case SpvOpSpecConstantComposite: {
+ unsigned elem_count = count - 3;
+ vtn_fail_if(elem_count != val->type->length,
+ "%s has %u constituents, expected %u",
+ spirv_op_to_string(opcode), elem_count, val->type->length);
+
+ vtn_assert(b->values[w[2]].value_type == vtn_value_type_ssa);
+ if (!b->values[w[2]].ssa) {
+ b->values[w[2]].value_type = vtn_value_type_invalid; /* Pretend not yet set */
+ vtn_handle_composite(b, SpvOpCompositeConstruct, w, count);
+ }
+ break;
+ }
+
+ case SpvOpSpecConstantOp: {
+ vtn_assert(val->value_type == vtn_value_type_ssa);
+ val->value_type = vtn_value_type_invalid;
+
+ unsigned count = (w[0] >> SpvWordCountShift) - 1;
+ uint32_t* sub_w = (uint32_t*)alloca(4 * count);
+ sub_w[0] = 0; /* Doesn't really matter */
+ sub_w[1] = val->type->id;
+ sub_w[2] = w[2];
+ SpvOp sub_opcode = w[3];
+ for (unsigned i = 0; i < count - 3; ++i)
+ sub_w[3 + i] = w[4 + i];
+ vtn_handle_body_instruction(b, sub_opcode, sub_w, count);
+ } break;
+
+ default:
+ return false; /* End of preamble */
+ }
+
+ return true;
+}
+
static bool
is_glslang(const struct vtn_builder *b)
{
@@ -7376,6 +7155,8 @@ spirv_to_nir(const uint32_t *words, size_t word_count,
/* Skip the SPIR-V header, handled at vtn_create_builder */
words+= 5;
+ const uint32_t *preamble_words = words;
+
/* Handle all the preamble instructions */
words = vtn_foreach_instruction(b, words, word_end,
vtn_handle_preamble_instruction);
@@ -7461,7 +7242,7 @@ spirv_to_nir(const uint32_t *words, size_t word_count,
vtn_foreach_function(func, &b->functions) {
if ((options->create_library || func->referenced) && !func->emitted) {
_mesa_hash_table_clear(b->strings, NULL);
- vtn_function_emit(b, func, vtn_handle_body_instruction);
+ vtn_function_emit(b, func, vtn_handle_spec_constant_instructions, preamble_words, vtn_handle_body_instruction);
progress = true;
}
}
diff --git a/godot-mesa/src/compiler/spirv/vtn_cfg.c b/godot-mesa/src/compiler/spirv/vtn_cfg.c
index 83d3228..78aee17 100644
--- a/godot-mesa/src/compiler/spirv/vtn_cfg.c
+++ b/godot-mesa/src/compiler/spirv/vtn_cfg.c
@@ -763,6 +763,8 @@ vtn_emit_cf_func_unstructured(struct vtn_builder *b, struct vtn_function *func,
void
vtn_function_emit(struct vtn_builder *b, struct vtn_function *func,
+ vtn_instruction_handler preamble_instruction_handler,
+ const uint32_t *preamble_words,
vtn_instruction_handler instruction_handler)
{
static int force_unstructured = -1;
@@ -777,6 +779,9 @@ vtn_function_emit(struct vtn_builder *b, struct vtn_function *func,
b->nb.exact = b->exact;
b->phi_table = _mesa_pointer_hash_table_create(b);
+ const uint32_t *word_end = b->spirv + b->spirv_word_count;
+ vtn_foreach_instruction(b, preamble_words, word_end, preamble_instruction_handler);
+
if (b->shader->info.stage == MESA_SHADER_KERNEL || force_unstructured) {
impl->structured = false;
vtn_emit_cf_func_unstructured(b, func, instruction_handler);
diff --git a/godot-mesa/src/compiler/spirv/vtn_private.h b/godot-mesa/src/compiler/spirv/vtn_private.h
index 5d601f9..4ab1a22 100644
--- a/godot-mesa/src/compiler/spirv/vtn_private.h
+++ b/godot-mesa/src/compiler/spirv/vtn_private.h
@@ -243,6 +243,8 @@ typedef bool (*vtn_instruction_handler)(struct vtn_builder *, SpvOp,
void vtn_build_cfg(struct vtn_builder *b, const uint32_t *words,
const uint32_t *end);
void vtn_function_emit(struct vtn_builder *b, struct vtn_function *func,
+ vtn_instruction_handler preamble_instruction_handler,
+ const uint32_t *preamble_words,
vtn_instruction_handler instruction_handler);
void vtn_handle_function_call(struct vtn_builder *b, SpvOp opcode,
const uint32_t *w, unsigned count);
@@ -521,6 +523,8 @@ struct vtn_variable {
unsigned descriptor_set;
unsigned binding;
bool explicit_binding;
+ unsigned orig_descriptor_set;
+ unsigned orig_binding;
unsigned offset;
unsigned input_attachment_index;
@@ -594,6 +598,9 @@ struct vtn_value {
struct vtn_ssa_value *ssa;
vtn_instruction_handler ext_handler;
};
+
+ bool is_sc;
+ uint32_t sc_id;
};
#define VTN_DEC_DECORATION -1
diff --git a/godot-mesa/src/compiler/spirv/vtn_variables.c b/godot-mesa/src/compiler/spirv/vtn_variables.c
index 55557e7..16866df 100644
--- a/godot-mesa/src/compiler/spirv/vtn_variables.c
+++ b/godot-mesa/src/compiler/spirv/vtn_variables.c
@@ -26,6 +26,8 @@
#include "nir_deref.h"
#include <vulkan/vulkan_core.h>
+#include "drivers/d3d12/d3d12_godot_nir_bridge.h"
+
static struct vtn_pointer*
vtn_align_pointer(struct vtn_builder *b, struct vtn_pointer *ptr,
unsigned alignment)
@@ -1620,13 +1622,17 @@ var_decoration_cb(struct vtn_builder *b, struct vtn_value *val, int member,
/* Handle decorations that apply to a vtn_variable as a whole */
switch (dec->decoration) {
+ case SpvDecorationDescriptorSet:
case SpvDecorationBinding:
- vtn_var->binding = dec->operands[0];
+ if (dec->decoration == SpvDecorationDescriptorSet) {
+ vtn_var->orig_descriptor_set = dec->operands[0];
+ } else {
+ vtn_var->orig_binding = dec->operands[0];
+ }
+ vtn_var->descriptor_set = 0;
+ vtn_var->binding = vtn_var->orig_descriptor_set * GODOT_NIR_DESCRIPTOR_SET_MULTIPLIER + vtn_var->orig_binding * GODOT_NIR_BINDING_MULTIPLIER;
vtn_var->explicit_binding = true;
return;
- case SpvDecorationDescriptorSet:
- vtn_var->descriptor_set = dec->operands[0];
- return;
case SpvDecorationInputAttachmentIndex:
vtn_var->input_attachment_index = dec->operands[0];
vtn_var->access |= ACCESS_NON_WRITEABLE;
diff --git a/godot-mesa/src/microsoft/compiler/dxil_container.c b/godot-mesa/src/microsoft/compiler/dxil_container.c
index 77a14cd..24d8362 100644
--- a/godot-mesa/src/microsoft/compiler/dxil_container.c
+++ b/godot-mesa/src/microsoft/compiler/dxil_container.c
@@ -331,7 +331,8 @@ dxil_container_add_state_validation(struct dxil_container *c,
bool
dxil_container_add_module(struct dxil_container *c,
- const struct dxil_module *m)
+ const struct dxil_module *m,
+ uint64_t *bitcode_bit_offset)
{
assert(m->buf.buf_bits == 0); // make sure the module is fully flushed
uint32_t version = (m->shader_kind << 16) |
@@ -345,18 +346,22 @@ dxil_container_add_module(struct dxil_container *c,
uint32_t bitcode_offset = 16;
uint32_t bitcode_size = m->buf.blob.size;
- return add_part_header(c, DXIL_DXIL, size) &&
+ if (!(add_part_header(c, DXIL_DXIL, size) &&
blob_write_bytes(&c->parts, &version, sizeof(version)) &&
blob_write_bytes(&c->parts, &uint32_size, sizeof(uint32_size)) &&
blob_write_bytes(&c->parts, &magic, sizeof(magic)) &&
blob_write_bytes(&c->parts, &dxil_version, sizeof(dxil_version)) &&
blob_write_bytes(&c->parts, &bitcode_offset, sizeof(bitcode_offset)) &&
- blob_write_bytes(&c->parts, &bitcode_size, sizeof(bitcode_size)) &&
- blob_write_bytes(&c->parts, m->buf.blob.data, m->buf.blob.size);
+ blob_write_bytes(&c->parts, &bitcode_size, sizeof(bitcode_size))))
+ return false;
+
+ *bitcode_bit_offset += c->parts.size * 8;
+
+ return blob_write_bytes(&c->parts, m->buf.blob.data, m->buf.blob.size);
}
bool
-dxil_container_write(struct dxil_container *c, struct blob *blob)
+dxil_container_write(struct dxil_container *c, struct blob *blob, uint64_t *bitcode_bit_offset)
{
assert(blob->size == 0);
if (!blob_write_bytes(blob, &DXIL_DXBC, sizeof(DXIL_DXBC)))
@@ -387,8 +392,12 @@ dxil_container_write(struct dxil_container *c, struct blob *blob)
}
if (!blob_write_bytes(blob, &c->num_parts, sizeof(c->num_parts)) ||
- !blob_write_bytes(blob, part_offsets, sizeof(uint32_t) * c->num_parts) ||
- !blob_write_bytes(blob, c->parts.data, c->parts.size))
+ !blob_write_bytes(blob, part_offsets, sizeof(uint32_t) * c->num_parts))
+ return false;
+
+ *bitcode_bit_offset += blob->size * 8;
+
+ if (!blob_write_bytes(blob, c->parts.data, c->parts.size))
return false;
return true;
diff --git a/godot-mesa/src/microsoft/compiler/dxil_container.h b/godot-mesa/src/microsoft/compiler/dxil_container.h
index 2c3f17c..3eb684f 100644
--- a/godot-mesa/src/microsoft/compiler/dxil_container.h
+++ b/godot-mesa/src/microsoft/compiler/dxil_container.h
@@ -123,10 +123,11 @@ dxil_container_add_state_validation(struct dxil_container *c,
bool
dxil_container_add_module(struct dxil_container *c,
- const struct dxil_module *m);
+ const struct dxil_module *m,
+ uint64_t *bitcode_bit_offset);
bool
-dxil_container_write(struct dxil_container *c, struct blob *blob);
+dxil_container_write(struct dxil_container *c, struct blob *blob, uint64_t *bitcode_bit_offset);
#ifdef __cplusplus
}
diff --git a/godot-mesa/src/microsoft/compiler/dxil_module.c b/godot-mesa/src/microsoft/compiler/dxil_module.c
index c70c67c..0634b52 100644
--- a/godot-mesa/src/microsoft/compiler/dxil_module.c
+++ b/godot-mesa/src/microsoft/compiler/dxil_module.c
@@ -32,6 +32,8 @@
#include <assert.h>
#include <stdio.h>
+#include "drivers/d3d12/d3d12_godot_nir_bridge.h"
+
void
dxil_module_init(struct dxil_module *m, void *ralloc_ctx)
{
@@ -2669,6 +2671,12 @@ emit_consts(struct dxil_module *m)
continue;
}
+ if (curr_type->type == TYPE_INTEGER && (c->int_value & GODOT_NIR_SC_SENTINEL_MAGIC_MASK) == GODOT_NIR_SC_SENTINEL_MAGIC) {
+ uint32_t sc_id = (uint32_t)(c->int_value & ~GODOT_NIR_SC_SENTINEL_MAGIC_MASK);
+ uint64_t sc_bit_offset = (uint64_t)m->buf.blob.size * 8 + m->buf.buf_bits + m->buf.abbrev_width;
+ m->godot_nir_callbacks->report_sc_bit_offset_fn(sc_id, sc_bit_offset, m->godot_nir_callbacks->data);
+ }
+
switch (curr_type->type) {
case TYPE_INTEGER:
if (!emit_int_value(m, c->int_value))
diff --git a/godot-mesa/src/microsoft/compiler/dxil_module.h b/godot-mesa/src/microsoft/compiler/dxil_module.h
index 08ba263..d4c412d 100644
--- a/godot-mesa/src/microsoft/compiler/dxil_module.h
+++ b/godot-mesa/src/microsoft/compiler/dxil_module.h
@@ -29,6 +29,8 @@
#ifndef DXIL_MODULE_H
#define DXIL_MODULE_H
+typedef struct GodotNirCallbacks GodotNirCallbacks;
+
#ifdef __cplusplus
extern "C" {
#endif
@@ -268,6 +270,8 @@ struct dxil_module {
struct rb_tree *functions;
struct dxil_func_def *cur_emitting_func;
+
+ const GodotNirCallbacks *godot_nir_callbacks;
};
struct dxil_instr;
diff --git a/godot-mesa/src/microsoft/compiler/nir_to_dxil.c b/godot-mesa/src/microsoft/compiler/nir_to_dxil.c
index f99cf69..fa5bec7 100644
--- a/godot-mesa/src/microsoft/compiler/nir_to_dxil.c
+++ b/godot-mesa/src/microsoft/compiler/nir_to_dxil.c
@@ -44,6 +44,8 @@
#include <stdint.h>
+#include "drivers/d3d12/d3d12_godot_nir_bridge.h"
+
int debug_dxil = 0;
static const struct debug_named_value
@@ -1226,6 +1228,8 @@ add_resource(struct ntd_context *ctx, enum dxil_resource_type type,
/* No flags supported yet */
resource_v1->resource_flags = 0;
}
+
+ ctx->opts->godot_nir_callbacks->report_resource(layout->binding, layout->space, (uint32_t)type, ctx->opts->godot_nir_callbacks->data);
}
static const struct dxil_value *
@@ -4963,6 +4967,12 @@ emit_intrinsic(struct ntd_context *ctx, nir_intrinsic_instr *intr)
return emit_load_unary_external_function(ctx, intr, "dx.op.startInstanceLocation",
DXIL_INTR_START_INSTANCE_LOCATION, nir_type_int);
+ case nir_intrinsic_load_constant_non_opt: {
+ const struct dxil_value* value = get_src(ctx, &intr->src[0], 0, nir_type_uint);
+ store_def(ctx, &intr->def, 0, value);
+ return true;
+ }
+
case nir_intrinsic_load_num_workgroups:
case nir_intrinsic_load_workgroup_size:
default:
@@ -6618,6 +6628,7 @@ nir_to_dxil(struct nir_shader *s, const struct nir_to_dxil_options *opts,
MIN2(opts->shader_model_max & 0xffff, validator_version & 0xffff);
ctx->mod.major_validator = validator_version >> 16;
ctx->mod.minor_validator = validator_version & 0xffff;
+ ctx->mod.godot_nir_callbacks = opts->godot_nir_callbacks;
if (s->info.stage <= MESA_SHADER_FRAGMENT) {
uint64_t in_mask =
@@ -6773,19 +6784,23 @@ nir_to_dxil(struct nir_shader *s, const struct nir_to_dxil_options *opts,
goto out;
}
- if (!dxil_container_add_module(&container, &ctx->mod)) {
+ uint64_t bitcode_bit_offset = 0;
+
+ if (!dxil_container_add_module(&container, &ctx->mod, &bitcode_bit_offset)) {
debug_printf("D3D12: failed to write module\n");
retval = false;
goto out;
}
- if (!dxil_container_write(&container, blob)) {
+ if (!dxil_container_write(&container, blob, &bitcode_bit_offset)) {
debug_printf("D3D12: dxil_container_write failed\n");
retval = false;
goto out;
}
dxil_container_finish(&container);
+ opts->godot_nir_callbacks->report_bitcode_bit_offset_fn(bitcode_bit_offset, opts->godot_nir_callbacks->data);
+
if (debug_dxil & DXIL_DEBUG_DUMP_BLOB) {
static int shader_id = 0;
char buffer[64];
diff --git a/godot-mesa/src/microsoft/compiler/nir_to_dxil.h b/godot-mesa/src/microsoft/compiler/nir_to_dxil.h
index 9890115..abe9e9d 100644
--- a/godot-mesa/src/microsoft/compiler/nir_to_dxil.h
+++ b/godot-mesa/src/microsoft/compiler/nir_to_dxil.h
@@ -29,6 +29,8 @@
#include "nir.h"
#include "dxil_versions.h"
+typedef struct GodotNirCallbacks GodotNirCallbacks;
+
#ifdef __cplusplus
extern "C" {
#endif
@@ -87,6 +89,7 @@ struct nir_to_dxil_options {
enum dxil_environment environment;
enum dxil_shader_model shader_model_max;
enum dxil_validator_version validator_version_max;
+ const GodotNirCallbacks *godot_nir_callbacks;
};
typedef void (*dxil_msg_callback)(void *priv, const char *msg);
diff --git a/godot-mesa/src/microsoft/spirv_to_dxil/dxil_spirv_nir.c b/godot-mesa/src/microsoft/spirv_to_dxil/dxil_spirv_nir.c
index 66a996e..4d8eff1 100644
--- a/godot-mesa/src/microsoft/spirv_to_dxil/dxil_spirv_nir.c
+++ b/godot-mesa/src/microsoft/spirv_to_dxil/dxil_spirv_nir.c
@@ -1058,30 +1058,7 @@ dxil_spirv_nir_passes(nir_shader *nir,
NIR_PASS(_, nir, nir_lower_alu_to_scalar, NULL, NULL);
NIR_PASS(_, nir, nir_opt_dce);
NIR_PASS(_, nir, dxil_nir_lower_double_math);
-
- {
- bool progress;
- do
- {
- progress = false;
- NIR_PASS(progress, nir, nir_copy_prop);
- NIR_PASS(progress, nir, nir_opt_copy_prop_vars);
- NIR_PASS(progress, nir, nir_opt_deref);
- NIR_PASS(progress, nir, nir_opt_dce);
- NIR_PASS(progress, nir, nir_opt_undef);
- NIR_PASS(progress, nir, nir_opt_constant_folding);
- NIR_PASS(progress, nir, nir_opt_cse);
- if (nir_opt_loop(nir)) {
- progress = true;
- NIR_PASS(progress, nir, nir_copy_prop);
- NIR_PASS(progress, nir, nir_opt_dce);
- }
- NIR_PASS(progress, nir, nir_lower_vars_to_ssa);
- NIR_PASS(progress, nir, nir_opt_algebraic);
- NIR_PASS(progress, nir, nir_opt_dead_cf);
- NIR_PASS(progress, nir, nir_opt_remove_phis);
- } while (progress);
- }
+ NIR_PASS(_, nir, nir_lower_vars_to_ssa);
NIR_PASS(_, nir, nir_remove_dead_variables, nir_var_function_temp, NULL);
NIR_PASS(_, nir, nir_split_struct_vars, nir_var_function_temp);
diff --git a/godot-mesa/src/microsoft/spirv_to_dxil/spirv_to_dxil.c b/godot-mesa/src/microsoft/spirv_to_dxil/spirv_to_dxil.c
index 1a8e6e2..1b3b9d3 100644
--- a/godot-mesa/src/microsoft/spirv_to_dxil/spirv_to_dxil.c
+++ b/godot-mesa/src/microsoft/spirv_to_dxil/spirv_to_dxil.c
@@ -32,6 +32,8 @@
#include "git_sha1.h"
#include "vulkan/vulkan.h"
+#include "drivers/d3d12/d3d12_godot_nir_bridge.h"
+
static_assert((mesa_shader_stage)DXIL_SPIRV_SHADER_NONE == MESA_SHADER_NONE, "must match");
static_assert((mesa_shader_stage)DXIL_SPIRV_SHADER_VERTEX == MESA_SHADER_VERTEX, "must match");
static_assert((mesa_shader_stage)DXIL_SPIRV_SHADER_TESS_CTRL == MESA_SHADER_TESS_CTRL, "must match");
@@ -50,6 +52,7 @@ spirv_to_dxil(const uint32_t *words, size_t word_count,
const struct dxil_spirv_debug_options *dgb_opts,
const struct dxil_spirv_runtime_conf *conf,
const struct dxil_spirv_logger *logger,
+ const GodotNirCallbacks *godot_nir_callbacks,
struct dxil_spirv_object *out_dxil)
{
if (stage == DXIL_SPIRV_SHADER_NONE || stage == DXIL_SPIRV_SHADER_KERNEL)
@@ -61,6 +64,7 @@ spirv_to_dxil(const uint32_t *words, size_t word_count,
.environment = DXIL_ENVIRONMENT_VULKAN,
.shader_model_max = conf->shader_model_max,
.validator_version_max = validator_version_max,
+ .godot_nir_callbacks = godot_nir_callbacks,
};
const struct spirv_to_nir_options *spirv_opts = dxil_spirv_nir_get_spirv_options();
diff --git a/godot-mesa/src/microsoft/spirv_to_dxil/spirv_to_dxil.h b/godot-mesa/src/microsoft/spirv_to_dxil/spirv_to_dxil.h
index 00b1884..0d10644 100644
--- a/godot-mesa/src/microsoft/spirv_to_dxil/spirv_to_dxil.h
+++ b/godot-mesa/src/microsoft/spirv_to_dxil/spirv_to_dxil.h
@@ -30,6 +30,8 @@
#include <stddef.h>
#include <stdint.h>
+typedef struct GodotNirCallbacks GodotNirCallbacks;
+
#ifdef __cplusplus
extern "C" {
#endif
@@ -234,6 +236,7 @@ spirv_to_dxil(const uint32_t *words, size_t word_count,
const struct dxil_spirv_debug_options *debug_options,
const struct dxil_spirv_runtime_conf *conf,
const struct dxil_spirv_logger *logger,
+ const GodotNirCallbacks *godot_nir_callbacks,
struct dxil_spirv_object *out_dxil);
/**
diff --git a/godot-mesa/src/util/macros.h b/godot-mesa/src/util/macros.h
index 0070067..72410ef 100644
--- a/godot-mesa/src/util/macros.h
+++ b/godot-mesa/src/util/macros.h
@@ -543,14 +543,14 @@ typedef int lock_cap_t;
* SWAP - swap value of @a and @b
*/
#if !defined(_MSC_VER) || _MSC_VER >= 1939 /* MSVC 17.9 or later for __typeof__ */
-#define SWAP(a, b) \
+#define MESA_SWAP(a, b) \
do { \
__typeof__(a) __tmp = (a); \
(a) = (b); \
(b) = __tmp; \
} while (0)
#else
-#define SWAP(a, b) \
+#define MESA_SWAP(a, b) \
do { \
/* NOLINTBEGIN(bugprone-sizeof-expression) */ \
char __tmp[sizeof(a) == sizeof(b) ? (ptrdiff_t)sizeof(a) : -1]; \
diff --git a/godot-mesa/src/util/ralloc.c b/godot-mesa/src/util/ralloc.c
index ba560c8..07379ae 100644
--- a/godot-mesa/src/util/ralloc.c
+++ b/godot-mesa/src/util/ralloc.c
@@ -35,6 +35,8 @@
#include "ralloc.h"
+#include "drivers/d3d12/d3d12_godot_nir_bridge.h"
+
#define CANARY 0x5A1106
#if defined(__LP64__) || defined(_WIN64)
@@ -115,8 +117,8 @@ ralloc_size(const void *ctx, size_t size)
* - Allocations of a size that rounds up to a multiple of 8 bytes and
* not 16 bytes, are only required to have at least 8 byte alignment.
*/
- void *block = malloc(align64(size + sizeof(ralloc_header),
- alignof(ralloc_header)));
+ void *block = godot_nir_malloc(align64(size + sizeof(ralloc_header),
+ alignof(ralloc_header)));
ralloc_header *info;
ralloc_header *parent;
@@ -164,8 +166,8 @@ resize(void *ptr, size_t size)
ralloc_header *child, *old, *info;
old = get_header(ptr);
- info = realloc(old, align64(size + sizeof(ralloc_header),
- alignof(ralloc_header)));
+ info = godot_nir_realloc(old, align64(size + sizeof(ralloc_header),
+ alignof(ralloc_header)));
if (info == NULL)
return NULL;
@@ -323,7 +325,7 @@ unsafe_free(ralloc_header *info)
if (info->destructor != NULL)
info->destructor(PTR_FROM_HEADER(info));
- free(info);
+ godot_nir_free(info);
}
void
diff --git a/godot-mesa/src/util/set.c b/godot-mesa/src/util/set.c
index d6c0623..1eac1e8 100644
--- a/godot-mesa/src/util/set.c
+++ b/godot-mesa/src/util/set.c
@@ -691,7 +691,7 @@ _mesa_set_intersects(struct set *a, struct set *b)
/* iterate over the set with less entries */
if (b->entries < a->entries) {
- SWAP(a, b);
+ MESA_SWAP(a, b);
}
set_foreach(a, entry) {
Reference in New Issue View Git Blame Copy Permalink