2020-03-03 00:47:11 +00:00
|
|
|
/*
|
|
|
|
* Copyright (C) 2020 Collabora Ltd.
|
|
|
|
*
|
|
|
|
* Permission is hereby granted, free of charge, to any person obtaining a
|
|
|
|
* copy of this software and associated documentation files (the "Software"),
|
|
|
|
* to deal in the Software without restriction, including without limitation
|
|
|
|
* the rights to use, copy, modify, merge, publish, distribute, sublicense,
|
|
|
|
* and/or sell copies of the Software, and to permit persons to whom the
|
|
|
|
* Software is furnished to do so, subject to the following conditions:
|
|
|
|
*
|
|
|
|
* The above copyright notice and this permission notice (including the next
|
|
|
|
* paragraph) shall be included in all copies or substantial portions of the
|
|
|
|
* Software.
|
|
|
|
*
|
|
|
|
* THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
|
|
|
|
* IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
|
|
|
* FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL
|
|
|
|
* THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
|
|
|
|
* LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
|
|
|
* OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
|
|
|
|
* SOFTWARE.
|
|
|
|
*
|
|
|
|
* Authors (Collabora):
|
|
|
|
* Alyssa Rosenzweig <alyssa.rosenzweig@collabora.com>
|
|
|
|
*/
|
|
|
|
|
|
|
|
#include "main/mtypes.h"
|
|
|
|
#include "compiler/glsl/glsl_to_nir.h"
|
|
|
|
#include "compiler/nir_types.h"
|
|
|
|
#include "compiler/nir/nir_builder.h"
|
2020-04-30 08:29:10 +01:00
|
|
|
#include "util/u_debug.h"
|
2020-03-03 00:47:11 +00:00
|
|
|
|
|
|
|
#include "disassemble.h"
|
|
|
|
#include "bifrost_compile.h"
|
|
|
|
#include "compiler.h"
|
2020-03-03 19:27:05 +00:00
|
|
|
#include "bi_quirks.h"
|
2020-12-16 17:50:07 +00:00
|
|
|
#include "bi_builder.h"
|
2020-03-05 15:25:19 +00:00
|
|
|
|
2020-12-17 16:07:30 +00:00
|
|
|
static const struct debug_named_value bifrost_debug_options[] = {
|
2020-04-30 08:29:10 +01:00
|
|
|
{"msgs", BIFROST_DBG_MSGS, "Print debug messages"},
|
|
|
|
{"shaders", BIFROST_DBG_SHADERS, "Dump shaders in NIR and MIR"},
|
2020-11-13 22:34:25 +00:00
|
|
|
{"shaderdb", BIFROST_DBG_SHADERDB, "Print statistics"},
|
2020-12-16 20:02:26 +00:00
|
|
|
{"verbose", BIFROST_DBG_VERBOSE, "Disassemble verbosely"},
|
2020-12-30 19:16:51 +00:00
|
|
|
{"internal", BIFROST_DBG_INTERNAL, "Dump even internal shaders"},
|
2021-01-20 22:36:03 +00:00
|
|
|
{"nosched", BIFROST_DBG_NOSCHED, "Force trivial scheduling"},
|
2020-04-30 08:29:10 +01:00
|
|
|
DEBUG_NAMED_VALUE_END
|
|
|
|
};
|
|
|
|
|
2020-12-17 16:07:30 +00:00
|
|
|
DEBUG_GET_ONCE_FLAGS_OPTION(bifrost_debug, "BIFROST_MESA_DEBUG", bifrost_debug_options, 0)
|
2020-04-30 08:29:10 +01:00
|
|
|
|
2021-01-28 08:41:10 +00:00
|
|
|
/* How many bytes are prefetched by the Bifrost shader core. From the final
|
|
|
|
* clause of the shader, this range must be valid instructions or zero. */
|
|
|
|
#define BIFROST_SHADER_PREFETCH 128
|
2021-01-12 09:21:39 +00:00
|
|
|
|
2020-04-30 08:29:10 +01:00
|
|
|
int bifrost_debug = 0;
|
|
|
|
|
|
|
|
#define DBG(fmt, ...) \
|
|
|
|
do { if (bifrost_debug & BIFROST_DBG_MSGS) \
|
|
|
|
fprintf(stderr, "%s:%d: "fmt, \
|
|
|
|
__FUNCTION__, __LINE__, ##__VA_ARGS__); } while (0)
|
|
|
|
|
2020-03-05 15:25:19 +00:00
|
|
|
static bi_block *emit_cf_list(bi_context *ctx, struct exec_list *list);
|
2020-03-05 22:10:46 +00:00
|
|
|
|
2020-12-12 03:33:09 +00:00
|
|
|
static void
|
|
|
|
bi_emit_jump(bi_builder *b, nir_jump_instr *instr)
|
|
|
|
{
|
2021-02-18 19:58:50 +00:00
|
|
|
bi_instr *branch = bi_jump(b, bi_zero());
|
2020-12-12 03:33:09 +00:00
|
|
|
|
|
|
|
switch (instr->type) {
|
|
|
|
case nir_jump_break:
|
|
|
|
branch->branch_target = b->shader->break_block;
|
|
|
|
break;
|
|
|
|
case nir_jump_continue:
|
|
|
|
branch->branch_target = b->shader->continue_block;
|
|
|
|
break;
|
|
|
|
default:
|
|
|
|
unreachable("Unhandled jump type");
|
|
|
|
}
|
|
|
|
|
|
|
|
pan_block_add_successor(&b->shader->current_block->base, &branch->branch_target->base);
|
|
|
|
b->shader->current_block->base.unconditional_jumps = true;
|
|
|
|
}
|
|
|
|
|
2021-02-16 23:58:22 +00:00
|
|
|
static bi_index
|
|
|
|
bi_varying_src0_for_barycentric(bi_builder *b, nir_intrinsic_instr *intr)
|
|
|
|
{
|
|
|
|
switch (intr->intrinsic) {
|
|
|
|
case nir_intrinsic_load_barycentric_centroid:
|
|
|
|
case nir_intrinsic_load_barycentric_sample:
|
|
|
|
return bi_register(61);
|
|
|
|
|
|
|
|
/* Need to put the sample ID in the top 16-bits */
|
|
|
|
case nir_intrinsic_load_barycentric_at_sample:
|
|
|
|
return bi_mkvec_v2i16(b, bi_half(bi_dontcare(), false),
|
|
|
|
bi_half(bi_src_index(&intr->src[0]), false));
|
|
|
|
|
2021-02-17 00:55:12 +00:00
|
|
|
/* Interpret as 8:8 signed fixed point positions in pixels along X and
|
|
|
|
* Y axes respectively, relative to top-left of pixel. In NIR, (0, 0)
|
|
|
|
* is the center of the pixel so we first fixup and then convert. For
|
|
|
|
* fp16 input:
|
|
|
|
*
|
|
|
|
* f2i16(((x, y) + (0.5, 0.5)) * 2**8) =
|
|
|
|
* f2i16((256 * (x, y)) + (128, 128)) =
|
|
|
|
* V2F16_TO_V2S16(FMA.v2f16((x, y), #256, #128))
|
|
|
|
*
|
|
|
|
* For fp32 input, that lacks enough precision for MSAA 16x, but the
|
|
|
|
* idea is the same. FIXME: still doesn't pass
|
|
|
|
*/
|
|
|
|
case nir_intrinsic_load_barycentric_at_offset: {
|
|
|
|
bi_index offset = bi_src_index(&intr->src[0]);
|
|
|
|
bi_index f16 = bi_null();
|
|
|
|
unsigned sz = nir_src_bit_size(intr->src[0]);
|
|
|
|
|
|
|
|
if (sz == 16) {
|
|
|
|
f16 = bi_fma_v2f16(b, offset, bi_imm_f16(256.0),
|
|
|
|
bi_imm_f16(128.0), BI_ROUND_NONE);
|
|
|
|
} else {
|
|
|
|
assert(sz == 32);
|
|
|
|
bi_index f[2];
|
|
|
|
for (unsigned i = 0; i < 2; ++i) {
|
|
|
|
f[i] = bi_fadd_rscale_f32(b,
|
|
|
|
bi_word(offset, i),
|
|
|
|
bi_imm_f32(0.5), bi_imm_u32(8),
|
|
|
|
BI_ROUND_NONE, BI_SPECIAL_NONE);
|
|
|
|
}
|
|
|
|
|
|
|
|
f16 = bi_v2f32_to_v2f16(b, f[0], f[1], BI_ROUND_NONE);
|
|
|
|
}
|
|
|
|
|
|
|
|
return bi_v2f16_to_v2s16(b, f16, BI_ROUND_RTZ);
|
|
|
|
}
|
|
|
|
|
2021-02-16 23:58:22 +00:00
|
|
|
case nir_intrinsic_load_barycentric_pixel:
|
|
|
|
default:
|
|
|
|
return bi_dontcare();
|
|
|
|
}
|
|
|
|
}
|
|
|
|
|
2020-11-27 15:43:06 +00:00
|
|
|
static enum bi_sample
|
2020-11-05 11:11:54 +00:00
|
|
|
bi_interp_for_intrinsic(nir_intrinsic_op op)
|
|
|
|
{
|
|
|
|
switch (op) {
|
|
|
|
case nir_intrinsic_load_barycentric_centroid:
|
2020-11-27 15:43:06 +00:00
|
|
|
return BI_SAMPLE_CENTROID;
|
2020-11-05 11:11:54 +00:00
|
|
|
case nir_intrinsic_load_barycentric_sample:
|
2021-02-16 23:58:22 +00:00
|
|
|
case nir_intrinsic_load_barycentric_at_sample:
|
2020-11-27 15:43:06 +00:00
|
|
|
return BI_SAMPLE_SAMPLE;
|
2021-02-17 00:55:12 +00:00
|
|
|
case nir_intrinsic_load_barycentric_at_offset:
|
|
|
|
return BI_SAMPLE_EXPLICIT;
|
2020-11-05 11:11:54 +00:00
|
|
|
case nir_intrinsic_load_barycentric_pixel:
|
|
|
|
default:
|
2020-11-27 15:43:06 +00:00
|
|
|
return BI_SAMPLE_CENTER;
|
2020-11-05 11:11:54 +00:00
|
|
|
}
|
|
|
|
}
|
|
|
|
|
2020-12-08 13:39:31 +00:00
|
|
|
/* auto, 64-bit omitted */
|
|
|
|
static enum bi_register_format
|
|
|
|
bi_reg_fmt_for_nir(nir_alu_type T)
|
|
|
|
{
|
|
|
|
switch (T) {
|
|
|
|
case nir_type_float16: return BI_REGISTER_FORMAT_F16;
|
|
|
|
case nir_type_float32: return BI_REGISTER_FORMAT_F32;
|
|
|
|
case nir_type_int16: return BI_REGISTER_FORMAT_S16;
|
|
|
|
case nir_type_uint16: return BI_REGISTER_FORMAT_U16;
|
|
|
|
case nir_type_int32: return BI_REGISTER_FORMAT_S32;
|
|
|
|
case nir_type_uint32: return BI_REGISTER_FORMAT_U32;
|
|
|
|
default: unreachable("Invalid type for register format");
|
|
|
|
}
|
|
|
|
}
|
|
|
|
|
2020-12-09 13:58:57 +00:00
|
|
|
/* Checks if the _IMM variant of an intrinsic can be used, returning in imm the
|
|
|
|
* immediate to be used (which applies even if _IMM can't be used) */
|
|
|
|
|
|
|
|
static bool
|
2021-01-11 19:02:38 +00:00
|
|
|
bi_is_intr_immediate(nir_intrinsic_instr *instr, unsigned *immediate, unsigned max)
|
2020-12-09 13:58:57 +00:00
|
|
|
{
|
|
|
|
nir_src *offset = nir_get_io_offset_src(instr);
|
|
|
|
|
|
|
|
if (!nir_src_is_const(*offset))
|
|
|
|
return false;
|
|
|
|
|
|
|
|
*immediate = nir_intrinsic_base(instr) + nir_src_as_uint(*offset);
|
2021-01-11 19:02:38 +00:00
|
|
|
return (*immediate) < max;
|
2020-12-09 13:58:57 +00:00
|
|
|
}
|
|
|
|
|
2020-12-21 22:35:28 +00:00
|
|
|
static void
|
|
|
|
bi_emit_load_attr(bi_builder *b, nir_intrinsic_instr *instr)
|
|
|
|
{
|
|
|
|
nir_alu_type T = nir_intrinsic_dest_type(instr);
|
|
|
|
enum bi_register_format regfmt = bi_reg_fmt_for_nir(T);
|
|
|
|
nir_src *offset = nir_get_io_offset_src(instr);
|
|
|
|
unsigned imm_index = 0;
|
|
|
|
unsigned base = nir_intrinsic_base(instr);
|
|
|
|
bool constant = nir_src_is_const(*offset);
|
2021-01-11 19:02:38 +00:00
|
|
|
bool immediate = bi_is_intr_immediate(instr, &imm_index, 16);
|
2020-12-21 22:35:28 +00:00
|
|
|
|
|
|
|
if (immediate) {
|
|
|
|
bi_ld_attr_imm_to(b, bi_dest_index(&instr->dest),
|
2021-03-04 23:20:48 +00:00
|
|
|
bi_register(61), bi_register(62),
|
2020-12-21 22:35:28 +00:00
|
|
|
regfmt, instr->num_components - 1, imm_index);
|
|
|
|
} else {
|
|
|
|
bi_index idx = bi_src_index(&instr->src[0]);
|
|
|
|
|
|
|
|
if (constant)
|
|
|
|
idx = bi_imm_u32(imm_index);
|
|
|
|
else if (base != 0)
|
|
|
|
idx = bi_iadd_u32(b, idx, bi_imm_u32(base), false);
|
|
|
|
|
|
|
|
bi_ld_attr_to(b, bi_dest_index(&instr->dest),
|
2021-03-04 23:20:48 +00:00
|
|
|
bi_register(61), bi_register(62),
|
2020-12-21 22:35:28 +00:00
|
|
|
idx, regfmt, instr->num_components - 1);
|
|
|
|
}
|
|
|
|
}
|
|
|
|
|
2020-12-08 13:39:20 +00:00
|
|
|
static void
|
|
|
|
bi_emit_load_vary(bi_builder *b, nir_intrinsic_instr *instr)
|
|
|
|
{
|
|
|
|
enum bi_sample sample = BI_SAMPLE_CENTER;
|
|
|
|
enum bi_update update = BI_UPDATE_STORE;
|
|
|
|
enum bi_register_format regfmt = BI_REGISTER_FORMAT_AUTO;
|
|
|
|
enum bi_vecsize vecsize = instr->num_components - 1;
|
|
|
|
bool smooth = instr->intrinsic == nir_intrinsic_load_interpolated_input;
|
2021-02-16 23:58:22 +00:00
|
|
|
bi_index src0 = bi_null();
|
2020-12-08 13:39:20 +00:00
|
|
|
|
|
|
|
if (smooth) {
|
|
|
|
nir_intrinsic_instr *parent = nir_src_as_intrinsic(instr->src[0]);
|
|
|
|
assert(parent);
|
|
|
|
|
|
|
|
sample = bi_interp_for_intrinsic(parent->intrinsic);
|
2021-02-16 23:58:22 +00:00
|
|
|
src0 = bi_varying_src0_for_barycentric(b, parent);
|
2020-12-31 23:17:50 +00:00
|
|
|
|
|
|
|
unsigned sz = nir_dest_bit_size(instr->dest);
|
|
|
|
assert(sz == 16 || sz == 32);
|
|
|
|
|
|
|
|
regfmt = (sz == 16) ? BI_REGISTER_FORMAT_F16
|
|
|
|
: BI_REGISTER_FORMAT_F32;
|
2020-12-08 13:39:20 +00:00
|
|
|
} else {
|
|
|
|
regfmt = bi_reg_fmt_for_nir(nir_intrinsic_dest_type(instr));
|
|
|
|
}
|
|
|
|
|
|
|
|
nir_src *offset = nir_get_io_offset_src(instr);
|
|
|
|
unsigned imm_index = 0;
|
2021-01-11 19:02:38 +00:00
|
|
|
bool immediate = bi_is_intr_immediate(instr, &imm_index, 20);
|
2020-12-08 13:39:20 +00:00
|
|
|
|
|
|
|
if (immediate && smooth) {
|
|
|
|
bi_ld_var_imm_to(b, bi_dest_index(&instr->dest),
|
|
|
|
src0, regfmt, sample, update, vecsize,
|
|
|
|
imm_index);
|
|
|
|
} else if (immediate && !smooth) {
|
|
|
|
bi_ld_var_flat_imm_to(b, bi_dest_index(&instr->dest),
|
|
|
|
BI_FUNCTION_NONE, regfmt, vecsize, imm_index);
|
|
|
|
} else {
|
|
|
|
bi_index idx = bi_src_index(offset);
|
|
|
|
unsigned base = nir_intrinsic_base(instr);
|
|
|
|
|
|
|
|
if (base != 0)
|
|
|
|
idx = bi_iadd_u32(b, idx, bi_imm_u32(base), false);
|
|
|
|
|
|
|
|
if (smooth) {
|
|
|
|
bi_ld_var_to(b, bi_dest_index(&instr->dest),
|
|
|
|
src0, idx, regfmt, sample, update,
|
|
|
|
vecsize);
|
|
|
|
} else {
|
|
|
|
bi_ld_var_flat_to(b, bi_dest_index(&instr->dest),
|
|
|
|
idx, BI_FUNCTION_NONE, regfmt,
|
|
|
|
vecsize);
|
|
|
|
}
|
|
|
|
}
|
|
|
|
}
|
|
|
|
|
2020-12-31 21:04:38 +00:00
|
|
|
static void
|
|
|
|
bi_make_vec16_to(bi_builder *b, bi_index dst, bi_index *src,
|
|
|
|
unsigned *channel, unsigned count)
|
|
|
|
{
|
|
|
|
for (unsigned i = 0; i < count; i += 2) {
|
|
|
|
bool next = (i + 1) < count;
|
|
|
|
|
|
|
|
unsigned chan = channel ? channel[i] : 0;
|
|
|
|
unsigned nextc = next && channel ? channel[i + 1] : 0;
|
|
|
|
|
|
|
|
bi_index w0 = bi_word(src[i], chan >> 1);
|
|
|
|
bi_index w1 = next ? bi_word(src[i + 1], nextc >> 1) : bi_zero();
|
|
|
|
|
|
|
|
bi_index h0 = bi_half(w0, chan & 1);
|
|
|
|
bi_index h1 = bi_half(w1, nextc & 1);
|
|
|
|
|
|
|
|
bi_index to = bi_word(dst, i >> 1);
|
|
|
|
|
|
|
|
if (bi_is_word_equiv(w0, w1) && (chan & 1) == 0 && ((nextc & 1) == 1))
|
|
|
|
bi_mov_i32_to(b, to, w0);
|
|
|
|
else if (bi_is_word_equiv(w0, w1))
|
|
|
|
bi_swz_v2i16_to(b, to, bi_swz_16(w0, chan & 1, nextc & 1));
|
|
|
|
else
|
|
|
|
bi_mkvec_v2i16_to(b, to, h0, h1);
|
|
|
|
}
|
|
|
|
}
|
|
|
|
|
2020-12-09 19:06:15 +00:00
|
|
|
static void
|
|
|
|
bi_make_vec_to(bi_builder *b, bi_index final_dst,
|
|
|
|
bi_index *src,
|
|
|
|
unsigned *channel,
|
|
|
|
unsigned count,
|
|
|
|
unsigned bitsize)
|
|
|
|
{
|
|
|
|
/* If we reads our own output, we need a temporary move to allow for
|
|
|
|
* swapping. TODO: Could do a bit better for pairwise swaps of 16-bit
|
|
|
|
* vectors */
|
|
|
|
bool reads_self = false;
|
|
|
|
|
|
|
|
for (unsigned i = 0; i < count; ++i)
|
|
|
|
reads_self |= bi_is_equiv(final_dst, src[i]);
|
|
|
|
|
|
|
|
/* SSA can't read itself */
|
|
|
|
assert(!reads_self || final_dst.reg);
|
|
|
|
|
|
|
|
bi_index dst = reads_self ? bi_temp(b->shader) : final_dst;
|
|
|
|
|
|
|
|
if (bitsize == 32) {
|
|
|
|
for (unsigned i = 0; i < count; ++i) {
|
|
|
|
bi_mov_i32_to(b, bi_word(dst, i),
|
|
|
|
bi_word(src[i], channel ? channel[i] : 0));
|
|
|
|
}
|
|
|
|
} else if (bitsize == 16) {
|
2020-12-31 21:04:38 +00:00
|
|
|
bi_make_vec16_to(b, dst, src, channel, count);
|
2021-05-03 17:11:26 +01:00
|
|
|
} else if (bitsize == 8 && count == 1) {
|
|
|
|
bi_swz_v4i8_to(b, dst, bi_byte(
|
|
|
|
bi_word(src[0], channel[0] >> 2),
|
|
|
|
channel[0] & 3));
|
2020-12-09 19:06:15 +00:00
|
|
|
} else {
|
|
|
|
unreachable("8-bit mkvec not yet supported");
|
|
|
|
}
|
|
|
|
|
|
|
|
/* Emit an explicit copy if needed */
|
|
|
|
if (!bi_is_equiv(dst, final_dst)) {
|
|
|
|
unsigned shift = (bitsize == 8) ? 2 : (bitsize == 16) ? 1 : 0;
|
|
|
|
unsigned vec = (1 << shift);
|
|
|
|
|
|
|
|
for (unsigned i = 0; i < count; i += vec) {
|
|
|
|
bi_mov_i32_to(b, bi_word(final_dst, i >> shift),
|
|
|
|
bi_word(dst, i >> shift));
|
|
|
|
}
|
|
|
|
}
|
|
|
|
}
|
|
|
|
|
2021-01-29 23:16:58 +00:00
|
|
|
static bi_instr *
|
|
|
|
bi_load_sysval_to(bi_builder *b, bi_index dest, int sysval,
|
|
|
|
unsigned nr_components, unsigned offset)
|
|
|
|
{
|
2021-03-05 12:20:03 +00:00
|
|
|
unsigned sysval_ubo =
|
|
|
|
MAX2(b->shader->inputs->sysval_ubo, b->shader->nir->info.num_ubos);
|
2021-02-12 16:49:23 +00:00
|
|
|
unsigned uniform =
|
2021-02-13 07:24:03 +00:00
|
|
|
pan_lookup_sysval(b->shader->sysval_to_id,
|
|
|
|
&b->shader->info->sysvals,
|
2021-02-12 16:49:23 +00:00
|
|
|
sysval);
|
2021-01-29 23:16:58 +00:00
|
|
|
unsigned idx = (uniform * 16) + offset;
|
|
|
|
|
|
|
|
return bi_load_to(b, nr_components * 32, dest,
|
|
|
|
bi_imm_u32(idx),
|
2021-03-05 12:20:03 +00:00
|
|
|
bi_imm_u32(sysval_ubo), BI_SEG_UBO);
|
2021-01-29 23:16:58 +00:00
|
|
|
}
|
|
|
|
|
|
|
|
static void
|
|
|
|
bi_load_sysval_nir(bi_builder *b, nir_intrinsic_instr *intr,
|
|
|
|
unsigned nr_components, unsigned offset)
|
|
|
|
{
|
|
|
|
bi_load_sysval_to(b, bi_dest_index(&intr->dest),
|
|
|
|
panfrost_sysval_for_instr(&intr->instr, NULL),
|
|
|
|
nr_components, offset);
|
|
|
|
}
|
|
|
|
|
|
|
|
static bi_index
|
|
|
|
bi_load_sysval(bi_builder *b, int sysval,
|
|
|
|
unsigned nr_components, unsigned offset)
|
|
|
|
{
|
|
|
|
bi_index tmp = bi_temp(b->shader);
|
|
|
|
bi_load_sysval_to(b, tmp, sysval, nr_components, offset);
|
|
|
|
return tmp;
|
|
|
|
}
|
|
|
|
|
2020-12-08 13:39:07 +00:00
|
|
|
static void
|
|
|
|
bi_emit_load_blend_input(bi_builder *b, nir_intrinsic_instr *instr)
|
|
|
|
{
|
|
|
|
ASSERTED nir_io_semantics sem = nir_intrinsic_io_semantics(instr);
|
|
|
|
|
2021-01-12 13:18:10 +00:00
|
|
|
/* Source color is passed through r0-r3, or r4-r7 for the second
|
|
|
|
* source when dual-source blending. TODO: Precolour instead */
|
2020-12-08 13:39:07 +00:00
|
|
|
bi_index srcs[] = {
|
|
|
|
bi_register(0), bi_register(1), bi_register(2), bi_register(3)
|
|
|
|
};
|
2021-01-12 13:18:10 +00:00
|
|
|
bi_index srcs2[] = {
|
|
|
|
bi_register(4), bi_register(5), bi_register(6), bi_register(7)
|
|
|
|
};
|
|
|
|
|
|
|
|
bool second_source = (sem.location == VARYING_SLOT_VAR0);
|
2020-12-08 13:39:07 +00:00
|
|
|
|
2021-01-12 13:18:10 +00:00
|
|
|
bi_make_vec_to(b, bi_dest_index(&instr->dest),
|
|
|
|
second_source ? srcs2 : srcs,
|
|
|
|
NULL, 4, 32);
|
2020-12-08 13:39:07 +00:00
|
|
|
}
|
|
|
|
|
2020-12-09 13:42:26 +00:00
|
|
|
static void
|
|
|
|
bi_emit_blend_op(bi_builder *b, bi_index rgba, nir_alu_type T, unsigned rt)
|
|
|
|
{
|
2021-02-22 19:47:42 +00:00
|
|
|
/* Reads 2 or 4 staging registers to cover the input */
|
|
|
|
unsigned sr_count = (nir_alu_type_get_type_size(T) <= 16) ? 2 : 4;
|
|
|
|
|
2021-02-12 17:11:06 +00:00
|
|
|
if (b->shader->inputs->is_blend) {
|
|
|
|
uint64_t blend_desc = b->shader->inputs->blend.bifrost_blend_desc;
|
|
|
|
|
2020-12-09 13:42:26 +00:00
|
|
|
/* Blend descriptor comes from the compile inputs */
|
|
|
|
/* Put the result in r0 */
|
|
|
|
bi_blend_to(b, bi_register(0), rgba,
|
2021-03-04 23:20:48 +00:00
|
|
|
bi_register(60),
|
2021-02-12 17:11:06 +00:00
|
|
|
bi_imm_u32(blend_desc & 0xffffffff),
|
2021-02-22 19:47:42 +00:00
|
|
|
bi_imm_u32(blend_desc >> 32), sr_count);
|
2020-12-09 13:42:26 +00:00
|
|
|
} else {
|
|
|
|
/* Blend descriptor comes from the FAU RAM. By convention, the
|
|
|
|
* return address is stored in r48 and will be used by the
|
|
|
|
* blend shader to jump back to the fragment shader after */
|
|
|
|
bi_blend_to(b, bi_register(48), rgba,
|
2021-03-04 23:20:48 +00:00
|
|
|
bi_register(60),
|
2020-12-09 13:42:26 +00:00
|
|
|
bi_fau(BIR_FAU_BLEND_0 + rt, false),
|
2021-02-22 19:47:42 +00:00
|
|
|
bi_fau(BIR_FAU_BLEND_0 + rt, true), sr_count);
|
2020-12-09 13:42:26 +00:00
|
|
|
}
|
|
|
|
|
|
|
|
assert(rt < 8);
|
2021-02-13 07:24:03 +00:00
|
|
|
b->shader->info->bifrost.blend[rt].type = T;
|
2020-12-09 13:42:26 +00:00
|
|
|
}
|
|
|
|
|
2021-02-12 02:33:02 +00:00
|
|
|
/* Blend shaders do not need to run ATEST since they are dependent on a
|
|
|
|
* fragment shader that runs it. Blit shaders may not need to run ATEST, since
|
|
|
|
* ATEST is not needed if early-z is forced, alpha-to-coverage is disabled, and
|
|
|
|
* there are no writes to the coverage mask. The latter two are satisfied for
|
|
|
|
* all blit shaders, so we just care about early-z, which blit shaders force
|
|
|
|
* iff they do not write depth or stencil */
|
|
|
|
|
|
|
|
static bool
|
|
|
|
bi_skip_atest(bi_context *ctx, bool emit_zs)
|
|
|
|
{
|
2021-02-12 17:11:06 +00:00
|
|
|
return (ctx->inputs->is_blit && !emit_zs) || ctx->inputs->is_blend;
|
2021-02-12 02:33:02 +00:00
|
|
|
}
|
|
|
|
|
2020-12-09 13:42:41 +00:00
|
|
|
static void
|
|
|
|
bi_emit_fragment_out(bi_builder *b, nir_intrinsic_instr *instr)
|
|
|
|
{
|
|
|
|
bool combined = instr->intrinsic ==
|
|
|
|
nir_intrinsic_store_combined_output_pan;
|
|
|
|
|
|
|
|
unsigned writeout = combined ? nir_intrinsic_component(instr) :
|
|
|
|
PAN_WRITEOUT_C;
|
|
|
|
|
|
|
|
bool emit_blend = writeout & (PAN_WRITEOUT_C);
|
|
|
|
bool emit_zs = writeout & (PAN_WRITEOUT_Z | PAN_WRITEOUT_S);
|
|
|
|
|
|
|
|
const nir_variable *var =
|
|
|
|
nir_find_variable_with_driver_location(b->shader->nir,
|
|
|
|
nir_var_shader_out, nir_intrinsic_base(instr));
|
|
|
|
assert(var);
|
|
|
|
|
2021-01-29 21:59:12 +00:00
|
|
|
unsigned loc = var->data.location;
|
|
|
|
bi_index src0 = bi_src_index(&instr->src[0]);
|
|
|
|
|
|
|
|
/* By ISA convention, the coverage mask is stored in R60. The store
|
|
|
|
* itself will be handled by a subsequent ATEST instruction */
|
|
|
|
if (loc == FRAG_RESULT_SAMPLE_MASK) {
|
|
|
|
bi_index orig = bi_register(60);
|
|
|
|
bi_index msaa = bi_load_sysval(b, PAN_SYSVAL_MULTISAMPLED, 1, 0);
|
|
|
|
bi_index new = bi_lshift_and_i32(b, orig, src0, bi_imm_u8(0));
|
|
|
|
bi_mux_i32_to(b, orig, orig, new, msaa, BI_MUX_INT_ZERO);
|
|
|
|
return;
|
|
|
|
}
|
|
|
|
|
|
|
|
|
2021-01-12 13:18:10 +00:00
|
|
|
/* Dual-source blending is implemented by putting the color in
|
|
|
|
* registers r4-r7. */
|
|
|
|
if (var->data.index) {
|
|
|
|
unsigned count = nir_src_num_components(instr->src[0]);
|
|
|
|
|
|
|
|
for (unsigned i = 0; i < count; ++i)
|
2021-03-04 23:20:48 +00:00
|
|
|
bi_mov_i32_to(b, bi_register(4 + i), bi_word(src0, i));
|
|
|
|
|
2021-04-21 20:01:49 +01:00
|
|
|
b->shader->info->bifrost.blend_src1_type =
|
|
|
|
nir_intrinsic_src_type(instr);
|
|
|
|
|
2021-01-12 13:18:10 +00:00
|
|
|
return;
|
|
|
|
}
|
|
|
|
|
2021-01-06 00:08:15 +00:00
|
|
|
/* Emit ATEST if we have to, note ATEST requires a floating-point alpha
|
|
|
|
* value, but render target #0 might not be floating point. However the
|
|
|
|
* alpha value is only used for alpha-to-coverage, a stage which is
|
|
|
|
* skipped for pure integer framebuffers, so the issue is moot. */
|
2021-02-12 02:33:02 +00:00
|
|
|
|
|
|
|
if (!b->shader->emitted_atest && !bi_skip_atest(b->shader, emit_zs)) {
|
2020-12-09 13:42:41 +00:00
|
|
|
nir_alu_type T = nir_intrinsic_src_type(instr);
|
|
|
|
|
|
|
|
bi_index rgba = bi_src_index(&instr->src[0]);
|
2021-01-06 00:08:15 +00:00
|
|
|
bi_index alpha =
|
|
|
|
(T == nir_type_float16) ? bi_half(bi_word(rgba, 1), true) :
|
|
|
|
(T == nir_type_float32) ? bi_word(rgba, 3) :
|
|
|
|
bi_dontcare();
|
2020-12-09 13:42:41 +00:00
|
|
|
|
2021-02-17 04:38:52 +00:00
|
|
|
/* Don't read out-of-bounds */
|
|
|
|
if (nir_src_num_components(instr->src[0]) < 4)
|
|
|
|
alpha = bi_imm_f32(1.0);
|
|
|
|
|
2021-01-19 00:14:23 +00:00
|
|
|
bi_instr *atest = bi_atest_to(b, bi_register(60),
|
|
|
|
bi_register(60), alpha);
|
2020-12-09 13:42:41 +00:00
|
|
|
b->shader->emitted_atest = true;
|
2021-01-19 00:14:23 +00:00
|
|
|
|
|
|
|
/* Pseudo-source to encode in the tuple */
|
|
|
|
atest->src[2] = bi_fau(BIR_FAU_ATEST_PARAM, false);
|
2020-12-09 13:42:41 +00:00
|
|
|
}
|
|
|
|
|
|
|
|
if (emit_zs) {
|
|
|
|
bi_index z = { 0 }, s = { 0 };
|
|
|
|
|
|
|
|
if (writeout & PAN_WRITEOUT_Z)
|
|
|
|
z = bi_src_index(&instr->src[2]);
|
|
|
|
|
|
|
|
if (writeout & PAN_WRITEOUT_S)
|
|
|
|
s = bi_src_index(&instr->src[3]);
|
|
|
|
|
2021-03-04 23:20:48 +00:00
|
|
|
bi_zs_emit_to(b, bi_register(60), z, s, bi_register(60),
|
2020-12-09 13:42:41 +00:00
|
|
|
writeout & PAN_WRITEOUT_S,
|
|
|
|
writeout & PAN_WRITEOUT_Z);
|
|
|
|
}
|
|
|
|
|
|
|
|
if (emit_blend) {
|
2021-04-21 19:13:41 +01:00
|
|
|
assert(loc >= FRAG_RESULT_DATA0);
|
2020-12-09 13:42:41 +00:00
|
|
|
|
2021-04-21 19:13:41 +01:00
|
|
|
unsigned rt = (loc - FRAG_RESULT_DATA0);
|
2021-01-07 14:28:11 +00:00
|
|
|
bi_index color = bi_src_index(&instr->src[0]);
|
|
|
|
|
|
|
|
/* Explicit copy since BLEND inputs are precoloured to R0-R3,
|
|
|
|
* TODO: maybe schedule around this or implement in RA as a
|
|
|
|
* spill */
|
2021-02-17 02:16:26 +00:00
|
|
|
bool has_mrt = false;
|
|
|
|
|
|
|
|
nir_foreach_shader_out_variable(var, b->shader->nir)
|
|
|
|
has_mrt |= (var->data.location > FRAG_RESULT_DATA0);
|
|
|
|
|
|
|
|
if (has_mrt) {
|
2021-01-07 14:28:11 +00:00
|
|
|
bi_index srcs[4] = { color, color, color, color };
|
|
|
|
unsigned channels[4] = { 0, 1, 2, 3 };
|
|
|
|
color = bi_temp(b->shader);
|
|
|
|
bi_make_vec_to(b, color, srcs, channels,
|
|
|
|
nir_src_num_components(instr->src[0]),
|
|
|
|
nir_alu_type_get_type_size(nir_intrinsic_src_type(instr)));
|
|
|
|
}
|
2020-12-09 13:42:41 +00:00
|
|
|
|
2021-01-07 14:28:11 +00:00
|
|
|
bi_emit_blend_op(b, color, nir_intrinsic_src_type(instr), rt);
|
2020-12-09 13:42:41 +00:00
|
|
|
}
|
|
|
|
|
2021-02-12 17:11:06 +00:00
|
|
|
if (b->shader->inputs->is_blend) {
|
2020-12-09 13:42:41 +00:00
|
|
|
/* Jump back to the fragment shader, return address is stored
|
|
|
|
* in r48 (see above).
|
|
|
|
*/
|
2021-02-18 19:58:50 +00:00
|
|
|
bi_jump(b, bi_register(48));
|
2020-12-09 13:42:41 +00:00
|
|
|
}
|
|
|
|
}
|
|
|
|
|
2020-12-09 14:00:04 +00:00
|
|
|
static void
|
|
|
|
bi_emit_store_vary(bi_builder *b, nir_intrinsic_instr *instr)
|
|
|
|
{
|
|
|
|
nir_alu_type T = nir_intrinsic_src_type(instr);
|
|
|
|
enum bi_register_format regfmt = bi_reg_fmt_for_nir(T);
|
|
|
|
|
|
|
|
unsigned imm_index = 0;
|
2021-01-11 19:02:38 +00:00
|
|
|
bool immediate = bi_is_intr_immediate(instr, &imm_index, 16);
|
2020-12-09 14:00:04 +00:00
|
|
|
|
2021-01-13 10:18:11 +00:00
|
|
|
bi_index address;
|
|
|
|
if (immediate) {
|
|
|
|
address = bi_lea_attr_imm(b,
|
2021-03-04 23:20:48 +00:00
|
|
|
bi_register(61), bi_register(62),
|
2021-01-13 10:18:11 +00:00
|
|
|
regfmt, imm_index);
|
|
|
|
} else {
|
|
|
|
bi_index idx =
|
|
|
|
bi_iadd_u32(b,
|
|
|
|
bi_src_index(nir_get_io_offset_src(instr)),
|
|
|
|
bi_imm_u32(nir_intrinsic_base(instr)),
|
|
|
|
false);
|
|
|
|
address = bi_lea_attr(b,
|
2021-03-04 23:20:48 +00:00
|
|
|
bi_register(61), bi_register(62),
|
2021-01-13 10:18:11 +00:00
|
|
|
idx, regfmt);
|
|
|
|
}
|
2020-12-09 14:00:04 +00:00
|
|
|
|
|
|
|
/* Only look at the total components needed. In effect, we fill in all
|
|
|
|
* the intermediate "holes" in the write mask, since we can't mask off
|
|
|
|
* stores. Since nir_lower_io_to_temporaries ensures each varying is
|
|
|
|
* written at most once, anything that's masked out is undefined, so it
|
|
|
|
* doesn't matter what we write there. So we may as well do the
|
|
|
|
* simplest thing possible. */
|
|
|
|
unsigned nr = util_last_bit(nir_intrinsic_write_mask(instr));
|
|
|
|
assert(nr > 0 && nr <= nir_intrinsic_src_components(instr, 0));
|
|
|
|
|
2021-02-18 19:58:50 +00:00
|
|
|
bi_st_cvt(b, bi_src_index(&instr->src[0]), address,
|
|
|
|
bi_word(address, 1), bi_word(address, 2),
|
2020-12-09 14:00:04 +00:00
|
|
|
regfmt, nr - 1);
|
|
|
|
}
|
|
|
|
|
2020-12-12 02:00:26 +00:00
|
|
|
static void
|
|
|
|
bi_emit_load_ubo(bi_builder *b, nir_intrinsic_instr *instr)
|
|
|
|
{
|
2020-12-22 10:26:21 +00:00
|
|
|
nir_src *offset = nir_get_io_offset_src(instr);
|
|
|
|
|
|
|
|
bool offset_is_const = nir_src_is_const(*offset);
|
|
|
|
bi_index dyn_offset = bi_src_index(offset);
|
2021-01-29 23:06:02 +00:00
|
|
|
uint32_t const_offset = offset_is_const ? nir_src_as_uint(*offset) : 0;
|
2020-12-22 10:26:21 +00:00
|
|
|
bool kernel_input = (instr->intrinsic == nir_intrinsic_load_kernel_input);
|
|
|
|
|
2021-02-11 10:51:10 +00:00
|
|
|
bi_load_to(b, instr->num_components * nir_dest_bit_size(instr->dest),
|
2020-12-12 02:00:26 +00:00
|
|
|
bi_dest_index(&instr->dest), offset_is_const ?
|
|
|
|
bi_imm_u32(const_offset) : dyn_offset,
|
2020-12-22 10:26:21 +00:00
|
|
|
kernel_input ? bi_zero() : bi_src_index(&instr->src[0]),
|
2020-12-12 02:00:26 +00:00
|
|
|
BI_SEG_UBO);
|
|
|
|
}
|
|
|
|
|
2021-01-04 09:05:06 +00:00
|
|
|
static bi_index
|
|
|
|
bi_addr_high(nir_src *src)
|
|
|
|
{
|
|
|
|
return (nir_src_bit_size(*src) == 64) ?
|
|
|
|
bi_word(bi_src_index(src), 1) : bi_zero();
|
|
|
|
}
|
|
|
|
|
|
|
|
static void
|
|
|
|
bi_emit_load(bi_builder *b, nir_intrinsic_instr *instr, enum bi_seg seg)
|
|
|
|
{
|
|
|
|
bi_load_to(b, instr->num_components * nir_dest_bit_size(instr->dest),
|
|
|
|
bi_dest_index(&instr->dest),
|
|
|
|
bi_src_index(&instr->src[0]), bi_addr_high(&instr->src[0]),
|
|
|
|
seg);
|
|
|
|
}
|
|
|
|
|
|
|
|
static void
|
|
|
|
bi_emit_store(bi_builder *b, nir_intrinsic_instr *instr, enum bi_seg seg)
|
|
|
|
{
|
2021-04-03 20:00:59 +01:00
|
|
|
/* Require contiguous masks, gauranteed by nir_lower_wrmasks */
|
|
|
|
assert(nir_intrinsic_write_mask(instr) ==
|
|
|
|
BITFIELD_MASK(instr->num_components));
|
|
|
|
|
2021-02-18 19:58:50 +00:00
|
|
|
bi_store(b, instr->num_components * nir_src_bit_size(instr->src[0]),
|
2021-01-04 09:05:06 +00:00
|
|
|
bi_src_index(&instr->src[0]),
|
|
|
|
bi_src_index(&instr->src[1]), bi_addr_high(&instr->src[1]),
|
|
|
|
seg);
|
|
|
|
}
|
|
|
|
|
2021-02-03 18:49:49 +00:00
|
|
|
/* Exchanges the staging register with memory */
|
|
|
|
|
|
|
|
static void
|
2021-02-16 14:50:39 +00:00
|
|
|
bi_emit_axchg_to(bi_builder *b, bi_index dst, bi_index addr, nir_src *arg, enum bi_seg seg)
|
2021-02-03 18:49:49 +00:00
|
|
|
{
|
|
|
|
assert(seg == BI_SEG_NONE || seg == BI_SEG_WLS);
|
|
|
|
|
2021-02-16 14:50:39 +00:00
|
|
|
unsigned sz = nir_src_bit_size(*arg);
|
2021-02-03 18:49:49 +00:00
|
|
|
assert(sz == 32 || sz == 64);
|
|
|
|
|
2021-02-16 14:50:39 +00:00
|
|
|
bi_index data = bi_src_index(arg);
|
|
|
|
|
2021-02-03 18:49:49 +00:00
|
|
|
bi_index data_words[] = {
|
|
|
|
bi_word(data, 0),
|
|
|
|
bi_word(data, 1),
|
|
|
|
};
|
|
|
|
|
|
|
|
bi_index inout = bi_temp_reg(b->shader);
|
|
|
|
bi_make_vec_to(b, inout, data_words, NULL, sz / 32, 32);
|
|
|
|
|
|
|
|
bi_axchg_to(b, sz, inout, inout,
|
|
|
|
bi_word(addr, 0),
|
|
|
|
(seg == BI_SEG_NONE) ? bi_word(addr, 1) : bi_zero(),
|
|
|
|
seg);
|
|
|
|
|
|
|
|
bi_index inout_words[] = {
|
|
|
|
bi_word(inout, 0),
|
|
|
|
bi_word(inout, 1),
|
|
|
|
};
|
|
|
|
|
2021-02-16 14:50:39 +00:00
|
|
|
bi_make_vec_to(b, dst, inout_words, NULL, sz / 32, 32);
|
2021-02-03 18:49:49 +00:00
|
|
|
}
|
|
|
|
|
2021-02-03 19:09:46 +00:00
|
|
|
/* Exchanges the second staging register with memory if comparison with first
|
|
|
|
* staging register passes */
|
|
|
|
|
|
|
|
static void
|
2021-02-16 14:54:19 +00:00
|
|
|
bi_emit_acmpxchg_to(bi_builder *b, bi_index dst, bi_index addr, nir_src *arg_1, nir_src *arg_2, enum bi_seg seg)
|
2021-02-03 19:09:46 +00:00
|
|
|
{
|
|
|
|
assert(seg == BI_SEG_NONE || seg == BI_SEG_WLS);
|
|
|
|
|
|
|
|
/* hardware is swapped from NIR */
|
2021-02-16 14:54:19 +00:00
|
|
|
bi_index src0 = bi_src_index(arg_2);
|
|
|
|
bi_index src1 = bi_src_index(arg_1);
|
2021-02-03 19:09:46 +00:00
|
|
|
|
2021-02-16 14:54:19 +00:00
|
|
|
unsigned sz = nir_src_bit_size(*arg_1);
|
2021-02-03 19:09:46 +00:00
|
|
|
assert(sz == 32 || sz == 64);
|
|
|
|
|
|
|
|
bi_index data_words[] = {
|
|
|
|
bi_word(src0, 0),
|
|
|
|
sz == 32 ? bi_word(src1, 0) : bi_word(src0, 1),
|
|
|
|
|
|
|
|
/* 64-bit */
|
|
|
|
bi_word(src1, 0),
|
|
|
|
bi_word(src1, 1),
|
|
|
|
};
|
|
|
|
|
|
|
|
bi_index inout = bi_temp_reg(b->shader);
|
|
|
|
bi_make_vec_to(b, inout, data_words, NULL, 2 * (sz / 32), 32);
|
|
|
|
|
|
|
|
bi_acmpxchg_to(b, sz, inout, inout,
|
|
|
|
bi_word(addr, 0),
|
|
|
|
(seg == BI_SEG_NONE) ? bi_word(addr, 1) : bi_zero(),
|
|
|
|
seg);
|
|
|
|
|
|
|
|
bi_index inout_words[] = {
|
|
|
|
bi_word(inout, 0),
|
|
|
|
bi_word(inout, 1),
|
|
|
|
};
|
|
|
|
|
2021-02-16 14:54:19 +00:00
|
|
|
bi_make_vec_to(b, dst, inout_words, NULL, sz / 32, 32);
|
2021-02-03 19:09:46 +00:00
|
|
|
}
|
|
|
|
|
2021-02-16 13:18:40 +00:00
|
|
|
/* Extracts an atomic opcode */
|
|
|
|
|
|
|
|
static enum bi_atom_opc
|
2021-03-16 19:09:12 +00:00
|
|
|
bi_atom_opc_for_nir(nir_intrinsic_op op)
|
2021-02-16 13:18:40 +00:00
|
|
|
{
|
|
|
|
switch (op) {
|
|
|
|
case nir_intrinsic_global_atomic_add:
|
|
|
|
case nir_intrinsic_shared_atomic_add:
|
|
|
|
case nir_intrinsic_image_atomic_add:
|
|
|
|
return BI_ATOM_OPC_AADD;
|
|
|
|
|
|
|
|
case nir_intrinsic_global_atomic_imin:
|
|
|
|
case nir_intrinsic_shared_atomic_imin:
|
|
|
|
case nir_intrinsic_image_atomic_imin:
|
|
|
|
return BI_ATOM_OPC_ASMIN;
|
|
|
|
|
|
|
|
case nir_intrinsic_global_atomic_umin:
|
|
|
|
case nir_intrinsic_shared_atomic_umin:
|
|
|
|
case nir_intrinsic_image_atomic_umin:
|
|
|
|
return BI_ATOM_OPC_AUMIN;
|
|
|
|
|
|
|
|
case nir_intrinsic_global_atomic_imax:
|
|
|
|
case nir_intrinsic_shared_atomic_imax:
|
|
|
|
case nir_intrinsic_image_atomic_imax:
|
|
|
|
return BI_ATOM_OPC_ASMAX;
|
|
|
|
|
|
|
|
case nir_intrinsic_global_atomic_umax:
|
|
|
|
case nir_intrinsic_shared_atomic_umax:
|
|
|
|
case nir_intrinsic_image_atomic_umax:
|
|
|
|
return BI_ATOM_OPC_AUMAX;
|
|
|
|
|
|
|
|
case nir_intrinsic_global_atomic_and:
|
|
|
|
case nir_intrinsic_shared_atomic_and:
|
|
|
|
case nir_intrinsic_image_atomic_and:
|
|
|
|
return BI_ATOM_OPC_AAND;
|
|
|
|
|
|
|
|
case nir_intrinsic_global_atomic_or:
|
|
|
|
case nir_intrinsic_shared_atomic_or:
|
|
|
|
case nir_intrinsic_image_atomic_or:
|
|
|
|
return BI_ATOM_OPC_AOR;
|
|
|
|
|
|
|
|
case nir_intrinsic_global_atomic_xor:
|
|
|
|
case nir_intrinsic_shared_atomic_xor:
|
|
|
|
case nir_intrinsic_image_atomic_xor:
|
|
|
|
return BI_ATOM_OPC_AXOR;
|
|
|
|
|
|
|
|
default:
|
|
|
|
unreachable("Unexpected computational atomic");
|
|
|
|
}
|
|
|
|
}
|
|
|
|
|
2021-02-17 18:51:55 +00:00
|
|
|
/* Optimized unary atomics are available with an implied #1 argument */
|
|
|
|
|
|
|
|
static bool
|
|
|
|
bi_promote_atom_c1(enum bi_atom_opc op, bi_index arg, enum bi_atom_opc *out)
|
|
|
|
{
|
|
|
|
/* Check we have a compatible constant */
|
|
|
|
if (arg.type != BI_INDEX_CONSTANT)
|
|
|
|
return false;
|
|
|
|
|
|
|
|
if (!(arg.value == 1 || (arg.value == -1 && op == BI_ATOM_OPC_AADD)))
|
|
|
|
return false;
|
|
|
|
|
|
|
|
/* Check for a compatible operation */
|
|
|
|
switch (op) {
|
|
|
|
case BI_ATOM_OPC_AADD:
|
|
|
|
*out = (arg.value == 1) ? BI_ATOM_OPC_AINC : BI_ATOM_OPC_ADEC;
|
|
|
|
return true;
|
|
|
|
case BI_ATOM_OPC_ASMAX:
|
|
|
|
*out = BI_ATOM_OPC_ASMAX1;
|
|
|
|
return true;
|
|
|
|
case BI_ATOM_OPC_AUMAX:
|
|
|
|
*out = BI_ATOM_OPC_AUMAX1;
|
|
|
|
return true;
|
|
|
|
case BI_ATOM_OPC_AOR:
|
|
|
|
*out = BI_ATOM_OPC_AOR1;
|
|
|
|
return true;
|
|
|
|
default:
|
|
|
|
return false;
|
|
|
|
}
|
|
|
|
}
|
|
|
|
|
2020-12-26 04:09:07 +00:00
|
|
|
/* Coordinates are 16-bit integers in Bifrost but 32-bit in NIR */
|
|
|
|
|
|
|
|
static bi_index
|
|
|
|
bi_emit_image_coord(bi_builder *b, bi_index coord)
|
|
|
|
{
|
|
|
|
return bi_mkvec_v2i16(b,
|
|
|
|
bi_half(bi_word(coord, 0), false),
|
|
|
|
bi_half(bi_word(coord, 1), false));
|
|
|
|
}
|
|
|
|
|
2021-06-01 23:50:30 +01:00
|
|
|
static bi_index
|
|
|
|
bi_emit_image_index(bi_builder *b, nir_intrinsic_instr *instr)
|
|
|
|
{
|
|
|
|
nir_src src = instr->src[0];
|
|
|
|
bi_index index = bi_src_index(&src);
|
|
|
|
bi_context *ctx = b->shader;
|
|
|
|
|
|
|
|
/* Images come after vertex attributes, so handle an explicit offset */
|
|
|
|
unsigned offset = (ctx->stage == MESA_SHADER_VERTEX) ?
|
|
|
|
util_bitcount64(ctx->nir->info.inputs_read) : 0;
|
|
|
|
|
|
|
|
if (offset == 0)
|
|
|
|
return index;
|
|
|
|
else if (nir_src_is_const(src))
|
|
|
|
return bi_imm_u32(nir_src_as_uint(src) + offset);
|
|
|
|
else
|
|
|
|
return bi_iadd_u32(b, index, bi_imm_u32(offset), false);
|
|
|
|
}
|
|
|
|
|
2020-12-26 04:09:07 +00:00
|
|
|
static void
|
|
|
|
bi_emit_image_load(bi_builder *b, nir_intrinsic_instr *instr)
|
|
|
|
{
|
|
|
|
enum glsl_sampler_dim dim = nir_intrinsic_image_dim(instr);
|
|
|
|
ASSERTED unsigned nr_dim = glsl_get_sampler_dim_coordinate_components(dim);
|
|
|
|
|
|
|
|
bi_index coords = bi_src_index(&instr->src[1]);
|
|
|
|
/* TODO: MSAA */
|
|
|
|
assert(nr_dim != GLSL_SAMPLER_DIM_MS && "MSAA'd images not supported");
|
|
|
|
|
|
|
|
bi_ld_attr_tex_to(b, bi_dest_index(&instr->dest),
|
|
|
|
bi_emit_image_coord(b, coords),
|
|
|
|
bi_emit_image_coord(b, bi_word(coords, 2)),
|
2021-06-01 23:50:30 +01:00
|
|
|
bi_emit_image_index(b, instr),
|
2020-12-26 04:09:07 +00:00
|
|
|
bi_reg_fmt_for_nir(nir_intrinsic_dest_type(instr)),
|
|
|
|
instr->num_components - 1);
|
|
|
|
}
|
|
|
|
|
|
|
|
static bi_index
|
|
|
|
bi_emit_lea_image(bi_builder *b, nir_intrinsic_instr *instr)
|
|
|
|
{
|
|
|
|
enum glsl_sampler_dim dim = nir_intrinsic_image_dim(instr);
|
|
|
|
ASSERTED unsigned nr_dim = glsl_get_sampler_dim_coordinate_components(dim);
|
|
|
|
|
|
|
|
/* TODO: MSAA */
|
|
|
|
assert(nr_dim != GLSL_SAMPLER_DIM_MS && "MSAA'd images not supported");
|
|
|
|
|
|
|
|
enum bi_register_format type = (instr->intrinsic == nir_intrinsic_image_store) ?
|
|
|
|
bi_reg_fmt_for_nir(nir_intrinsic_src_type(instr)) :
|
|
|
|
BI_REGISTER_FORMAT_AUTO;
|
|
|
|
|
|
|
|
bi_index coords = bi_src_index(&instr->src[1]);
|
|
|
|
bi_index xy = bi_emit_image_coord(b, coords);
|
|
|
|
bi_index zw = bi_emit_image_coord(b, bi_word(coords, 2));
|
|
|
|
|
|
|
|
bi_instr *I = bi_lea_attr_tex_to(b, bi_temp(b->shader), xy, zw,
|
2021-06-01 23:50:30 +01:00
|
|
|
bi_emit_image_index(b, instr), type);
|
2020-12-26 04:09:07 +00:00
|
|
|
|
|
|
|
/* LEA_ATTR_TEX defaults to the secondary attribute table, but our ABI
|
|
|
|
* has all images in the primary attribute table */
|
|
|
|
I->table = BI_TABLE_ATTRIBUTE_1;
|
|
|
|
|
|
|
|
return I->dest[0];
|
|
|
|
}
|
|
|
|
|
|
|
|
static void
|
|
|
|
bi_emit_image_store(bi_builder *b, nir_intrinsic_instr *instr)
|
|
|
|
{
|
|
|
|
bi_index addr = bi_emit_lea_image(b, instr);
|
|
|
|
|
2021-02-18 19:58:50 +00:00
|
|
|
bi_st_cvt(b, bi_src_index(&instr->src[3]),
|
2020-12-26 04:09:07 +00:00
|
|
|
addr, bi_word(addr, 1), bi_word(addr, 2),
|
|
|
|
bi_reg_fmt_for_nir(nir_intrinsic_src_type(instr)),
|
|
|
|
instr->num_components - 1);
|
|
|
|
}
|
|
|
|
|
2021-02-16 16:02:24 +00:00
|
|
|
static void
|
|
|
|
bi_emit_atomic_i32_to(bi_builder *b, bi_index dst,
|
2021-03-16 19:09:12 +00:00
|
|
|
bi_index addr, bi_index arg, nir_intrinsic_op intrinsic)
|
2021-02-16 16:02:24 +00:00
|
|
|
{
|
|
|
|
/* ATOM_C.i32 takes a vector with {arg, coalesced}, ATOM_C1.i32 doesn't
|
|
|
|
* take any vector but can still output in RETURN mode */
|
|
|
|
bi_index sr = bi_temp_reg(b->shader);
|
|
|
|
|
|
|
|
enum bi_atom_opc opc = bi_atom_opc_for_nir(intrinsic);
|
|
|
|
enum bi_atom_opc post_opc = opc;
|
|
|
|
|
|
|
|
/* Generate either ATOM_C or ATOM_C1 as required */
|
|
|
|
if (bi_promote_atom_c1(opc, arg, &opc)) {
|
2021-05-03 15:57:40 +01:00
|
|
|
bi_patom_c1_i32_to(b, sr, bi_word(addr, 0),
|
|
|
|
bi_word(addr, 1), opc, 2);
|
2021-02-16 16:02:24 +00:00
|
|
|
} else {
|
|
|
|
bi_mov_i32_to(b, sr, arg);
|
2021-05-03 15:57:40 +01:00
|
|
|
bi_patom_c_i32_to(b, sr, sr, bi_word(addr, 0),
|
|
|
|
bi_word(addr, 1), opc, 2);
|
2021-02-16 16:02:24 +00:00
|
|
|
}
|
|
|
|
|
|
|
|
/* Post-process it */
|
|
|
|
bi_atom_post_i32_to(b, dst, bi_word(sr, 0), bi_word(sr, 1), post_opc);
|
|
|
|
}
|
|
|
|
|
2020-12-12 03:49:54 +00:00
|
|
|
/* gl_FragCoord.xy = u16_to_f32(R59.xy) + 0.5
|
|
|
|
* gl_FragCoord.z = ld_vary(fragz)
|
|
|
|
* gl_FragCoord.w = ld_vary(fragw)
|
|
|
|
*/
|
|
|
|
|
|
|
|
static void
|
|
|
|
bi_emit_load_frag_coord(bi_builder *b, nir_intrinsic_instr *instr)
|
|
|
|
{
|
|
|
|
bi_index src[4] = {};
|
|
|
|
|
|
|
|
for (unsigned i = 0; i < 2; ++i) {
|
|
|
|
src[i] = bi_fadd_f32(b,
|
|
|
|
bi_u16_to_f32(b, bi_half(bi_register(59), i)),
|
|
|
|
bi_imm_f32(0.5f), BI_ROUND_NONE);
|
|
|
|
}
|
|
|
|
|
|
|
|
for (unsigned i = 0; i < 2; ++i) {
|
|
|
|
src[2 + i] = bi_ld_var_special(b, bi_zero(),
|
|
|
|
BI_REGISTER_FORMAT_F32, BI_SAMPLE_CENTER,
|
|
|
|
BI_UPDATE_CLOBBER,
|
|
|
|
(i == 0) ? BI_VARYING_NAME_FRAG_Z :
|
|
|
|
BI_VARYING_NAME_FRAG_W,
|
|
|
|
BI_VECSIZE_NONE);
|
|
|
|
}
|
|
|
|
|
|
|
|
bi_make_vec_to(b, bi_dest_index(&instr->dest), src, NULL, 4, 32);
|
|
|
|
}
|
|
|
|
|
2021-02-16 23:12:30 +00:00
|
|
|
static void
|
|
|
|
bi_emit_ld_tile(bi_builder *b, nir_intrinsic_instr *instr)
|
|
|
|
{
|
|
|
|
unsigned rt = b->shader->inputs->blend.rt;
|
2021-04-21 22:42:54 +01:00
|
|
|
unsigned size = nir_dest_bit_size(instr->dest);
|
2021-02-16 23:12:30 +00:00
|
|
|
|
|
|
|
/* Get the render target */
|
|
|
|
if (!b->shader->inputs->is_blend) {
|
|
|
|
const nir_variable *var =
|
|
|
|
nir_find_variable_with_driver_location(b->shader->nir,
|
|
|
|
nir_var_shader_out, nir_intrinsic_base(instr));
|
|
|
|
unsigned loc = var->data.location;
|
2021-04-21 19:13:41 +01:00
|
|
|
assert(loc >= FRAG_RESULT_DATA0);
|
|
|
|
rt = (loc - FRAG_RESULT_DATA0);
|
2021-02-16 23:12:30 +00:00
|
|
|
}
|
|
|
|
|
|
|
|
/* We want to load the current pixel.
|
|
|
|
* FIXME: The sample to load is currently hardcoded to 0. This should
|
|
|
|
* be addressed for multi-sample FBs.
|
|
|
|
*/
|
|
|
|
struct bifrost_pixel_indices pix = {
|
|
|
|
.y = BIFROST_CURRENT_PIXEL,
|
|
|
|
.rt = rt
|
|
|
|
};
|
|
|
|
|
|
|
|
bi_index desc = b->shader->inputs->is_blend ?
|
|
|
|
bi_imm_u32(b->shader->inputs->blend.bifrost_blend_desc >> 32) :
|
2021-04-21 22:42:54 +01:00
|
|
|
bi_load_sysval(b, PAN_SYSVAL(RT_CONVERSION, rt | (size << 4)), 1, 0);
|
2021-02-16 23:12:30 +00:00
|
|
|
|
|
|
|
uint32_t indices = 0;
|
|
|
|
memcpy(&indices, &pix, sizeof(indices));
|
|
|
|
|
|
|
|
bi_ld_tile_to(b, bi_dest_index(&instr->dest), bi_imm_u32(indices),
|
|
|
|
bi_register(60), desc, (instr->num_components - 1));
|
|
|
|
}
|
|
|
|
|
2020-05-01 19:55:04 +01:00
|
|
|
static void
|
2020-12-16 19:24:18 +00:00
|
|
|
bi_emit_intrinsic(bi_builder *b, nir_intrinsic_instr *instr)
|
2020-05-01 19:55:04 +01:00
|
|
|
{
|
2020-12-16 19:24:18 +00:00
|
|
|
bi_index dst = nir_intrinsic_infos[instr->intrinsic].has_dest ?
|
|
|
|
bi_dest_index(&instr->dest) : bi_null();
|
|
|
|
gl_shader_stage stage = b->shader->stage;
|
2020-05-01 19:55:04 +01:00
|
|
|
|
2020-12-16 19:24:18 +00:00
|
|
|
switch (instr->intrinsic) {
|
|
|
|
case nir_intrinsic_load_barycentric_pixel:
|
|
|
|
case nir_intrinsic_load_barycentric_centroid:
|
|
|
|
case nir_intrinsic_load_barycentric_sample:
|
2021-02-16 23:58:22 +00:00
|
|
|
case nir_intrinsic_load_barycentric_at_sample:
|
2021-02-17 00:55:12 +00:00
|
|
|
case nir_intrinsic_load_barycentric_at_offset:
|
2020-12-16 19:24:18 +00:00
|
|
|
/* handled later via load_vary */
|
|
|
|
break;
|
|
|
|
case nir_intrinsic_load_interpolated_input:
|
|
|
|
case nir_intrinsic_load_input:
|
2021-02-12 17:11:06 +00:00
|
|
|
if (b->shader->inputs->is_blend)
|
2020-12-16 19:24:18 +00:00
|
|
|
bi_emit_load_blend_input(b, instr);
|
|
|
|
else if (stage == MESA_SHADER_FRAGMENT)
|
|
|
|
bi_emit_load_vary(b, instr);
|
|
|
|
else if (stage == MESA_SHADER_VERTEX)
|
|
|
|
bi_emit_load_attr(b, instr);
|
|
|
|
else
|
|
|
|
unreachable("Unsupported shader stage");
|
|
|
|
break;
|
2020-05-01 19:55:04 +01:00
|
|
|
|
2020-12-16 19:24:18 +00:00
|
|
|
case nir_intrinsic_store_output:
|
|
|
|
if (stage == MESA_SHADER_FRAGMENT)
|
|
|
|
bi_emit_fragment_out(b, instr);
|
|
|
|
else if (stage == MESA_SHADER_VERTEX)
|
|
|
|
bi_emit_store_vary(b, instr);
|
|
|
|
else
|
|
|
|
unreachable("Unsupported shader stage");
|
|
|
|
break;
|
2020-05-01 19:55:04 +01:00
|
|
|
|
2020-12-16 19:24:18 +00:00
|
|
|
case nir_intrinsic_store_combined_output_pan:
|
|
|
|
assert(stage == MESA_SHADER_FRAGMENT);
|
|
|
|
bi_emit_fragment_out(b, instr);
|
|
|
|
break;
|
2020-05-01 19:55:04 +01:00
|
|
|
|
2020-12-16 19:24:18 +00:00
|
|
|
case nir_intrinsic_load_ubo:
|
2020-12-22 10:26:21 +00:00
|
|
|
case nir_intrinsic_load_kernel_input:
|
2020-12-16 19:24:18 +00:00
|
|
|
bi_emit_load_ubo(b, instr);
|
|
|
|
break;
|
2020-05-01 19:55:04 +01:00
|
|
|
|
2021-01-04 09:05:06 +00:00
|
|
|
case nir_intrinsic_load_global:
|
|
|
|
case nir_intrinsic_load_global_constant:
|
|
|
|
bi_emit_load(b, instr, BI_SEG_NONE);
|
|
|
|
break;
|
|
|
|
|
|
|
|
case nir_intrinsic_store_global:
|
|
|
|
bi_emit_store(b, instr, BI_SEG_NONE);
|
|
|
|
break;
|
|
|
|
|
|
|
|
case nir_intrinsic_load_scratch:
|
|
|
|
bi_emit_load(b, instr, BI_SEG_TL);
|
|
|
|
break;
|
|
|
|
|
|
|
|
case nir_intrinsic_store_scratch:
|
|
|
|
bi_emit_store(b, instr, BI_SEG_TL);
|
|
|
|
break;
|
|
|
|
|
|
|
|
case nir_intrinsic_load_shared:
|
|
|
|
bi_emit_load(b, instr, BI_SEG_WLS);
|
|
|
|
break;
|
|
|
|
|
|
|
|
case nir_intrinsic_store_shared:
|
|
|
|
bi_emit_store(b, instr, BI_SEG_WLS);
|
|
|
|
break;
|
|
|
|
|
2021-02-16 19:18:10 +00:00
|
|
|
/* Blob doesn't seem to do anything for memory barriers, note +BARRIER
|
|
|
|
* is illegal in fragment shaders */
|
|
|
|
case nir_intrinsic_memory_barrier:
|
|
|
|
case nir_intrinsic_memory_barrier_buffer:
|
|
|
|
case nir_intrinsic_memory_barrier_image:
|
|
|
|
case nir_intrinsic_memory_barrier_shared:
|
|
|
|
case nir_intrinsic_group_memory_barrier:
|
|
|
|
break;
|
|
|
|
|
|
|
|
case nir_intrinsic_control_barrier:
|
|
|
|
assert(b->shader->stage != MESA_SHADER_FRAGMENT);
|
2021-02-18 19:58:50 +00:00
|
|
|
bi_barrier(b);
|
2021-02-16 19:18:10 +00:00
|
|
|
break;
|
|
|
|
|
2021-02-16 16:02:24 +00:00
|
|
|
case nir_intrinsic_shared_atomic_add:
|
|
|
|
case nir_intrinsic_shared_atomic_imin:
|
|
|
|
case nir_intrinsic_shared_atomic_umin:
|
|
|
|
case nir_intrinsic_shared_atomic_imax:
|
|
|
|
case nir_intrinsic_shared_atomic_umax:
|
|
|
|
case nir_intrinsic_shared_atomic_and:
|
|
|
|
case nir_intrinsic_shared_atomic_or:
|
|
|
|
case nir_intrinsic_shared_atomic_xor: {
|
|
|
|
assert(nir_src_bit_size(instr->src[1]) == 32);
|
|
|
|
|
|
|
|
bi_index addr = bi_seg_add_i64(b, bi_src_index(&instr->src[0]),
|
|
|
|
bi_zero(), false, BI_SEG_WLS);
|
|
|
|
|
|
|
|
bi_emit_atomic_i32_to(b, dst, addr, bi_src_index(&instr->src[1]),
|
|
|
|
instr->intrinsic);
|
|
|
|
break;
|
|
|
|
}
|
|
|
|
|
|
|
|
case nir_intrinsic_image_atomic_add:
|
|
|
|
case nir_intrinsic_image_atomic_imin:
|
|
|
|
case nir_intrinsic_image_atomic_umin:
|
|
|
|
case nir_intrinsic_image_atomic_imax:
|
|
|
|
case nir_intrinsic_image_atomic_umax:
|
|
|
|
case nir_intrinsic_image_atomic_and:
|
|
|
|
case nir_intrinsic_image_atomic_or:
|
|
|
|
case nir_intrinsic_image_atomic_xor:
|
|
|
|
assert(nir_src_bit_size(instr->src[3]) == 32);
|
|
|
|
|
|
|
|
bi_emit_atomic_i32_to(b, dst,
|
|
|
|
bi_emit_lea_image(b, instr),
|
|
|
|
bi_src_index(&instr->src[3]),
|
|
|
|
instr->intrinsic);
|
|
|
|
break;
|
|
|
|
|
|
|
|
case nir_intrinsic_global_atomic_add:
|
|
|
|
case nir_intrinsic_global_atomic_imin:
|
|
|
|
case nir_intrinsic_global_atomic_umin:
|
|
|
|
case nir_intrinsic_global_atomic_imax:
|
|
|
|
case nir_intrinsic_global_atomic_umax:
|
|
|
|
case nir_intrinsic_global_atomic_and:
|
|
|
|
case nir_intrinsic_global_atomic_or:
|
|
|
|
case nir_intrinsic_global_atomic_xor:
|
|
|
|
assert(nir_src_bit_size(instr->src[1]) == 32);
|
|
|
|
|
|
|
|
bi_emit_atomic_i32_to(b, dst,
|
|
|
|
bi_src_index(&instr->src[0]),
|
|
|
|
bi_src_index(&instr->src[1]),
|
|
|
|
instr->intrinsic);
|
|
|
|
break;
|
|
|
|
|
2020-12-26 04:09:07 +00:00
|
|
|
case nir_intrinsic_image_load:
|
|
|
|
bi_emit_image_load(b, instr);
|
|
|
|
break;
|
|
|
|
|
|
|
|
case nir_intrinsic_image_store:
|
|
|
|
bi_emit_image_store(b, instr);
|
|
|
|
break;
|
|
|
|
|
2021-02-03 18:49:49 +00:00
|
|
|
case nir_intrinsic_global_atomic_exchange:
|
2021-02-16 14:50:39 +00:00
|
|
|
bi_emit_axchg_to(b, dst, bi_src_index(&instr->src[0]),
|
|
|
|
&instr->src[1], BI_SEG_NONE);
|
|
|
|
break;
|
|
|
|
|
|
|
|
case nir_intrinsic_image_atomic_exchange:
|
|
|
|
bi_emit_axchg_to(b, dst, bi_emit_lea_image(b, instr),
|
|
|
|
&instr->src[3], BI_SEG_NONE);
|
2021-02-03 18:49:49 +00:00
|
|
|
break;
|
|
|
|
|
|
|
|
case nir_intrinsic_shared_atomic_exchange:
|
2021-02-16 14:50:39 +00:00
|
|
|
bi_emit_axchg_to(b, dst, bi_src_index(&instr->src[0]),
|
|
|
|
&instr->src[1], BI_SEG_WLS);
|
2021-02-03 18:49:49 +00:00
|
|
|
break;
|
|
|
|
|
2021-02-03 19:09:46 +00:00
|
|
|
case nir_intrinsic_global_atomic_comp_swap:
|
2021-02-16 14:54:19 +00:00
|
|
|
bi_emit_acmpxchg_to(b, dst, bi_src_index(&instr->src[0]),
|
|
|
|
&instr->src[1], &instr->src[2], BI_SEG_NONE);
|
|
|
|
break;
|
|
|
|
|
|
|
|
case nir_intrinsic_image_atomic_comp_swap:
|
|
|
|
bi_emit_acmpxchg_to(b, dst, bi_emit_lea_image(b, instr),
|
|
|
|
&instr->src[3], &instr->src[4], BI_SEG_NONE);
|
2021-02-03 19:09:46 +00:00
|
|
|
break;
|
|
|
|
|
|
|
|
case nir_intrinsic_shared_atomic_comp_swap:
|
2021-02-16 14:54:19 +00:00
|
|
|
bi_emit_acmpxchg_to(b, dst, bi_src_index(&instr->src[0]),
|
|
|
|
&instr->src[1], &instr->src[2], BI_SEG_WLS);
|
2021-02-03 19:09:46 +00:00
|
|
|
break;
|
|
|
|
|
2020-12-16 19:24:18 +00:00
|
|
|
case nir_intrinsic_load_frag_coord:
|
|
|
|
bi_emit_load_frag_coord(b, instr);
|
|
|
|
break;
|
2020-12-16 18:35:42 +00:00
|
|
|
|
|
|
|
case nir_intrinsic_load_output:
|
|
|
|
bi_emit_ld_tile(b, instr);
|
|
|
|
break;
|
|
|
|
|
|
|
|
case nir_intrinsic_discard_if: {
|
|
|
|
bi_index src = bi_src_index(&instr->src[0]);
|
2021-05-04 00:00:35 +01:00
|
|
|
assert(nir_src_bit_size(instr->src[0]) == 1);
|
|
|
|
bi_discard_f32(b, bi_half(src, false), bi_imm_u16(0), BI_CMPF_NE);
|
2020-12-16 18:35:42 +00:00
|
|
|
break;
|
|
|
|
}
|
|
|
|
|
|
|
|
case nir_intrinsic_discard:
|
2021-02-18 19:58:50 +00:00
|
|
|
bi_discard_f32(b, bi_zero(), bi_zero(), BI_CMPF_EQ);
|
2020-12-16 18:35:42 +00:00
|
|
|
break;
|
|
|
|
|
|
|
|
case nir_intrinsic_load_ssbo_address:
|
2021-01-29 23:16:58 +00:00
|
|
|
bi_load_sysval_nir(b, instr, 2, 0);
|
2021-01-08 19:13:16 +00:00
|
|
|
break;
|
|
|
|
|
2020-12-28 11:47:54 +00:00
|
|
|
case nir_intrinsic_load_work_dim:
|
2021-01-29 23:16:58 +00:00
|
|
|
bi_load_sysval_nir(b, instr, 1, 0);
|
2020-12-16 18:35:42 +00:00
|
|
|
break;
|
|
|
|
|
2021-04-23 00:35:56 +01:00
|
|
|
case nir_intrinsic_load_first_vertex:
|
|
|
|
bi_load_sysval_nir(b, instr, 1, 0);
|
|
|
|
break;
|
|
|
|
|
2021-05-11 10:16:46 +01:00
|
|
|
case nir_intrinsic_load_base_vertex:
|
|
|
|
bi_load_sysval_nir(b, instr, 1, 4);
|
|
|
|
break;
|
|
|
|
|
|
|
|
case nir_intrinsic_load_base_instance:
|
|
|
|
bi_load_sysval_nir(b, instr, 1, 8);
|
|
|
|
break;
|
|
|
|
|
2021-05-11 11:21:52 +01:00
|
|
|
case nir_intrinsic_load_draw_id:
|
|
|
|
bi_load_sysval_nir(b, instr, 1, 0);
|
|
|
|
break;
|
|
|
|
|
2020-12-16 18:35:42 +00:00
|
|
|
case nir_intrinsic_get_ssbo_size:
|
2021-01-29 23:16:58 +00:00
|
|
|
bi_load_sysval_nir(b, instr, 1, 8);
|
2020-12-16 18:35:42 +00:00
|
|
|
break;
|
|
|
|
|
|
|
|
case nir_intrinsic_load_viewport_scale:
|
|
|
|
case nir_intrinsic_load_viewport_offset:
|
2021-06-04 20:04:15 +01:00
|
|
|
case nir_intrinsic_load_num_workgroups:
|
2021-05-27 22:44:54 +01:00
|
|
|
case nir_intrinsic_load_workgroup_size:
|
2021-01-29 23:16:58 +00:00
|
|
|
bi_load_sysval_nir(b, instr, 3, 0);
|
2020-12-16 18:35:42 +00:00
|
|
|
break;
|
2020-12-28 12:05:05 +00:00
|
|
|
|
|
|
|
case nir_intrinsic_image_size:
|
|
|
|
bi_load_sysval_nir(b, instr,
|
|
|
|
nir_dest_num_components(instr->dest), 0);
|
|
|
|
break;
|
|
|
|
|
2021-01-27 21:18:38 +00:00
|
|
|
case nir_intrinsic_load_sample_positions_pan:
|
|
|
|
bi_mov_i32_to(b, bi_word(dst, 0),
|
|
|
|
bi_fau(BIR_FAU_SAMPLE_POS_ARRAY, false));
|
|
|
|
bi_mov_i32_to(b, bi_word(dst, 1),
|
|
|
|
bi_fau(BIR_FAU_SAMPLE_POS_ARRAY, true));
|
|
|
|
break;
|
|
|
|
|
2021-01-27 20:44:00 +00:00
|
|
|
case nir_intrinsic_load_sample_mask_in:
|
|
|
|
/* r61[0:15] contains the coverage bitmap */
|
|
|
|
bi_u16_to_u32_to(b, dst, bi_half(bi_register(61), false));
|
|
|
|
break;
|
|
|
|
|
2021-01-27 19:01:23 +00:00
|
|
|
case nir_intrinsic_load_sample_id: {
|
|
|
|
/* r61[16:23] contains the sampleID, mask it out. Upper bits
|
|
|
|
* seem to read garbage (despite being architecturally defined
|
|
|
|
* as zero), so use a 5-bit mask instead of 8-bits */
|
2020-12-16 18:35:42 +00:00
|
|
|
|
2021-01-27 19:01:23 +00:00
|
|
|
bi_rshift_and_i32_to(b, dst, bi_register(61), bi_imm_u32(0x1f),
|
2020-12-16 18:35:42 +00:00
|
|
|
bi_imm_u8(16));
|
|
|
|
break;
|
|
|
|
}
|
|
|
|
|
|
|
|
case nir_intrinsic_load_front_face:
|
|
|
|
/* r58 == 0 means primitive is front facing */
|
|
|
|
bi_icmp_i32_to(b, dst, bi_register(58), bi_zero(), BI_CMPF_EQ,
|
2021-05-06 23:43:00 +01:00
|
|
|
BI_RESULT_TYPE_M1);
|
2020-12-16 18:35:42 +00:00
|
|
|
break;
|
|
|
|
|
|
|
|
case nir_intrinsic_load_point_coord:
|
|
|
|
bi_ld_var_special_to(b, dst, bi_zero(), BI_REGISTER_FORMAT_F32,
|
|
|
|
BI_SAMPLE_CENTER, BI_UPDATE_CLOBBER,
|
|
|
|
BI_VARYING_NAME_POINT, BI_VECSIZE_V2);
|
|
|
|
break;
|
|
|
|
|
2021-04-23 00:35:56 +01:00
|
|
|
case nir_intrinsic_load_vertex_id_zero_base:
|
2020-12-16 18:35:42 +00:00
|
|
|
bi_mov_i32_to(b, dst, bi_register(61));
|
|
|
|
break;
|
|
|
|
|
|
|
|
case nir_intrinsic_load_instance_id:
|
|
|
|
bi_mov_i32_to(b, dst, bi_register(62));
|
|
|
|
break;
|
|
|
|
|
2021-05-06 21:48:31 +01:00
|
|
|
case nir_intrinsic_load_subgroup_invocation:
|
|
|
|
bi_mov_i32_to(b, dst, bi_fau(BIR_FAU_LANE_ID, false));
|
|
|
|
break;
|
|
|
|
|
2020-12-22 12:28:57 +00:00
|
|
|
case nir_intrinsic_load_local_invocation_id:
|
|
|
|
for (unsigned i = 0; i < 3; ++i)
|
|
|
|
bi_u16_to_u32_to(b, bi_word(dst, i),
|
|
|
|
bi_half(bi_register(55 + i / 2), i % 2));
|
|
|
|
break;
|
|
|
|
|
2021-06-04 20:04:15 +01:00
|
|
|
case nir_intrinsic_load_workgroup_id:
|
2020-12-22 12:28:57 +00:00
|
|
|
for (unsigned i = 0; i < 3; ++i)
|
|
|
|
bi_mov_i32_to(b, bi_word(dst, i), bi_register(57 + i));
|
|
|
|
break;
|
|
|
|
|
|
|
|
case nir_intrinsic_load_global_invocation_id:
|
|
|
|
case nir_intrinsic_load_global_invocation_id_zero_base:
|
|
|
|
for (unsigned i = 0; i < 3; ++i)
|
|
|
|
bi_mov_i32_to(b, bi_word(dst, i), bi_register(60 + i));
|
|
|
|
break;
|
|
|
|
|
2021-02-16 18:48:57 +00:00
|
|
|
case nir_intrinsic_shader_clock:
|
|
|
|
bi_ld_gclk_u64_to(b, dst, BI_SOURCE_CYCLE_COUNTER);
|
|
|
|
break;
|
|
|
|
|
2020-12-16 18:35:42 +00:00
|
|
|
default:
|
2020-12-23 00:27:37 +00:00
|
|
|
fprintf(stderr, "Unhandled intrinsic %s\n", nir_intrinsic_infos[instr->intrinsic].name);
|
|
|
|
assert(0);
|
2020-12-16 18:35:42 +00:00
|
|
|
}
|
|
|
|
}
|
|
|
|
|
2020-12-08 00:15:24 +00:00
|
|
|
static void
|
|
|
|
bi_emit_load_const(bi_builder *b, nir_load_const_instr *instr)
|
|
|
|
{
|
|
|
|
/* Make sure we've been lowered */
|
|
|
|
assert(instr->def.num_components <= (32 / instr->def.bit_size));
|
|
|
|
|
|
|
|
/* Accumulate all the channels of the constant, as if we did an
|
|
|
|
* implicit SEL over them */
|
|
|
|
uint32_t acc = 0;
|
|
|
|
|
|
|
|
for (unsigned i = 0; i < instr->def.num_components; ++i) {
|
|
|
|
unsigned v = nir_const_value_as_uint(instr->value[i], instr->def.bit_size);
|
|
|
|
acc |= (v << (i * instr->def.bit_size));
|
|
|
|
}
|
|
|
|
|
|
|
|
bi_mov_i32_to(b, bi_get_index(instr->def.index, false, 0), bi_imm_u32(acc));
|
|
|
|
}
|
|
|
|
|
2020-12-12 01:30:19 +00:00
|
|
|
static bi_index
|
|
|
|
bi_alu_src_index(nir_alu_src src, unsigned comps)
|
|
|
|
{
|
|
|
|
/* we don't lower modifiers until the backend */
|
|
|
|
assert(!(src.negate || src.abs));
|
|
|
|
|
|
|
|
unsigned bitsize = nir_src_bit_size(src.src);
|
|
|
|
|
2021-05-04 00:00:35 +01:00
|
|
|
/* TODO: Do we need to do something more clever with 1-bit bools? */
|
|
|
|
if (bitsize == 1)
|
|
|
|
bitsize = 16;
|
|
|
|
|
2020-12-12 01:30:19 +00:00
|
|
|
/* the bi_index carries the 32-bit (word) offset separate from the
|
|
|
|
* subword swizzle, first handle the offset */
|
|
|
|
|
|
|
|
unsigned offset = 0;
|
|
|
|
|
|
|
|
assert(bitsize == 8 || bitsize == 16 || bitsize == 32);
|
|
|
|
unsigned subword_shift = (bitsize == 32) ? 0 : (bitsize == 16) ? 1 : 2;
|
|
|
|
|
|
|
|
for (unsigned i = 0; i < comps; ++i) {
|
|
|
|
unsigned new_offset = (src.swizzle[i] >> subword_shift);
|
|
|
|
|
|
|
|
if (i > 0)
|
2021-04-13 18:21:19 +01:00
|
|
|
assert(offset == new_offset && "wrong vectorization");
|
2020-12-12 01:30:19 +00:00
|
|
|
|
|
|
|
offset = new_offset;
|
|
|
|
}
|
|
|
|
|
|
|
|
bi_index idx = bi_word(bi_src_index(&src.src), offset);
|
|
|
|
|
|
|
|
/* Compose the subword swizzle with existing (identity) swizzle */
|
|
|
|
assert(idx.swizzle == BI_SWIZZLE_H01);
|
|
|
|
|
|
|
|
/* Bigger vectors should have been lowered */
|
2021-01-13 01:13:47 +00:00
|
|
|
assert(comps <= (1 << subword_shift));
|
2020-12-12 01:30:19 +00:00
|
|
|
|
|
|
|
if (bitsize == 16) {
|
|
|
|
unsigned c0 = src.swizzle[0] & 1;
|
|
|
|
unsigned c1 = (comps > 1) ? src.swizzle[1] & 1 : c0;
|
|
|
|
idx.swizzle = BI_SWIZZLE_H00 + c1 + (c0 << 1);
|
|
|
|
} else if (bitsize == 8) {
|
2021-01-06 00:21:26 +00:00
|
|
|
/* 8-bit vectors not yet supported */
|
|
|
|
assert(comps == 1 && "8-bit vectors not supported");
|
|
|
|
assert(src.swizzle[0] == 0 && "8-bit vectors not supported");
|
|
|
|
idx.swizzle = BI_SWIZZLE_B0000;
|
2020-12-12 01:30:19 +00:00
|
|
|
}
|
|
|
|
|
|
|
|
return idx;
|
|
|
|
}
|
|
|
|
|
2020-12-12 02:24:59 +00:00
|
|
|
static enum bi_round
|
|
|
|
bi_nir_round(nir_op op)
|
|
|
|
{
|
|
|
|
switch (op) {
|
|
|
|
case nir_op_fround_even: return BI_ROUND_NONE;
|
|
|
|
case nir_op_ftrunc: return BI_ROUND_RTZ;
|
|
|
|
case nir_op_fceil: return BI_ROUND_RTP;
|
|
|
|
case nir_op_ffloor: return BI_ROUND_RTN;
|
|
|
|
default: unreachable("invalid nir round op");
|
|
|
|
}
|
|
|
|
}
|
|
|
|
|
2021-02-06 02:36:44 +00:00
|
|
|
/* Convenience for lowered transcendentals */
|
|
|
|
|
|
|
|
static bi_index
|
|
|
|
bi_fmul_f32(bi_builder *b, bi_index s0, bi_index s1)
|
|
|
|
{
|
|
|
|
return bi_fma_f32(b, s0, s1, bi_imm_f32(-0.0f), BI_ROUND_NONE);
|
|
|
|
}
|
|
|
|
|
2021-01-15 20:38:35 +00:00
|
|
|
/* Approximate with FRCP_APPROX.f32 and apply a single iteration of
|
|
|
|
* Newton-Raphson to improve precision */
|
|
|
|
|
|
|
|
static void
|
|
|
|
bi_lower_frcp_32(bi_builder *b, bi_index dst, bi_index s0)
|
|
|
|
{
|
|
|
|
bi_index x1 = bi_frcp_approx_f32(b, s0);
|
|
|
|
bi_index m = bi_frexpm_f32(b, s0, false, false);
|
|
|
|
bi_index e = bi_frexpe_f32(b, bi_neg(s0), false, false);
|
|
|
|
bi_index t1 = bi_fma_rscale_f32(b, m, bi_neg(x1), bi_imm_f32(1.0),
|
|
|
|
bi_zero(), BI_ROUND_NONE, BI_SPECIAL_N);
|
|
|
|
bi_fma_rscale_f32_to(b, dst, t1, x1, x1, e,
|
|
|
|
BI_ROUND_NONE, BI_SPECIAL_NONE);
|
|
|
|
}
|
|
|
|
|
2021-01-15 20:46:39 +00:00
|
|
|
static void
|
|
|
|
bi_lower_frsq_32(bi_builder *b, bi_index dst, bi_index s0)
|
|
|
|
{
|
|
|
|
bi_index x1 = bi_frsq_approx_f32(b, s0);
|
|
|
|
bi_index m = bi_frexpm_f32(b, s0, false, true);
|
|
|
|
bi_index e = bi_frexpe_f32(b, bi_neg(s0), false, true);
|
|
|
|
bi_index t1 = bi_fmul_f32(b, x1, x1);
|
|
|
|
bi_index t2 = bi_fma_rscale_f32(b, m, bi_neg(t1), bi_imm_f32(1.0),
|
|
|
|
bi_imm_u32(-1), BI_ROUND_NONE, BI_SPECIAL_N);
|
|
|
|
bi_fma_rscale_f32_to(b, dst, t2, x1, x1, e,
|
|
|
|
BI_ROUND_NONE, BI_SPECIAL_N);
|
|
|
|
}
|
|
|
|
|
2021-01-15 21:02:46 +00:00
|
|
|
/* More complex transcendentals, see
|
|
|
|
* https://gitlab.freedesktop.org/panfrost/mali-isa-docs/-/blob/master/Bifrost.adoc
|
|
|
|
* for documentation */
|
|
|
|
|
|
|
|
static void
|
|
|
|
bi_lower_fexp2_32(bi_builder *b, bi_index dst, bi_index s0)
|
|
|
|
{
|
|
|
|
bi_index t1 = bi_temp(b->shader);
|
|
|
|
bi_instr *t1_instr = bi_fadd_f32_to(b, t1,
|
|
|
|
s0, bi_imm_u32(0x49400000), BI_ROUND_NONE);
|
|
|
|
t1_instr->clamp = BI_CLAMP_CLAMP_0_INF;
|
|
|
|
|
|
|
|
bi_index t2 = bi_fadd_f32(b, t1, bi_imm_u32(0xc9400000), BI_ROUND_NONE);
|
|
|
|
|
|
|
|
bi_instr *a2 = bi_fadd_f32_to(b, bi_temp(b->shader),
|
|
|
|
s0, bi_neg(t2), BI_ROUND_NONE);
|
|
|
|
a2->clamp = BI_CLAMP_CLAMP_M1_1;
|
|
|
|
|
|
|
|
bi_index a1t = bi_fexp_table_u4(b, t1, BI_ADJ_NONE);
|
|
|
|
bi_index t3 = bi_isub_u32(b, t1, bi_imm_u32(0x49400000), false);
|
|
|
|
bi_index a1i = bi_arshift_i32(b, t3, bi_null(), bi_imm_u8(4));
|
|
|
|
bi_index p1 = bi_fma_f32(b, a2->dest[0], bi_imm_u32(0x3d635635),
|
|
|
|
bi_imm_u32(0x3e75fffa), BI_ROUND_NONE);
|
|
|
|
bi_index p2 = bi_fma_f32(b, p1, a2->dest[0],
|
|
|
|
bi_imm_u32(0x3f317218), BI_ROUND_NONE);
|
|
|
|
bi_index p3 = bi_fmul_f32(b, a2->dest[0], p2);
|
|
|
|
bi_instr *x = bi_fma_rscale_f32_to(b, bi_temp(b->shader),
|
|
|
|
p3, a1t, a1t, a1i, BI_ROUND_NONE, BI_SPECIAL_NONE);
|
|
|
|
x->clamp = BI_CLAMP_CLAMP_0_INF;
|
|
|
|
|
|
|
|
bi_instr *max = bi_fmax_f32_to(b, dst, x->dest[0], s0);
|
|
|
|
max->sem = BI_SEM_NAN_PROPAGATE;
|
|
|
|
}
|
|
|
|
|
2021-01-15 21:34:41 +00:00
|
|
|
static void
|
|
|
|
bi_lower_flog2_32(bi_builder *b, bi_index dst, bi_index s0)
|
|
|
|
{
|
|
|
|
/* s0 = a1 * 2^e, with a1 in [0.75, 1.5) */
|
|
|
|
bi_index a1 = bi_frexpm_f32(b, s0, true, false);
|
|
|
|
bi_index ei = bi_frexpe_f32(b, s0, true, false);
|
|
|
|
bi_index ef = bi_s32_to_f32(b, ei, BI_ROUND_RTZ);
|
|
|
|
|
|
|
|
/* xt estimates -log(r1), a coarse approximation of log(a1) */
|
|
|
|
bi_index r1 = bi_flog_table_f32(b, s0, BI_MODE_RED, BI_PRECISION_NONE);
|
|
|
|
bi_index xt = bi_flog_table_f32(b, s0, BI_MODE_BASE2, BI_PRECISION_NONE);
|
|
|
|
|
|
|
|
/* log(s0) = log(a1 * 2^e) = e + log(a1) = e + log(a1 * r1) -
|
|
|
|
* log(r1), so let x1 = e - log(r1) ~= e + xt and x2 = log(a1 * r1),
|
|
|
|
* and then log(s0) = x1 + x2 */
|
|
|
|
bi_index x1 = bi_fadd_f32(b, ef, xt, BI_ROUND_NONE);
|
|
|
|
|
|
|
|
/* Since a1 * r1 is close to 1, x2 = log(a1 * r1) may be computed by
|
|
|
|
* polynomial approximation around 1. The series is expressed around
|
|
|
|
* 1, so set y = (a1 * r1) - 1.0 */
|
|
|
|
bi_index y = bi_fma_f32(b, a1, r1, bi_imm_f32(-1.0), BI_ROUND_NONE);
|
|
|
|
|
|
|
|
/* x2 = log_2(1 + y) = log_e(1 + y) * (1/log_e(2)), so approximate
|
|
|
|
* log_e(1 + y) by the Taylor series (lower precision than the blob):
|
|
|
|
* y - y^2/2 + O(y^3) = y(1 - y/2) + O(y^3) */
|
|
|
|
bi_index loge = bi_fmul_f32(b, y,
|
|
|
|
bi_fma_f32(b, y, bi_imm_f32(-0.5), bi_imm_f32(1.0), BI_ROUND_NONE));
|
|
|
|
|
|
|
|
bi_index x2 = bi_fmul_f32(b, loge, bi_imm_f32(1.0 / logf(2.0)));
|
|
|
|
|
|
|
|
/* log(s0) = x1 + x2 */
|
|
|
|
bi_fadd_f32_to(b, dst, x1, x2, BI_ROUND_NONE);
|
|
|
|
}
|
|
|
|
|
2021-03-05 02:19:22 +00:00
|
|
|
/* Bifrost has extremely coarse tables for approximating sin/cos, accessible as
|
|
|
|
* FSIN/COS_TABLE.u6, which multiplies the bottom 6-bits by pi/32 and
|
|
|
|
* calculates the results. We use them to calculate sin/cos via a Taylor
|
|
|
|
* approximation:
|
|
|
|
*
|
|
|
|
* f(x + e) = f(x) + e f'(x) + (e^2)/2 f''(x)
|
|
|
|
* sin(x + e) = sin(x) + e cos(x) - (e^2)/2 sin(x)
|
|
|
|
* cos(x + e) = cos(x) - e sin(x) - (e^2)/2 cos(x)
|
|
|
|
*/
|
|
|
|
|
|
|
|
#define TWO_OVER_PI bi_imm_f32(2.0f / 3.14159f)
|
|
|
|
#define MPI_OVER_TWO bi_imm_f32(-3.14159f / 2.0)
|
|
|
|
#define SINCOS_BIAS bi_imm_u32(0x49400000)
|
|
|
|
|
|
|
|
static void
|
|
|
|
bi_lower_fsincos_32(bi_builder *b, bi_index dst, bi_index s0, bool cos)
|
|
|
|
{
|
|
|
|
/* bottom 6-bits of result times pi/32 approximately s0 mod 2pi */
|
|
|
|
bi_index x_u6 = bi_fma_f32(b, s0, TWO_OVER_PI, SINCOS_BIAS, BI_ROUND_NONE);
|
|
|
|
|
|
|
|
/* Approximate domain error (small) */
|
|
|
|
bi_index e = bi_fma_f32(b, bi_fadd_f32(b, x_u6, bi_neg(SINCOS_BIAS),
|
|
|
|
BI_ROUND_NONE),
|
|
|
|
MPI_OVER_TWO, s0, BI_ROUND_NONE);
|
|
|
|
|
|
|
|
/* Lookup sin(x), cos(x) */
|
|
|
|
bi_index sinx = bi_fsin_table_u6(b, x_u6, false);
|
|
|
|
bi_index cosx = bi_fcos_table_u6(b, x_u6, false);
|
|
|
|
|
|
|
|
/* e^2 / 2 */
|
2021-04-13 19:36:08 +01:00
|
|
|
bi_index e2_over_2 = bi_fma_rscale_f32(b, e, e, bi_negzero(),
|
2021-03-05 02:19:22 +00:00
|
|
|
bi_imm_u32(-1), BI_ROUND_NONE, BI_SPECIAL_NONE);
|
|
|
|
|
|
|
|
/* (-e^2)/2 f''(x) */
|
|
|
|
bi_index quadratic = bi_fma_f32(b, bi_neg(e2_over_2),
|
|
|
|
cos ? cosx : sinx,
|
2021-04-13 19:36:08 +01:00
|
|
|
bi_negzero(), BI_ROUND_NONE);
|
2021-03-05 02:19:22 +00:00
|
|
|
|
|
|
|
/* e f'(x) - (e^2/2) f''(x) */
|
|
|
|
bi_instr *I = bi_fma_f32_to(b, bi_temp(b->shader), e,
|
|
|
|
cos ? bi_neg(sinx) : cosx,
|
|
|
|
quadratic, BI_ROUND_NONE);
|
|
|
|
I->clamp = BI_CLAMP_CLAMP_M1_1;
|
|
|
|
|
|
|
|
/* f(x) + e f'(x) - (e^2/2) f''(x) */
|
|
|
|
bi_fadd_f32_to(b, dst, I->dest[0], cos ? cosx : sinx, BI_ROUND_NONE);
|
|
|
|
}
|
|
|
|
|
2021-05-04 00:00:35 +01:00
|
|
|
static bi_instr *
|
|
|
|
bi_emit_alu_bool(bi_builder *b, unsigned sz, nir_op op,
|
|
|
|
bi_index dst, bi_index s0, bi_index s1, bi_index s2)
|
|
|
|
{
|
2021-05-06 23:43:00 +01:00
|
|
|
/* Handle 1-bit bools as 0/~0 by default and let the optimizer deal
|
|
|
|
* with the bit patterns later. 0/~0 has the nice property of being
|
|
|
|
* independent of replicated vectorization. */
|
2021-05-04 00:00:35 +01:00
|
|
|
if (sz == 1) sz = 16;
|
|
|
|
bi_index f = bi_zero();
|
2021-05-06 23:43:00 +01:00
|
|
|
bi_index t = bi_imm_u16(0xFFFF);
|
2021-05-04 00:00:35 +01:00
|
|
|
|
|
|
|
switch (op) {
|
|
|
|
case nir_op_feq:
|
2021-05-06 23:43:00 +01:00
|
|
|
return bi_fcmp_to(b, sz, dst, s0, s1, BI_CMPF_EQ, BI_RESULT_TYPE_M1);
|
2021-05-04 00:00:35 +01:00
|
|
|
case nir_op_flt:
|
2021-05-06 23:43:00 +01:00
|
|
|
return bi_fcmp_to(b, sz, dst, s0, s1, BI_CMPF_LT, BI_RESULT_TYPE_M1);
|
2021-05-04 00:00:35 +01:00
|
|
|
case nir_op_fge:
|
2021-05-06 23:43:00 +01:00
|
|
|
return bi_fcmp_to(b, sz, dst, s0, s1, BI_CMPF_GE, BI_RESULT_TYPE_M1);
|
2021-05-04 00:00:35 +01:00
|
|
|
case nir_op_fneu:
|
2021-05-06 23:43:00 +01:00
|
|
|
return bi_fcmp_to(b, sz, dst, s0, s1, BI_CMPF_NE, BI_RESULT_TYPE_M1);
|
2021-05-04 00:00:35 +01:00
|
|
|
|
|
|
|
case nir_op_ieq:
|
2021-05-06 23:43:00 +01:00
|
|
|
return bi_icmp_to(b, nir_type_int, sz, dst, s0, s1, BI_CMPF_EQ, BI_RESULT_TYPE_M1);
|
2021-05-04 00:00:35 +01:00
|
|
|
case nir_op_ine:
|
2021-05-06 23:43:00 +01:00
|
|
|
return bi_icmp_to(b, nir_type_int, sz, dst, s0, s1, BI_CMPF_NE, BI_RESULT_TYPE_M1);
|
2021-05-04 00:00:35 +01:00
|
|
|
case nir_op_ilt:
|
2021-05-06 23:43:00 +01:00
|
|
|
return bi_icmp_to(b, nir_type_int, sz, dst, s0, s1, BI_CMPF_LT, BI_RESULT_TYPE_M1);
|
2021-05-04 00:00:35 +01:00
|
|
|
case nir_op_ige:
|
2021-05-06 23:43:00 +01:00
|
|
|
return bi_icmp_to(b, nir_type_int, sz, dst, s0, s1, BI_CMPF_GE, BI_RESULT_TYPE_M1);
|
2021-05-04 00:00:35 +01:00
|
|
|
case nir_op_ult:
|
2021-05-06 23:43:00 +01:00
|
|
|
return bi_icmp_to(b, nir_type_uint, sz, dst, s0, s1, BI_CMPF_LT, BI_RESULT_TYPE_M1);
|
2021-05-04 00:00:35 +01:00
|
|
|
case nir_op_uge:
|
2021-05-06 23:43:00 +01:00
|
|
|
return bi_icmp_to(b, nir_type_uint, sz, dst, s0, s1, BI_CMPF_GE, BI_RESULT_TYPE_M1);
|
2021-05-04 00:00:35 +01:00
|
|
|
|
|
|
|
case nir_op_iand:
|
|
|
|
return bi_lshift_and_to(b, sz, dst, s0, s1, bi_imm_u8(0));
|
|
|
|
case nir_op_ior:
|
|
|
|
return bi_lshift_or_to(b, sz, dst, s0, s1, bi_imm_u8(0));
|
|
|
|
case nir_op_ixor:
|
|
|
|
return bi_lshift_xor_to(b, sz, dst, s0, s1, bi_imm_u8(0));
|
|
|
|
case nir_op_inot:
|
2021-05-06 23:43:00 +01:00
|
|
|
return bi_lshift_or_to(b, sz, dst, bi_zero(), bi_not(s0), bi_imm_u8(0));
|
2021-05-04 00:00:35 +01:00
|
|
|
|
|
|
|
case nir_op_f2b1:
|
|
|
|
return bi_csel_to(b, nir_type_int, sz, dst, s0, f, f, t, BI_CMPF_EQ);
|
|
|
|
case nir_op_i2b1:
|
|
|
|
return bi_csel_to(b, nir_type_int, sz, dst, s0, f, f, t, BI_CMPF_EQ);
|
|
|
|
case nir_op_b2b1:
|
|
|
|
return bi_csel_to(b, nir_type_int, sz, dst, s0, f, f, t, BI_CMPF_EQ);
|
|
|
|
|
|
|
|
case nir_op_bcsel:
|
|
|
|
return bi_csel_to(b, nir_type_int, sz, dst, s0, f, s1, s2, BI_CMPF_NE);
|
|
|
|
|
|
|
|
default:
|
|
|
|
fprintf(stderr, "Unhandled ALU op %s\n", nir_op_infos[op].name);
|
|
|
|
unreachable("Unhandled boolean ALU instruction");
|
|
|
|
}
|
|
|
|
}
|
|
|
|
|
2020-12-16 18:58:12 +00:00
|
|
|
static void
|
|
|
|
bi_emit_alu(bi_builder *b, nir_alu_instr *instr)
|
|
|
|
{
|
|
|
|
bi_index dst = bi_dest_index(&instr->dest.dest);
|
|
|
|
unsigned srcs = nir_op_infos[instr->op].num_inputs;
|
2021-05-03 17:11:54 +01:00
|
|
|
unsigned sz = nir_dest_bit_size(instr->dest.dest);
|
2020-12-16 18:58:12 +00:00
|
|
|
unsigned comps = nir_dest_num_components(instr->dest.dest);
|
2020-12-31 21:03:38 +00:00
|
|
|
unsigned src_sz = srcs > 0 ? nir_src_bit_size(instr->src[0].src) : 0;
|
2021-05-04 00:00:35 +01:00
|
|
|
unsigned src1_sz = srcs > 1 ? nir_src_bit_size(instr->src[1].src) : 0;
|
2021-05-06 22:57:17 +01:00
|
|
|
bool is_bool = (sz == 1);
|
|
|
|
|
|
|
|
/* TODO: Anything else? */
|
|
|
|
if (sz == 1)
|
|
|
|
sz = 16;
|
2020-12-16 18:58:12 +00:00
|
|
|
|
2021-04-14 01:21:03 +01:00
|
|
|
/* Indicate scalarness */
|
2021-05-06 22:57:17 +01:00
|
|
|
if (sz == 16 && comps == 1)
|
2021-04-14 01:21:03 +01:00
|
|
|
dst.swizzle = BI_SWIZZLE_H00;
|
|
|
|
|
2020-12-16 18:58:12 +00:00
|
|
|
if (!instr->dest.dest.is_ssa) {
|
|
|
|
for (unsigned i = 0; i < comps; ++i)
|
|
|
|
assert(instr->dest.write_mask);
|
|
|
|
}
|
|
|
|
|
|
|
|
/* First, match against the various moves in NIR. These are
|
|
|
|
* special-cased because they can operate on vectors even after
|
|
|
|
* lowering ALU to scalar. For Bifrost, bi_alu_src_index assumes the
|
|
|
|
* instruction is no "bigger" than SIMD-within-a-register. These moves
|
|
|
|
* are the exceptions that need to handle swizzles specially. */
|
|
|
|
|
|
|
|
switch (instr->op) {
|
2021-01-03 08:45:56 +00:00
|
|
|
case nir_op_pack_32_2x16:
|
2020-12-16 18:58:12 +00:00
|
|
|
case nir_op_vec2:
|
|
|
|
case nir_op_vec3:
|
|
|
|
case nir_op_vec4: {
|
|
|
|
bi_index unoffset_srcs[4] = {
|
|
|
|
srcs > 0 ? bi_src_index(&instr->src[0].src) : bi_null(),
|
|
|
|
srcs > 1 ? bi_src_index(&instr->src[1].src) : bi_null(),
|
|
|
|
srcs > 2 ? bi_src_index(&instr->src[2].src) : bi_null(),
|
|
|
|
srcs > 3 ? bi_src_index(&instr->src[3].src) : bi_null(),
|
|
|
|
};
|
|
|
|
|
|
|
|
unsigned channels[4] = {
|
|
|
|
instr->src[0].swizzle[0],
|
|
|
|
instr->src[1].swizzle[0],
|
|
|
|
srcs > 2 ? instr->src[2].swizzle[0] : 0,
|
|
|
|
srcs > 3 ? instr->src[3].swizzle[0] : 0,
|
|
|
|
};
|
|
|
|
|
|
|
|
bi_make_vec_to(b, dst, unoffset_srcs, channels, srcs, sz);
|
|
|
|
return;
|
|
|
|
}
|
|
|
|
|
|
|
|
case nir_op_vec8:
|
|
|
|
case nir_op_vec16:
|
|
|
|
unreachable("should've been lowered");
|
|
|
|
|
2021-01-03 08:45:56 +00:00
|
|
|
case nir_op_unpack_32_2x16:
|
2020-12-22 10:25:30 +00:00
|
|
|
case nir_op_unpack_64_2x32_split_x:
|
|
|
|
bi_mov_i32_to(b, dst, bi_src_index(&instr->src[0].src));
|
|
|
|
return;
|
|
|
|
|
|
|
|
case nir_op_unpack_64_2x32_split_y:
|
|
|
|
bi_mov_i32_to(b, dst, bi_word(bi_src_index(&instr->src[0].src), 1));
|
|
|
|
return;
|
|
|
|
|
|
|
|
case nir_op_pack_64_2x32_split:
|
|
|
|
bi_mov_i32_to(b, bi_word(dst, 0), bi_src_index(&instr->src[0].src));
|
|
|
|
bi_mov_i32_to(b, bi_word(dst, 1), bi_src_index(&instr->src[1].src));
|
|
|
|
return;
|
|
|
|
|
|
|
|
case nir_op_pack_64_2x32:
|
|
|
|
bi_mov_i32_to(b, bi_word(dst, 0), bi_word(bi_src_index(&instr->src[0].src), 0));
|
|
|
|
bi_mov_i32_to(b, bi_word(dst, 1), bi_word(bi_src_index(&instr->src[0].src), 1));
|
|
|
|
return;
|
|
|
|
|
2020-12-16 18:58:12 +00:00
|
|
|
case nir_op_mov: {
|
|
|
|
bi_index idx = bi_src_index(&instr->src[0].src);
|
|
|
|
bi_index unoffset_srcs[4] = { idx, idx, idx, idx };
|
|
|
|
|
|
|
|
unsigned channels[4] = {
|
|
|
|
comps > 0 ? instr->src[0].swizzle[0] : 0,
|
|
|
|
comps > 1 ? instr->src[0].swizzle[1] : 0,
|
|
|
|
comps > 2 ? instr->src[0].swizzle[2] : 0,
|
|
|
|
comps > 3 ? instr->src[0].swizzle[3] : 0,
|
|
|
|
};
|
|
|
|
|
2021-05-04 00:00:35 +01:00
|
|
|
if (sz == 1) sz = 16;
|
2020-12-16 18:58:12 +00:00
|
|
|
bi_make_vec_to(b, dst, unoffset_srcs, channels, comps, sz);
|
|
|
|
return;
|
|
|
|
}
|
|
|
|
|
2020-12-31 21:03:38 +00:00
|
|
|
case nir_op_f2f16:
|
|
|
|
assert(src_sz == 32);
|
|
|
|
bi_index idx = bi_src_index(&instr->src[0].src);
|
|
|
|
bi_index s0 = bi_word(idx, instr->src[0].swizzle[0]);
|
|
|
|
bi_index s1 = comps > 1 ?
|
|
|
|
bi_word(idx, instr->src[0].swizzle[1]) : s0;
|
|
|
|
|
|
|
|
bi_v2f32_to_v2f16_to(b, dst, s0, s1, BI_ROUND_NONE);
|
|
|
|
return;
|
|
|
|
|
2021-04-13 18:22:01 +01:00
|
|
|
/* Vectorized downcasts */
|
|
|
|
case nir_op_u2u16:
|
|
|
|
case nir_op_i2i16: {
|
|
|
|
if (!(src_sz == 32 && comps == 2))
|
|
|
|
break;
|
|
|
|
|
|
|
|
bi_index idx = bi_src_index(&instr->src[0].src);
|
|
|
|
bi_index s0 = bi_word(idx, instr->src[0].swizzle[0]);
|
|
|
|
bi_index s1 = bi_word(idx, instr->src[0].swizzle[1]);
|
|
|
|
|
|
|
|
bi_mkvec_v2i16_to(b, dst,
|
|
|
|
bi_half(s0, false), bi_half(s1, false));
|
|
|
|
return;
|
|
|
|
}
|
|
|
|
|
2021-05-03 17:11:54 +01:00
|
|
|
case nir_op_i2i8:
|
|
|
|
case nir_op_u2u8:
|
|
|
|
{
|
|
|
|
/* Acts like an 8-bit swizzle */
|
|
|
|
bi_index idx = bi_src_index(&instr->src[0].src);
|
|
|
|
unsigned factor = src_sz / 8;
|
|
|
|
unsigned chan[4] = { 0 };
|
|
|
|
|
|
|
|
for (unsigned i = 0; i < comps; ++i)
|
|
|
|
chan[i] = instr->src[0].swizzle[i] * factor;
|
|
|
|
|
|
|
|
bi_make_vec_to(b, dst, &idx, chan, comps, 8);
|
|
|
|
return;
|
|
|
|
}
|
|
|
|
|
2020-12-16 18:58:12 +00:00
|
|
|
default:
|
|
|
|
break;
|
|
|
|
}
|
|
|
|
|
|
|
|
bi_index s0 = srcs > 0 ? bi_alu_src_index(instr->src[0], comps) : bi_null();
|
|
|
|
bi_index s1 = srcs > 1 ? bi_alu_src_index(instr->src[1], comps) : bi_null();
|
|
|
|
bi_index s2 = srcs > 2 ? bi_alu_src_index(instr->src[2], comps) : bi_null();
|
|
|
|
|
2021-05-06 22:57:17 +01:00
|
|
|
if (is_bool) {
|
2021-05-04 00:00:35 +01:00
|
|
|
bi_emit_alu_bool(b, src_sz, instr->op, dst, s0, s1, s2);
|
|
|
|
return;
|
|
|
|
}
|
|
|
|
|
2020-12-16 18:58:12 +00:00
|
|
|
switch (instr->op) {
|
|
|
|
case nir_op_ffma:
|
|
|
|
bi_fma_to(b, sz, dst, s0, s1, s2, BI_ROUND_NONE);
|
|
|
|
break;
|
|
|
|
|
|
|
|
case nir_op_fmul:
|
2021-04-13 19:36:08 +01:00
|
|
|
bi_fma_to(b, sz, dst, s0, s1, bi_negzero(), BI_ROUND_NONE);
|
2020-12-16 18:58:12 +00:00
|
|
|
break;
|
|
|
|
|
|
|
|
case nir_op_fsub:
|
|
|
|
s1 = bi_neg(s1);
|
2021-04-10 16:11:58 +01:00
|
|
|
FALLTHROUGH;
|
2020-12-16 18:58:12 +00:00
|
|
|
case nir_op_fadd:
|
|
|
|
bi_fadd_to(b, sz, dst, s0, s1, BI_ROUND_NONE);
|
|
|
|
break;
|
|
|
|
|
|
|
|
case nir_op_fsat: {
|
2021-04-13 19:36:08 +01:00
|
|
|
bi_instr *I = bi_fadd_to(b, sz, dst, s0, bi_negzero(), BI_ROUND_NONE);
|
2020-12-16 18:58:12 +00:00
|
|
|
I->clamp = BI_CLAMP_CLAMP_0_1;
|
|
|
|
break;
|
|
|
|
}
|
|
|
|
|
|
|
|
case nir_op_fneg:
|
2021-04-13 19:36:08 +01:00
|
|
|
bi_fadd_to(b, sz, dst, bi_neg(s0), bi_negzero(), BI_ROUND_NONE);
|
2020-12-16 18:58:12 +00:00
|
|
|
break;
|
|
|
|
|
|
|
|
case nir_op_fabs:
|
2021-04-13 19:36:08 +01:00
|
|
|
bi_fadd_to(b, sz, dst, bi_abs(s0), bi_negzero(), BI_ROUND_NONE);
|
2020-12-16 18:58:12 +00:00
|
|
|
break;
|
|
|
|
|
2021-03-05 02:19:22 +00:00
|
|
|
case nir_op_fsin:
|
|
|
|
bi_lower_fsincos_32(b, dst, s0, false);
|
|
|
|
break;
|
|
|
|
|
|
|
|
case nir_op_fcos:
|
|
|
|
bi_lower_fsincos_32(b, dst, s0, true);
|
|
|
|
break;
|
|
|
|
|
2020-12-16 18:58:12 +00:00
|
|
|
case nir_op_fexp2: {
|
|
|
|
assert(sz == 32); /* should've been lowered */
|
|
|
|
|
2021-01-15 21:39:58 +00:00
|
|
|
if (b->shader->quirks & BIFROST_NO_FP32_TRANSCENDENTALS) {
|
|
|
|
bi_lower_fexp2_32(b, dst, s0);
|
|
|
|
break;
|
|
|
|
}
|
|
|
|
|
2020-12-16 18:58:12 +00:00
|
|
|
/* multiply by 1.0 * 2*24 */
|
|
|
|
bi_index scale = bi_fma_rscale_f32(b, s0, bi_imm_f32(1.0f),
|
2021-04-13 19:36:08 +01:00
|
|
|
bi_negzero(), bi_imm_u32(24), BI_ROUND_NONE,
|
2020-12-16 18:58:12 +00:00
|
|
|
BI_SPECIAL_NONE);
|
|
|
|
|
|
|
|
bi_fexp_f32_to(b, dst, bi_f32_to_s32(b, scale, BI_ROUND_NONE), s0);
|
|
|
|
break;
|
|
|
|
}
|
|
|
|
|
|
|
|
case nir_op_flog2: {
|
|
|
|
assert(sz == 32); /* should've been lowered */
|
2021-01-15 21:39:58 +00:00
|
|
|
|
|
|
|
if (b->shader->quirks & BIFROST_NO_FP32_TRANSCENDENTALS) {
|
|
|
|
bi_lower_flog2_32(b, dst, s0);
|
|
|
|
break;
|
|
|
|
}
|
|
|
|
|
2020-12-16 18:58:12 +00:00
|
|
|
bi_index frexp = bi_frexpe_f32(b, s0, true, false);
|
|
|
|
bi_index frexpi = bi_s32_to_f32(b, frexp, BI_ROUND_RTZ);
|
|
|
|
bi_index add = bi_fadd_lscale_f32(b, bi_imm_f32(-1.0f), s0);
|
|
|
|
bi_fma_f32_to(b, dst, bi_flogd_f32(b, s0), add, frexpi,
|
|
|
|
BI_ROUND_NONE);
|
|
|
|
break;
|
|
|
|
}
|
|
|
|
|
2021-05-04 00:00:35 +01:00
|
|
|
case nir_op_bcsel:
|
|
|
|
if (src1_sz == 8)
|
2020-12-16 18:58:12 +00:00
|
|
|
bi_mux_v4i8_to(b, dst, s2, s1, s0, BI_MUX_INT_ZERO);
|
|
|
|
else
|
2021-05-04 00:00:35 +01:00
|
|
|
bi_csel_to(b, nir_type_int, src1_sz,
|
|
|
|
dst, s0, bi_zero(), s1, s2, BI_CMPF_NE);
|
2020-12-16 18:58:12 +00:00
|
|
|
break;
|
|
|
|
|
|
|
|
case nir_op_ishl:
|
|
|
|
bi_lshift_or_to(b, sz, dst, s0, bi_zero(), bi_byte(s1, 0));
|
|
|
|
break;
|
|
|
|
case nir_op_ushr:
|
|
|
|
bi_rshift_or_to(b, sz, dst, s0, bi_zero(), bi_byte(s1, 0));
|
|
|
|
break;
|
|
|
|
|
|
|
|
case nir_op_ishr:
|
|
|
|
bi_arshift_to(b, sz, dst, s0, bi_null(), bi_byte(s1, 0));
|
|
|
|
break;
|
|
|
|
|
2021-05-03 23:23:54 +01:00
|
|
|
case nir_op_imin:
|
|
|
|
case nir_op_umin:
|
|
|
|
bi_csel_to(b, nir_op_infos[instr->op].input_types[0], sz, dst,
|
|
|
|
s0, s1, s0, s1, BI_CMPF_LT);
|
|
|
|
break;
|
|
|
|
|
|
|
|
case nir_op_imax:
|
|
|
|
case nir_op_umax:
|
|
|
|
bi_csel_to(b, nir_op_infos[instr->op].input_types[0], sz, dst,
|
|
|
|
s0, s1, s0, s1, BI_CMPF_GT);
|
|
|
|
break;
|
|
|
|
|
2020-12-16 18:58:12 +00:00
|
|
|
case nir_op_fddx:
|
|
|
|
case nir_op_fddy: {
|
2021-02-10 16:44:37 +00:00
|
|
|
bi_index lane1 = bi_lshift_and_i32(b,
|
|
|
|
bi_fau(BIR_FAU_LANE_ID, false),
|
2020-12-16 18:58:12 +00:00
|
|
|
bi_imm_u32(instr->op == nir_op_fddx ? 2 : 1),
|
2021-02-10 16:44:37 +00:00
|
|
|
bi_imm_u8(0));
|
2020-12-16 18:58:12 +00:00
|
|
|
|
|
|
|
bi_index lane2 = bi_iadd_u32(b, lane1,
|
|
|
|
bi_imm_u32(instr->op == nir_op_fddx ? 1 : 2),
|
|
|
|
false);
|
|
|
|
|
|
|
|
bi_index left, right;
|
|
|
|
|
|
|
|
if (b->shader->arch == 6) {
|
|
|
|
left = bi_clper_v6_i32(b, s0, lane1);
|
|
|
|
right = bi_clper_v6_i32(b, s0, lane2);
|
|
|
|
} else {
|
|
|
|
left = bi_clper_v7_i32(b, s0, lane1,
|
|
|
|
BI_INACTIVE_RESULT_ZERO, BI_LANE_OP_NONE,
|
|
|
|
BI_SUBGROUP_SUBGROUP4);
|
|
|
|
|
|
|
|
right = bi_clper_v7_i32(b, s0, lane2,
|
|
|
|
BI_INACTIVE_RESULT_ZERO, BI_LANE_OP_NONE,
|
|
|
|
BI_SUBGROUP_SUBGROUP4);
|
|
|
|
}
|
|
|
|
|
|
|
|
bi_fadd_to(b, sz, dst, right, bi_neg(left), BI_ROUND_NONE);
|
|
|
|
break;
|
|
|
|
}
|
|
|
|
|
|
|
|
case nir_op_f2f32:
|
|
|
|
bi_f16_to_f32_to(b, dst, s0);
|
|
|
|
break;
|
|
|
|
|
|
|
|
case nir_op_f2i32:
|
|
|
|
if (src_sz == 32)
|
|
|
|
bi_f32_to_s32_to(b, dst, s0, BI_ROUND_RTZ);
|
|
|
|
else
|
|
|
|
bi_f16_to_s32_to(b, dst, s0, BI_ROUND_RTZ);
|
|
|
|
break;
|
|
|
|
|
2021-01-06 22:13:34 +00:00
|
|
|
/* Note 32-bit sources => no vectorization, so 32-bit works */
|
2020-12-16 18:58:12 +00:00
|
|
|
case nir_op_f2u16:
|
2021-01-06 22:13:34 +00:00
|
|
|
if (src_sz == 32)
|
|
|
|
bi_f32_to_u32_to(b, dst, s0, BI_ROUND_RTZ);
|
|
|
|
else
|
2020-12-16 18:58:12 +00:00
|
|
|
bi_v2f16_to_v2u16_to(b, dst, s0, BI_ROUND_RTZ);
|
|
|
|
break;
|
|
|
|
|
|
|
|
case nir_op_f2i16:
|
2021-01-06 22:13:34 +00:00
|
|
|
if (src_sz == 32)
|
|
|
|
bi_f32_to_s32_to(b, dst, s0, BI_ROUND_RTZ);
|
|
|
|
else
|
2020-12-16 18:58:12 +00:00
|
|
|
bi_v2f16_to_v2s16_to(b, dst, s0, BI_ROUND_RTZ);
|
|
|
|
break;
|
|
|
|
|
|
|
|
case nir_op_f2u32:
|
|
|
|
if (src_sz == 32)
|
|
|
|
bi_f32_to_u32_to(b, dst, s0, BI_ROUND_RTZ);
|
|
|
|
else
|
|
|
|
bi_f16_to_u32_to(b, dst, s0, BI_ROUND_RTZ);
|
|
|
|
break;
|
|
|
|
|
|
|
|
case nir_op_u2f16:
|
|
|
|
if (src_sz == 32)
|
2021-01-01 03:14:42 +00:00
|
|
|
bi_v2u16_to_v2f16_to(b, dst, bi_half(s0, false), BI_ROUND_RTZ);
|
2020-12-16 18:58:12 +00:00
|
|
|
else if (src_sz == 16)
|
|
|
|
bi_v2u16_to_v2f16_to(b, dst, s0, BI_ROUND_RTZ);
|
|
|
|
else if (src_sz == 8)
|
|
|
|
bi_v2u8_to_v2f16_to(b, dst, s0);
|
|
|
|
break;
|
|
|
|
|
|
|
|
case nir_op_u2f32:
|
|
|
|
if (src_sz == 32)
|
|
|
|
bi_u32_to_f32_to(b, dst, s0, BI_ROUND_RTZ);
|
|
|
|
else if (src_sz == 16)
|
|
|
|
bi_u16_to_f32_to(b, dst, s0);
|
|
|
|
else
|
2021-04-13 18:22:18 +01:00
|
|
|
bi_u8_to_f32_to(b, dst, s0);
|
2020-12-16 18:58:12 +00:00
|
|
|
break;
|
|
|
|
|
|
|
|
case nir_op_i2f16:
|
|
|
|
if (src_sz == 32)
|
2021-01-01 03:14:42 +00:00
|
|
|
bi_v2s16_to_v2f16_to(b, dst, bi_half(s0, false), BI_ROUND_RTZ);
|
|
|
|
else if (src_sz == 16)
|
2020-12-16 18:58:12 +00:00
|
|
|
bi_v2s16_to_v2f16_to(b, dst, s0, BI_ROUND_RTZ);
|
2021-01-01 03:14:42 +00:00
|
|
|
else if (src_sz == 8)
|
|
|
|
bi_v2s8_to_v2f16_to(b, dst, s0);
|
2020-12-16 18:58:12 +00:00
|
|
|
break;
|
|
|
|
|
|
|
|
case nir_op_i2f32:
|
|
|
|
if (src_sz == 32)
|
|
|
|
bi_s32_to_f32_to(b, dst, s0, BI_ROUND_RTZ);
|
2021-04-13 18:22:18 +01:00
|
|
|
else if (src_sz == 16)
|
2020-12-16 18:58:12 +00:00
|
|
|
bi_s16_to_f32_to(b, dst, s0);
|
2021-04-13 18:22:18 +01:00
|
|
|
else if (src_sz == 8)
|
|
|
|
bi_s8_to_f32_to(b, dst, s0);
|
2020-12-16 18:58:12 +00:00
|
|
|
break;
|
|
|
|
|
|
|
|
case nir_op_i2i32:
|
|
|
|
if (src_sz == 16)
|
|
|
|
bi_s16_to_s32_to(b, dst, s0);
|
|
|
|
else
|
|
|
|
bi_s8_to_s32_to(b, dst, s0);
|
|
|
|
break;
|
|
|
|
|
|
|
|
case nir_op_u2u32:
|
|
|
|
if (src_sz == 16)
|
|
|
|
bi_u16_to_u32_to(b, dst, s0);
|
|
|
|
else
|
|
|
|
bi_u8_to_u32_to(b, dst, s0);
|
|
|
|
break;
|
|
|
|
|
|
|
|
case nir_op_i2i16:
|
|
|
|
assert(src_sz == 8 || src_sz == 32);
|
|
|
|
|
|
|
|
if (src_sz == 8)
|
|
|
|
bi_v2s8_to_v2s16_to(b, dst, s0);
|
|
|
|
else
|
2021-01-01 03:14:42 +00:00
|
|
|
bi_mov_i32_to(b, dst, s0);
|
2020-12-16 18:58:12 +00:00
|
|
|
break;
|
|
|
|
|
|
|
|
case nir_op_u2u16:
|
|
|
|
assert(src_sz == 8 || src_sz == 32);
|
|
|
|
|
|
|
|
if (src_sz == 8)
|
|
|
|
bi_v2u8_to_v2u16_to(b, dst, s0);
|
|
|
|
else
|
2021-01-01 03:14:42 +00:00
|
|
|
bi_mov_i32_to(b, dst, s0);
|
2020-12-16 18:58:12 +00:00
|
|
|
break;
|
|
|
|
|
2021-05-03 23:33:57 +01:00
|
|
|
case nir_op_b2f16:
|
|
|
|
case nir_op_b2f32:
|
|
|
|
bi_csel_to(b, nir_type_int, sz, dst, s0, bi_zero(),
|
|
|
|
(sz == 16) ? bi_imm_f16(1.0) : bi_imm_f32(1.0),
|
|
|
|
(sz == 16) ? bi_imm_f16(0.0) : bi_imm_f32(0.0),
|
|
|
|
BI_CMPF_NE);
|
|
|
|
break;
|
|
|
|
|
2021-05-04 00:00:35 +01:00
|
|
|
case nir_op_b2b32:
|
|
|
|
bi_csel_to(b, nir_type_int, sz, dst, s0, bi_zero(),
|
|
|
|
bi_imm_u32(~0), bi_zero(), BI_CMPF_NE);
|
|
|
|
break;
|
|
|
|
|
2021-05-03 23:37:41 +01:00
|
|
|
case nir_op_b2i8:
|
|
|
|
case nir_op_b2i16:
|
|
|
|
case nir_op_b2i32:
|
|
|
|
bi_lshift_and_to(b, sz, dst, s0, bi_imm_uintN(1, sz), bi_imm_u8(0));
|
|
|
|
break;
|
|
|
|
|
2020-12-16 18:58:12 +00:00
|
|
|
case nir_op_fround_even:
|
|
|
|
case nir_op_fceil:
|
|
|
|
case nir_op_ffloor:
|
|
|
|
case nir_op_ftrunc:
|
|
|
|
bi_fround_to(b, sz, dst, s0, bi_nir_round(instr->op));
|
|
|
|
break;
|
|
|
|
|
|
|
|
case nir_op_fmin:
|
|
|
|
bi_fmin_to(b, sz, dst, s0, s1);
|
|
|
|
break;
|
|
|
|
|
|
|
|
case nir_op_fmax:
|
|
|
|
bi_fmax_to(b, sz, dst, s0, s1);
|
|
|
|
break;
|
|
|
|
|
|
|
|
case nir_op_iadd:
|
2021-03-02 12:08:05 +00:00
|
|
|
bi_iadd_to(b, nir_type_int, sz, dst, s0, s1, false);
|
2020-12-16 18:58:12 +00:00
|
|
|
break;
|
|
|
|
|
2020-12-29 12:08:55 +00:00
|
|
|
case nir_op_iadd_sat:
|
2021-03-02 12:08:05 +00:00
|
|
|
bi_iadd_to(b, nir_type_int, sz, dst, s0, s1, true);
|
2020-12-29 12:08:55 +00:00
|
|
|
break;
|
|
|
|
|
2021-03-02 12:14:27 +00:00
|
|
|
case nir_op_uadd_sat:
|
|
|
|
bi_iadd_to(b, nir_type_uint, sz, dst, s0, s1, true);
|
|
|
|
break;
|
|
|
|
|
2020-12-31 13:05:34 +00:00
|
|
|
case nir_op_ihadd:
|
2021-03-02 12:08:05 +00:00
|
|
|
bi_hadd_to(b, nir_type_int, sz, dst, s0, s1, BI_ROUND_RTN);
|
2020-12-31 13:05:34 +00:00
|
|
|
break;
|
|
|
|
|
|
|
|
case nir_op_irhadd:
|
2021-03-02 12:08:05 +00:00
|
|
|
bi_hadd_to(b, nir_type_int, sz, dst, s0, s1, BI_ROUND_RTP);
|
2020-12-31 13:05:34 +00:00
|
|
|
break;
|
|
|
|
|
2021-05-03 23:29:35 +01:00
|
|
|
case nir_op_ineg:
|
|
|
|
bi_isub_to(b, nir_type_int, sz, dst, bi_zero(), s0, false);
|
|
|
|
break;
|
|
|
|
|
2020-12-16 18:58:12 +00:00
|
|
|
case nir_op_isub:
|
2021-03-02 12:08:05 +00:00
|
|
|
bi_isub_to(b, nir_type_int, sz, dst, s0, s1, false);
|
2020-12-16 18:58:12 +00:00
|
|
|
break;
|
|
|
|
|
2020-12-29 12:08:55 +00:00
|
|
|
case nir_op_isub_sat:
|
2021-03-02 12:08:05 +00:00
|
|
|
bi_isub_to(b, nir_type_int, sz, dst, s0, s1, true);
|
2020-12-29 12:08:55 +00:00
|
|
|
break;
|
|
|
|
|
2021-03-02 12:14:27 +00:00
|
|
|
case nir_op_usub_sat:
|
|
|
|
bi_isub_to(b, nir_type_uint, sz, dst, s0, s1, true);
|
|
|
|
break;
|
|
|
|
|
2020-12-16 18:58:12 +00:00
|
|
|
case nir_op_imul:
|
|
|
|
bi_imul_to(b, sz, dst, s0, s1);
|
|
|
|
break;
|
|
|
|
|
|
|
|
case nir_op_iabs:
|
|
|
|
bi_iabs_to(b, sz, dst, s0);
|
|
|
|
break;
|
|
|
|
|
|
|
|
case nir_op_iand:
|
|
|
|
bi_lshift_and_to(b, sz, dst, s0, s1, bi_imm_u8(0));
|
|
|
|
break;
|
|
|
|
|
|
|
|
case nir_op_ior:
|
|
|
|
bi_lshift_or_to(b, sz, dst, s0, s1, bi_imm_u8(0));
|
|
|
|
break;
|
|
|
|
|
|
|
|
case nir_op_ixor:
|
|
|
|
bi_lshift_xor_to(b, sz, dst, s0, s1, bi_imm_u8(0));
|
|
|
|
break;
|
|
|
|
|
|
|
|
case nir_op_inot:
|
|
|
|
bi_lshift_or_to(b, sz, dst, bi_zero(), bi_not(s0), bi_imm_u8(0));
|
|
|
|
break;
|
|
|
|
|
|
|
|
case nir_op_frsq:
|
2021-01-15 21:39:58 +00:00
|
|
|
if (sz == 32 && b->shader->quirks & BIFROST_NO_FP32_TRANSCENDENTALS)
|
|
|
|
bi_lower_frsq_32(b, dst, s0);
|
|
|
|
else
|
|
|
|
bi_frsq_to(b, sz, dst, s0);
|
2020-12-16 18:58:12 +00:00
|
|
|
break;
|
|
|
|
|
|
|
|
case nir_op_frcp:
|
2021-01-15 21:39:58 +00:00
|
|
|
if (sz == 32 && b->shader->quirks & BIFROST_NO_FP32_TRANSCENDENTALS)
|
|
|
|
bi_lower_frcp_32(b, dst, s0);
|
|
|
|
else
|
|
|
|
bi_frcp_to(b, sz, dst, s0);
|
2020-12-16 18:58:12 +00:00
|
|
|
break;
|
|
|
|
|
2021-01-25 12:21:28 +00:00
|
|
|
case nir_op_uclz:
|
|
|
|
bi_clz_to(b, sz, dst, s0, false);
|
|
|
|
break;
|
|
|
|
|
2021-01-25 12:22:05 +00:00
|
|
|
case nir_op_bit_count:
|
|
|
|
bi_popcount_i32_to(b, dst, s0);
|
|
|
|
break;
|
|
|
|
|
2021-01-27 17:17:26 +00:00
|
|
|
case nir_op_bitfield_reverse:
|
|
|
|
bi_bitrev_i32_to(b, dst, s0);
|
|
|
|
break;
|
|
|
|
|
2021-01-27 17:31:45 +00:00
|
|
|
case nir_op_ufind_msb: {
|
|
|
|
bi_index clz = bi_clz(b, src_sz, s0, false);
|
|
|
|
|
|
|
|
if (sz == 8)
|
|
|
|
clz = bi_byte(clz, 0);
|
|
|
|
else if (sz == 16)
|
|
|
|
clz = bi_half(clz, false);
|
|
|
|
|
|
|
|
bi_isub_u32_to(b, dst, bi_imm_u32(src_sz - 1), clz, false);
|
|
|
|
break;
|
|
|
|
}
|
|
|
|
|
2020-12-16 18:58:12 +00:00
|
|
|
default:
|
|
|
|
fprintf(stderr, "Unhandled ALU op %s\n", nir_op_infos[instr->op].name);
|
|
|
|
unreachable("Unknown ALU op");
|
|
|
|
}
|
|
|
|
}
|
|
|
|
|
2020-10-06 16:18:16 +01:00
|
|
|
/* Returns dimension with 0 special casing cubemaps. Shamelessly copied from Midgard */
|
|
|
|
static unsigned
|
|
|
|
bifrost_tex_format(enum glsl_sampler_dim dim)
|
|
|
|
{
|
|
|
|
switch (dim) {
|
|
|
|
case GLSL_SAMPLER_DIM_1D:
|
|
|
|
case GLSL_SAMPLER_DIM_BUF:
|
|
|
|
return 1;
|
|
|
|
|
|
|
|
case GLSL_SAMPLER_DIM_2D:
|
|
|
|
case GLSL_SAMPLER_DIM_MS:
|
|
|
|
case GLSL_SAMPLER_DIM_EXTERNAL:
|
|
|
|
case GLSL_SAMPLER_DIM_RECT:
|
|
|
|
return 2;
|
|
|
|
|
|
|
|
case GLSL_SAMPLER_DIM_3D:
|
|
|
|
return 3;
|
|
|
|
|
|
|
|
case GLSL_SAMPLER_DIM_CUBE:
|
|
|
|
return 0;
|
|
|
|
|
|
|
|
default:
|
|
|
|
DBG("Unknown sampler dim type\n");
|
|
|
|
assert(0);
|
|
|
|
return 0;
|
|
|
|
}
|
|
|
|
}
|
|
|
|
|
|
|
|
static enum bifrost_texture_format_full
|
2020-11-27 15:33:19 +00:00
|
|
|
bi_texture_format(nir_alu_type T, enum bi_clamp clamp)
|
2020-10-06 16:18:16 +01:00
|
|
|
{
|
|
|
|
switch (T) {
|
2020-11-27 15:33:19 +00:00
|
|
|
case nir_type_float16: return BIFROST_TEXTURE_FORMAT_F16 + clamp;
|
|
|
|
case nir_type_float32: return BIFROST_TEXTURE_FORMAT_F32 + clamp;
|
2020-10-06 16:18:16 +01:00
|
|
|
case nir_type_uint16: return BIFROST_TEXTURE_FORMAT_U16;
|
|
|
|
case nir_type_int16: return BIFROST_TEXTURE_FORMAT_S16;
|
|
|
|
case nir_type_uint32: return BIFROST_TEXTURE_FORMAT_U32;
|
|
|
|
case nir_type_int32: return BIFROST_TEXTURE_FORMAT_S32;
|
|
|
|
default: unreachable("Invalid type for texturing");
|
|
|
|
}
|
|
|
|
}
|
|
|
|
|
2020-11-03 13:22:42 +00:00
|
|
|
/* Array indices are specified as 32-bit uints, need to convert. In .z component from NIR */
|
2020-12-16 18:01:45 +00:00
|
|
|
static bi_index
|
|
|
|
bi_emit_texc_array_index(bi_builder *b, bi_index idx, nir_alu_type T)
|
|
|
|
{
|
|
|
|
/* For (u)int we can just passthrough */
|
|
|
|
nir_alu_type base = nir_alu_type_get_base_type(T);
|
|
|
|
if (base == nir_type_int || base == nir_type_uint)
|
|
|
|
return idx;
|
|
|
|
|
|
|
|
/* Otherwise we convert */
|
|
|
|
assert(T == nir_type_float32);
|
|
|
|
|
|
|
|
/* OpenGL ES 3.2 specification section 8.14.2 ("Coordinate Wrapping and
|
|
|
|
* Texel Selection") defines the layer to be taken from clamp(RNE(r),
|
|
|
|
* 0, dt - 1). So we use round RTE, clamping is handled at the data
|
|
|
|
* structure level */
|
|
|
|
|
|
|
|
return bi_f32_to_u32(b, idx, BI_ROUND_NONE);
|
|
|
|
}
|
|
|
|
|
2020-10-06 16:46:50 +01:00
|
|
|
/* TEXC's explicit and bias LOD modes requires the LOD to be transformed to a
|
|
|
|
* 16-bit 8:8 fixed-point format. We lower as:
|
|
|
|
*
|
|
|
|
* F32_TO_S32(clamp(x, -16.0, +16.0) * 256.0) & 0xFFFF =
|
|
|
|
* MKVEC(F32_TO_S32(clamp(x * 1.0/16.0, -1.0, 1.0) * (16.0 * 256.0)), #0)
|
|
|
|
*/
|
|
|
|
|
2020-12-16 18:01:45 +00:00
|
|
|
static bi_index
|
|
|
|
bi_emit_texc_lod_88(bi_builder *b, bi_index lod, bool fp16)
|
|
|
|
{
|
|
|
|
/* Sort of arbitrary. Must be less than 128.0, greater than or equal to
|
|
|
|
* the max LOD (16 since we cap at 2^16 texture dimensions), and
|
|
|
|
* preferably small to minimize precision loss */
|
|
|
|
const float max_lod = 16.0;
|
|
|
|
|
|
|
|
bi_instr *fsat = bi_fma_f32_to(b, bi_temp(b->shader),
|
|
|
|
fp16 ? bi_half(lod, false) : lod,
|
2021-04-13 19:36:08 +01:00
|
|
|
bi_imm_f32(1.0f / max_lod), bi_negzero(), BI_ROUND_NONE);
|
2020-12-16 18:01:45 +00:00
|
|
|
|
|
|
|
fsat->clamp = BI_CLAMP_CLAMP_M1_1;
|
|
|
|
|
|
|
|
bi_index fmul = bi_fma_f32(b, fsat->dest[0], bi_imm_f32(max_lod * 256.0f),
|
2021-04-13 19:36:08 +01:00
|
|
|
bi_negzero(), BI_ROUND_NONE);
|
2020-12-16 18:01:45 +00:00
|
|
|
|
|
|
|
return bi_mkvec_v2i16(b,
|
|
|
|
bi_half(bi_f32_to_s32(b, fmul, BI_ROUND_RTZ), false),
|
|
|
|
bi_imm_u16(0));
|
|
|
|
}
|
|
|
|
|
2020-10-09 00:31:41 +01:00
|
|
|
/* FETCH takes a 32-bit staging register containing the LOD as an integer in
|
|
|
|
* the bottom 16-bits and (if present) the cube face index in the top 16-bits.
|
|
|
|
* TODO: Cube face.
|
|
|
|
*/
|
|
|
|
|
2020-12-16 18:01:45 +00:00
|
|
|
static bi_index
|
|
|
|
bi_emit_texc_lod_cube(bi_builder *b, bi_index lod)
|
|
|
|
{
|
|
|
|
return bi_lshift_or_i32(b, lod, bi_zero(), bi_imm_u8(8));
|
|
|
|
}
|
|
|
|
|
2020-10-17 10:32:41 +01:00
|
|
|
/* The hardware specifies texel offsets and multisample indices together as a
|
|
|
|
* u8vec4 <offset, ms index>. By default all are zero, so if have either a
|
|
|
|
* nonzero texel offset or a nonzero multisample index, we build a u8vec4 with
|
|
|
|
* the bits we need and return that to be passed as a staging register. Else we
|
|
|
|
* return 0 to avoid allocating a data register when everything is zero. */
|
|
|
|
|
2020-12-16 18:01:45 +00:00
|
|
|
static bi_index
|
|
|
|
bi_emit_texc_offset_ms_index(bi_builder *b, nir_tex_instr *instr)
|
|
|
|
{
|
|
|
|
bi_index dest = bi_zero();
|
|
|
|
|
|
|
|
int offs_idx = nir_tex_instr_src_index(instr, nir_tex_src_offset);
|
|
|
|
if (offs_idx >= 0 &&
|
|
|
|
(!nir_src_is_const(instr->src[offs_idx].src) ||
|
|
|
|
nir_src_as_uint(instr->src[offs_idx].src) != 0)) {
|
|
|
|
unsigned nr = nir_src_num_components(instr->src[offs_idx].src);
|
|
|
|
bi_index idx = bi_src_index(&instr->src[offs_idx].src);
|
|
|
|
dest = bi_mkvec_v4i8(b,
|
|
|
|
(nr > 0) ? bi_byte(bi_word(idx, 0), 0) : bi_imm_u8(0),
|
|
|
|
(nr > 1) ? bi_byte(bi_word(idx, 1), 0) : bi_imm_u8(0),
|
|
|
|
(nr > 2) ? bi_byte(bi_word(idx, 2), 0) : bi_imm_u8(0),
|
|
|
|
bi_imm_u8(0));
|
|
|
|
}
|
|
|
|
|
|
|
|
int ms_idx = nir_tex_instr_src_index(instr, nir_tex_src_ms_index);
|
|
|
|
if (ms_idx >= 0 &&
|
|
|
|
(!nir_src_is_const(instr->src[ms_idx].src) ||
|
|
|
|
nir_src_as_uint(instr->src[ms_idx].src) != 0)) {
|
2020-12-16 19:24:18 +00:00
|
|
|
dest = bi_lshift_or_i32(b,
|
|
|
|
bi_src_index(&instr->src[ms_idx].src), dest,
|
|
|
|
bi_imm_u8(24));
|
2020-10-17 10:32:41 +01:00
|
|
|
}
|
|
|
|
|
|
|
|
return dest;
|
|
|
|
}
|
|
|
|
|
2020-12-16 18:01:45 +00:00
|
|
|
static void
|
|
|
|
bi_emit_cube_coord(bi_builder *b, bi_index coord,
|
|
|
|
bi_index *face, bi_index *s, bi_index *t)
|
|
|
|
{
|
|
|
|
/* Compute max { |x|, |y|, |z| } */
|
2021-02-18 19:58:50 +00:00
|
|
|
bi_instr *cubeface = bi_cubeface_to(b, bi_temp(b->shader),
|
|
|
|
bi_temp(b->shader), coord,
|
2020-12-16 18:01:45 +00:00
|
|
|
bi_word(coord, 1), bi_word(coord, 2));
|
|
|
|
|
|
|
|
/* Select coordinates */
|
|
|
|
|
|
|
|
bi_index ssel = bi_cube_ssel(b, bi_word(coord, 2), coord,
|
2021-01-06 20:02:28 +00:00
|
|
|
cubeface->dest[1]);
|
2020-12-16 18:01:45 +00:00
|
|
|
|
|
|
|
bi_index tsel = bi_cube_tsel(b, bi_word(coord, 1), bi_word(coord, 2),
|
2021-01-06 20:02:28 +00:00
|
|
|
cubeface->dest[1]);
|
2020-12-16 18:01:45 +00:00
|
|
|
|
|
|
|
/* The OpenGL ES specification requires us to transform an input vector
|
|
|
|
* (x, y, z) to the coordinate, given the selected S/T:
|
|
|
|
*
|
|
|
|
* (1/2 ((s / max{x,y,z}) + 1), 1/2 ((t / max{x, y, z}) + 1))
|
|
|
|
*
|
|
|
|
* We implement (s shown, t similar) in a form friendlier to FMA
|
|
|
|
* instructions, and clamp coordinates at the end for correct
|
|
|
|
* NaN/infinity handling:
|
|
|
|
*
|
|
|
|
* fsat(s * (0.5 * (1 / max{x, y, z})) + 0.5)
|
|
|
|
*
|
|
|
|
* Take the reciprocal of max{x, y, z}
|
|
|
|
*/
|
|
|
|
|
2021-01-06 20:02:28 +00:00
|
|
|
bi_index rcp = bi_frcp_f32(b, cubeface->dest[0]);
|
2020-12-16 18:01:45 +00:00
|
|
|
|
|
|
|
/* Calculate 0.5 * (1.0 / max{x, y, z}) */
|
2021-04-13 19:36:08 +01:00
|
|
|
bi_index fma1 = bi_fma_f32(b, rcp, bi_imm_f32(0.5f), bi_negzero(),
|
2020-12-16 18:01:45 +00:00
|
|
|
BI_ROUND_NONE);
|
|
|
|
|
|
|
|
/* Transform the coordinates */
|
|
|
|
*s = bi_temp(b->shader);
|
|
|
|
*t = bi_temp(b->shader);
|
|
|
|
|
|
|
|
bi_instr *S = bi_fma_f32_to(b, *s, fma1, ssel, bi_imm_f32(0.5f),
|
|
|
|
BI_ROUND_NONE);
|
|
|
|
bi_instr *T = bi_fma_f32_to(b, *t, fma1, tsel, bi_imm_f32(0.5f),
|
|
|
|
BI_ROUND_NONE);
|
|
|
|
|
|
|
|
S->clamp = BI_CLAMP_CLAMP_0_1;
|
|
|
|
T->clamp = BI_CLAMP_CLAMP_0_1;
|
|
|
|
|
|
|
|
/* Cube face is stored in bit[29:31], we don't apply the shift here
|
|
|
|
* because the TEXS_CUBE and TEXC instructions expect the face index to
|
|
|
|
* be at this position.
|
|
|
|
*/
|
2021-01-06 20:02:28 +00:00
|
|
|
*face = cubeface->dest[1];
|
2020-12-16 18:01:45 +00:00
|
|
|
}
|
|
|
|
|
|
|
|
/* Emits a cube map descriptor, returning lower 32-bits and putting upper
|
|
|
|
* 32-bits in passed pointer t */
|
|
|
|
|
|
|
|
static bi_index
|
|
|
|
bi_emit_texc_cube_coord(bi_builder *b, bi_index coord, bi_index *t)
|
|
|
|
{
|
|
|
|
bi_index face, s;
|
|
|
|
bi_emit_cube_coord(b, coord, &face, &s, t);
|
|
|
|
|
|
|
|
bi_index and1 = bi_lshift_and_i32(b, face, bi_imm_u32(0xe0000000),
|
|
|
|
bi_imm_u8(0));
|
|
|
|
|
|
|
|
bi_index and2 = bi_lshift_and_i32(b, s, bi_imm_u32(0x1fffffff),
|
|
|
|
bi_imm_u8(0));
|
|
|
|
|
|
|
|
return bi_lshift_or_i32(b, and1, and2, bi_imm_u8(0));
|
|
|
|
}
|
|
|
|
|
2020-10-09 00:01:29 +01:00
|
|
|
/* Map to the main texture op used. Some of these (txd in particular) will
|
|
|
|
* lower to multiple texture ops with different opcodes (GRDESC_DER + TEX in
|
|
|
|
* sequence). We assume that lowering is handled elsewhere.
|
|
|
|
*/
|
|
|
|
|
|
|
|
static enum bifrost_tex_op
|
|
|
|
bi_tex_op(nir_texop op)
|
|
|
|
{
|
|
|
|
switch (op) {
|
|
|
|
case nir_texop_tex:
|
|
|
|
case nir_texop_txb:
|
|
|
|
case nir_texop_txl:
|
|
|
|
case nir_texop_txd:
|
|
|
|
case nir_texop_tex_prefetch:
|
|
|
|
return BIFROST_TEX_OP_TEX;
|
|
|
|
case nir_texop_txf:
|
|
|
|
case nir_texop_txf_ms:
|
|
|
|
case nir_texop_txf_ms_fb:
|
|
|
|
case nir_texop_txf_ms_mcs:
|
|
|
|
case nir_texop_tg4:
|
|
|
|
return BIFROST_TEX_OP_FETCH;
|
|
|
|
case nir_texop_txs:
|
|
|
|
case nir_texop_lod:
|
|
|
|
case nir_texop_query_levels:
|
|
|
|
case nir_texop_texture_samples:
|
|
|
|
case nir_texop_samples_identical:
|
|
|
|
unreachable("should've been lowered");
|
|
|
|
default:
|
|
|
|
unreachable("unsupported tex op");
|
|
|
|
}
|
|
|
|
}
|
|
|
|
|
2020-10-06 16:32:55 +01:00
|
|
|
/* Data registers required by texturing in the order they appear. All are
|
|
|
|
* optional, the texture operation descriptor determines which are present.
|
|
|
|
* Note since 3D arrays are not permitted at an API level, Z_COORD and
|
|
|
|
* ARRAY/SHADOW are exlusive, so TEXC in practice reads at most 8 registers */
|
|
|
|
|
|
|
|
enum bifrost_tex_dreg {
|
|
|
|
BIFROST_TEX_DREG_Z_COORD = 0,
|
|
|
|
BIFROST_TEX_DREG_Y_DELTAS = 1,
|
|
|
|
BIFROST_TEX_DREG_LOD = 2,
|
|
|
|
BIFROST_TEX_DREG_GRDESC_HI = 3,
|
|
|
|
BIFROST_TEX_DREG_SHADOW = 4,
|
|
|
|
BIFROST_TEX_DREG_ARRAY = 5,
|
|
|
|
BIFROST_TEX_DREG_OFFSETMS = 6,
|
|
|
|
BIFROST_TEX_DREG_SAMPLER = 7,
|
|
|
|
BIFROST_TEX_DREG_TEXTURE = 8,
|
|
|
|
BIFROST_TEX_DREG_COUNT,
|
|
|
|
};
|
|
|
|
|
2020-12-16 18:02:03 +00:00
|
|
|
static void
|
|
|
|
bi_emit_texc(bi_builder *b, nir_tex_instr *instr)
|
|
|
|
{
|
|
|
|
struct bifrost_texture_operation desc = {
|
|
|
|
.op = bi_tex_op(instr->op),
|
|
|
|
.offset_or_bias_disable = false, /* TODO */
|
|
|
|
.shadow_or_clamp_disable = instr->is_shadow,
|
|
|
|
.array = instr->is_array,
|
|
|
|
.dimension = bifrost_tex_format(instr->sampler_dim),
|
|
|
|
.format = bi_texture_format(instr->dest_type | nir_dest_bit_size(instr->dest), BI_CLAMP_NONE), /* TODO */
|
|
|
|
.mask = 0xF,
|
|
|
|
};
|
|
|
|
|
|
|
|
switch (desc.op) {
|
|
|
|
case BIFROST_TEX_OP_TEX:
|
|
|
|
desc.lod_or_fetch = BIFROST_LOD_MODE_COMPUTE;
|
|
|
|
break;
|
|
|
|
case BIFROST_TEX_OP_FETCH:
|
2021-01-27 16:22:10 +00:00
|
|
|
desc.lod_or_fetch = instr->op == nir_texop_tg4 ?
|
|
|
|
BIFROST_TEXTURE_FETCH_GATHER4_R + instr->component :
|
|
|
|
BIFROST_TEXTURE_FETCH_TEXEL;
|
2020-12-16 18:02:03 +00:00
|
|
|
break;
|
|
|
|
default:
|
|
|
|
unreachable("texture op unsupported");
|
|
|
|
}
|
|
|
|
|
|
|
|
/* 32-bit indices to be allocated as consecutive staging registers */
|
|
|
|
bi_index dregs[BIFROST_TEX_DREG_COUNT] = { };
|
|
|
|
bi_index cx = bi_null(), cy = bi_null();
|
|
|
|
|
|
|
|
for (unsigned i = 0; i < instr->num_srcs; ++i) {
|
|
|
|
bi_index index = bi_src_index(&instr->src[i].src);
|
|
|
|
unsigned sz = nir_src_bit_size(instr->src[i].src);
|
|
|
|
ASSERTED nir_alu_type base = nir_tex_instr_src_type(instr, i);
|
|
|
|
nir_alu_type T = base | sz;
|
|
|
|
|
|
|
|
switch (instr->src[i].src_type) {
|
|
|
|
case nir_tex_src_coord:
|
|
|
|
if (instr->sampler_dim == GLSL_SAMPLER_DIM_CUBE) {
|
|
|
|
cx = bi_emit_texc_cube_coord(b, index, &cy);
|
|
|
|
} else {
|
|
|
|
unsigned components = nir_src_num_components(instr->src[i].src);
|
|
|
|
|
|
|
|
/* Copy XY (for 2D+) or XX (for 1D) */
|
|
|
|
cx = index;
|
|
|
|
cy = bi_word(index, MIN2(1, components - 1));
|
|
|
|
|
|
|
|
assert(components >= 1 && components <= 3);
|
|
|
|
|
|
|
|
if (components < 3) {
|
|
|
|
/* nothing to do */
|
|
|
|
} else if (desc.array) {
|
|
|
|
/* 2D array */
|
|
|
|
dregs[BIFROST_TEX_DREG_ARRAY] =
|
|
|
|
bi_emit_texc_array_index(b,
|
|
|
|
bi_word(index, 2), T);
|
|
|
|
} else {
|
|
|
|
/* 3D */
|
|
|
|
dregs[BIFROST_TEX_DREG_Z_COORD] =
|
|
|
|
bi_word(index, 2);
|
|
|
|
}
|
|
|
|
}
|
|
|
|
break;
|
|
|
|
|
|
|
|
case nir_tex_src_lod:
|
|
|
|
if (desc.op == BIFROST_TEX_OP_TEX &&
|
|
|
|
nir_src_is_const(instr->src[i].src) &&
|
|
|
|
nir_src_as_uint(instr->src[i].src) == 0) {
|
|
|
|
desc.lod_or_fetch = BIFROST_LOD_MODE_ZERO;
|
|
|
|
} else if (desc.op == BIFROST_TEX_OP_TEX) {
|
|
|
|
assert(base == nir_type_float);
|
|
|
|
|
|
|
|
assert(sz == 16 || sz == 32);
|
|
|
|
dregs[BIFROST_TEX_DREG_LOD] =
|
|
|
|
bi_emit_texc_lod_88(b, index, sz == 16);
|
|
|
|
desc.lod_or_fetch = BIFROST_LOD_MODE_EXPLICIT;
|
|
|
|
} else {
|
|
|
|
assert(desc.op == BIFROST_TEX_OP_FETCH);
|
|
|
|
assert(base == nir_type_uint || base == nir_type_int);
|
|
|
|
assert(sz == 16 || sz == 32);
|
|
|
|
|
|
|
|
dregs[BIFROST_TEX_DREG_LOD] =
|
|
|
|
bi_emit_texc_lod_cube(b, index);
|
|
|
|
}
|
|
|
|
|
|
|
|
break;
|
|
|
|
|
|
|
|
case nir_tex_src_bias:
|
|
|
|
/* Upper 16-bits interpreted as a clamp, leave zero */
|
|
|
|
assert(desc.op == BIFROST_TEX_OP_TEX);
|
|
|
|
assert(base == nir_type_float);
|
|
|
|
assert(sz == 16 || sz == 32);
|
|
|
|
dregs[BIFROST_TEX_DREG_LOD] =
|
|
|
|
bi_emit_texc_lod_88(b, index, sz == 16);
|
|
|
|
desc.lod_or_fetch = BIFROST_LOD_MODE_BIAS;
|
|
|
|
break;
|
|
|
|
|
|
|
|
case nir_tex_src_ms_index:
|
|
|
|
case nir_tex_src_offset:
|
|
|
|
if (desc.offset_or_bias_disable)
|
|
|
|
break;
|
|
|
|
|
|
|
|
dregs[BIFROST_TEX_DREG_OFFSETMS] =
|
|
|
|
bi_emit_texc_offset_ms_index(b, instr);
|
|
|
|
if (!bi_is_equiv(dregs[BIFROST_TEX_DREG_OFFSETMS], bi_zero()))
|
|
|
|
desc.offset_or_bias_disable = true;
|
|
|
|
break;
|
|
|
|
|
|
|
|
case nir_tex_src_comparator:
|
|
|
|
dregs[BIFROST_TEX_DREG_SHADOW] = index;
|
|
|
|
break;
|
|
|
|
|
2021-02-17 01:40:42 +00:00
|
|
|
case nir_tex_src_texture_offset:
|
|
|
|
assert(instr->texture_index == 0);
|
|
|
|
dregs[BIFROST_TEX_DREG_TEXTURE] = index;
|
|
|
|
break;
|
|
|
|
|
|
|
|
case nir_tex_src_sampler_offset:
|
|
|
|
assert(instr->sampler_index == 0);
|
|
|
|
dregs[BIFROST_TEX_DREG_SAMPLER] = index;
|
|
|
|
break;
|
|
|
|
|
2020-12-16 18:02:03 +00:00
|
|
|
default:
|
|
|
|
unreachable("Unhandled src type in texc emit");
|
|
|
|
}
|
|
|
|
}
|
|
|
|
|
|
|
|
if (desc.op == BIFROST_TEX_OP_FETCH && bi_is_null(dregs[BIFROST_TEX_DREG_LOD])) {
|
|
|
|
dregs[BIFROST_TEX_DREG_LOD] =
|
|
|
|
bi_emit_texc_lod_cube(b, bi_zero());
|
|
|
|
}
|
|
|
|
|
2021-02-17 01:40:42 +00:00
|
|
|
/* Choose an index mode */
|
|
|
|
|
|
|
|
bool direct_tex = bi_is_null(dregs[BIFROST_TEX_DREG_TEXTURE]);
|
|
|
|
bool direct_samp = bi_is_null(dregs[BIFROST_TEX_DREG_SAMPLER]);
|
|
|
|
bool direct = direct_tex && direct_samp;
|
|
|
|
|
|
|
|
desc.immediate_indices = direct && (instr->sampler_index < 16);
|
|
|
|
|
|
|
|
if (desc.immediate_indices) {
|
|
|
|
desc.sampler_index_or_mode = instr->sampler_index;
|
|
|
|
desc.index = instr->texture_index;
|
|
|
|
} else {
|
|
|
|
enum bifrost_index mode = 0;
|
|
|
|
|
|
|
|
if (direct && instr->sampler_index == instr->texture_index) {
|
|
|
|
mode = BIFROST_INDEX_IMMEDIATE_SHARED;
|
|
|
|
desc.index = instr->texture_index;
|
|
|
|
} else if (direct) {
|
|
|
|
mode = BIFROST_INDEX_IMMEDIATE_SAMPLER;
|
|
|
|
desc.index = instr->sampler_index;
|
|
|
|
dregs[BIFROST_TEX_DREG_TEXTURE] = bi_mov_i32(b,
|
|
|
|
bi_imm_u32(instr->texture_index));
|
|
|
|
} else if (direct_tex) {
|
|
|
|
assert(!direct_samp);
|
|
|
|
mode = BIFROST_INDEX_IMMEDIATE_TEXTURE;
|
|
|
|
desc.index = instr->texture_index;
|
|
|
|
} else if (direct_samp) {
|
|
|
|
assert(!direct_tex);
|
|
|
|
mode = BIFROST_INDEX_IMMEDIATE_SAMPLER;
|
|
|
|
desc.index = instr->sampler_index;
|
|
|
|
} else {
|
|
|
|
mode = BIFROST_INDEX_REGISTER;
|
|
|
|
}
|
|
|
|
|
|
|
|
desc.sampler_index_or_mode = mode | (0x3 << 2);
|
|
|
|
}
|
|
|
|
|
2020-12-16 18:02:03 +00:00
|
|
|
/* Allocate staging registers contiguously by compacting the array.
|
|
|
|
* Index is not SSA (tied operands) */
|
|
|
|
|
|
|
|
unsigned sr_count = 0;
|
|
|
|
|
|
|
|
for (unsigned i = 0; i < ARRAY_SIZE(dregs); ++i) {
|
|
|
|
if (!bi_is_null(dregs[i]))
|
|
|
|
dregs[sr_count++] = dregs[i];
|
|
|
|
}
|
|
|
|
|
2021-05-05 20:48:35 +01:00
|
|
|
bi_index idx = sr_count ? bi_temp_reg(b->shader) : bi_null();
|
|
|
|
|
2020-12-16 18:02:03 +00:00
|
|
|
if (sr_count)
|
|
|
|
bi_make_vec_to(b, idx, dregs, NULL, sr_count, 32);
|
|
|
|
|
|
|
|
uint32_t desc_u = 0;
|
|
|
|
memcpy(&desc_u, &desc, sizeof(desc_u));
|
2021-05-05 20:48:35 +01:00
|
|
|
bi_texc_to(b, sr_count ? idx : bi_dest_index(&instr->dest),
|
|
|
|
idx, cx, cy, bi_imm_u32(desc_u), sr_count);
|
2020-12-16 18:02:03 +00:00
|
|
|
|
|
|
|
/* Explicit copy to facilitate tied operands */
|
2021-05-05 20:48:35 +01:00
|
|
|
if (sr_count) {
|
|
|
|
bi_index srcs[4] = { idx, idx, idx, idx };
|
|
|
|
unsigned channels[4] = { 0, 1, 2, 3 };
|
|
|
|
bi_make_vec_to(b, bi_dest_index(&instr->dest), srcs, channels, 4, 32);
|
|
|
|
}
|
2020-12-16 18:02:03 +00:00
|
|
|
}
|
|
|
|
|
2020-12-31 19:34:50 +00:00
|
|
|
/* Simple textures ops correspond to NIR tex or txl with LOD = 0 on 2D/cube
|
|
|
|
* textures with sufficiently small immediate indices. Anything else
|
2021-01-11 18:52:53 +00:00
|
|
|
* needs a complete texture op. */
|
2020-05-27 16:41:42 +01:00
|
|
|
|
2020-12-31 19:34:50 +00:00
|
|
|
static void
|
|
|
|
bi_emit_texs(bi_builder *b, nir_tex_instr *instr)
|
|
|
|
{
|
|
|
|
int coord_idx = nir_tex_instr_src_index(instr, nir_tex_src_coord);
|
|
|
|
assert(coord_idx >= 0);
|
|
|
|
bi_index coords = bi_src_index(&instr->src[coord_idx].src);
|
|
|
|
|
|
|
|
if (instr->sampler_dim == GLSL_SAMPLER_DIM_CUBE) {
|
|
|
|
bi_index face, s, t;
|
|
|
|
bi_emit_cube_coord(b, coords, &face, &s, &t);
|
|
|
|
|
|
|
|
bi_texs_cube_to(b, nir_dest_bit_size(instr->dest),
|
|
|
|
bi_dest_index(&instr->dest),
|
|
|
|
s, t, face,
|
|
|
|
instr->sampler_index, instr->texture_index);
|
|
|
|
} else {
|
|
|
|
bi_texs_2d_to(b, nir_dest_bit_size(instr->dest),
|
|
|
|
bi_dest_index(&instr->dest),
|
|
|
|
coords, bi_word(coords, 1),
|
|
|
|
instr->op != nir_texop_tex, /* zero LOD */
|
|
|
|
instr->sampler_index, instr->texture_index);
|
|
|
|
}
|
|
|
|
}
|
|
|
|
|
2020-05-27 16:41:42 +01:00
|
|
|
static bool
|
2020-12-16 18:02:22 +00:00
|
|
|
bi_is_simple_tex(nir_tex_instr *instr)
|
2020-05-27 16:41:42 +01:00
|
|
|
{
|
2020-11-12 17:18:57 +00:00
|
|
|
if (instr->op != nir_texop_tex && instr->op != nir_texop_txl)
|
2020-05-27 16:41:42 +01:00
|
|
|
return false;
|
|
|
|
|
2020-12-08 14:15:35 +00:00
|
|
|
if (instr->dest_type != nir_type_float32 &&
|
|
|
|
instr->dest_type != nir_type_float16)
|
2020-12-31 19:34:50 +00:00
|
|
|
return false;
|
|
|
|
|
|
|
|
if (instr->is_shadow || instr->is_array)
|
|
|
|
return false;
|
|
|
|
|
|
|
|
switch (instr->sampler_dim) {
|
|
|
|
case GLSL_SAMPLER_DIM_2D:
|
|
|
|
case GLSL_SAMPLER_DIM_EXTERNAL:
|
2021-05-04 22:42:13 +01:00
|
|
|
case GLSL_SAMPLER_DIM_RECT:
|
2020-12-31 19:34:50 +00:00
|
|
|
break;
|
|
|
|
|
|
|
|
case GLSL_SAMPLER_DIM_CUBE:
|
|
|
|
/* LOD can't be specified with TEXS_CUBE */
|
|
|
|
if (instr->op == nir_texop_txl)
|
|
|
|
return false;
|
|
|
|
break;
|
|
|
|
|
|
|
|
default:
|
|
|
|
return false;
|
|
|
|
}
|
|
|
|
|
2020-11-12 17:18:57 +00:00
|
|
|
for (unsigned i = 0; i < instr->num_srcs; ++i) {
|
|
|
|
if (instr->src[i].src_type != nir_tex_src_lod &&
|
|
|
|
instr->src[i].src_type != nir_tex_src_coord)
|
|
|
|
return false;
|
|
|
|
}
|
|
|
|
|
2020-12-31 19:34:50 +00:00
|
|
|
/* Indices need to fit in provided bits */
|
|
|
|
unsigned idx_bits = instr->sampler_dim == GLSL_SAMPLER_DIM_CUBE ? 2 : 3;
|
|
|
|
if (MAX2(instr->sampler_index, instr->texture_index) >= (1 << idx_bits))
|
2021-01-11 18:52:53 +00:00
|
|
|
return false;
|
|
|
|
|
2020-10-06 15:24:56 +01:00
|
|
|
int lod_idx = nir_tex_instr_src_index(instr, nir_tex_src_lod);
|
|
|
|
if (lod_idx < 0)
|
|
|
|
return true;
|
2020-05-27 16:41:42 +01:00
|
|
|
|
2020-10-06 15:24:56 +01:00
|
|
|
nir_src lod = instr->src[lod_idx].src;
|
|
|
|
return nir_src_is_const(lod) && nir_src_as_uint(lod) == 0;
|
2020-05-27 16:41:42 +01:00
|
|
|
}
|
|
|
|
|
2020-12-16 18:03:04 +00:00
|
|
|
static void
|
|
|
|
bi_emit_tex(bi_builder *b, nir_tex_instr *instr)
|
|
|
|
{
|
|
|
|
switch (instr->op) {
|
|
|
|
case nir_texop_txs:
|
2021-01-29 23:16:58 +00:00
|
|
|
bi_load_sysval_to(b, bi_dest_index(&instr->dest),
|
|
|
|
panfrost_sysval_for_instr(&instr->instr, NULL),
|
|
|
|
4, 0);
|
2020-12-16 18:03:04 +00:00
|
|
|
return;
|
|
|
|
case nir_texop_tex:
|
|
|
|
case nir_texop_txl:
|
|
|
|
case nir_texop_txb:
|
|
|
|
case nir_texop_txf:
|
|
|
|
case nir_texop_txf_ms:
|
2021-01-27 16:22:10 +00:00
|
|
|
case nir_texop_tg4:
|
2020-12-16 18:03:04 +00:00
|
|
|
break;
|
|
|
|
default:
|
|
|
|
unreachable("Invalid texture operation");
|
|
|
|
}
|
|
|
|
|
2020-12-31 19:34:50 +00:00
|
|
|
if (bi_is_simple_tex(instr))
|
2020-12-16 18:03:04 +00:00
|
|
|
bi_emit_texs(b, instr);
|
|
|
|
else
|
|
|
|
bi_emit_texc(b, instr);
|
|
|
|
}
|
2020-04-21 17:15:29 +01:00
|
|
|
|
2020-03-05 22:10:46 +00:00
|
|
|
static void
|
2020-12-16 18:35:51 +00:00
|
|
|
bi_emit_instr(bi_builder *b, struct nir_instr *instr)
|
|
|
|
{
|
|
|
|
switch (instr->type) {
|
|
|
|
case nir_instr_type_load_const:
|
|
|
|
bi_emit_load_const(b, nir_instr_as_load_const(instr));
|
|
|
|
break;
|
|
|
|
|
|
|
|
case nir_instr_type_intrinsic:
|
|
|
|
bi_emit_intrinsic(b, nir_instr_as_intrinsic(instr));
|
|
|
|
break;
|
|
|
|
|
|
|
|
case nir_instr_type_alu:
|
|
|
|
bi_emit_alu(b, nir_instr_as_alu(instr));
|
|
|
|
break;
|
|
|
|
|
|
|
|
case nir_instr_type_tex:
|
|
|
|
bi_emit_tex(b, nir_instr_as_tex(instr));
|
|
|
|
break;
|
|
|
|
|
|
|
|
case nir_instr_type_jump:
|
|
|
|
bi_emit_jump(b, nir_instr_as_jump(instr));
|
|
|
|
break;
|
|
|
|
|
|
|
|
default:
|
|
|
|
unreachable("should've been lowered");
|
|
|
|
}
|
|
|
|
}
|
2020-03-05 15:25:19 +00:00
|
|
|
|
|
|
|
static bi_block *
|
|
|
|
create_empty_block(bi_context *ctx)
|
|
|
|
{
|
|
|
|
bi_block *blk = rzalloc(ctx, bi_block);
|
|
|
|
|
2020-03-11 18:35:38 +00:00
|
|
|
blk->base.predecessors = _mesa_set_create(blk,
|
2020-03-05 15:25:19 +00:00
|
|
|
_mesa_hash_pointer,
|
|
|
|
_mesa_key_pointer_equal);
|
|
|
|
|
|
|
|
return blk;
|
|
|
|
}
|
|
|
|
|
|
|
|
static bi_block *
|
|
|
|
emit_block(bi_context *ctx, nir_block *block)
|
|
|
|
{
|
2020-03-05 21:45:16 +00:00
|
|
|
if (ctx->after_block) {
|
|
|
|
ctx->current_block = ctx->after_block;
|
|
|
|
ctx->after_block = NULL;
|
|
|
|
} else {
|
|
|
|
ctx->current_block = create_empty_block(ctx);
|
|
|
|
}
|
|
|
|
|
2020-03-11 18:35:38 +00:00
|
|
|
list_addtail(&ctx->current_block->base.link, &ctx->blocks);
|
|
|
|
list_inithead(&ctx->current_block->base.instructions);
|
2020-03-05 15:25:19 +00:00
|
|
|
|
2021-01-05 18:26:56 +00:00
|
|
|
bi_builder _b = bi_init_builder(ctx, bi_after_block(ctx->current_block));
|
2020-12-30 20:50:50 +00:00
|
|
|
|
2020-03-05 15:25:19 +00:00
|
|
|
nir_foreach_instr(instr, block) {
|
2020-12-30 20:50:50 +00:00
|
|
|
bi_emit_instr(&_b, instr);
|
2020-03-05 15:25:19 +00:00
|
|
|
++ctx->instruction_count;
|
|
|
|
}
|
|
|
|
|
|
|
|
return ctx->current_block;
|
|
|
|
}
|
|
|
|
|
2020-03-05 21:45:16 +00:00
|
|
|
static void
|
|
|
|
emit_if(bi_context *ctx, nir_if *nif)
|
|
|
|
{
|
|
|
|
bi_block *before_block = ctx->current_block;
|
|
|
|
|
|
|
|
/* Speculatively emit the branch, but we can't fill it in until later */
|
2021-01-05 18:26:56 +00:00
|
|
|
bi_builder _b = bi_init_builder(ctx, bi_after_block(ctx->current_block));
|
2021-05-05 17:51:30 +01:00
|
|
|
bi_instr *then_branch = bi_branchz_i16(&_b,
|
|
|
|
bi_half(bi_src_index(&nif->condition), false),
|
|
|
|
bi_zero(), BI_CMPF_EQ);
|
2020-03-05 21:45:16 +00:00
|
|
|
|
|
|
|
/* Emit the two subblocks. */
|
|
|
|
bi_block *then_block = emit_cf_list(ctx, &nif->then_list);
|
|
|
|
bi_block *end_then_block = ctx->current_block;
|
|
|
|
|
|
|
|
/* Emit second block, and check if it's empty */
|
|
|
|
|
|
|
|
int count_in = ctx->instruction_count;
|
|
|
|
bi_block *else_block = emit_cf_list(ctx, &nif->else_list);
|
|
|
|
bi_block *end_else_block = ctx->current_block;
|
|
|
|
ctx->after_block = create_empty_block(ctx);
|
|
|
|
|
|
|
|
/* Now that we have the subblocks emitted, fix up the branches */
|
|
|
|
|
|
|
|
assert(then_block);
|
|
|
|
assert(else_block);
|
|
|
|
|
|
|
|
if (ctx->instruction_count == count_in) {
|
2020-05-01 23:13:54 +01:00
|
|
|
then_branch->branch_target = ctx->after_block;
|
2020-05-27 23:27:08 +01:00
|
|
|
pan_block_add_successor(&end_then_block->base, &ctx->after_block->base); /* fallthrough */
|
2020-03-05 21:45:16 +00:00
|
|
|
} else {
|
2020-05-01 23:13:54 +01:00
|
|
|
then_branch->branch_target = else_block;
|
2020-12-30 20:50:50 +00:00
|
|
|
|
|
|
|
/* Emit a jump from the end of the then block to the end of the else */
|
|
|
|
_b.cursor = bi_after_block(end_then_block);
|
2021-02-18 19:58:50 +00:00
|
|
|
bi_instr *then_exit = bi_jump(&_b, bi_zero());
|
|
|
|
then_exit->branch_target = ctx->after_block;
|
2020-12-30 20:50:50 +00:00
|
|
|
|
2020-05-01 23:13:54 +01:00
|
|
|
pan_block_add_successor(&end_then_block->base, &then_exit->branch_target->base);
|
2020-05-27 23:27:08 +01:00
|
|
|
pan_block_add_successor(&end_else_block->base, &ctx->after_block->base); /* fallthrough */
|
2020-03-05 21:45:16 +00:00
|
|
|
}
|
|
|
|
|
2020-05-01 23:13:54 +01:00
|
|
|
pan_block_add_successor(&before_block->base, &then_branch->branch_target->base); /* then_branch */
|
2020-03-11 18:35:38 +00:00
|
|
|
pan_block_add_successor(&before_block->base, &then_block->base); /* fallthrough */
|
2020-03-05 21:45:16 +00:00
|
|
|
}
|
|
|
|
|
2020-03-05 22:03:53 +00:00
|
|
|
static void
|
|
|
|
emit_loop(bi_context *ctx, nir_loop *nloop)
|
|
|
|
{
|
|
|
|
/* Remember where we are */
|
|
|
|
bi_block *start_block = ctx->current_block;
|
|
|
|
|
|
|
|
bi_block *saved_break = ctx->break_block;
|
|
|
|
bi_block *saved_continue = ctx->continue_block;
|
|
|
|
|
|
|
|
ctx->continue_block = create_empty_block(ctx);
|
|
|
|
ctx->break_block = create_empty_block(ctx);
|
|
|
|
ctx->after_block = ctx->continue_block;
|
|
|
|
|
|
|
|
/* Emit the body itself */
|
|
|
|
emit_cf_list(ctx, &nloop->body);
|
|
|
|
|
|
|
|
/* Branch back to loop back */
|
2021-01-05 18:26:56 +00:00
|
|
|
bi_builder _b = bi_init_builder(ctx, bi_after_block(ctx->current_block));
|
2021-02-18 19:58:50 +00:00
|
|
|
bi_instr *I = bi_jump(&_b, bi_zero());
|
|
|
|
I->branch_target = ctx->continue_block;
|
2020-03-11 18:35:38 +00:00
|
|
|
pan_block_add_successor(&start_block->base, &ctx->continue_block->base);
|
|
|
|
pan_block_add_successor(&ctx->current_block->base, &ctx->continue_block->base);
|
2020-03-05 22:03:53 +00:00
|
|
|
|
|
|
|
ctx->after_block = ctx->break_block;
|
|
|
|
|
|
|
|
/* Pop off */
|
|
|
|
ctx->break_block = saved_break;
|
|
|
|
ctx->continue_block = saved_continue;
|
|
|
|
++ctx->loop_count;
|
|
|
|
}
|
|
|
|
|
2020-03-05 15:25:19 +00:00
|
|
|
static bi_block *
|
|
|
|
emit_cf_list(bi_context *ctx, struct exec_list *list)
|
|
|
|
{
|
|
|
|
bi_block *start_block = NULL;
|
|
|
|
|
|
|
|
foreach_list_typed(nir_cf_node, node, node, list) {
|
|
|
|
switch (node->type) {
|
|
|
|
case nir_cf_node_block: {
|
|
|
|
bi_block *block = emit_block(ctx, nir_cf_node_as_block(node));
|
|
|
|
|
|
|
|
if (!start_block)
|
|
|
|
start_block = block;
|
|
|
|
|
|
|
|
break;
|
|
|
|
}
|
|
|
|
|
|
|
|
case nir_cf_node_if:
|
|
|
|
emit_if(ctx, nir_cf_node_as_if(node));
|
|
|
|
break;
|
|
|
|
|
|
|
|
case nir_cf_node_loop:
|
|
|
|
emit_loop(ctx, nir_cf_node_as_loop(node));
|
|
|
|
break;
|
|
|
|
|
|
|
|
default:
|
|
|
|
unreachable("Unknown control flow");
|
|
|
|
}
|
|
|
|
}
|
|
|
|
|
|
|
|
return start_block;
|
|
|
|
}
|
2020-03-03 00:47:11 +00:00
|
|
|
|
2020-11-13 22:34:25 +00:00
|
|
|
/* shader-db stuff */
|
|
|
|
|
|
|
|
static void
|
2021-02-06 01:25:39 +00:00
|
|
|
bi_print_stats(bi_context *ctx, unsigned size, FILE *fp)
|
2020-11-13 22:34:25 +00:00
|
|
|
{
|
|
|
|
unsigned nr_clauses = 0, nr_tuples = 0, nr_ins = 0;
|
|
|
|
|
|
|
|
/* Count instructions, clauses, and tuples */
|
|
|
|
bi_foreach_block(ctx, _block) {
|
|
|
|
bi_block *block = (bi_block *) _block;
|
|
|
|
|
|
|
|
bi_foreach_clause_in_block(block, clause) {
|
|
|
|
nr_clauses++;
|
2020-12-22 20:36:20 +00:00
|
|
|
nr_tuples += clause->tuple_count;
|
2020-11-13 22:34:25 +00:00
|
|
|
|
2020-12-22 20:36:20 +00:00
|
|
|
for (unsigned i = 0; i < clause->tuple_count; ++i) {
|
|
|
|
if (clause->tuples[i].fma)
|
2020-11-13 22:34:25 +00:00
|
|
|
nr_ins++;
|
|
|
|
|
2020-12-22 20:36:20 +00:00
|
|
|
if (clause->tuples[i].add)
|
2020-11-13 22:34:25 +00:00
|
|
|
nr_ins++;
|
|
|
|
}
|
|
|
|
}
|
|
|
|
}
|
|
|
|
|
|
|
|
/* tuples = ((# of instructions) + (# of nops)) / 2 */
|
|
|
|
unsigned nr_nops = (2 * nr_tuples) - nr_ins;
|
|
|
|
|
|
|
|
/* In the future, we'll calculate thread count for v7. For now we
|
|
|
|
* always use fewer threads than we should (v6 style) due to missing
|
|
|
|
* piping, TODO: fix that for a nice perf win */
|
|
|
|
unsigned nr_threads = 1;
|
|
|
|
|
|
|
|
/* Dump stats */
|
|
|
|
|
2021-02-11 00:40:38 +00:00
|
|
|
fprintf(stderr, "%s - %s shader: "
|
2020-11-13 22:34:25 +00:00
|
|
|
"%u inst, %u nops, %u clauses, "
|
2021-02-06 01:25:39 +00:00
|
|
|
"%u quadwords, %u threads, %u loops, "
|
2020-11-13 22:34:25 +00:00
|
|
|
"%u:%u spills:fills\n",
|
2020-12-17 01:39:58 +00:00
|
|
|
ctx->nir->info.label ?: "",
|
2021-02-12 17:11:06 +00:00
|
|
|
ctx->inputs->is_blend ? "PAN_SHADER_BLEND" :
|
2020-11-13 22:34:25 +00:00
|
|
|
gl_shader_stage_name(ctx->stage),
|
|
|
|
nr_ins, nr_nops, nr_clauses,
|
2021-02-06 01:25:39 +00:00
|
|
|
size / 16, nr_threads,
|
2020-11-13 22:34:25 +00:00
|
|
|
ctx->loop_count,
|
|
|
|
ctx->spills, ctx->fills);
|
|
|
|
}
|
|
|
|
|
2020-03-05 15:11:39 +00:00
|
|
|
static int
|
|
|
|
glsl_type_size(const struct glsl_type *type, bool bindless)
|
|
|
|
{
|
|
|
|
return glsl_count_attribute_slots(type, false);
|
|
|
|
}
|
|
|
|
|
2021-04-03 20:00:59 +01:00
|
|
|
/* Split stores to memory. We don't split stores to vertex outputs, since
|
|
|
|
* nir_lower_io_to_temporaries will ensure there's only a single write.
|
|
|
|
*/
|
|
|
|
|
|
|
|
static bool
|
|
|
|
should_split_wrmask(const nir_instr *instr, UNUSED const void *data)
|
|
|
|
{
|
|
|
|
nir_intrinsic_instr *intr = nir_instr_as_intrinsic(instr);
|
|
|
|
|
|
|
|
switch (intr->intrinsic) {
|
|
|
|
case nir_intrinsic_store_ssbo:
|
|
|
|
case nir_intrinsic_store_shared:
|
|
|
|
case nir_intrinsic_store_global:
|
|
|
|
case nir_intrinsic_store_scratch:
|
|
|
|
return true;
|
|
|
|
default:
|
|
|
|
return false;
|
|
|
|
}
|
|
|
|
}
|
|
|
|
|
2021-05-04 00:39:59 +01:00
|
|
|
/* Bifrost wants transcendentals as FP32 */
|
|
|
|
|
|
|
|
static unsigned
|
|
|
|
bi_lower_bit_size(const nir_instr *instr, UNUSED void *data)
|
|
|
|
{
|
|
|
|
if (instr->type != nir_instr_type_alu)
|
|
|
|
return 0;
|
|
|
|
|
|
|
|
nir_alu_instr *alu = nir_instr_as_alu(instr);
|
|
|
|
|
|
|
|
switch (alu->op) {
|
|
|
|
case nir_op_fexp2:
|
|
|
|
case nir_op_flog2:
|
|
|
|
case nir_op_fsin:
|
|
|
|
case nir_op_fcos:
|
|
|
|
return (nir_dest_bit_size(alu->dest.dest) == 32) ? 0 : 32;
|
|
|
|
default:
|
|
|
|
return 0;
|
|
|
|
}
|
|
|
|
}
|
|
|
|
|
2020-12-31 20:00:24 +00:00
|
|
|
/* Although Bifrost generally supports packed 16-bit vec2 and 8-bit vec4,
|
|
|
|
* transcendentals are an exception. Also shifts because of lane size mismatch
|
|
|
|
* (8-bit in Bifrost, 32-bit in NIR TODO - workaround!). Some conversions need
|
|
|
|
* to be scalarized due to type size. */
|
|
|
|
|
|
|
|
static bool
|
|
|
|
bi_vectorize_filter(const nir_instr *instr, void *data)
|
|
|
|
{
|
|
|
|
/* Defaults work for everything else */
|
|
|
|
if (instr->type != nir_instr_type_alu)
|
|
|
|
return true;
|
|
|
|
|
|
|
|
const nir_alu_instr *alu = nir_instr_as_alu(instr);
|
|
|
|
|
|
|
|
switch (alu->op) {
|
|
|
|
case nir_op_frcp:
|
|
|
|
case nir_op_frsq:
|
|
|
|
case nir_op_ishl:
|
|
|
|
case nir_op_ishr:
|
|
|
|
case nir_op_ushr:
|
|
|
|
case nir_op_f2i16:
|
|
|
|
case nir_op_f2u16:
|
2021-04-13 18:22:18 +01:00
|
|
|
case nir_op_i2f16:
|
|
|
|
case nir_op_u2f16:
|
2020-12-31 20:00:24 +00:00
|
|
|
return false;
|
|
|
|
default:
|
|
|
|
return true;
|
|
|
|
}
|
|
|
|
}
|
|
|
|
|
2021-04-03 21:06:04 +01:00
|
|
|
/* XXX: This is a kludge to workaround NIR's lack of divergence metadata. If we
|
|
|
|
* keep divergence info around after we consume it for indirect lowering,
|
|
|
|
* nir_convert_from_ssa will regress code quality since it will avoid
|
|
|
|
* coalescing divergent with non-divergent nodes. */
|
|
|
|
|
|
|
|
static bool
|
|
|
|
nir_invalidate_divergence_ssa(nir_ssa_def *ssa, UNUSED void *data)
|
|
|
|
{
|
|
|
|
ssa->divergent = false;
|
|
|
|
return true;
|
|
|
|
}
|
|
|
|
|
|
|
|
static bool
|
|
|
|
nir_invalidate_divergence(struct nir_builder *b, nir_instr *instr,
|
|
|
|
UNUSED void *data)
|
|
|
|
{
|
|
|
|
return nir_foreach_ssa_def(instr, nir_invalidate_divergence_ssa, NULL);
|
|
|
|
}
|
|
|
|
|
2020-03-05 15:11:39 +00:00
|
|
|
static void
|
2021-04-03 21:06:04 +01:00
|
|
|
bi_optimize_nir(nir_shader *nir, unsigned gpu_id, bool is_blend)
|
2020-03-05 15:11:39 +00:00
|
|
|
{
|
|
|
|
bool progress;
|
|
|
|
unsigned lower_flrp = 16 | 32 | 64;
|
|
|
|
|
|
|
|
NIR_PASS(progress, nir, nir_lower_regs_to_ssa);
|
|
|
|
|
|
|
|
nir_lower_tex_options lower_tex_options = {
|
|
|
|
.lower_txs_lod = true,
|
|
|
|
.lower_txp = ~0,
|
|
|
|
.lower_tex_without_implicit_lod = true,
|
2021-01-27 16:22:10 +00:00
|
|
|
.lower_tg4_broadcom_swizzle = true,
|
2020-03-05 15:11:39 +00:00
|
|
|
.lower_txd = true,
|
|
|
|
};
|
|
|
|
|
2021-01-05 05:15:47 +00:00
|
|
|
NIR_PASS(progress, nir, pan_nir_lower_64bit_intrin);
|
2021-01-27 18:32:56 +00:00
|
|
|
NIR_PASS(progress, nir, pan_lower_helper_invocation);
|
2021-01-05 05:15:47 +00:00
|
|
|
|
2020-12-22 10:20:10 +00:00
|
|
|
NIR_PASS(progress, nir, nir_lower_int64);
|
|
|
|
|
2021-04-07 19:17:46 +01:00
|
|
|
nir_lower_idiv_options idiv_options = {
|
|
|
|
.imprecise_32bit_lowering = true,
|
|
|
|
.allow_fp16 = true,
|
|
|
|
};
|
|
|
|
NIR_PASS(progress, nir, nir_lower_idiv, &idiv_options);
|
2021-01-25 16:09:25 +00:00
|
|
|
|
2020-03-05 15:11:39 +00:00
|
|
|
NIR_PASS(progress, nir, nir_lower_tex, &lower_tex_options);
|
2020-03-06 21:29:35 +00:00
|
|
|
NIR_PASS(progress, nir, nir_lower_alu_to_scalar, NULL, NULL);
|
|
|
|
NIR_PASS(progress, nir, nir_lower_load_const_to_scalar);
|
2020-03-05 15:11:39 +00:00
|
|
|
|
|
|
|
do {
|
|
|
|
progress = false;
|
|
|
|
|
|
|
|
NIR_PASS(progress, nir, nir_lower_var_copies);
|
|
|
|
NIR_PASS(progress, nir, nir_lower_vars_to_ssa);
|
2021-04-03 20:00:59 +01:00
|
|
|
NIR_PASS(progress, nir, nir_lower_wrmasks, should_split_wrmask, NULL);
|
2020-03-05 15:11:39 +00:00
|
|
|
|
|
|
|
NIR_PASS(progress, nir, nir_copy_prop);
|
|
|
|
NIR_PASS(progress, nir, nir_opt_remove_phis);
|
|
|
|
NIR_PASS(progress, nir, nir_opt_dce);
|
|
|
|
NIR_PASS(progress, nir, nir_opt_dead_cf);
|
|
|
|
NIR_PASS(progress, nir, nir_opt_cse);
|
|
|
|
NIR_PASS(progress, nir, nir_opt_peephole_select, 64, false, true);
|
|
|
|
NIR_PASS(progress, nir, nir_opt_algebraic);
|
|
|
|
NIR_PASS(progress, nir, nir_opt_constant_folding);
|
|
|
|
|
2020-12-28 22:21:27 +00:00
|
|
|
NIR_PASS(progress, nir, nir_lower_alu);
|
|
|
|
|
2020-03-05 15:11:39 +00:00
|
|
|
if (lower_flrp != 0) {
|
|
|
|
bool lower_flrp_progress = false;
|
|
|
|
NIR_PASS(lower_flrp_progress,
|
|
|
|
nir,
|
|
|
|
nir_lower_flrp,
|
|
|
|
lower_flrp,
|
2020-07-23 03:13:16 +01:00
|
|
|
false /* always_precise */);
|
2020-03-05 15:11:39 +00:00
|
|
|
if (lower_flrp_progress) {
|
|
|
|
NIR_PASS(progress, nir,
|
|
|
|
nir_opt_constant_folding);
|
|
|
|
progress = true;
|
|
|
|
}
|
|
|
|
|
|
|
|
/* Nothing should rematerialize any flrps, so we only
|
|
|
|
* need to do this lowering once.
|
|
|
|
*/
|
|
|
|
lower_flrp = 0;
|
|
|
|
}
|
|
|
|
|
|
|
|
NIR_PASS(progress, nir, nir_opt_undef);
|
2020-06-02 16:30:35 +01:00
|
|
|
NIR_PASS(progress, nir, nir_lower_undef_to_zero);
|
2020-10-15 13:47:40 +01:00
|
|
|
|
2020-03-05 15:11:39 +00:00
|
|
|
NIR_PASS(progress, nir, nir_opt_loop_unroll,
|
|
|
|
nir_var_shader_in |
|
|
|
|
nir_var_shader_out |
|
|
|
|
nir_var_function_temp);
|
|
|
|
} while (progress);
|
|
|
|
|
2020-12-22 23:33:56 +00:00
|
|
|
/* We need to cleanup after each iteration of late algebraic
|
|
|
|
* optimizations, since otherwise NIR can produce weird edge cases
|
|
|
|
* (like fneg of a constant) which we don't handle */
|
|
|
|
bool late_algebraic = true;
|
|
|
|
while (late_algebraic) {
|
|
|
|
late_algebraic = false;
|
|
|
|
NIR_PASS(late_algebraic, nir, nir_opt_algebraic_late);
|
|
|
|
NIR_PASS(progress, nir, nir_opt_constant_folding);
|
|
|
|
NIR_PASS(progress, nir, nir_copy_prop);
|
|
|
|
NIR_PASS(progress, nir, nir_opt_dce);
|
|
|
|
NIR_PASS(progress, nir, nir_opt_cse);
|
|
|
|
}
|
|
|
|
|
2020-03-06 21:29:35 +00:00
|
|
|
NIR_PASS(progress, nir, nir_lower_alu_to_scalar, NULL, NULL);
|
2020-12-31 20:00:24 +00:00
|
|
|
NIR_PASS(progress, nir, nir_opt_vectorize, bi_vectorize_filter, NULL);
|
|
|
|
NIR_PASS(progress, nir, nir_lower_load_const_to_scalar);
|
|
|
|
NIR_PASS(progress, nir, nir_opt_dce);
|
2021-02-18 20:56:43 +00:00
|
|
|
|
|
|
|
/* Backend scheduler is purely local, so do some global optimizations
|
2021-05-25 18:52:56 +01:00
|
|
|
* to reduce register pressure. */
|
2021-03-02 16:54:33 +00:00
|
|
|
nir_move_options move_all =
|
|
|
|
nir_move_const_undef | nir_move_load_ubo | nir_move_load_input |
|
|
|
|
nir_move_comparisons | nir_move_copies | nir_move_load_ssbo;
|
|
|
|
|
2021-05-25 18:52:56 +01:00
|
|
|
NIR_PASS_V(nir, nir_opt_sink, move_all);
|
|
|
|
NIR_PASS_V(nir, nir_opt_move, move_all);
|
2021-02-18 20:56:43 +00:00
|
|
|
|
2021-04-03 21:06:04 +01:00
|
|
|
/* We might lower attribute, varying, and image indirects. Use the
|
|
|
|
* gathered info to skip the extra analysis in the happy path. */
|
|
|
|
bool any_indirects =
|
|
|
|
nir->info.inputs_read_indirectly ||
|
|
|
|
nir->info.outputs_accessed_indirectly ||
|
|
|
|
nir->info.patch_inputs_read_indirectly ||
|
|
|
|
nir->info.patch_outputs_accessed_indirectly ||
|
|
|
|
nir->info.images_used;
|
|
|
|
|
|
|
|
if (any_indirects) {
|
|
|
|
nir_convert_to_lcssa(nir, true, true);
|
|
|
|
NIR_PASS_V(nir, nir_divergence_analysis);
|
|
|
|
NIR_PASS_V(nir, bi_lower_divergent_indirects,
|
|
|
|
bifrost_lanes_per_warp(gpu_id));
|
|
|
|
NIR_PASS_V(nir, nir_shader_instructions_pass,
|
|
|
|
nir_invalidate_divergence, nir_metadata_all, NULL);
|
|
|
|
}
|
|
|
|
|
2020-03-05 15:11:39 +00:00
|
|
|
/* Take us out of SSA */
|
|
|
|
NIR_PASS(progress, nir, nir_lower_locals_to_regs);
|
2020-03-09 23:56:35 +00:00
|
|
|
NIR_PASS(progress, nir, nir_move_vec_src_uses_to_dest);
|
2020-03-22 21:31:23 +00:00
|
|
|
NIR_PASS(progress, nir, nir_convert_from_ssa, true);
|
2020-03-19 21:21:49 +00:00
|
|
|
}
|
|
|
|
|
2020-11-20 15:57:46 +00:00
|
|
|
/* The cmdstream lowers 8-bit fragment output as 16-bit, so we need to do the
|
|
|
|
* same lowering here to zero-extend correctly */
|
|
|
|
|
|
|
|
static bool
|
|
|
|
bifrost_nir_lower_i8_fragout_impl(struct nir_builder *b,
|
2021-05-03 17:12:21 +01:00
|
|
|
nir_intrinsic_instr *intr, UNUSED void *data)
|
2020-11-20 15:57:46 +00:00
|
|
|
{
|
|
|
|
if (nir_src_bit_size(intr->src[0]) != 8)
|
|
|
|
return false;
|
|
|
|
|
|
|
|
nir_alu_type type =
|
|
|
|
nir_alu_type_get_base_type(nir_intrinsic_src_type(intr));
|
|
|
|
|
2021-01-14 04:16:30 +00:00
|
|
|
assert(type == nir_type_int || type == nir_type_uint);
|
2020-11-20 15:57:46 +00:00
|
|
|
|
2021-05-03 17:12:21 +01:00
|
|
|
b->cursor = nir_before_instr(&intr->instr);
|
|
|
|
nir_ssa_def *cast = nir_convert_to_bit_size(b, intr->src[0].ssa, type, 16);
|
2020-11-20 15:57:46 +00:00
|
|
|
|
|
|
|
nir_intrinsic_set_src_type(intr, type | 16);
|
2021-05-03 17:12:21 +01:00
|
|
|
nir_instr_rewrite_src_ssa(&intr->instr, &intr->src[0], cast);
|
2020-11-20 15:57:46 +00:00
|
|
|
return true;
|
|
|
|
}
|
|
|
|
|
|
|
|
static bool
|
2021-05-03 17:12:21 +01:00
|
|
|
bifrost_nir_lower_i8_fragin_impl(struct nir_builder *b,
|
|
|
|
nir_intrinsic_instr *intr, UNUSED void *data)
|
2020-11-20 15:57:46 +00:00
|
|
|
{
|
2021-05-03 17:12:21 +01:00
|
|
|
if (nir_dest_bit_size(intr->dest) != 8)
|
2020-11-20 15:57:46 +00:00
|
|
|
return false;
|
|
|
|
|
2021-05-03 17:12:21 +01:00
|
|
|
nir_alu_type type =
|
|
|
|
nir_alu_type_get_base_type(nir_intrinsic_dest_type(intr));
|
|
|
|
|
|
|
|
assert(type == nir_type_int || type == nir_type_uint);
|
|
|
|
|
|
|
|
b->cursor = nir_before_instr(&intr->instr);
|
|
|
|
nir_ssa_def *out =
|
|
|
|
nir_load_output(b, intr->num_components, 16, intr->src[0].ssa,
|
|
|
|
.base = nir_intrinsic_base(intr),
|
|
|
|
.component = nir_intrinsic_component(intr),
|
|
|
|
.dest_type = type | 16,
|
|
|
|
.io_semantics = nir_intrinsic_io_semantics(intr));
|
|
|
|
|
|
|
|
nir_ssa_def *cast = nir_convert_to_bit_size(b, out, type, 8);
|
|
|
|
nir_ssa_def_rewrite_uses(&intr->dest.ssa, cast);
|
|
|
|
return true;
|
|
|
|
}
|
|
|
|
|
|
|
|
static bool
|
|
|
|
bifrost_nir_lower_i8_frag(struct nir_builder *b,
|
|
|
|
nir_instr *instr, UNUSED void *data)
|
|
|
|
{
|
|
|
|
if (instr->type != nir_instr_type_intrinsic)
|
|
|
|
return false;
|
|
|
|
|
|
|
|
nir_intrinsic_instr *intr = nir_instr_as_intrinsic(instr);
|
|
|
|
if (intr->intrinsic == nir_intrinsic_load_output)
|
|
|
|
return bifrost_nir_lower_i8_fragin_impl(b, intr, data);
|
|
|
|
else if (intr->intrinsic == nir_intrinsic_store_output)
|
|
|
|
return bifrost_nir_lower_i8_fragout_impl(b, intr, data);
|
|
|
|
else
|
|
|
|
return false;
|
2020-11-20 15:57:46 +00:00
|
|
|
}
|
|
|
|
|
2021-05-04 22:49:24 +01:00
|
|
|
static void
|
|
|
|
bi_opt_post_ra(bi_context *ctx)
|
|
|
|
{
|
|
|
|
bi_foreach_instr_global_safe(ctx, ins) {
|
|
|
|
if (ins->op == BI_OPCODE_MOV_I32 && bi_is_equiv(ins->dest[0], ins->src[0]))
|
|
|
|
bi_remove_instruction(ins);
|
|
|
|
}
|
|
|
|
}
|
|
|
|
|
2021-01-21 00:09:34 +00:00
|
|
|
/* Dead code elimination for branches at the end of a block - only one branch
|
2021-02-11 20:23:01 +00:00
|
|
|
* per block is legal semantically, but unreachable jumps can be generated.
|
|
|
|
* Likewise we can generate jumps to the terminal block which need to be
|
|
|
|
* lowered away to a jump to #0x0, which induces successful termination. */
|
2021-01-21 00:09:34 +00:00
|
|
|
|
|
|
|
static void
|
2021-02-11 20:23:01 +00:00
|
|
|
bi_lower_branch(bi_block *block)
|
2021-01-21 00:09:34 +00:00
|
|
|
{
|
|
|
|
bool branched = false;
|
|
|
|
ASSERTED bool was_jump = false;
|
|
|
|
|
|
|
|
bi_foreach_instr_in_block_safe(block, ins) {
|
|
|
|
if (!ins->branch_target) continue;
|
|
|
|
|
|
|
|
if (branched) {
|
2021-02-11 20:23:01 +00:00
|
|
|
assert(was_jump && (ins->op == BI_OPCODE_JUMP));
|
2021-01-21 00:09:34 +00:00
|
|
|
bi_remove_instruction(ins);
|
2021-02-18 18:54:30 +00:00
|
|
|
continue;
|
2021-01-21 00:09:34 +00:00
|
|
|
}
|
|
|
|
|
|
|
|
branched = true;
|
|
|
|
was_jump = ins->op == BI_OPCODE_JUMP;
|
2021-02-11 20:23:01 +00:00
|
|
|
|
|
|
|
if (bi_is_terminal_block(ins->branch_target))
|
|
|
|
ins->branch_target = NULL;
|
2021-01-21 00:09:34 +00:00
|
|
|
}
|
|
|
|
}
|
|
|
|
|
2021-02-13 07:24:03 +00:00
|
|
|
void
|
|
|
|
bifrost_compile_shader_nir(nir_shader *nir,
|
|
|
|
const struct panfrost_compile_inputs *inputs,
|
|
|
|
struct util_dynarray *binary,
|
|
|
|
struct pan_shader_info *info)
|
2020-03-03 00:47:11 +00:00
|
|
|
{
|
2020-04-30 08:29:10 +01:00
|
|
|
bifrost_debug = debug_get_option_bifrost_debug();
|
|
|
|
|
2020-03-03 00:47:11 +00:00
|
|
|
bi_context *ctx = rzalloc(NULL, bi_context);
|
2021-02-13 07:24:03 +00:00
|
|
|
ctx->sysval_to_id = panfrost_init_sysvals(&info->sysvals, ctx);
|
2021-01-29 22:16:07 +00:00
|
|
|
|
2021-02-12 17:11:06 +00:00
|
|
|
ctx->inputs = inputs;
|
2020-03-03 00:47:11 +00:00
|
|
|
ctx->nir = nir;
|
2021-02-13 07:24:03 +00:00
|
|
|
ctx->info = info;
|
2020-03-05 15:11:39 +00:00
|
|
|
ctx->stage = nir->info.stage;
|
2020-10-08 09:09:56 +01:00
|
|
|
ctx->quirks = bifrost_get_quirks(inputs->gpu_id);
|
2020-11-12 15:13:08 +00:00
|
|
|
ctx->arch = inputs->gpu_id >> 12;
|
2020-03-05 15:25:19 +00:00
|
|
|
list_inithead(&ctx->blocks);
|
2020-03-03 00:47:11 +00:00
|
|
|
|
2020-03-05 15:11:39 +00:00
|
|
|
/* Lower gl_Position pre-optimisation, but after lowering vars to ssa
|
|
|
|
* (so we don't accidentally duplicate the epilogue since mesa/st has
|
|
|
|
* messed with our I/O quite a bit already) */
|
|
|
|
|
|
|
|
NIR_PASS_V(nir, nir_lower_vars_to_ssa);
|
|
|
|
|
|
|
|
if (ctx->stage == MESA_SHADER_VERTEX) {
|
|
|
|
NIR_PASS_V(nir, nir_lower_viewport_transform);
|
|
|
|
NIR_PASS_V(nir, nir_lower_point_size, 1.0, 1024.0);
|
|
|
|
}
|
|
|
|
|
2021-03-05 21:55:18 +00:00
|
|
|
/* Lower large arrays to scratch and small arrays to bcsel (TODO: tune
|
|
|
|
* threshold, but not until addresses / csel is optimized better) */
|
|
|
|
NIR_PASS_V(nir, nir_lower_vars_to_scratch, nir_var_function_temp, 16,
|
|
|
|
glsl_get_natural_size_align_bytes);
|
|
|
|
NIR_PASS_V(nir, nir_lower_indirect_derefs, nir_var_function_temp, ~0);
|
|
|
|
|
2020-03-05 15:11:39 +00:00
|
|
|
NIR_PASS_V(nir, nir_split_var_copies);
|
|
|
|
NIR_PASS_V(nir, nir_lower_global_vars_to_local);
|
|
|
|
NIR_PASS_V(nir, nir_lower_var_copies);
|
|
|
|
NIR_PASS_V(nir, nir_lower_vars_to_ssa);
|
2020-06-10 23:54:25 +01:00
|
|
|
NIR_PASS_V(nir, nir_lower_io, nir_var_shader_in | nir_var_shader_out,
|
|
|
|
glsl_type_size, 0);
|
2021-05-04 21:05:38 +01:00
|
|
|
|
|
|
|
if (ctx->stage == MESA_SHADER_FRAGMENT) {
|
|
|
|
NIR_PASS_V(nir, nir_lower_mediump_io, nir_var_shader_out,
|
|
|
|
~0, false);
|
|
|
|
}
|
|
|
|
|
2020-03-05 15:11:39 +00:00
|
|
|
NIR_PASS_V(nir, nir_lower_ssbo);
|
2020-11-04 14:05:57 +00:00
|
|
|
NIR_PASS_V(nir, pan_nir_lower_zs_store);
|
2021-01-27 21:19:04 +00:00
|
|
|
NIR_PASS_V(nir, pan_lower_sample_pos);
|
2021-05-04 00:39:59 +01:00
|
|
|
NIR_PASS_V(nir, nir_lower_bit_size, bi_lower_bit_size, NULL);
|
2021-05-03 17:12:21 +01:00
|
|
|
|
|
|
|
if (nir->info.stage == MESA_SHADER_FRAGMENT) {
|
|
|
|
NIR_PASS_V(nir, nir_shader_instructions_pass,
|
|
|
|
bifrost_nir_lower_i8_frag,
|
|
|
|
nir_metadata_block_index | nir_metadata_dominance,
|
|
|
|
NULL);
|
|
|
|
}
|
2020-03-05 15:11:39 +00:00
|
|
|
|
2021-04-03 21:06:04 +01:00
|
|
|
bi_optimize_nir(nir, ctx->inputs->gpu_id, ctx->inputs->is_blend);
|
2020-04-30 08:29:10 +01:00
|
|
|
|
2020-11-04 14:05:57 +00:00
|
|
|
NIR_PASS_V(nir, pan_nir_reorder_writeout);
|
|
|
|
|
2020-12-30 19:16:51 +00:00
|
|
|
bool skip_internal = nir->info.internal;
|
|
|
|
skip_internal &= !(bifrost_debug & BIFROST_DBG_INTERNAL);
|
|
|
|
|
|
|
|
if (bifrost_debug & BIFROST_DBG_SHADERS && !skip_internal) {
|
2020-04-30 08:29:10 +01:00
|
|
|
nir_print_shader(nir, stdout);
|
|
|
|
}
|
2020-03-03 00:47:11 +00:00
|
|
|
|
2021-02-13 07:24:03 +00:00
|
|
|
info->tls_size = nir->scratch_size;
|
2020-03-10 20:20:18 +00:00
|
|
|
|
2020-03-05 15:25:19 +00:00
|
|
|
nir_foreach_function(func, nir) {
|
|
|
|
if (!func->impl)
|
|
|
|
continue;
|
|
|
|
|
2020-12-21 20:44:51 +00:00
|
|
|
ctx->ssa_alloc += func->impl->ssa_alloc;
|
|
|
|
ctx->reg_alloc += func->impl->reg_alloc;
|
|
|
|
|
2020-03-05 15:25:19 +00:00
|
|
|
emit_cf_list(ctx, &func->impl->body);
|
|
|
|
break; /* TODO: Multi-function shaders */
|
|
|
|
}
|
|
|
|
|
2020-05-28 19:44:33 +01:00
|
|
|
unsigned block_source_count = 0;
|
|
|
|
|
2020-03-19 21:21:49 +00:00
|
|
|
bi_foreach_block(ctx, _block) {
|
|
|
|
bi_block *block = (bi_block *) _block;
|
2020-05-28 19:44:33 +01:00
|
|
|
|
|
|
|
/* Name blocks now that we're done emitting so the order is
|
|
|
|
* consistent */
|
|
|
|
block->base.name = block_source_count++;
|
2020-03-19 21:21:49 +00:00
|
|
|
}
|
|
|
|
|
2021-04-13 23:47:04 +01:00
|
|
|
/* Runs before constant folding */
|
|
|
|
bi_lower_swizzle(ctx);
|
|
|
|
|
2021-02-09 23:23:06 +00:00
|
|
|
/* Runs before copy prop */
|
|
|
|
bi_opt_push_ubo(ctx);
|
2021-04-13 23:47:04 +01:00
|
|
|
bi_opt_constant_fold(ctx);
|
2021-03-02 17:59:49 +00:00
|
|
|
bi_opt_copy_prop(ctx);
|
2021-05-04 22:27:07 +01:00
|
|
|
bi_opt_dead_code_eliminate(ctx);
|
2020-03-11 19:10:32 +00:00
|
|
|
|
2021-02-18 18:46:41 +00:00
|
|
|
bi_foreach_block(ctx, _block) {
|
|
|
|
bi_block *block = (bi_block *) _block;
|
|
|
|
bi_lower_branch(block);
|
|
|
|
}
|
|
|
|
|
2021-05-04 22:27:07 +01:00
|
|
|
if (bifrost_debug & BIFROST_DBG_SHADERS && !skip_internal)
|
|
|
|
bi_print_shader(ctx, stdout);
|
|
|
|
bi_lower_fau(ctx);
|
|
|
|
bi_register_allocate(ctx);
|
2021-05-04 22:49:24 +01:00
|
|
|
bi_opt_post_ra(ctx);
|
2020-12-30 19:16:51 +00:00
|
|
|
if (bifrost_debug & BIFROST_DBG_SHADERS && !skip_internal)
|
2020-04-30 08:29:10 +01:00
|
|
|
bi_print_shader(ctx, stdout);
|
2020-03-07 00:25:00 +00:00
|
|
|
bi_schedule(ctx);
|
2021-02-12 21:50:08 +00:00
|
|
|
bi_assign_scoreboard(ctx);
|
2020-12-30 19:16:51 +00:00
|
|
|
if (bifrost_debug & BIFROST_DBG_SHADERS && !skip_internal)
|
2020-04-30 08:29:10 +01:00
|
|
|
bi_print_shader(ctx, stdout);
|
2020-10-17 11:08:17 +01:00
|
|
|
|
2021-02-13 07:24:03 +00:00
|
|
|
unsigned final_clause = bi_pack(ctx, binary);
|
2020-04-30 08:29:10 +01:00
|
|
|
|
2021-01-09 03:58:09 +00:00
|
|
|
/* If we need to wait for ATEST or BLEND in the first clause, pass the
|
|
|
|
* corresponding bits through to the renderer state descriptor */
|
|
|
|
pan_block *first_block = list_first_entry(&ctx->blocks, pan_block, link);
|
|
|
|
bi_clause *first_clause = bi_next_clause(ctx, first_block, NULL);
|
|
|
|
|
2021-02-03 19:11:44 +00:00
|
|
|
unsigned first_deps = first_clause ? first_clause->dependencies : 0;
|
2021-02-13 07:24:03 +00:00
|
|
|
info->bifrost.wait_6 = (first_deps & (1 << 6));
|
|
|
|
info->bifrost.wait_7 = (first_deps & (1 << 7));
|
2020-10-13 11:26:11 +01:00
|
|
|
|
2021-01-19 14:18:03 +00:00
|
|
|
if (bifrost_debug & BIFROST_DBG_SHADERS && !skip_internal) {
|
2021-02-13 07:24:03 +00:00
|
|
|
disassemble_bifrost(stdout, binary->data, binary->size,
|
|
|
|
bifrost_debug & BIFROST_DBG_VERBOSE);
|
2021-04-07 12:51:01 +01:00
|
|
|
fflush(stdout);
|
2020-12-16 20:02:26 +00:00
|
|
|
}
|
2020-03-05 15:11:39 +00:00
|
|
|
|
2021-02-12 13:21:45 +00:00
|
|
|
/* Pad the shader with enough zero bytes to trick the prefetcher,
|
|
|
|
* unless we're compiling an empty shader (in which case we don't pad
|
|
|
|
* so the size remains 0) */
|
2021-01-28 08:41:10 +00:00
|
|
|
unsigned prefetch_size = BIFROST_SHADER_PREFETCH - final_clause;
|
|
|
|
|
2021-02-13 07:24:03 +00:00
|
|
|
if (binary->size) {
|
|
|
|
memset(util_dynarray_grow(binary, uint8_t, prefetch_size),
|
2021-02-12 13:21:45 +00:00
|
|
|
0, prefetch_size);
|
|
|
|
}
|
2021-01-12 09:31:55 +00:00
|
|
|
|
2020-11-13 22:34:25 +00:00
|
|
|
if ((bifrost_debug & BIFROST_DBG_SHADERDB || inputs->shaderdb) &&
|
2021-01-19 14:18:03 +00:00
|
|
|
!skip_internal) {
|
2021-02-13 07:24:03 +00:00
|
|
|
bi_print_stats(ctx, binary->size, stderr);
|
2020-11-13 22:34:25 +00:00
|
|
|
}
|
|
|
|
|
2020-03-03 00:47:11 +00:00
|
|
|
ralloc_free(ctx);
|
|
|
|
}
|