[Mesa-dev] [PATCH 0/8] spirv: Types, types, and OpSwitch

Discussion:

Jason Ekstrand

2017-12-07 16:12:07 UTC

When I started working on switching spirv_to_nir from having piles of
assert() to vtn_assert/fail, Ian and I both agreed that we should start
moving in a direction where we had vtn_fail with reasonable error messages
rather than vtn_assert() with some compiler-internal garbage message.
However some error checking is better than none, so we let it through with
the plan to come back and improve things. This series starts us down that
road.

The first two patches add an autogenerated pre-pass which records the types
of all SPIR-V values. Previously, we only really recorded types for
OpType*, OpConstant*, and OpSpecConstant*. Now, we have types for
everything.

The next four add some very basic parameter validation for OpTypeVector and
OpTypeMatrix and type validation for OpLoad, OpStore, OpCopyMemory, and
OpSelect. Now that we have the types recorded everywhere, it's fairly easy
to provide validation that looks a lot more like the SPIR-V spec than our
asserts and provides reasonable error messages. The reason for choosing to
start with those four is that they (along with OpPhi) are the hot-spots for
variable pointers and type mismatches are death there.

The last two patches fix a long-standing bug where we didn't properly
implement switch statements with 64-bit selectors. Given the infrequency
of 64-bit integers in shaders combined with the infrequency of switch
statements, the chances of hitting this in the wild are low (hence no CC to
stable) but we should make it work correctly. I haven't done this until
now because it relies on the type pre-pass because we need to know the type
of the OpSwitch selector while we're in the CFG pre-pass which happens
before we emit actual instructions.

Cc: Ian Romanick <***@intel.com>

Jason Ekstrand (8):
spirv: Add a vtn_type field to all vtn_values
spirv: Add a prepass to set types on vtn_values
spirv: Add basic type validation for OpLoad, OpStore, and OpCopyMemory
spirv: Set lengths on scalar and vector types
spirv: Add type validation for OpSelect
spirv: Add better parameter validation for vector and matrix types
spirv: Restructure the case loop in OpSwitch handling
spirv: Add support for all bit sizes in OpSwitch

src/compiler/Makefile.nir.am | 4 +
src/compiler/nir/meson.build | 11 ++-
src/compiler/spirv/spirv_to_nir.c | 117 +++++++++++++++++++----------
src/compiler/spirv/vtn_cfg.c | 44 +++++++++--
src/compiler/spirv/vtn_gather_types_c.py | 125 +++++++++++++++++++++++++++++++
src/compiler/spirv/vtn_private.h | 11 +--
src/compiler/spirv/vtn_variables.c | 18 ++++-
7 files changed, 273 insertions(+), 57 deletions(-)
create mode 100644 src/compiler/spirv/vtn_gather_types_c.py

--
2.5.0.400.gff86faf

Jason Ekstrand

2017-12-07 16:12:08 UTC

Permalink

At the moment, this just lets us drop the const_type for constants and
unify things a bit. Eventually, we will use this to store the types of
all SPIR-V SSA values.
---
src/compiler/spirv/spirv_to_nir.c | 63 +++++++++++++++++----------------------
src/compiler/spirv/vtn_private.h | 7 ++---
2 files changed, 29 insertions(+), 41 deletions(-)

diff --git a/src/compiler/spirv/spirv_to_nir.c b/src/compiler/spirv/spirv_to_nir.c
index d321d1a..a50b14d 100644
--- a/src/compiler/spirv/spirv_to_nir.c
+++ b/src/compiler/spirv/spirv_to_nir.c
@@ -257,7 +257,7 @@ vtn_ssa_value(struct vtn_builder *b, uint32_t value_id)
return vtn_undef_ssa_value(b, val->type->type);

case vtn_value_type_constant:
- return vtn_const_ssa_value(b, val->constant, val->const_type);
+ return vtn_const_ssa_value(b, val->constant, val->type->type);

case vtn_value_type_ssa:
return val->ssa;
@@ -1249,7 +1249,7 @@ handle_workgroup_size_decoration_cb(struct vtn_builder *b,
dec->literals[0] != SpvBuiltInWorkgroupSize)
return;

- vtn_assert(val->const_type == glsl_vector_type(GLSL_TYPE_UINT, 3));
+ vtn_assert(val->type->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[1] = val->constant->values[0].u32[1];
@@ -1261,21 +1261,21 @@ 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);
- val->const_type = vtn_value(b, w[1], vtn_value_type_type)->type->type;
+ val->type = vtn_value(b, w[1], vtn_value_type_type)->type;
val->constant = rzalloc(b, nir_constant);
switch (opcode) {
case SpvOpConstantTrue:
- vtn_assert(val->const_type == glsl_bool_type());
+ vtn_assert(val->type->type == glsl_bool_type());
val->constant->values[0].u32[0] = NIR_TRUE;
break;
case SpvOpConstantFalse:
- vtn_assert(val->const_type == glsl_bool_type());
+ vtn_assert(val->type->type == glsl_bool_type());
val->constant->values[0].u32[0] = NIR_FALSE;
break;

case SpvOpSpecConstantTrue:
case SpvOpSpecConstantFalse: {
- vtn_assert(val->const_type == glsl_bool_type());
+ vtn_assert(val->type->type == glsl_bool_type());
uint32_t int_val =
get_specialization(b, val, (opcode == SpvOpSpecConstantTrue));
val->constant->values[0].u32[0] = int_val ? NIR_TRUE : NIR_FALSE;
@@ -1283,8 +1283,8 @@ vtn_handle_constant(struct vtn_builder *b, SpvOp opcode,
}

case SpvOpConstant: {
- vtn_assert(glsl_type_is_scalar(val->const_type));
- int bit_size = glsl_get_bit_size(val->const_type);
+ vtn_assert(glsl_type_is_scalar(val->type->type));
+ int bit_size = glsl_get_bit_size(val->type->type);
switch (bit_size) {
case 64:
val->constant->values->u64[0] = vtn_u64_literal(&w[3]);
@@ -1301,9 +1301,9 @@ vtn_handle_constant(struct vtn_builder *b, SpvOp opcode,
break;
}
case SpvOpSpecConstant: {
- vtn_assert(glsl_type_is_scalar(val->const_type));
+ vtn_assert(glsl_type_is_scalar(val->type->type));
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->type->type);
switch (bit_size) {
case 64:
val->constant->values[0].u64[0] =
@@ -1327,7 +1327,7 @@ vtn_handle_constant(struct vtn_builder *b, SpvOp opcode,
for (unsigned i = 0; i < elem_count; i++)
elems[i] = vtn_value(b, w[i + 3], vtn_value_type_constant)->constant;

- switch (glsl_get_base_type(val->const_type)) {
+ switch (glsl_get_base_type(val->type->type)) {
case GLSL_TYPE_UINT:
case GLSL_TYPE_INT:
case GLSL_TYPE_UINT16:
@@ -1338,14 +1338,14 @@ vtn_handle_constant(struct vtn_builder *b, SpvOp opcode,
case GLSL_TYPE_FLOAT16:
case GLSL_TYPE_BOOL:
case GLSL_TYPE_DOUBLE: {
- int bit_size = glsl_get_bit_size(val->const_type);
- if (glsl_type_is_matrix(val->const_type)) {
- vtn_assert(glsl_get_matrix_columns(val->const_type) == elem_count);
+ int bit_size = glsl_get_bit_size(val->type->type);
+ if (glsl_type_is_matrix(val->type->type)) {
+ vtn_assert(glsl_get_matrix_columns(val->type->type) == elem_count);
for (unsigned i = 0; i < elem_count; i++)
val->constant->values[i] = elems[i]->values[0];
} else {
- vtn_assert(glsl_type_is_vector(val->const_type));
- vtn_assert(glsl_get_vector_elements(val->const_type) == elem_count);
+ vtn_assert(glsl_type_is_vector(val->type->type));
+ vtn_assert(glsl_get_vector_elements(val->type->type) == elem_count);
for (unsigned i = 0; i < elem_count; i++) {
switch (bit_size) {
case 64:
@@ -1390,22 +1390,14 @@ vtn_handle_constant(struct vtn_builder *b, SpvOp opcode,
vtn_assert(v1->value_type == vtn_value_type_constant ||
v1->value_type == vtn_value_type_undef);

- unsigned len0 = v0->value_type == vtn_value_type_constant ?
- glsl_get_vector_elements(v0->const_type) :
- glsl_get_vector_elements(v0->type->type);
- unsigned len1 = v1->value_type == vtn_value_type_constant ?
- glsl_get_vector_elements(v1->const_type) :
- glsl_get_vector_elements(v1->type->type);
+ 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->const_type);
- unsigned bit_size0 = v0->value_type == vtn_value_type_constant ?
- glsl_get_bit_size(v0->const_type) :
- glsl_get_bit_size(v0->type->type);
- unsigned bit_size1 = v1->value_type == vtn_value_type_constant ?
- glsl_get_bit_size(v1->const_type) :
- glsl_get_bit_size(v1->type->type);
+ 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;
@@ -1476,7 +1468,7 @@ vtn_handle_constant(struct vtn_builder *b, SpvOp opcode,

int elem = -1;
int col = 0;
- const struct glsl_type *type = comp->const_type;
+ const struct glsl_type *type = comp->type->type;
for (unsigned i = deref_start; i < count; i++) {
switch (glsl_get_base_type(type)) {
case GLSL_TYPE_UINT:
@@ -1541,7 +1533,7 @@ vtn_handle_constant(struct vtn_builder *b, SpvOp opcode,
} else {
struct vtn_value *insert =
vtn_value(b, w[4], vtn_value_type_constant);
- vtn_assert(insert->const_type == type);
+ vtn_assert(insert->type->type == type);
if (elem == -1) {
*c = insert->constant;
} else {
@@ -1568,15 +1560,14 @@ vtn_handle_constant(struct vtn_builder *b, SpvOp opcode,

default: {
bool swap;
- nir_alu_type dst_alu_type = nir_get_nir_type_for_glsl_type(val->const_type);
+ nir_alu_type dst_alu_type = nir_get_nir_type_for_glsl_type(val->type->type);
nir_alu_type src_alu_type = dst_alu_type;
nir_op op = vtn_nir_alu_op_for_spirv_opcode(b, opcode, &swap,
src_alu_type,
dst_alu_type);

- unsigned num_components = glsl_get_vector_elements(val->const_type);
- unsigned bit_size =
- glsl_get_bit_size(val->const_type);
+ unsigned num_components = glsl_get_vector_elements(val->type->type);
+ unsigned bit_size = glsl_get_bit_size(val->type->type);

nir_const_value src[4];
vtn_assert(count <= 7);
@@ -1598,7 +1589,7 @@ vtn_handle_constant(struct vtn_builder *b, SpvOp opcode,
}

case SpvOpConstantNull:
- val->constant = vtn_null_constant(b, val->const_type);
+ val->constant = vtn_null_constant(b, val->type->type);
break;

case SpvOpConstantSampler:
diff --git a/src/compiler/spirv/vtn_private.h b/src/compiler/spirv/vtn_private.h
index 5f140b4..6d4ad3c 100644
--- a/src/compiler/spirv/vtn_private.h
+++ b/src/compiler/spirv/vtn_private.h
@@ -480,14 +480,11 @@ struct vtn_value {
enum vtn_value_type value_type;
const char *name;
struct vtn_decoration *decoration;
+ struct vtn_type *type;
union {
void *ptr;
char *str;
- struct vtn_type *type;
- struct {
- nir_constant *constant;
- const struct glsl_type *const_type;
- };
+ nir_constant *constant;
struct vtn_pointer *pointer;
struct vtn_image_pointer *image;
struct vtn_sampled_image *sampled_image;

--
2.5.0.400.gff86faf

Samuel Iglesias Gonsálvez

2017-12-11 12:51:25 UTC

Permalink

Patches 1, 3, 4, 5, 6, 7 are,

Reviewed-by: Samuel Iglesias GonsÃ¡lvez <***@igalia.com>

Sam

Post by Jason Ekstrand
At the moment, this just lets us drop the const_type for constants and
unify things a bit. Eventually, we will use this to store the types of
all SPIR-V SSA values.
---
src/compiler/spirv/spirv_to_nir.c | 63 +++++++++++++++++----------------------
src/compiler/spirv/vtn_private.h | 7 ++---
2 files changed, 29 insertions(+), 41 deletions(-)
diff --git a/src/compiler/spirv/spirv_to_nir.c b/src/compiler/spirv/spirv_to_nir.c
index d321d1a..a50b14d 100644
--- a/src/compiler/spirv/spirv_to_nir.c
+++ b/src/compiler/spirv/spirv_to_nir.c
@@ -257,7 +257,7 @@ vtn_ssa_value(struct vtn_builder *b, uint32_t value_id)
return vtn_undef_ssa_value(b, val->type->type);
- return vtn_const_ssa_value(b, val->constant, val->const_type);
+ return vtn_const_ssa_value(b, val->constant, val->type->type);
return val->ssa;
@@ -1249,7 +1249,7 @@ handle_workgroup_size_decoration_cb(struct vtn_builder *b,
dec->literals[0] != SpvBuiltInWorkgroupSize)
return;
- vtn_assert(val->const_type == glsl_vector_type(GLSL_TYPE_UINT, 3));
+ vtn_assert(val->type->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[1] = val->constant->values[0].u32[1];
@@ -1261,21 +1261,21 @@ 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);
- val->const_type = vtn_value(b, w[1], vtn_value_type_type)->type->type;
+ val->type = vtn_value(b, w[1], vtn_value_type_type)->type;
val->constant = rzalloc(b, nir_constant);
switch (opcode) {
- vtn_assert(val->const_type == glsl_bool_type());
+ vtn_assert(val->type->type == glsl_bool_type());
val->constant->values[0].u32[0] = NIR_TRUE;
break;
- vtn_assert(val->const_type == glsl_bool_type());
+ vtn_assert(val->type->type == glsl_bool_type());
val->constant->values[0].u32[0] = NIR_FALSE;
break;
case SpvOpSpecConstantFalse: {
- vtn_assert(val->const_type == glsl_bool_type());
+ vtn_assert(val->type->type == glsl_bool_type());
uint32_t int_val =
get_specialization(b, val, (opcode == SpvOpSpecConstantTrue));
val->constant->values[0].u32[0] = int_val ? NIR_TRUE : NIR_FALSE;
@@ -1283,8 +1283,8 @@ vtn_handle_constant(struct vtn_builder *b, SpvOp opcode,
}
case SpvOpConstant: {
- vtn_assert(glsl_type_is_scalar(val->const_type));
- int bit_size = glsl_get_bit_size(val->const_type);
+ vtn_assert(glsl_type_is_scalar(val->type->type));
+ int bit_size = glsl_get_bit_size(val->type->type);
switch (bit_size) {
val->constant->values->u64[0] = vtn_u64_literal(&w[3]);
@@ -1301,9 +1301,9 @@ vtn_handle_constant(struct vtn_builder *b, SpvOp opcode,
break;
}
case SpvOpSpecConstant: {
- vtn_assert(glsl_type_is_scalar(val->const_type));
+ vtn_assert(glsl_type_is_scalar(val->type->type));
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->type->type);
switch (bit_size) {
val->constant->values[0].u64[0] =
@@ -1327,7 +1327,7 @@ vtn_handle_constant(struct vtn_builder *b, SpvOp opcode,
for (unsigned i = 0; i < elem_count; i++)
elems[i] = vtn_value(b, w[i + 3], vtn_value_type_constant)->constant;
- switch (glsl_get_base_type(val->const_type)) {
+ switch (glsl_get_base_type(val->type->type)) {
@@ -1338,14 +1338,14 @@ vtn_handle_constant(struct vtn_builder *b, SpvOp opcode,
case GLSL_TYPE_DOUBLE: {
- int bit_size = glsl_get_bit_size(val->const_type);
- if (glsl_type_is_matrix(val->const_type)) {
- vtn_assert(glsl_get_matrix_columns(val->const_type) == elem_count);
+ int bit_size = glsl_get_bit_size(val->type->type);
+ if (glsl_type_is_matrix(val->type->type)) {
+ vtn_assert(glsl_get_matrix_columns(val->type->type) == elem_count);
for (unsigned i = 0; i < elem_count; i++)
val->constant->values[i] = elems[i]->values[0];
} else {
- vtn_assert(glsl_type_is_vector(val->const_type));
- vtn_assert(glsl_get_vector_elements(val->const_type) == elem_count);
+ vtn_assert(glsl_type_is_vector(val->type->type));
+ vtn_assert(glsl_get_vector_elements(val->type->type) == elem_count);
for (unsigned i = 0; i < elem_count; i++) {
switch (bit_size) {
@@ -1390,22 +1390,14 @@ vtn_handle_constant(struct vtn_builder *b, SpvOp opcode,
vtn_assert(v1->value_type == vtn_value_type_constant ||
v1->value_type == vtn_value_type_undef);
- unsigned len0 = v0->value_type == vtn_value_type_constant ?
- glsl_get_vector_elements(v0->type->type);
- unsigned len1 = v1->value_type == vtn_value_type_constant ?
- glsl_get_vector_elements(v1->type->type);
+ 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->const_type);
- unsigned bit_size0 = v0->value_type == vtn_value_type_constant ?
- glsl_get_bit_size(v0->type->type);
- unsigned bit_size1 = v1->value_type == vtn_value_type_constant ?
- glsl_get_bit_size(v1->type->type);
+ 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;
@@ -1476,7 +1468,7 @@ vtn_handle_constant(struct vtn_builder *b, SpvOp opcode,
int elem = -1;
int col = 0;
- const struct glsl_type *type = comp->const_type;
+ const struct glsl_type *type = comp->type->type;
for (unsigned i = deref_start; i < count; i++) {
switch (glsl_get_base_type(type)) {
@@ -1541,7 +1533,7 @@ vtn_handle_constant(struct vtn_builder *b, SpvOp opcode,
} else {
struct vtn_value *insert =
vtn_value(b, w[4], vtn_value_type_constant);
- vtn_assert(insert->const_type == type);
+ vtn_assert(insert->type->type == type);
if (elem == -1) {
*c = insert->constant;
} else {
@@ -1568,15 +1560,14 @@ vtn_handle_constant(struct vtn_builder *b, SpvOp opcode,
default: {
bool swap;
- nir_alu_type dst_alu_type = nir_get_nir_type_for_glsl_type(val->const_type);
+ nir_alu_type dst_alu_type = nir_get_nir_type_for_glsl_type(val->type->type);
nir_alu_type src_alu_type = dst_alu_type;
nir_op op = vtn_nir_alu_op_for_spirv_opcode(b, opcode, &swap,
src_alu_type,
dst_alu_type);
- unsigned num_components = glsl_get_vector_elements(val->const_type);
- unsigned bit_size =
- glsl_get_bit_size(val->const_type);
+ unsigned num_components = glsl_get_vector_elements(val->type->type);
+ unsigned bit_size = glsl_get_bit_size(val->type->type);
nir_const_value src[4];
vtn_assert(count <= 7);
@@ -1598,7 +1589,7 @@ vtn_handle_constant(struct vtn_builder *b, SpvOp opcode,
}
- val->constant = vtn_null_constant(b, val->const_type);
+ val->constant = vtn_null_constant(b, val->type->type);
break;
diff --git a/src/compiler/spirv/vtn_private.h b/src/compiler/spirv/vtn_private.h
index 5f140b4..6d4ad3c 100644
--- a/src/compiler/spirv/vtn_private.h
+++ b/src/compiler/spirv/vtn_private.h
@@ -480,14 +480,11 @@ struct vtn_value {
enum vtn_value_type value_type;
const char *name;
struct vtn_decoration *decoration;
+ struct vtn_type *type;
union {
void *ptr;
char *str;
- struct vtn_type *type;
- struct {
- nir_constant *constant;
- const struct glsl_type *const_type;
- };
+ nir_constant *constant;
struct vtn_pointer *pointer;
struct vtn_image_pointer *image;
struct vtn_sampled_image *sampled_image;

Ian Romanick

2017-12-11 18:00:51 UTC

Permalink

This patch is

Jason Ekstrand

2017-12-07 16:12:09 UTC

Permalink

This autogenerated pass will automatically find and set the type field
on all vtn_values. This way we always have the type and can use it for
validation and other checks.
---
src/compiler/Makefile.nir.am | 4 +
src/compiler/nir/meson.build | 11 ++-
src/compiler/spirv/spirv_to_nir.c | 6 +-
src/compiler/spirv/vtn_gather_types_c.py | 125 +++++++++++++++++++++++++++++++
src/compiler/spirv/vtn_private.h | 4 +
5 files changed, 148 insertions(+), 2 deletions(-)
create mode 100644 src/compiler/spirv/vtn_gather_types_c.py

diff --git a/src/compiler/Makefile.nir.am b/src/compiler/Makefile.nir.am
index 1533ee5..dd38c45 100644
--- a/src/compiler/Makefile.nir.am
+++ b/src/compiler/Makefile.nir.am
@@ -56,6 +56,10 @@ spirv/spirv_info.c: spirv/spirv_info_c.py spirv/spirv.core.grammar.json
$(MKDIR_GEN)
$(PYTHON_GEN) $(srcdir)/spirv/spirv_info_c.py $(srcdir)/spirv/spirv.core.grammar.json $@ || ($(RM) $@; false)

+spirv/vtn_gather_types.c: spirv/vtn_gather_types_c.py spirv/spirv.core.grammar.json
+ $(MKDIR_GEN)
+ $(PYTHON_GEN) $(srcdir)/spirv/vtn_gather_types_c.py $(srcdir)/spirv/spirv.core.grammar.json $@ || ($(RM) $@; false)
+
noinst_PROGRAMS += spirv2nir

spirv2nir_SOURCES = \
diff --git a/src/compiler/nir/meson.build b/src/compiler/nir/meson.build
index b61a077..5dd21e6 100644
--- a/src/compiler/nir/meson.build
+++ b/src/compiler/nir/meson.build
@@ -72,6 +72,14 @@ spirv_info_c = custom_target(
command : [prog_python2, '@INPUT0@', '@INPUT1@', '@OUTPUT@'],
)

+vtn_gather_types_c = custom_target(
+ 'vtn_gather_types.c',
+ input : files('../spirv/vtn_gather_types_c.py',
+ '../spirv/spirv.core.grammar.json'),
+ output : 'vtn_gather_types.c',
+ command : [prog_python2, '@INPUT0@', '@INPUT1@', '@OUTPUT@'],
+)
+
files_libnir = files(
'nir.c',
'nir.h',
@@ -189,7 +197,8 @@ files_libnir = files(
libnir = static_library(
'nir',
[files_libnir, spirv_info_c, nir_opt_algebraic_c, nir_opcodes_c,
- nir_opcodes_h, nir_constant_expressions_c, nir_builder_opcodes_h],
+ nir_opcodes_h, nir_constant_expressions_c, nir_builder_opcodes_h,
+ vtn_gather_types_c],
include_directories : [inc_common, inc_compiler, include_directories('../spirv')],
c_args : [c_vis_args, c_msvc_compat_args, no_override_init_args],
link_with : libcompiler,
diff --git a/src/compiler/spirv/spirv_to_nir.c b/src/compiler/spirv/spirv_to_nir.c
index a50b14d..a2426bc 100644
--- a/src/compiler/spirv/spirv_to_nir.c
+++ b/src/compiler/spirv/spirv_to_nir.c
@@ -1261,7 +1261,6 @@ 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);
- val->type = vtn_value(b, w[1], vtn_value_type_type)->type;
val->constant = rzalloc(b, nir_constant);
switch (opcode) {
case SpvOpConstantTrue:
@@ -3268,6 +3267,8 @@ static bool
vtn_handle_variable_or_type_instruction(struct vtn_builder *b, SpvOp opcode,
const uint32_t *w, unsigned count)
{
+ vtn_set_instruction_result_type(b, opcode, w, count);
+
switch (opcode) {
case SpvOpSource:
case SpvOpSourceContinued:
@@ -3658,6 +3659,9 @@ spirv_to_nir(const uint32_t *words, size_t word_count,
words = vtn_foreach_instruction(b, words, word_end,
vtn_handle_variable_or_type_instruction);

+ /* Set types on all vtn_values */
+ vtn_foreach_instruction(b, words, word_end, vtn_set_instruction_result_type);
+
vtn_build_cfg(b, words, word_end);

assert(b->entry_point->value_type == vtn_value_type_function);
diff --git a/src/compiler/spirv/vtn_gather_types_c.py b/src/compiler/spirv/vtn_gather_types_c.py
new file mode 100644
index 0000000..8cd8d9f
--- /dev/null
+++ b/src/compiler/spirv/vtn_gather_types_c.py
@@ -0,0 +1,125 @@
+COPYRIGHT = """\
+/*
+ * Copyright (C) 2017 Intel Corporation
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a
+ * copy of this software and associated documentation files (the "Software"),
+ * to deal in the Software without restriction, including without limitation
+ * the rights to use, copy, modify, merge, publish, distribute, sublicense,
+ * and/or sell copies of the Software, and to permit persons to whom the
+ * Software is furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice (including the next
+ * paragraph) shall be included in all copies or substantial portions of the
+ * Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL
+ * THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING
+ * FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER
+ * DEALINGS IN THE SOFTWARE.
+ */
+"""
+
+import argparse
+import json
+from sys import stdout
+from mako.template import Template
+
+def find_result_types(spirv):
+ for inst in spirv['instructions']:
+ name = inst['opname']
+
+ if 'operands' not in inst:
+ continue
+
+ res_arg_idx = -1
+ res_type_arg_idx = -1
+ for idx, arg in enumerate(inst['operands']):
+ if arg['kind'] == 'IdResult':
+ res_arg_idx = idx
+ elif arg['kind'] == 'IdResultType':
+ res_type_arg_idx = idx
+
+ if res_type_arg_idx >= 0:
+ assert res_arg_idx >= 0
+ elif res_arg_idx >= 0:
+ untyped_insts = [
+ 'OpString',
+ 'OpExtInstImport',
+ 'OpDecorationGroup',
+ 'OpLabel',
+ ]
+ assert name.startswith('OpType') or name in untyped_insts
+
+ if res_arg_idx >= 0 or res_type_arg_idx >= 0:
+ yield (name, res_arg_idx, res_type_arg_idx)
+
+TEMPLATE = Template(COPYRIGHT + """\
+
+/* DO NOT EDIT - This file is generated automatically by the
+ * vtn_gather_types_c.py script
+ */
+
+#include "vtn_private.h"
+
+struct type_args {
+ int res_idx;
+ int res_type_idx;
+};
+
+static struct type_args
+result_type_args_for_opcode(SpvOp opcode)
+{
+ switch (opcode) {
+% for opcode in opcodes:
+ case Spv${opcode[0]}: return (struct type_args){ ${opcode[1]}, ${opcode[2]} };
+% endfor
+ default: return (struct type_args){ -1, -1 };
+ }
+}
+
+bool
+vtn_set_instruction_result_type(struct vtn_builder *b, SpvOp opcode,
+ const uint32_t *w, unsigned count)
+{
+ struct type_args args = result_type_args_for_opcode(opcode);
+
+ if (args.res_idx < 0 || args.res_type_idx < 0)
+ return true;
+
+ struct vtn_value *val = vtn_untyped_value(b, w[1 + args.res_idx]);
+ val->type = vtn_value(b, w[1 + args.res_type_idx],
+ vtn_value_type_type)->type;
+
+ return true;
+}
+
+""")
+
+if __name__ == "__main__":
+ p = argparse.ArgumentParser()
+ p.add_argument("json")
+ p.add_argument("out")
+ args = p.parse_args()
+
+ spirv_info = json.JSONDecoder().decode(open(args.json, "r").read())
+
+ opcodes = list(find_result_types(spirv_info))
+
+ try:
+ with open(args.out, 'w') as f:
+ f.write(TEMPLATE.render(opcodes=opcodes))
+ except Exception:
+ # In the even there's an error this imports some helpers from mako
+ # to print a useful stack trace and prints it, then exits with
+ # status 1, if python is run with debug; otherwise it just raises
+ # the exception
+ if __debug__:
+ import sys
+ from mako import exceptions
+ sys.stderr.write(exceptions.text_error_template().render() + '\n')
+ sys.exit(1)
+ raise
diff --git a/src/compiler/spirv/vtn_private.h b/src/compiler/spirv/vtn_private.h
index 6d4ad3c..a0a4f3a 100644
--- a/src/compiler/spirv/vtn_private.h
+++ b/src/compiler/spirv/vtn_private.h
@@ -620,6 +620,10 @@ vtn_value(struct vtn_builder *b, uint32_t value_id,
return val;
}

+bool
+vtn_set_instruction_result_type(struct vtn_builder *b, SpvOp opcode,
+ const uint32_t *w, unsigned count);
+
struct vtn_ssa_value *vtn_ssa_value(struct vtn_builder *b, uint32_t value_id);

struct vtn_ssa_value *vtn_create_ssa_value(struct vtn_builder *b,

--
2.5.0.400.gff86faf

Juan A. Suarez Romero

2017-12-11 16:19:28 UTC

Permalink

spirv/vtn_gather_types.c must be included also in
src/compiler/Makefile.sources, under the SPIRV_GENERATED_FILES.

Otherwise the file won't be generated.

J.A.

Post by Jason Ekstrand
noinst_PROGRAMS += spirv2nir
spirv2nir_SOURCES = \
diff --git a/src/compiler/nir/meson.build b/src/compiler/nir/meson.build
index b61a077..5dd21e6 100644
--- a/src/compiler/nir/meson.build
+++ b/src/compiler/nir/meson.build
@@ -72,6 +72,14 @@ spirv_info_c = custom_target(
)
+vtn_gather_types_c = custom_target(
+ 'vtn_gather_types.c',
+ input : files('../spirv/vtn_gather_types_c.py',
+ '../spirv/spirv.core.grammar.json'),
+ output : 'vtn_gather_types.c',
+)
+
files_libnir = files(
'nir.c',
'nir.h',
@@ -189,7 +197,8 @@ files_libnir = files(
libnir = static_library(
'nir',
[files_libnir, spirv_info_c, nir_opt_algebraic_c, nir_opcodes_c,
- nir_opcodes_h, nir_constant_expressions_c, nir_builder_opcodes_h],
+ nir_opcodes_h, nir_constant_expressions_c, nir_builder_opcodes_h,
+ vtn_gather_types_c],
include_directories : [inc_common, inc_compiler, include_directories('../spirv')],
c_args : [c_vis_args, c_msvc_compat_args, no_override_init_args],
link_with : libcompiler,
diff --git a/src/compiler/spirv/spirv_to_nir.c b/src/compiler/spirv/spirv_to_nir.c
index a50b14d..a2426bc 100644
--- a/src/compiler/spirv/spirv_to_nir.c
+++ b/src/compiler/spirv/spirv_to_nir.c
@@ -1261,7 +1261,6 @@ 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);
- val->type = vtn_value(b, w[1], vtn_value_type_type)->type;
val->constant = rzalloc(b, nir_constant);
switch (opcode) {
@@ -3268,6 +3267,8 @@ static bool
vtn_handle_variable_or_type_instruction(struct vtn_builder *b, SpvOp opcode,
const uint32_t *w, unsigned count)
{
+ vtn_set_instruction_result_type(b, opcode, w, count);
+
switch (opcode) {
@@ -3658,6 +3659,9 @@ spirv_to_nir(const uint32_t *words, size_t word_count,
words = vtn_foreach_instruction(b, words, word_end,
vtn_handle_variable_or_type_instruction);
+ /* Set types on all vtn_values */
+ vtn_foreach_instruction(b, words, word_end, vtn_set_instruction_result_type);
+
vtn_build_cfg(b, words, word_end);
assert(b->entry_point->value_type == vtn_value_type_function);
diff --git a/src/compiler/spirv/vtn_gather_types_c.py b/src/compiler/spirv/vtn_gather_types_c.py
new file mode 100644
index 0000000..8cd8d9f
--- /dev/null
+++ b/src/compiler/spirv/vtn_gather_types_c.py
@@ -0,0 +1,125 @@
+COPYRIGHT = """\
+/*
+ * Copyright (C) 2017 Intel Corporation
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a
+ * copy of this software and associated documentation files (the "Software"),
+ * to deal in the Software without restriction, including without limitation
+ * the rights to use, copy, modify, merge, publish, distribute, sublicense,
+ * and/or sell copies of the Software, and to permit persons to whom the
+ *
+ * The above copyright notice and this permission notice (including the next
+ * paragraph) shall be included in all copies or substantial portions of the
+ * Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL
+ * THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING
+ * FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER
+ * DEALINGS IN THE SOFTWARE.
+ */
+"""
+
+import argparse
+import json
+from sys import stdout
+from mako.template import Template
+
+ name = inst['opname']
+
+ continue
+
+ res_arg_idx = -1
+ res_type_arg_idx = -1
+ res_arg_idx = idx
+ res_type_arg_idx = idx
+
+ assert res_arg_idx >= 0
+ untyped_insts = [
+ 'OpString',
+ 'OpExtInstImport',
+ 'OpDecorationGroup',
+ 'OpLabel',
+ ]
+ assert name.startswith('OpType') or name in untyped_insts
+
+ yield (name, res_arg_idx, res_type_arg_idx)
+
+TEMPLATE = Template(COPYRIGHT + """\
+
+/* DO NOT EDIT - This file is generated automatically by the
+ * vtn_gather_types_c.py script
+ */
+
+#include "vtn_private.h"
+
+struct type_args {
+ int res_idx;
+ int res_type_idx;
+};
+
+static struct type_args
+result_type_args_for_opcode(SpvOp opcode)
+{
+ switch (opcode) {
+ case Spv${opcode[0]}: return (struct type_args){ ${opcode[1]}, ${opcode[2]} };
+% endfor
+ default: return (struct type_args){ -1, -1 };
+ }
+}
+
+bool
+vtn_set_instruction_result_type(struct vtn_builder *b, SpvOp opcode,
+ const uint32_t *w, unsigned count)
+{
+ struct type_args args = result_type_args_for_opcode(opcode);
+
+ if (args.res_idx < 0 || args.res_type_idx < 0)
+ return true;
+
+ struct vtn_value *val = vtn_untyped_value(b, w[1 + args.res_idx]);
+ val->type = vtn_value(b, w[1 + args.res_type_idx],
+ vtn_value_type_type)->type;
+
+ return true;
+}
+
+""")
+
+ p = argparse.ArgumentParser()
+ p.add_argument("json")
+ p.add_argument("out")
+ args = p.parse_args()
+
+ spirv_info = json.JSONDecoder().decode(open(args.json, "r").read())
+
+ opcodes = list(find_result_types(spirv_info))
+
+ f.write(TEMPLATE.render(opcodes=opcodes))
+ # In the even there's an error this imports some helpers from mako
+ # to print a useful stack trace and prints it, then exits with
+ # status 1, if python is run with debug; otherwise it just raises
+ # the exception
+ import sys
+ from mako import exceptions
+ sys.stderr.write(exceptions.text_error_template().render() + '\n')
+ sys.exit(1)
+ raise
diff --git a/src/compiler/spirv/vtn_private.h b/src/compiler/spirv/vtn_private.h
index 6d4ad3c..a0a4f3a 100644
--- a/src/compiler/spirv/vtn_private.h
+++ b/src/compiler/spirv/vtn_private.h
@@ -620,6 +620,10 @@ vtn_value(struct vtn_builder *b, uint32_t value_id,
return val;
}
+bool
+vtn_set_instruction_result_type(struct vtn_builder *b, SpvOp opcode,
+ const uint32_t *w, unsigned count);
+
struct vtn_ssa_value *vtn_ssa_value(struct vtn_builder *b, uint32_t value_id);
struct vtn_ssa_value *vtn_create_ssa_value(struct vtn_builder *b,

Jason Ekstrand

2017-12-11 16:29:03 UTC

Permalink

Post by Jason Ekstrand

+++++++++++++++++++++++++++++++

Post by Jason Ekstrand
src/compiler/spirv/vtn_private.h | 4 +
5 files changed, 148 insertions(+), 2 deletions(-)
create mode 100644 src/compiler/spirv/vtn_gather_types_c.py
diff --git a/src/compiler/Makefile.nir.am b/src/compiler/Makefile.nir.am
index 1533ee5..dd38c45 100644
--- a/src/compiler/Makefile.nir.am
+++ b/src/compiler/Makefile.nir.am
@@ -56,6 +56,10 @@ spirv/spirv_info.c: spirv/spirv_info_c.py

spirv/spirv.core.grammar.json

Post by Jason Ekstrand
$(MKDIR_GEN)
$(PYTHON_GEN) $(srcdir)/spirv/spirv_info_c.py
+spirv/vtn_gather_types.c: spirv/vtn_gather_types_c.py

spirv/spirv.core.grammar.json

Post by Jason Ekstrand
+ $(MKDIR_GEN)
+ $(PYTHON_GEN) $(srcdir)/spirv/vtn_gather_types_c.py
+

spirv/vtn_gather_types.c must be included also in
src/compiler/Makefile.sources, under the SPIRV_GENERATED_FILES.

Thanks! Clearly, I don't use autotools anymore. :-)

Post by Jason Ekstrand
Otherwise the file won't be generated.
J.A.

include_directories('../spirv')],

Post by Jason Ekstrand
c_args : [c_vis_args, c_msvc_compat_args, no_override_init_args],
link_with : libcompiler,
diff --git a/src/compiler/spirv/spirv_to_nir.c

b/src/compiler/spirv/spirv_to_nir.c

Post by Jason Ekstrand
index a50b14d..a2426bc 100644
--- a/src/compiler/spirv/spirv_to_nir.c
+++ b/src/compiler/spirv/spirv_to_nir.c
@@ -1261,7 +1261,6 @@ vtn_handle_constant(struct vtn_builder *b, SpvOp

opcode,

Post by Jason Ekstrand
const uint32_t *w, unsigned count)
{
struct vtn_value *val = vtn_push_value(b, w[2],

vtn_value_type_constant);

Post by Jason Ekstrand
- val->type = vtn_value(b, w[1], vtn_value_type_type)->type;
val->constant = rzalloc(b, nir_constant);
switch (opcode) {
@@ -3268,6 +3267,8 @@ static bool
vtn_handle_variable_or_type_instruction(struct vtn_builder *b, SpvOp

opcode,

Post by Jason Ekstrand
const uint32_t *w, unsigned

count)

Post by Jason Ekstrand
{
+ vtn_set_instruction_result_type(b, opcode, w, count);
+
switch (opcode) {
@@ -3658,6 +3659,9 @@ spirv_to_nir(const uint32_t *words, size_t

word_count,

Post by Jason Ekstrand
words = vtn_foreach_instruction(b, words, word_end,
vtn_handle_variable_or_type_

instruction);

Post by Jason Ekstrand
+ /* Set types on all vtn_values */
+ vtn_foreach_instruction(b, words, word_end,

vtn_set_instruction_result_type);

Post by Jason Ekstrand
+
vtn_build_cfg(b, words, word_end);
assert(b->entry_point->value_type == vtn_value_type_function);
diff --git a/src/compiler/spirv/vtn_gather_types_c.py

b/src/compiler/spirv/vtn_gather_types_c.py

Post by Jason Ekstrand
new file mode 100644
index 0000000..8cd8d9f
--- /dev/null
+++ b/src/compiler/spirv/vtn_gather_types_c.py
@@ -0,0 +1,125 @@
+COPYRIGHT = """\
+/*
+ * Copyright (C) 2017 Intel Corporation
+ *
+ * Permission is hereby granted, free of charge, to any person

obtaining a

Post by Jason Ekstrand
+ * copy of this software and associated documentation files (the

"Software"),

Post by Jason Ekstrand
+ * to deal in the Software without restriction, including without

limitation

Post by Jason Ekstrand
+ * the rights to use, copy, modify, merge, publish, distribute,

sublicense,

Post by Jason Ekstrand
+ * and/or sell copies of the Software, and to permit persons to whom the
+ *
+ * The above copyright notice and this permission notice (including the

Post by Jason Ekstrand
+ * paragraph) shall be included in all copies or substantial portions

of the

Post by Jason Ekstrand
+ * Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND,

EXPRESS OR

Post by Jason Ekstrand
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF

MERCHANTABILITY,

Post by Jason Ekstrand
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT

SHALL

Post by Jason Ekstrand
+ * THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR

OTHER

Post by Jason Ekstrand
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE,

ARISING

Post by Jason Ekstrand
+ * FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER
+ * DEALINGS IN THE SOFTWARE.
+ */
+"""
+
+import argparse
+import json
+from sys import stdout
+from mako.template import Template
+
+ name = inst['opname']
+
+ continue
+
+ res_arg_idx = -1
+ res_type_arg_idx = -1
+ res_arg_idx = idx
+ res_type_arg_idx = idx
+
+ assert res_arg_idx >= 0
+ untyped_insts = [
+ 'OpString',
+ 'OpExtInstImport',
+ 'OpDecorationGroup',
+ 'OpLabel',
+ ]
+ assert name.startswith('OpType') or name in untyped_insts
+
+ yield (name, res_arg_idx, res_type_arg_idx)
+
+TEMPLATE = Template(COPYRIGHT + """\
+
+/* DO NOT EDIT - This file is generated automatically by the
+ * vtn_gather_types_c.py script
+ */
+
+#include "vtn_private.h"
+
+struct type_args {
+ int res_idx;
+ int res_type_idx;
+};
+
+static struct type_args
+result_type_args_for_opcode(SpvOp opcode)
+{
+ switch (opcode) {
+ case Spv${opcode[0]}: return (struct type_args){ ${opcode[1]},

${opcode[2]} };

Post by Jason Ekstrand
+% endfor
+ default: return (struct type_args){ -1, -1 };
+ }
+}
+
+bool
+vtn_set_instruction_result_type(struct vtn_builder *b, SpvOp opcode,
+ const uint32_t *w, unsigned count)
+{
+ struct type_args args = result_type_args_for_opcode(opcode);
+
+ if (args.res_idx < 0 || args.res_type_idx < 0)
+ return true;
+
+ struct vtn_value *val = vtn_untyped_value(b, w[1 + args.res_idx]);
+ val->type = vtn_value(b, w[1 + args.res_type_idx],
+ vtn_value_type_type)->type;
+
+ return true;
+}
+
+""")
+
+ p = argparse.ArgumentParser()
+ p.add_argument("json")
+ p.add_argument("out")
+ args = p.parse_args()
+
+ spirv_info = json.JSONDecoder().decode(open(args.json, "r").read())
+
+ opcodes = list(find_result_types(spirv_info))
+
+ f.write(TEMPLATE.render(opcodes=opcodes))
+ # In the even there's an error this imports some helpers from

mako

Post by Jason Ekstrand
+ # to print a useful stack trace and prints it, then exits with
+ # status 1, if python is run with debug; otherwise it just

raises

Post by Jason Ekstrand
+ # the exception
+ import sys
+ from mako import exceptions
+ sys.stderr.write(exceptions.text_error_template().render()

+ '\n')

Post by Jason Ekstrand
+ sys.exit(1)
+ raise
diff --git a/src/compiler/spirv/vtn_private.h b/src/compiler/spirv/vtn_

private.h

Post by Jason Ekstrand
index 6d4ad3c..a0a4f3a 100644
--- a/src/compiler/spirv/vtn_private.h
+++ b/src/compiler/spirv/vtn_private.h
@@ -620,6 +620,10 @@ vtn_value(struct vtn_builder *b, uint32_t value_id,
return val;
}
+bool
+vtn_set_instruction_result_type(struct vtn_builder *b, SpvOp opcode,
+ const uint32_t *w, unsigned count);
+
struct vtn_ssa_value *vtn_ssa_value(struct vtn_builder *b, uint32_t

value_id);

Post by Jason Ekstrand
struct vtn_ssa_value *vtn_create_ssa_value(struct vtn_builder *b,

_______________________________________________
mesa-dev mailing list
https://lists.freedesktop.org/mailman/listinfo/mesa-dev

Ian Romanick

2017-12-11 17:50:39 UTC

Permalink

Post by Jason Ekstrand
This autogenerated pass will automatically find and set the type field
on all vtn_values. This way we always have the type and can use it for
validation and other checks.
---
src/compiler/Makefile.nir.am | 4 +
src/compiler/nir/meson.build | 11 ++-
src/compiler/spirv/spirv_to_nir.c | 6 +-
src/compiler/spirv/vtn_gather_types_c.py | 125 +++++++++++++++++++++++++++++++
src/compiler/spirv/vtn_private.h | 4 +
5 files changed, 148 insertions(+), 2 deletions(-)
create mode 100644 src/compiler/spirv/vtn_gather_types_c.py
diff --git a/src/compiler/Makefile.nir.am b/src/compiler/Makefile.nir.am
index 1533ee5..dd38c45 100644
--- a/src/compiler/Makefile.nir.am
+++ b/src/compiler/Makefile.nir.am
@@ -56,6 +56,10 @@ spirv/spirv_info.c: spirv/spirv_info_c.py spirv/spirv.core.grammar.json
$(MKDIR_GEN)
+spirv/vtn_gather_types.c: spirv/vtn_gather_types_c.py spirv/spirv.core.grammar.json
+ $(MKDIR_GEN)
+
noinst_PROGRAMS += spirv2nir
spirv2nir_SOURCES = \
diff --git a/src/compiler/nir/meson.build b/src/compiler/nir/meson.build
index b61a077..5dd21e6 100644
--- a/src/compiler/nir/meson.build
+++ b/src/compiler/nir/meson.build
@@ -72,6 +72,14 @@ spirv_info_c = custom_target(
)
+vtn_gather_types_c = custom_target(
+ 'vtn_gather_types.c',
+ input : files('../spirv/vtn_gather_types_c.py',
+ '../spirv/spirv.core.grammar.json'),
+ output : 'vtn_gather_types.c',
+)
+
files_libnir = files(
'nir.c',
'nir.h',
@@ -189,7 +197,8 @@ files_libnir = files(
libnir = static_library(
'nir',
[files_libnir, spirv_info_c, nir_opt_algebraic_c, nir_opcodes_c,
- nir_opcodes_h, nir_constant_expressions_c, nir_builder_opcodes_h],
+ nir_opcodes_h, nir_constant_expressions_c, nir_builder_opcodes_h,
+ vtn_gather_types_c],
include_directories : [inc_common, inc_compiler, include_directories('../spirv')],
c_args : [c_vis_args, c_msvc_compat_args, no_override_init_args],
link_with : libcompiler,
diff --git a/src/compiler/spirv/spirv_to_nir.c b/src/compiler/spirv/spirv_to_nir.c
index a50b14d..a2426bc 100644
--- a/src/compiler/spirv/spirv_to_nir.c
+++ b/src/compiler/spirv/spirv_to_nir.c
@@ -1261,7 +1261,6 @@ 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);
- val->type = vtn_value(b, w[1], vtn_value_type_type)->type;
val->constant = rzalloc(b, nir_constant);
switch (opcode) {
@@ -3268,6 +3267,8 @@ static bool
vtn_handle_variable_or_type_instruction(struct vtn_builder *b, SpvOp opcode,
const uint32_t *w, unsigned count)
{
+ vtn_set_instruction_result_type(b, opcode, w, count);
+
switch (opcode) {
@@ -3658,6 +3659,9 @@ spirv_to_nir(const uint32_t *words, size_t word_count,
words = vtn_foreach_instruction(b, words, word_end,
vtn_handle_variable_or_type_instruction);
+ /* Set types on all vtn_values */
+ vtn_foreach_instruction(b, words, word_end, vtn_set_instruction_result_type);
+
vtn_build_cfg(b, words, word_end);
assert(b->entry_point->value_type == vtn_value_type_function);
diff --git a/src/compiler/spirv/vtn_gather_types_c.py b/src/compiler/spirv/vtn_gather_types_c.py
new file mode 100644
index 0000000..8cd8d9f
--- /dev/null
+++ b/src/compiler/spirv/vtn_gather_types_c.py
@@ -0,0 +1,125 @@
+COPYRIGHT = """\
+/*
+ * Copyright (C) 2017 Intel Corporation
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a
+ * copy of this software and associated documentation files (the "Software"),
+ * to deal in the Software without restriction, including without limitation
+ * the rights to use, copy, modify, merge, publish, distribute, sublicense,
+ * and/or sell copies of the Software, and to permit persons to whom the
+ *
+ * The above copyright notice and this permission notice (including the next
+ * paragraph) shall be included in all copies or substantial portions of the
+ * Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL
+ * THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING
+ * FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER
+ * DEALINGS IN THE SOFTWARE.
+ */
+"""
+
+import argparse
+import json
+from sys import stdout
+from mako.template import Template
+
+ name = inst['opname']
+
+ continue
+
+ res_arg_idx = -1
+ res_type_arg_idx = -1
+ res_arg_idx = idx
+ res_type_arg_idx = idx
+
+ assert res_arg_idx >= 0
+ untyped_insts = [
+ 'OpString',
+ 'OpExtInstImport',
+ 'OpDecorationGroup',
+ 'OpLabel',
+ ]
+ assert name.startswith('OpType') or name in untyped_insts
+
+ yield (name, res_arg_idx, res_type_arg_idx)
+
+TEMPLATE = Template(COPYRIGHT + """\
+
+/* DO NOT EDIT - This file is generated automatically by the
+ * vtn_gather_types_c.py script
+ */
+
+#include "vtn_private.h"
+
+struct type_args {
+ int res_idx;
+ int res_type_idx;
+};
+
+static struct type_args
+result_type_args_for_opcode(SpvOp opcode)
+{
+ switch (opcode) {
+ case Spv${opcode[0]}: return (struct type_args){ ${opcode[1]}, ${opcode[2]} };
+% endfor
+ default: return (struct type_args){ -1, -1 };
+ }
+}
+
+bool
+vtn_set_instruction_result_type(struct vtn_builder *b, SpvOp opcode,
+ const uint32_t *w, unsigned count)
+{
+ struct type_args args = result_type_args_for_opcode(opcode);
+
+ if (args.res_idx < 0 || args.res_type_idx < 0)
+ return true;
+
+ struct vtn_value *val = vtn_untyped_value(b, w[1 + args.res_idx]);
+ val->type = vtn_value(b, w[1 + args.res_type_idx],
+ vtn_value_type_type)->type;
+
+ return true;

It seems like one of the returns in this function should be false. The
first one?

Post by Jason Ekstrand
+}
+
+""")
+
+ p = argparse.ArgumentParser()
+ p.add_argument("json")
+ p.add_argument("out")
+ args = p.parse_args()
+
+ spirv_info = json.JSONDecoder().decode(open(args.json, "r").read())
+
+ opcodes = list(find_result_types(spirv_info))
+
+ f.write(TEMPLATE.render(opcodes=opcodes))
+ # In the even there's an error this imports some helpers from mako
+ # to print a useful stack trace and prints it, then exits with
+ # status 1, if python is run with debug; otherwise it just raises
+ # the exception
+ import sys
+ from mako import exceptions
+ sys.stderr.write(exceptions.text_error_template().render() + '\n')
+ sys.exit(1)
+ raise
diff --git a/src/compiler/spirv/vtn_private.h b/src/compiler/spirv/vtn_private.h
index 6d4ad3c..a0a4f3a 100644
--- a/src/compiler/spirv/vtn_private.h
+++ b/src/compiler/spirv/vtn_private.h
@@ -620,6 +620,10 @@ vtn_value(struct vtn_builder *b, uint32_t value_id,
return val;
}
+bool
+vtn_set_instruction_result_type(struct vtn_builder *b, SpvOp opcode,
+ const uint32_t *w, unsigned count);
+
struct vtn_ssa_value *vtn_ssa_value(struct vtn_builder *b, uint32_t value_id);
struct vtn_ssa_value *vtn_create_ssa_value(struct vtn_builder *b,

Jason Ekstrand

2017-12-11 18:50:38 UTC

Permalink

Post by Jason Ekstrand

+++++++++++++++++++++++++++++++

spirv/spirv.core.grammar.json

Post by Jason Ekstrand
$(MKDIR_GEN)
$(PYTHON_GEN) $(srcdir)/spirv/spirv_info_c.py
+spirv/vtn_gather_types.c: spirv/vtn_gather_types_c.py

spirv/spirv.core.grammar.json

Post by Jason Ekstrand
+ $(MKDIR_GEN)
+ $(PYTHON_GEN) $(srcdir)/spirv/vtn_gather_types_c.py
+
noinst_PROGRAMS += spirv2nir
spirv2nir_SOURCES = \
diff --git a/src/compiler/nir/meson.build b/src/compiler/nir/meson.build
index b61a077..5dd21e6 100644
--- a/src/compiler/nir/meson.build
+++ b/src/compiler/nir/meson.build
@@ -72,6 +72,14 @@ spirv_info_c = custom_target(
)
+vtn_gather_types_c = custom_target(
+ 'vtn_gather_types.c',
+ input : files('../spirv/vtn_gather_types_c.py',
+ '../spirv/spirv.core.grammar.json'),
+ output : 'vtn_gather_types.c',
+)
+
files_libnir = files(
'nir.c',
'nir.h',
@@ -189,7 +197,8 @@ files_libnir = files(
libnir = static_library(
'nir',
[files_libnir, spirv_info_c, nir_opt_algebraic_c, nir_opcodes_c,
- nir_opcodes_h, nir_constant_expressions_c, nir_builder_opcodes_h],
+ nir_opcodes_h, nir_constant_expressions_c, nir_builder_opcodes_h,
+ vtn_gather_types_c],
include_directories : [inc_common, inc_compiler,

include_directories('../spirv')],

Post by Jason Ekstrand
c_args : [c_vis_args, c_msvc_compat_args, no_override_init_args],
link_with : libcompiler,
diff --git a/src/compiler/spirv/spirv_to_nir.c

b/src/compiler/spirv/spirv_to_nir.c

opcode,

Post by Jason Ekstrand
const uint32_t *w, unsigned count)
{
struct vtn_value *val = vtn_push_value(b, w[2],

vtn_value_type_constant);

opcode,

Post by Jason Ekstrand
const uint32_t *w, unsigned

count)

Post by Jason Ekstrand
{
+ vtn_set_instruction_result_type(b, opcode, w, count);
+
switch (opcode) {
@@ -3658,6 +3659,9 @@ spirv_to_nir(const uint32_t *words, size_t

word_count,

Post by Jason Ekstrand
words = vtn_foreach_instruction(b, words, word_end,
vtn_handle_variable_or_type_

instruction);

Post by Jason Ekstrand
+ /* Set types on all vtn_values */
+ vtn_foreach_instruction(b, words, word_end,

vtn_set_instruction_result_type);

Post by Jason Ekstrand
+
vtn_build_cfg(b, words, word_end);
assert(b->entry_point->value_type == vtn_value_type_function);
diff --git a/src/compiler/spirv/vtn_gather_types_c.py

b/src/compiler/spirv/vtn_gather_types_c.py

obtaining a

Post by Jason Ekstrand
+ * copy of this software and associated documentation files (the

"Software"),

Post by Jason Ekstrand
+ * to deal in the Software without restriction, including without

limitation

Post by Jason Ekstrand
+ * the rights to use, copy, modify, merge, publish, distribute,

sublicense,

Post by Jason Ekstrand
+ * and/or sell copies of the Software, and to permit persons to whom the
+ *
+ * The above copyright notice and this permission notice (including the

Post by Jason Ekstrand
+ * paragraph) shall be included in all copies or substantial portions

of the

Post by Jason Ekstrand
+ * Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND,

EXPRESS OR

Post by Jason Ekstrand
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF

MERCHANTABILITY,

Post by Jason Ekstrand
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT

SHALL

Post by Jason Ekstrand
+ * THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR

OTHER

Post by Jason Ekstrand
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE,

ARISING

${opcode[2]} };

It seems like one of the returns in this function should be false. The
first one?

No, they should both return true. This function uses the callback
interface for vtn_foreach_instruction where returning false means "stop
iterating". We want to continue on even if we see an instruction which has
no result type. I suppose we could at some point replace that bool with an
enum if that would be more clear.

Post by Jason Ekstrand

mako

Post by Jason Ekstrand
+ # to print a useful stack trace and prints it, then exits with
+ # status 1, if python is run with debug; otherwise it just

raises

Post by Jason Ekstrand
+ # the exception
+ import sys
+ from mako import exceptions
+ sys.stderr.write(exceptions.text_error_template().render()

+ '\n')

Post by Jason Ekstrand
+ sys.exit(1)
+ raise
diff --git a/src/compiler/spirv/vtn_private.h b/src/compiler/spirv/vtn_

private.h

value_id);

Post by Jason Ekstrand
struct vtn_ssa_value *vtn_create_ssa_value(struct vtn_builder *b,

Ian Romanick

2017-12-11 19:13:26 UTC

Permalink

Post by Jason Ekstrand

Post by Jason Ekstrand
This autogenerated pass will automatically find and set the type field
on all vtn_values. This way we always have the type and can use

it for

Post by Jason Ekstrand
validation and other checks.
---
src/compiler/Makefile.nir.am <http://Makefile.nir.am>

| 4 +

Post by Jason Ekstrand
src/compiler/nir/meson.build | 11 ++-
src/compiler/spirv/spirv_to_nir.c | 6 +-
src/compiler/spirv/vtn_gather_types_c.py | 125

+++++++++++++++++++++++++++++++

b/src/compiler/Makefile.nir.am <http://Makefile.nir.am>

Post by Jason Ekstrand
index 1533ee5..dd38c45 100644
--- a/src/compiler/Makefile.nir.am <http://Makefile.nir.am>
+++ b/src/compiler/Makefile.nir.am <http://Makefile.nir.am>
@@ -56,6 +56,10 @@ spirv/spirv_info.c: spirv/spirv_info_c.py

spirv/spirv.core.grammar.json

Post by Jason Ekstrand
$(MKDIR_GEN)
$(PYTHON_GEN) $(srcdir)/spirv/spirv_info_c.py
+spirv/vtn_gather_types.c: spirv/vtn_gather_types_c.py

spirv/spirv.core.grammar.json

Post by Jason Ekstrand
+ $(MKDIR_GEN)
+ $(PYTHON_GEN) $(srcdir)/spirv/vtn_gather_types_c.py
+
noinst_PROGRAMS += spirv2nir
spirv2nir_SOURCES = \
diff --git a/src/compiler/nir/meson.build

b/src/compiler/nir/meson.build

Post by Jason Ekstrand
index b61a077..5dd21e6 100644
--- a/src/compiler/nir/meson.build
+++ b/src/compiler/nir/meson.build
@@ -72,6 +72,14 @@ spirv_info_c = custom_target(
)
+vtn_gather_types_c = custom_target(
+ 'vtn_gather_types.c',
+ input : files('../spirv/vtn_gather_types_c.py',
+ '../spirv/spirv.core.grammar.json'),
+ output : 'vtn_gather_types.c',
+)
+
files_libnir = files(
'nir.c',
'nir.h',
@@ -189,7 +197,8 @@ files_libnir = files(
libnir = static_library(
'nir',
[files_libnir, spirv_info_c, nir_opt_algebraic_c, nir_opcodes_c,
- nir_opcodes_h, nir_constant_expressions_c, nir_builder_opcodes_h],
+ nir_opcodes_h, nir_constant_expressions_c, nir_builder_opcodes_h,
+ vtn_gather_types_c],
include_directories : [inc_common, inc_compiler,

include_directories('../spirv')],

Post by Jason Ekstrand
c_args : [c_vis_args, c_msvc_compat_args, no_override_init_args],
link_with : libcompiler,
diff --git a/src/compiler/spirv/spirv_to_nir.c

b/src/compiler/spirv/spirv_to_nir.c

Post by Jason Ekstrand
index a50b14d..a2426bc 100644
--- a/src/compiler/spirv/spirv_to_nir.c
+++ b/src/compiler/spirv/spirv_to_nir.c
@@ -1261,7 +1261,6 @@ vtn_handle_constant(struct vtn_builder *b,

SpvOp opcode,

Post by Jason Ekstrand
const uint32_t *w, unsigned count)
{
struct vtn_value *val = vtn_push_value(b, w[2],

vtn_value_type_constant);

SpvOp opcode,

Post by Jason Ekstrand
const uint32_t *w,

unsigned count)

Post by Jason Ekstrand
{
+ vtn_set_instruction_result_type(b, opcode, w, count);
+
switch (opcode) {
@@ -3658,6 +3659,9 @@ spirv_to_nir(const uint32_t *words, size_t

word_count,

Post by Jason Ekstrand
words = vtn_foreach_instruction(b, words, word_end,

vtn_handle_variable_or_type_instruction);

Post by Jason Ekstrand
+ /* Set types on all vtn_values */
+ vtn_foreach_instruction(b, words, word_end,

vtn_set_instruction_result_type);

Post by Jason Ekstrand
+
vtn_build_cfg(b, words, word_end);
assert(b->entry_point->value_type == vtn_value_type_function);
diff --git a/src/compiler/spirv/vtn_gather_types_c.py

b/src/compiler/spirv/vtn_gather_types_c.py

obtaining a

Post by Jason Ekstrand
+ * copy of this software and associated documentation files (the

"Software"),

Post by Jason Ekstrand
+ * to deal in the Software without restriction, including without

limitation

Post by Jason Ekstrand
+ * the rights to use, copy, modify, merge, publish, distribute,

sublicense,

Post by Jason Ekstrand
+ * and/or sell copies of the Software, and to permit persons to

whom the

Post by Jason Ekstrand
+ * Software is furnished to do so, subject to the following
+ *
+ * The above copyright notice and this permission notice

(including the next

Post by Jason Ekstrand
+ * paragraph) shall be included in all copies or substantial

portions of the

Post by Jason Ekstrand
+ * Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY

KIND, EXPRESS OR

Post by Jason Ekstrand
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF

MERCHANTABILITY,

Post by Jason Ekstrand
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO

EVENT SHALL

Post by Jason Ekstrand
+ * THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM,

DAMAGES OR OTHER

Post by Jason Ekstrand
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR

OTHERWISE, ARISING

Post by Jason Ekstrand
+ * FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR

OTHER

Post by Jason Ekstrand
+ * DEALINGS IN THE SOFTWARE.
+ */
+"""
+
+import argparse
+import json
+from sys import stdout
+from mako.template import Template
+
+ name = inst['opname']
+
+ continue
+
+ res_arg_idx = -1
+ res_type_arg_idx = -1
+ res_arg_idx = idx
+ res_type_arg_idx = idx
+
+ assert res_arg_idx >= 0
+ untyped_insts = [
+ 'OpString',
+ 'OpExtInstImport',
+ 'OpDecorationGroup',
+ 'OpLabel',
+ ]
+ assert name.startswith('OpType') or name in untyped_insts
+
+ yield (name, res_arg_idx, res_type_arg_idx)
+
+TEMPLATE = Template(COPYRIGHT + """\
+
+/* DO NOT EDIT - This file is generated automatically by the
+ * vtn_gather_types_c.py script
+ */
+
+#include "vtn_private.h"
+
+struct type_args {
+ int res_idx;
+ int res_type_idx;
+};
+
+static struct type_args
+result_type_args_for_opcode(SpvOp opcode)
+{
+ switch (opcode) {
+ case Spv${opcode[0]}: return (struct type_args){ ${opcode[1]},

${opcode[2]} };

It seems like one of the returns in this function should be false. The
first one?
No, they should both return true. This function uses the callback
interface for vtn_foreach_instruction where returning false means "stop
iterating". We want to continue on even if we see an instruction which
has no result type. I suppose we could at some point replace that bool
with an enum if that would be more clear.

Ah. For this function, if you restructure it as

if (args.res_index >= 0 && args.res_type_idx >= 0) {
...
}

return true;

the next person to look at the code won't think it's weird. With that
or a comment describing why all returns are true, this patch is

Post by Jason Ekstrand

"r").read())

Post by Jason Ekstrand
+
+ opcodes = list(find_result_types(spirv_info))
+
+ f.write(TEMPLATE.render(opcodes=opcodes))
+ # In the even there's an error this imports some helpers

from mako

Post by Jason Ekstrand
+ # to print a useful stack trace and prints it, then exits

with

Post by Jason Ekstrand
+ # status 1, if python is run with debug; otherwise it

just raises

Post by Jason Ekstrand
+ # the exception
+ import sys
+ from mako import exceptions
+

sys.stderr.write(exceptions.text_error_template().render() + '\n')

Post by Jason Ekstrand
+ sys.exit(1)
+ raise
diff --git a/src/compiler/spirv/vtn_private.h

b/src/compiler/spirv/vtn_private.h

Post by Jason Ekstrand
index 6d4ad3c..a0a4f3a 100644
--- a/src/compiler/spirv/vtn_private.h
+++ b/src/compiler/spirv/vtn_private.h
@@ -620,6 +620,10 @@ vtn_value(struct vtn_builder *b, uint32_t

value_id,

Post by Jason Ekstrand
return val;
}
+bool
+vtn_set_instruction_result_type(struct vtn_builder *b, SpvOp opcode,
+ const uint32_t *w, unsigned count);
+
struct vtn_ssa_value *vtn_ssa_value(struct vtn_builder *b,

uint32_t value_id);

Post by Jason Ekstrand
struct vtn_ssa_value *vtn_create_ssa_value(struct vtn_builder *b,

Jason Ekstrand

2017-12-11 22:44:21 UTC

Permalink

Post by Jason Ekstrand

Post by Jason Ekstrand
This autogenerated pass will automatically find and set the type

field

Post by Jason Ekstrand

Post by Jason Ekstrand
on all vtn_values. This way we always have the type and can use

it for

Post by Jason Ekstrand
validation and other checks.
---
src/compiler/Makefile.nir.am <http://Makefile.nir.am>

| 4 +

Post by Jason Ekstrand
src/compiler/nir/meson.build | 11 ++-
src/compiler/spirv/spirv_to_nir.c | 6 +-
src/compiler/spirv/vtn_gather_types_c.py | 125

+++++++++++++++++++++++++++++++

b/src/compiler/Makefile.nir.am <http://Makefile.nir.am>

spirv/spirv.core.grammar.json

Post by Jason Ekstrand
$(MKDIR_GEN)
$(PYTHON_GEN) $(srcdir)/spirv/spirv_info_c.py
+spirv/vtn_gather_types.c: spirv/vtn_gather_types_c.py

spirv/spirv.core.grammar.json

Post by Jason Ekstrand
+ $(MKDIR_GEN)
+ $(PYTHON_GEN) $(srcdir)/spirv/vtn_gather_types_c.py
+
noinst_PROGRAMS += spirv2nir
spirv2nir_SOURCES = \
diff --git a/src/compiler/nir/meson.build

b/src/compiler/nir/meson.build

nir_builder_opcodes_h],

Post by Jason Ekstrand

Post by Jason Ekstrand
+ nir_opcodes_h, nir_constant_expressions_c,

nir_builder_opcodes_h,

Post by Jason Ekstrand

Post by Jason Ekstrand
+ vtn_gather_types_c],
include_directories : [inc_common, inc_compiler,

include_directories('../spirv')],

Post by Jason Ekstrand
c_args : [c_vis_args, c_msvc_compat_args,

no_override_init_args],

Post by Jason Ekstrand

Post by Jason Ekstrand
link_with : libcompiler,
diff --git a/src/compiler/spirv/spirv_to_nir.c

b/src/compiler/spirv/spirv_to_nir.c

Post by Jason Ekstrand
index a50b14d..a2426bc 100644
--- a/src/compiler/spirv/spirv_to_nir.c
+++ b/src/compiler/spirv/spirv_to_nir.c
@@ -1261,7 +1261,6 @@ vtn_handle_constant(struct vtn_builder *b,

SpvOp opcode,

Post by Jason Ekstrand
const uint32_t *w, unsigned count)
{
struct vtn_value *val = vtn_push_value(b, w[2],

vtn_value_type_constant);

SpvOp opcode,

Post by Jason Ekstrand
const uint32_t *w,

unsigned count)

Post by Jason Ekstrand
{
+ vtn_set_instruction_result_type(b, opcode, w, count);
+
switch (opcode) {
@@ -3658,6 +3659,9 @@ spirv_to_nir(const uint32_t *words, size_t

word_count,

Post by Jason Ekstrand
words = vtn_foreach_instruction(b, words, word_end,

vtn_handle_variable_or_type_instruction);

Post by Jason Ekstrand
+ /* Set types on all vtn_values */
+ vtn_foreach_instruction(b, words, word_end,

vtn_set_instruction_result_type);

Post by Jason Ekstrand
+
vtn_build_cfg(b, words, word_end);
assert(b->entry_point->value_type == vtn_value_type_function);
diff --git a/src/compiler/spirv/vtn_gather_types_c.py

b/src/compiler/spirv/vtn_gather_types_c.py

obtaining a

Post by Jason Ekstrand
+ * copy of this software and associated documentation files (the

"Software"),

Post by Jason Ekstrand
+ * to deal in the Software without restriction, including without

limitation

Post by Jason Ekstrand
+ * the rights to use, copy, modify, merge, publish, distribute,

sublicense,

Post by Jason Ekstrand
+ * and/or sell copies of the Software, and to permit persons to

whom the

Post by Jason Ekstrand
+ * Software is furnished to do so, subject to the following
+ *
+ * The above copyright notice and this permission notice

(including the next

Post by Jason Ekstrand
+ * paragraph) shall be included in all copies or substantial

portions of the

Post by Jason Ekstrand
+ * Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY

KIND, EXPRESS OR

Post by Jason Ekstrand
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF

MERCHANTABILITY,

Post by Jason Ekstrand
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO

EVENT SHALL

Post by Jason Ekstrand
+ * THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM,

DAMAGES OR OTHER

Post by Jason Ekstrand
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR

OTHERWISE, ARISING

Post by Jason Ekstrand
+ * FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR

OTHER

untyped_insts

Post by Jason Ekstrand

Post by Jason Ekstrand
+
+ yield (name, res_arg_idx, res_type_arg_idx)
+
+TEMPLATE = Template(COPYRIGHT + """\
+
+/* DO NOT EDIT - This file is generated automatically by the
+ * vtn_gather_types_c.py script
+ */
+
+#include "vtn_private.h"
+
+struct type_args {
+ int res_idx;
+ int res_type_idx;
+};
+
+static struct type_args
+result_type_args_for_opcode(SpvOp opcode)
+{
+ switch (opcode) {
+ case Spv${opcode[0]}: return (struct type_args){ ${opcode[1]},

${opcode[2]} };

Post by Jason Ekstrand
+% endfor
+ default: return (struct type_args){ -1, -1 };
+ }
+}
+
+bool
+vtn_set_instruction_result_type(struct vtn_builder *b, SpvOp

opcode,

Post by Jason Ekstrand

Post by Jason Ekstrand
+ const uint32_t *w, unsigned count)
+{
+ struct type_args args = result_type_args_for_opcode(opcode);
+
+ if (args.res_idx < 0 || args.res_type_idx < 0)
+ return true;
+
+ struct vtn_value *val = vtn_untyped_value(b, w[1 +

args.res_idx]);

Post by Jason Ekstrand

Post by Jason Ekstrand
+ val->type = vtn_value(b, w[1 + args.res_type_idx],
+ vtn_value_type_type)->type;
+
+ return true;

It seems like one of the returns in this function should be false.

The

Post by Jason Ekstrand
first one?
No, they should both return true. This function uses the callback
interface for vtn_foreach_instruction where returning false means "stop
iterating". We want to continue on even if we see an instruction which
has no result type. I suppose we could at some point replace that bool
with an enum if that would be more clear.

Ah. For this function, if you restructure it as
if (args.res_index >= 0 && args.res_type_idx >= 0) {
...
}
return true;
the next person to look at the code won't think it's weird.

Done.

Post by Jason Ekstrand
With that
or a comment describing why all returns are true, this patch is

Thanks!

Post by Jason Ekstrand

"r").read())

Post by Jason Ekstrand
+
+ opcodes = list(find_result_types(spirv_info))
+
+ f.write(TEMPLATE.render(opcodes=opcodes))
+ # In the even there's an error this imports some helpers

from mako

Post by Jason Ekstrand
+ # to print a useful stack trace and prints it, then exits

with

Post by Jason Ekstrand
+ # status 1, if python is run with debug; otherwise it

just raises

Post by Jason Ekstrand
+ # the exception
+ import sys
+ from mako import exceptions
+

sys.stderr.write(exceptions.text_error_template().render() + '\n')

Post by Jason Ekstrand
+ sys.exit(1)
+ raise
diff --git a/src/compiler/spirv/vtn_private.h

b/src/compiler/spirv/vtn_private.h

Post by Jason Ekstrand
index 6d4ad3c..a0a4f3a 100644
--- a/src/compiler/spirv/vtn_private.h
+++ b/src/compiler/spirv/vtn_private.h
@@ -620,6 +620,10 @@ vtn_value(struct vtn_builder *b, uint32_t

value_id,

Post by Jason Ekstrand
return val;
}
+bool
+vtn_set_instruction_result_type(struct vtn_builder *b, SpvOp

opcode,

Post by Jason Ekstrand

Post by Jason Ekstrand
+ const uint32_t *w, unsigned

count);

Post by Jason Ekstrand

Post by Jason Ekstrand
+
struct vtn_ssa_value *vtn_ssa_value(struct vtn_builder *b,

uint32_t value_id);

Post by Jason Ekstrand
struct vtn_ssa_value *vtn_create_ssa_value(struct vtn_builder *b,

Rob Herring

2017-12-12 14:46:15 UTC

Permalink

This is missing the necessary android changes.

Rob

Jason Ekstrand

2017-12-07 16:12:10 UTC

Permalink

---
src/compiler/spirv/vtn_variables.c | 18 ++++++++++++++----
1 file changed, 14 insertions(+), 4 deletions(-)

diff --git a/src/compiler/spirv/vtn_variables.c b/src/compiler/spirv/vtn_variables.c
index cf44ed3..8ce19ff 100644
--- a/src/compiler/spirv/vtn_variables.c
+++ b/src/compiler/spirv/vtn_variables.c
@@ -1969,6 +1969,9 @@ vtn_handle_variables(struct vtn_builder *b, SpvOp opcode,
struct vtn_value *dest = vtn_value(b, w[1], vtn_value_type_pointer);
struct vtn_value *src = vtn_value(b, w[2], vtn_value_type_pointer);

+ vtn_fail_if(dest->type->deref != src->type->deref,
+ "Result and pointer types of OpLoad do not match");
+
vtn_variable_copy(b, dest->pointer, src->pointer);
break;
}
@@ -1976,8 +1979,11 @@ vtn_handle_variables(struct vtn_builder *b, SpvOp opcode,
case SpvOpLoad: {
struct vtn_type *res_type =
vtn_value(b, w[1], vtn_value_type_type)->type;
- struct vtn_pointer *src =
- vtn_value(b, w[3], vtn_value_type_pointer)->pointer;
+ struct vtn_value *src_val = vtn_value(b, w[3], vtn_value_type_pointer);
+ struct vtn_pointer *src = src_val->pointer;
+
+ vtn_fail_if(res_type != src_val->type->deref,
+ "Result and pointer types of OpLoad do not match");

if (src->mode == vtn_variable_mode_image ||
src->mode == vtn_variable_mode_sampler) {
@@ -1990,8 +1996,12 @@ vtn_handle_variables(struct vtn_builder *b, SpvOp opcode,
}

case SpvOpStore: {
- struct vtn_pointer *dest =
- vtn_value(b, w[1], vtn_value_type_pointer)->pointer;
+ struct vtn_value *dest_val = vtn_value(b, w[1], vtn_value_type_pointer);
+ struct vtn_pointer *dest = dest_val->pointer;
+ struct vtn_value *src_val = vtn_untyped_value(b, w[2]);
+
+ vtn_fail_if(dest_val->type->deref != src_val->type,
+ "Value and pointer types of OpStore do not match");

if (glsl_type_is_sampler(dest->type->type)) {
vtn_warn("OpStore of a sampler detected. Doing on-the-fly copy "

--
2.5.0.400.gff86faf

Jason Ekstrand

2017-12-07 22:26:27 UTC

Permalink

On Thu, Dec 7, 2017 at 11:54 AM, Michael Schellenberger Costa <

Hi Jason,

This should be OpCopyMemory?

Oops. Fixed locally.

On a more general side: As you want to cover every OpCode, why not
overload vtn_fail_if() so that it takes the OpCode and then prepends it to
vtn_fail_if(dest->type->deref != src->type->deref, opcode,"Result and
pointer types of do not match");
Would extend to
"OpCodeMemory: Result and pointer types of do not match"
That way there is no chance to really mess op the opcodes.

I'm not sure what I think about that. There may be cases where we want to
use vtn_fail_if where we don't have ready access to the opcode. One option
would be to use spirv_opcode_to_string instead of putting it in the string
but it may be hard to make that happen in such a way that we only call that
function if the vtn_fail condition is true.

All the best
Michael
+

Post by Jason Ekstrand
vtn_variable_copy(b, dest->pointer, src->pointer);
break;
}
@@ -1976,8 +1979,11 @@ vtn_handle_variables(struct vtn_builder *b, SpvOp opcode,
case SpvOpLoad: {
struct vtn_type *res_type =
vtn_value(b, w[1], vtn_value_type_type)->type;
- struct vtn_pointer *src =
- vtn_value(b, w[3], vtn_value_type_pointer)->pointer;
+ struct vtn_value *src_val = vtn_value(b, w[3],
vtn_value_type_pointer);
+ struct vtn_pointer *src = src_val->pointer;
+
+ vtn_fail_if(res_type != src_val->type->deref,
+ "Result and pointer types of OpLoad do not match");
if (src->mode == vtn_variable_mode_image ||
src->mode == vtn_variable_mode_sampler) {
@@ -1990,8 +1996,12 @@ vtn_handle_variables(struct vtn_builder *b, SpvOp opcode,
}
case SpvOpStore: {
- struct vtn_pointer *dest =
- vtn_value(b, w[1], vtn_value_type_pointer)->pointer;
+ struct vtn_value *dest_val = vtn_value(b, w[1],
vtn_value_type_pointer);
+ struct vtn_pointer *dest = dest_val->pointer;
+ struct vtn_value *src_val = vtn_untyped_value(b, w[2]);
+
+ vtn_fail_if(dest_val->type->deref != src_val->type,
+ "Value and pointer types of OpStore do not match");
if (glsl_type_is_sampler(dest->type->type)) {
vtn_warn("OpStore of a sampler detected. Doing on-the-fly copy "

Ian Romanick

2017-12-11 17:52:45 UTC

Permalink

Post by Jason Ekstrand
On Thu, Dec 7, 2017 at 11:54 AM, Michael Schellenberger Costa
Hi Jason,
---
src/compiler/spirv/vtn_variables.c | 18 ++++++++++++++----
1 file changed, 14 insertions(+), 4 deletions(-)
diff --git a/src/compiler/spirv/vtn_variables.c
b/src/compiler/spirv/vtn_variables.c
index cf44ed3..8ce19ff 100644
--- a/src/compiler/spirv/vtn_variables.c
+++ b/src/compiler/spirv/vtn_variables.c
@@ -1969,6 +1969,9 @@ vtn_handle_variables(struct vtn_builder
*b, SpvOp opcode,
struct vtn_value *dest = vtn_value(b, w[1],
vtn_value_type_pointer);
struct vtn_value *src = vtn_value(b, w[2],
vtn_value_type_pointer);
+ vtn_fail_if(dest->type->deref != src->type->deref,
+ "Result and pointer types of OpLoad do not
match");
This should be OpCopyMemory?
Oops. Fixed locally.

On a more general side: As you want to cover every OpCode, why not
overload vtn_fail_if() so that it takes the OpCode and then prepends
vtn_fail_if(dest->type->deref != src->type->deref, opcode,"Result
and pointer types of do not match");
Would extend to
"OpCodeMemory: Result and pointer types of do not match"
That way there is no chance to really mess op the opcodes.
I'm not sure what I think about that. There may be cases where we want
to use vtn_fail_if where we don't have ready access to the opcode. One
option would be to use spirv_opcode_to_string instead of putting it in
the string but it may be hard to make that happen in such a way that we
only call that function if the vtn_fail condition is true.

You could make a vtn_fail_opcode_if that takes the opcode as a
parameter. I'm not sure if that would strictly be an improvement. *shrug*

Post by Jason Ekstrand

All the best
Michael
+
vtn_variable_copy(b, dest->pointer, src->pointer);
break;
}
@@ -1976,8 +1979,11 @@ vtn_handle_variables(struct vtn_builder
*b, SpvOp opcode,
case SpvOpLoad: {
struct vtn_type *res_type =
vtn_value(b, w[1], vtn_value_type_type)->type;
- struct vtn_pointer *src =
- vtn_value(b, w[3], vtn_value_type_pointer)->pointer;
+ struct vtn_value *src_val = vtn_value(b, w[3],
vtn_value_type_pointer);
+ struct vtn_pointer *src = src_val->pointer;
+
+ vtn_fail_if(res_type != src_val->type->deref,
+ "Result and pointer types of OpLoad do not
match");
if (src->mode == vtn_variable_mode_image ||
src->mode == vtn_variable_mode_sampler) {
@@ -1990,8 +1996,12 @@ vtn_handle_variables(struct vtn_builder
*b, SpvOp opcode,
}
case SpvOpStore: {
- struct vtn_pointer *dest =
- vtn_value(b, w[1], vtn_value_type_pointer)->pointer;
+ struct vtn_value *dest_val = vtn_value(b, w[1],
vtn_value_type_pointer);
+ struct vtn_pointer *dest = dest_val->pointer;
+ struct vtn_value *src_val = vtn_untyped_value(b, w[2]);
+
+ vtn_fail_if(dest_val->type->deref != src_val->type,
+ "Value and pointer types of OpStore do not
match");
if (glsl_type_is_sampler(dest->type->type)) {
vtn_warn("OpStore of a sampler detected. Doing
on-the-fly copy "
_______________________________________________
mesa-dev mailing list
https://lists.freedesktop.org/mailman/listinfo/mesa-dev

Grazvydas Ignotas

2017-12-30 23:57:06 UTC

Permalink

Hi,

I don't know if it's the game's fault, but it appears this change broke DOOM.
here is the offending spirv binary:
https://people.freedesktop.org/~notaz/doom_compute_spirv

Gražvydas

Post by Jason Ekstrand
---
src/compiler/spirv/vtn_variables.c | 18 ++++++++++++++----
1 file changed, 14 insertions(+), 4 deletions(-)
diff --git a/src/compiler/spirv/vtn_variables.c b/src/compiler/spirv/vtn_variables.c
index cf44ed3..8ce19ff 100644
--- a/src/compiler/spirv/vtn_variables.c
+++ b/src/compiler/spirv/vtn_variables.c
@@ -1969,6 +1969,9 @@ vtn_handle_variables(struct vtn_builder *b, SpvOp opcode,
struct vtn_value *dest = vtn_value(b, w[1], vtn_value_type_pointer);
struct vtn_value *src = vtn_value(b, w[2], vtn_value_type_pointer);
+ vtn_fail_if(dest->type->deref != src->type->deref,
+ "Result and pointer types of OpLoad do not match");
+
vtn_variable_copy(b, dest->pointer, src->pointer);
break;
}
@@ -1976,8 +1979,11 @@ vtn_handle_variables(struct vtn_builder *b, SpvOp opcode,
case SpvOpLoad: {
struct vtn_type *res_type =
vtn_value(b, w[1], vtn_value_type_type)->type;
- struct vtn_pointer *src =
- vtn_value(b, w[3], vtn_value_type_pointer)->pointer;
+ struct vtn_value *src_val = vtn_value(b, w[3], vtn_value_type_pointer);
+ struct vtn_pointer *src = src_val->pointer;
+
+ vtn_fail_if(res_type != src_val->type->deref,
+ "Result and pointer types of OpLoad do not match");
if (src->mode == vtn_variable_mode_image ||
src->mode == vtn_variable_mode_sampler) {
@@ -1990,8 +1996,12 @@ vtn_handle_variables(struct vtn_builder *b, SpvOp opcode,
}
case SpvOpStore: {
- struct vtn_pointer *dest =
- vtn_value(b, w[1], vtn_value_type_pointer)->pointer;
+ struct vtn_value *dest_val = vtn_value(b, w[1], vtn_value_type_pointer);
+ struct vtn_pointer *dest = dest_val->pointer;
+ struct vtn_value *src_val = vtn_untyped_value(b, w[2]);
+
+ vtn_fail_if(dest_val->type->deref != src_val->type,
+ "Value and pointer types of OpStore do not match");
if (glsl_type_is_sampler(dest->type->type)) {
vtn_warn("OpStore of a sampler detected. Doing on-the-fly copy "
--
2.5.0.400.gff86faf
_______________________________________________
mesa-dev mailing list
https://lists.freedesktop.org/mailman/listinfo/mesa-dev

Jason Ekstrand

2017-12-31 02:31:54 UTC

Permalink

Post by Grazvydas Ignotas
Hi,
I don't know if it's the game's fault, but it appears this change broke DOOM.
https://people.freedesktop.org/~notaz/doom_compute_spirv

Have you filed a bug? Please do and assign it to me. I'll take a look at
it in a week or so.

--Jason

Post by Grazvydas Ignotas
GraÅŸvydas

Post by Jason Ekstrand
---
src/compiler/spirv/vtn_variables.c | 18 ++++++++++++++----
1 file changed, 14 insertions(+), 4 deletions(-)
diff --git a/src/compiler/spirv/vtn_variables.c

b/src/compiler/spirv/vtn_variables.c

Post by Jason Ekstrand
index cf44ed3..8ce19ff 100644
--- a/src/compiler/spirv/vtn_variables.c
+++ b/src/compiler/spirv/vtn_variables.c
@@ -1969,6 +1969,9 @@ vtn_handle_variables(struct vtn_builder *b, SpvOp

opcode,

Post by Jason Ekstrand
struct vtn_value *dest = vtn_value(b, w[1],

vtn_value_type_pointer);

Post by Jason Ekstrand
struct vtn_value *src = vtn_value(b, w[2],

vtn_value_type_pointer);

Post by Jason Ekstrand
+ vtn_fail_if(dest->type->deref != src->type->deref,
+ "Result and pointer types of OpLoad do not match");
+
vtn_variable_copy(b, dest->pointer, src->pointer);
break;
}
@@ -1976,8 +1979,11 @@ vtn_handle_variables(struct vtn_builder *b, SpvOp

opcode,

Post by Jason Ekstrand
case SpvOpLoad: {
struct vtn_type *res_type =
vtn_value(b, w[1], vtn_value_type_type)->type;
- struct vtn_pointer *src =
- vtn_value(b, w[3], vtn_value_type_pointer)->pointer;
+ struct vtn_value *src_val = vtn_value(b, w[3],

vtn_value_type_pointer);

Post by Jason Ekstrand
+ struct vtn_pointer *src = src_val->pointer;
+
+ vtn_fail_if(res_type != src_val->type->deref,
+ "Result and pointer types of OpLoad do not match");
if (src->mode == vtn_variable_mode_image ||
src->mode == vtn_variable_mode_sampler) {
@@ -1990,8 +1996,12 @@ vtn_handle_variables(struct vtn_builder *b, SpvOp

opcode,

Post by Jason Ekstrand
}
case SpvOpStore: {
- struct vtn_pointer *dest =
- vtn_value(b, w[1], vtn_value_type_pointer)->pointer;
+ struct vtn_value *dest_val = vtn_value(b, w[1],

vtn_value_type_pointer);

Post by Jason Ekstrand
+ struct vtn_pointer *dest = dest_val->pointer;
+ struct vtn_value *src_val = vtn_untyped_value(b, w[2]);
+
+ vtn_fail_if(dest_val->type->deref != src_val->type,
+ "Value and pointer types of OpStore do not match");
if (glsl_type_is_sampler(dest->type->type)) {
vtn_warn("OpStore of a sampler detected. Doing on-the-fly

copy "

Post by Jason Ekstrand
--
2.5.0.400.gff86faf
_______________________________________________
mesa-dev mailing list
https://lists.freedesktop.org/mailman/listinfo/mesa-dev

Jason Ekstrand

2017-12-07 16:12:11 UTC

Permalink

---
src/compiler/spirv/spirv_to_nir.c | 4 ++++
1 file changed, 4 insertions(+)

diff --git a/src/compiler/spirv/spirv_to_nir.c b/src/compiler/spirv/spirv_to_nir.c
index a2426bc..253a012 100644
--- a/src/compiler/spirv/spirv_to_nir.c
+++ b/src/compiler/spirv/spirv_to_nir.c
@@ -887,6 +887,7 @@ vtn_handle_type(struct vtn_builder *b, SpvOp opcode,
case SpvOpTypeBool:
val->type->base_type = vtn_base_type_scalar;
val->type->type = glsl_bool_type();
+ val->type->length = 1;
break;
case SpvOpTypeInt: {
int bit_size = w[2];
@@ -905,6 +906,7 @@ vtn_handle_type(struct vtn_builder *b, SpvOp opcode,
default:
vtn_fail("Invalid int bit size");
}
+ val->type->length = 1;
break;
}

@@ -924,6 +926,7 @@ vtn_handle_type(struct vtn_builder *b, SpvOp opcode,
default:
vtn_fail("Invalid float bit size");
}
+ val->type->length = 1;
break;
}

@@ -934,6 +937,7 @@ vtn_handle_type(struct vtn_builder *b, SpvOp opcode,
vtn_assert(glsl_type_is_scalar(base->type));
val->type->base_type = vtn_base_type_vector;
val->type->type = glsl_vector_type(glsl_get_base_type(base->type), elems);
+ val->type->length = elems;
val->type->stride = glsl_get_bit_size(base->type) / 8;
val->type->array_element = base;
break;

--
2.5.0.400.gff86faf

Ian Romanick

2017-12-11 18:01:03 UTC

Permalink

This patch is

Post by Jason Ekstrand
---
src/compiler/spirv/spirv_to_nir.c | 4 ++++
1 file changed, 4 insertions(+)
diff --git a/src/compiler/spirv/spirv_to_nir.c b/src/compiler/spirv/spirv_to_nir.c
index a2426bc..253a012 100644
--- a/src/compiler/spirv/spirv_to_nir.c
+++ b/src/compiler/spirv/spirv_to_nir.c
@@ -887,6 +887,7 @@ vtn_handle_type(struct vtn_builder *b, SpvOp opcode,
val->type->base_type = vtn_base_type_scalar;
val->type->type = glsl_bool_type();
+ val->type->length = 1;
break;
case SpvOpTypeInt: {
int bit_size = w[2];
@@ -905,6 +906,7 @@ vtn_handle_type(struct vtn_builder *b, SpvOp opcode,
vtn_fail("Invalid int bit size");
}
+ val->type->length = 1;
break;
}
@@ -924,6 +926,7 @@ vtn_handle_type(struct vtn_builder *b, SpvOp opcode,
vtn_fail("Invalid float bit size");
}
+ val->type->length = 1;
break;
}
@@ -934,6 +937,7 @@ vtn_handle_type(struct vtn_builder *b, SpvOp opcode,
vtn_assert(glsl_type_is_scalar(base->type));
val->type->base_type = vtn_base_type_vector;
val->type->type = glsl_vector_type(glsl_get_base_type(base->type), elems);
+ val->type->length = elems;
val->type->stride = glsl_get_bit_size(base->type) / 8;
val->type->array_element = base;
break;

Jason Ekstrand

2017-12-07 16:12:12 UTC

Permalink

---
src/compiler/spirv/spirv_to_nir.c | 32 ++++++++++++++++++++++++++++++++
1 file changed, 32 insertions(+)

diff --git a/src/compiler/spirv/spirv_to_nir.c b/src/compiler/spirv/spirv_to_nir.c
index 253a012..11c8c2a 100644
--- a/src/compiler/spirv/spirv_to_nir.c
+++ b/src/compiler/spirv/spirv_to_nir.c
@@ -3460,6 +3460,38 @@ vtn_handle_body_instruction(struct vtn_builder *b, SpvOp opcode,
/* Handle OpSelect up-front here because it needs to be able to handle
* pointers and not just regular vectors and scalars.
*/
+ struct vtn_value *res_val = vtn_untyped_value(b, w[2]);
+ struct vtn_value *sel_val = vtn_untyped_value(b, w[3]);
+ struct vtn_value *obj1_val = vtn_untyped_value(b, w[4]);
+ struct vtn_value *obj2_val = vtn_untyped_value(b, w[5]);
+
+ const struct glsl_type *sel_type;
+ switch (res_val->type->base_type) {
+ case vtn_base_type_scalar:
+ sel_type = glsl_bool_type();
+ break;
+ case vtn_base_type_vector:
+ sel_type = glsl_vector_type(GLSL_TYPE_BOOL, res_val->type->length);
+ break;
+ case vtn_base_type_pointer:
+ /* We need to have actual storage for pointer types */
+ vtn_fail_if(res_val->type->type == NULL,
+ "Invalid pointer result type for OpSelect");
+ sel_type = glsl_bool_type();
+ break;
+ default:
+ vtn_fail("Result type of OpSelect must be a scalar, vector, or pointer");
+ }
+
+ vtn_fail_if(sel_val->type->type != sel_type,
+ "Condition type of ObSelect must be a scalar or vector of "
+ "Boolean type. It must have the same number of components "
+ "as Result Type");
+
+ vtn_fail_if(obj1_val->type != res_val->type ||
+ obj2_val->type != res_val->type,
+ "Object types must match the result type in OpSelect");
+
struct vtn_type *res_type = vtn_value(b, w[1], vtn_value_type_type)->type;
struct vtn_ssa_value *ssa = vtn_create_ssa_value(b, res_type->type);
ssa->def = nir_bcsel(&b->nb, vtn_ssa_value(b, w[3])->def,

--
2.5.0.400.gff86faf

Ian Romanick

2017-12-11 18:01:22 UTC

Permalink

Post by Jason Ekstrand
---
src/compiler/spirv/spirv_to_nir.c | 32 ++++++++++++++++++++++++++++++++
1 file changed, 32 insertions(+)
diff --git a/src/compiler/spirv/spirv_to_nir.c b/src/compiler/spirv/spirv_to_nir.c
index 253a012..11c8c2a 100644
--- a/src/compiler/spirv/spirv_to_nir.c
+++ b/src/compiler/spirv/spirv_to_nir.c
@@ -3460,6 +3460,38 @@ vtn_handle_body_instruction(struct vtn_builder *b, SpvOp opcode,
/* Handle OpSelect up-front here because it needs to be able to handle
* pointers and not just regular vectors and scalars.
*/
+ struct vtn_value *res_val = vtn_untyped_value(b, w[2]);
+ struct vtn_value *sel_val = vtn_untyped_value(b, w[3]);
+ struct vtn_value *obj1_val = vtn_untyped_value(b, w[4]);
+ struct vtn_value *obj2_val = vtn_untyped_value(b, w[5]);
+
+ const struct glsl_type *sel_type;
+ switch (res_val->type->base_type) {
+ sel_type = glsl_bool_type();
+ break;
+ sel_type = glsl_vector_type(GLSL_TYPE_BOOL, res_val->type->length);
+ break;
+ /* We need to have actual storage for pointer types */
+ vtn_fail_if(res_val->type->type == NULL,
+ "Invalid pointer result type for OpSelect");
+ sel_type = glsl_bool_type();
+ break;
+ vtn_fail("Result type of OpSelect must be a scalar, vector, or pointer");
+ }
+
+ vtn_fail_if(sel_val->type->type != sel_type,
+ "Condition type of ObSelect must be a scalar or vector of "

OpSelect

With that fixed, this patch is

Post by Jason Ekstrand
+ "Boolean type. It must have the same number of components "
+ "as Result Type");
+
+ vtn_fail_if(obj1_val->type != res_val->type ||
+ obj2_val->type != res_val->type,
+ "Object types must match the result type in OpSelect");
+
struct vtn_type *res_type = vtn_value(b, w[1], vtn_value_type_type)->type;
struct vtn_ssa_value *ssa = vtn_create_ssa_value(b, res_type->type);
ssa->def = nir_bcsel(&b->nb, vtn_ssa_value(b, w[3])->def,

Jason Ekstrand

2017-12-07 16:12:13 UTC

Permalink

---
src/compiler/spirv/spirv_to_nir.c | 14 +++++++++++---
1 file changed, 11 insertions(+), 3 deletions(-)

diff --git a/src/compiler/spirv/spirv_to_nir.c b/src/compiler/spirv/spirv_to_nir.c
index 11c8c2a..4b93b11 100644
--- a/src/compiler/spirv/spirv_to_nir.c
+++ b/src/compiler/spirv/spirv_to_nir.c
@@ -934,7 +934,11 @@ vtn_handle_type(struct vtn_builder *b, SpvOp opcode,
struct vtn_type *base = vtn_value(b, w[2], vtn_value_type_type)->type;
unsigned elems = w[3];

- vtn_assert(glsl_type_is_scalar(base->type));
+ vtn_fail_if(base->base_type != vtn_base_type_scalar,
+ "Base type for OpTypeVector must be a scalar");
+ vtn_fail_if(elems < 2 || elems > 4,
+ "Invalid component count for OpTypeVector");
+
val->type->base_type = vtn_base_type_vector;
val->type->type = glsl_vector_type(glsl_get_base_type(base->type), elems);
val->type->length = elems;
@@ -947,12 +951,16 @@ vtn_handle_type(struct vtn_builder *b, SpvOp opcode,
struct vtn_type *base = vtn_value(b, w[2], vtn_value_type_type)->type;
unsigned columns = w[3];

- vtn_assert(glsl_type_is_vector(base->type));
+ vtn_fail_if(columns < 2 || columns > 4,
+ "Invalid column count for OpTypeMatrix");
+
val->type->base_type = vtn_base_type_matrix;
val->type->type = glsl_matrix_type(glsl_get_base_type(base->type),
glsl_get_vector_elements(base->type),
columns);
- vtn_assert(!glsl_type_is_error(val->type->type));
+ vtn_fail_if(glsl_type_is_error(val->type->type),
+ "Unsupported base type for OpTypeMatrix");
+ assert(!glsl_type_is_error(val->type->type));
val->type->length = columns;
val->type->array_element = base;
val->type->row_major = false;

--
2.5.0.400.gff86faf

Ian Romanick

2017-12-11 18:05:38 UTC

Permalink

Post by Jason Ekstrand
---
src/compiler/spirv/spirv_to_nir.c | 14 +++++++++++---
1 file changed, 11 insertions(+), 3 deletions(-)
diff --git a/src/compiler/spirv/spirv_to_nir.c b/src/compiler/spirv/spirv_to_nir.c
index 11c8c2a..4b93b11 100644
--- a/src/compiler/spirv/spirv_to_nir.c
+++ b/src/compiler/spirv/spirv_to_nir.c
@@ -934,7 +934,11 @@ vtn_handle_type(struct vtn_builder *b, SpvOp opcode,
struct vtn_type *base = vtn_value(b, w[2], vtn_value_type_type)->type;
unsigned elems = w[3];
- vtn_assert(glsl_type_is_scalar(base->type));
+ vtn_fail_if(base->base_type != vtn_base_type_scalar,
+ "Base type for OpTypeVector must be a scalar");
+ vtn_fail_if(elems < 2 || elems > 4,
+ "Invalid component count for OpTypeVector");
+
val->type->base_type = vtn_base_type_vector;
val->type->type = glsl_vector_type(glsl_get_base_type(base->type), elems);
val->type->length = elems;
@@ -947,12 +951,16 @@ vtn_handle_type(struct vtn_builder *b, SpvOp opcode,
struct vtn_type *base = vtn_value(b, w[2], vtn_value_type_type)->type;
unsigned columns = w[3];
- vtn_assert(glsl_type_is_vector(base->type));
+ vtn_fail_if(columns < 2 || columns > 4,
+ "Invalid column count for OpTypeMatrix");
+

I think we should still check glsl_type_is_vector here. The
glsl_type_is_error check below will catch that case, but we'll give a
better error message with the explicit check here.

Post by Jason Ekstrand
val->type->base_type = vtn_base_type_matrix;
val->type->type = glsl_matrix_type(glsl_get_base_type(base->type),
glsl_get_vector_elements(base->type),
columns);
- vtn_assert(!glsl_type_is_error(val->type->type));
+ vtn_fail_if(glsl_type_is_error(val->type->type),
+ "Unsupported base type for OpTypeMatrix");
+ assert(!glsl_type_is_error(val->type->type));
val->type->length = columns;
val->type->array_element = base;
val->type->row_major = false;

Jason Ekstrand

2017-12-11 22:46:17 UTC

Permalink

Post by Jason Ekstrand

Post by Jason Ekstrand
---
src/compiler/spirv/spirv_to_nir.c | 14 +++++++++++---
1 file changed, 11 insertions(+), 3 deletions(-)
diff --git a/src/compiler/spirv/spirv_to_nir.c

b/src/compiler/spirv/spirv_to_nir.c

Post by Jason Ekstrand
index 11c8c2a..4b93b11 100644
--- a/src/compiler/spirv/spirv_to_nir.c
+++ b/src/compiler/spirv/spirv_to_nir.c
@@ -934,7 +934,11 @@ vtn_handle_type(struct vtn_builder *b, SpvOp opcode,
struct vtn_type *base = vtn_value(b, w[2],

vtn_value_type_type)->type;

Post by Jason Ekstrand
unsigned elems = w[3];
- vtn_assert(glsl_type_is_scalar(base->type));
+ vtn_fail_if(base->base_type != vtn_base_type_scalar,
+ "Base type for OpTypeVector must be a scalar");
+ vtn_fail_if(elems < 2 || elems > 4,
+ "Invalid component count for OpTypeVector");
+
val->type->base_type = vtn_base_type_vector;
val->type->type = glsl_vector_type(glsl_get_base_type(base->type),

elems);

Post by Jason Ekstrand
val->type->length = elems;
@@ -947,12 +951,16 @@ vtn_handle_type(struct vtn_builder *b, SpvOp

opcode,

Post by Jason Ekstrand
struct vtn_type *base = vtn_value(b, w[2],

vtn_value_type_type)->type;

Post by Jason Ekstrand
unsigned columns = w[3];
- vtn_assert(glsl_type_is_vector(base->type));
+ vtn_fail_if(columns < 2 || columns > 4,
+ "Invalid column count for OpTypeMatrix");
+

I think we should still check glsl_type_is_vector here. The
glsl_type_is_error check below will catch that case, but we'll give a
better error message with the explicit check here.

Sure. I'll add something and send a v2.

Post by Jason Ekstrand

Post by Jason Ekstrand
val->type->base_type = vtn_base_type_matrix;
val->type->type = glsl_matrix_type(glsl_get_

base_type(base->type),

Post by Jason Ekstrand
glsl_get_vector_elements(base-
type),
columns);
- vtn_assert(!glsl_type_is_error(val->type->type));
+ vtn_fail_if(glsl_type_is_error(val->type->type),
+ "Unsupported base type for OpTypeMatrix");
+ assert(!glsl_type_is_error(val->type->type));
val->type->length = columns;
val->type->array_element = base;
val->type->row_major = false;

Jason Ekstrand

2017-12-11 23:13:59 UTC

Permalink

---
src/compiler/spirv/spirv_to_nir.c | 16 +++++++++++++---
1 file changed, 13 insertions(+), 3 deletions(-)

diff --git a/src/compiler/spirv/spirv_to_nir.c b/src/compiler/spirv/spirv_to_nir.c
index 82e5c8c..c5d6131 100644
--- a/src/compiler/spirv/spirv_to_nir.c
+++ b/src/compiler/spirv/spirv_to_nir.c
@@ -934,7 +934,11 @@ vtn_handle_type(struct vtn_builder *b, SpvOp opcode,
struct vtn_type *base = vtn_value(b, w[2], vtn_value_type_type)->type;
unsigned elems = w[3];

- vtn_assert(glsl_type_is_scalar(base->type));
+ vtn_fail_if(base->base_type != vtn_base_type_scalar,
+ "Base type for OpTypeVector must be a scalar");
+ vtn_fail_if(elems < 2 || elems > 4,
+ "Invalid component count for OpTypeVector");
+
val->type->base_type = vtn_base_type_vector;
val->type->type = glsl_vector_type(glsl_get_base_type(base->type), elems);
val->type->length = elems;
@@ -947,12 +951,18 @@ vtn_handle_type(struct vtn_builder *b, SpvOp opcode,
struct vtn_type *base = vtn_value(b, w[2], vtn_value_type_type)->type;
unsigned columns = w[3];

- vtn_assert(glsl_type_is_vector(base->type));
+ vtn_fail_if(base->base_type != vtn_base_type_vector,
+ "Base type for OpTypeMatrix must be a vector");
+ vtn_fail_if(columns < 2 || columns > 4,
+ "Invalid column count for OpTypeMatrix");
+
val->type->base_type = vtn_base_type_matrix;
val->type->type = glsl_matrix_type(glsl_get_base_type(base->type),
glsl_get_vector_elements(base->type),
columns);
- vtn_assert(!glsl_type_is_error(val->type->type));
+ vtn_fail_if(glsl_type_is_error(val->type->type),
+ "Unsupported base type for OpTypeMatrix");
+ assert(!glsl_type_is_error(val->type->type));
val->type->length = columns;
val->type->array_element = base;
val->type->row_major = false;

--
2.5.0.400.gff86faf

Ian Romanick

2017-12-12 00:10:47 UTC

Permalink

Post by Jason Ekstrand
---
src/compiler/spirv/spirv_to_nir.c | 16 +++++++++++++---
1 file changed, 13 insertions(+), 3 deletions(-)
diff --git a/src/compiler/spirv/spirv_to_nir.c b/src/compiler/spirv/spirv_to_nir.c
index 82e5c8c..c5d6131 100644
--- a/src/compiler/spirv/spirv_to_nir.c
+++ b/src/compiler/spirv/spirv_to_nir.c
@@ -934,7 +934,11 @@ vtn_handle_type(struct vtn_builder *b, SpvOp opcode,
struct vtn_type *base = vtn_value(b, w[2], vtn_value_type_type)->type;
unsigned elems = w[3];
- vtn_assert(glsl_type_is_scalar(base->type));
+ vtn_fail_if(base->base_type != vtn_base_type_scalar,
+ "Base type for OpTypeVector must be a scalar");
+ vtn_fail_if(elems < 2 || elems > 4,
+ "Invalid component count for OpTypeVector");
+
val->type->base_type = vtn_base_type_vector;
val->type->type = glsl_vector_type(glsl_get_base_type(base->type), elems);
val->type->length = elems;
@@ -947,12 +951,18 @@ vtn_handle_type(struct vtn_builder *b, SpvOp opcode,
struct vtn_type *base = vtn_value(b, w[2], vtn_value_type_type)->type;
unsigned columns = w[3];
- vtn_assert(glsl_type_is_vector(base->type));
+ vtn_fail_if(base->base_type != vtn_base_type_vector,
+ "Base type for OpTypeMatrix must be a vector");
+ vtn_fail_if(columns < 2 || columns > 4,
+ "Invalid column count for OpTypeMatrix");
+
val->type->base_type = vtn_base_type_matrix;
val->type->type = glsl_matrix_type(glsl_get_base_type(base->type),
glsl_get_vector_elements(base->type),
columns);
- vtn_assert(!glsl_type_is_error(val->type->type));
+ vtn_fail_if(glsl_type_is_error(val->type->type),
+ "Unsupported base type for OpTypeMatrix");
+ assert(!glsl_type_is_error(val->type->type));
val->type->length = columns;
val->type->array_element = base;
val->type->row_major = false;

Jason Ekstrand

2017-12-07 16:12:14 UTC

Permalink

Instead of calling vtn_add_case for the default case and then looping,
add an is_default variable and do everything inside the loop. This will
make the next commit easier.
---
src/compiler/spirv/vtn_cfg.c | 17 ++++++++++++++---
1 file changed, 14 insertions(+), 3 deletions(-)

diff --git a/src/compiler/spirv/vtn_cfg.c b/src/compiler/spirv/vtn_cfg.c
index 25140ff..9d1ca84 100644
--- a/src/compiler/spirv/vtn_cfg.c
+++ b/src/compiler/spirv/vtn_cfg.c
@@ -425,9 +425,20 @@ vtn_cfg_walk_blocks(struct vtn_builder *b, struct list_head *cf_list,
const uint32_t *branch_end =
block->branch + (block->branch[0] >> SpvWordCountShift);

- vtn_add_case(b, swtch, break_block, block->branch[2], 0, true);
- for (const uint32_t *w = block->branch + 3; w < branch_end; w += 2)
- vtn_add_case(b, swtch, break_block, w[1], w[0], false);
+ bool is_default = true;
+ for (const uint32_t *w = block->branch + 2; w < branch_end;) {
+ uint32_t literal = 0;
+ if (!is_default) {
+ literal = *w;
+ w++;
+ }
+
+ uint32_t block_id = *w;
+ w++;
+
+ vtn_add_case(b, swtch, break_block, block_id, literal, is_default);
+ is_default = false;
+ }

/* Now, we go through and walk the blocks. While we walk through
* the blocks, we also gather the much-needed fall-through

--
2.5.0.400.gff86faf

Ian Romanick

2017-12-11 18:08:38 UTC

Permalink

Post by Jason Ekstrand
Instead of calling vtn_add_case for the default case and then looping,
add an is_default variable and do everything inside the loop. This will
make the next commit easier.
---
src/compiler/spirv/vtn_cfg.c | 17 ++++++++++++++---
1 file changed, 14 insertions(+), 3 deletions(-)
diff --git a/src/compiler/spirv/vtn_cfg.c b/src/compiler/spirv/vtn_cfg.c
index 25140ff..9d1ca84 100644
--- a/src/compiler/spirv/vtn_cfg.c
+++ b/src/compiler/spirv/vtn_cfg.c
@@ -425,9 +425,20 @@ vtn_cfg_walk_blocks(struct vtn_builder *b, struct list_head *cf_list,
const uint32_t *branch_end =
block->branch + (block->branch[0] >> SpvWordCountShift);
- vtn_add_case(b, swtch, break_block, block->branch[2], 0, true);
- for (const uint32_t *w = block->branch + 3; w < branch_end; w += 2)
- vtn_add_case(b, swtch, break_block, w[1], w[0], false);
+ bool is_default = true;
+ for (const uint32_t *w = block->branch + 2; w < branch_end;) {
+ uint32_t literal = 0;
+ if (!is_default) {
+ literal = *w;
+ w++;
+ }
+
+ uint32_t block_id = *w;
+ w++;

In other parts of Mesa, this would be

const uint32_t block_id = *(w++);

Is that not the preferred style here too? Having looked ahead at the
next patch, I can see why the other dereference of w is not like this.

Post by Jason Ekstrand
+
+ vtn_add_case(b, swtch, break_block, block_id, literal, is_default);
+ is_default = false;
+ }
/* Now, we go through and walk the blocks. While we walk through
* the blocks, we also gather the much-needed fall-through

Jason Ekstrand

2017-12-11 22:50:59 UTC

Permalink

Post by Jason Ekstrand

list_head *cf_list,

Post by Jason Ekstrand
const uint32_t *branch_end =
block->branch + (block->branch[0] >> SpvWordCountShift);
- vtn_add_case(b, swtch, break_block, block->branch[2], 0, true);
- for (const uint32_t *w = block->branch + 3; w < branch_end; w

+= 2)

Post by Jason Ekstrand
- vtn_add_case(b, swtch, break_block, w[1], w[0], false);
+ bool is_default = true;
+ for (const uint32_t *w = block->branch + 2; w < branch_end;) {
+ uint32_t literal = 0;
+ if (!is_default) {
+ literal = *w;
+ w++;
+ }
+
+ uint32_t block_id = *w;
+ w++;

In other parts of Mesa, this would be
const uint32_t block_id = *(w++);
Is that not the preferred style here too? Having looked ahead at the
next patch, I can see why the other dereference of w is not like this.

Yeah, that's cleaner. I've switched to that and rebased patch 8 on it.

Post by Jason Ekstrand

Post by Jason Ekstrand
+
+ vtn_add_case(b, swtch, break_block, block_id, literal,

is_default);

Post by Jason Ekstrand
+ is_default = false;
+ }
/* Now, we go through and walk the blocks. While we walk

through

Post by Jason Ekstrand
* the blocks, we also gather the much-needed fall-through

Ian Romanick

2017-12-11 22:56:07 UTC

Permalink

Post by Jason Ekstrand

Post by Jason Ekstrand
Instead of calling vtn_add_case for the default case and then looping,
add an is_default variable and do everything inside the loop.

This will

Post by Jason Ekstrand
make the next commit easier.
---
src/compiler/spirv/vtn_cfg.c | 17 ++++++++++++++---
1 file changed, 14 insertions(+), 3 deletions(-)
diff --git a/src/compiler/spirv/vtn_cfg.c

b/src/compiler/spirv/vtn_cfg.c

Post by Jason Ekstrand
index 25140ff..9d1ca84 100644
--- a/src/compiler/spirv/vtn_cfg.c
+++ b/src/compiler/spirv/vtn_cfg.c
@@ -425,9 +425,20 @@ vtn_cfg_walk_blocks(struct vtn_builder *b,

struct list_head *cf_list,

Post by Jason Ekstrand
const uint32_t *branch_end =
block->branch + (block->branch[0] >> SpvWordCountShift);
- vtn_add_case(b, swtch, break_block, block->branch[2], 0,

true);

Post by Jason Ekstrand
- for (const uint32_t *w = block->branch + 3; w <

branch_end; w += 2)

Post by Jason Ekstrand
- vtn_add_case(b, swtch, break_block, w[1], w[0], false);
+ bool is_default = true;
+ for (const uint32_t *w = block->branch + 2; w <

branch_end;) {

Post by Jason Ekstrand
+ uint32_t literal = 0;
+ if (!is_default) {
+ literal = *w;
+ w++;
+ }
+
+ uint32_t block_id = *w;
+ w++;

In other parts of Mesa, this would be
const uint32_t block_id = *(w++);
Is that not the preferred style here too? Having looked ahead at the
next patch, I can see why the other dereference of w is not like this.
Yeah, that's cleaner. I've switched to that and rebased patch 8 on it.

Post by Jason Ekstrand
+
+ vtn_add_case(b, swtch, break_block, block_id,

literal, is_default);

Post by Jason Ekstrand
+ is_default = false;
+ }
/* Now, we go through and walk the blocks. While we

walk through

Post by Jason Ekstrand
* the blocks, we also gather the much-needed fall-through

Jason Ekstrand

2017-12-07 16:12:15 UTC

Permalink

Bugzilla: https://bugs.freedesktop.org/show_bug.cgi?id=101560
Cc: Juan A. Suarez Romero <***@igalia.com>
---
src/compiler/spirv/vtn_cfg.c | 33 +++++++++++++++++++++++++--------
1 file changed, 25 insertions(+), 8 deletions(-)

diff --git a/src/compiler/spirv/vtn_cfg.c b/src/compiler/spirv/vtn_cfg.c
index 9d1ca84..3befc1f 100644
--- a/src/compiler/spirv/vtn_cfg.c
+++ b/src/compiler/spirv/vtn_cfg.c
@@ -170,7 +170,7 @@ vtn_cfg_handle_prepass_instruction(struct vtn_builder *b, SpvOp opcode,
static void
vtn_add_case(struct vtn_builder *b, struct vtn_switch *swtch,
struct vtn_block *break_block,
- uint32_t block_id, uint32_t val, bool is_default)
+ uint32_t block_id, uint64_t val, bool is_default)
{
struct vtn_block *case_block =
vtn_value(b, block_id, vtn_value_type_block)->block;
@@ -197,7 +197,7 @@ vtn_add_case(struct vtn_builder *b, struct vtn_switch *swtch,
if (is_default) {
case_block->switch_case->is_default = true;
} else {
- util_dynarray_append(&case_block->switch_case->values, uint32_t, val);
+ util_dynarray_append(&case_block->switch_case->values, uint64_t, val);
}
}

@@ -425,12 +425,29 @@ vtn_cfg_walk_blocks(struct vtn_builder *b, struct list_head *cf_list,
const uint32_t *branch_end =
block->branch + (block->branch[0] >> SpvWordCountShift);

+ struct vtn_value *cond_val = vtn_untyped_value(b, block->branch[1]);
+ vtn_fail_if(!cond_val->type ||
+ cond_val->type->base_type != vtn_base_type_scalar,
+ "Selector of OpSelect must have a type of OpTypeInt");
+
+ nir_alu_type cond_type =
+ nir_get_nir_type_for_glsl_type(cond_val->type->type);
+ vtn_fail_if(nir_alu_type_get_base_type(cond_type) != nir_type_int &&
+ nir_alu_type_get_base_type(cond_type) != nir_type_uint,
+ "Selector of OpSelect must have a type of OpTypeInt");
+
bool is_default = true;
for (const uint32_t *w = block->branch + 2; w < branch_end;) {
- uint32_t literal = 0;
+ uint64_t literal = 0;
if (!is_default) {
- literal = *w;
- w++;
+ if (nir_alu_type_get_type_size(cond_type) <= 32) {
+ literal = *w;
+ w++;
+ } else {
+ assert(nir_alu_type_get_type_size(cond_type) == 64);
+ literal = vtn_u64_literal(w);
+ w += 2;
+ }
}

uint32_t block_id = *w;
@@ -730,9 +747,9 @@ vtn_emit_cf_list(struct vtn_builder *b, struct list_head *cf_list,
}

nir_ssa_def *cond = NULL;
- util_dynarray_foreach(&cse->values, uint32_t, val) {
- nir_ssa_def *is_val =
- nir_ieq(&b->nb, sel, nir_imm_int(&b->nb, *val));
+ util_dynarray_foreach(&cse->values, uint64_t, val) {
+ nir_ssa_def *imm = nir_imm_intN_t(&b->nb, *val, sel->bit_size);
+ nir_ssa_def *is_val = nir_ieq(&b->nb, sel, imm);

cond = cond ? nir_ior(&b->nb, cond, is_val) : is_val;
}

--
2.5.0.400.gff86faf

Ian Romanick

2017-12-11 18:12:52 UTC

Permalink

This patch is

Post by Jason Ekstrand
Bugzilla: https://bugs.freedesktop.org/show_bug.cgi?id=101560
---
src/compiler/spirv/vtn_cfg.c | 33 +++++++++++++++++++++++++--------
1 file changed, 25 insertions(+), 8 deletions(-)
diff --git a/src/compiler/spirv/vtn_cfg.c b/src/compiler/spirv/vtn_cfg.c
index 9d1ca84..3befc1f 100644
--- a/src/compiler/spirv/vtn_cfg.c
+++ b/src/compiler/spirv/vtn_cfg.c
@@ -170,7 +170,7 @@ vtn_cfg_handle_prepass_instruction(struct vtn_builder *b, SpvOp opcode,
static void
vtn_add_case(struct vtn_builder *b, struct vtn_switch *swtch,
struct vtn_block *break_block,
- uint32_t block_id, uint32_t val, bool is_default)
+ uint32_t block_id, uint64_t val, bool is_default)
{
struct vtn_block *case_block =
vtn_value(b, block_id, vtn_value_type_block)->block;
@@ -197,7 +197,7 @@ vtn_add_case(struct vtn_builder *b, struct vtn_switch *swtch,
if (is_default) {
case_block->switch_case->is_default = true;
} else {
- util_dynarray_append(&case_block->switch_case->values, uint32_t, val);
+ util_dynarray_append(&case_block->switch_case->values, uint64_t, val);
}
}
@@ -425,12 +425,29 @@ vtn_cfg_walk_blocks(struct vtn_builder *b, struct list_head *cf_list,
const uint32_t *branch_end =
block->branch + (block->branch[0] >> SpvWordCountShift);
+ struct vtn_value *cond_val = vtn_untyped_value(b, block->branch[1]);
+ vtn_fail_if(!cond_val->type ||
+ cond_val->type->base_type != vtn_base_type_scalar,
+ "Selector of OpSelect must have a type of OpTypeInt");
+
+ nir_alu_type cond_type =
+ nir_get_nir_type_for_glsl_type(cond_val->type->type);
+ vtn_fail_if(nir_alu_type_get_base_type(cond_type) != nir_type_int &&
+ nir_alu_type_get_base_type(cond_type) != nir_type_uint,
+ "Selector of OpSelect must have a type of OpTypeInt");
+
bool is_default = true;
for (const uint32_t *w = block->branch + 2; w < branch_end;) {
- uint32_t literal = 0;
+ uint64_t literal = 0;
if (!is_default) {
- literal = *w;
- w++;
+ if (nir_alu_type_get_type_size(cond_type) <= 32) {
+ literal = *w;
+ w++;
+ } else {
+ assert(nir_alu_type_get_type_size(cond_type) == 64);
+ literal = vtn_u64_literal(w);
+ w += 2;
+ }
}
uint32_t block_id = *w;
@@ -730,9 +747,9 @@ vtn_emit_cf_list(struct vtn_builder *b, struct list_head *cf_list,
}
nir_ssa_def *cond = NULL;
- util_dynarray_foreach(&cse->values, uint32_t, val) {
- nir_ssa_def *is_val =
- nir_ieq(&b->nb, sel, nir_imm_int(&b->nb, *val));
+ util_dynarray_foreach(&cse->values, uint64_t, val) {
+ nir_ssa_def *imm = nir_imm_intN_t(&b->nb, *val, sel->bit_size);
+ nir_ssa_def *is_val = nir_ieq(&b->nb, sel, imm);
cond = cond ? nir_ior(&b->nb, cond, is_val) : is_val;
}