OpenHarmony/third_party_mesa3d
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

spirv: Replace assert with vtn_assert

Browse Source 
Reviewed-by: Tapani Pälli <tapani.palli@intel.com>
Reviewed-by: Ian Romanick <idr@freedesktop.org>
This commit is contained in:
Jason Ekstrand
2017-08-16 17:38:13 -07:00
parent 94ca8e04ad
commit b7ef60d846
5 changed files with 251 additions and 249 deletions
Download Patch File Download Diff File Expand all files Collapse all files

183
src/compiler/spirv/spirv_to_nir.c
View File

@@ -260,7 +260,7 @@ vtn_ssa_value(struct vtn_builder *b, uint32_t value_id)
return val->ssa; return val->ssa;
case vtn_value_type_pointer: case vtn_value_type_pointer:
assert(val->pointer->ptr_type && val->pointer->ptr_type->type); vtn_assert(val->pointer->ptr_type && val->pointer->ptr_type->type);
struct vtn_ssa_value *ssa = struct vtn_ssa_value *ssa =
vtn_create_ssa_value(b, val->pointer->ptr_type->type); vtn_create_ssa_value(b, val->pointer->ptr_type->type);
ssa->def = vtn_pointer_to_ssa(b, val->pointer); ssa->def = vtn_pointer_to_ssa(b, val->pointer);
@@ -296,7 +296,7 @@ vtn_foreach_instruction(struct vtn_builder *b, const uint32_t *start,
while (w < end) { while (w < end) {
SpvOp opcode = w[0] & SpvOpCodeMask; SpvOp opcode = w[0] & SpvOpCodeMask;
unsigned count = w[0] >> SpvWordCountShift; unsigned count = w[0] >> SpvWordCountShift;
assert(count >= 1 && w + count <= end); vtn_assert(count >= 1 && w + count <= end);
b->spirv_offset = (uint8_t *)w - (uint8_t *)b->spirv; b->spirv_offset = (uint8_t *)w - (uint8_t *)b->spirv;
@@ -352,8 +352,7 @@ vtn_handle_extension(struct vtn_builder *b, SpvOp opcode,
case SpvOpExtInst: { case SpvOpExtInst: {
struct vtn_value *val = vtn_value(b, w[3], vtn_value_type_extension); struct vtn_value *val = vtn_value(b, w[3], vtn_value_type_extension);
bool handled = val->ext_handler(b, w[4], w, count); bool handled = val->ext_handler(b, w[4], w, count);
(void)handled; vtn_assert(handled);
assert(handled);
break; break;
} }
@@ -374,7 +373,7 @@ _foreach_decoration_helper(struct vtn_builder *b,
if (dec->scope == VTN_DEC_DECORATION) { if (dec->scope == VTN_DEC_DECORATION) {
member = parent_member; member = parent_member;
} else if (dec->scope >= VTN_DEC_STRUCT_MEMBER0) { } else if (dec->scope >= VTN_DEC_STRUCT_MEMBER0) {
assert(parent_member == -1); vtn_assert(parent_member == -1);
member = dec->scope - VTN_DEC_STRUCT_MEMBER0; member = dec->scope - VTN_DEC_STRUCT_MEMBER0;
} else { } else {
/* Not a decoration */ /* Not a decoration */
@@ -382,7 +381,7 @@ _foreach_decoration_helper(struct vtn_builder *b,
} }
if (dec->group) { if (dec->group) {
assert(dec->group->value_type == vtn_value_type_decoration_group); vtn_assert(dec->group->value_type == vtn_value_type_decoration_group);
_foreach_decoration_helper(b, base_value, member, dec->group, _foreach_decoration_helper(b, base_value, member, dec->group,
cb, data); cb, data);
} else { } else {
@@ -412,7 +411,7 @@ vtn_foreach_execution_mode(struct vtn_builder *b, struct vtn_value *value,
if (dec->scope != VTN_DEC_EXECUTION_MODE) if (dec->scope != VTN_DEC_EXECUTION_MODE)
continue; continue;
assert(dec->group == NULL); vtn_assert(dec->group == NULL);
cb(b, value, dec, data); cb(b, value, dec, data);
} }
} }
@@ -543,7 +542,7 @@ mutable_matrix_member(struct vtn_builder *b, struct vtn_type *type, int member)
type = type->array_element; type = type->array_element;
} }
assert(glsl_type_is_matrix(type->type)); vtn_assert(glsl_type_is_matrix(type->type));
return type; return type;
} }
@@ -558,7 +557,7 @@ struct_member_decoration_cb(struct vtn_builder *b,
if (member < 0) if (member < 0)
return; return;
assert(member < ctx->num_fields); vtn_assert(member < ctx->num_fields);
switch (dec->decoration) { switch (dec->decoration) {
case SpvDecorationNonWritable: case SpvDecorationNonWritable:
@@ -582,7 +581,7 @@ struct_member_decoration_cb(struct vtn_builder *b,
break; break;
case SpvDecorationStream: case SpvDecorationStream:
/* Vulkan only allows one GS stream */ /* Vulkan only allows one GS stream */
assert(dec->literals[0] == 0); vtn_assert(dec->literals[0] == 0);
break; break;
case SpvDecorationLocation: case SpvDecorationLocation:
ctx->fields[member].location = dec->literals[0]; ctx->fields[member].location = dec->literals[0];
@@ -661,7 +660,7 @@ struct_member_matrix_stride_cb(struct vtn_builder *b,
{ {
if (dec->decoration != SpvDecorationMatrixStride) if (dec->decoration != SpvDecorationMatrixStride)
return; return;
assert(member >= 0); vtn_assert(member >= 0);
struct member_decoration_ctx *ctx = void_ctx; struct member_decoration_ctx *ctx = void_ctx;
@@ -671,7 +670,7 @@ struct_member_matrix_stride_cb(struct vtn_builder *b,
mat_type->stride = mat_type->array_element->stride; mat_type->stride = mat_type->array_element->stride;
mat_type->array_element->stride = dec->literals[0]; mat_type->array_element->stride = dec->literals[0];
} else { } else {
assert(mat_type->array_element->stride > 0); vtn_assert(mat_type->array_element->stride > 0);
mat_type->stride = dec->literals[0]; mat_type->stride = dec->literals[0];
} }
} }
@@ -688,17 +687,17 @@ type_decoration_cb(struct vtn_builder *b,
switch (dec->decoration) { switch (dec->decoration) {
case SpvDecorationArrayStride: case SpvDecorationArrayStride:
assert(type->base_type == vtn_base_type_matrix || vtn_assert(type->base_type == vtn_base_type_matrix ||
type->base_type == vtn_base_type_array || type->base_type == vtn_base_type_array ||
type->base_type == vtn_base_type_pointer); type->base_type == vtn_base_type_pointer);
type->stride = dec->literals[0]; type->stride = dec->literals[0];
break; break;
case SpvDecorationBlock: case SpvDecorationBlock:
assert(type->base_type == vtn_base_type_struct); vtn_assert(type->base_type == vtn_base_type_struct);
type->block = true; type->block = true;
break; break;
case SpvDecorationBufferBlock: case SpvDecorationBufferBlock:
assert(type->base_type == vtn_base_type_struct); vtn_assert(type->base_type == vtn_base_type_struct);
type->buffer_block = true; type->buffer_block = true;
break; break;
case SpvDecorationGLSLShared: case SpvDecorationGLSLShared:
@@ -850,7 +849,7 @@ vtn_handle_type(struct vtn_builder *b, SpvOp opcode,
struct vtn_type *base = vtn_value(b, w[2], vtn_value_type_type)->type; struct vtn_type *base = vtn_value(b, w[2], vtn_value_type_type)->type;
unsigned elems = w[3]; unsigned elems = w[3];
assert(glsl_type_is_scalar(base->type)); vtn_assert(glsl_type_is_scalar(base->type));
val->type->base_type = vtn_base_type_vector; val->type->base_type = vtn_base_type_vector;
val->type->type = glsl_vector_type(glsl_get_base_type(base->type), elems); val->type->type = glsl_vector_type(glsl_get_base_type(base->type), elems);
val->type->stride = glsl_get_bit_size(base->type) / 8; val->type->stride = glsl_get_bit_size(base->type) / 8;
@@ -862,12 +861,12 @@ vtn_handle_type(struct vtn_builder *b, SpvOp opcode,
struct vtn_type *base = vtn_value(b, w[2], vtn_value_type_type)->type; struct vtn_type *base = vtn_value(b, w[2], vtn_value_type_type)->type;
unsigned columns = w[3]; unsigned columns = w[3];
assert(glsl_type_is_vector(base->type)); vtn_assert(glsl_type_is_vector(base->type));
val->type->base_type = vtn_base_type_matrix; val->type->base_type = vtn_base_type_matrix;
val->type->type = glsl_matrix_type(glsl_get_base_type(base->type), val->type->type = glsl_matrix_type(glsl_get_base_type(base->type),
glsl_get_vector_elements(base->type), glsl_get_vector_elements(base->type),
columns); columns);
assert(!glsl_type_is_error(val->type->type)); vtn_assert(!glsl_type_is_error(val->type->type));
val->type->length = columns; val->type->length = columns;
val->type->array_element = base; val->type->array_element = base;
val->type->row_major = false; val->type->row_major = false;
@@ -969,7 +968,7 @@ vtn_handle_type(struct vtn_builder *b, SpvOp opcode,
const struct glsl_type *sampled_type = const struct glsl_type *sampled_type =
vtn_value(b, w[2], vtn_value_type_type)->type->type; vtn_value(b, w[2], vtn_value_type_type)->type->type;
assert(glsl_type_is_vector_or_scalar(sampled_type)); vtn_assert(glsl_type_is_vector_or_scalar(sampled_type));
enum glsl_sampler_dim dim; enum glsl_sampler_dim dim;
switch ((SpvDim)w[3]) { switch ((SpvDim)w[3]) {
@@ -1011,7 +1010,7 @@ vtn_handle_type(struct vtn_builder *b, SpvOp opcode,
val->type->type = glsl_sampler_type(dim, is_shadow, is_array, val->type->type = glsl_sampler_type(dim, is_shadow, is_array,
glsl_get_base_type(sampled_type)); glsl_get_base_type(sampled_type));
} else if (sampled == 2) { } else if (sampled == 2) {
assert(!is_shadow); vtn_assert(!is_shadow);
val->type->sampled = false; val->type->sampled = false;
val->type->type = glsl_image_type(dim, is_array, val->type->type = glsl_image_type(dim, is_array,
glsl_get_base_type(sampled_type)); glsl_get_base_type(sampled_type));
@@ -1071,7 +1070,7 @@ vtn_null_constant(struct vtn_builder *b, const struct glsl_type *type)
break; break;
case GLSL_TYPE_ARRAY: case GLSL_TYPE_ARRAY:
assert(glsl_get_length(type) > 0); vtn_assert(glsl_get_length(type) > 0);
c->num_elements = glsl_get_length(type); c->num_elements = glsl_get_length(type);
c->elements = ralloc_array(b, nir_constant *, c->num_elements); c->elements = ralloc_array(b, nir_constant *, c->num_elements);
@@ -1101,7 +1100,7 @@ spec_constant_decoration_cb(struct vtn_builder *b, struct vtn_value *v,
int member, const struct vtn_decoration *dec, int member, const struct vtn_decoration *dec,
void *data) void *data)
{ {
assert(member == -1); vtn_assert(member == -1);
if (dec->decoration != SpvDecorationSpecId) if (dec->decoration != SpvDecorationSpecId)
return; return;
@@ -1147,12 +1146,12 @@ handle_workgroup_size_decoration_cb(struct vtn_builder *b,
const struct vtn_decoration *dec, const struct vtn_decoration *dec,
void *data) void *data)
{ {
assert(member == -1); vtn_assert(member == -1);
if (dec->decoration != SpvDecorationBuiltIn || if (dec->decoration != SpvDecorationBuiltIn ||
dec->literals[0] != SpvBuiltInWorkgroupSize) dec->literals[0] != SpvBuiltInWorkgroupSize)
return; return;
assert(val->const_type == glsl_vector_type(GLSL_TYPE_UINT, 3)); vtn_assert(val->const_type == glsl_vector_type(GLSL_TYPE_UINT, 3));
b->shader->info.cs.local_size[0] = val->constant->values[0].u32[0]; b->shader->info.cs.local_size[0] = val->constant->values[0].u32[0];
b->shader->info.cs.local_size[1] = val->constant->values[0].u32[1]; b->shader->info.cs.local_size[1] = val->constant->values[0].u32[1];
@@ -1168,17 +1167,17 @@ vtn_handle_constant(struct vtn_builder *b, SpvOp opcode,
val->constant = rzalloc(b, nir_constant); val->constant = rzalloc(b, nir_constant);
switch (opcode) { switch (opcode) {
case SpvOpConstantTrue: case SpvOpConstantTrue:
assert(val->const_type == glsl_bool_type()); vtn_assert(val->const_type == glsl_bool_type());
val->constant->values[0].u32[0] = NIR_TRUE; val->constant->values[0].u32[0] = NIR_TRUE;
break; break;
case SpvOpConstantFalse: case SpvOpConstantFalse:
assert(val->const_type == glsl_bool_type()); vtn_assert(val->const_type == glsl_bool_type());
val->constant->values[0].u32[0] = NIR_FALSE; val->constant->values[0].u32[0] = NIR_FALSE;
break; break;
case SpvOpSpecConstantTrue: case SpvOpSpecConstantTrue:
case SpvOpSpecConstantFalse: { case SpvOpSpecConstantFalse: {
assert(val->const_type == glsl_bool_type()); vtn_assert(val->const_type == glsl_bool_type());
uint32_t int_val = uint32_t int_val =
get_specialization(b, val, (opcode == SpvOpSpecConstantTrue)); get_specialization(b, val, (opcode == SpvOpSpecConstantTrue));
val->constant->values[0].u32[0] = int_val ? NIR_TRUE : NIR_FALSE; val->constant->values[0].u32[0] = int_val ? NIR_TRUE : NIR_FALSE;
@@ -1186,19 +1185,19 @@ vtn_handle_constant(struct vtn_builder *b, SpvOp opcode,
} }
case SpvOpConstant: { case SpvOpConstant: {
assert(glsl_type_is_scalar(val->const_type)); vtn_assert(glsl_type_is_scalar(val->const_type));
int bit_size = glsl_get_bit_size(val->const_type); int bit_size = glsl_get_bit_size(val->const_type);
if (bit_size == 64) { if (bit_size == 64) {
val->constant->values->u32[0] = w[3]; val->constant->values->u32[0] = w[3];
val->constant->values->u32[1] = w[4]; val->constant->values->u32[1] = w[4];
} else { } else {
assert(bit_size == 32); vtn_assert(bit_size == 32);
val->constant->values->u32[0] = w[3]; val->constant->values->u32[0] = w[3];
} }
break; break;
} }
case SpvOpSpecConstant: { case SpvOpSpecConstant: {
assert(glsl_type_is_scalar(val->const_type)); vtn_assert(glsl_type_is_scalar(val->const_type));
val->constant->values[0].u32[0] = get_specialization(b, val, w[3]); val->constant->values[0].u32[0] = get_specialization(b, val, w[3]);
int bit_size = glsl_get_bit_size(val->const_type); int bit_size = glsl_get_bit_size(val->const_type);
if (bit_size == 64) if (bit_size == 64)
@@ -1225,17 +1224,17 @@ vtn_handle_constant(struct vtn_builder *b, SpvOp opcode,
case GLSL_TYPE_DOUBLE: { case GLSL_TYPE_DOUBLE: {
int bit_size = glsl_get_bit_size(val->const_type); int bit_size = glsl_get_bit_size(val->const_type);
if (glsl_type_is_matrix(val->const_type)) { if (glsl_type_is_matrix(val->const_type)) {
assert(glsl_get_matrix_columns(val->const_type) == elem_count); vtn_assert(glsl_get_matrix_columns(val->const_type) == elem_count);
for (unsigned i = 0; i < elem_count; i++) for (unsigned i = 0; i < elem_count; i++)
val->constant->values[i] = elems[i]->values[0]; val->constant->values[i] = elems[i]->values[0];
} else { } else {
assert(glsl_type_is_vector(val->const_type)); vtn_assert(glsl_type_is_vector(val->const_type));
assert(glsl_get_vector_elements(val->const_type) == elem_count); vtn_assert(glsl_get_vector_elements(val->const_type) == elem_count);
for (unsigned i = 0; i < elem_count; i++) { for (unsigned i = 0; i < elem_count; i++) {
if (bit_size == 64) { if (bit_size == 64) {
val->constant->values[0].u64[i] = elems[i]->values[0].u64[0]; val->constant->values[0].u64[i] = elems[i]->values[0].u64[0];
} else { } else {
assert(bit_size == 32); vtn_assert(bit_size == 32);
val->constant->values[0].u32[i] = elems[i]->values[0].u32[0]; val->constant->values[0].u32[i] = elems[i]->values[0].u32[0];
} }
} }
@@ -1263,9 +1262,9 @@ vtn_handle_constant(struct vtn_builder *b, SpvOp opcode,
struct vtn_value *v0 = &b->values[w[4]]; struct vtn_value *v0 = &b->values[w[4]];
struct vtn_value *v1 = &b->values[w[5]]; struct vtn_value *v1 = &b->values[w[5]];
assert(v0->value_type == vtn_value_type_constant || vtn_assert(v0->value_type == vtn_value_type_constant ||
v0->value_type == vtn_value_type_undef); v0->value_type == vtn_value_type_undef);
assert(v1->value_type == vtn_value_type_constant || vtn_assert(v1->value_type == vtn_value_type_constant ||
v1->value_type == vtn_value_type_undef); v1->value_type == vtn_value_type_undef);
unsigned len0 = v0->value_type == vtn_value_type_constant ? unsigned len0 = v0->value_type == vtn_value_type_constant ?
@@ -1275,7 +1274,7 @@ vtn_handle_constant(struct vtn_builder *b, SpvOp opcode,
glsl_get_vector_elements(v1->const_type) : glsl_get_vector_elements(v1->const_type) :
glsl_get_vector_elements(v1->type->type); glsl_get_vector_elements(v1->type->type);
assert(len0 + len1 < 16); vtn_assert(len0 + len1 < 16);
unsigned bit_size = glsl_get_bit_size(val->const_type); unsigned bit_size = glsl_get_bit_size(val->const_type);
unsigned bit_size0 = v0->value_type == vtn_value_type_constant ? unsigned bit_size0 = v0->value_type == vtn_value_type_constant ?
@@ -1285,7 +1284,7 @@ vtn_handle_constant(struct vtn_builder *b, SpvOp opcode,
glsl_get_bit_size(v1->const_type) : glsl_get_bit_size(v1->const_type) :
glsl_get_bit_size(v1->type->type); glsl_get_bit_size(v1->type->type);
assert(bit_size == bit_size0 && bit_size == bit_size1); vtn_assert(bit_size == bit_size0 && bit_size == bit_size1);
(void)bit_size0; (void)bit_size1; (void)bit_size0; (void)bit_size1;
if (bit_size == 64) { if (bit_size == 64) {
@@ -1365,12 +1364,12 @@ vtn_handle_constant(struct vtn_builder *b, SpvOp opcode,
case GLSL_TYPE_BOOL: case GLSL_TYPE_BOOL:
/* If we hit this granularity, we're picking off an element */ /* If we hit this granularity, we're picking off an element */
if (glsl_type_is_matrix(type)) { if (glsl_type_is_matrix(type)) {
assert(col == 0 && elem == -1); vtn_assert(col == 0 && elem == -1);
col = w[i]; col = w[i];
elem = 0; elem = 0;
type = glsl_get_column_type(type); type = glsl_get_column_type(type);
} else { } else {
assert(elem <= 0 && glsl_type_is_vector(type)); vtn_assert(elem <= 0 && glsl_type_is_vector(type));
elem = w[i]; elem = w[i];
type = glsl_scalar_type(glsl_get_base_type(type)); type = glsl_scalar_type(glsl_get_base_type(type));
} }
@@ -1401,14 +1400,14 @@ vtn_handle_constant(struct vtn_builder *b, SpvOp opcode,
if (bit_size == 64) { if (bit_size == 64) {
val->constant->values[0].u64[i] = (*c)->values[col].u64[elem + i]; val->constant->values[0].u64[i] = (*c)->values[col].u64[elem + i];
} else { } else {
assert(bit_size == 32); vtn_assert(bit_size == 32);
val->constant->values[0].u32[i] = (*c)->values[col].u32[elem + i]; val->constant->values[0].u32[i] = (*c)->values[col].u32[elem + i];
} }
} }
} else { } else {
struct vtn_value *insert = struct vtn_value *insert =
vtn_value(b, w[4], vtn_value_type_constant); vtn_value(b, w[4], vtn_value_type_constant);
assert(insert->const_type == type); vtn_assert(insert->const_type == type);
if (elem == -1) { if (elem == -1) {
*c = insert->constant; *c = insert->constant;
} else { } else {
@@ -1418,7 +1417,7 @@ vtn_handle_constant(struct vtn_builder *b, SpvOp opcode,
if (bit_size == 64) { if (bit_size == 64) {
(*c)->values[col].u64[elem + i] = insert->constant->values[0].u64[i]; (*c)->values[col].u64[elem + i] = insert->constant->values[0].u64[i];
} else { } else {
assert(bit_size == 32); vtn_assert(bit_size == 32);
(*c)->values[col].u32[elem + i] = insert->constant->values[0].u32[i]; (*c)->values[col].u32[elem + i] = insert->constant->values[0].u32[i];
} }
} }
@@ -1437,13 +1436,13 @@ vtn_handle_constant(struct vtn_builder *b, SpvOp opcode,
glsl_get_bit_size(val->const_type); glsl_get_bit_size(val->const_type);
nir_const_value src[4]; nir_const_value src[4];
assert(count <= 7); vtn_assert(count <= 7);
for (unsigned i = 0; i < count - 4; i++) { for (unsigned i = 0; i < count - 4; i++) {
nir_constant *c = nir_constant *c =
vtn_value(b, w[4 + i], vtn_value_type_constant)->constant; vtn_value(b, w[4 + i], vtn_value_type_constant)->constant;
unsigned j = swap ? 1 - i : i; unsigned j = swap ? 1 - i : i;
assert(bit_size == 32); vtn_assert(bit_size == 32);
src[j] = c->values[0]; src[j] = c->values[0];
} }
@@ -1503,7 +1502,7 @@ vtn_handle_function_call(struct vtn_builder *b, SpvOp opcode,
} }
nir_variable *out_tmp = NULL; nir_variable *out_tmp = NULL;
assert(res_type->type == callee->return_type); vtn_assert(res_type->type == callee->return_type);
if (!glsl_type_is_void(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 = nir_local_variable_create(b->nb.impl, callee->return_type,
"out_tmp"); "out_tmp");
@@ -1588,7 +1587,7 @@ vtn_handle_texture(struct vtn_builder *b, SpvOp opcode,
if (src_val->value_type == vtn_value_type_sampled_image) { if (src_val->value_type == vtn_value_type_sampled_image) {
val->pointer = src_val->sampled_image->image; val->pointer = src_val->sampled_image->image;
} else { } else {
assert(src_val->value_type == vtn_value_type_pointer); vtn_assert(src_val->value_type == vtn_value_type_pointer);
val->pointer = src_val->pointer; val->pointer = src_val->pointer;
} }
return; return;
@@ -1602,7 +1601,7 @@ vtn_handle_texture(struct vtn_builder *b, SpvOp opcode,
if (sampled_val->value_type == vtn_value_type_sampled_image) { if (sampled_val->value_type == vtn_value_type_sampled_image) {
sampled = *sampled_val->sampled_image; sampled = *sampled_val->sampled_image;
} else { } else {
assert(sampled_val->value_type == vtn_value_type_pointer); vtn_assert(sampled_val->value_type == vtn_value_type_pointer);
sampled.type = sampled_val->pointer->type; sampled.type = sampled_val->pointer->type;
sampled.image = NULL; sampled.image = NULL;
sampled.sampler = sampled_val->pointer; sampled.sampler = sampled_val->pointer;
@@ -1766,19 +1765,19 @@ vtn_handle_texture(struct vtn_builder *b, SpvOp opcode,
uint32_t operands = w[idx++]; uint32_t operands = w[idx++];
if (operands & SpvImageOperandsBiasMask) { if (operands & SpvImageOperandsBiasMask) {
assert(texop == nir_texop_tex); vtn_assert(texop == nir_texop_tex);
texop = nir_texop_txb; texop = nir_texop_txb;
(*p++) = vtn_tex_src(b, w[idx++], nir_tex_src_bias); (*p++) = vtn_tex_src(b, w[idx++], nir_tex_src_bias);
} }
if (operands & SpvImageOperandsLodMask) { if (operands & SpvImageOperandsLodMask) {
assert(texop == nir_texop_txl || texop == nir_texop_txf || vtn_assert(texop == nir_texop_txl || texop == nir_texop_txf ||
texop == nir_texop_txs); texop == nir_texop_txs);
(*p++) = vtn_tex_src(b, w[idx++], nir_tex_src_lod); (*p++) = vtn_tex_src(b, w[idx++], nir_tex_src_lod);
} }
if (operands & SpvImageOperandsGradMask) { if (operands & SpvImageOperandsGradMask) {
assert(texop == nir_texop_txl); vtn_assert(texop == nir_texop_txl);
texop = nir_texop_txd; texop = nir_texop_txd;
(*p++) = vtn_tex_src(b, w[idx++], nir_tex_src_ddx); (*p++) = vtn_tex_src(b, w[idx++], nir_tex_src_ddx);
(*p++) = vtn_tex_src(b, w[idx++], nir_tex_src_ddy); (*p++) = vtn_tex_src(b, w[idx++], nir_tex_src_ddy);
@@ -1794,13 +1793,13 @@ vtn_handle_texture(struct vtn_builder *b, SpvOp opcode,
} }
if (operands & SpvImageOperandsSampleMask) { if (operands & SpvImageOperandsSampleMask) {
assert(texop == nir_texop_txf_ms); vtn_assert(texop == nir_texop_txf_ms);
texop = nir_texop_txf_ms; texop = nir_texop_txf_ms;
(*p++) = vtn_tex_src(b, w[idx++], nir_tex_src_ms_index); (*p++) = vtn_tex_src(b, w[idx++], nir_tex_src_ms_index);
} }
} }
/* We should have now consumed exactly all of the arguments */ /* We should have now consumed exactly all of the arguments */
assert(idx == count); vtn_assert(idx == count);
nir_tex_instr *instr = nir_tex_instr_create(b->shader, p - srcs); nir_tex_instr *instr = nir_tex_instr_create(b->shader, p - srcs);
instr->op = texop; instr->op = texop;
@@ -1861,14 +1860,14 @@ vtn_handle_texture(struct vtn_builder *b, SpvOp opcode,
nir_ssa_dest_init(&instr->instr, &instr->dest, nir_ssa_dest_init(&instr->instr, &instr->dest,
nir_tex_instr_dest_size(instr), 32, NULL); nir_tex_instr_dest_size(instr), 32, NULL);
assert(glsl_get_vector_elements(ret_type->type) == vtn_assert(glsl_get_vector_elements(ret_type->type) ==
nir_tex_instr_dest_size(instr)); nir_tex_instr_dest_size(instr));
nir_ssa_def *def; nir_ssa_def *def;
nir_instr *instruction; nir_instr *instruction;
if (gather_offsets) { if (gather_offsets) {
assert(glsl_get_base_type(gather_offsets->type) == GLSL_TYPE_ARRAY); vtn_assert(glsl_get_base_type(gather_offsets->type) == GLSL_TYPE_ARRAY);
assert(glsl_get_length(gather_offsets->type) == 4); vtn_assert(glsl_get_length(gather_offsets->type) == 4);
nir_tex_instr *instrs[4] = {instr, NULL, NULL, NULL}; nir_tex_instr *instrs[4] = {instr, NULL, NULL, NULL};
/* Copy the current instruction 4x */ /* Copy the current instruction 4x */
@@ -2031,7 +2030,7 @@ vtn_handle_image(struct vtn_builder *b, SpvOp opcode,
image.coord = get_image_coord(b, w[4]); image.coord = get_image_coord(b, w[4]);
if (count > 5 && (w[5] & SpvImageOperandsSampleMask)) { if (count > 5 && (w[5] & SpvImageOperandsSampleMask)) {
assert(w[5] == SpvImageOperandsSampleMask); vtn_assert(w[5] == SpvImageOperandsSampleMask);
image.sample = vtn_ssa_value(b, w[6])->def; image.sample = vtn_ssa_value(b, w[6])->def;
} else { } else {
image.sample = nir_ssa_undef(&b->nb, 1, 32); image.sample = nir_ssa_undef(&b->nb, 1, 32);
@@ -2045,7 +2044,7 @@ vtn_handle_image(struct vtn_builder *b, SpvOp opcode,
/* texel = w[3] */ /* texel = w[3] */
if (count > 4 && (w[4] & SpvImageOperandsSampleMask)) { if (count > 4 && (w[4] & SpvImageOperandsSampleMask)) {
assert(w[4] == SpvImageOperandsSampleMask); vtn_assert(w[4] == SpvImageOperandsSampleMask);
image.sample = vtn_ssa_value(b, w[5])->def; image.sample = vtn_ssa_value(b, w[5])->def;
} else { } else {
image.sample = nir_ssa_undef(&b->nb, 1, 32); image.sample = nir_ssa_undef(&b->nb, 1, 32);
@@ -2286,7 +2285,7 @@ vtn_handle_ssbo_or_shared_atomic(struct vtn_builder *b, SpvOp opcode,
} }
} else { } else {
assert(ptr->mode == vtn_variable_mode_ssbo); vtn_assert(ptr->mode == vtn_variable_mode_ssbo);
nir_ssa_def *offset, *index; nir_ssa_def *offset, *index;
offset = vtn_pointer_to_offset(b, ptr, &index, NULL); offset = vtn_pointer_to_offset(b, ptr, &index, NULL);
@@ -2493,12 +2492,12 @@ vtn_vector_construct(struct vtn_builder *b, unsigned num_components,
* "When constructing a vector, there must be at least two Constituent * "When constructing a vector, there must be at least two Constituent
* operands." * operands."
*/ */
assert(num_srcs >= 2); vtn_assert(num_srcs >= 2);
unsigned dest_idx = 0; unsigned dest_idx = 0;
for (unsigned i = 0; i < num_srcs; i++) { for (unsigned i = 0; i < num_srcs; i++) {
nir_ssa_def *src = srcs[i]; nir_ssa_def *src = srcs[i];
assert(dest_idx + src->num_components <= num_components); vtn_assert(dest_idx + src->num_components <= num_components);
for (unsigned j = 0; j < src->num_components; j++) { for (unsigned j = 0; j < src->num_components; j++) {
vec->src[dest_idx].src = nir_src_for_ssa(src); vec->src[dest_idx].src = nir_src_for_ssa(src);
vec->src[dest_idx].swizzle[0] = j; vec->src[dest_idx].swizzle[0] = j;
@@ -2511,7 +2510,7 @@ vtn_vector_construct(struct vtn_builder *b, unsigned num_components,
* "When constructing a vector, the total number of components in all * "When constructing a vector, the total number of components in all
* the operands must equal the number of components in Result Type." * the operands must equal the number of components in Result Type."
*/ */
assert(dest_idx == num_components); vtn_assert(dest_idx == num_components);
nir_builder_instr_insert(&b->nb, &vec->instr); nir_builder_instr_insert(&b->nb, &vec->instr);
@@ -2571,7 +2570,7 @@ vtn_composite_extract(struct vtn_builder *b, struct vtn_ssa_value *src,
struct vtn_ssa_value *cur = src; struct vtn_ssa_value *cur = src;
for (unsigned i = 0; i < num_indices; i++) { for (unsigned i = 0; i < num_indices; i++) {
if (glsl_type_is_vector_or_scalar(cur->type)) { if (glsl_type_is_vector_or_scalar(cur->type)) {
assert(i == num_indices - 1); vtn_assert(i == num_indices - 1);
/* According to the SPIR-V spec, OpCompositeExtract may work down to /* According to the SPIR-V spec, OpCompositeExtract may work down to
* the component granularity. The last index will be the index of the * the component granularity. The last index will be the index of the
* vector to extract. * vector to extract.
@@ -2904,8 +2903,8 @@ vtn_handle_preamble_instruction(struct vtn_builder *b, SpvOp opcode,
break; break;
case SpvOpMemoryModel: case SpvOpMemoryModel:
assert(w[1] == SpvAddressingModelLogical); vtn_assert(w[1] == SpvAddressingModelLogical);
assert(w[2] == SpvMemoryModelSimple || vtn_assert(w[2] == SpvMemoryModelSimple ||
w[2] == SpvMemoryModelGLSL450); w[2] == SpvMemoryModelGLSL450);
break; break;
@@ -2919,7 +2918,7 @@ vtn_handle_preamble_instruction(struct vtn_builder *b, SpvOp opcode,
stage_for_execution_model(w[1]) != b->entry_point_stage) stage_for_execution_model(w[1]) != b->entry_point_stage)
break; break;
assert(b->entry_point == NULL); vtn_assert(b->entry_point == NULL);
b->entry_point = entry_point; b->entry_point = entry_point;
break; break;
} }
@@ -2957,7 +2956,7 @@ static void
vtn_handle_execution_mode(struct vtn_builder *b, struct vtn_value *entry_point, vtn_handle_execution_mode(struct vtn_builder *b, struct vtn_value *entry_point,
const struct vtn_decoration *mode, void *data) const struct vtn_decoration *mode, void *data)
{ {
assert(b->entry_point == entry_point); vtn_assert(b->entry_point == entry_point);
switch(mode->exec_mode) { switch(mode->exec_mode) {
case SpvExecutionModeOriginUpperLeft: case SpvExecutionModeOriginUpperLeft:
@@ -2967,34 +2966,34 @@ vtn_handle_execution_mode(struct vtn_builder *b, struct vtn_value *entry_point,
break; break;
case SpvExecutionModeEarlyFragmentTests: case SpvExecutionModeEarlyFragmentTests:
assert(b->shader->info.stage == MESA_SHADER_FRAGMENT); vtn_assert(b->shader->info.stage == MESA_SHADER_FRAGMENT);
b->shader->info.fs.early_fragment_tests = true; b->shader->info.fs.early_fragment_tests = true;
break; break;
case SpvExecutionModeInvocations: case SpvExecutionModeInvocations:
assert(b->shader->info.stage == MESA_SHADER_GEOMETRY); vtn_assert(b->shader->info.stage == MESA_SHADER_GEOMETRY);
b->shader->info.gs.invocations = MAX2(1, mode->literals[0]); b->shader->info.gs.invocations = MAX2(1, mode->literals[0]);
break; break;
case SpvExecutionModeDepthReplacing: case SpvExecutionModeDepthReplacing:
assert(b->shader->info.stage == MESA_SHADER_FRAGMENT); vtn_assert(b->shader->info.stage == MESA_SHADER_FRAGMENT);
b->shader->info.fs.depth_layout = FRAG_DEPTH_LAYOUT_ANY; b->shader->info.fs.depth_layout = FRAG_DEPTH_LAYOUT_ANY;
break; break;
case SpvExecutionModeDepthGreater: case SpvExecutionModeDepthGreater:
assert(b->shader->info.stage == MESA_SHADER_FRAGMENT); vtn_assert(b->shader->info.stage == MESA_SHADER_FRAGMENT);
b->shader->info.fs.depth_layout = FRAG_DEPTH_LAYOUT_GREATER; b->shader->info.fs.depth_layout = FRAG_DEPTH_LAYOUT_GREATER;
break; break;
case SpvExecutionModeDepthLess: case SpvExecutionModeDepthLess:
assert(b->shader->info.stage == MESA_SHADER_FRAGMENT); vtn_assert(b->shader->info.stage == MESA_SHADER_FRAGMENT);
b->shader->info.fs.depth_layout = FRAG_DEPTH_LAYOUT_LESS; b->shader->info.fs.depth_layout = FRAG_DEPTH_LAYOUT_LESS;
break; break;
case SpvExecutionModeDepthUnchanged: case SpvExecutionModeDepthUnchanged:
assert(b->shader->info.stage == MESA_SHADER_FRAGMENT); vtn_assert(b->shader->info.stage == MESA_SHADER_FRAGMENT);
b->shader->info.fs.depth_layout = FRAG_DEPTH_LAYOUT_UNCHANGED; b->shader->info.fs.depth_layout = FRAG_DEPTH_LAYOUT_UNCHANGED;
break; break;
case SpvExecutionModeLocalSize: case SpvExecutionModeLocalSize:
assert(b->shader->info.stage == MESA_SHADER_COMPUTE); vtn_assert(b->shader->info.stage == MESA_SHADER_COMPUTE);
b->shader->info.cs.local_size[0] = mode->literals[0]; b->shader->info.cs.local_size[0] = mode->literals[0];
b->shader->info.cs.local_size[1] = mode->literals[1]; b->shader->info.cs.local_size[1] = mode->literals[1];
b->shader->info.cs.local_size[2] = mode->literals[2]; b->shader->info.cs.local_size[2] = mode->literals[2];
@@ -3007,7 +3006,7 @@ vtn_handle_execution_mode(struct vtn_builder *b, struct vtn_value *entry_point,
b->shader->info.stage == MESA_SHADER_TESS_EVAL) { b->shader->info.stage == MESA_SHADER_TESS_EVAL) {
b->shader->info.tess.tcs_vertices_out = mode->literals[0]; b->shader->info.tess.tcs_vertices_out = mode->literals[0];
} else { } else {
assert(b->shader->info.stage == MESA_SHADER_GEOMETRY); vtn_assert(b->shader->info.stage == MESA_SHADER_GEOMETRY);
b->shader->info.gs.vertices_out = mode->literals[0]; b->shader->info.gs.vertices_out = mode->literals[0];
} }
break; break;
@@ -3024,7 +3023,7 @@ vtn_handle_execution_mode(struct vtn_builder *b, struct vtn_value *entry_point,
b->shader->info.tess.primitive_mode = b->shader->info.tess.primitive_mode =
gl_primitive_from_spv_execution_mode(mode->exec_mode); gl_primitive_from_spv_execution_mode(mode->exec_mode);
} else { } else {
assert(b->shader->info.stage == MESA_SHADER_GEOMETRY); vtn_assert(b->shader->info.stage == MESA_SHADER_GEOMETRY);
b->shader->info.gs.vertices_in = b->shader->info.gs.vertices_in =
vertices_in_from_spv_execution_mode(mode->exec_mode); vertices_in_from_spv_execution_mode(mode->exec_mode);
} }
@@ -3033,38 +3032,38 @@ vtn_handle_execution_mode(struct vtn_builder *b, struct vtn_value *entry_point,
case SpvExecutionModeOutputPoints: case SpvExecutionModeOutputPoints:
case SpvExecutionModeOutputLineStrip: case SpvExecutionModeOutputLineStrip:
case SpvExecutionModeOutputTriangleStrip: case SpvExecutionModeOutputTriangleStrip:
assert(b->shader->info.stage == MESA_SHADER_GEOMETRY); vtn_assert(b->shader->info.stage == MESA_SHADER_GEOMETRY);
b->shader->info.gs.output_primitive = b->shader->info.gs.output_primitive =
gl_primitive_from_spv_execution_mode(mode->exec_mode); gl_primitive_from_spv_execution_mode(mode->exec_mode);
break; break;
case SpvExecutionModeSpacingEqual: case SpvExecutionModeSpacingEqual:
assert(b->shader->info.stage == MESA_SHADER_TESS_CTRL || vtn_assert(b->shader->info.stage == MESA_SHADER_TESS_CTRL ||
b->shader->info.stage == MESA_SHADER_TESS_EVAL); b->shader->info.stage == MESA_SHADER_TESS_EVAL);
b->shader->info.tess.spacing = TESS_SPACING_EQUAL; b->shader->info.tess.spacing = TESS_SPACING_EQUAL;
break; break;
case SpvExecutionModeSpacingFractionalEven: case SpvExecutionModeSpacingFractionalEven:
assert(b->shader->info.stage == MESA_SHADER_TESS_CTRL || vtn_assert(b->shader->info.stage == MESA_SHADER_TESS_CTRL ||
b->shader->info.stage == MESA_SHADER_TESS_EVAL); b->shader->info.stage == MESA_SHADER_TESS_EVAL);
b->shader->info.tess.spacing = TESS_SPACING_FRACTIONAL_EVEN; b->shader->info.tess.spacing = TESS_SPACING_FRACTIONAL_EVEN;
break; break;
case SpvExecutionModeSpacingFractionalOdd: case SpvExecutionModeSpacingFractionalOdd:
assert(b->shader->info.stage == MESA_SHADER_TESS_CTRL || vtn_assert(b->shader->info.stage == MESA_SHADER_TESS_CTRL ||
b->shader->info.stage == MESA_SHADER_TESS_EVAL); b->shader->info.stage == MESA_SHADER_TESS_EVAL);
b->shader->info.tess.spacing = TESS_SPACING_FRACTIONAL_ODD; b->shader->info.tess.spacing = TESS_SPACING_FRACTIONAL_ODD;
break; break;
case SpvExecutionModeVertexOrderCw: case SpvExecutionModeVertexOrderCw:
assert(b->shader->info.stage == MESA_SHADER_TESS_CTRL || vtn_assert(b->shader->info.stage == MESA_SHADER_TESS_CTRL ||
b->shader->info.stage == MESA_SHADER_TESS_EVAL); b->shader->info.stage == MESA_SHADER_TESS_EVAL);
b->shader->info.tess.ccw = false; b->shader->info.tess.ccw = false;
break; break;
case SpvExecutionModeVertexOrderCcw: case SpvExecutionModeVertexOrderCcw:
assert(b->shader->info.stage == MESA_SHADER_TESS_CTRL || vtn_assert(b->shader->info.stage == MESA_SHADER_TESS_CTRL ||
b->shader->info.stage == MESA_SHADER_TESS_EVAL); b->shader->info.stage == MESA_SHADER_TESS_EVAL);
b->shader->info.tess.ccw = true; b->shader->info.tess.ccw = true;
break; break;
case SpvExecutionModePointMode: case SpvExecutionModePointMode:
assert(b->shader->info.stage == MESA_SHADER_TESS_CTRL || vtn_assert(b->shader->info.stage == MESA_SHADER_TESS_CTRL ||
b->shader->info.stage == MESA_SHADER_TESS_EVAL); b->shader->info.stage == MESA_SHADER_TESS_EVAL);
b->shader->info.tess.point_mode = true; b->shader->info.tess.point_mode = true;
break; break;
@@ -3231,7 +3230,7 @@ vtn_handle_body_instruction(struct vtn_builder *b, SpvOp opcode,
if (image->mode == vtn_variable_mode_image) { if (image->mode == vtn_variable_mode_image) {
vtn_handle_image(b, opcode, w, count); vtn_handle_image(b, opcode, w, count);
} else { } else {
assert(image->mode == vtn_variable_mode_sampler); vtn_assert(image->mode == vtn_variable_mode_sampler);
vtn_handle_texture(b, opcode, w, count); vtn_handle_texture(b, opcode, w, count);
} }
break; break;
@@ -3256,7 +3255,7 @@ vtn_handle_body_instruction(struct vtn_builder *b, SpvOp opcode,
if (pointer->value_type == vtn_value_type_image_pointer) { if (pointer->value_type == vtn_value_type_image_pointer) {
vtn_handle_image(b, opcode, w, count); vtn_handle_image(b, opcode, w, count);
} else { } else {
assert(pointer->value_type == vtn_value_type_pointer); vtn_assert(pointer->value_type == vtn_value_type_pointer);
vtn_handle_ssbo_or_shared_atomic(b, opcode, w, count); vtn_handle_ssbo_or_shared_atomic(b, opcode, w, count);
} }
break; break;
@@ -3267,7 +3266,7 @@ vtn_handle_body_instruction(struct vtn_builder *b, SpvOp opcode,
if (pointer->value_type == vtn_value_type_image_pointer) { if (pointer->value_type == vtn_value_type_image_pointer) {
vtn_handle_image(b, opcode, w, count); vtn_handle_image(b, opcode, w, count);
} else { } else {
assert(pointer->value_type == vtn_value_type_pointer); vtn_assert(pointer->value_type == vtn_value_type_pointer);
vtn_handle_ssbo_or_shared_atomic(b, opcode, w, count); vtn_handle_ssbo_or_shared_atomic(b, opcode, w, count);
} }
break; break;
@@ -3441,13 +3440,13 @@ spirv_to_nir(const uint32_t *words, size_t word_count,
const uint32_t *word_end = words + word_count; const uint32_t *word_end = words + word_count;
/* Handle the SPIR-V header (first 4 dwords) */ /* Handle the SPIR-V header (first 4 dwords) */
assert(word_count > 5); vtn_assert(word_count > 5);
assert(words[0] == SpvMagicNumber); vtn_assert(words[0] == SpvMagicNumber);
assert(words[1] >= 0x10000); vtn_assert(words[1] >= 0x10000);
/* words[2] == generator magic */ /* words[2] == generator magic */
unsigned value_id_bound = words[3]; unsigned value_id_bound = words[3];
assert(words[4] == 0); vtn_assert(words[4] == 0);
words+= 5; words+= 5;
@@ -3499,9 +3498,9 @@ spirv_to_nir(const uint32_t *words, size_t word_count,
} }
} while (progress); } while (progress);
assert(b->entry_point->value_type == vtn_value_type_function); vtn_assert(b->entry_point->value_type == vtn_value_type_function);
nir_function *entry_point = b->entry_point->func->impl->function; nir_function *entry_point = b->entry_point->func->impl->function;
assert(entry_point); vtn_assert(entry_point);
/* Unparent the shader from the vtn_builder before we delete the builder */ /* Unparent the shader from the vtn_builder before we delete the builder */
ralloc_steal(NULL, b->shader); ralloc_steal(NULL, b->shader);

26
src/compiler/spirv/vtn_alu.c
View File

@@ -243,29 +243,29 @@ vtn_handle_bitcast(struct vtn_builder *b, struct vtn_ssa_value *dest,
unsigned dest_bit_size = glsl_get_bit_size(dest->type); unsigned dest_bit_size = glsl_get_bit_size(dest->type);
unsigned src_components = src->num_components; unsigned src_components = src->num_components;
unsigned dest_components = glsl_get_vector_elements(dest->type); unsigned dest_components = glsl_get_vector_elements(dest->type);
assert(src_bit_size * src_components == dest_bit_size * dest_components); vtn_assert(src_bit_size * src_components == dest_bit_size * dest_components);
nir_ssa_def *dest_chan[4]; nir_ssa_def *dest_chan[4];
if (src_bit_size > dest_bit_size) { if (src_bit_size > dest_bit_size) {
assert(src_bit_size % dest_bit_size == 0); vtn_assert(src_bit_size % dest_bit_size == 0);
unsigned divisor = src_bit_size / dest_bit_size; unsigned divisor = src_bit_size / dest_bit_size;
for (unsigned comp = 0; comp < src_components; comp++) { for (unsigned comp = 0; comp < src_components; comp++) {
assert(src_bit_size == 64); vtn_assert(src_bit_size == 64);
assert(dest_bit_size == 32); vtn_assert(dest_bit_size == 32);
nir_ssa_def *split = nir_ssa_def *split =
nir_unpack_64_2x32(&b->nb, nir_channel(&b->nb, src, comp)); nir_unpack_64_2x32(&b->nb, nir_channel(&b->nb, src, comp));
for (unsigned i = 0; i < divisor; i++) for (unsigned i = 0; i < divisor; i++)
dest_chan[divisor * comp + i] = nir_channel(&b->nb, split, i); dest_chan[divisor * comp + i] = nir_channel(&b->nb, split, i);
} }
} else { } else {
assert(dest_bit_size % src_bit_size == 0); vtn_assert(dest_bit_size % src_bit_size == 0);
unsigned divisor = dest_bit_size / src_bit_size; unsigned divisor = dest_bit_size / src_bit_size;
for (unsigned comp = 0; comp < dest_components; comp++) { for (unsigned comp = 0; comp < dest_components; comp++) {
unsigned channels = ((1 << divisor) - 1) << (comp * divisor); unsigned channels = ((1 << divisor) - 1) << (comp * divisor);
nir_ssa_def *src_chan = nir_ssa_def *src_chan =
nir_channels(&b->nb, src, channels); nir_channels(&b->nb, src, channels);
assert(dest_bit_size == 64); vtn_assert(dest_bit_size == 64);
assert(src_bit_size == 32); vtn_assert(src_bit_size == 32);
dest_chan[comp] = nir_pack_64_2x32(&b->nb, src_chan); dest_chan[comp] = nir_pack_64_2x32(&b->nb, src_chan);
} }
} }
@@ -374,7 +374,7 @@ static void
handle_no_contraction(struct vtn_builder *b, struct vtn_value *val, int member, handle_no_contraction(struct vtn_builder *b, struct vtn_value *val, int member,
const struct vtn_decoration *dec, void *_void) const struct vtn_decoration *dec, void *_void)
{ {
assert(dec->scope == VTN_DEC_DECORATION); vtn_assert(dec->scope == VTN_DEC_DECORATION);
if (dec->decoration != SpvDecorationNoContraction) if (dec->decoration != SpvDecorationNoContraction)
return; return;
@@ -407,7 +407,7 @@ vtn_handle_alu(struct vtn_builder *b, SpvOp opcode,
val->ssa = vtn_create_ssa_value(b, type); val->ssa = vtn_create_ssa_value(b, type);
nir_ssa_def *src[4] = { NULL, }; nir_ssa_def *src[4] = { NULL, };
for (unsigned i = 0; i < num_inputs; i++) { for (unsigned i = 0; i < num_inputs; i++) {
assert(glsl_type_is_vector_or_scalar(vtn_src[i]->type)); vtn_assert(glsl_type_is_vector_or_scalar(vtn_src[i]->type));
src[i] = vtn_src[i]->def; src[i] = vtn_src[i]->def;
} }
@@ -459,25 +459,25 @@ vtn_handle_alu(struct vtn_builder *b, SpvOp opcode,
break; break;
case SpvOpIAddCarry: case SpvOpIAddCarry:
assert(glsl_type_is_struct(val->ssa->type)); vtn_assert(glsl_type_is_struct(val->ssa->type));
val->ssa->elems[0]->def = nir_iadd(&b->nb, src[0], src[1]); val->ssa->elems[0]->def = nir_iadd(&b->nb, src[0], src[1]);
val->ssa->elems[1]->def = nir_uadd_carry(&b->nb, src[0], src[1]); val->ssa->elems[1]->def = nir_uadd_carry(&b->nb, src[0], src[1]);
break; break;
case SpvOpISubBorrow: case SpvOpISubBorrow:
assert(glsl_type_is_struct(val->ssa->type)); vtn_assert(glsl_type_is_struct(val->ssa->type));
val->ssa->elems[0]->def = nir_isub(&b->nb, src[0], src[1]); val->ssa->elems[0]->def = nir_isub(&b->nb, src[0], src[1]);
val->ssa->elems[1]->def = nir_usub_borrow(&b->nb, src[0], src[1]); val->ssa->elems[1]->def = nir_usub_borrow(&b->nb, src[0], src[1]);
break; break;
case SpvOpUMulExtended: case SpvOpUMulExtended:
assert(glsl_type_is_struct(val->ssa->type)); vtn_assert(glsl_type_is_struct(val->ssa->type));
val->ssa->elems[0]->def = nir_imul(&b->nb, src[0], src[1]); val->ssa->elems[0]->def = nir_imul(&b->nb, src[0], src[1]);
val->ssa->elems[1]->def = nir_umul_high(&b->nb, src[0], src[1]); val->ssa->elems[1]->def = nir_umul_high(&b->nb, src[0], src[1]);
break; break;
case SpvOpSMulExtended: case SpvOpSMulExtended:
assert(glsl_type_is_struct(val->ssa->type)); vtn_assert(glsl_type_is_struct(val->ssa->type));
val->ssa->elems[0]->def = nir_imul(&b->nb, src[0], src[1]); val->ssa->elems[0]->def = nir_imul(&b->nb, src[0], src[1]);
val->ssa->elems[1]->def = nir_imul_high(&b->nb, src[0], src[1]); val->ssa->elems[1]->def = nir_imul_high(&b->nb, src[0], src[1]);
break; break;

49
src/compiler/spirv/vtn_cfg.c
View File

@@ -30,7 +30,7 @@ vtn_cfg_handle_prepass_instruction(struct vtn_builder *b, SpvOp opcode,
{ {
switch (opcode) { switch (opcode) {
case SpvOpFunction: { case SpvOpFunction: {
assert(b->func == NULL); vtn_assert(b->func == NULL);
b->func = rzalloc(b, struct vtn_function); b->func = rzalloc(b, struct vtn_function);
list_inithead(&b->func->body); list_inithead(&b->func->body);
@@ -44,7 +44,7 @@ vtn_cfg_handle_prepass_instruction(struct vtn_builder *b, SpvOp opcode,
const struct vtn_type *func_type = const struct vtn_type *func_type =
vtn_value(b, w[4], vtn_value_type_type)->type; vtn_value(b, w[4], vtn_value_type_type)->type;
assert(func_type->return_type->type == result_type); vtn_assert(func_type->return_type->type == result_type);
nir_function *func = nir_function *func =
nir_function_create(b->shader, ralloc_strdup(b->shader, val->name)); nir_function_create(b->shader, ralloc_strdup(b->shader, val->name));
@@ -80,7 +80,7 @@ vtn_cfg_handle_prepass_instruction(struct vtn_builder *b, SpvOp opcode,
case SpvOpFunctionParameter: { case SpvOpFunctionParameter: {
struct vtn_type *type = vtn_value(b, w[1], vtn_value_type_type)->type; struct vtn_type *type = vtn_value(b, w[1], vtn_value_type_type)->type;
assert(b->func_param_idx < b->func->impl->num_params); vtn_assert(b->func_param_idx < b->func->impl->num_params);
nir_variable *param = b->func->impl->params[b->func_param_idx++]; nir_variable *param = b->func->impl->params[b->func_param_idx++];
if (type->base_type == vtn_base_type_pointer && type->type == NULL) { if (type->base_type == vtn_base_type_pointer && type->type == NULL) {
@@ -88,7 +88,7 @@ vtn_cfg_handle_prepass_instruction(struct vtn_builder *b, SpvOp opcode,
vtn_var->type = type->deref; vtn_var->type = type->deref;
vtn_var->var = param; vtn_var->var = param;
assert(vtn_var->type->type == param->type); vtn_assert(vtn_var->type->type == param->type);
struct vtn_type *without_array = vtn_var->type; struct vtn_type *without_array = vtn_var->type;
while(glsl_type_is_array(without_array->type)) while(glsl_type_is_array(without_array->type))
@@ -124,7 +124,7 @@ vtn_cfg_handle_prepass_instruction(struct vtn_builder *b, SpvOp opcode,
} }
case SpvOpLabel: { case SpvOpLabel: {
assert(b->block == NULL); vtn_assert(b->block == NULL);
b->block = rzalloc(b, struct vtn_block); b->block = rzalloc(b, struct vtn_block);
b->block->node.type = vtn_cf_node_type_block; b->block->node.type = vtn_cf_node_type_block;
b->block->label = w; b->block->label = w;
@@ -143,7 +143,7 @@ vtn_cfg_handle_prepass_instruction(struct vtn_builder *b, SpvOp opcode,
case SpvOpSelectionMerge: case SpvOpSelectionMerge:
case SpvOpLoopMerge: case SpvOpLoopMerge:
assert(b->block && b->block->merge == NULL); vtn_assert(b->block && b->block->merge == NULL);
b->block->merge = w; b->block->merge = w;
break; break;
@@ -154,7 +154,7 @@ vtn_cfg_handle_prepass_instruction(struct vtn_builder *b, SpvOp opcode,
case SpvOpReturn: case SpvOpReturn:
case SpvOpReturnValue: case SpvOpReturnValue:
case SpvOpUnreachable: case SpvOpUnreachable:
assert(b->block && b->block->branch == NULL); vtn_assert(b->block && b->block->branch == NULL);
b->block->branch = w; b->block->branch = w;
b->block = NULL; b->block = NULL;
break; break;
@@ -231,13 +231,14 @@ vtn_order_case(struct vtn_switch *swtch, struct vtn_case *cse)
} }
static enum vtn_branch_type static enum vtn_branch_type
vtn_get_branch_type(struct vtn_block *block, vtn_get_branch_type(struct vtn_builder *b,
struct vtn_block *block,
struct vtn_case *swcase, struct vtn_block *switch_break, struct vtn_case *swcase, struct vtn_block *switch_break,
struct vtn_block *loop_break, struct vtn_block *loop_cont) struct vtn_block *loop_break, struct vtn_block *loop_cont)
{ {
if (block->switch_case) { if (block->switch_case) {
/* This branch is actually a fallthrough */ /* This branch is actually a fallthrough */
assert(swcase->fallthrough == NULL || vtn_assert(swcase->fallthrough == NULL ||
swcase->fallthrough == block->switch_case); swcase->fallthrough == block->switch_case);
swcase->fallthrough = block->switch_case; swcase->fallthrough = block->switch_case;
return vtn_branch_type_switch_fallthrough; return vtn_branch_type_switch_fallthrough;
@@ -301,7 +302,7 @@ vtn_cfg_walk_blocks(struct vtn_builder *b, struct list_head *cf_list,
continue; continue;
} }
assert(block->node.link.next == NULL); vtn_assert(block->node.link.next == NULL);
list_addtail(&block->node.link, cf_list); list_addtail(&block->node.link, cf_list);
switch (*block->branch & SpvOpCodeMask) { switch (*block->branch & SpvOpCodeMask) {
@@ -309,7 +310,7 @@ vtn_cfg_walk_blocks(struct vtn_builder *b, struct list_head *cf_list,
struct vtn_block *branch_block = struct vtn_block *branch_block =
vtn_value(b, block->branch[1], vtn_value_type_block)->block; vtn_value(b, block->branch[1], vtn_value_type_block)->block;
block->branch_type = vtn_get_branch_type(branch_block, block->branch_type = vtn_get_branch_type(b, branch_block,
switch_case, switch_break, switch_case, switch_break,
loop_break, loop_cont); loop_break, loop_cont);
@@ -349,10 +350,10 @@ vtn_cfg_walk_blocks(struct vtn_builder *b, struct list_head *cf_list,
if_stmt->control = block->merge[2]; if_stmt->control = block->merge[2];
} }
if_stmt->then_type = vtn_get_branch_type(then_block, if_stmt->then_type = vtn_get_branch_type(b, then_block,
switch_case, switch_break, switch_case, switch_break,
loop_break, loop_cont); loop_break, loop_cont);
if_stmt->else_type = vtn_get_branch_type(else_block, if_stmt->else_type = vtn_get_branch_type(b, else_block,
switch_case, switch_break, switch_case, switch_break,
loop_break, loop_cont); loop_break, loop_cont);
@@ -367,7 +368,7 @@ vtn_cfg_walk_blocks(struct vtn_builder *b, struct list_head *cf_list,
} else if (if_stmt->then_type == vtn_branch_type_none && } else if (if_stmt->then_type == vtn_branch_type_none &&
if_stmt->else_type == vtn_branch_type_none) { if_stmt->else_type == vtn_branch_type_none) {
/* Neither side of the if is something we can short-circuit. */ /* Neither side of the if is something we can short-circuit. */
assert((*block->merge & SpvOpCodeMask) == SpvOpSelectionMerge); vtn_assert((*block->merge & SpvOpCodeMask) == SpvOpSelectionMerge);
struct vtn_block *merge_block = struct vtn_block *merge_block =
vtn_value(b, block->merge[1], vtn_value_type_block)->block; vtn_value(b, block->merge[1], vtn_value_type_block)->block;
@@ -379,7 +380,7 @@ vtn_cfg_walk_blocks(struct vtn_builder *b, struct list_head *cf_list,
loop_break, loop_cont, merge_block); loop_break, loop_cont, merge_block);
enum vtn_branch_type merge_type = enum vtn_branch_type merge_type =
vtn_get_branch_type(merge_block, switch_case, switch_break, vtn_get_branch_type(b, merge_block, switch_case, switch_break,
loop_break, loop_cont); loop_break, loop_cont);
if (merge_type == vtn_branch_type_none) { if (merge_type == vtn_branch_type_none) {
block = merge_block; block = merge_block;
@@ -408,7 +409,7 @@ vtn_cfg_walk_blocks(struct vtn_builder *b, struct list_head *cf_list,
} }
case SpvOpSwitch: { case SpvOpSwitch: {
assert((*block->merge & SpvOpCodeMask) == SpvOpSelectionMerge); vtn_assert((*block->merge & SpvOpCodeMask) == SpvOpSelectionMerge);
struct vtn_block *break_block = struct vtn_block *break_block =
vtn_value(b, block->merge[1], vtn_value_type_block)->block; vtn_value(b, block->merge[1], vtn_value_type_block)->block;
@@ -433,7 +434,7 @@ vtn_cfg_walk_blocks(struct vtn_builder *b, struct list_head *cf_list,
* information. * information.
*/ */
list_for_each_entry(struct vtn_case, cse, &swtch->cases, link) { list_for_each_entry(struct vtn_case, cse, &swtch->cases, link) {
assert(cse->start_block != break_block); vtn_assert(cse->start_block != break_block);
vtn_cfg_walk_blocks(b, &cse->body, cse->start_block, cse, vtn_cfg_walk_blocks(b, &cse->body, cse->start_block, cse,
break_block, loop_break, loop_cont, NULL); break_block, loop_break, loop_cont, NULL);
} }
@@ -448,13 +449,13 @@ vtn_cfg_walk_blocks(struct vtn_builder *b, struct list_head *cf_list,
if (case_block == break_block) if (case_block == break_block)
continue; continue;
assert(case_block->switch_case); vtn_assert(case_block->switch_case);
vtn_order_case(swtch, case_block->switch_case); vtn_order_case(swtch, case_block->switch_case);
} }
enum vtn_branch_type branch_type = enum vtn_branch_type branch_type =
vtn_get_branch_type(break_block, switch_case, NULL, vtn_get_branch_type(b, break_block, switch_case, NULL,
loop_break, loop_cont); loop_break, loop_cont);
if (branch_type != vtn_branch_type_none) { if (branch_type != vtn_branch_type_none) {
@@ -462,7 +463,7 @@ vtn_cfg_walk_blocks(struct vtn_builder *b, struct list_head *cf_list,
* for the containing loop. In this case, we need to bail and let * for the containing loop. In this case, we need to bail and let
* the loop parsing code handle the continue properly. * the loop parsing code handle the continue properly.
*/ */
assert(branch_type == vtn_branch_type_loop_continue); vtn_assert(branch_type == vtn_branch_type_loop_continue);
return; return;
} }
@@ -532,7 +533,7 @@ vtn_handle_phi_second_pass(struct vtn_builder *b, SpvOp opcode,
return true; return true;
struct hash_entry *phi_entry = _mesa_hash_table_search(b->phi_table, w); struct hash_entry *phi_entry = _mesa_hash_table_search(b->phi_table, w);
assert(phi_entry); vtn_assert(phi_entry);
nir_variable *phi_var = phi_entry->data; nir_variable *phi_var = phi_entry->data;
for (unsigned i = 3; i < count; i += 2) { for (unsigned i = 3; i < count; i += 2) {
@@ -728,7 +729,7 @@ vtn_emit_cf_list(struct vtn_builder *b, struct list_head *cf_list,
any = any ? nir_ior(&b->nb, any, cond) : cond; any = any ? nir_ior(&b->nb, any, cond) : cond;
conditions[i++] = cond; conditions[i++] = cond;
} }
assert(i == num_cases); vtn_assert(i == num_cases);
/* Now we can walk the list of cases and actually emit code */ /* Now we can walk the list of cases and actually emit code */
i = 0; i = 0;
@@ -736,7 +737,7 @@ vtn_emit_cf_list(struct vtn_builder *b, struct list_head *cf_list,
/* Figure out the condition */ /* Figure out the condition */
nir_ssa_def *cond = conditions[i++]; nir_ssa_def *cond = conditions[i++];
if (cse->is_default) { if (cse->is_default) {
assert(cond == NULL); vtn_assert(cond == NULL);
cond = nir_inot(&b->nb, any); cond = nir_inot(&b->nb, any);
} }
/* Take fallthrough into account */ /* Take fallthrough into account */
@@ -751,7 +752,7 @@ vtn_emit_cf_list(struct vtn_builder *b, struct list_head *cf_list,
nir_pop_if(&b->nb, case_if); nir_pop_if(&b->nb, case_if);
} }
assert(i == num_cases); vtn_assert(i == num_cases);
break; break;
} }

4
src/compiler/spirv/vtn_glsl450.c
View File

@@ -515,7 +515,7 @@ handle_glsl450_alu(struct vtn_builder *b, enum GLSLstd450 entrypoint,
case GLSLstd450ModfStruct: { case GLSLstd450ModfStruct: {
nir_ssa_def *sign = nir_fsign(nb, src[0]); nir_ssa_def *sign = nir_fsign(nb, src[0]);
nir_ssa_def *abs = nir_fabs(nb, src[0]); nir_ssa_def *abs = nir_fabs(nb, src[0]);
assert(glsl_type_is_struct(val->ssa->type)); vtn_assert(glsl_type_is_struct(val->ssa->type));
val->ssa->elems[0]->def = nir_fmul(nb, sign, nir_ffract(nb, abs)); val->ssa->elems[0]->def = nir_fmul(nb, sign, nir_ffract(nb, abs));
val->ssa->elems[1]->def = nir_fmul(nb, sign, nir_ffloor(nb, abs)); val->ssa->elems[1]->def = nir_fmul(nb, sign, nir_ffloor(nb, abs));
return; return;
@@ -690,7 +690,7 @@ handle_glsl450_alu(struct vtn_builder *b, enum GLSLstd450 entrypoint,
} }
case GLSLstd450FrexpStruct: { case GLSLstd450FrexpStruct: {
assert(glsl_type_is_struct(val->ssa->type)); vtn_assert(glsl_type_is_struct(val->ssa->type));
val->ssa->elems[0]->def = build_frexp(nb, src[0], val->ssa->elems[0]->def = build_frexp(nb, src[0],
&val->ssa->elems[1]->def); &val->ssa->elems[1]->def);
return; return;

202
src/compiler/spirv/vtn_variables.c
View File

@@ -71,14 +71,14 @@ vtn_access_chain_pointer_dereference(struct vtn_builder *b,
* pointers. For everything else, the client is expected to just pass us * pointers. For everything else, the client is expected to just pass us
* the right access chain. * the right access chain.
*/ */
assert(!deref_chain->ptr_as_array); vtn_assert(!deref_chain->ptr_as_array);
unsigned start = base->chain ? base->chain->length : 0; unsigned start = base->chain ? base->chain->length : 0;
for (unsigned i = 0; i < deref_chain->length; i++) { for (unsigned i = 0; i < deref_chain->length; i++) {
chain->link[start + i] = deref_chain->link[i]; chain->link[start + i] = deref_chain->link[i];
if (glsl_type_is_struct(type->type)) { if (glsl_type_is_struct(type->type)) {
assert(deref_chain->link[i].mode == vtn_access_mode_literal); vtn_assert(deref_chain->link[i].mode == vtn_access_mode_literal);
type = type->members[deref_chain->link[i].id]; type = type->members[deref_chain->link[i].id];
} else { } else {
type = type->array_element; type = type->array_element;
@@ -98,7 +98,7 @@ static nir_ssa_def *
vtn_access_link_as_ssa(struct vtn_builder *b, struct vtn_access_link link, vtn_access_link_as_ssa(struct vtn_builder *b, struct vtn_access_link link,
unsigned stride) unsigned stride)
{ {
assert(stride > 0); vtn_assert(stride > 0);
if (link.mode == vtn_access_mode_literal) { if (link.mode == vtn_access_mode_literal) {
return nir_imm_int(&b->nb, link.id * stride); return nir_imm_int(&b->nb, link.id * stride);
} else if (stride == 1) { } else if (stride == 1) {
@@ -119,7 +119,7 @@ vtn_variable_resource_index(struct vtn_builder *b, struct vtn_variable *var,
nir_ssa_def *desc_array_index) nir_ssa_def *desc_array_index)
{ {
if (!desc_array_index) { if (!desc_array_index) {
assert(glsl_type_is_struct(var->type->type)); vtn_assert(glsl_type_is_struct(var->type->type));
desc_array_index = nir_imm_int(&b->nb, 0); desc_array_index = nir_imm_int(&b->nb, 0);
} }
@@ -148,11 +148,11 @@ vtn_ssa_offset_pointer_dereference(struct vtn_builder *b,
unsigned idx = 0; unsigned idx = 0;
if (deref_chain->ptr_as_array) { if (deref_chain->ptr_as_array) {
/* We need ptr_type for the stride */ /* We need ptr_type for the stride */
assert(base->ptr_type); vtn_assert(base->ptr_type);
/* This must be a pointer to an actual element somewhere */ /* This must be a pointer to an actual element somewhere */
assert(block_index && offset); vtn_assert(block_index && offset);
/* We need at least one element in the chain */ /* We need at least one element in the chain */
assert(deref_chain->length >= 1); vtn_assert(deref_chain->length >= 1);
nir_ssa_def *elem_offset = nir_ssa_def *elem_offset =
vtn_access_link_as_ssa(b, deref_chain->link[idx], vtn_access_link_as_ssa(b, deref_chain->link[idx],
@@ -162,10 +162,10 @@ vtn_ssa_offset_pointer_dereference(struct vtn_builder *b,
} }
if (!block_index) { if (!block_index) {
assert(base->var); vtn_assert(base->var);
if (glsl_type_is_array(type->type)) { if (glsl_type_is_array(type->type)) {
/* We need at least one element in the chain */ /* We need at least one element in the chain */
assert(deref_chain->length >= 1); vtn_assert(deref_chain->length >= 1);
nir_ssa_def *desc_arr_idx = nir_ssa_def *desc_arr_idx =
vtn_access_link_as_ssa(b, deref_chain->link[0], 1); vtn_access_link_as_ssa(b, deref_chain->link[0], 1);
@@ -177,10 +177,10 @@ vtn_ssa_offset_pointer_dereference(struct vtn_builder *b,
} }
/* This is the first access chain so we also need an offset */ /* This is the first access chain so we also need an offset */
assert(!offset); vtn_assert(!offset);
offset = nir_imm_int(&b->nb, 0); offset = nir_imm_int(&b->nb, 0);
} }
assert(offset); vtn_assert(offset);
for (; idx < deref_chain->length; idx++) { for (; idx < deref_chain->length; idx++) {
switch (glsl_get_base_type(type->type)) { switch (glsl_get_base_type(type->type)) {
@@ -200,7 +200,7 @@ vtn_ssa_offset_pointer_dereference(struct vtn_builder *b,
} }
case GLSL_TYPE_STRUCT: { case GLSL_TYPE_STRUCT: {
assert(deref_chain->link[idx].mode == vtn_access_mode_literal); vtn_assert(deref_chain->link[idx].mode == vtn_access_mode_literal);
unsigned member = deref_chain->link[idx].id; unsigned member = deref_chain->link[idx].id;
nir_ssa_def *mem_offset = nir_imm_int(&b->nb, type->offsets[member]); nir_ssa_def *mem_offset = nir_imm_int(&b->nb, type->offsets[member]);
offset = nir_iadd(&b->nb, offset, mem_offset); offset = nir_iadd(&b->nb, offset, mem_offset);
@@ -240,13 +240,14 @@ vtn_pointer_dereference(struct vtn_builder *b,
* tail_type. This is useful for split structures. * tail_type. This is useful for split structures.
*/ */
static void static void
rewrite_deref_types(nir_deref *deref, const struct glsl_type *type) rewrite_deref_types(struct vtn_builder *b, nir_deref *deref,
const struct glsl_type *type)
{ {
deref->type = type; deref->type = type;
if (deref->child) { if (deref->child) {
assert(deref->child->deref_type == nir_deref_type_array); vtn_assert(deref->child->deref_type == nir_deref_type_array);
assert(glsl_type_is_array(deref->type)); vtn_assert(glsl_type_is_array(deref->type));
rewrite_deref_types(deref->child, glsl_get_array_element(type)); rewrite_deref_types(b, deref->child, glsl_get_array_element(type));
} }
} }
@@ -258,8 +259,8 @@ vtn_pointer_for_variable(struct vtn_builder *b,
pointer->mode = var->mode; pointer->mode = var->mode;
pointer->type = var->type; pointer->type = var->type;
assert(ptr_type->base_type == vtn_base_type_pointer); vtn_assert(ptr_type->base_type == vtn_base_type_pointer);
assert(ptr_type->deref->type == var->type->type); vtn_assert(ptr_type->deref->type == var->type->type);
pointer->ptr_type = ptr_type; pointer->ptr_type = ptr_type;
pointer->var = var; pointer->var = var;
@@ -280,14 +281,14 @@ vtn_pointer_to_deref(struct vtn_builder *b, struct vtn_pointer *ptr)
if (!ptr->chain) if (!ptr->chain)
return deref_var; return deref_var;
} else { } else {
assert(ptr->var->members); vtn_assert(ptr->var->members);
/* Create the deref_var manually. It will get filled out later. */ /* Create the deref_var manually. It will get filled out later. */
deref_var = rzalloc(b, nir_deref_var); deref_var = rzalloc(b, nir_deref_var);
deref_var->deref.deref_type = nir_deref_type_var; deref_var->deref.deref_type = nir_deref_type_var;
} }
struct vtn_access_chain *chain = ptr->chain; struct vtn_access_chain *chain = ptr->chain;
assert(chain); vtn_assert(chain);
struct vtn_type *deref_type = ptr->var->type; struct vtn_type *deref_type = ptr->var->type;
nir_deref *tail = &deref_var->deref; nir_deref *tail = &deref_var->deref;
@@ -313,7 +314,7 @@ vtn_pointer_to_deref(struct vtn_builder *b, struct vtn_pointer *ptr)
deref_arr->deref_array_type = nir_deref_array_type_direct; deref_arr->deref_array_type = nir_deref_array_type_direct;
deref_arr->base_offset = chain->link[i].id; deref_arr->base_offset = chain->link[i].id;
} else { } else {
assert(chain->link[i].mode == vtn_access_mode_id); vtn_assert(chain->link[i].mode == vtn_access_mode_id);
deref_arr->deref_array_type = nir_deref_array_type_indirect; deref_arr->deref_array_type = nir_deref_array_type_indirect;
deref_arr->base_offset = 0; deref_arr->base_offset = 0;
deref_arr->indirect = deref_arr->indirect =
@@ -325,14 +326,14 @@ vtn_pointer_to_deref(struct vtn_builder *b, struct vtn_pointer *ptr)
} }
case GLSL_TYPE_STRUCT: { case GLSL_TYPE_STRUCT: {
assert(chain->link[i].mode == vtn_access_mode_literal); vtn_assert(chain->link[i].mode == vtn_access_mode_literal);
unsigned idx = chain->link[i].id; unsigned idx = chain->link[i].id;
deref_type = deref_type->members[idx]; deref_type = deref_type->members[idx];
if (members) { if (members) {
/* This is a pre-split structure. */ /* This is a pre-split structure. */
deref_var->var = members[idx]; deref_var->var = members[idx];
rewrite_deref_types(&deref_var->deref, members[idx]->type); rewrite_deref_types(b, &deref_var->deref, members[idx]->type);
assert(tail->type == deref_type->type); vtn_assert(tail->type == deref_type->type);
members = NULL; members = NULL;
} else { } else {
nir_deref_struct *deref_struct = nir_deref_struct_create(b, idx); nir_deref_struct *deref_struct = nir_deref_struct_create(b, idx);
@@ -347,7 +348,7 @@ vtn_pointer_to_deref(struct vtn_builder *b, struct vtn_pointer *ptr)
} }
} }
assert(members == NULL); vtn_assert(members == NULL);
return deref_var; return deref_var;
} }
@@ -398,7 +399,7 @@ _vtn_local_load_store(struct vtn_builder *b, bool load, nir_deref_var *deref,
_vtn_local_load_store(b, load, deref, tail->child, inout->elems[i]); _vtn_local_load_store(b, load, deref, tail->child, inout->elems[i]);
} }
} else { } else {
assert(glsl_get_base_type(tail->type) == GLSL_TYPE_STRUCT); vtn_assert(glsl_get_base_type(tail->type) == GLSL_TYPE_STRUCT);
unsigned elems = glsl_get_length(tail->type); unsigned elems = glsl_get_length(tail->type);
nir_deref_struct *deref_struct = nir_deref_struct_create(b, 0); nir_deref_struct *deref_struct = nir_deref_struct_create(b, 0);
tail->child = &deref_struct->deref; tail->child = &deref_struct->deref;
@@ -443,7 +444,7 @@ vtn_local_load(struct vtn_builder *b, nir_deref_var *src)
if (src_tail->child) { if (src_tail->child) {
nir_deref_array *vec_deref = nir_deref_as_array(src_tail->child); nir_deref_array *vec_deref = nir_deref_as_array(src_tail->child);
assert(vec_deref->deref.child == NULL); vtn_assert(vec_deref->deref.child == NULL);
val->type = vec_deref->deref.type; val->type = vec_deref->deref.type;
if (vec_deref->deref_array_type == nir_deref_array_type_direct) if (vec_deref->deref_array_type == nir_deref_array_type_direct)
val->def = vtn_vector_extract(b, val->def, vec_deref->base_offset); val->def = vtn_vector_extract(b, val->def, vec_deref->base_offset);
@@ -465,7 +466,7 @@ vtn_local_store(struct vtn_builder *b, struct vtn_ssa_value *src,
struct vtn_ssa_value *val = vtn_create_ssa_value(b, dest_tail->type); struct vtn_ssa_value *val = vtn_create_ssa_value(b, dest_tail->type);
_vtn_local_load_store(b, true, dest, dest_tail, val); _vtn_local_load_store(b, true, dest, dest_tail, val);
nir_deref_array *deref = nir_deref_as_array(dest_tail->child); nir_deref_array *deref = nir_deref_as_array(dest_tail->child);
assert(deref->deref.child == NULL); vtn_assert(deref->deref.child == NULL);
if (deref->deref_array_type == nir_deref_array_type_direct) if (deref->deref_array_type == nir_deref_array_type_direct)
val->def = vtn_vector_insert(b, val->def, src->def, val->def = vtn_vector_insert(b, val->def, src->def,
deref->base_offset); deref->base_offset);
@@ -490,7 +491,7 @@ get_vulkan_resource_index(struct vtn_builder *b, struct vtn_pointer *ptr,
} }
if (glsl_type_is_array(ptr->var->type->type)) { if (glsl_type_is_array(ptr->var->type->type)) {
assert(ptr->chain->length > 0); vtn_assert(ptr->chain->length > 0);
nir_ssa_def *desc_array_index = nir_ssa_def *desc_array_index =
vtn_access_link_as_ssa(b, ptr->chain->link[0], 1); vtn_access_link_as_ssa(b, ptr->chain->link[0], 1);
*chain_idx = 1; *chain_idx = 1;
@@ -508,7 +509,7 @@ vtn_pointer_to_offset(struct vtn_builder *b, struct vtn_pointer *ptr,
nir_ssa_def **index_out, unsigned *end_idx_out) nir_ssa_def **index_out, unsigned *end_idx_out)
{ {
if (ptr->offset) { if (ptr->offset) {
assert(ptr->block_index); vtn_assert(ptr->block_index);
*index_out = ptr->block_index; *index_out = ptr->block_index;
return ptr->offset; return ptr->offset;
} }
@@ -537,7 +538,7 @@ vtn_pointer_to_offset(struct vtn_builder *b, struct vtn_pointer *ptr,
break; break;
case GLSL_TYPE_STRUCT: { case GLSL_TYPE_STRUCT: {
assert(ptr->chain->link[idx].mode == vtn_access_mode_literal); vtn_assert(ptr->chain->link[idx].mode == vtn_access_mode_literal);
unsigned member = ptr->chain->link[idx].id; unsigned member = ptr->chain->link[idx].id;
offset = nir_iadd(&b->nb, offset, offset = nir_iadd(&b->nb, offset,
nir_imm_int(&b->nb, type->offsets[member])); nir_imm_int(&b->nb, type->offsets[member]));
@@ -550,7 +551,7 @@ vtn_pointer_to_offset(struct vtn_builder *b, struct vtn_pointer *ptr,
} }
} }
assert(type == ptr->type); vtn_assert(type == ptr->type);
if (end_idx_out) if (end_idx_out)
*end_idx_out = idx; *end_idx_out = idx;
@@ -561,7 +562,7 @@ vtn_pointer_to_offset(struct vtn_builder *b, struct vtn_pointer *ptr,
* offsets that are provided to us in the SPIR-V source. * offsets that are provided to us in the SPIR-V source.
*/ */
static unsigned static unsigned
vtn_type_block_size(struct vtn_type *type) vtn_type_block_size(struct vtn_builder *b, struct vtn_type *type)
{ {
enum glsl_base_type base_type = glsl_get_base_type(type->type); enum glsl_base_type base_type = glsl_get_base_type(type->type);
switch (base_type) { switch (base_type) {
@@ -575,7 +576,7 @@ vtn_type_block_size(struct vtn_type *type)
unsigned cols = type->row_major ? glsl_get_vector_elements(type->type) : unsigned cols = type->row_major ? glsl_get_vector_elements(type->type) :
glsl_get_matrix_columns(type->type); glsl_get_matrix_columns(type->type);
if (cols > 1) { if (cols > 1) {
assert(type->stride > 0); vtn_assert(type->stride > 0);
return type->stride * cols; return type->stride * cols;
} else if (base_type == GLSL_TYPE_DOUBLE || } else if (base_type == GLSL_TYPE_DOUBLE ||
base_type == GLSL_TYPE_UINT64 || base_type == GLSL_TYPE_UINT64 ||
@@ -592,15 +593,15 @@ vtn_type_block_size(struct vtn_type *type)
unsigned num_fields = glsl_get_length(type->type); unsigned num_fields = glsl_get_length(type->type);
for (unsigned f = 0; f < num_fields; f++) { for (unsigned f = 0; f < num_fields; f++) {
unsigned field_end = type->offsets[f] + unsigned field_end = type->offsets[f] +
vtn_type_block_size(type->members[f]); vtn_type_block_size(b, type->members[f]);
size = MAX2(size, field_end); size = MAX2(size, field_end);
} }
return size; return size;
} }
case GLSL_TYPE_ARRAY: case GLSL_TYPE_ARRAY:
assert(type->stride > 0); vtn_assert(type->stride > 0);
assert(glsl_get_length(type->type) > 0); vtn_assert(glsl_get_length(type->type) > 0);
return type->stride * glsl_get_length(type->type); return type->stride * glsl_get_length(type->type);
default: default:
@@ -610,7 +611,8 @@ vtn_type_block_size(struct vtn_type *type)
} }
static void static void
vtn_access_chain_get_offset_size(struct vtn_access_chain *chain, vtn_access_chain_get_offset_size(struct vtn_builder *b,
struct vtn_access_chain *chain,
struct vtn_type *type, struct vtn_type *type,
unsigned *access_offset, unsigned *access_offset,
unsigned *access_size) unsigned *access_size)
@@ -630,7 +632,7 @@ vtn_access_chain_get_offset_size(struct vtn_access_chain *chain,
} }
} }
*access_size = vtn_type_block_size(type); *access_size = vtn_type_block_size(b, type);
} }
static void static void
@@ -649,7 +651,7 @@ _vtn_load_store_tail(struct vtn_builder *b, nir_intrinsic_op op, bool load,
} }
if (op == nir_intrinsic_load_push_constant) { if (op == nir_intrinsic_load_push_constant) {
assert(access_offset % 4 == 0); vtn_assert(access_offset % 4 == 0);
nir_intrinsic_set_base(instr, access_offset); nir_intrinsic_set_base(instr, access_offset);
nir_intrinsic_set_range(instr, access_size); nir_intrinsic_set_range(instr, access_size);
@@ -744,7 +746,7 @@ _vtn_block_load_store(struct vtn_builder *b, nir_intrinsic_op op, bool load,
unsigned type_size = glsl_get_bit_size(type->type) / 8; unsigned type_size = glsl_get_bit_size(type->type) / 8;
if (elems == 1 || type->stride == type_size) { if (elems == 1 || type->stride == type_size) {
/* This is a tightly-packed normal scalar or vector load */ /* This is a tightly-packed normal scalar or vector load */
assert(glsl_type_is_vector_or_scalar(type->type)); vtn_assert(glsl_type_is_vector_or_scalar(type->type));
_vtn_load_store_tail(b, op, load, index, offset, _vtn_load_store_tail(b, op, load, index, offset,
access_offset, access_size, access_offset, access_size,
inout, type->type); inout, type->type);
@@ -752,8 +754,8 @@ _vtn_block_load_store(struct vtn_builder *b, nir_intrinsic_op op, bool load,
/* This is a strided load. We have to load N things separately. /* This is a strided load. We have to load N things separately.
* This is the single column of a row-major matrix case. * This is the single column of a row-major matrix case.
*/ */
assert(type->stride > type_size); vtn_assert(type->stride > type_size);
assert(type->stride % type_size == 0); vtn_assert(type->stride % type_size == 0);
nir_ssa_def *per_comp[4]; nir_ssa_def *per_comp[4];
for (unsigned i = 0; i < elems; i++) { for (unsigned i = 0; i < elems; i++) {
@@ -826,7 +828,7 @@ vtn_block_load(struct vtn_builder *b, struct vtn_pointer *src)
break; break;
case vtn_variable_mode_push_constant: case vtn_variable_mode_push_constant:
op = nir_intrinsic_load_push_constant; op = nir_intrinsic_load_push_constant;
vtn_access_chain_get_offset_size(src->chain, src->var->type, vtn_access_chain_get_offset_size(b, src->chain, src->var->type,
&access_offset, &access_size); &access_offset, &access_size);
break; break;
default: default:
@@ -895,7 +897,7 @@ _vtn_variable_load_store(struct vtn_builder *b, bool load,
case GLSL_TYPE_STRUCT: { case GLSL_TYPE_STRUCT: {
unsigned elems = glsl_get_length(ptr->type->type); unsigned elems = glsl_get_length(ptr->type->type);
if (load) { if (load) {
assert(*inout == NULL); vtn_assert(*inout == NULL);
*inout = rzalloc(b, struct vtn_ssa_value); *inout = rzalloc(b, struct vtn_ssa_value);
(*inout)->type = ptr->type->type; (*inout)->type = ptr->type->type;
(*inout)->elems = rzalloc_array(b, struct vtn_ssa_value *, elems); (*inout)->elems = rzalloc_array(b, struct vtn_ssa_value *, elems);
@@ -937,7 +939,7 @@ vtn_variable_store(struct vtn_builder *b, struct vtn_ssa_value *src,
struct vtn_pointer *dest) struct vtn_pointer *dest)
{ {
if (vtn_pointer_is_external_block(dest)) { if (vtn_pointer_is_external_block(dest)) {
assert(dest->mode == vtn_variable_mode_ssbo); vtn_assert(dest->mode == vtn_variable_mode_ssbo);
vtn_block_store(b, src, dest); vtn_block_store(b, src, dest);
} else { } else {
_vtn_variable_load_store(b, false, dest, &src); _vtn_variable_load_store(b, false, dest, &src);
@@ -948,7 +950,7 @@ static void
_vtn_variable_copy(struct vtn_builder *b, struct vtn_pointer *dest, _vtn_variable_copy(struct vtn_builder *b, struct vtn_pointer *dest,
struct vtn_pointer *src) struct vtn_pointer *src)
{ {
assert(src->type->type == dest->type->type); vtn_assert(src->type->type == dest->type->type);
enum glsl_base_type base_type = glsl_get_base_type(src->type->type); enum glsl_base_type base_type = glsl_get_base_type(src->type->type);
switch (base_type) { switch (base_type) {
case GLSL_TYPE_UINT: case GLSL_TYPE_UINT:
@@ -1004,9 +1006,9 @@ vtn_variable_copy(struct vtn_builder *b, struct vtn_pointer *dest,
} }
static void static void
set_mode_system_value(nir_variable_mode *mode) set_mode_system_value(struct vtn_builder *b, nir_variable_mode *mode)
{ {
assert(*mode == nir_var_system_value || *mode == nir_var_shader_in); vtn_assert(*mode == nir_var_system_value || *mode == nir_var_shader_in);
*mode = nir_var_system_value; *mode = nir_var_system_value;
} }
@@ -1030,37 +1032,37 @@ vtn_get_builtin_location(struct vtn_builder *b,
break; break;
case SpvBuiltInVertexIndex: case SpvBuiltInVertexIndex:
*location = SYSTEM_VALUE_VERTEX_ID; *location = SYSTEM_VALUE_VERTEX_ID;
set_mode_system_value(mode); set_mode_system_value(b, mode);
break; break;
case SpvBuiltInVertexId: case SpvBuiltInVertexId:
/* Vulkan defines VertexID to be zero-based and reserves the new /* Vulkan defines VertexID to be zero-based and reserves the new
* builtin keyword VertexIndex to indicate the non-zero-based value. * builtin keyword VertexIndex to indicate the non-zero-based value.
*/ */
*location = SYSTEM_VALUE_VERTEX_ID_ZERO_BASE; *location = SYSTEM_VALUE_VERTEX_ID_ZERO_BASE;
set_mode_system_value(mode); set_mode_system_value(b, mode);
break; break;
case SpvBuiltInInstanceIndex: case SpvBuiltInInstanceIndex:
*location = SYSTEM_VALUE_INSTANCE_INDEX; *location = SYSTEM_VALUE_INSTANCE_INDEX;
set_mode_system_value(mode); set_mode_system_value(b, mode);
break; break;
case SpvBuiltInInstanceId: case SpvBuiltInInstanceId:
*location = SYSTEM_VALUE_INSTANCE_ID; *location = SYSTEM_VALUE_INSTANCE_ID;
set_mode_system_value(mode); set_mode_system_value(b, mode);
break; break;
case SpvBuiltInPrimitiveId: case SpvBuiltInPrimitiveId:
if (b->shader->info.stage == MESA_SHADER_FRAGMENT) { if (b->shader->info.stage == MESA_SHADER_FRAGMENT) {
assert(*mode == nir_var_shader_in); vtn_assert(*mode == nir_var_shader_in);
*location = VARYING_SLOT_PRIMITIVE_ID; *location = VARYING_SLOT_PRIMITIVE_ID;
} else if (*mode == nir_var_shader_out) { } else if (*mode == nir_var_shader_out) {
*location = VARYING_SLOT_PRIMITIVE_ID; *location = VARYING_SLOT_PRIMITIVE_ID;
} else { } else {
*location = SYSTEM_VALUE_PRIMITIVE_ID; *location = SYSTEM_VALUE_PRIMITIVE_ID;
set_mode_system_value(mode); set_mode_system_value(b, mode);
} }
break; break;
case SpvBuiltInInvocationId: case SpvBuiltInInvocationId:
*location = SYSTEM_VALUE_INVOCATION_ID; *location = SYSTEM_VALUE_INVOCATION_ID;
set_mode_system_value(mode); set_mode_system_value(b, mode);
break; break;
case SpvBuiltInLayer: case SpvBuiltInLayer:
*location = VARYING_SLOT_LAYER; *location = VARYING_SLOT_LAYER;
@@ -1088,51 +1090,51 @@ vtn_get_builtin_location(struct vtn_builder *b,
break; break;
case SpvBuiltInTessCoord: case SpvBuiltInTessCoord:
*location = SYSTEM_VALUE_TESS_COORD; *location = SYSTEM_VALUE_TESS_COORD;
set_mode_system_value(mode); set_mode_system_value(b, mode);
break; break;
case SpvBuiltInPatchVertices: case SpvBuiltInPatchVertices:
*location = SYSTEM_VALUE_VERTICES_IN; *location = SYSTEM_VALUE_VERTICES_IN;
set_mode_system_value(mode); set_mode_system_value(b, mode);
break; break;
case SpvBuiltInFragCoord: case SpvBuiltInFragCoord:
*location = VARYING_SLOT_POS; *location = VARYING_SLOT_POS;
assert(*mode == nir_var_shader_in); vtn_assert(*mode == nir_var_shader_in);
break; break;
case SpvBuiltInPointCoord: case SpvBuiltInPointCoord:
*location = VARYING_SLOT_PNTC; *location = VARYING_SLOT_PNTC;
assert(*mode == nir_var_shader_in); vtn_assert(*mode == nir_var_shader_in);
break; break;
case SpvBuiltInFrontFacing: case SpvBuiltInFrontFacing:
*location = SYSTEM_VALUE_FRONT_FACE; *location = SYSTEM_VALUE_FRONT_FACE;
set_mode_system_value(mode); set_mode_system_value(b, mode);
break; break;
case SpvBuiltInSampleId: case SpvBuiltInSampleId:
*location = SYSTEM_VALUE_SAMPLE_ID; *location = SYSTEM_VALUE_SAMPLE_ID;
set_mode_system_value(mode); set_mode_system_value(b, mode);
break; break;
case SpvBuiltInSamplePosition: case SpvBuiltInSamplePosition:
*location = SYSTEM_VALUE_SAMPLE_POS; *location = SYSTEM_VALUE_SAMPLE_POS;
set_mode_system_value(mode); set_mode_system_value(b, mode);
break; break;
case SpvBuiltInSampleMask: case SpvBuiltInSampleMask:
if (*mode == nir_var_shader_out) { if (*mode == nir_var_shader_out) {
*location = FRAG_RESULT_SAMPLE_MASK; *location = FRAG_RESULT_SAMPLE_MASK;
} else { } else {
*location = SYSTEM_VALUE_SAMPLE_MASK_IN; *location = SYSTEM_VALUE_SAMPLE_MASK_IN;
set_mode_system_value(mode); set_mode_system_value(b, mode);
} }
break; break;
case SpvBuiltInFragDepth: case SpvBuiltInFragDepth:
*location = FRAG_RESULT_DEPTH; *location = FRAG_RESULT_DEPTH;
assert(*mode == nir_var_shader_out); vtn_assert(*mode == nir_var_shader_out);
break; break;
case SpvBuiltInHelperInvocation: case SpvBuiltInHelperInvocation:
*location = SYSTEM_VALUE_HELPER_INVOCATION; *location = SYSTEM_VALUE_HELPER_INVOCATION;
set_mode_system_value(mode); set_mode_system_value(b, mode);
break; break;
case SpvBuiltInNumWorkgroups: case SpvBuiltInNumWorkgroups:
*location = SYSTEM_VALUE_NUM_WORK_GROUPS; *location = SYSTEM_VALUE_NUM_WORK_GROUPS;
set_mode_system_value(mode); set_mode_system_value(b, mode);
break; break;
case SpvBuiltInWorkgroupSize: case SpvBuiltInWorkgroupSize:
/* This should already be handled */ /* This should already be handled */
@@ -1140,35 +1142,35 @@ vtn_get_builtin_location(struct vtn_builder *b,
break; break;
case SpvBuiltInWorkgroupId: case SpvBuiltInWorkgroupId:
*location = SYSTEM_VALUE_WORK_GROUP_ID; *location = SYSTEM_VALUE_WORK_GROUP_ID;
set_mode_system_value(mode); set_mode_system_value(b, mode);
break; break;
case SpvBuiltInLocalInvocationId: case SpvBuiltInLocalInvocationId:
*location = SYSTEM_VALUE_LOCAL_INVOCATION_ID; *location = SYSTEM_VALUE_LOCAL_INVOCATION_ID;
set_mode_system_value(mode); set_mode_system_value(b, mode);
break; break;
case SpvBuiltInLocalInvocationIndex: case SpvBuiltInLocalInvocationIndex:
*location = SYSTEM_VALUE_LOCAL_INVOCATION_INDEX; *location = SYSTEM_VALUE_LOCAL_INVOCATION_INDEX;
set_mode_system_value(mode); set_mode_system_value(b, mode);
break; break;
case SpvBuiltInGlobalInvocationId: case SpvBuiltInGlobalInvocationId:
*location = SYSTEM_VALUE_GLOBAL_INVOCATION_ID; *location = SYSTEM_VALUE_GLOBAL_INVOCATION_ID;
set_mode_system_value(mode); set_mode_system_value(b, mode);
break; break;
case SpvBuiltInBaseVertex: case SpvBuiltInBaseVertex:
*location = SYSTEM_VALUE_BASE_VERTEX; *location = SYSTEM_VALUE_BASE_VERTEX;
set_mode_system_value(mode); set_mode_system_value(b, mode);
break; break;
case SpvBuiltInBaseInstance: case SpvBuiltInBaseInstance:
*location = SYSTEM_VALUE_BASE_INSTANCE; *location = SYSTEM_VALUE_BASE_INSTANCE;
set_mode_system_value(mode); set_mode_system_value(b, mode);
break; break;
case SpvBuiltInDrawIndex: case SpvBuiltInDrawIndex:
*location = SYSTEM_VALUE_DRAW_ID; *location = SYSTEM_VALUE_DRAW_ID;
set_mode_system_value(mode); set_mode_system_value(b, mode);
break; break;
case SpvBuiltInViewIndex: case SpvBuiltInViewIndex:
*location = SYSTEM_VALUE_VIEW_INDEX; *location = SYSTEM_VALUE_VIEW_INDEX;
set_mode_system_value(mode); set_mode_system_value(b, mode);
break; break;
default: default:
unreachable("unsupported builtin"); unreachable("unsupported builtin");
@@ -1198,7 +1200,7 @@ apply_var_decoration(struct vtn_builder *b, nir_variable *nir_var,
nir_var->data.invariant = true; nir_var->data.invariant = true;
break; break;
case SpvDecorationConstant: case SpvDecorationConstant:
assert(nir_var->constant_initializer != NULL); vtn_assert(nir_var->constant_initializer != NULL);
nir_var->data.read_only = true; nir_var->data.read_only = true;
break; break;
case SpvDecorationNonReadable: case SpvDecorationNonReadable:
@@ -1341,11 +1343,11 @@ var_decoration_cb(struct vtn_builder *b, struct vtn_value *val, int member,
} }
if (val->value_type == vtn_value_type_pointer) { if (val->value_type == vtn_value_type_pointer) {
assert(val->pointer->var == void_var); vtn_assert(val->pointer->var == void_var);
assert(val->pointer->chain == NULL); vtn_assert(val->pointer->chain == NULL);
assert(member == -1); vtn_assert(member == -1);
} else { } else {
assert(val->value_type == vtn_value_type_type); vtn_assert(val->value_type == vtn_value_type_type);
} }
/* Location is odd. If applied to a split structure, we have to walk the /* Location is odd. If applied to a split structure, we have to walk the
@@ -1377,7 +1379,7 @@ var_decoration_cb(struct vtn_builder *b, struct vtn_value *val, int member,
vtn_var->var->data.location = location; vtn_var->var->data.location = location;
} else { } else {
/* This handles the structure member case */ /* This handles the structure member case */
assert(vtn_var->members); vtn_assert(vtn_var->members);
unsigned length = unsigned length =
glsl_get_length(glsl_without_array(vtn_var->type->type)); glsl_get_length(glsl_without_array(vtn_var->type->type));
for (unsigned i = 0; i < length; i++) { for (unsigned i = 0; i < length; i++) {
@@ -1390,11 +1392,11 @@ var_decoration_cb(struct vtn_builder *b, struct vtn_value *val, int member,
return; return;
} else { } else {
if (vtn_var->var) { if (vtn_var->var) {
assert(member <= 0); vtn_assert(member <= 0);
apply_var_decoration(b, vtn_var->var, dec); apply_var_decoration(b, vtn_var->var, dec);
} else if (vtn_var->members) { } else if (vtn_var->members) {
if (member >= 0) { if (member >= 0) {
assert(vtn_var->members); vtn_assert(vtn_var->members);
apply_var_decoration(b, vtn_var->members[member], dec); apply_var_decoration(b, vtn_var->members[member], dec);
} else { } else {
unsigned length = unsigned length =
@@ -1407,7 +1409,7 @@ var_decoration_cb(struct vtn_builder *b, struct vtn_value *val, int member,
* nir_variables associated with them. Fortunately, all decorations * nir_variables associated with them. Fortunately, all decorations
* we care about for those variables are on the type only. * we care about for those variables are on the type only.
*/ */
assert(vtn_var->mode == vtn_variable_mode_ubo || vtn_assert(vtn_var->mode == vtn_variable_mode_ubo ||
vtn_var->mode == vtn_variable_mode_ssbo || vtn_var->mode == vtn_variable_mode_ssbo ||
vtn_var->mode == vtn_variable_mode_push_constant); vtn_var->mode == vtn_variable_mode_push_constant);
} }
@@ -1489,8 +1491,8 @@ nir_ssa_def *
vtn_pointer_to_ssa(struct vtn_builder *b, struct vtn_pointer *ptr) vtn_pointer_to_ssa(struct vtn_builder *b, struct vtn_pointer *ptr)
{ {
/* This pointer needs to have a pointer type with actual storage */ /* This pointer needs to have a pointer type with actual storage */
assert(ptr->ptr_type); vtn_assert(ptr->ptr_type);
assert(ptr->ptr_type->type); vtn_assert(ptr->ptr_type->type);
if (ptr->offset && ptr->block_index) { if (ptr->offset && ptr->block_index) {
return nir_vec2(&b->nb, ptr->block_index, ptr->offset); return nir_vec2(&b->nb, ptr->block_index, ptr->offset);
@@ -1498,13 +1500,13 @@ vtn_pointer_to_ssa(struct vtn_builder *b, struct vtn_pointer *ptr)
/* If we don't have an offset or block index, then we must be a pointer /* If we don't have an offset or block index, then we must be a pointer
* to the variable itself. * to the variable itself.
*/ */
assert(!ptr->offset && !ptr->block_index); vtn_assert(!ptr->offset && !ptr->block_index);
/* We can't handle a pointer to an array of descriptors because we have /* We can't handle a pointer to an array of descriptors because we have
* no way of knowing later on that we need to add to update the block * no way of knowing later on that we need to add to update the block
* index when dereferencing. * index when dereferencing.
*/ */
assert(ptr->var && ptr->var->type->base_type == vtn_base_type_struct); vtn_assert(ptr->var && ptr->var->type->base_type == vtn_base_type_struct);
return nir_vec2(&b->nb, vtn_variable_resource_index(b, ptr->var, NULL), return nir_vec2(&b->nb, vtn_variable_resource_index(b, ptr->var, NULL),
nir_imm_int(&b->nb, 0)); nir_imm_int(&b->nb, 0));
@@ -1515,11 +1517,11 @@ struct vtn_pointer *
vtn_pointer_from_ssa(struct vtn_builder *b, nir_ssa_def *ssa, vtn_pointer_from_ssa(struct vtn_builder *b, nir_ssa_def *ssa,
struct vtn_type *ptr_type) struct vtn_type *ptr_type)
{ {
assert(ssa->num_components == 2 && ssa->bit_size == 32); vtn_assert(ssa->num_components == 2 && ssa->bit_size == 32);
assert(ptr_type->base_type == vtn_base_type_pointer); vtn_assert(ptr_type->base_type == vtn_base_type_pointer);
assert(ptr_type->deref->base_type != vtn_base_type_pointer); vtn_assert(ptr_type->deref->base_type != vtn_base_type_pointer);
/* This pointer type needs to have actual storage */ /* This pointer type needs to have actual storage */
assert(ptr_type->type); vtn_assert(ptr_type->type);
struct vtn_pointer *ptr = rzalloc(b, struct vtn_pointer); struct vtn_pointer *ptr = rzalloc(b, struct vtn_pointer);
ptr->mode = vtn_storage_class_to_mode(ptr_type->storage_class, ptr->mode = vtn_storage_class_to_mode(ptr_type->storage_class,
@@ -1555,7 +1557,7 @@ vtn_create_variable(struct vtn_builder *b, struct vtn_value *val,
struct vtn_type *ptr_type, SpvStorageClass storage_class, struct vtn_type *ptr_type, SpvStorageClass storage_class,
nir_constant *initializer) nir_constant *initializer)
{ {
assert(ptr_type->base_type == vtn_base_type_pointer); vtn_assert(ptr_type->base_type == vtn_base_type_pointer);
struct vtn_type *type = ptr_type->deref; struct vtn_type *type = ptr_type->deref;
struct vtn_type *without_array = type; struct vtn_type *without_array = type;
@@ -1580,7 +1582,7 @@ vtn_create_variable(struct vtn_builder *b, struct vtn_value *val,
b->shader->info.num_textures++; b->shader->info.num_textures++;
break; break;
case vtn_variable_mode_push_constant: case vtn_variable_mode_push_constant:
b->shader->num_uniforms = vtn_type_block_size(type); b->shader->num_uniforms = vtn_type_block_size(b, type);
break; break;
default: default:
/* No tallying is needed */ /* No tallying is needed */
@@ -1591,7 +1593,7 @@ vtn_create_variable(struct vtn_builder *b, struct vtn_value *val,
var->type = type; var->type = type;
var->mode = mode; var->mode = mode;
assert(val->value_type == vtn_value_type_pointer); vtn_assert(val->value_type == vtn_value_type_pointer);
val->pointer = vtn_pointer_for_variable(b, var, ptr_type); val->pointer = vtn_pointer_for_variable(b, var, ptr_type);
switch (var->mode) { switch (var->mode) {
@@ -1730,18 +1732,18 @@ vtn_create_variable(struct vtn_builder *b, struct vtn_value *val,
} }
if (var->mode == vtn_variable_mode_local) { if (var->mode == vtn_variable_mode_local) {
assert(var->members == NULL && var->var != NULL); vtn_assert(var->members == NULL && var->var != NULL);
nir_function_impl_add_variable(b->nb.impl, var->var); nir_function_impl_add_variable(b->nb.impl, var->var);
} else if (var->var) { } else if (var->var) {
nir_shader_add_variable(b->shader, var->var); nir_shader_add_variable(b->shader, var->var);
} else if (var->members) { } else if (var->members) {
unsigned count = glsl_get_length(without_array->type); unsigned count = glsl_get_length(without_array->type);
for (unsigned i = 0; i < count; i++) { for (unsigned i = 0; i < count; i++) {
assert(var->members[i]->data.mode != nir_var_local); vtn_assert(var->members[i]->data.mode != nir_var_local);
nir_shader_add_variable(b->shader, var->members[i]); nir_shader_add_variable(b->shader, var->members[i]);
} }
} else { } else {
assert(var->mode == vtn_variable_mode_ubo || vtn_assert(var->mode == vtn_variable_mode_ubo ||
var->mode == vtn_variable_mode_ssbo || var->mode == vtn_variable_mode_ssbo ||
var->mode == vtn_variable_mode_push_constant); var->mode == vtn_variable_mode_push_constant);
} }
@@ -1810,7 +1812,7 @@ vtn_handle_variables(struct vtn_builder *b, SpvOp opcode,
vtn_pointer_dereference(b, base_val->sampled_image->image, chain); vtn_pointer_dereference(b, base_val->sampled_image->image, chain);
val->sampled_image->sampler = base_val->sampled_image->sampler; val->sampled_image->sampler = base_val->sampled_image->sampler;
} else { } else {
assert(base_val->value_type == vtn_value_type_pointer); vtn_assert(base_val->value_type == vtn_value_type_pointer);
struct vtn_value *val = struct vtn_value *val =
vtn_push_value(b, w[2], vtn_value_type_pointer); vtn_push_value(b, w[2], vtn_value_type_pointer);
val->pointer = vtn_pointer_dereference(b, base_val->pointer, chain); val->pointer = vtn_pointer_dereference(b, base_val->pointer, chain);
@@ -1850,7 +1852,7 @@ vtn_handle_variables(struct vtn_builder *b, SpvOp opcode,
if (glsl_type_is_sampler(dest->type->type)) { if (glsl_type_is_sampler(dest->type->type)) {
vtn_warn("OpStore of a sampler detected. Doing on-the-fly copy " vtn_warn("OpStore of a sampler detected. Doing on-the-fly copy "
"propagation to workaround the problem."); "propagation to workaround the problem.");
assert(dest->var->copy_prop_sampler == NULL); vtn_assert(dest->var->copy_prop_sampler == NULL);
dest->var->copy_prop_sampler = dest->var->copy_prop_sampler =
vtn_value(b, w[2], vtn_value_type_pointer)->pointer; vtn_value(b, w[2], vtn_value_type_pointer)->pointer;
break; break;