fdo-mirrors/mesa
1
0
Fork 0
mirror of https://gitlab.freedesktop.org/mesa/mesa.git synced 2025-12-21 22:20:14 +01:00
Code Issues Projects Releases Packages Wiki Activity Actions 1

radv: use nir_op_imm helpers

Browse source 
Signed-off-by: Rhys Perry <pendingchaos02@gmail.com>
Acked-by: Samuel Pitoiset <samuel.pitoiset@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/15854>
This commit is contained in:
Rhys Perry 2022-03-25 19:53:38 +00:00 • committed by Marge Bot
parent df994ecc8f
commit b62e90ad43
15 changed files with 505 additions and 662 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

306
src/amd/vulkan/radv_acceleration_structure.c
View file

@ -752,14 +752,13 @@ get_indices(nir_builder *b, nir_ssa_def *addr, nir_ssa_def *type, nir_ssa_def *i
nir_variable_create(b->shader, nir_var_shader_temp, uvec3_type, "indices");
nir_push_if(b, nir_ult(b, type, nir_imm_int(b, 2)));
nir_push_if(b, nir_ieq(b, type, nir_imm_int(b, VK_INDEX_TYPE_UINT16)));
nir_push_if(b, nir_ieq_imm(b, type, VK_INDEX_TYPE_UINT16));
{
nir_ssa_def *index_id = nir_umul24(b, id, nir_imm_int(b, 6));
nir_ssa_def *indices[3];
for (unsigned i = 0; i < 3; ++i) {
indices[i] = nir_build_load_global(
b, 1, 16,
nir_iadd(b, addr, nir_u2u64(b, nir_iadd(b, index_id, nir_imm_int(b, 2 * i)))));
b, 1, 16, nir_iadd(b, addr, nir_u2u64(b, nir_iadd_imm(b, index_id, 2 * i))));
}
nir_store_var(b, result, nir_u2u32(b, nir_vec(b, indices, 3)), 7);
}
@ -776,11 +775,11 @@ get_indices(nir_builder *b, nir_ssa_def *addr, nir_ssa_def *type, nir_ssa_def *i
nir_ssa_def *index_id = nir_umul24(b, id, nir_imm_int(b, 3));
nir_ssa_def *indices[] = {
index_id,
nir_iadd(b, index_id, nir_imm_int(b, 1)),
nir_iadd(b, index_id, nir_imm_int(b, 2)),
nir_iadd_imm(b, index_id, 1),
nir_iadd_imm(b, index_id, 2),
};
nir_push_if(b, nir_ieq(b, type, nir_imm_int(b, VK_INDEX_TYPE_NONE_KHR)));
nir_push_if(b, nir_ieq_imm(b, type, VK_INDEX_TYPE_NONE_KHR));
{
nir_store_var(b, result, nir_vec(b, indices, 3), 7);
}
@ -827,7 +826,7 @@ get_vertices(nir_builder *b, nir_ssa_def *addresses, nir_ssa_def *format, nir_ss
for (unsigned f = 0; f < ARRAY_SIZE(formats); ++f) {
if (f + 1 < ARRAY_SIZE(formats))
nir_push_if(b, nir_ieq(b, format, nir_imm_int(b, formats[f])));
nir_push_if(b, nir_ieq_imm(b, format, formats[f]));
for (unsigned i = 0; i < 3; ++i) {
switch (formats[f]) {
@ -863,8 +862,8 @@ get_vertices(nir_builder *b, nir_ssa_def *addresses, nir_ssa_def *format, nir_ss
values[j] = nir_ubfe(b, val, nir_imm_int(b, j * 10), nir_imm_int(b, 10));
} else {
for (unsigned j = 0; j < components; ++j)
values[j] = nir_build_load_global(
b, 1, comp_bits, nir_iadd(b, addr, nir_imm_int64(b, j * comp_bytes)));
values[j] =
nir_build_load_global(b, 1, comp_bits, nir_iadd_imm(b, addr, j * comp_bytes));
for (unsigned j = components; j < 3; ++j)
values[j] = nir_imm_intN_t(b, 0, comp_bits);
@ -1007,7 +1006,7 @@ build_leaf_shader(struct radv_device *dev)
&b, scratch_addr,
nir_u2u64(&b, nir_iadd(&b, scratch_offset, nir_umul24(&b, global_id, nir_imm_int(&b, 4)))));
nir_push_if(&b, nir_ieq(&b, geom_type, nir_imm_int(&b, VK_GEOMETRY_TYPE_TRIANGLES_KHR)));
nir_push_if(&b, nir_ieq_imm(&b, geom_type, VK_GEOMETRY_TYPE_TRIANGLES_KHR));
{ /* Triangles */
nir_ssa_def *vertex_addr = nir_pack_64_2x32(&b, nir_channels(&b, pconst2, 3));
nir_ssa_def *index_addr = nir_pack_64_2x32(&b, nir_channels(&b, pconst2, 12));
@ -1040,18 +1039,13 @@ build_leaf_shader(struct radv_device *dev)
nir_store_var(&b, transform[1], nir_imm_vec4(&b, 0.0, 1.0, 0.0, 0.0), 0xf);
nir_store_var(&b, transform[2], nir_imm_vec4(&b, 0.0, 0.0, 1.0, 0.0), 0xf);
nir_push_if(&b, nir_ine(&b, transform_addr, nir_imm_int64(&b, 0)));
nir_store_var(
&b, transform[0],
nir_build_load_global(&b, 4, 32, nir_iadd(&b, transform_addr, nir_imm_int64(&b, 0))), 0xf);
nir_store_var(
&b, transform[1],
nir_build_load_global(&b, 4, 32, nir_iadd(&b, transform_addr, nir_imm_int64(&b, 16))),
0xf);
nir_store_var(
&b, transform[2],
nir_build_load_global(&b, 4, 32, nir_iadd(&b, transform_addr, nir_imm_int64(&b, 32))),
0xf);
nir_push_if(&b, nir_ine_imm(&b, transform_addr, 0));
nir_store_var(&b, transform[0],
nir_build_load_global(&b, 4, 32, nir_iadd_imm(&b, transform_addr, 0)), 0xf);
nir_store_var(&b, transform[1],
nir_build_load_global(&b, 4, 32, nir_iadd_imm(&b, transform_addr, 16)), 0xf);
nir_store_var(&b, transform[2],
nir_build_load_global(&b, 4, 32, nir_iadd_imm(&b, transform_addr, 32)), 0xf);
nir_pop_if(&b, NULL);
for (unsigned i = 0; i < 3; ++i)
@ -1067,15 +1061,14 @@ build_leaf_shader(struct radv_device *dev)
for (unsigned i = 0; i < 4; ++i) {
nir_build_store_global(&b, nir_vec(&b, node_data + i * 4, 4),
nir_iadd(&b, triangle_node_dst_addr, nir_imm_int64(&b, i * 16)),
.align_mul = 16);
nir_iadd_imm(&b, triangle_node_dst_addr, i * 16), .align_mul = 16);
}
nir_ssa_def *node_id = nir_ushr(&b, node_offset, nir_imm_int(&b, 3));
nir_ssa_def *node_id = nir_ushr_imm(&b, node_offset, 3);
nir_build_store_global(&b, node_id, scratch_addr);
}
nir_push_else(&b, NULL);
nir_push_if(&b, nir_ieq(&b, geom_type, nir_imm_int(&b, VK_GEOMETRY_TYPE_AABBS_KHR)));
nir_push_if(&b, nir_ieq_imm(&b, geom_type, VK_GEOMETRY_TYPE_AABBS_KHR));
{ /* AABBs */
nir_ssa_def *aabb_addr = nir_pack_64_2x32(&b, nir_channels(&b, pconst2, 3));
nir_ssa_def *aabb_stride = nir_channel(&b, pconst2, 2);
@ -1083,16 +1076,13 @@ build_leaf_shader(struct radv_device *dev)
nir_ssa_def *node_offset =
nir_iadd(&b, node_dst_offset, nir_umul24(&b, global_id, nir_imm_int(&b, 64)));
nir_ssa_def *aabb_node_dst_addr = nir_iadd(&b, node_dst_addr, nir_u2u64(&b, node_offset));
nir_ssa_def *node_id =
nir_iadd(&b, nir_ushr(&b, node_offset, nir_imm_int(&b, 3)), nir_imm_int(&b, 7));
nir_ssa_def *node_id = nir_iadd_imm(&b, nir_ushr_imm(&b, node_offset, 3), 7);
nir_build_store_global(&b, node_id, scratch_addr);
aabb_addr = nir_iadd(&b, aabb_addr, nir_u2u64(&b, nir_imul(&b, aabb_stride, global_id)));
nir_ssa_def *min_bound =
nir_build_load_global(&b, 3, 32, nir_iadd(&b, aabb_addr, nir_imm_int64(&b, 0)));
nir_ssa_def *max_bound =
nir_build_load_global(&b, 3, 32, nir_iadd(&b, aabb_addr, nir_imm_int64(&b, 12)));
nir_ssa_def *min_bound = nir_build_load_global(&b, 3, 32, nir_iadd_imm(&b, aabb_addr, 0));
nir_ssa_def *max_bound = nir_build_load_global(&b, 3, 32, nir_iadd_imm(&b, aabb_addr, 12));
nir_ssa_def *values[] = {nir_channel(&b, min_bound, 0),
nir_channel(&b, min_bound, 1),
@ -1104,21 +1094,19 @@ build_leaf_shader(struct radv_device *dev)
geometry_id};
nir_build_store_global(&b, nir_vec(&b, values + 0, 4),
nir_iadd(&b, aabb_node_dst_addr, nir_imm_int64(&b, 0)),
.align_mul = 16);
nir_iadd_imm(&b, aabb_node_dst_addr, 0), .align_mul = 16);
nir_build_store_global(&b, nir_vec(&b, values + 4, 4),
nir_iadd(&b, aabb_node_dst_addr, nir_imm_int64(&b, 16)),
.align_mul = 16);
nir_iadd_imm(&b, aabb_node_dst_addr, 16), .align_mul = 16);
}
nir_push_else(&b, NULL);
{ /* Instances */
nir_variable *instance_addr_var =
nir_variable_create(b.shader, nir_var_shader_temp, glsl_uint64_t_type(), "instance_addr");
nir_push_if(&b, nir_ine(&b, nir_channel(&b, pconst2, 2), nir_imm_int(&b, 0)));
nir_push_if(&b, nir_ine_imm(&b, nir_channel(&b, pconst2, 2), 0));
{
nir_ssa_def *ptr = nir_iadd(&b, nir_pack_64_2x32(&b, nir_channels(&b, pconst2, 3)),
nir_u2u64(&b, nir_imul(&b, global_id, nir_imm_int(&b, 8))));
nir_u2u64(&b, nir_imul_imm(&b, global_id, 8)));
nir_ssa_def *addr =
nir_pack_64_2x32(&b, nir_build_load_global(&b, 2, 32, ptr, .align_mul = 8));
nir_store_var(&b, instance_addr_var, addr, 1);
@ -1126,24 +1114,22 @@ build_leaf_shader(struct radv_device *dev)
nir_push_else(&b, NULL);
{
nir_ssa_def *addr = nir_iadd(&b, nir_pack_64_2x32(&b, nir_channels(&b, pconst2, 3)),
nir_u2u64(&b, nir_imul(&b, global_id, nir_imm_int(&b, 64))));
nir_u2u64(&b, nir_imul_imm(&b, global_id, 64)));
nir_store_var(&b, instance_addr_var, addr, 1);
}
nir_pop_if(&b, NULL);
nir_ssa_def *instance_addr = nir_load_var(&b, instance_addr_var);
nir_ssa_def *inst_transform[] = {
nir_build_load_global(&b, 4, 32, nir_iadd(&b, instance_addr, nir_imm_int64(&b, 0))),
nir_build_load_global(&b, 4, 32, nir_iadd(&b, instance_addr, nir_imm_int64(&b, 16))),
nir_build_load_global(&b, 4, 32, nir_iadd(&b, instance_addr, nir_imm_int64(&b, 32)))};
nir_ssa_def *inst3 =
nir_build_load_global(&b, 4, 32, nir_iadd(&b, instance_addr, nir_imm_int64(&b, 48)));
nir_build_load_global(&b, 4, 32, nir_iadd_imm(&b, instance_addr, 0)),
nir_build_load_global(&b, 4, 32, nir_iadd_imm(&b, instance_addr, 16)),
nir_build_load_global(&b, 4, 32, nir_iadd_imm(&b, instance_addr, 32))};
nir_ssa_def *inst3 = nir_build_load_global(&b, 4, 32, nir_iadd_imm(&b, instance_addr, 48));
nir_ssa_def *node_offset =
nir_iadd(&b, node_dst_offset, nir_umul24(&b, global_id, nir_imm_int(&b, 128)));
node_dst_addr = nir_iadd(&b, node_dst_addr, nir_u2u64(&b, node_offset));
nir_ssa_def *node_id =
nir_iadd(&b, nir_ushr(&b, node_offset, nir_imm_int(&b, 3)), nir_imm_int(&b, 6));
nir_ssa_def *node_id = nir_iadd_imm(&b, nir_ushr_imm(&b, node_offset, 3), 6);
nir_build_store_global(&b, node_id, scratch_addr);
nir_variable *bounds[2] = {
@ -1155,13 +1141,11 @@ build_leaf_shader(struct radv_device *dev)
nir_store_var(&b, bounds[1], nir_channels(&b, nir_imm_vec4(&b, NAN, NAN, NAN, NAN), 7), 7);
nir_ssa_def *header_addr = nir_pack_64_2x32(&b, nir_channels(&b, inst3, 12));
nir_push_if(&b, nir_ine(&b, header_addr, nir_imm_int64(&b, 0)));
nir_push_if(&b, nir_ine_imm(&b, header_addr, 0));
nir_ssa_def *header_root_offset =
nir_build_load_global(&b, 1, 32, nir_iadd(&b, header_addr, nir_imm_int64(&b, 0)));
nir_ssa_def *header_min =
nir_build_load_global(&b, 3, 32, nir_iadd(&b, header_addr, nir_imm_int64(&b, 8)));
nir_ssa_def *header_max =
nir_build_load_global(&b, 3, 32, nir_iadd(&b, header_addr, nir_imm_int64(&b, 20)));
nir_build_load_global(&b, 1, 32, nir_iadd_imm(&b, header_addr, 0));
nir_ssa_def *header_min = nir_build_load_global(&b, 3, 32, nir_iadd_imm(&b, header_addr, 8));
nir_ssa_def *header_max = nir_build_load_global(&b, 3, 32, nir_iadd_imm(&b, header_addr, 20));
nir_ssa_def *bound_defs[2][3];
for (unsigned i = 0; i < 3; ++i) {
@ -1187,7 +1171,7 @@ build_leaf_shader(struct radv_device *dev)
vals[j] = nir_channel(&b, inst_transform[j], i);
nir_build_store_global(&b, nir_vec(&b, vals, 3),
nir_iadd(&b, node_dst_addr, nir_imm_int64(&b, 92 + 12 * i)));
nir_iadd_imm(&b, node_dst_addr, 92 + 12 * i));
}
nir_ssa_def *m_in[3][3], *m_out[3][3], *m_vec[3][4];
@ -1203,21 +1187,18 @@ build_leaf_shader(struct radv_device *dev)
for (unsigned i = 0; i < 3; ++i) {
nir_build_store_global(&b, nir_vec(&b, m_vec[i], 4),
nir_iadd(&b, node_dst_addr, nir_imm_int64(&b, 16 + 16 * i)));
nir_iadd_imm(&b, node_dst_addr, 16 + 16 * i));
}
nir_ssa_def *out0[4] = {
nir_ior(&b, nir_channel(&b, nir_unpack_64_2x32(&b, header_addr), 0), header_root_offset),
nir_channel(&b, nir_unpack_64_2x32(&b, header_addr), 1), nir_channel(&b, inst3, 0),
nir_channel(&b, inst3, 1)};
nir_build_store_global(&b, nir_vec(&b, out0, 4),
nir_iadd(&b, node_dst_addr, nir_imm_int64(&b, 0)));
nir_build_store_global(&b, global_id, nir_iadd(&b, node_dst_addr, nir_imm_int64(&b, 88)));
nir_build_store_global(&b, nir_vec(&b, out0, 4), nir_iadd_imm(&b, node_dst_addr, 0));
nir_build_store_global(&b, global_id, nir_iadd_imm(&b, node_dst_addr, 88));
nir_pop_if(&b, NULL);
nir_build_store_global(&b, nir_load_var(&b, bounds[0]),
nir_iadd(&b, node_dst_addr, nir_imm_int64(&b, 64)));
nir_build_store_global(&b, nir_load_var(&b, bounds[1]),
nir_iadd(&b, node_dst_addr, nir_imm_int64(&b, 76)));
nir_build_store_global(&b, nir_load_var(&b, bounds[0]), nir_iadd_imm(&b, node_dst_addr, 64));
nir_build_store_global(&b, nir_load_var(&b, bounds[1]), nir_iadd_imm(&b, node_dst_addr, 76));
}
nir_pop_if(&b, NULL);
nir_pop_if(&b, NULL);
@ -1229,17 +1210,15 @@ static void
determine_bounds(nir_builder *b, nir_ssa_def *node_addr, nir_ssa_def *node_id,
nir_variable *bounds_vars[2])
{
nir_ssa_def *node_type = nir_iand(b, node_id, nir_imm_int(b, 7));
node_addr = nir_iadd(
b, node_addr,
nir_u2u64(b, nir_ishl(b, nir_iand(b, node_id, nir_imm_int(b, ~7u)), nir_imm_int(b, 3))));
nir_ssa_def *node_type = nir_iand_imm(b, node_id, 7);
node_addr =
nir_iadd(b, node_addr, nir_u2u64(b, nir_ishl_imm(b, nir_iand_imm(b, node_id, ~7u), 3)));
nir_push_if(b, nir_ieq(b, node_type, nir_imm_int(b, 0)));
nir_push_if(b, nir_ieq_imm(b, node_type, 0));
{
nir_ssa_def *positions[3];
for (unsigned i = 0; i < 3; ++i)
positions[i] =
nir_build_load_global(b, 3, 32, nir_iadd(b, node_addr, nir_imm_int64(b, i * 12)));
positions[i] = nir_build_load_global(b, 3, 32, nir_iadd_imm(b, node_addr, i * 12));
nir_ssa_def *bounds[] = {positions[0], positions[0]};
for (unsigned i = 1; i < 3; ++i) {
bounds[0] = nir_fmin(b, bounds[0], positions[i]);
@ -1249,13 +1228,13 @@ determine_bounds(nir_builder *b, nir_ssa_def *node_addr, nir_ssa_def *node_id,
nir_store_var(b, bounds_vars[1], bounds[1], 7);
}
nir_push_else(b, NULL);
nir_push_if(b, nir_ieq(b, node_type, nir_imm_int(b, 5)));
nir_push_if(b, nir_ieq_imm(b, node_type, 5));
{
nir_ssa_def *input_bounds[4][2];
for (unsigned i = 0; i < 4; ++i)
for (unsigned j = 0; j < 2; ++j)
input_bounds[i][j] = nir_build_load_global(
b, 3, 32, nir_iadd(b, node_addr, nir_imm_int64(b, 16 + i * 24 + j * 12)));
input_bounds[i][j] =
nir_build_load_global(b, 3, 32, nir_iadd_imm(b, node_addr, 16 + i * 24 + j * 12));
nir_ssa_def *bounds[] = {input_bounds[0][0], input_bounds[0][1]};
for (unsigned i = 1; i < 4; ++i) {
bounds[0] = nir_fmin(b, bounds[0], input_bounds[i][0]);
@ -1266,12 +1245,11 @@ determine_bounds(nir_builder *b, nir_ssa_def *node_addr, nir_ssa_def *node_id,
nir_store_var(b, bounds_vars[1], bounds[1], 7);
}
nir_push_else(b, NULL);
nir_push_if(b, nir_ieq(b, node_type, nir_imm_int(b, 6)));
nir_push_if(b, nir_ieq_imm(b, node_type, 6));
{ /* Instances */
nir_ssa_def *bounds[2];
for (unsigned i = 0; i < 2; ++i)
bounds[i] =
nir_build_load_global(b, 3, 32, nir_iadd(b, node_addr, nir_imm_int64(b, 64 + i * 12)));
bounds[i] = nir_build_load_global(b, 3, 32, nir_iadd_imm(b, node_addr, 64 + i * 12));
nir_store_var(b, bounds_vars[0], bounds[0], 7);
nir_store_var(b, bounds_vars[1], bounds[1], 7);
}
@ -1279,8 +1257,7 @@ determine_bounds(nir_builder *b, nir_ssa_def *node_addr, nir_ssa_def *node_id,
{ /* AABBs */
nir_ssa_def *bounds[2];
for (unsigned i = 0; i < 2; ++i)
bounds[i] =
nir_build_load_global(b, 3, 32, nir_iadd(b, node_addr, nir_imm_int64(b, i * 12)));
bounds[i] = nir_build_load_global(b, 3, 32, nir_iadd_imm(b, node_addr, i * 12));
nir_store_var(b, bounds_vars[0], bounds[0], 7);
nir_store_var(b, bounds_vars[1], bounds[1], 7);
}
@ -1316,30 +1293,26 @@ build_internal_shader(struct radv_device *dev)
nir_ssa_def *node_dst_offset = nir_channel(&b, pconst1, 0);
nir_ssa_def *dst_scratch_offset = nir_channel(&b, pconst1, 1);
nir_ssa_def *src_scratch_offset = nir_channel(&b, pconst1, 2);
nir_ssa_def *src_node_count =
nir_iand(&b, nir_channel(&b, pconst1, 3), nir_imm_int(&b, 0x7FFFFFFFU));
nir_ssa_def *src_node_count = nir_iand_imm(&b, nir_channel(&b, pconst1, 3), 0x7FFFFFFFU);
nir_ssa_def *fill_header =
nir_ine(&b, nir_iand(&b, nir_channel(&b, pconst1, 3), nir_imm_int(&b, 0x80000000U)),
nir_imm_int(&b, 0));
nir_ine_imm(&b, nir_iand_imm(&b, nir_channel(&b, pconst1, 3), 0x80000000U), 0);
nir_ssa_def *global_id =
nir_iadd(&b,
nir_umul24(&b, nir_channels(&b, nir_load_workgroup_id(&b, 32), 1),
nir_imm_int(&b, b.shader->info.workgroup_size[0])),
nir_channels(&b, nir_load_local_invocation_id(&b), 1));
nir_ssa_def *src_idx = nir_imul(&b, global_id, nir_imm_int(&b, 4));
nir_ssa_def *src_idx = nir_imul_imm(&b, global_id, 4);
nir_ssa_def *src_count = nir_umin(&b, nir_imm_int(&b, 4), nir_isub(&b, src_node_count, src_idx));
nir_ssa_def *node_offset =
nir_iadd(&b, node_dst_offset, nir_ishl(&b, global_id, nir_imm_int(&b, 7)));
nir_ssa_def *node_offset = nir_iadd(&b, node_dst_offset, nir_ishl_imm(&b, global_id, 7));
nir_ssa_def *node_dst_addr = nir_iadd(&b, node_addr, nir_u2u64(&b, node_offset));
nir_ssa_def *src_nodes = nir_build_load_global(
&b, 4, 32,
nir_iadd(&b, scratch_addr,
nir_u2u64(&b, nir_iadd(&b, src_scratch_offset,
nir_ishl(&b, global_id, nir_imm_int(&b, 4))))));
nir_u2u64(&b, nir_iadd(&b, src_scratch_offset, nir_ishl_imm(&b, global_id, 4)))));
nir_build_store_global(&b, src_nodes, nir_iadd(&b, node_dst_addr, nir_imm_int64(&b, 0)));
nir_build_store_global(&b, src_nodes, nir_iadd_imm(&b, node_dst_addr, 0));
nir_ssa_def *total_bounds[2] = {
nir_channels(&b, nir_imm_vec4(&b, NAN, NAN, NAN, NAN), 7),
@ -1358,24 +1331,23 @@ build_internal_shader(struct radv_device *dev)
determine_bounds(&b, node_addr, nir_channel(&b, src_nodes, i), bounds);
nir_pop_if(&b, NULL);
nir_build_store_global(&b, nir_load_var(&b, bounds[0]),
nir_iadd(&b, node_dst_addr, nir_imm_int64(&b, 16 + 24 * i)));
nir_iadd_imm(&b, node_dst_addr, 16 + 24 * i));
nir_build_store_global(&b, nir_load_var(&b, bounds[1]),
nir_iadd(&b, node_dst_addr, nir_imm_int64(&b, 28 + 24 * i)));
nir_iadd_imm(&b, node_dst_addr, 28 + 24 * i));
total_bounds[0] = nir_fmin(&b, total_bounds[0], nir_load_var(&b, bounds[0]));
total_bounds[1] = nir_fmax(&b, total_bounds[1], nir_load_var(&b, bounds[1]));
}
nir_ssa_def *node_id =
nir_iadd(&b, nir_ushr(&b, node_offset, nir_imm_int(&b, 3)), nir_imm_int(&b, 5));
nir_ssa_def *dst_scratch_addr = nir_iadd(
&b, scratch_addr,
nir_u2u64(&b, nir_iadd(&b, dst_scratch_offset, nir_ishl(&b, global_id, nir_imm_int(&b, 2)))));
nir_ssa_def *node_id = nir_iadd_imm(&b, nir_ushr_imm(&b, node_offset, 3), 5);
nir_ssa_def *dst_scratch_addr =
nir_iadd(&b, scratch_addr,
nir_u2u64(&b, nir_iadd(&b, dst_scratch_offset, nir_ishl_imm(&b, global_id, 2))));
nir_build_store_global(&b, node_id, dst_scratch_addr);
nir_push_if(&b, fill_header);
nir_build_store_global(&b, node_id, node_addr);
nir_build_store_global(&b, total_bounds[0], nir_iadd(&b, node_addr, nir_imm_int64(&b, 8)));
nir_build_store_global(&b, total_bounds[1], nir_iadd(&b, node_addr, nir_imm_int64(&b, 20)));
nir_build_store_global(&b, total_bounds[0], nir_iadd_imm(&b, node_addr, 8));
nir_build_store_global(&b, total_bounds[1], nir_iadd_imm(&b, node_addr, 20));
nir_pop_if(&b, NULL);
return b.shader;
}
@ -1409,11 +1381,11 @@ build_copy_shader(struct radv_device *dev)
nir_variable *offset_var =
nir_variable_create(b.shader, nir_var_shader_temp, glsl_uint_type(), "offset");
nir_ssa_def *offset = nir_imul(&b, global_id, nir_imm_int(&b, 16));
nir_ssa_def *offset = nir_imul_imm(&b, global_id, 16);
nir_store_var(&b, offset_var, offset, 1);
nir_ssa_def *increment = nir_imul(&b, nir_channel(&b, nir_load_num_workgroups(&b, 32), 0),
nir_imm_int(&b, b.shader->info.workgroup_size[0] * 16));
nir_ssa_def *increment = nir_imul_imm(&b, nir_channel(&b, nir_load_num_workgroups(&b, 32), 0),
b.shader->info.workgroup_size[0] * 16);
nir_ssa_def *pconst0 =
nir_load_push_constant(&b, 4, 32, nir_imm_int(&b, 0), .base = 0, .range = 16);
@ -1436,98 +1408,87 @@ build_copy_shader(struct radv_device *dev)
nir_variable *value_var =
nir_variable_create(b.shader, nir_var_shader_temp, glsl_vec4_type(), "value");
nir_push_if(&b, nir_ieq(&b, mode, nir_imm_int(&b, COPY_MODE_SERIALIZE)));
nir_push_if(&b, nir_ieq_imm(&b, mode, COPY_MODE_SERIALIZE));
{
nir_ssa_def *instance_count = nir_build_load_global(
&b, 1, 32,
nir_iadd(&b, src_base_addr,
nir_imm_int64(&b, offsetof(struct radv_accel_struct_header, instance_count))));
nir_iadd_imm(&b, src_base_addr,
offsetof(struct radv_accel_struct_header, instance_count)));
nir_ssa_def *compacted_size = nir_build_load_global(
&b, 1, 64,
nir_iadd(&b, src_base_addr,
nir_imm_int64(&b, offsetof(struct radv_accel_struct_header, compacted_size))));
nir_iadd_imm(&b, src_base_addr,
offsetof(struct radv_accel_struct_header, compacted_size)));
nir_ssa_def *serialization_size = nir_build_load_global(
&b, 1, 64,
nir_iadd(
&b, src_base_addr,
nir_imm_int64(&b, offsetof(struct radv_accel_struct_header, serialization_size))));
nir_iadd_imm(&b, src_base_addr,
offsetof(struct radv_accel_struct_header, serialization_size)));
nir_store_var(&b, compacted_size_var, compacted_size, 1);
nir_store_var(
&b, instance_offset_var,
nir_build_load_global(&b, 1, 32,
nir_iadd(&b, src_base_addr,
nir_imm_int64(&b, offsetof(struct radv_accel_struct_header,
instance_offset)))),
1);
nir_store_var(&b, instance_offset_var,
nir_build_load_global(
&b, 1, 32,
nir_iadd_imm(&b, src_base_addr,
offsetof(struct radv_accel_struct_header, instance_offset))),
1);
nir_store_var(&b, instance_count_var, instance_count, 1);
nir_ssa_def *dst_offset =
nir_iadd(&b, nir_imm_int(&b, sizeof(struct radv_accel_struct_serialization_header)),
nir_imul(&b, instance_count, nir_imm_int(&b, sizeof(uint64_t))));
nir_ssa_def *dst_offset = nir_iadd_imm(&b, nir_imul_imm(&b, instance_count, sizeof(uint64_t)),
sizeof(struct radv_accel_struct_serialization_header));
nir_store_var(&b, src_offset_var, nir_imm_int(&b, 0), 1);
nir_store_var(&b, dst_offset_var, dst_offset, 1);
nir_push_if(&b, nir_ieq(&b, global_id, nir_imm_int(&b, 0)));
nir_push_if(&b, nir_ieq_imm(&b, global_id, 0));
{
nir_build_store_global(
&b, serialization_size,
nir_iadd(&b, dst_base_addr,
nir_imm_int64(&b, offsetof(struct radv_accel_struct_serialization_header,
serialization_size))));
nir_build_store_global(&b, serialization_size,
nir_iadd_imm(&b, dst_base_addr,
offsetof(struct radv_accel_struct_serialization_header,
serialization_size)));
nir_build_store_global(
&b, compacted_size,
nir_iadd(&b, dst_base_addr,
nir_imm_int64(&b, offsetof(struct radv_accel_struct_serialization_header,
compacted_size))));
nir_iadd_imm(&b, dst_base_addr,
offsetof(struct radv_accel_struct_serialization_header, compacted_size)));
nir_build_store_global(
&b, nir_u2u64(&b, instance_count),
nir_iadd(&b, dst_base_addr,
nir_imm_int64(&b, offsetof(struct radv_accel_struct_serialization_header,
instance_count))));
nir_iadd_imm(&b, dst_base_addr,
offsetof(struct radv_accel_struct_serialization_header, instance_count)));
}
nir_pop_if(&b, NULL);
}
nir_push_else(&b, NULL);
nir_push_if(&b, nir_ieq(&b, mode, nir_imm_int(&b, COPY_MODE_DESERIALIZE)));
nir_push_if(&b, nir_ieq_imm(&b, mode, COPY_MODE_DESERIALIZE));
{
nir_ssa_def *instance_count = nir_build_load_global(
&b, 1, 32,
nir_iadd(&b, src_base_addr,
nir_imm_int64(
&b, offsetof(struct radv_accel_struct_serialization_header, instance_count))));
nir_ssa_def *src_offset =
nir_iadd(&b, nir_imm_int(&b, sizeof(struct radv_accel_struct_serialization_header)),
nir_imul(&b, instance_count, nir_imm_int(&b, sizeof(uint64_t))));
nir_iadd_imm(&b, src_base_addr,
offsetof(struct radv_accel_struct_serialization_header, instance_count)));
nir_ssa_def *src_offset = nir_iadd_imm(&b, nir_imul_imm(&b, instance_count, sizeof(uint64_t)),
sizeof(struct radv_accel_struct_serialization_header));
nir_ssa_def *header_addr = nir_iadd(&b, src_base_addr, nir_u2u64(&b, src_offset));
nir_store_var(
&b, compacted_size_var,
nir_build_load_global(
&b, 1, 64,
nir_iadd(&b, header_addr,
nir_imm_int64(&b, offsetof(struct radv_accel_struct_header, compacted_size)))),
1);
nir_store_var(
&b, instance_offset_var,
nir_build_load_global(&b, 1, 32,
nir_iadd(&b, header_addr,
nir_imm_int64(&b, offsetof(struct radv_accel_struct_header,
instance_offset)))),
1);
nir_store_var(&b, compacted_size_var,
nir_build_load_global(
&b, 1, 64,
nir_iadd_imm(&b, header_addr,
offsetof(struct radv_accel_struct_header, compacted_size))),
1);
nir_store_var(&b, instance_offset_var,
nir_build_load_global(
&b, 1, 32,
nir_iadd_imm(&b, header_addr,
offsetof(struct radv_accel_struct_header, instance_offset))),
1);
nir_store_var(&b, instance_count_var, instance_count, 1);
nir_store_var(&b, src_offset_var, src_offset, 1);
nir_store_var(&b, dst_offset_var, nir_imm_int(&b, 0), 1);
}
nir_push_else(&b, NULL); /* COPY_MODE_COPY */
{
nir_store_var(
&b, compacted_size_var,
nir_build_load_global(
&b, 1, 64,
nir_iadd(&b, src_base_addr,
nir_imm_int64(&b, offsetof(struct radv_accel_struct_header, compacted_size)))),
1);
nir_store_var(&b, compacted_size_var,
nir_build_load_global(
&b, 1, 64,
nir_iadd_imm(&b, src_base_addr,
offsetof(struct radv_accel_struct_header, compacted_size))),
1);
nir_store_var(&b, src_offset_var, nir_imm_int(&b, 0), 1);
nir_store_var(&b, dst_offset_var, nir_imm_int(&b, 0), 1);
@ -1538,12 +1499,10 @@ build_copy_shader(struct radv_device *dev)
nir_pop_if(&b, NULL);
nir_ssa_def *instance_bound =
nir_imul(&b, nir_imm_int(&b, sizeof(struct radv_bvh_instance_node)),
nir_load_var(&b, instance_count_var));
nir_imul_imm(&b, nir_load_var(&b, instance_count_var), sizeof(struct radv_bvh_instance_node));
nir_ssa_def *compacted_size = nir_build_load_global(
&b, 1, 32,
nir_iadd(&b, src_base_addr,
nir_imm_int64(&b, offsetof(struct radv_accel_struct_header, compacted_size))));
nir_iadd_imm(&b, src_base_addr, offsetof(struct radv_accel_struct_header, compacted_size)));
nir_push_loop(&b);
{
@ -1562,23 +1521,18 @@ build_copy_shader(struct radv_device *dev)
nir_ssa_def *in_instance_bound =
nir_iand(&b, nir_uge(&b, offset, nir_load_var(&b, instance_offset_var)),
nir_ult(&b, instance_offset, instance_bound));
nir_ssa_def *instance_start =
nir_ieq(&b,
nir_iand(&b, instance_offset,
nir_imm_int(&b, sizeof(struct radv_bvh_instance_node) - 1)),
nir_imm_int(&b, 0));
nir_ssa_def *instance_start = nir_ieq_imm(
&b, nir_iand_imm(&b, instance_offset, sizeof(struct radv_bvh_instance_node) - 1), 0);
nir_push_if(&b, nir_iand(&b, in_instance_bound, instance_start));
{
nir_ssa_def *instance_id = nir_ushr(&b, instance_offset, nir_imm_int(&b, 7));
nir_ssa_def *instance_id = nir_ushr_imm(&b, instance_offset, 7);
nir_push_if(&b, nir_ieq(&b, mode, nir_imm_int(&b, COPY_MODE_SERIALIZE)));
nir_push_if(&b, nir_ieq_imm(&b, mode, COPY_MODE_SERIALIZE));
{
nir_ssa_def *instance_addr =
nir_imul(&b, instance_id, nir_imm_int(&b, sizeof(uint64_t)));
instance_addr =
nir_iadd(&b, instance_addr,
nir_imm_int(&b, sizeof(struct radv_accel_struct_serialization_header)));
nir_ssa_def *instance_addr = nir_imul_imm(&b, instance_id, sizeof(uint64_t));
instance_addr = nir_iadd_imm(&b, instance_addr,
sizeof(struct radv_accel_struct_serialization_header));
instance_addr = nir_iadd(&b, dst_base_addr, nir_u2u64(&b, instance_addr));
nir_build_store_global(&b, nir_channels(&b, value, 3), instance_addr,
@ -1586,11 +1540,9 @@ build_copy_shader(struct radv_device *dev)
}
nir_push_else(&b, NULL);
{
nir_ssa_def *instance_addr =
nir_imul(&b, instance_id, nir_imm_int(&b, sizeof(uint64_t)));
instance_addr =
nir_iadd(&b, instance_addr,
nir_imm_int(&b, sizeof(struct radv_accel_struct_serialization_header)));
nir_ssa_def *instance_addr = nir_imul_imm(&b, instance_id, sizeof(uint64_t));
instance_addr = nir_iadd_imm(&b, instance_addr,
sizeof(struct radv_accel_struct_serialization_header));
instance_addr = nir_iadd(&b, src_base_addr, nir_u2u64(&b, instance_addr));
nir_ssa_def *instance_value =

6
src/amd/vulkan/radv_meta.c
View file

@ -690,8 +690,8 @@ radv_meta_gen_rect_vertices_comp2(nir_builder *vs_b, nir_ssa_def *comp2)
/* so channel 0 is vertex_id != 2 ? -1.0 : 1.0
channel 1 is vertex id != 1 ? -1.0 : 1.0 */
nir_ssa_def *c0cmp = nir_ine(vs_b, vertex_id, nir_imm_int(vs_b, 2));
nir_ssa_def *c1cmp = nir_ine(vs_b, vertex_id, nir_imm_int(vs_b, 1));
nir_ssa_def *c0cmp = nir_ine_imm(vs_b, vertex_id, 2);
nir_ssa_def *c1cmp = nir_ine_imm(vs_b, vertex_id, 1);
nir_ssa_def *comp[4];
comp[0] = nir_bcsel(vs_b, c0cmp, nir_imm_float(vs_b, -1.0), nir_imm_float(vs_b, 1.0));
@ -846,6 +846,6 @@ radv_break_on_count(nir_builder *b, nir_variable *var, nir_ssa_def *count)
nir_jump(b, nir_jump_break);
nir_pop_if(b, NULL);
counter = nir_iadd(b, counter, nir_imm_int(b, 1));
counter = nir_iadd_imm(b, counter, 1);
nir_store_var(b, var, counter, 0x1);
}

4
src/amd/vulkan/radv_meta_blit.c
View file

@ -64,8 +64,8 @@ build_nir_vertex_shader(struct radv_device *dev)
/* so channel 0 is vertex_id != 2 ? src_x : src_x + w
channel 1 is vertex id != 1 ? src_y : src_y + w */
nir_ssa_def *c0cmp = nir_ine(&b, vertex_id, nir_imm_int(&b, 2));
nir_ssa_def *c1cmp = nir_ine(&b, vertex_id, nir_imm_int(&b, 1));
nir_ssa_def *c0cmp = nir_ine_imm(&b, vertex_id, 2);
nir_ssa_def *c1cmp = nir_ine_imm(&b, vertex_id, 1);
nir_ssa_def *comp[4];
comp[0] = nir_bcsel(&b, c0cmp, nir_channel(&b, src_box, 0), nir_channel(&b, src_box, 2));

4
src/amd/vulkan/radv_meta_blit2d.c
View file

@ -411,8 +411,8 @@ build_nir_vertex_shader(struct radv_device *device)
/* so channel 0 is vertex_id != 2 ? src_x : src_x + w
channel 1 is vertex id != 1 ? src_y : src_y + w */
nir_ssa_def *c0cmp = nir_ine(&b, vertex_id, nir_imm_int(&b, 2));
nir_ssa_def *c1cmp = nir_ine(&b, vertex_id, nir_imm_int(&b, 1));
nir_ssa_def *c0cmp = nir_ine_imm(&b, vertex_id, 2);
nir_ssa_def *c1cmp = nir_ine_imm(&b, vertex_id, 1);
nir_ssa_def *comp[2];
comp[0] = nir_bcsel(&b, c0cmp, nir_channel(&b, src_box, 0), nir_channel(&b, src_box, 2));

4
src/amd/vulkan/radv_meta_buffer.c
View file

@ -12,7 +12,7 @@ build_buffer_fill_shader(struct radv_device *dev)
nir_ssa_def *global_id = get_global_ids(&b, 1);
nir_ssa_def *offset = nir_imul(&b, global_id, nir_imm_int(&b, 16));
nir_ssa_def *offset = nir_imul_imm(&b, global_id, 16);
offset = nir_channel(&b, offset, 0);
nir_ssa_def *dst_buf = radv_meta_load_descriptor(&b, 0, 0);
@ -34,7 +34,7 @@ build_buffer_copy_shader(struct radv_device *dev)
nir_ssa_def *global_id = get_global_ids(&b, 1);
nir_ssa_def *offset = nir_imul(&b, global_id, nir_imm_int(&b, 16));
nir_ssa_def *offset = nir_imul_imm(&b, global_id, 16);
offset = nir_channel(&b, offset, 0);
nir_ssa_def *dst_buf = radv_meta_load_descriptor(&b, 0, 0);

19
src/amd/vulkan/radv_meta_bufimage.c
View file

@ -423,9 +423,8 @@ build_nir_btoi_r32g32b32_compute_shader(struct radv_device *dev)
nir_ssa_def *img_coord = nir_iadd(&b, global_id, offset);
nir_ssa_def *global_pos =
nir_iadd(&b, nir_imul(&b, nir_channel(&b, img_coord, 1), pitch),
nir_imul(&b, nir_channel(&b, img_coord, 0), nir_imm_int(&b, 3)));
nir_ssa_def *global_pos = nir_iadd(&b, nir_imul(&b, nir_channel(&b, img_coord, 1), pitch),
nir_imul_imm(&b, nir_channel(&b, img_coord, 0), 3));
nir_ssa_def *input_img_deref = &nir_build_deref_var(&b, input_img)->dest.ssa;
@ -447,7 +446,7 @@ build_nir_btoi_r32g32b32_compute_shader(struct radv_device *dev)
nir_ssa_def *outval = &tex->dest.ssa;
for (int chan = 0; chan < 3; chan++) {
nir_ssa_def *local_pos = nir_iadd(&b, global_pos, nir_imm_int(&b, chan));
nir_ssa_def *local_pos = nir_iadd_imm(&b, global_pos, chan);
nir_ssa_def *coord = nir_vec4(&b, local_pos, local_pos, local_pos, local_pos);
@ -774,15 +773,15 @@ build_nir_itoi_r32g32b32_compute_shader(struct radv_device *dev)
nir_ssa_def *src_global_pos =
nir_iadd(&b, nir_imul(&b, nir_channel(&b, src_img_coord, 1), src_stride),
nir_imul(&b, nir_channel(&b, src_img_coord, 0), nir_imm_int(&b, 3)));
nir_imul_imm(&b, nir_channel(&b, src_img_coord, 0), 3));
nir_ssa_def *dst_global_pos =
nir_iadd(&b, nir_imul(&b, nir_channel(&b, dst_img_coord, 1), dst_stride),
nir_imul(&b, nir_channel(&b, dst_img_coord, 0), nir_imm_int(&b, 3)));
nir_imul_imm(&b, nir_channel(&b, dst_img_coord, 0), 3));
for (int chan = 0; chan < 3; chan++) {
/* src */
nir_ssa_def *src_local_pos = nir_iadd(&b, src_global_pos, nir_imm_int(&b, chan));
nir_ssa_def *src_local_pos = nir_iadd_imm(&b, src_global_pos, chan);
nir_ssa_def *input_img_deref = &nir_build_deref_var(&b, input_img)->dest.ssa;
nir_tex_instr *tex = nir_tex_instr_create(b.shader, 3);
@ -803,7 +802,7 @@ build_nir_itoi_r32g32b32_compute_shader(struct radv_device *dev)
nir_ssa_def *outval = &tex->dest.ssa;
/* dst */
nir_ssa_def *dst_local_pos = nir_iadd(&b, dst_global_pos, nir_imm_int(&b, chan));
nir_ssa_def *dst_local_pos = nir_iadd_imm(&b, dst_global_pos, chan);
nir_ssa_def *dst_coord =
nir_vec4(&b, dst_local_pos, dst_local_pos, dst_local_pos, dst_local_pos);
@ -1081,10 +1080,10 @@ build_nir_cleari_r32g32b32_compute_shader(struct radv_device *dev)
nir_ssa_def *global_y = nir_channel(&b, global_id, 1);
nir_ssa_def *global_pos =
nir_iadd(&b, nir_imul(&b, global_y, stride), nir_imul(&b, global_x, nir_imm_int(&b, 3)));
nir_iadd(&b, nir_imul(&b, global_y, stride), nir_imul_imm(&b, global_x, 3));
for (unsigned chan = 0; chan < 3; chan++) {
nir_ssa_def *local_pos = nir_iadd(&b, global_pos, nir_imm_int(&b, chan));
nir_ssa_def *local_pos = nir_iadd_imm(&b, global_pos, chan);
nir_ssa_def *coord = nir_vec4(&b, local_pos, local_pos, local_pos, local_pos);

2
src/amd/vulkan/radv_meta_clear.c
View file

@ -912,7 +912,7 @@ build_clear_htile_mask_shader(struct radv_device *dev)
nir_ssa_def *global_id = get_global_ids(&b, 1);
nir_ssa_def *offset = nir_imul(&b, global_id, nir_imm_int(&b, 16));
nir_ssa_def *offset = nir_imul_imm(&b, global_id, 16);
offset = nir_channel(&b, offset, 0);
nir_ssa_def *buf = radv_meta_load_descriptor(&b, 0, 0);

13
src/amd/vulkan/radv_meta_copy_vrs_htile.c
View file

@ -52,7 +52,7 @@ build_copy_vrs_htile_shader(struct radv_device *device, struct radeon_surf *surf
nir_ssa_def *global_id = get_global_ids(&b, 2);
/* Multiply the coordinates by the HTILE block size. */
nir_ssa_def *coord = nir_imul(&b, global_id, nir_imm_ivec2(&b, 8, 8));
nir_ssa_def *coord = nir_imul_imm(&b, global_id, 8);
/* Load constants. */
nir_ssa_def *constants = nir_load_push_constant(&b, 3, 32, nir_imm_int(&b, 0), .range = 12);
@ -99,15 +99,14 @@ build_copy_vrs_htile_shader(struct radv_device *device, struct radeon_surf *surf
* VRS rate X = min(value >> 2, 1)
* VRS rate Y = min(value & 3, 1)
*/
nir_ssa_def *x_rate = nir_ushr(&b, nir_channel(&b, &tex->dest.ssa, 0), nir_imm_int(&b, 2));
nir_ssa_def *x_rate = nir_ushr_imm(&b, nir_channel(&b, &tex->dest.ssa, 0), 2);
x_rate = nir_umin(&b, x_rate, nir_imm_int(&b, 1));
nir_ssa_def *y_rate = nir_iand(&b, nir_channel(&b, &tex->dest.ssa, 0), nir_imm_int(&b, 3));
nir_ssa_def *y_rate = nir_iand_imm(&b, nir_channel(&b, &tex->dest.ssa, 0), 3);
y_rate = nir_umin(&b, y_rate, nir_imm_int(&b, 1));
/* Compute the final VRS rate. */
nir_ssa_def *vrs_rates = nir_ior(&b, nir_ishl(&b, y_rate, nir_imm_int(&b, 10)),
nir_ishl(&b, x_rate, nir_imm_int(&b, 6)));
nir_ssa_def *vrs_rates = nir_ior(&b, nir_ishl_imm(&b, y_rate, 10), nir_ishl_imm(&b, x_rate, 6));
/* Load the HTILE buffer descriptor. */
nir_ssa_def *htile_buf = radv_meta_load_descriptor(&b, 0, 1);
@ -115,13 +114,13 @@ build_copy_vrs_htile_shader(struct radv_device *device, struct radeon_surf *surf
/* Load the HTILE value if requested, otherwise use the default value. */
nir_variable *htile_value = nir_local_variable_create(b.impl, glsl_int_type(), "htile_value");
nir_push_if(&b, nir_ieq(&b, read_htile_value, nir_imm_int(&b, 1)));
nir_push_if(&b, nir_ieq_imm(&b, read_htile_value, 1));
{
/* Load the existing HTILE 32-bit value for this 8x8 pixels area. */
nir_ssa_def *input_value = nir_load_ssbo(&b, 1, 32, htile_buf, htile_addr);
/* Clear the 4-bit VRS rates. */
nir_store_var(&b, htile_value, nir_iand(&b, input_value, nir_imm_int(&b, 0xfffff33f)), 0x1);
nir_store_var(&b, htile_value, nir_iand_imm(&b, input_value, 0xfffff33f), 0x1);
}
nir_push_else(&b, NULL);
{

240
src/amd/vulkan/radv_meta_etc_decode.c
View file

@ -49,12 +49,10 @@ flip_endian(nir_builder *b, nir_ssa_def *src, unsigned cnt)
nir_ssa_def *intermediate[4];
nir_ssa_def *chan = cnt == 1 ? src : nir_channel(b, src, i);
for (unsigned j = 0; j < 4; ++j)
intermediate[j] = nir_ubfe(b, chan, nir_imm_int(b, 8 * j), nir_imm_int(b, 8));
v[i] = nir_ior(b,
nir_ior(b, nir_ishl(b, intermediate[0], nir_imm_int(b, 24)),
nir_ishl(b, intermediate[1], nir_imm_int(b, 16))),
nir_ior(b, nir_ishl(b, intermediate[2], nir_imm_int(b, 8)),
nir_ishl(b, intermediate[3], nir_imm_int(b, 0))));
intermediate[j] = nir_ubfe_imm(b, chan, 8 * j, 8);
v[i] = nir_ior(
b, nir_ior(b, nir_ishl_imm(b, intermediate[0], 24), nir_ishl_imm(b, intermediate[1], 16)),
nir_ior(b, nir_ishl_imm(b, intermediate[2], 8), nir_ishl_imm(b, intermediate[3], 0)));
}
return cnt == 1 ? v[0] : nir_vec(b, v, cnt);
}
@ -64,13 +62,13 @@ etc1_color_modifier_lookup(nir_builder *b, nir_ssa_def *x, nir_ssa_def *y)
{
const unsigned table[8][2] = {{2, 8}, {5, 17}, {9, 29}, {13, 42},
{18, 60}, {24, 80}, {33, 106}, {47, 183}};
nir_ssa_def *upper = nir_ieq(b, y, nir_imm_int(b, 1));
nir_ssa_def *upper = nir_ieq_imm(b, y, 1);
nir_ssa_def *result = NULL;
for (unsigned i = 0; i < 8; ++i) {
nir_ssa_def *tmp =
nir_bcsel(b, upper, nir_imm_int(b, table[i][1]), nir_imm_int(b, table[i][0]));
if (result)
result = nir_bcsel(b, nir_ieq(b, x, nir_imm_int(b, i)), tmp, result);
result = nir_bcsel(b, nir_ieq_imm(b, x, i), tmp, result);
else
result = tmp;
}
@ -84,7 +82,7 @@ etc2_distance_lookup(nir_builder *b, nir_ssa_def *x)
nir_ssa_def *result = NULL;
for (unsigned i = 0; i < 8; ++i) {
if (result)
result = nir_bcsel(b, nir_ieq(b, x, nir_imm_int(b, i)), nir_imm_int(b, table[i]), result);
result = nir_bcsel(b, nir_ieq_imm(b, x, i), nir_imm_int(b, table[i]), result);
else
result = nir_imm_int(b, table[i]);
}
@ -100,20 +98,19 @@ etc1_alpha_modifier_lookup(nir_builder *b, nir_ssa_def *x, nir_ssa_def *y)
for (unsigned i = 0; i < 16; ++i) {
nir_ssa_def *tmp = nir_imm_int(b, table[i]);
if (result)
result = nir_bcsel(b, nir_ieq(b, x, nir_imm_int(b, i)), tmp, result);
result = nir_bcsel(b, nir_ieq_imm(b, x, i), tmp, result);
else
result = tmp;
}
return nir_ubfe(b, result, nir_imul(b, y, nir_imm_int(b, 4)), nir_imm_int(b, 4));
return nir_ubfe(b, result, nir_imul_imm(b, y, 4), nir_imm_int(b, 4));
}
static nir_ssa_def *
etc_extend(nir_builder *b, nir_ssa_def *v, int bits)
{
if (bits == 4)
return nir_imul(b, v, nir_imm_int(b, 0x11));
return nir_ior(b, nir_ishl(b, v, nir_imm_int(b, 8 - bits)),
nir_ushr(b, v, nir_imm_int(b, bits - (8 - bits))));
return nir_imul_imm(b, v, 0x11);
return nir_ior(b, nir_ishl_imm(b, v, 8 - bits), nir_ushr_imm(b, v, bits - (8 - bits)));
}
static nir_ssa_def *
@ -123,29 +120,28 @@ decode_etc2_alpha(struct nir_builder *b, nir_ssa_def *alpha_payload, nir_ssa_def
alpha_payload = flip_endian(b, alpha_payload, 2);
nir_ssa_def *alpha_x = nir_channel(b, alpha_payload, 1);
nir_ssa_def *alpha_y = nir_channel(b, alpha_payload, 0);
nir_ssa_def *bit_offset =
nir_isub(b, nir_imm_int(b, 45), nir_imul(b, nir_imm_int(b, 3), linear_pixel));
nir_ssa_def *base = nir_ubfe(b, alpha_y, nir_imm_int(b, 24), nir_imm_int(b, 8));
nir_ssa_def *multiplier = nir_ubfe(b, alpha_y, nir_imm_int(b, 20), nir_imm_int(b, 4));
nir_ssa_def *table = nir_ubfe(b, alpha_y, nir_imm_int(b, 16), nir_imm_int(b, 4));
nir_ssa_def *bit_offset = nir_isub_imm(b, 45, nir_imul_imm(b, linear_pixel, 3));
nir_ssa_def *base = nir_ubfe_imm(b, alpha_y, 24, 8);
nir_ssa_def *multiplier = nir_ubfe_imm(b, alpha_y, 20, 4);
nir_ssa_def *table = nir_ubfe_imm(b, alpha_y, 16, 4);
if (eac) {
nir_ssa_def *signed_base = nir_ibfe(b, alpha_y, nir_imm_int(b, 24), nir_imm_int(b, 8));
signed_base = nir_imul(b, signed_base, nir_imm_int(b, 8));
base = nir_iadd(b, nir_imul(b, base, nir_imm_int(b, 8)), nir_imm_int(b, 4));
nir_ssa_def *signed_base = nir_ibfe_imm(b, alpha_y, 24, 8);
signed_base = nir_imul_imm(b, signed_base, 8);
base = nir_iadd_imm(b, nir_imul_imm(b, base, 8), 4);
base = nir_bcsel(b, is_signed, signed_base, base);
multiplier = nir_imax(b, nir_imul(b, multiplier, nir_imm_int(b, 8)), nir_imm_int(b, 1));
multiplier = nir_imax(b, nir_imul_imm(b, multiplier, 8), nir_imm_int(b, 1));
}
nir_ssa_def *lsb_index =
nir_ubfe(b, nir_bcsel(b, nir_uge(b, bit_offset, nir_imm_int(b, 32)), alpha_y, alpha_x),
nir_iand(b, bit_offset, nir_imm_int(b, 31)), nir_imm_int(b, 2));
bit_offset = nir_iadd(b, bit_offset, nir_imm_int(b, 2));
nir_iand_imm(b, bit_offset, 31), nir_imm_int(b, 2));
bit_offset = nir_iadd_imm(b, bit_offset, 2);
nir_ssa_def *msb =
nir_ubfe(b, nir_bcsel(b, nir_uge(b, bit_offset, nir_imm_int(b, 32)), alpha_y, alpha_x),
nir_iand(b, bit_offset, nir_imm_int(b, 31)), nir_imm_int(b, 1));
nir_ssa_def *mod = nir_ixor(b, etc1_alpha_modifier_lookup(b, table, lsb_index),
nir_isub(b, msb, nir_imm_int(b, 1)));
nir_iand_imm(b, bit_offset, 31), nir_imm_int(b, 1));
nir_ssa_def *mod =
nir_ixor(b, etc1_alpha_modifier_lookup(b, table, lsb_index), nir_iadd_imm(b, msb, -1));
nir_ssa_def *a = nir_iadd(b, base, nir_imul(b, mod, multiplier));
nir_ssa_def *low_bound = nir_imm_int(b, 0);
@ -204,7 +200,7 @@ build_shader(struct radv_device *dev)
nir_ssa_def *offset = nir_channels(&b, consts, 7);
nir_ssa_def *format = nir_channel(&b, consts, 3);
nir_ssa_def *image_type = nir_channel(&b, consts2, 0);
nir_ssa_def *is_3d = nir_ieq(&b, image_type, nir_imm_int(&b, VK_IMAGE_TYPE_3D));
nir_ssa_def *is_3d = nir_ieq_imm(&b, image_type, VK_IMAGE_TYPE_3D);
nir_ssa_def *coord = nir_iadd(&b, global_id, offset);
nir_ssa_def *src_coord =
nir_vec3(&b, nir_ushr_imm(&b, nir_channel(&b, coord, 0), 2),
@ -256,10 +252,9 @@ build_shader(struct radv_device *dev)
}
nir_pop_if(&b, NULL);
nir_ssa_def *pixel_coord = nir_iand(&b, nir_channels(&b, coord, 3), nir_imm_ivec2(&b, 3, 3));
nir_ssa_def *linear_pixel =
nir_iadd(&b, nir_imul(&b, nir_channel(&b, pixel_coord, 0), nir_imm_int(&b, 4)),
nir_channel(&b, pixel_coord, 1));
nir_ssa_def *pixel_coord = nir_iand_imm(&b, nir_channels(&b, coord, 3), 3);
nir_ssa_def *linear_pixel = nir_iadd(&b, nir_imul_imm(&b, nir_channel(&b, pixel_coord, 0), 4),
nir_channel(&b, pixel_coord, 1));
nir_ssa_def *payload = nir_load_var(&b, payload_var);
nir_variable *color =
@ -278,8 +273,7 @@ build_shader(struct radv_device *dev)
color_payload = flip_endian(&b, color_payload, 2);
nir_ssa_def *color_y = nir_channel(&b, color_payload, 0);
nir_ssa_def *color_x = nir_channel(&b, color_payload, 1);
nir_ssa_def *flip =
nir_ine(&b, nir_iand(&b, color_y, nir_imm_int(&b, 1)), nir_imm_int(&b, 0));
nir_ssa_def *flip = nir_ine_imm(&b, nir_iand_imm(&b, color_y, 1), 0);
nir_ssa_def *subblock = nir_ushr_imm(
&b, nir_bcsel(&b, flip, nir_channel(&b, pixel_coord, 1), nir_channel(&b, pixel_coord, 0)),
1);
@ -287,8 +281,7 @@ build_shader(struct radv_device *dev)
nir_variable *punchthrough =
nir_variable_create(b.shader, nir_var_shader_temp, glsl_bool_type(), "punchthrough");
nir_ssa_def *punchthrough_init =
nir_iand(&b, alpha_bits_1,
nir_ieq(&b, nir_iand(&b, color_y, nir_imm_int(&b, 2)), nir_imm_int(&b, 0)));
nir_iand(&b, alpha_bits_1, nir_ieq_imm(&b, nir_iand_imm(&b, color_y, 2), 0));
nir_store_var(&b, punchthrough, punchthrough_init, 0x1);
nir_variable *etc1_compat =
@ -317,72 +310,63 @@ build_shader(struct radv_device *dev)
nir_store_var(&b, rgb_result, nir_imm_ivec3(&b, 255, 0, 0), 0x7);
nir_ssa_def *msb =
nir_iand(&b, nir_ushr(&b, color_x, nir_iadd(&b, nir_imm_int(&b, 15), linear_pixel)),
nir_imm_int(&b, 2));
nir_ssa_def *lsb = nir_iand(&b, nir_ushr(&b, color_x, linear_pixel), nir_imm_int(&b, 1));
nir_iand_imm(&b, nir_ushr(&b, color_x, nir_iadd_imm(&b, linear_pixel, 15)), 2);
nir_ssa_def *lsb = nir_iand_imm(&b, nir_ushr(&b, color_x, linear_pixel), 1);
nir_push_if(
&b, nir_iand(&b, nir_inot(&b, alpha_bits_1),
nir_ieq(&b, nir_iand(&b, color_y, nir_imm_int(&b, 2)), nir_imm_int(&b, 0))));
nir_push_if(&b, nir_iand(&b, nir_inot(&b, alpha_bits_1),
nir_ieq_imm(&b, nir_iand_imm(&b, color_y, 2), 0)));
{
nir_store_var(&b, etc1_compat, nir_imm_bool(&b, true), 1);
nir_ssa_def *tmp[3];
for (unsigned i = 0; i < 3; ++i)
tmp[i] =
etc_extend(&b,
nir_iand(&b,
nir_ushr(&b, color_y,
nir_isub(&b, nir_imm_int(&b, 28 - 8 * i),
nir_imul(&b, subblock, nir_imm_int(&b, 4)))),
nir_imm_int(&b, 0xf)),
4);
tmp[i] = etc_extend(
&b,
nir_iand_imm(&b,
nir_ushr(&b, color_y,
nir_isub_imm(&b, 28 - 8 * i, nir_imul_imm(&b, subblock, 4))),
0xf),
4);
nir_store_var(&b, base_rgb, nir_vec(&b, tmp, 3), 0x7);
}
nir_push_else(&b, NULL);
{
nir_ssa_def *rb = nir_ubfe(&b, color_y, nir_imm_int(&b, 27), nir_imm_int(&b, 5));
nir_ssa_def *rd = nir_ibfe(&b, color_y, nir_imm_int(&b, 24), nir_imm_int(&b, 3));
nir_ssa_def *gb = nir_ubfe(&b, color_y, nir_imm_int(&b, 19), nir_imm_int(&b, 5));
nir_ssa_def *gd = nir_ibfe(&b, color_y, nir_imm_int(&b, 16), nir_imm_int(&b, 3));
nir_ssa_def *bb = nir_ubfe(&b, color_y, nir_imm_int(&b, 11), nir_imm_int(&b, 5));
nir_ssa_def *bd = nir_ibfe(&b, color_y, nir_imm_int(&b, 8), nir_imm_int(&b, 3));
nir_ssa_def *rb = nir_ubfe_imm(&b, color_y, 27, 5);
nir_ssa_def *rd = nir_ibfe_imm(&b, color_y, 24, 3);
nir_ssa_def *gb = nir_ubfe_imm(&b, color_y, 19, 5);
nir_ssa_def *gd = nir_ibfe_imm(&b, color_y, 16, 3);
nir_ssa_def *bb = nir_ubfe_imm(&b, color_y, 11, 5);
nir_ssa_def *bd = nir_ibfe_imm(&b, color_y, 8, 3);
nir_ssa_def *r1 = nir_iadd(&b, rb, rd);
nir_ssa_def *g1 = nir_iadd(&b, gb, gd);
nir_ssa_def *b1 = nir_iadd(&b, bb, bd);
nir_push_if(&b, nir_ult(&b, nir_imm_int(&b, 31), r1));
{
nir_ssa_def *r0 =
nir_ior(&b, nir_ubfe(&b, color_y, nir_imm_int(&b, 24), nir_imm_int(&b, 2)),
nir_ishl(&b, nir_ubfe(&b, color_y, nir_imm_int(&b, 27), nir_imm_int(&b, 2)),
nir_imm_int(&b, 2)));
nir_ssa_def *g0 = nir_ubfe(&b, color_y, nir_imm_int(&b, 20), nir_imm_int(&b, 4));
nir_ssa_def *b0 = nir_ubfe(&b, color_y, nir_imm_int(&b, 16), nir_imm_int(&b, 4));
nir_ssa_def *r2 = nir_ubfe(&b, color_y, nir_imm_int(&b, 12), nir_imm_int(&b, 4));
nir_ssa_def *g2 = nir_ubfe(&b, color_y, nir_imm_int(&b, 8), nir_imm_int(&b, 4));
nir_ssa_def *b2 = nir_ubfe(&b, color_y, nir_imm_int(&b, 4), nir_imm_int(&b, 4));
nir_ssa_def *da =
nir_ior(&b,
nir_ishl(&b, nir_ubfe(&b, color_y, nir_imm_int(&b, 2), nir_imm_int(&b, 2)),
nir_imm_int(&b, 1)),
nir_iand(&b, color_y, nir_imm_int(&b, 1)));
nir_ssa_def *r0 = nir_ior(&b, nir_ubfe_imm(&b, color_y, 24, 2),
nir_ishl_imm(&b, nir_ubfe_imm(&b, color_y, 27, 2), 2));
nir_ssa_def *g0 = nir_ubfe_imm(&b, color_y, 20, 4);
nir_ssa_def *b0 = nir_ubfe_imm(&b, color_y, 16, 4);
nir_ssa_def *r2 = nir_ubfe_imm(&b, color_y, 12, 4);
nir_ssa_def *g2 = nir_ubfe_imm(&b, color_y, 8, 4);
nir_ssa_def *b2 = nir_ubfe_imm(&b, color_y, 4, 4);
nir_ssa_def *da = nir_ior(&b, nir_ishl_imm(&b, nir_ubfe_imm(&b, color_y, 2, 2), 1),
nir_iand_imm(&b, color_y, 1));
nir_ssa_def *dist = etc2_distance_lookup(&b, da);
nir_ssa_def *index = nir_ior(&b, lsb, msb);
nir_store_var(&b, punchthrough,
nir_iand(&b, nir_load_var(&b, punchthrough),
nir_ieq(&b, nir_iadd(&b, lsb, msb), nir_imm_int(&b, 2))),
nir_ieq_imm(&b, nir_iadd(&b, lsb, msb), 2)),
0x1);
nir_push_if(&b, nir_ieq(&b, index, nir_imm_int(&b, 0)));
nir_push_if(&b, nir_ieq_imm(&b, index, 0));
{
nir_store_var(&b, rgb_result, etc_extend(&b, nir_vec3(&b, r0, g0, b0), 4), 0x7);
}
nir_push_else(&b, NULL);
{
nir_ssa_def *tmp =
nir_iadd(&b, etc_extend(&b, nir_vec3(&b, r2, g2, b2), 4),
nir_imul(&b, dist, nir_isub(&b, nir_imm_int(&b, 2), index)));
nir_ssa_def *tmp = nir_iadd(&b, etc_extend(&b, nir_vec3(&b, r2, g2, b2), 4),
nir_imul(&b, dist, nir_isub_imm(&b, 2, index)));
nir_store_var(&b, rgb_result, tmp, 0x7);
}
nir_pop_if(&b, NULL);
@ -390,64 +374,51 @@ build_shader(struct radv_device *dev)
nir_push_else(&b, NULL);
nir_push_if(&b, nir_ult(&b, nir_imm_int(&b, 31), g1));
{
nir_ssa_def *r0 = nir_ubfe(&b, color_y, nir_imm_int(&b, 27), nir_imm_int(&b, 4));
nir_ssa_def *g0 = nir_ior(
&b,
nir_ishl(&b, nir_ubfe(&b, color_y, nir_imm_int(&b, 24), nir_imm_int(&b, 3)),
nir_imm_int(&b, 1)),
nir_iand(&b, nir_ushr(&b, color_y, nir_imm_int(&b, 20)), nir_imm_int(&b, 1)));
nir_ssa_def *b0 = nir_ior(
&b, nir_ubfe(&b, color_y, nir_imm_int(&b, 15), nir_imm_int(&b, 3)),
nir_iand(&b, nir_ushr(&b, color_y, nir_imm_int(&b, 16)), nir_imm_int(&b, 8)));
nir_ssa_def *r2 = nir_ubfe(&b, color_y, nir_imm_int(&b, 11), nir_imm_int(&b, 4));
nir_ssa_def *g2 = nir_ubfe(&b, color_y, nir_imm_int(&b, 7), nir_imm_int(&b, 4));
nir_ssa_def *b2 = nir_ubfe(&b, color_y, nir_imm_int(&b, 3), nir_imm_int(&b, 4));
nir_ssa_def *da = nir_iand(&b, color_y, nir_imm_int(&b, 4));
nir_ssa_def *db = nir_iand(&b, color_y, nir_imm_int(&b, 1));
nir_ssa_def *d = nir_iadd(&b, da, nir_imul(&b, db, nir_imm_int(&b, 2)));
nir_ssa_def *d0 = nir_iadd(&b, nir_ishl(&b, r0, nir_imm_int(&b, 16)),
nir_iadd(&b, nir_ishl(&b, g0, nir_imm_int(&b, 8)), b0));
nir_ssa_def *d2 = nir_iadd(&b, nir_ishl(&b, r2, nir_imm_int(&b, 16)),
nir_iadd(&b, nir_ishl(&b, g2, nir_imm_int(&b, 8)), b2));
d = nir_bcsel(&b, nir_uge(&b, d0, d2), nir_iadd(&b, d, nir_imm_int(&b, 1)), d);
nir_ssa_def *r0 = nir_ubfe_imm(&b, color_y, 27, 4);
nir_ssa_def *g0 = nir_ior(&b, nir_ishl_imm(&b, nir_ubfe_imm(&b, color_y, 24, 3), 1),
nir_iand_imm(&b, nir_ushr_imm(&b, color_y, 20), 1));
nir_ssa_def *b0 = nir_ior(&b, nir_ubfe_imm(&b, color_y, 15, 3),
nir_iand_imm(&b, nir_ushr_imm(&b, color_y, 16), 8));
nir_ssa_def *r2 = nir_ubfe_imm(&b, color_y, 11, 4);
nir_ssa_def *g2 = nir_ubfe_imm(&b, color_y, 7, 4);
nir_ssa_def *b2 = nir_ubfe_imm(&b, color_y, 3, 4);
nir_ssa_def *da = nir_iand_imm(&b, color_y, 4);
nir_ssa_def *db = nir_iand_imm(&b, color_y, 1);
nir_ssa_def *d = nir_iadd(&b, da, nir_imul_imm(&b, db, 2));
nir_ssa_def *d0 =
nir_iadd(&b, nir_ishl_imm(&b, r0, 16), nir_iadd(&b, nir_ishl_imm(&b, g0, 8), b0));
nir_ssa_def *d2 =
nir_iadd(&b, nir_ishl_imm(&b, r2, 16), nir_iadd(&b, nir_ishl_imm(&b, g2, 8), b2));
d = nir_bcsel(&b, nir_uge(&b, d0, d2), nir_iadd_imm(&b, d, 1), d);
nir_ssa_def *dist = etc2_distance_lookup(&b, d);
nir_ssa_def *base = nir_bcsel(&b, nir_ine(&b, msb, nir_imm_int(&b, 0)),
nir_vec3(&b, r2, g2, b2), nir_vec3(&b, r0, g0, b0));
nir_ssa_def *base = nir_bcsel(&b, nir_ine_imm(&b, msb, 0), nir_vec3(&b, r2, g2, b2),
nir_vec3(&b, r0, g0, b0));
base = etc_extend(&b, base, 4);
base = nir_iadd(
&b, base,
nir_imul(&b, dist,
nir_isub(&b, nir_imm_int(&b, 1), nir_imul(&b, lsb, nir_imm_int(&b, 2)))));
base = nir_iadd(&b, base,
nir_imul(&b, dist, nir_isub_imm(&b, 1, nir_imul_imm(&b, lsb, 2))));
nir_store_var(&b, rgb_result, base, 0x7);
nir_store_var(&b, punchthrough,
nir_iand(&b, nir_load_var(&b, punchthrough),
nir_ieq(&b, nir_iadd(&b, lsb, msb), nir_imm_int(&b, 2))),
nir_ieq_imm(&b, nir_iadd(&b, lsb, msb), 2)),
0x1);
}
nir_push_else(&b, NULL);
nir_push_if(&b, nir_ult(&b, nir_imm_int(&b, 31), b1));
{
nir_ssa_def *r0 = nir_ubfe(&b, color_y, nir_imm_int(&b, 25), nir_imm_int(&b, 6));
nir_ssa_def *g0 = nir_ior(
&b, nir_ubfe(&b, color_y, nir_imm_int(&b, 17), nir_imm_int(&b, 6)),
nir_iand(&b, nir_ushr(&b, color_y, nir_imm_int(&b, 18)), nir_imm_int(&b, 0x40)));
nir_ssa_def *b0 = nir_ior(
&b,
nir_ishl(&b, nir_ubfe(&b, color_y, nir_imm_int(&b, 11), nir_imm_int(&b, 2)),
nir_imm_int(&b, 3)),
nir_ior(
&b,
nir_iand(&b, nir_ushr(&b, color_y, nir_imm_int(&b, 11)), nir_imm_int(&b, 0x20)),
nir_ubfe(&b, color_y, nir_imm_int(&b, 7), nir_imm_int(&b, 3))));
nir_ssa_def *rh =
nir_ior(&b, nir_iand(&b, color_y, nir_imm_int(&b, 1)),
nir_ishl(&b, nir_ubfe(&b, color_y, nir_imm_int(&b, 2), nir_imm_int(&b, 5)),
nir_imm_int(&b, 1)));
nir_ssa_def *rv = nir_ubfe(&b, color_x, nir_imm_int(&b, 13), nir_imm_int(&b, 6));
nir_ssa_def *gh = nir_ubfe(&b, color_x, nir_imm_int(&b, 25), nir_imm_int(&b, 7));
nir_ssa_def *gv = nir_ubfe(&b, color_x, nir_imm_int(&b, 6), nir_imm_int(&b, 7));
nir_ssa_def *bh = nir_ubfe(&b, color_x, nir_imm_int(&b, 19), nir_imm_int(&b, 6));
nir_ssa_def *bv = nir_ubfe(&b, color_x, nir_imm_int(&b, 0), nir_imm_int(&b, 6));
nir_ssa_def *r0 = nir_ubfe_imm(&b, color_y, 25, 6);
nir_ssa_def *g0 = nir_ior(&b, nir_ubfe_imm(&b, color_y, 17, 6),
nir_iand_imm(&b, nir_ushr_imm(&b, color_y, 18), 0x40));
nir_ssa_def *b0 =
nir_ior(&b, nir_ishl_imm(&b, nir_ubfe_imm(&b, color_y, 11, 2), 3),
nir_ior(&b, nir_iand_imm(&b, nir_ushr_imm(&b, color_y, 11), 0x20),
nir_ubfe_imm(&b, color_y, 7, 3)));
nir_ssa_def *rh = nir_ior(&b, nir_iand_imm(&b, color_y, 1),
nir_ishl_imm(&b, nir_ubfe_imm(&b, color_y, 2, 5), 1));
nir_ssa_def *rv = nir_ubfe_imm(&b, color_x, 13, 6);
nir_ssa_def *gh = nir_ubfe_imm(&b, color_x, 25, 7);
nir_ssa_def *gv = nir_ubfe_imm(&b, color_x, 6, 7);
nir_ssa_def *bh = nir_ubfe_imm(&b, color_x, 19, 6);
nir_ssa_def *bv = nir_ubfe_imm(&b, color_x, 0, 6);
r0 = etc_extend(&b, r0, 6);
g0 = etc_extend(&b, g0, 7);
@ -464,16 +435,14 @@ build_shader(struct radv_device *dev)
nir_channel(&b, pixel_coord, 0));
nir_ssa_def *dy = nir_imul(&b, nir_isub(&b, nir_vec3(&b, rv, gv, bv), rgb),
nir_channel(&b, pixel_coord, 1));
rgb = nir_iadd(&b, rgb,
nir_ishr(&b, nir_iadd(&b, nir_iadd(&b, dx, dy), nir_imm_int(&b, 2)),
nir_imm_int(&b, 2)));
rgb = nir_iadd(&b, rgb, nir_ishr_imm(&b, nir_iadd_imm(&b, nir_iadd(&b, dx, dy), 2), 2));
nir_store_var(&b, rgb_result, rgb, 0x7);
nir_store_var(&b, punchthrough, nir_imm_bool(&b, false), 0x1);
}
nir_push_else(&b, NULL);
{
nir_store_var(&b, etc1_compat, nir_imm_bool(&b, true), 1);
nir_ssa_def *subblock_b = nir_ine(&b, subblock, nir_imm_int(&b, 0));
nir_ssa_def *subblock_b = nir_ine_imm(&b, subblock, 0);
nir_ssa_def *tmp[] = {
nir_bcsel(&b, subblock_b, r1, rb),
nir_bcsel(&b, subblock_b, g1, gb),
@ -488,15 +457,13 @@ build_shader(struct radv_device *dev)
nir_pop_if(&b, NULL);
nir_push_if(&b, nir_load_var(&b, etc1_compat));
{
nir_ssa_def *etc1_table_index =
nir_ubfe(&b, color_y,
nir_isub(&b, nir_imm_int(&b, 5), nir_imul(&b, nir_imm_int(&b, 3), subblock)),
nir_imm_int(&b, 3));
nir_ssa_def *sgn = nir_isub(&b, nir_imm_int(&b, 1), msb);
nir_ssa_def *etc1_table_index = nir_ubfe(
&b, color_y, nir_isub_imm(&b, 5, nir_imul_imm(&b, subblock, 3)), nir_imm_int(&b, 3));
nir_ssa_def *sgn = nir_isub_imm(&b, 1, msb);
sgn = nir_bcsel(&b, nir_load_var(&b, punchthrough), nir_imul(&b, sgn, lsb), sgn);
nir_store_var(&b, punchthrough,
nir_iand(&b, nir_load_var(&b, punchthrough),
nir_ieq(&b, nir_iadd(&b, lsb, msb), nir_imm_int(&b, 2))),
nir_ieq_imm(&b, nir_iadd(&b, lsb, msb), 2)),
0x1);
nir_ssa_def *off =
nir_imul(&b, etc1_color_modifier_lookup(&b, etc1_table_index, lsb), sgn);
@ -519,9 +486,8 @@ build_shader(struct radv_device *dev)
}
nir_push_else(&b, NULL);
{ /* EAC */
nir_ssa_def *is_signed =
nir_ior(&b, nir_ieq(&b, format, nir_imm_int(&b, VK_FORMAT_EAC_R11_SNORM_BLOCK)),
nir_ieq(&b, format, nir_imm_int(&b, VK_FORMAT_EAC_R11G11_SNORM_BLOCK)));
nir_ssa_def *is_signed = nir_ior(&b, nir_ieq_imm(&b, format, VK_FORMAT_EAC_R11_SNORM_BLOCK),
nir_ieq_imm(&b, format, VK_FORMAT_EAC_R11G11_SNORM_BLOCK));
nir_ssa_def *val[4];
for (int i = 0; i < 2; ++i) {
val[i] = decode_etc2_alpha(&b, nir_channels(&b, payload, 3 << (2 * i)), linear_pixel, true,

166
src/amd/vulkan/radv_nir_lower_ray_queries.c
View file

@ -298,11 +298,9 @@ static void
insert_terminate_on_first_hit(nir_builder *b, nir_ssa_def *index, struct ray_query_vars *vars,
bool break_on_terminate)
{
nir_ssa_def *terminate_on_first_hit =
nir_ine(b,
nir_iand(b, rq_load_var(b, index, vars->flags),
nir_imm_int(b, SpvRayFlagsTerminateOnFirstHitKHRMask)),
nir_imm_int(b, 0));
nir_ssa_def *terminate_on_first_hit = nir_ine_imm(
b, nir_iand_imm(b, rq_load_var(b, index, vars->flags), SpvRayFlagsTerminateOnFirstHitKHRMask),
0);
nir_push_if(b, terminate_on_first_hit);
{
rq_store_var(b, index, vars->incomplete, nir_imm_bool(b, false), 0x1);
@ -346,8 +344,7 @@ lower_rq_initialize(nir_builder *b, nir_ssa_def *index, nir_intrinsic_instr *ins
{
rq_store_var(b, index, vars->accel_struct, instr->src[1].ssa, 0x1);
rq_store_var(b, index, vars->flags, instr->src[2].ssa, 0x1);
rq_store_var(b, index, vars->cull_mask, nir_iand(b, instr->src[3].ssa, nir_imm_int(b, 0xff)),
0x1);
rq_store_var(b, index, vars->cull_mask, nir_iand_imm(b, instr->src[3].ssa, 0xff), 0x1);
rq_store_var(b, index, vars->origin, instr->src[4].ssa, 0x7);
rq_store_var(b, index, vars->trav.origin, instr->src[4].ssa, 0x7);
@ -366,7 +363,7 @@ lower_rq_initialize(nir_builder *b, nir_ssa_def *index, nir_intrinsic_instr *ins
nir_ssa_def *accel_struct = rq_load_var(b, index, vars->accel_struct);
nir_push_if(b, nir_ine(b, accel_struct, nir_imm_int64(b, 0)));
nir_push_if(b, nir_ine_imm(b, accel_struct, 0));
{
rq_store_var(b, index, vars->trav.bvh_base, build_addr_to_node(b, accel_struct), 1);
@ -400,32 +397,32 @@ lower_rq_load(nir_builder *b, nir_ssa_def *index, struct ray_query_vars *vars,
rq_load_var(b, index, vars->candidate.barycentrics));
case nir_ray_query_value_intersection_candidate_aabb_opaque:
return nir_iand(b, rq_load_var(b, index, vars->candidate.opaque),
nir_ieq(b, rq_load_var(b, index, vars->candidate.intersection_type),
nir_imm_int(b, intersection_type_aabb)));
nir_ieq_imm(b, rq_load_var(b, index, vars->candidate.intersection_type),
intersection_type_aabb));
case nir_ray_query_value_intersection_front_face:
return nir_bcsel(b, committed, rq_load_var(b, index, vars->closest.frontface),
rq_load_var(b, index, vars->candidate.frontface));
case nir_ray_query_value_intersection_geometry_index:
return nir_iand(
return nir_iand_imm(
b,
nir_bcsel(b, committed, rq_load_var(b, index, vars->closest.geometry_id_and_flags),
rq_load_var(b, index, vars->candidate.geometry_id_and_flags)),
nir_imm_int(b, 0xFFFFFF));
0xFFFFFF);
case nir_ray_query_value_intersection_instance_custom_index:
return nir_iand(
return nir_iand_imm(
b,
nir_bcsel(b, committed, rq_load_var(b, index, vars->closest.custom_instance_and_mask),
rq_load_var(b, index, vars->candidate.custom_instance_and_mask)),
nir_imm_int(b, 0xFFFFFF));
0xFFFFFF);
case nir_ray_query_value_intersection_instance_id:
return nir_bcsel(b, committed, rq_load_var(b, index, vars->closest.instance_id),
rq_load_var(b, index, vars->candidate.instance_id));
case nir_ray_query_value_intersection_instance_sbt_index:
return nir_iand(
return nir_iand_imm(
b,
nir_bcsel(b, committed, rq_load_var(b, index, vars->closest.sbt_offset_and_flags),
rq_load_var(b, index, vars->candidate.sbt_offset_and_flags)),
nir_imm_int(b, 0xFFFFFF));
0xFFFFFF);
case nir_ray_query_value_intersection_object_ray_direction: {
nir_ssa_def *instance_node_addr =
nir_bcsel(b, committed, rq_load_var(b, index, vars->closest.instance_addr),
@ -439,12 +436,12 @@ lower_rq_load(nir_builder *b, nir_ssa_def *index, struct ray_query_vars *vars,
nir_bcsel(b, committed, rq_load_var(b, index, vars->closest.instance_addr),
rq_load_var(b, index, vars->candidate.instance_addr));
nir_ssa_def *wto_matrix[] = {
nir_build_load_global(b, 4, 32, nir_iadd(b, instance_node_addr, nir_imm_int64(b, 16)),
.align_mul = 64, .align_offset = 16),
nir_build_load_global(b, 4, 32, nir_iadd(b, instance_node_addr, nir_imm_int64(b, 32)),
.align_mul = 64, .align_offset = 32),
nir_build_load_global(b, 4, 32, nir_iadd(b, instance_node_addr, nir_imm_int64(b, 48)),
.align_mul = 64, .align_offset = 48)};
nir_build_load_global(b, 4, 32, nir_iadd_imm(b, instance_node_addr, 16), .align_mul = 64,
.align_offset = 16),
nir_build_load_global(b, 4, 32, nir_iadd_imm(b, instance_node_addr, 32), .align_mul = 64,
.align_offset = 32),
nir_build_load_global(b, 4, 32, nir_iadd_imm(b, instance_node_addr, 48), .align_mul = 64,
.align_offset = 48)};
return nir_build_vec3_mat_mult_pre(b, rq_load_var(b, index, vars->origin), wto_matrix);
}
case nir_ray_query_value_intersection_object_to_world: {
@ -463,8 +460,7 @@ lower_rq_load(nir_builder *b, nir_ssa_def *index, struct ray_query_vars *vars,
return nir_vec(b, vals, 3);
}
return nir_build_load_global(
b, 3, 32, nir_iadd(b, instance_node_addr, nir_imm_int64(b, 92 + column * 12)));
return nir_build_load_global(b, 3, 32, nir_iadd_imm(b, instance_node_addr, 92 + column * 12));
}
case nir_ray_query_value_intersection_primitive_index:
return nir_bcsel(b, committed, rq_load_var(b, index, vars->closest.primitive_id),
@ -475,7 +471,7 @@ lower_rq_load(nir_builder *b, nir_ssa_def *index, struct ray_query_vars *vars,
case nir_ray_query_value_intersection_type:
return nir_bcsel(
b, committed, rq_load_var(b, index, vars->closest.intersection_type),
nir_isub(b, rq_load_var(b, index, vars->candidate.intersection_type), nir_imm_int(b, 1)));
nir_iadd_imm(b, rq_load_var(b, index, vars->candidate.intersection_type), -1));
case nir_ray_query_value_intersection_world_to_object: {
nir_ssa_def *instance_node_addr =
nir_bcsel(b, committed, rq_load_var(b, index, vars->closest.instance_addr),
@ -515,44 +511,41 @@ insert_traversal_triangle_case(struct radv_device *device, nir_builder *b, nir_s
nir_ssa_def *div = nir_vector_extract(b, result, nir_imm_int(b, 1));
dist = nir_fdiv(b, dist, div);
nir_ssa_def *frontface = nir_flt(b, nir_imm_float(b, 0), div);
nir_ssa_def *switch_ccw = nir_ine(
b,
nir_iand(b, rq_load_var(b, index, vars->candidate.sbt_offset_and_flags),
nir_imm_int(b, VK_GEOMETRY_INSTANCE_TRIANGLE_FRONT_COUNTERCLOCKWISE_BIT_KHR << 24)),
nir_imm_int(b, 0));
nir_ssa_def *switch_ccw =
nir_ine_imm(b,
nir_iand_imm(b, rq_load_var(b, index, vars->candidate.sbt_offset_and_flags),
VK_GEOMETRY_INSTANCE_TRIANGLE_FRONT_COUNTERCLOCKWISE_BIT_KHR << 24),
0);
frontface = nir_ixor(b, frontface, switch_ccw);
rq_store_var(b, index, vars->candidate.frontface, frontface, 0x1);
nir_ssa_def *not_cull = nir_ieq(b,
nir_iand(b, rq_load_var(b, index, vars->flags),
nir_imm_int(b, SpvRayFlagsSkipTrianglesKHRMask)),
nir_imm_int(b, 0));
nir_ssa_def *not_facing_cull = nir_ieq(
nir_ssa_def *not_cull = nir_ieq_imm(
b, nir_iand_imm(b, rq_load_var(b, index, vars->flags), SpvRayFlagsSkipTrianglesKHRMask), 0);
nir_ssa_def *not_facing_cull = nir_ieq_imm(
b,
nir_iand(b, rq_load_var(b, index, vars->flags),
nir_bcsel(b, frontface, nir_imm_int(b, SpvRayFlagsCullFrontFacingTrianglesKHRMask),
nir_imm_int(b, SpvRayFlagsCullBackFacingTrianglesKHRMask))),
nir_imm_int(b, 0));
0);
not_cull = nir_iand(
b, not_cull,
nir_ior(
b, not_facing_cull,
nir_ine(b,
nir_iand(b, rq_load_var(b, index, vars->candidate.sbt_offset_and_flags),
nir_imm_int(
b, VK_GEOMETRY_INSTANCE_TRIANGLE_FACING_CULL_DISABLE_BIT_KHR << 24)),
nir_imm_int(b, 0))));
nir_ine_imm(b,
nir_iand_imm(b, rq_load_var(b, index, vars->candidate.sbt_offset_and_flags),
VK_GEOMETRY_INSTANCE_TRIANGLE_FACING_CULL_DISABLE_BIT_KHR << 24),
0)));
nir_push_if(b, nir_iand(b,
nir_iand(b, nir_flt(b, dist, rq_load_var(b, index, vars->closest.t)),
nir_fge(b, dist, rq_load_var(b, index, vars->tmin))),
not_cull));
{
nir_ssa_def *triangle_info = nir_build_load_global(
b, 2, 32,
nir_iadd(b, build_node_to_addr(device, b, bvh_node),
nir_imm_int64(b, offsetof(struct radv_bvh_triangle_node, triangle_id))));
nir_ssa_def *triangle_info =
nir_build_load_global(b, 2, 32,
nir_iadd_imm(b, build_node_to_addr(device, b, bvh_node),
offsetof(struct radv_bvh_triangle_node, triangle_id)));
nir_ssa_def *primitive_id = nir_channel(b, triangle_info, 0);
nir_ssa_def *geometry_id_and_flags = nir_channel(b, triangle_info, 1);
nir_ssa_def *is_opaque =
@ -560,11 +553,11 @@ insert_traversal_triangle_case(struct radv_device *device, nir_builder *b, nir_s
rq_load_var(b, index, vars->flags), geometry_id_and_flags);
not_cull =
nir_ieq(b,
nir_iand(b, rq_load_var(b, index, vars->flags),
nir_bcsel(b, is_opaque, nir_imm_int(b, SpvRayFlagsCullOpaqueKHRMask),
nir_imm_int(b, SpvRayFlagsCullNoOpaqueKHRMask))),
nir_imm_int(b, 0));
nir_ieq_imm(b,
nir_iand(b, rq_load_var(b, index, vars->flags),
nir_bcsel(b, is_opaque, nir_imm_int(b, SpvRayFlagsCullOpaqueKHRMask),
nir_imm_int(b, SpvRayFlagsCullNoOpaqueKHRMask))),
0);
nir_push_if(b, not_cull);
{
nir_ssa_def *divs[2] = {div, div};
@ -599,35 +592,30 @@ insert_traversal_aabb_case(struct radv_device *device, nir_builder *b, nir_ssa_d
struct ray_query_vars *vars, nir_ssa_def *bvh_node)
{
nir_ssa_def *node_addr = build_node_to_addr(device, b, bvh_node);
nir_ssa_def *triangle_info =
nir_build_load_global(b, 2, 32, nir_iadd(b, node_addr, nir_imm_int64(b, 24)));
nir_ssa_def *triangle_info = nir_build_load_global(b, 2, 32, nir_iadd_imm(b, node_addr, 24));
nir_ssa_def *primitive_id = nir_channel(b, triangle_info, 0);
nir_ssa_def *geometry_id_and_flags = nir_channel(b, triangle_info, 1);
nir_ssa_def *is_opaque =
hit_is_opaque(b, rq_load_var(b, index, vars->candidate.sbt_offset_and_flags),
rq_load_var(b, index, vars->flags), geometry_id_and_flags);
nir_ssa_def *not_skip_aabb = nir_ieq(
b,
nir_iand(b, rq_load_var(b, index, vars->flags), nir_imm_int(b, SpvRayFlagsSkipAABBsKHRMask)),
nir_imm_int(b, 0));
nir_ssa_def *not_skip_aabb = nir_ieq_imm(
b, nir_iand_imm(b, rq_load_var(b, index, vars->flags), SpvRayFlagsSkipAABBsKHRMask), 0);
nir_ssa_def *not_cull = nir_iand(
b, not_skip_aabb,
nir_ieq(b,
nir_iand(b, rq_load_var(b, index, vars->flags),
nir_bcsel(b, is_opaque, nir_imm_int(b, SpvRayFlagsCullOpaqueKHRMask),
nir_imm_int(b, SpvRayFlagsCullNoOpaqueKHRMask))),
nir_imm_int(b, 0)));
nir_ieq_imm(b,
nir_iand(b, rq_load_var(b, index, vars->flags),
nir_bcsel(b, is_opaque, nir_imm_int(b, SpvRayFlagsCullOpaqueKHRMask),
nir_imm_int(b, SpvRayFlagsCullNoOpaqueKHRMask))),
0));
nir_push_if(b, not_cull);
{
nir_ssa_def *vec3_zero = nir_channels(b, nir_imm_vec4(b, 0, 0, 0, 0), 0x7);
nir_ssa_def *vec3_inf =
nir_channels(b, nir_imm_vec4(b, INFINITY, INFINITY, INFINITY, 0), 0x7);
nir_ssa_def *bvh_lo =
nir_build_load_global(b, 3, 32, nir_iadd(b, node_addr, nir_imm_int64(b, 0)));
nir_ssa_def *bvh_hi =
nir_build_load_global(b, 3, 32, nir_iadd(b, node_addr, nir_imm_int64(b, 12)));
nir_ssa_def *bvh_lo = nir_build_load_global(b, 3, 32, nir_iadd_imm(b, node_addr, 0));
nir_ssa_def *bvh_hi = nir_build_load_global(b, 3, 32, nir_iadd_imm(b, node_addr, 12));
bvh_lo = nir_fsub(b, bvh_lo, rq_load_var(b, index, vars->trav.origin));
bvh_hi = nir_fsub(b, bvh_hi, rq_load_var(b, index, vars->trav.origin));
@ -685,7 +673,7 @@ lower_rq_proceed(nir_builder *b, nir_ssa_def *index, struct ray_query_vars *vars
nir_push_if(b, nir_uge(b, rq_load_var(b, index, vars->trav.top_stack),
rq_load_var(b, index, vars->trav.stack)));
{
nir_push_if(b, nir_ieq(b, rq_load_var(b, index, vars->trav.stack), nir_imm_int(b, 0)));
nir_push_if(b, nir_ieq_imm(b, rq_load_var(b, index, vars->trav.stack), 0));
{
rq_store_var(b, index, vars->incomplete, nir_imm_bool(b, false), 0x1);
nir_jump(b, nir_jump_break);
@ -703,7 +691,7 @@ lower_rq_proceed(nir_builder *b, nir_ssa_def *index, struct ray_query_vars *vars
nir_pop_if(b, NULL);
rq_store_var(b, index, vars->trav.stack,
nir_isub(b, rq_load_var(b, index, vars->trav.stack), nir_imm_int(b, 1)), 1);
nir_iadd_imm(b, rq_load_var(b, index, vars->trav.stack), 1), 1);
nir_ssa_def *bvh_node =
rq_load_array(b, index, vars->stack, rq_load_var(b, index, vars->trav.stack));
@ -721,16 +709,13 @@ lower_rq_proceed(nir_builder *b, nir_ssa_def *index, struct ray_query_vars *vars
}
/* if (node.type_flags & aabb) */
nir_push_if(b,
nir_ine(b, nir_iand(b, bvh_node_type, nir_imm_int(b, 4)), nir_imm_int(b, 0)));
nir_push_if(b, nir_ine_imm(b, nir_iand_imm(b, bvh_node_type, 4), 0));
{
/* if (node.type_flags & leaf) */
nir_push_if(
b, nir_ine(b, nir_iand(b, bvh_node_type, nir_imm_int(b, 2)), nir_imm_int(b, 0)));
nir_push_if(b, nir_ine_imm(b, nir_iand_imm(b, bvh_node_type, 2), 0));
{
/* custom */
nir_push_if(
b, nir_ine(b, nir_iand(b, bvh_node_type, nir_imm_int(b, 1)), nir_imm_int(b, 0)));
nir_push_if(b, nir_ine_imm(b, nir_iand_imm(b, bvh_node_type, 1), 0));
{
insert_traversal_aabb_case(device, b, index, vars, bvh_node);
}
@ -741,30 +726,26 @@ lower_rq_proceed(nir_builder *b, nir_ssa_def *index, struct ray_query_vars *vars
nir_ssa_def *instance_data = nir_build_load_global(
b, 4, 32, instance_node_addr, .align_mul = 64, .align_offset = 0);
nir_ssa_def *instance_and_mask = nir_channel(b, instance_data, 2);
nir_ssa_def *instance_mask = nir_ushr(b, instance_and_mask, nir_imm_int(b, 24));
nir_ssa_def *instance_mask = nir_ushr_imm(b, instance_and_mask, 24);
nir_push_if(
b,
nir_ieq(b, nir_iand(b, instance_mask, rq_load_var(b, index, vars->cull_mask)),
nir_imm_int(b, 0)));
nir_ieq_imm(
b, nir_iand(b, instance_mask, rq_load_var(b, index, vars->cull_mask)), 0));
{
nir_jump(b, nir_jump_continue);
}
nir_pop_if(b, NULL);
nir_ssa_def *wto_matrix[] = {
nir_build_load_global(b, 4, 32,
nir_iadd(b, instance_node_addr, nir_imm_int64(b, 16)),
nir_build_load_global(b, 4, 32, nir_iadd_imm(b, instance_node_addr, 16),
.align_mul = 64, .align_offset = 16),
nir_build_load_global(b, 4, 32,
nir_iadd(b, instance_node_addr, nir_imm_int64(b, 32)),
nir_build_load_global(b, 4, 32, nir_iadd_imm(b, instance_node_addr, 32),
.align_mul = 64, .align_offset = 32),
nir_build_load_global(b, 4, 32,
nir_iadd(b, instance_node_addr, nir_imm_int64(b, 48)),
nir_build_load_global(b, 4, 32, nir_iadd_imm(b, instance_node_addr, 48),
.align_mul = 64, .align_offset = 48)};
nir_ssa_def *instance_id = nir_build_load_global(
b, 1, 32, nir_iadd(b, instance_node_addr, nir_imm_int64(b, 88)),
.align_mul = 4, .align_offset = 0);
nir_ssa_def *instance_id =
nir_build_load_global(b, 1, 32, nir_iadd_imm(b, instance_node_addr, 88));
rq_store_var(b, index, vars->trav.top_stack,
rq_load_var(b, index, vars->trav.stack), 1);
@ -774,11 +755,9 @@ lower_rq_proceed(nir_builder *b, nir_ssa_def *index, struct ray_query_vars *vars
1);
rq_store_array(b, index, vars->stack, rq_load_var(b, index, vars->trav.stack),
nir_iand(b, nir_channel(b, instance_data, 0), nir_imm_int(b, 63)),
0x1);
rq_store_var(
b, index, vars->trav.stack,
nir_iadd(b, rq_load_var(b, index, vars->trav.stack), nir_imm_int(b, 1)), 1);
nir_iand_imm(b, nir_channel(b, instance_data, 0), 63), 0x1);
rq_store_var(b, index, vars->trav.stack,
nir_iadd_imm(b, rq_load_var(b, index, vars->trav.stack), 1), 1);
rq_store_var(b, index, vars->trav.origin,
nir_build_vec3_mat_mult_pre(b, rq_load_var(b, index, vars->origin),
@ -817,13 +796,12 @@ lower_rq_proceed(nir_builder *b, nir_ssa_def *index, struct ray_query_vars *vars
/* box */
for (unsigned i = 4; i-- > 0;) {
nir_ssa_def *new_node = nir_vector_extract(b, result, nir_imm_int(b, i));
nir_push_if(b, nir_ine(b, new_node, nir_imm_int(b, 0xffffffff)));
nir_push_if(b, nir_ine_imm(b, new_node, 0xffffffff));
{
rq_store_array(b, index, vars->stack, rq_load_var(b, index, vars->trav.stack),
new_node, 0x1);
rq_store_var(
b, index, vars->trav.stack,
nir_iadd(b, rq_load_var(b, index, vars->trav.stack), nir_imm_int(b, 1)), 1);
rq_store_var(b, index, vars->trav.stack,
nir_iadd_imm(b, rq_load_var(b, index, vars->trav.stack), 1), 1);
}
nir_pop_if(b, NULL);
}

2
src/amd/vulkan/radv_pipeline.c
View file

@ -3824,7 +3824,7 @@ radv_adjust_vertex_fetch_alpha(nir_builder *b,
*/
unsigned offset = alpha_adjust == ALPHA_ADJUST_SNORM ? 23u : 0u;
alpha = nir_ibfe(b, alpha, nir_imm_int(b, offset), nir_imm_int(b, 2u));
alpha = nir_ibfe_imm(b, alpha, offset, 2u);
/* Convert back to the right type. */
if (alpha_adjust == ALPHA_ADJUST_SNORM) {

295
src/amd/vulkan/radv_pipeline_rt.c
View file

@ -314,8 +314,7 @@ const uint32_t RADV_HIT_ATTRIB_OFFSET = -(16 + RADV_MAX_HIT_ATTRIB_SIZE);
static void
insert_rt_return(nir_builder *b, const struct rt_variables *vars)
{
nir_store_var(b, vars->stack_ptr,
nir_iadd(b, nir_load_var(b, vars->stack_ptr), nir_imm_int(b, -16)), 1);
nir_store_var(b, vars->stack_ptr, nir_iadd_imm(b, nir_load_var(b, vars->stack_ptr), -16), 1);
nir_store_var(b, vars->idx,
nir_load_scratch(b, 1, 32, nir_load_var(b, vars->stack_ptr), .align_mul = 16), 1);
}
@ -346,14 +345,12 @@ load_sbt_entry(nir_builder *b, const struct rt_variables *vars, nir_ssa_def *idx
{
nir_ssa_def *addr = get_sbt_ptr(b, idx, binding);
nir_ssa_def *load_addr = addr;
if (offset)
load_addr = nir_iadd(b, load_addr, nir_imm_int64(b, offset));
nir_ssa_def *load_addr = nir_iadd_imm(b, addr, offset);
nir_ssa_def *v_idx = nir_build_load_global(b, 1, 32, load_addr);
nir_store_var(b, vars->idx, v_idx, 1);
nir_ssa_def *record_addr = nir_iadd(b, addr, nir_imm_int64(b, RADV_RT_HANDLE_SIZE));
nir_ssa_def *record_addr = nir_iadd_imm(b, addr, RADV_RT_HANDLE_SIZE);
nir_store_var(b, vars->shader_record_ptr, record_addr, 1);
}
@ -376,22 +373,19 @@ lower_rt_instructions(nir_shader *shader, struct rt_variables *vars, unsigned ca
uint32_t ret = call_idx_base + nir_intrinsic_call_idx(intr) + 1;
b_shader.cursor = nir_instr_remove(instr);
nir_store_var(&b_shader, vars->stack_ptr,
nir_iadd(&b_shader, nir_load_var(&b_shader, vars->stack_ptr),
nir_imm_int(&b_shader, size)),
1);
nir_store_var(
&b_shader, vars->stack_ptr,
nir_iadd_imm(&b_shader, nir_load_var(&b_shader, vars->stack_ptr), size), 1);
nir_store_scratch(&b_shader, nir_imm_int(&b_shader, ret),
nir_load_var(&b_shader, vars->stack_ptr), .align_mul = 16);
nir_store_var(&b_shader, vars->stack_ptr,
nir_iadd(&b_shader, nir_load_var(&b_shader, vars->stack_ptr),
nir_imm_int(&b_shader, 16)),
nir_iadd_imm(&b_shader, nir_load_var(&b_shader, vars->stack_ptr), 16),
1);
load_sbt_entry(&b_shader, vars, intr->src[0].ssa, SBT_CALLABLE, 0);
nir_store_var(
&b_shader, vars->arg,
nir_isub(&b_shader, intr->src[1].ssa, nir_imm_int(&b_shader, size + 16)), 1);
nir_store_var(&b_shader, vars->arg,
nir_iadd_imm(&b_shader, intr->src[1].ssa, -size - 16), 1);
vars->stack_sizes[vars->group_idx].recursive_size =
MAX2(vars->stack_sizes[vars->group_idx].recursive_size, size + 16);
@ -402,22 +396,19 @@ lower_rt_instructions(nir_shader *shader, struct rt_variables *vars, unsigned ca
uint32_t ret = call_idx_base + nir_intrinsic_call_idx(intr) + 1;
b_shader.cursor = nir_instr_remove(instr);
nir_store_var(&b_shader, vars->stack_ptr,
nir_iadd(&b_shader, nir_load_var(&b_shader, vars->stack_ptr),
nir_imm_int(&b_shader, size)),
1);
nir_store_var(
&b_shader, vars->stack_ptr,
nir_iadd_imm(&b_shader, nir_load_var(&b_shader, vars->stack_ptr), size), 1);
nir_store_scratch(&b_shader, nir_imm_int(&b_shader, ret),
nir_load_var(&b_shader, vars->stack_ptr), .align_mul = 16);
nir_store_var(&b_shader, vars->stack_ptr,
nir_iadd(&b_shader, nir_load_var(&b_shader, vars->stack_ptr),
nir_imm_int(&b_shader, 16)),
nir_iadd_imm(&b_shader, nir_load_var(&b_shader, vars->stack_ptr), 16),
1);
nir_store_var(&b_shader, vars->idx, nir_imm_int(&b_shader, 1), 1);
nir_store_var(
&b_shader, vars->arg,
nir_isub(&b_shader, intr->src[10].ssa, nir_imm_int(&b_shader, size + 16)), 1);
nir_store_var(&b_shader, vars->arg,
nir_iadd_imm(&b_shader, intr->src[10].ssa, -size - 16), 1);
vars->stack_sizes[vars->group_idx].recursive_size =
MAX2(vars->stack_sizes[vars->group_idx].recursive_size, size + 16);
@ -426,17 +417,13 @@ lower_rt_instructions(nir_shader *shader, struct rt_variables *vars, unsigned ca
nir_store_var(&b_shader, vars->accel_struct, intr->src[0].ssa, 0x1);
nir_store_var(&b_shader, vars->flags, intr->src[1].ssa, 0x1);
nir_store_var(&b_shader, vars->cull_mask,
nir_iand(&b_shader, intr->src[2].ssa, nir_imm_int(&b_shader, 0xff)),
0x1);
nir_iand_imm(&b_shader, intr->src[2].ssa, 0xff), 0x1);
nir_store_var(&b_shader, vars->sbt_offset,
nir_iand(&b_shader, intr->src[3].ssa, nir_imm_int(&b_shader, 0xf)),
0x1);
nir_iand_imm(&b_shader, intr->src[3].ssa, 0xf), 0x1);
nir_store_var(&b_shader, vars->sbt_stride,
nir_iand(&b_shader, intr->src[4].ssa, nir_imm_int(&b_shader, 0xf)),
0x1);
nir_iand_imm(&b_shader, intr->src[4].ssa, 0xf), 0x1);
nir_store_var(&b_shader, vars->miss_index,
nir_iand(&b_shader, intr->src[5].ssa, nir_imm_int(&b_shader, 0xffff)),
0x1);
nir_iand_imm(&b_shader, intr->src[5].ssa, 0xffff), 0x1);
nir_store_var(&b_shader, vars->origin, intr->src[6].ssa, 0x7);
nir_store_var(&b_shader, vars->tmin, intr->src[7].ssa, 0x1);
nir_store_var(&b_shader, vars->direction, intr->src[8].ssa, 0x7);
@ -447,10 +434,9 @@ lower_rt_instructions(nir_shader *shader, struct rt_variables *vars, unsigned ca
uint32_t size = align(nir_intrinsic_stack_size(intr), 16) + RADV_MAX_HIT_ATTRIB_SIZE;
b_shader.cursor = nir_instr_remove(instr);
nir_store_var(&b_shader, vars->stack_ptr,
nir_iadd(&b_shader, nir_load_var(&b_shader, vars->stack_ptr),
nir_imm_int(&b_shader, -size)),
1);
nir_store_var(
&b_shader, vars->stack_ptr,
nir_iadd_imm(&b_shader, nir_load_var(&b_shader, vars->stack_ptr), -size), 1);
break;
}
case nir_intrinsic_rt_return_amd: {
@ -522,7 +508,7 @@ lower_rt_instructions(nir_shader *shader, struct rt_variables *vars, unsigned ca
case nir_intrinsic_load_ray_instance_custom_index: {
b_shader.cursor = nir_instr_remove(instr);
nir_ssa_def *ret = nir_load_var(&b_shader, vars->custom_instance_and_mask);
ret = nir_iand(&b_shader, ret, nir_imm_int(&b_shader, 0xFFFFFF));
ret = nir_iand_imm(&b_shader, ret, 0xFFFFFF);
nir_ssa_def_rewrite_uses(&intr->dest.ssa, ret);
break;
}
@ -535,7 +521,7 @@ lower_rt_instructions(nir_shader *shader, struct rt_variables *vars, unsigned ca
case nir_intrinsic_load_ray_geometry_index: {
b_shader.cursor = nir_instr_remove(instr);
nir_ssa_def *ret = nir_load_var(&b_shader, vars->geometry_id_and_flags);
ret = nir_iand(&b_shader, ret, nir_imm_int(&b_shader, 0xFFFFFFF));
ret = nir_iand_imm(&b_shader, ret, 0xFFFFFFF);
nir_ssa_def_rewrite_uses(&intr->dest.ssa, ret);
break;
}
@ -589,9 +575,8 @@ lower_rt_instructions(nir_shader *shader, struct rt_variables *vars, unsigned ca
val = nir_vec(&b_shader, vals, 3);
} else {
val = nir_build_load_global(&b_shader, 3, 32,
nir_iadd(&b_shader, instance_node_addr,
nir_imm_int64(&b_shader, 92 + c * 12)));
val = nir_build_load_global(
&b_shader, 3, 32, nir_iadd_imm(&b_shader, instance_node_addr, 92 + c * 12));
}
b_shader.cursor = nir_instr_remove(instr);
nir_ssa_def_rewrite_uses(&intr->dest.ssa, val);
@ -600,18 +585,15 @@ lower_rt_instructions(nir_shader *shader, struct rt_variables *vars, unsigned ca
case nir_intrinsic_load_ray_object_origin: {
nir_ssa_def *instance_node_addr = nir_load_var(&b_shader, vars->instance_addr);
nir_ssa_def *wto_matrix[] = {
nir_build_load_global(
&b_shader, 4, 32,
nir_iadd(&b_shader, instance_node_addr, nir_imm_int64(&b_shader, 16)),
.align_mul = 64, .align_offset = 16),
nir_build_load_global(
&b_shader, 4, 32,
nir_iadd(&b_shader, instance_node_addr, nir_imm_int64(&b_shader, 32)),
.align_mul = 64, .align_offset = 32),
nir_build_load_global(
&b_shader, 4, 32,
nir_iadd(&b_shader, instance_node_addr, nir_imm_int64(&b_shader, 48)),
.align_mul = 64, .align_offset = 48)};
nir_build_load_global(&b_shader, 4, 32,
nir_iadd_imm(&b_shader, instance_node_addr, 16),
.align_mul = 64, .align_offset = 16),
nir_build_load_global(&b_shader, 4, 32,
nir_iadd_imm(&b_shader, instance_node_addr, 32),
.align_mul = 64, .align_offset = 32),
nir_build_load_global(&b_shader, 4, 32,
nir_iadd_imm(&b_shader, instance_node_addr, 48),
.align_mul = 64, .align_offset = 48)};
nir_ssa_def *val = nir_build_vec3_mat_mult_pre(
&b_shader, nir_load_var(&b_shader, vars->origin), wto_matrix);
b_shader.cursor = nir_instr_remove(instr);
@ -718,7 +700,7 @@ insert_rt_case(nir_builder *b, nir_shader *shader, const struct rt_variables *va
MAX2(src_vars.stack_sizes[src_vars.group_idx].recursive_size, shader->scratch_size);
}
nir_push_if(b, nir_ieq(b, idx, nir_imm_int(b, call_idx)));
nir_push_if(b, nir_ieq_imm(b, idx, call_idx));
nir_store_var(b, vars->main_loop_case_visited, nir_imm_bool(b, true), 1);
nir_inline_function_impl(b, nir_shader_get_entrypoint(shader), NULL, var_remap);
nir_pop_if(b, NULL);
@ -1070,7 +1052,7 @@ visit_any_hit_shaders(struct radv_device *device,
{
nir_ssa_def *sbt_idx = nir_load_var(b, vars->idx);
nir_push_if(b, nir_ine(b, sbt_idx, nir_imm_int(b, 0)));
nir_push_if(b, nir_ine_imm(b, sbt_idx, 0));
for (unsigned i = 0; i < pCreateInfo->groupCount; ++i) {
const VkRayTracingShaderGroupCreateInfoKHR *group_info = &pCreateInfo->pGroups[i];
uint32_t shader_id = VK_SHADER_UNUSED_KHR;
@ -1104,34 +1086,30 @@ insert_traversal_triangle_case(struct radv_device *device,
nir_ssa_def *div = nir_vector_extract(b, result, nir_imm_int(b, 1));
dist = nir_fdiv(b, dist, div);
nir_ssa_def *frontface = nir_flt(b, nir_imm_float(b, 0), div);
nir_ssa_def *switch_ccw = nir_ine(
b,
nir_iand(
b, nir_load_var(b, trav_vars->sbt_offset_and_flags),
nir_imm_int(b, VK_GEOMETRY_INSTANCE_TRIANGLE_FRONT_COUNTERCLOCKWISE_BIT_KHR << 24)),
nir_imm_int(b, 0));
nir_ssa_def *switch_ccw =
nir_ine_imm(b,
nir_iand_imm(b, nir_load_var(b, trav_vars->sbt_offset_and_flags),
VK_GEOMETRY_INSTANCE_TRIANGLE_FRONT_COUNTERCLOCKWISE_BIT_KHR << 24),
0);
frontface = nir_ixor(b, frontface, switch_ccw);
nir_ssa_def *not_cull = nir_ieq(
b, nir_iand(b, nir_load_var(b, vars->flags), nir_imm_int(b, SpvRayFlagsSkipTrianglesKHRMask)),
nir_imm_int(b, 0));
nir_ssa_def *not_facing_cull = nir_ieq(
nir_ssa_def *not_cull = nir_ieq_imm(
b, nir_iand_imm(b, nir_load_var(b, vars->flags), SpvRayFlagsSkipTrianglesKHRMask), 0);
nir_ssa_def *not_facing_cull = nir_ieq_imm(
b,
nir_iand(b, nir_load_var(b, vars->flags),
nir_bcsel(b, frontface, nir_imm_int(b, SpvRayFlagsCullFrontFacingTrianglesKHRMask),
nir_imm_int(b, SpvRayFlagsCullBackFacingTrianglesKHRMask))),
nir_imm_int(b, 0));
0);
not_cull = nir_iand(
b, not_cull,
nir_ior(
b, not_facing_cull,
nir_ine(
b,
nir_iand(
b, nir_load_var(b, trav_vars->sbt_offset_and_flags),
nir_imm_int(b, VK_GEOMETRY_INSTANCE_TRIANGLE_FACING_CULL_DISABLE_BIT_KHR << 24)),
nir_imm_int(b, 0))));
nir_ine_imm(b,
nir_iand_imm(b, nir_load_var(b, trav_vars->sbt_offset_and_flags),
VK_GEOMETRY_INSTANCE_TRIANGLE_FACING_CULL_DISABLE_BIT_KHR << 24),
0)));
nir_push_if(b, nir_iand(b,
nir_iand(b, nir_flt(b, dist, nir_load_var(b, vars->tmax)),
@ -1139,38 +1117,36 @@ insert_traversal_triangle_case(struct radv_device *device,
not_cull));
{
nir_ssa_def *triangle_info = nir_build_load_global(
b, 2, 32,
nir_iadd(b, build_node_to_addr(device, b, bvh_node),
nir_imm_int64(b, offsetof(struct radv_bvh_triangle_node, triangle_id))));
nir_ssa_def *triangle_info =
nir_build_load_global(b, 2, 32,
nir_iadd_imm(b, build_node_to_addr(device, b, bvh_node),
offsetof(struct radv_bvh_triangle_node, triangle_id)));
nir_ssa_def *primitive_id = nir_channel(b, triangle_info, 0);
nir_ssa_def *geometry_id_and_flags = nir_channel(b, triangle_info, 1);
nir_ssa_def *geometry_id = nir_iand(b, geometry_id_and_flags, nir_imm_int(b, 0xfffffff));
nir_ssa_def *geometry_id = nir_iand_imm(b, geometry_id_and_flags, 0xfffffff);
nir_ssa_def *is_opaque = hit_is_opaque(b, nir_load_var(b, trav_vars->sbt_offset_and_flags),
nir_load_var(b, vars->flags), geometry_id_and_flags);
not_cull =
nir_ieq(b,
nir_iand(b, nir_load_var(b, vars->flags),
nir_bcsel(b, is_opaque, nir_imm_int(b, SpvRayFlagsCullOpaqueKHRMask),
nir_imm_int(b, SpvRayFlagsCullNoOpaqueKHRMask))),
nir_imm_int(b, 0));
nir_ieq_imm(b,
nir_iand(b, nir_load_var(b, vars->flags),
nir_bcsel(b, is_opaque, nir_imm_int(b, SpvRayFlagsCullOpaqueKHRMask),
nir_imm_int(b, SpvRayFlagsCullNoOpaqueKHRMask))),
0);
nir_push_if(b, not_cull);
{
nir_ssa_def *sbt_idx =
nir_iadd(b,
nir_iadd(b, nir_load_var(b, vars->sbt_offset),
nir_iand(b, nir_load_var(b, trav_vars->sbt_offset_and_flags),
nir_imm_int(b, 0xffffff))),
nir_imul(b, nir_load_var(b, vars->sbt_stride), geometry_id));
nir_ssa_def *sbt_idx = nir_iadd(
b,
nir_iadd(b, nir_load_var(b, vars->sbt_offset),
nir_iand_imm(b, nir_load_var(b, trav_vars->sbt_offset_and_flags), 0xffffff)),
nir_imul(b, nir_load_var(b, vars->sbt_stride), geometry_id));
nir_ssa_def *divs[2] = {div, div};
nir_ssa_def *ij = nir_fdiv(b, nir_channels(b, result, 0xc), nir_vec(b, divs, 2));
nir_ssa_def *hit_kind =
nir_bcsel(b, frontface, nir_imm_int(b, 0xFE), nir_imm_int(b, 0xFF));
nir_store_scratch(
b, ij,
nir_iadd(b, nir_load_var(b, vars->stack_ptr), nir_imm_int(b, RADV_HIT_ATTRIB_OFFSET)),
b, ij, nir_iadd_imm(b, nir_load_var(b, vars->stack_ptr), RADV_HIT_ATTRIB_OFFSET),
.align_mul = 16);
nir_store_var(b, vars->ahit_status, nir_imm_int(b, 0), 1);
@ -1193,7 +1169,7 @@ insert_traversal_triangle_case(struct radv_device *device,
visit_any_hit_shaders(device, pCreateInfo, b, &inner_vars);
nir_push_if(b, nir_ieq(b, nir_load_var(b, vars->ahit_status), nir_imm_int(b, 1)));
nir_push_if(b, nir_ieq_imm(b, nir_load_var(b, vars->ahit_status), 1));
{
nir_jump(b, nir_jump_continue);
}
@ -1214,20 +1190,17 @@ insert_traversal_triangle_case(struct radv_device *device,
nir_store_var(b, trav_vars->should_return,
nir_ior(b,
nir_ine(b,
nir_iand(b, nir_load_var(b, vars->flags),
nir_imm_int(b, SpvRayFlagsSkipClosestHitShaderKHRMask)),
nir_imm_int(b, 0)),
nir_ieq(b, nir_load_var(b, vars->idx), nir_imm_int(b, 0))),
nir_ine_imm(b,
nir_iand_imm(b, nir_load_var(b, vars->flags),
SpvRayFlagsSkipClosestHitShaderKHRMask),
0),
nir_ieq_imm(b, nir_load_var(b, vars->idx), 0)),
1);
nir_ssa_def *terminate_on_first_hit =
nir_ine(b,
nir_iand(b, nir_load_var(b, vars->flags),
nir_imm_int(b, SpvRayFlagsTerminateOnFirstHitKHRMask)),
nir_imm_int(b, 0));
nir_ssa_def *ray_terminated =
nir_ieq(b, nir_load_var(b, vars->ahit_status), nir_imm_int(b, 2));
nir_ssa_def *terminate_on_first_hit = nir_ine_imm(
b, nir_iand_imm(b, nir_load_var(b, vars->flags), SpvRayFlagsTerminateOnFirstHitKHRMask),
0);
nir_ssa_def *ray_terminated = nir_ieq_imm(b, nir_load_var(b, vars->ahit_status), 2);
nir_push_if(b, nir_ior(b, terminate_on_first_hit, ray_terminated));
{
nir_jump(b, nir_jump_break);
@ -1246,31 +1219,29 @@ insert_traversal_aabb_case(struct radv_device *device,
const struct rt_traversal_vars *trav_vars, nir_ssa_def *bvh_node)
{
nir_ssa_def *node_addr = build_node_to_addr(device, b, bvh_node);
nir_ssa_def *triangle_info =
nir_build_load_global(b, 2, 32, nir_iadd(b, node_addr, nir_imm_int64(b, 24)));
nir_ssa_def *triangle_info = nir_build_load_global(b, 2, 32, nir_iadd_imm(b, node_addr, 24));
nir_ssa_def *primitive_id = nir_channel(b, triangle_info, 0);
nir_ssa_def *geometry_id_and_flags = nir_channel(b, triangle_info, 1);
nir_ssa_def *geometry_id = nir_iand(b, geometry_id_and_flags, nir_imm_int(b, 0xfffffff));
nir_ssa_def *geometry_id = nir_iand_imm(b, geometry_id_and_flags, 0xfffffff);
nir_ssa_def *is_opaque = hit_is_opaque(b, nir_load_var(b, trav_vars->sbt_offset_and_flags),
nir_load_var(b, vars->flags), geometry_id_and_flags);
nir_ssa_def *not_skip_aabb = nir_ieq(
b, nir_iand(b, nir_load_var(b, vars->flags), nir_imm_int(b, SpvRayFlagsSkipAABBsKHRMask)),
nir_imm_int(b, 0));
nir_ssa_def *not_cull =
nir_iand(b, not_skip_aabb, nir_ieq(b,
nir_iand(b, nir_load_var(b, vars->flags),
nir_bcsel(b, is_opaque, nir_imm_int(b, SpvRayFlagsCullOpaqueKHRMask),
nir_imm_int(b, SpvRayFlagsCullNoOpaqueKHRMask))),
nir_imm_int(b, 0)));
nir_ssa_def *not_skip_aabb =
nir_ieq_imm(b, nir_iand_imm(b, nir_load_var(b, vars->flags), SpvRayFlagsSkipAABBsKHRMask), 0);
nir_ssa_def *not_cull = nir_iand(
b, not_skip_aabb,
nir_ieq_imm(b,
nir_iand(b, nir_load_var(b, vars->flags),
nir_bcsel(b, is_opaque, nir_imm_int(b, SpvRayFlagsCullOpaqueKHRMask),
nir_imm_int(b, SpvRayFlagsCullNoOpaqueKHRMask))),
0));
nir_push_if(b, not_cull);
{
nir_ssa_def *sbt_idx =
nir_iadd(b,
nir_iadd(b, nir_load_var(b, vars->sbt_offset),
nir_iand(b, nir_load_var(b, trav_vars->sbt_offset_and_flags),
nir_imm_int(b, 0xffffff))),
nir_imul(b, nir_load_var(b, vars->sbt_stride), geometry_id));
nir_ssa_def *sbt_idx = nir_iadd(
b,
nir_iadd(b, nir_load_var(b, vars->sbt_offset),
nir_iand_imm(b, nir_load_var(b, trav_vars->sbt_offset_and_flags), 0xffffff)),
nir_imul(b, nir_load_var(b, vars->sbt_stride), geometry_id));
struct rt_variables inner_vars = create_inner_vars(b, vars);
@ -1291,7 +1262,7 @@ insert_traversal_aabb_case(struct radv_device *device,
nir_store_var(b, vars->ahit_status, nir_imm_int(b, 1), 1);
nir_push_if(b, nir_ine(b, nir_load_var(b, inner_vars.idx), nir_imm_int(b, 0)));
nir_push_if(b, nir_ine_imm(b, nir_load_var(b, inner_vars.idx), 0));
for (unsigned i = 0; i < pCreateInfo->groupCount; ++i) {
const VkRayTracingShaderGroupCreateInfoKHR *group_info = &pCreateInfo->pGroups[i];
uint32_t shader_id = VK_SHADER_UNUSED_KHR;
@ -1329,10 +1300,8 @@ insert_traversal_aabb_case(struct radv_device *device,
nir_ssa_def *vec3_inf =
nir_channels(b, nir_imm_vec4(b, INFINITY, INFINITY, INFINITY, 0), 0x7);
nir_ssa_def *bvh_lo =
nir_build_load_global(b, 3, 32, nir_iadd(b, node_addr, nir_imm_int64(b, 0)));
nir_ssa_def *bvh_hi =
nir_build_load_global(b, 3, 32, nir_iadd(b, node_addr, nir_imm_int64(b, 12)));
nir_ssa_def *bvh_lo = nir_build_load_global(b, 3, 32, nir_iadd_imm(b, node_addr, 0));
nir_ssa_def *bvh_hi = nir_build_load_global(b, 3, 32, nir_iadd_imm(b, node_addr, 12));
bvh_lo = nir_fsub(b, bvh_lo, nir_load_var(b, trav_vars->origin));
bvh_hi = nir_fsub(b, bvh_hi, nir_load_var(b, trav_vars->origin));
@ -1360,7 +1329,7 @@ insert_traversal_aabb_case(struct radv_device *device,
}
nir_pop_if(b, NULL);
nir_push_if(b, nir_ine(b, nir_load_var(b, vars->ahit_status), nir_imm_int(b, 1)));
nir_push_if(b, nir_ine_imm(b, nir_load_var(b, vars->ahit_status), 1));
{
nir_store_var(b, vars->primitive_id, primitive_id, 1);
nir_store_var(b, vars->geometry_id_and_flags, geometry_id_and_flags, 1);
@ -1374,20 +1343,17 @@ insert_traversal_aabb_case(struct radv_device *device,
nir_store_var(b, trav_vars->should_return,
nir_ior(b,
nir_ine(b,
nir_iand(b, nir_load_var(b, vars->flags),
nir_imm_int(b, SpvRayFlagsSkipClosestHitShaderKHRMask)),
nir_imm_int(b, 0)),
nir_ieq(b, nir_load_var(b, vars->idx), nir_imm_int(b, 0))),
nir_ine_imm(b,
nir_iand_imm(b, nir_load_var(b, vars->flags),
SpvRayFlagsSkipClosestHitShaderKHRMask),
0),
nir_ieq_imm(b, nir_load_var(b, vars->idx), 0)),
1);
nir_ssa_def *terminate_on_first_hit =
nir_ine(b,
nir_iand(b, nir_load_var(b, vars->flags),
nir_imm_int(b, SpvRayFlagsTerminateOnFirstHitKHRMask)),
nir_imm_int(b, 0));
nir_ssa_def *ray_terminated =
nir_ieq(b, nir_load_var(b, vars->ahit_status), nir_imm_int(b, 2));
nir_ssa_def *terminate_on_first_hit = nir_ine_imm(
b, nir_iand_imm(b, nir_load_var(b, vars->flags), SpvRayFlagsTerminateOnFirstHitKHRMask),
0);
nir_ssa_def *ray_terminated = nir_ieq_imm(b, nir_load_var(b, vars->ahit_status), 2);
nir_push_if(b, nir_ior(b, terminate_on_first_hit, ray_terminated));
{
nir_jump(b, nir_jump_break);
@ -1409,8 +1375,8 @@ insert_traversal(struct radv_device *device, const VkRayTracingPipelineCreateInf
unsigned stack_entry_stride = stack_entry_size * lanes;
nir_ssa_def *stack_entry_stride_def = nir_imm_int(b, stack_entry_stride);
nir_ssa_def *stack_base =
nir_iadd(b, nir_imm_int(b, b->shader->info.shared_size),
nir_imul(b, nir_load_local_invocation_index(b), nir_imm_int(b, stack_entry_size)));
nir_iadd_imm(b, nir_imul_imm(b, nir_load_local_invocation_index(b), stack_entry_size),
b->shader->info.shared_size);
b->shader->info.shared_size += stack_entry_stride * MAX_STACK_ENTRY_COUNT;
assert(b->shader->info.shared_size <= 32768);
@ -1425,7 +1391,7 @@ insert_traversal(struct radv_device *device, const VkRayTracingPipelineCreateInf
nir_store_var(b, trav_vars.should_return, nir_imm_bool(b, false), 1);
nir_push_if(b, nir_ine(b, accel_struct, nir_imm_int64(b, 0)));
nir_push_if(b, nir_ine_imm(b, accel_struct, 0));
{
nir_store_var(b, trav_vars.bvh_base, build_addr_to_node(b, accel_struct), 1);
@ -1469,7 +1435,7 @@ insert_traversal(struct radv_device *device, const VkRayTracingPipelineCreateInf
nir_ssa_def *bvh_node = nir_load_shared(b, 1, 32, nir_load_var(b, trav_vars.stack), .base = 0,
.align_mul = stack_entry_size);
nir_ssa_def *bvh_node_type = nir_iand(b, bvh_node, nir_imm_int(b, 7));
nir_ssa_def *bvh_node_type = nir_iand_imm(b, bvh_node, 7);
bvh_node = nir_iadd(b, nir_load_var(b, trav_vars.bvh_base), nir_u2u(b, bvh_node, 64));
nir_ssa_def *intrinsic_result = NULL;
@ -1480,14 +1446,12 @@ insert_traversal(struct radv_device *device, const VkRayTracingPipelineCreateInf
nir_load_var(b, trav_vars.inv_dir));
}
nir_push_if(b, nir_ine(b, nir_iand(b, bvh_node_type, nir_imm_int(b, 4)), nir_imm_int(b, 0)));
nir_push_if(b, nir_ine_imm(b, nir_iand_imm(b, bvh_node_type, 4), 0));
{
nir_push_if(b,
nir_ine(b, nir_iand(b, bvh_node_type, nir_imm_int(b, 2)), nir_imm_int(b, 0)));
nir_push_if(b, nir_ine_imm(b, nir_iand_imm(b, bvh_node_type, 2), 0));
{
/* custom */
nir_push_if(
b, nir_ine(b, nir_iand(b, bvh_node_type, nir_imm_int(b, 1)), nir_imm_int(b, 0)));
nir_push_if(b, nir_ine_imm(b, nir_iand_imm(b, bvh_node_type, 1), 0));
if (!(pCreateInfo->flags & VK_PIPELINE_CREATE_RAY_TRACING_SKIP_AABBS_BIT_KHR)) {
insert_traversal_aabb_case(device, pCreateInfo, b, vars, &trav_vars, bvh_node);
}
@ -1498,23 +1462,20 @@ insert_traversal(struct radv_device *device, const VkRayTracingPipelineCreateInf
nir_ssa_def *instance_data =
nir_build_load_global(b, 4, 32, instance_node_addr, .align_mul = 64);
nir_ssa_def *wto_matrix[] = {
nir_build_load_global(b, 4, 32,
nir_iadd(b, instance_node_addr, nir_imm_int64(b, 16)),
nir_build_load_global(b, 4, 32, nir_iadd_imm(b, instance_node_addr, 16),
.align_mul = 64, .align_offset = 16),
nir_build_load_global(b, 4, 32,
nir_iadd(b, instance_node_addr, nir_imm_int64(b, 32)),
nir_build_load_global(b, 4, 32, nir_iadd_imm(b, instance_node_addr, 32),
.align_mul = 64, .align_offset = 32),
nir_build_load_global(b, 4, 32,
nir_iadd(b, instance_node_addr, nir_imm_int64(b, 48)),
nir_build_load_global(b, 4, 32, nir_iadd_imm(b, instance_node_addr, 48),
.align_mul = 64, .align_offset = 48)};
nir_ssa_def *instance_id = nir_build_load_global(
b, 1, 32, nir_iadd(b, instance_node_addr, nir_imm_int64(b, 88)));
nir_ssa_def *instance_id =
nir_build_load_global(b, 1, 32, nir_iadd_imm(b, instance_node_addr, 88));
nir_ssa_def *instance_and_mask = nir_channel(b, instance_data, 2);
nir_ssa_def *instance_mask = nir_ushr(b, instance_and_mask, nir_imm_int(b, 24));
nir_ssa_def *instance_mask = nir_ushr_imm(b, instance_and_mask, 24);
nir_push_if(b,
nir_ieq(b, nir_iand(b, instance_mask, nir_load_var(b, vars->cull_mask)),
nir_imm_int(b, 0)));
nir_push_if(
b,
nir_ieq_imm(b, nir_iand(b, instance_mask, nir_load_var(b, vars->cull_mask)), 0));
nir_jump(b, nir_jump_continue);
nir_pop_if(b, NULL);
@ -1523,9 +1484,9 @@ insert_traversal(struct radv_device *device, const VkRayTracingPipelineCreateInf
build_addr_to_node(
b, nir_pack_64_2x32(b, nir_channels(b, instance_data, 0x3))),
1);
nir_store_shared(
b, nir_iand(b, nir_channel(b, instance_data, 0), nir_imm_int(b, 63)),
nir_load_var(b, trav_vars.stack), .base = 0, .align_mul = stack_entry_size);
nir_store_shared(b, nir_iand_imm(b, nir_channel(b, instance_data, 0), 63),
nir_load_var(b, trav_vars.stack), .base = 0,
.align_mul = stack_entry_size);
nir_store_var(b, trav_vars.stack,
nir_iadd(b, nir_load_var(b, trav_vars.stack), stack_entry_stride_def),
1);
@ -1561,7 +1522,7 @@ insert_traversal(struct radv_device *device, const VkRayTracingPipelineCreateInf
for (unsigned i = 4; i-- > 0; ) {
nir_ssa_def *new_node = nir_vector_extract(b, result, nir_imm_int(b, i));
nir_push_if(b, nir_ine(b, new_node, nir_imm_int(b, 0xffffffff)));
nir_push_if(b, nir_ine_imm(b, new_node, 0xffffffff));
{
nir_store_shared(b, new_node, nir_load_var(b, trav_vars.stack), .base = 0,
.align_mul = stack_entry_size);
@ -1603,7 +1564,7 @@ insert_traversal(struct radv_device *device, const VkRayTracingPipelineCreateInf
/* Only load the miss shader if we actually miss, which we determining by not having set
* a closest hit shader. It is valid to not specify an SBT pointer for miss shaders if none
* of the rays miss. */
nir_push_if(b, nir_ieq(b, nir_load_var(b, vars->idx), nir_imm_int(b, 0)));
nir_push_if(b, nir_ieq_imm(b, nir_load_var(b, vars->idx), 0));
{
load_sbt_entry(b, vars, nir_load_var(b, vars->miss_index), SBT_MISS, 0);
}
@ -1704,7 +1665,7 @@ create_rt_shader(struct radv_device *device, const VkRayTracingPipelineCreateInf
nir_loop *loop = nir_push_loop(&b);
nir_push_if(&b, nir_ior(&b, nir_ieq(&b, nir_load_var(&b, vars.idx), nir_imm_int(&b, 0)),
nir_push_if(&b, nir_ior(&b, nir_ieq_imm(&b, nir_load_var(&b, vars.idx), 0),
nir_ine(&b, nir_load_var(&b, vars.main_loop_case_visited),
nir_imm_bool(&b, true))));
nir_jump(&b, nir_jump_break);
@ -1712,7 +1673,7 @@ create_rt_shader(struct radv_device *device, const VkRayTracingPipelineCreateInf
nir_store_var(&b, vars.main_loop_case_visited, nir_imm_bool(&b, false), 1);
nir_push_if(&b, nir_ieq(&b, nir_load_var(&b, vars.idx), nir_imm_int(&b, 1)));
nir_push_if(&b, nir_ieq_imm(&b, nir_load_var(&b, vars.idx), 1));
nir_store_var(&b, vars.main_loop_case_visited, nir_imm_bool(&b, true), 1);
insert_traversal(device, pCreateInfo, &b, &vars);
nir_pop_if(&b, NULL);

27
src/amd/vulkan/radv_query.c
View file

@ -52,7 +52,7 @@ radv_get_pipeline_statistics_index(const VkQueryPipelineStatisticFlagBits flag)
static nir_ssa_def *
nir_test_flag(nir_builder *b, nir_ssa_def *flags, uint32_t flag)
{
return nir_i2b(b, nir_iand(b, flags, nir_imm_int(b, flag)));
return nir_i2b(b, nir_iand_imm(b, flags, flag));
}
static void
@ -149,12 +149,12 @@ build_occlusion_query_shader(struct radv_device *device)
nir_ssa_def *current_outer_count = nir_load_var(&b, outer_counter);
radv_break_on_count(&b, outer_counter, nir_imm_int(&b, db_count));
nir_ssa_def *enabled_cond = nir_iand(&b, nir_imm_int(&b, enabled_rb_mask),
nir_ishl(&b, nir_imm_int(&b, 1), current_outer_count));
nir_ssa_def *enabled_cond =
nir_iand_imm(&b, nir_ishl(&b, nir_imm_int(&b, 1), current_outer_count), enabled_rb_mask);
nir_push_if(&b, nir_i2b(&b, enabled_cond));
nir_ssa_def *load_offset = nir_imul(&b, current_outer_count, nir_imm_int(&b, 16));
nir_ssa_def *load_offset = nir_imul_imm(&b, current_outer_count, 16);
load_offset = nir_iadd(&b, input_base, load_offset);
nir_ssa_def *load = nir_load_ssbo(&b, 2, 64, src_buf, load_offset, .align_mul = 16);
@ -271,13 +271,13 @@ build_pipeline_statistics_query_shader(struct radv_device *device)
nir_ssa_def *output_stride = nir_load_push_constant(&b, 1, 32, nir_imm_int(&b, 4), .range = 16);
nir_ssa_def *output_base = nir_imul(&b, output_stride, global_id);
avail_offset = nir_iadd(&b, avail_offset, nir_imul(&b, global_id, nir_imm_int(&b, 4)));
avail_offset = nir_iadd(&b, avail_offset, nir_imul_imm(&b, global_id, 4));
nir_ssa_def *available32 = nir_load_ssbo(&b, 1, 32, src_buf, avail_offset);
nir_ssa_def *result_is_64bit = nir_test_flag(&b, flags, VK_QUERY_RESULT_64_BIT);
nir_ssa_def *elem_size = nir_bcsel(&b, result_is_64bit, nir_imm_int(&b, 8), nir_imm_int(&b, 4));
nir_ssa_def *elem_count = nir_ushr(&b, stats_mask, nir_imm_int(&b, 16));
nir_ssa_def *elem_count = nir_ushr_imm(&b, stats_mask, 16);
radv_store_availability(&b, flags, dst_buf,
nir_iadd(&b, output_base, nir_imul(&b, elem_count, elem_size)),
@ -289,13 +289,11 @@ build_pipeline_statistics_query_shader(struct radv_device *device)
for (int i = 0; i < ARRAY_SIZE(pipeline_statistics_indices); ++i) {
nir_push_if(&b, nir_test_flag(&b, stats_mask, 1u << i));
nir_ssa_def *start_offset =
nir_iadd(&b, input_base, nir_imm_int(&b, pipeline_statistics_indices[i] * 8));
nir_ssa_def *start_offset = nir_iadd_imm(&b, input_base, pipeline_statistics_indices[i] * 8);
nir_ssa_def *start = nir_load_ssbo(&b, 1, 64, src_buf, start_offset);
nir_ssa_def *end_offset =
nir_iadd(&b, input_base,
nir_imm_int(&b, pipeline_statistics_indices[i] * 8 + pipelinestat_block_size));
nir_iadd_imm(&b, input_base, pipeline_statistics_indices[i] * 8 + pipelinestat_block_size);
nir_ssa_def *end = nir_load_ssbo(&b, 1, 64, src_buf, end_offset);
nir_ssa_def *result = nir_isub(&b, end, start);
@ -414,15 +412,15 @@ build_tfb_query_shader(struct radv_device *device)
/* Load data from the query pool. */
nir_ssa_def *load1 = nir_load_ssbo(&b, 4, 32, src_buf, input_base, .align_mul = 32);
nir_ssa_def *load2 = nir_load_ssbo(
&b, 4, 32, src_buf, nir_iadd(&b, input_base, nir_imm_int(&b, 16)), .align_mul = 16);
nir_ssa_def *load2 =
nir_load_ssbo(&b, 4, 32, src_buf, nir_iadd_imm(&b, input_base, 16), .align_mul = 16);
/* Check if result is available. */
nir_ssa_def *avails[2];
avails[0] = nir_iand(&b, nir_channel(&b, load1, 1), nir_channel(&b, load1, 3));
avails[1] = nir_iand(&b, nir_channel(&b, load2, 1), nir_channel(&b, load2, 3));
nir_ssa_def *result_is_available =
nir_i2b(&b, nir_iand(&b, nir_iand(&b, avails[0], avails[1]), nir_imm_int(&b, 0x80000000)));
nir_i2b(&b, nir_iand_imm(&b, nir_iand(&b, avails[0], avails[1]), 0x80000000));
/* Only compute result if available. */
nir_push_if(&b, result_is_available);
@ -541,8 +539,7 @@ build_timestamp_query_shader(struct radv_device *device)
nir_pack_64_2x32(&b, nir_vec2(&b, nir_channel(&b, load, 0), nir_channel(&b, load, 1)));
/* Check if result is available. */
nir_ssa_def *result_is_available =
nir_i2b(&b, nir_ine(&b, timestamp, nir_imm_int64(&b, TIMESTAMP_NOT_READY)));
nir_ssa_def *result_is_available = nir_i2b(&b, nir_ine_imm(&b, timestamp, TIMESTAMP_NOT_READY));
/* Only store result if available. */
nir_push_if(&b, result_is_available);

61
src/amd/vulkan/radv_rt_common.c
View file

@ -101,13 +101,13 @@ intersect_ray_amd_software_box(struct radv_device *device, nir_builder *b, nir_s
/* node->children[i] -> uint */
nir_ssa_def *child_index =
nir_build_load_global(b, 1, 32, nir_iadd(b, node_addr, nir_imm_int64(b, child_offset)),
.align_mul = 64, .align_offset = child_offset % 64);
nir_build_load_global(b, 1, 32, nir_iadd_imm(b, node_addr, child_offset), .align_mul = 64,
.align_offset = child_offset % 64);
/* node->coords[i][0], node->coords[i][1] -> vec3 */
nir_ssa_def *node_coords[2] = {
nir_build_load_global(b, 3, 32, nir_iadd(b, node_addr, nir_imm_int64(b, coord_offsets[0])),
nir_build_load_global(b, 3, 32, nir_iadd_imm(b, node_addr, coord_offsets[0]),
.align_mul = 64, .align_offset = coord_offsets[0] % 64),
nir_build_load_global(b, 3, 32, nir_iadd(b, node_addr, nir_imm_int64(b, coord_offsets[1])),
nir_build_load_global(b, 3, 32, nir_iadd_imm(b, node_addr, coord_offsets[1]),
.align_mul = 64, .align_offset = coord_offsets[1] % 64),
};
@ -185,12 +185,12 @@ intersect_ray_amd_software_tri(struct radv_device *device, nir_builder *b, nir_s
/* node->coords[0], node->coords[1], node->coords[2] -> vec3 */
nir_ssa_def *node_coords[3] = {
nir_build_load_global(b, 3, 32, nir_iadd(b, node_addr, nir_imm_int64(b, coord_offsets[0])),
.align_mul = 64, .align_offset = coord_offsets[0] % 64),
nir_build_load_global(b, 3, 32, nir_iadd(b, node_addr, nir_imm_int64(b, coord_offsets[1])),
.align_mul = 64, .align_offset = coord_offsets[1] % 64),
nir_build_load_global(b, 3, 32, nir_iadd(b, node_addr, nir_imm_int64(b, coord_offsets[2])),
.align_mul = 64, .align_offset = coord_offsets[2] % 64),
nir_build_load_global(b, 3, 32, nir_iadd_imm(b, node_addr, coord_offsets[0]), .align_mul = 64,
.align_offset = coord_offsets[0] % 64),
nir_build_load_global(b, 3, 32, nir_iadd_imm(b, node_addr, coord_offsets[1]), .align_mul = 64,
.align_offset = coord_offsets[1] % 64),
nir_build_load_global(b, 3, 32, nir_iadd_imm(b, node_addr, coord_offsets[2]), .align_mul = 64,
.align_offset = coord_offsets[2] % 64),
};
nir_variable *result = nir_variable_create(b->shader, nir_var_shader_temp, vec4_type, "result");
@ -212,8 +212,8 @@ intersect_ray_amd_software_tri(struct radv_device *device, nir_builder *b, nir_s
b, nir_fge(b, abs_dirs[0], abs_dirs[1]),
nir_bcsel(b, nir_fge(b, abs_dirs[0], abs_dirs[2]), nir_imm_int(b, 0), nir_imm_int(b, 2)),
nir_bcsel(b, nir_fge(b, abs_dirs[1], abs_dirs[2]), nir_imm_int(b, 1), nir_imm_int(b, 2)));
nir_ssa_def *kx = nir_imod(b, nir_iadd(b, kz, nir_imm_int(b, 1)), nir_imm_int(b, 3));
nir_ssa_def *ky = nir_imod(b, nir_iadd(b, kx, nir_imm_int(b, 1)), nir_imm_int(b, 3));
nir_ssa_def *kx = nir_imod(b, nir_iadd_imm(b, kz, 1), nir_imm_int(b, 3));
nir_ssa_def *ky = nir_imod(b, nir_iadd_imm(b, kx, 1), nir_imm_int(b, 3));
nir_ssa_def *k_indices[3] = {kx, ky, kz};
nir_ssa_def *k = nir_vec(b, k_indices, 3);
@ -337,19 +337,19 @@ nir_ssa_def *
build_addr_to_node(nir_builder *b, nir_ssa_def *addr)
{
const uint64_t bvh_size = 1ull << 42;
nir_ssa_def *node = nir_ushr(b, addr, nir_imm_int(b, 3));
return nir_iand(b, node, nir_imm_int64(b, (bvh_size - 1) << 3));
nir_ssa_def *node = nir_ushr_imm(b, addr, 3);
return nir_iand_imm(b, node, (bvh_size - 1) << 3);
}
nir_ssa_def *
build_node_to_addr(struct radv_device *device, nir_builder *b, nir_ssa_def *node)
{
nir_ssa_def *addr = nir_iand(b, node, nir_imm_int64(b, ~7ull));
addr = nir_ishl(b, addr, nir_imm_int(b, 3));
nir_ssa_def *addr = nir_iand_imm(b, node, ~7ull);
addr = nir_ishl_imm(b, addr, 3);
/* Assumes everything is in the top half of address space, which is true in
* GFX9+ for now. */
return device->physical_device->rad_info.chip_class >= GFX9
? nir_ior(b, addr, nir_imm_int64(b, 0xffffull << 48))
? nir_ior_imm(b, addr, 0xffffull << 48)
: addr;
}
@ -388,8 +388,7 @@ nir_build_wto_matrix_load(nir_builder *b, nir_ssa_def *instance_addr, nir_ssa_de
{
unsigned offset = offsetof(struct radv_bvh_instance_node, wto_matrix);
for (unsigned i = 0; i < 3; ++i) {
out[i] = nir_build_load_global(b, 4, 32,
nir_iadd(b, instance_addr, nir_imm_int64(b, offset + i * 16)),
out[i] = nir_build_load_global(b, 4, 32, nir_iadd_imm(b, instance_addr, offset + i * 16),
.align_mul = 64, .align_offset = offset + i * 16);
}
}
@ -400,28 +399,22 @@ nir_ssa_def *
hit_is_opaque(nir_builder *b, nir_ssa_def *sbt_offset_and_flags, nir_ssa_def *flags,
nir_ssa_def *geometry_id_and_flags)
{
nir_ssa_def *geom_force_opaque = nir_ine(
b, nir_iand(b, geometry_id_and_flags, nir_imm_int(b, VK_GEOMETRY_OPAQUE_BIT_KHR << 28)),
nir_imm_int(b, 0));
nir_ssa_def *instance_force_opaque =
nir_ine(b,
nir_iand(b, sbt_offset_and_flags,
nir_imm_int(b, VK_GEOMETRY_INSTANCE_FORCE_OPAQUE_BIT_KHR << 24)),
nir_imm_int(b, 0));
nir_ssa_def *instance_force_non_opaque =
nir_ine(b,
nir_iand(b, sbt_offset_and_flags,
nir_imm_int(b, VK_GEOMETRY_INSTANCE_FORCE_NO_OPAQUE_BIT_KHR << 24)),
nir_imm_int(b, 0));
nir_ssa_def *geom_force_opaque =
nir_ine_imm(b, nir_iand_imm(b, geometry_id_and_flags, VK_GEOMETRY_OPAQUE_BIT_KHR << 28), 0);
nir_ssa_def *instance_force_opaque = nir_ine_imm(
b, nir_iand_imm(b, sbt_offset_and_flags, VK_GEOMETRY_INSTANCE_FORCE_OPAQUE_BIT_KHR << 24), 0);
nir_ssa_def *instance_force_non_opaque = nir_ine_imm(
b, nir_iand_imm(b, sbt_offset_and_flags, VK_GEOMETRY_INSTANCE_FORCE_NO_OPAQUE_BIT_KHR << 24),
0);
nir_ssa_def *opaque = geom_force_opaque;
opaque = nir_bcsel(b, instance_force_opaque, nir_imm_bool(b, true), opaque);
opaque = nir_bcsel(b, instance_force_non_opaque, nir_imm_bool(b, false), opaque);
nir_ssa_def *ray_force_opaque =
nir_ine(b, nir_iand(b, flags, nir_imm_int(b, SpvRayFlagsOpaqueKHRMask)), nir_imm_int(b, 0));
nir_ine_imm(b, nir_iand_imm(b, flags, SpvRayFlagsOpaqueKHRMask), 0);
nir_ssa_def *ray_force_non_opaque =
nir_ine(b, nir_iand(b, flags, nir_imm_int(b, SpvRayFlagsNoOpaqueKHRMask)), nir_imm_int(b, 0));
nir_ine_imm(b, nir_iand_imm(b, flags, SpvRayFlagsNoOpaqueKHRMask), 0);
opaque = nir_bcsel(b, ray_force_opaque, nir_imm_bool(b, true), opaque);
opaque = nir_bcsel(b, ray_force_non_opaque, nir_imm_bool(b, false), opaque);

18
src/amd/vulkan/radv_shader.c
View file

@ -365,12 +365,12 @@ radv_lower_primitive_shading_rate(nir_shader *nir)
nir_ssa_def *val = nir_ssa_for_src(&b, intr->src[1], 1);
/* x_rate = (shadingRate & (Horizontal2Pixels | Horizontal4Pixels)) ? 0x1 : 0x0; */
nir_ssa_def *x_rate = nir_iand(&b, val, nir_imm_int(&b, 12));
x_rate = nir_b2i32(&b, nir_ine(&b, x_rate, nir_imm_int(&b, 0)));
nir_ssa_def *x_rate = nir_iand_imm(&b, val, 12);
x_rate = nir_b2i32(&b, nir_ine_imm(&b, x_rate, 0));
/* y_rate = (shadingRate & (Vertical2Pixels | Vertical4Pixels)) ? 0x1 : 0x0; */
nir_ssa_def *y_rate = nir_iand(&b, val, nir_imm_int(&b, 3));
y_rate = nir_b2i32(&b, nir_ine(&b, y_rate, nir_imm_int(&b, 0)));
nir_ssa_def *y_rate = nir_iand_imm(&b, val, 3);
y_rate = nir_b2i32(&b, nir_ine_imm(&b, y_rate, 0));
nir_ssa_def *out = NULL;
@ -383,8 +383,7 @@ radv_lower_primitive_shading_rate(nir_shader *nir)
* Bits [30:31] = VRS rate Y
* This will be added to the other bits of that channel in the backend.
*/
out = nir_ior(&b, nir_ishl(&b, x_rate, nir_imm_int(&b, 28)),
nir_ishl(&b, y_rate, nir_imm_int(&b, 30)));
out = nir_ior(&b, nir_ishl_imm(&b, x_rate, 28), nir_ishl_imm(&b, y_rate, 30));
} else {
/* VS, TES, GS:
* Primitive shading rate is a per-vertex output pos export.
@ -393,8 +392,7 @@ radv_lower_primitive_shading_rate(nir_shader *nir)
* Bits [4:5] = VRS rate Y
* HW shading rate = (xRate << 2) | (yRate << 4)
*/
out = nir_ior(&b, nir_ishl(&b, x_rate, nir_imm_int(&b, 2)),
nir_ishl(&b, y_rate, nir_imm_int(&b, 4)));
out = nir_ior(&b, nir_ishl_imm(&b, x_rate, 2), nir_ishl_imm(&b, y_rate, 4));
}
nir_instr_rewrite_src(&intr->instr, &intr->src[1], nir_src_for_ssa(out));
@ -531,10 +529,10 @@ radv_lower_fs_intrinsics(nir_shader *nir, const struct radv_pipeline_stage *fs_s
/* VRS Rate X = Ancillary[2:3] */
nir_ssa_def *ancillary =
nir_load_vector_arg_amd(&b, 1, .base = args->ac.ancillary.arg_index);
nir_ssa_def *x_rate = nir_ubfe(&b, ancillary, nir_imm_int(&b, 2), nir_imm_int(&b, 2));
nir_ssa_def *x_rate = nir_ubfe_imm(&b, ancillary, 2, 2);
/* xRate = xRate == 0x1 ? adjusted_frag_z : frag_z. */
nir_ssa_def *cond = nir_ieq(&b, x_rate, nir_imm_int(&b, 1));
nir_ssa_def *cond = nir_ieq_imm(&b, x_rate, 1);
frag_z = nir_bcsel(&b, cond, adjusted_frag_z, frag_z);
nir_ssa_def *new_dest = nir_vector_insert_imm(&b, &intrin->dest.ssa, frag_z, 2);