mirror of
https://gitlab.freedesktop.org/mesa/mesa.git
synced 2026-06-15 19:58:31 +02:00
vulkan/runtime: add a separate function to build ETC2 decode core shader
This will allow us to re-use it for indirect. Signed-off-by: Samuel Pitoiset <samuel.pitoiset@gmail.com> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/39908>
This commit is contained in:
Samuel Pitoiset
Marge Bot
parent
0bfad39f15
commit
b6f292e146
1 changed files with 269 additions and 261 deletions
|
|
@ -144,292 +144,300 @@ get_global_ids(nir_builder *b, unsigned num_components)
|
|||
}
|
||||
|
||||
static nir_shader *
|
||||
etc2_build_shader(struct vk_device *dev, const struct nir_shader_compiler_options *nir_options)
|
||||
etc2_build_shader_core(nir_builder *b, nir_def *offset, nir_def *format, nir_def *image_type)
|
||||
{
|
||||
const struct glsl_type *sampler_type_2d = glsl_sampler_type(GLSL_SAMPLER_DIM_2D, false, true, GLSL_TYPE_UINT);
|
||||
const struct glsl_type *sampler_type_3d = glsl_sampler_type(GLSL_SAMPLER_DIM_3D, false, false, GLSL_TYPE_UINT);
|
||||
const struct glsl_type *img_type_2d = glsl_image_type(GLSL_SAMPLER_DIM_2D, true, GLSL_TYPE_FLOAT);
|
||||
const struct glsl_type *img_type_3d = glsl_image_type(GLSL_SAMPLER_DIM_3D, false, GLSL_TYPE_FLOAT);
|
||||
nir_builder b = nir_builder_init_simple_shader(MESA_SHADER_COMPUTE, nir_options, "meta_decode_etc");
|
||||
b.shader->info.workgroup_size[0] = 8;
|
||||
b.shader->info.workgroup_size[1] = 8;
|
||||
b->shader->info.workgroup_size[0] = 8;
|
||||
b->shader->info.workgroup_size[1] = 8;
|
||||
|
||||
nir_variable *input_img_2d = nir_variable_create(b.shader, nir_var_uniform, sampler_type_2d, "s_tex_2d");
|
||||
nir_variable *input_img_2d = nir_variable_create(b->shader, nir_var_uniform, sampler_type_2d, "s_tex_2d");
|
||||
input_img_2d->data.descriptor_set = 0;
|
||||
input_img_2d->data.binding = 0;
|
||||
|
||||
nir_variable *input_img_3d = nir_variable_create(b.shader, nir_var_uniform, sampler_type_3d, "s_tex_3d");
|
||||
nir_variable *input_img_3d = nir_variable_create(b->shader, nir_var_uniform, sampler_type_3d, "s_tex_3d");
|
||||
input_img_3d->data.descriptor_set = 0;
|
||||
input_img_3d->data.binding = 0;
|
||||
|
||||
nir_variable *output_img_2d = nir_variable_create(b.shader, nir_var_image, img_type_2d, "out_img_2d");
|
||||
nir_variable *output_img_2d = nir_variable_create(b->shader, nir_var_image, img_type_2d, "out_img_2d");
|
||||
output_img_2d->data.descriptor_set = 0;
|
||||
output_img_2d->data.binding = 1;
|
||||
|
||||
nir_variable *output_img_3d = nir_variable_create(b.shader, nir_var_image, img_type_3d, "out_img_3d");
|
||||
nir_variable *output_img_3d = nir_variable_create(b->shader, nir_var_image, img_type_3d, "out_img_3d");
|
||||
output_img_3d->data.descriptor_set = 0;
|
||||
output_img_3d->data.binding = 1;
|
||||
|
||||
nir_def *global_id = get_global_ids(&b, 3);
|
||||
nir_def *global_id = get_global_ids(b, 3);
|
||||
|
||||
nir_def *is_3d = nir_ieq_imm(b, image_type, VK_IMAGE_TYPE_3D);
|
||||
nir_def *coord = nir_iadd(b, global_id, offset);
|
||||
nir_def *src_coord = nir_vec3(b, nir_ushr_imm(b, nir_channel(b, coord, 0), 2),
|
||||
nir_ushr_imm(b, nir_channel(b, coord, 1), 2), nir_channel(b, coord, 2));
|
||||
|
||||
nir_variable *payload_var = nir_variable_create(b->shader, nir_var_shader_temp, glsl_vec4_type(), "payload");
|
||||
nir_push_if(b, is_3d);
|
||||
{
|
||||
nir_def *color = nir_txf(b, src_coord, .texture_deref = nir_build_deref_var(b, input_img_3d), .lod = nir_imm_int(b, 0));
|
||||
nir_store_var(b, payload_var, color, 0xf);
|
||||
}
|
||||
nir_push_else(b, NULL);
|
||||
{
|
||||
nir_def *color = nir_txf(b, src_coord, .texture_deref = nir_build_deref_var(b, input_img_2d), .lod = nir_imm_int(b, 0));
|
||||
nir_store_var(b, payload_var, color, 0xf);
|
||||
}
|
||||
nir_pop_if(b, NULL);
|
||||
|
||||
nir_def *pixel_coord = nir_iand_imm(b, nir_channels(b, coord, 3), 3);
|
||||
nir_def *linear_pixel =
|
||||
nir_iadd(b, nir_imul_imm(b, nir_channel(b, pixel_coord, 0), 4), nir_channel(b, pixel_coord, 1));
|
||||
|
||||
nir_def *payload = nir_load_var(b, payload_var);
|
||||
nir_variable *color = nir_variable_create(b->shader, nir_var_shader_temp, glsl_vec4_type(), "color");
|
||||
nir_store_var(b, color, nir_imm_vec4(b, 1.0, 0.0, 0.0, 1.0), 0xf);
|
||||
nir_push_if(b, nir_ilt_imm(b, format, VK_FORMAT_EAC_R11_UNORM_BLOCK));
|
||||
{
|
||||
nir_def *alpha_bits_8 = nir_ige_imm(b, format, VK_FORMAT_ETC2_R8G8B8A8_UNORM_BLOCK);
|
||||
nir_def *alpha_bits_1 = nir_iand(b, nir_ige_imm(b, format, VK_FORMAT_ETC2_R8G8B8A1_UNORM_BLOCK),
|
||||
nir_ilt_imm(b, format, VK_FORMAT_ETC2_R8G8B8A8_UNORM_BLOCK));
|
||||
|
||||
nir_def *color_payload =
|
||||
nir_bcsel(b, alpha_bits_8, nir_channels(b, payload, 0xC), nir_channels(b, payload, 3));
|
||||
color_payload = flip_endian(b, color_payload, 2);
|
||||
nir_def *color_y = nir_channel(b, color_payload, 0);
|
||||
nir_def *color_x = nir_channel(b, color_payload, 1);
|
||||
nir_def *flip = nir_test_mask(b, color_y, 1);
|
||||
nir_def *subblock =
|
||||
nir_ushr_imm(b, nir_bcsel(b, flip, nir_channel(b, pixel_coord, 1), nir_channel(b, pixel_coord, 0)), 1);
|
||||
|
||||
nir_variable *punchthrough = nir_variable_create(b->shader, nir_var_shader_temp, glsl_bool_type(), "punchthrough");
|
||||
nir_def *punchthrough_init = nir_iand(b, alpha_bits_1, nir_inot(b, nir_test_mask(b, color_y, 2)));
|
||||
nir_store_var(b, punchthrough, punchthrough_init, 0x1);
|
||||
|
||||
nir_variable *etc1_compat = nir_variable_create(b->shader, nir_var_shader_temp, glsl_bool_type(), "etc1_compat");
|
||||
nir_store_var(b, etc1_compat, nir_imm_false(b), 0x1);
|
||||
|
||||
nir_variable *alpha_result =
|
||||
nir_variable_create(b->shader, nir_var_shader_temp, glsl_float_type(), "alpha_result");
|
||||
nir_push_if(b, alpha_bits_8);
|
||||
{
|
||||
nir_store_var(b, alpha_result, decode_etc2_alpha(b, nir_channels(b, payload, 3), linear_pixel, false, NULL),
|
||||
1);
|
||||
}
|
||||
nir_push_else(b, NULL);
|
||||
{
|
||||
nir_store_var(b, alpha_result, nir_imm_float(b, 1.0), 1);
|
||||
}
|
||||
nir_pop_if(b, NULL);
|
||||
|
||||
const struct glsl_type *uvec3_type = glsl_vector_type(GLSL_TYPE_UINT, 3);
|
||||
nir_variable *rgb_result = nir_variable_create(b->shader, nir_var_shader_temp, uvec3_type, "rgb_result");
|
||||
nir_variable *base_rgb = nir_variable_create(b->shader, nir_var_shader_temp, uvec3_type, "base_rgb");
|
||||
nir_store_var(b, rgb_result, nir_imm_ivec3(b, 255, 0, 0), 0x7);
|
||||
|
||||
nir_def *msb = nir_iand_imm(b, nir_ushr(b, color_x, nir_iadd_imm(b, linear_pixel, 15)), 2);
|
||||
nir_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_inot(b, nir_test_mask(b, color_y, 2))));
|
||||
{
|
||||
nir_store_var(b, etc1_compat, nir_imm_true(b), 1);
|
||||
nir_def *tmp[3];
|
||||
for (unsigned i = 0; i < 3; ++i)
|
||||
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_def *rb = nir_ubfe_imm(b, color_y, 27, 5);
|
||||
nir_def *rd = nir_ibfe_imm(b, color_y, 24, 3);
|
||||
nir_def *gb = nir_ubfe_imm(b, color_y, 19, 5);
|
||||
nir_def *gd = nir_ibfe_imm(b, color_y, 16, 3);
|
||||
nir_def *bb = nir_ubfe_imm(b, color_y, 11, 5);
|
||||
nir_def *bd = nir_ibfe_imm(b, color_y, 8, 3);
|
||||
nir_def *r1 = nir_iadd(b, rb, rd);
|
||||
nir_def *g1 = nir_iadd(b, gb, gd);
|
||||
nir_def *b1 = nir_iadd(b, bb, bd);
|
||||
|
||||
nir_push_if(b, nir_ugt_imm(b, r1, 31));
|
||||
{
|
||||
nir_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_def *g0 = nir_ubfe_imm(b, color_y, 20, 4);
|
||||
nir_def *b0 = nir_ubfe_imm(b, color_y, 16, 4);
|
||||
nir_def *r2 = nir_ubfe_imm(b, color_y, 12, 4);
|
||||
nir_def *g2 = nir_ubfe_imm(b, color_y, 8, 4);
|
||||
nir_def *b2 = nir_ubfe_imm(b, color_y, 4, 4);
|
||||
nir_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_def *dist = etc2_distance_lookup(b, da);
|
||||
nir_def *index = nir_ior(b, lsb, msb);
|
||||
|
||||
nir_store_var(b, punchthrough,
|
||||
nir_iand(b, nir_load_var(b, punchthrough), nir_ieq_imm(b, nir_iadd(b, lsb, msb), 2)),
|
||||
0x1);
|
||||
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_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);
|
||||
}
|
||||
nir_push_else(b, NULL);
|
||||
nir_push_if(b, nir_ugt_imm(b, g1, 31));
|
||||
{
|
||||
nir_def *r0 = nir_ubfe_imm(b, color_y, 27, 4);
|
||||
nir_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_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_def *r2 = nir_ubfe_imm(b, color_y, 11, 4);
|
||||
nir_def *g2 = nir_ubfe_imm(b, color_y, 7, 4);
|
||||
nir_def *b2 = nir_ubfe_imm(b, color_y, 3, 4);
|
||||
nir_def *da = nir_iand_imm(b, color_y, 4);
|
||||
nir_def *db = nir_iand_imm(b, color_y, 1);
|
||||
nir_def *d = nir_iadd(b, da, nir_imul_imm(b, db, 2));
|
||||
nir_def *d0 = nir_iadd(b, nir_ishl_imm(b, r0, 16), nir_iadd(b, nir_ishl_imm(b, g0, 8), b0));
|
||||
nir_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_def *dist = etc2_distance_lookup(b, d);
|
||||
nir_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_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_imm(b, nir_iadd(b, lsb, msb), 2)),
|
||||
0x1);
|
||||
}
|
||||
nir_push_else(b, NULL);
|
||||
nir_push_if(b, nir_ugt_imm(b, b1, 31));
|
||||
{
|
||||
nir_def *r0 = nir_ubfe_imm(b, color_y, 25, 6);
|
||||
nir_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_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_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_def *rv = nir_ubfe_imm(b, color_x, 13, 6);
|
||||
nir_def *gh = nir_ubfe_imm(b, color_x, 25, 7);
|
||||
nir_def *gv = nir_ubfe_imm(b, color_x, 6, 7);
|
||||
nir_def *bh = nir_ubfe_imm(b, color_x, 19, 6);
|
||||
nir_def *bv = nir_ubfe_imm(b, color_x, 0, 6);
|
||||
|
||||
r0 = etc_extend(b, r0, 6);
|
||||
g0 = etc_extend(b, g0, 7);
|
||||
b0 = etc_extend(b, b0, 6);
|
||||
rh = etc_extend(b, rh, 6);
|
||||
rv = etc_extend(b, rv, 6);
|
||||
gh = etc_extend(b, gh, 7);
|
||||
gv = etc_extend(b, gv, 7);
|
||||
bh = etc_extend(b, bh, 6);
|
||||
bv = etc_extend(b, bv, 6);
|
||||
|
||||
nir_def *rgb = nir_vec3(b, r0, g0, b0);
|
||||
nir_def *dx = nir_imul(b, nir_isub(b, nir_vec3(b, rh, gh, bh), rgb), nir_channel(b, pixel_coord, 0));
|
||||
nir_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_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_false(b), 0x1);
|
||||
}
|
||||
nir_push_else(b, NULL);
|
||||
{
|
||||
nir_store_var(b, etc1_compat, nir_imm_true(b), 1);
|
||||
nir_def *subblock_b = nir_ine_imm(b, subblock, 0);
|
||||
nir_def *tmp[] = {
|
||||
nir_bcsel(b, subblock_b, r1, rb),
|
||||
nir_bcsel(b, subblock_b, g1, gb),
|
||||
nir_bcsel(b, subblock_b, b1, bb),
|
||||
};
|
||||
nir_store_var(b, base_rgb, etc_extend(b, nir_vec(b, tmp, 3), 5), 0x7);
|
||||
}
|
||||
nir_pop_if(b, NULL);
|
||||
nir_pop_if(b, NULL);
|
||||
nir_pop_if(b, NULL);
|
||||
}
|
||||
nir_pop_if(b, NULL);
|
||||
nir_push_if(b, nir_load_var(b, etc1_compat));
|
||||
{
|
||||
nir_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_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_imm(b, nir_iadd(b, lsb, msb), 2)), 0x1);
|
||||
nir_def *off = nir_imul(b, etc1_color_modifier_lookup(b, etc1_table_index, lsb), sgn);
|
||||
nir_def *result = nir_iadd(b, nir_load_var(b, base_rgb), off);
|
||||
nir_store_var(b, rgb_result, result, 0x7);
|
||||
}
|
||||
nir_pop_if(b, NULL);
|
||||
nir_push_if(b, nir_load_var(b, punchthrough));
|
||||
{
|
||||
nir_store_var(b, alpha_result, nir_imm_float(b, 0), 0x1);
|
||||
nir_store_var(b, rgb_result, nir_imm_ivec3(b, 0, 0, 0), 0x7);
|
||||
}
|
||||
nir_pop_if(b, NULL);
|
||||
nir_def *col[4];
|
||||
for (unsigned i = 0; i < 3; ++i)
|
||||
col[i] = nir_fdiv_imm(b, nir_i2f32(b, nir_channel(b, nir_load_var(b, rgb_result), i)), 255.0);
|
||||
col[3] = nir_load_var(b, alpha_result);
|
||||
nir_store_var(b, color, nir_vec(b, col, 4), 0xf);
|
||||
}
|
||||
nir_push_else(b, NULL);
|
||||
{ /* EAC */
|
||||
nir_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_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, is_signed);
|
||||
}
|
||||
val[2] = nir_imm_float(b, 0.0);
|
||||
val[3] = nir_imm_float(b, 1.0);
|
||||
nir_store_var(b, color, nir_vec(b, val, 4), 0xf);
|
||||
}
|
||||
nir_pop_if(b, NULL);
|
||||
|
||||
nir_def *outval = nir_load_var(b, color);
|
||||
nir_def *img_coord = nir_vec4(b, nir_channel(b, coord, 0), nir_channel(b, coord, 1), nir_channel(b, coord, 2),
|
||||
nir_undef(b, 1, 32));
|
||||
|
||||
nir_push_if(b, is_3d);
|
||||
{
|
||||
nir_image_deref_store(b, &nir_build_deref_var(b, output_img_3d)->def, img_coord, nir_undef(b, 1, 32), outval,
|
||||
nir_imm_int(b, 0), .image_dim = GLSL_SAMPLER_DIM_3D);
|
||||
}
|
||||
nir_push_else(b, NULL);
|
||||
{
|
||||
nir_image_deref_store(b, &nir_build_deref_var(b, output_img_2d)->def, img_coord, nir_undef(b, 1, 32), outval,
|
||||
nir_imm_int(b, 0), .image_dim = GLSL_SAMPLER_DIM_2D, .image_array = true);
|
||||
}
|
||||
nir_pop_if(b, NULL);
|
||||
return b->shader;
|
||||
}
|
||||
|
||||
static nir_shader *
|
||||
etc2_build_shader(struct vk_device *dev, const struct nir_shader_compiler_options *nir_options)
|
||||
{
|
||||
nir_builder b = nir_builder_init_simple_shader(MESA_SHADER_COMPUTE, nir_options, "meta_decode_etc");
|
||||
|
||||
nir_def *consts = nir_load_push_constant(&b, 4, 32, nir_imm_int(&b, 0), .range = 16);
|
||||
nir_def *consts2 = nir_load_push_constant(&b, 1, 32, nir_imm_int(&b, 0), .base = 16, .range = 4);
|
||||
nir_def *offset = nir_channels(&b, consts, 7);
|
||||
nir_def *format = nir_channel(&b, consts, 3);
|
||||
nir_def *image_type = nir_channel(&b, consts2, 0);
|
||||
nir_def *is_3d = nir_ieq_imm(&b, image_type, VK_IMAGE_TYPE_3D);
|
||||
nir_def *coord = nir_iadd(&b, global_id, offset);
|
||||
nir_def *src_coord = nir_vec3(&b, nir_ushr_imm(&b, nir_channel(&b, coord, 0), 2),
|
||||
nir_ushr_imm(&b, nir_channel(&b, coord, 1), 2), nir_channel(&b, coord, 2));
|
||||
|
||||
nir_variable *payload_var = nir_variable_create(b.shader, nir_var_shader_temp, glsl_vec4_type(), "payload");
|
||||
nir_push_if(&b, is_3d);
|
||||
{
|
||||
nir_def *color = nir_txf(&b, src_coord, .texture_deref = nir_build_deref_var(&b, input_img_3d), .lod = nir_imm_int(&b, 0));
|
||||
nir_store_var(&b, payload_var, color, 0xf);
|
||||
}
|
||||
nir_push_else(&b, NULL);
|
||||
{
|
||||
nir_def *color = nir_txf(&b, src_coord, .texture_deref = nir_build_deref_var(&b, input_img_2d), .lod = nir_imm_int(&b, 0));
|
||||
nir_store_var(&b, payload_var, color, 0xf);
|
||||
}
|
||||
nir_pop_if(&b, NULL);
|
||||
|
||||
nir_def *pixel_coord = nir_iand_imm(&b, nir_channels(&b, coord, 3), 3);
|
||||
nir_def *linear_pixel =
|
||||
nir_iadd(&b, nir_imul_imm(&b, nir_channel(&b, pixel_coord, 0), 4), nir_channel(&b, pixel_coord, 1));
|
||||
|
||||
nir_def *payload = nir_load_var(&b, payload_var);
|
||||
nir_variable *color = nir_variable_create(b.shader, nir_var_shader_temp, glsl_vec4_type(), "color");
|
||||
nir_store_var(&b, color, nir_imm_vec4(&b, 1.0, 0.0, 0.0, 1.0), 0xf);
|
||||
nir_push_if(&b, nir_ilt_imm(&b, format, VK_FORMAT_EAC_R11_UNORM_BLOCK));
|
||||
{
|
||||
nir_def *alpha_bits_8 = nir_ige_imm(&b, format, VK_FORMAT_ETC2_R8G8B8A8_UNORM_BLOCK);
|
||||
nir_def *alpha_bits_1 = nir_iand(&b, nir_ige_imm(&b, format, VK_FORMAT_ETC2_R8G8B8A1_UNORM_BLOCK),
|
||||
nir_ilt_imm(&b, format, VK_FORMAT_ETC2_R8G8B8A8_UNORM_BLOCK));
|
||||
|
||||
nir_def *color_payload =
|
||||
nir_bcsel(&b, alpha_bits_8, nir_channels(&b, payload, 0xC), nir_channels(&b, payload, 3));
|
||||
color_payload = flip_endian(&b, color_payload, 2);
|
||||
nir_def *color_y = nir_channel(&b, color_payload, 0);
|
||||
nir_def *color_x = nir_channel(&b, color_payload, 1);
|
||||
nir_def *flip = nir_test_mask(&b, color_y, 1);
|
||||
nir_def *subblock =
|
||||
nir_ushr_imm(&b, nir_bcsel(&b, flip, nir_channel(&b, pixel_coord, 1), nir_channel(&b, pixel_coord, 0)), 1);
|
||||
|
||||
nir_variable *punchthrough = nir_variable_create(b.shader, nir_var_shader_temp, glsl_bool_type(), "punchthrough");
|
||||
nir_def *punchthrough_init = nir_iand(&b, alpha_bits_1, nir_inot(&b, nir_test_mask(&b, color_y, 2)));
|
||||
nir_store_var(&b, punchthrough, punchthrough_init, 0x1);
|
||||
|
||||
nir_variable *etc1_compat = nir_variable_create(b.shader, nir_var_shader_temp, glsl_bool_type(), "etc1_compat");
|
||||
nir_store_var(&b, etc1_compat, nir_imm_false(&b), 0x1);
|
||||
|
||||
nir_variable *alpha_result =
|
||||
nir_variable_create(b.shader, nir_var_shader_temp, glsl_float_type(), "alpha_result");
|
||||
nir_push_if(&b, alpha_bits_8);
|
||||
{
|
||||
nir_store_var(&b, alpha_result, decode_etc2_alpha(&b, nir_channels(&b, payload, 3), linear_pixel, false, NULL),
|
||||
1);
|
||||
}
|
||||
nir_push_else(&b, NULL);
|
||||
{
|
||||
nir_store_var(&b, alpha_result, nir_imm_float(&b, 1.0), 1);
|
||||
}
|
||||
nir_pop_if(&b, NULL);
|
||||
|
||||
const struct glsl_type *uvec3_type = glsl_vector_type(GLSL_TYPE_UINT, 3);
|
||||
nir_variable *rgb_result = nir_variable_create(b.shader, nir_var_shader_temp, uvec3_type, "rgb_result");
|
||||
nir_variable *base_rgb = nir_variable_create(b.shader, nir_var_shader_temp, uvec3_type, "base_rgb");
|
||||
nir_store_var(&b, rgb_result, nir_imm_ivec3(&b, 255, 0, 0), 0x7);
|
||||
|
||||
nir_def *msb = nir_iand_imm(&b, nir_ushr(&b, color_x, nir_iadd_imm(&b, linear_pixel, 15)), 2);
|
||||
nir_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_inot(&b, nir_test_mask(&b, color_y, 2))));
|
||||
{
|
||||
nir_store_var(&b, etc1_compat, nir_imm_true(&b), 1);
|
||||
nir_def *tmp[3];
|
||||
for (unsigned i = 0; i < 3; ++i)
|
||||
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_def *rb = nir_ubfe_imm(&b, color_y, 27, 5);
|
||||
nir_def *rd = nir_ibfe_imm(&b, color_y, 24, 3);
|
||||
nir_def *gb = nir_ubfe_imm(&b, color_y, 19, 5);
|
||||
nir_def *gd = nir_ibfe_imm(&b, color_y, 16, 3);
|
||||
nir_def *bb = nir_ubfe_imm(&b, color_y, 11, 5);
|
||||
nir_def *bd = nir_ibfe_imm(&b, color_y, 8, 3);
|
||||
nir_def *r1 = nir_iadd(&b, rb, rd);
|
||||
nir_def *g1 = nir_iadd(&b, gb, gd);
|
||||
nir_def *b1 = nir_iadd(&b, bb, bd);
|
||||
|
||||
nir_push_if(&b, nir_ugt_imm(&b, r1, 31));
|
||||
{
|
||||
nir_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_def *g0 = nir_ubfe_imm(&b, color_y, 20, 4);
|
||||
nir_def *b0 = nir_ubfe_imm(&b, color_y, 16, 4);
|
||||
nir_def *r2 = nir_ubfe_imm(&b, color_y, 12, 4);
|
||||
nir_def *g2 = nir_ubfe_imm(&b, color_y, 8, 4);
|
||||
nir_def *b2 = nir_ubfe_imm(&b, color_y, 4, 4);
|
||||
nir_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_def *dist = etc2_distance_lookup(&b, da);
|
||||
nir_def *index = nir_ior(&b, lsb, msb);
|
||||
|
||||
nir_store_var(&b, punchthrough,
|
||||
nir_iand(&b, nir_load_var(&b, punchthrough), nir_ieq_imm(&b, nir_iadd(&b, lsb, msb), 2)),
|
||||
0x1);
|
||||
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_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);
|
||||
}
|
||||
nir_push_else(&b, NULL);
|
||||
nir_push_if(&b, nir_ugt_imm(&b, g1, 31));
|
||||
{
|
||||
nir_def *r0 = nir_ubfe_imm(&b, color_y, 27, 4);
|
||||
nir_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_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_def *r2 = nir_ubfe_imm(&b, color_y, 11, 4);
|
||||
nir_def *g2 = nir_ubfe_imm(&b, color_y, 7, 4);
|
||||
nir_def *b2 = nir_ubfe_imm(&b, color_y, 3, 4);
|
||||
nir_def *da = nir_iand_imm(&b, color_y, 4);
|
||||
nir_def *db = nir_iand_imm(&b, color_y, 1);
|
||||
nir_def *d = nir_iadd(&b, da, nir_imul_imm(&b, db, 2));
|
||||
nir_def *d0 = nir_iadd(&b, nir_ishl_imm(&b, r0, 16), nir_iadd(&b, nir_ishl_imm(&b, g0, 8), b0));
|
||||
nir_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_def *dist = etc2_distance_lookup(&b, d);
|
||||
nir_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_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_imm(&b, nir_iadd(&b, lsb, msb), 2)),
|
||||
0x1);
|
||||
}
|
||||
nir_push_else(&b, NULL);
|
||||
nir_push_if(&b, nir_ugt_imm(&b, b1, 31));
|
||||
{
|
||||
nir_def *r0 = nir_ubfe_imm(&b, color_y, 25, 6);
|
||||
nir_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_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_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_def *rv = nir_ubfe_imm(&b, color_x, 13, 6);
|
||||
nir_def *gh = nir_ubfe_imm(&b, color_x, 25, 7);
|
||||
nir_def *gv = nir_ubfe_imm(&b, color_x, 6, 7);
|
||||
nir_def *bh = nir_ubfe_imm(&b, color_x, 19, 6);
|
||||
nir_def *bv = nir_ubfe_imm(&b, color_x, 0, 6);
|
||||
|
||||
r0 = etc_extend(&b, r0, 6);
|
||||
g0 = etc_extend(&b, g0, 7);
|
||||
b0 = etc_extend(&b, b0, 6);
|
||||
rh = etc_extend(&b, rh, 6);
|
||||
rv = etc_extend(&b, rv, 6);
|
||||
gh = etc_extend(&b, gh, 7);
|
||||
gv = etc_extend(&b, gv, 7);
|
||||
bh = etc_extend(&b, bh, 6);
|
||||
bv = etc_extend(&b, bv, 6);
|
||||
|
||||
nir_def *rgb = nir_vec3(&b, r0, g0, b0);
|
||||
nir_def *dx = nir_imul(&b, nir_isub(&b, nir_vec3(&b, rh, gh, bh), rgb), nir_channel(&b, pixel_coord, 0));
|
||||
nir_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_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_false(&b), 0x1);
|
||||
}
|
||||
nir_push_else(&b, NULL);
|
||||
{
|
||||
nir_store_var(&b, etc1_compat, nir_imm_true(&b), 1);
|
||||
nir_def *subblock_b = nir_ine_imm(&b, subblock, 0);
|
||||
nir_def *tmp[] = {
|
||||
nir_bcsel(&b, subblock_b, r1, rb),
|
||||
nir_bcsel(&b, subblock_b, g1, gb),
|
||||
nir_bcsel(&b, subblock_b, b1, bb),
|
||||
};
|
||||
nir_store_var(&b, base_rgb, etc_extend(&b, nir_vec(&b, tmp, 3), 5), 0x7);
|
||||
}
|
||||
nir_pop_if(&b, NULL);
|
||||
nir_pop_if(&b, NULL);
|
||||
nir_pop_if(&b, NULL);
|
||||
}
|
||||
nir_pop_if(&b, NULL);
|
||||
nir_push_if(&b, nir_load_var(&b, etc1_compat));
|
||||
{
|
||||
nir_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_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_imm(&b, nir_iadd(&b, lsb, msb), 2)), 0x1);
|
||||
nir_def *off = nir_imul(&b, etc1_color_modifier_lookup(&b, etc1_table_index, lsb), sgn);
|
||||
nir_def *result = nir_iadd(&b, nir_load_var(&b, base_rgb), off);
|
||||
nir_store_var(&b, rgb_result, result, 0x7);
|
||||
}
|
||||
nir_pop_if(&b, NULL);
|
||||
nir_push_if(&b, nir_load_var(&b, punchthrough));
|
||||
{
|
||||
nir_store_var(&b, alpha_result, nir_imm_float(&b, 0), 0x1);
|
||||
nir_store_var(&b, rgb_result, nir_imm_ivec3(&b, 0, 0, 0), 0x7);
|
||||
}
|
||||
nir_pop_if(&b, NULL);
|
||||
nir_def *col[4];
|
||||
for (unsigned i = 0; i < 3; ++i)
|
||||
col[i] = nir_fdiv_imm(&b, nir_i2f32(&b, nir_channel(&b, nir_load_var(&b, rgb_result), i)), 255.0);
|
||||
col[3] = nir_load_var(&b, alpha_result);
|
||||
nir_store_var(&b, color, nir_vec(&b, col, 4), 0xf);
|
||||
}
|
||||
nir_push_else(&b, NULL);
|
||||
{ /* EAC */
|
||||
nir_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_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, is_signed);
|
||||
}
|
||||
val[2] = nir_imm_float(&b, 0.0);
|
||||
val[3] = nir_imm_float(&b, 1.0);
|
||||
nir_store_var(&b, color, nir_vec(&b, val, 4), 0xf);
|
||||
}
|
||||
nir_pop_if(&b, NULL);
|
||||
|
||||
nir_def *outval = nir_load_var(&b, color);
|
||||
nir_def *img_coord = nir_vec4(&b, nir_channel(&b, coord, 0), nir_channel(&b, coord, 1), nir_channel(&b, coord, 2),
|
||||
nir_undef(&b, 1, 32));
|
||||
|
||||
nir_push_if(&b, is_3d);
|
||||
{
|
||||
nir_image_deref_store(&b, &nir_build_deref_var(&b, output_img_3d)->def, img_coord, nir_undef(&b, 1, 32), outval,
|
||||
nir_imm_int(&b, 0), .image_dim = GLSL_SAMPLER_DIM_3D);
|
||||
}
|
||||
nir_push_else(&b, NULL);
|
||||
{
|
||||
nir_image_deref_store(&b, &nir_build_deref_var(&b, output_img_2d)->def, img_coord, nir_undef(&b, 1, 32), outval,
|
||||
nir_imm_int(&b, 0), .image_dim = GLSL_SAMPLER_DIM_2D, .image_array = true);
|
||||
}
|
||||
nir_pop_if(&b, NULL);
|
||||
return b.shader;
|
||||
return etc2_build_shader_core(&b, offset, format, image_type);
|
||||
}
|
||||
|
||||
static VkResult
|
||||
|
|
|
|||
Loading…
Add table
Reference in a new issue