mesa/src/compiler/nir/nir_range_analysis.c

Ignoring revisions in .git-blame-ignore-revs. Click here to bypass and see the normal blame view.

2535 lines
80 KiB
C
Raw Normal View History

/*
* Copyright © 2018 Intel Corporation
*
* Permission is hereby granted, free of charge, to any person obtaining a
* copy of this software and associated documentation files (the "Software"),
* to deal in the Software without restriction, including without limitation
* the rights to use, copy, modify, merge, publish, distribute, sublicense,
* and/or sell copies of the Software, and to permit persons to whom the
* Software is furnished to do so, subject to the following conditions:
*
* The above copyright notice and this permission notice (including the next
* paragraph) shall be included in all copies or substantial portions of the
* Software.
*
* THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
* IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
* FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL
* THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
* LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING
* FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS
* IN THE SOFTWARE.
*/
#include "nir.h"
#include "nir_range_analysis.h"
#include <float.h>
#include <math.h>
#include "util/hash_table.h"
#include "util/u_dynarray.h"
#include "util/u_math.h"
#include "c99_alloca.h"
/**
* Analyzes a sequence of operations to determine some aspects of the range of
* the result.
*/
struct analysis_query {
uint32_t pushed_queries;
uint32_t result_index;
};
struct analysis_state {
nir_shader *shader;
void *range_ht;
struct util_dynarray query_stack;
struct util_dynarray result_stack;
size_t query_size;
uint32_t (*get_key)(struct analysis_query *q);
bool (*lookup)(void *table, uint32_t key, uint32_t *value);
void (*insert)(void *table, uint32_t key, uint32_t value);
void (*process_query)(struct analysis_state *state, struct analysis_query *q,
uint32_t *result, const uint32_t *src);
};
static void *
push_analysis_query(struct analysis_state *state, size_t size)
{
struct analysis_query *q = util_dynarray_grow_bytes(&state->query_stack, 1, size);
q->pushed_queries = 0;
q->result_index = util_dynarray_num_elements(&state->result_stack, uint32_t);
util/dynarray: infer type in append Most of the time, we can infer the type to append in util_dynarray_append using __typeof__, which is standardized in C23 and support in Jesse's MSMSVCV. This patch drops the type argument most of the time, making util_dynarray a little more ergonomic to use. This is done in four steps. First, rename util_dynarray_append -> util_dynarray_append_typed bash -c "find . -type f -exec sed -i -e 's/util_dynarray_append(/util_dynarray_append_typed(/g' \{} \;" Then, add a new append that infers the type. This is much more ergonomic for what you want most of the time. Next, use type-inferred append as much as possible, via Coccinelle patch (plus manual fixup): @@ expression dynarray, element; type type; @@ -util_dynarray_append_typed(dynarray, type, element); +util_dynarray_append(dynarray, element); Finally, hand fixup cases that Coccinelle missed or incorrectly translated, of which there were several because we can't used the untyped append with a literal (since the sizeof won't do what you want). All four steps are squashed to produce a single patch changing every util_dynarray_append call site in tree to either drop a type parameter (if possible) or insert a _typed suffix (if we can't infer). As such, the final patch is best reviewed by hand even though it was tool-assisted. No Long Linguine Meals were involved in the making of this patch. Signed-off-by: Alyssa Rosenzweig <alyssa.rosenzweig@intel.com> Acked-by: Faith Ekstrand <faith.ekstrand@collabora.com> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/38038>
2025-10-23 15:36:13 -04:00
util_dynarray_append_typed(&state->result_stack, uint32_t, 0);
return q;
}
/* Helper for performing range analysis without recursion. */
static uint32_t
perform_analysis(struct analysis_state *state)
{
while (state->query_stack.size) {
struct analysis_query *cur =
(struct analysis_query *)((char *)util_dynarray_end(&state->query_stack) - state->query_size);
uint32_t *result = util_dynarray_element(&state->result_stack, uint32_t, cur->result_index);
uint32_t key = state->get_key(cur);
/* There might be a cycle-resolving entry for loop header phis. Ignore this when finishing
* them by testing pushed_queries.
*/
if (cur->pushed_queries == 0 && key != UINT32_MAX &&
state->lookup(state->range_ht, key, result)) {
state->query_stack.size -= state->query_size;
continue;
}
uint32_t *src = (uint32_t *)util_dynarray_end(&state->result_stack) - cur->pushed_queries;
state->result_stack.size -= sizeof(uint32_t) * cur->pushed_queries;
uint32_t prev_num_queries = state->query_stack.size;
state->process_query(state, cur, result, src);
uint32_t num_queries = state->query_stack.size;
if (num_queries > prev_num_queries) {
cur = (struct analysis_query *)util_dynarray_element(&state->query_stack, char,
prev_num_queries - state->query_size);
cur->pushed_queries = (num_queries - prev_num_queries) / state->query_size;
continue;
}
if (key != UINT32_MAX)
state->insert(state->range_ht, key, *result);
state->query_stack.size -= state->query_size;
}
assert(state->result_stack.size == sizeof(uint32_t));
uint32_t res = util_dynarray_top(&state->result_stack, uint32_t);
util_dynarray_fini(&state->query_stack);
util_dynarray_fini(&state->result_stack);
return res;
}
static fp_class_mask
analyze_fp_constant(const nir_load_const_instr *const load)
{
fp_class_mask result = 0;
for (unsigned i = 0; i < load->def.num_components; ++i) {
const double v = nir_const_value_as_float(load->value[i],
load->def.bit_size);
if (!isnan(v) && floor(v) != v)
result |= FP_CLASS_NON_INTEGRAL;
if (isnan(v))
result |= FP_CLASS_NAN;
else if (v == -INFINITY)
result |= FP_CLASS_NEG_INF;
else if (v < -1.0)
result |= FP_CLASS_LT_NEG_ONE;
else if (v == -1.0)
result |= FP_CLASS_NEG_ONE;
else if (v < 0.0)
result |= FP_CLASS_LT_ZERO_GT_NEG_ONE;
else if (dui(v) == 0)
result |= FP_CLASS_POS_ZERO;
else if (v == 0.0)
result |= FP_CLASS_NEG_ZERO;
else if (v < 1.0)
result |= FP_CLASS_GT_ZERO_LT_POS_ONE;
else if (v == 1.0)
result |= FP_CLASS_POS_ONE;
else if (v < INFINITY)
result |= FP_CLASS_GT_POS_ONE;
else
result |= FP_CLASS_POS_INF;
if (v != 0) {
/* handle potential denorm flushing. */
bool is_denorm = false;
switch (load->def.bit_size) {
case 64:
is_denorm = fabs(v) < DBL_MIN;
break;
case 32:
is_denorm = fabs(v) < FLT_MIN;
break;
case 16:
is_denorm = fabs(v) < ldexp(1.0, -14);
break;
default:
UNREACHABLE("unsupported float size");
}
if (is_denorm)
result |= v < 0.0 ? FP_CLASS_NEG_ZERO : FP_CLASS_POS_ZERO;
}
}
return result;
}
struct fp_query {
struct analysis_query head;
const nir_def *def;
};
static void
push_fp_query(struct analysis_state *state, const nir_def *def)
{
struct fp_query *pushed_q = push_analysis_query(state, sizeof(struct fp_query));
pushed_q->def = def;
}
static uint32_t
get_fp_key(struct analysis_query *q)
{
struct fp_query *fp_q = (struct fp_query *)q;
if (!nir_def_is_alu(fp_q->def))
return UINT32_MAX;
return fp_q->def->index;
}
static bool
fp_lookup(void *table, uint32_t key, uint32_t *value)
{
nir_fp_analysis_state *state = table;
if (BITSET_TEST(state->bitset, key)) {
*value = *(uint16_t *)util_sparse_array_get(&state->arr, key);
return true;
} else {
return false;
}
}
static void
fp_insert(void *table, uint32_t key, uint32_t value)
{
nir_fp_analysis_state *state = table;
BITSET_SET(state->bitset, key);
state->max = MAX2(state->max, (int)key);
*(uint16_t *)util_sparse_array_get(&state->arr, key) = value;
}
static fp_class_mask
fneg_fp_class(fp_class_mask src)
{
fp_class_mask result = src & (FP_CLASS_NAN | FP_CLASS_NON_INTEGRAL);
#define NEG_BIT(neg, pos) \
if (src & FP_CLASS_##pos) \
result |= FP_CLASS_##neg; \
if (src & FP_CLASS_##neg) \
result |= FP_CLASS_##pos;
NEG_BIT(NEG_INF, POS_INF);
NEG_BIT(LT_NEG_ONE, GT_POS_ONE);
NEG_BIT(NEG_ONE, POS_ONE);
NEG_BIT(LT_ZERO_GT_NEG_ONE, GT_ZERO_LT_POS_ONE);
NEG_BIT(NEG_ZERO, POS_ZERO);
#undef NEG_BIT
return result;
}
static fp_class_mask
fmul_fp_class(fp_class_mask left, fp_class_mask right, bool mulz, bool src_eq, bool src_neg_eq)
{
/* For runtime performance, shortcut the common completely unknown case. */
if (left == FP_CLASS_UNKNOWN && right == FP_CLASS_UNKNOWN && !src_eq && !src_neg_eq)
return FP_CLASS_UNKNOWN;
fp_class_mask result = 0;
if (left & FP_CLASS_NAN) {
if (right & FP_CLASS_ANY_ZERO)
result |= mulz ? FP_CLASS_POS_ZERO : FP_CLASS_NAN;
if (right & (FP_CLASS_ANY_NEG | FP_CLASS_ANY_POS | FP_CLASS_NAN))
result |= FP_CLASS_NAN;
}
if (left & FP_CLASS_NEG_INF) {
if (right & FP_CLASS_ANY_ZERO)
result |= mulz ? FP_CLASS_POS_ZERO : FP_CLASS_NAN;
if (right & FP_CLASS_ANY_NEG)
result |= FP_CLASS_POS_INF;
if (right & FP_CLASS_ANY_POS)
result |= FP_CLASS_NEG_INF;
if (right & FP_CLASS_NAN)
result |= FP_CLASS_NAN;
}
if (left & FP_CLASS_POS_INF) {
if (right & FP_CLASS_ANY_ZERO)
result |= mulz ? FP_CLASS_POS_ZERO : FP_CLASS_NAN;
if (right & FP_CLASS_ANY_POS)
result |= FP_CLASS_POS_INF;
if (right & FP_CLASS_ANY_NEG)
result |= FP_CLASS_NEG_INF;
if (right & FP_CLASS_NAN)
result |= FP_CLASS_NAN;
}
if (left & FP_CLASS_ANY_NEG_FINITE) {
if (right & FP_CLASS_POS_ZERO)
result |= mulz ? FP_CLASS_POS_ZERO : FP_CLASS_NEG_ZERO;
result |= fneg_fp_class(right & (FP_CLASS_NEG_ZERO | FP_CLASS_POS_INF | FP_CLASS_NEG_INF | FP_CLASS_NAN));
}
if (left & FP_CLASS_ANY_POS_FINITE) {
if (right & FP_CLASS_NEG_ZERO)
result |= mulz ? FP_CLASS_POS_ZERO : FP_CLASS_NEG_ZERO;
result |= right & (FP_CLASS_POS_ZERO | FP_CLASS_POS_INF | FP_CLASS_NEG_INF | FP_CLASS_NAN);
}
if (left & FP_CLASS_LT_NEG_ONE) {
if (right & FP_CLASS_LT_NEG_ONE)
result |= FP_CLASS_POS_INF | FP_CLASS_GT_POS_ONE;
if (right & FP_CLASS_NEG_ONE)
result |= FP_CLASS_GT_POS_ONE;
if (right & FP_CLASS_LT_ZERO_GT_NEG_ONE)
result |= FP_CLASS_GT_ZERO_LT_POS_ONE | FP_CLASS_POS_ONE | FP_CLASS_GT_POS_ONE;
if (right & FP_CLASS_GT_POS_ONE)
result |= FP_CLASS_NEG_INF | FP_CLASS_LT_NEG_ONE;
if (right & FP_CLASS_POS_ONE)
result |= FP_CLASS_LT_NEG_ONE;
if (right & FP_CLASS_GT_ZERO_LT_POS_ONE)
result |= FP_CLASS_LT_ZERO_GT_NEG_ONE | FP_CLASS_NEG_ONE | FP_CLASS_LT_NEG_ONE;
}
if (left & FP_CLASS_GT_POS_ONE) {
if (right & FP_CLASS_GT_POS_ONE)
result |= FP_CLASS_POS_INF | FP_CLASS_GT_POS_ONE;
if (right & FP_CLASS_POS_ONE)
result |= FP_CLASS_GT_POS_ONE;
if (right & FP_CLASS_GT_ZERO_LT_POS_ONE)
result |= FP_CLASS_GT_ZERO_LT_POS_ONE | FP_CLASS_POS_ONE | FP_CLASS_GT_POS_ONE;
if (right & FP_CLASS_LT_NEG_ONE)
result |= FP_CLASS_NEG_INF | FP_CLASS_LT_NEG_ONE;
if (right & FP_CLASS_NEG_ONE)
result |= FP_CLASS_LT_NEG_ONE;
if (right & FP_CLASS_LT_ZERO_GT_NEG_ONE)
result |= FP_CLASS_LT_ZERO_GT_NEG_ONE | FP_CLASS_NEG_ONE | FP_CLASS_LT_NEG_ONE;
}
if (left & FP_CLASS_NEG_ONE)
result |= fneg_fp_class(right & ~FP_CLASS_ANY_ZERO);
if (left & FP_CLASS_POS_ONE)
result |= right & ~FP_CLASS_ANY_ZERO;
if (left & FP_CLASS_LT_ZERO_GT_NEG_ONE) {
if (right & FP_CLASS_LT_NEG_ONE)
result |= FP_CLASS_GT_ZERO_LT_POS_ONE | FP_CLASS_POS_ONE | FP_CLASS_GT_POS_ONE;
if (right & FP_CLASS_NEG_ONE)
result |= FP_CLASS_GT_ZERO_LT_POS_ONE;
if (right & FP_CLASS_LT_ZERO_GT_NEG_ONE)
result |= FP_CLASS_POS_ZERO | FP_CLASS_GT_ZERO_LT_POS_ONE;
if (right & FP_CLASS_GT_POS_ONE)
result |= FP_CLASS_LT_ZERO_GT_NEG_ONE | FP_CLASS_NEG_ONE | FP_CLASS_LT_NEG_ONE;
if (right & FP_CLASS_POS_ONE)
result |= FP_CLASS_LT_ZERO_GT_NEG_ONE;
if (right & FP_CLASS_GT_ZERO_LT_POS_ONE)
result |= FP_CLASS_NEG_ZERO | FP_CLASS_LT_ZERO_GT_NEG_ONE;
}
if (left & FP_CLASS_GT_ZERO_LT_POS_ONE) {
if (right & FP_CLASS_GT_POS_ONE)
result |= FP_CLASS_GT_ZERO_LT_POS_ONE | FP_CLASS_POS_ONE | FP_CLASS_GT_POS_ONE;
if (right & FP_CLASS_POS_ONE)
result |= FP_CLASS_GT_ZERO_LT_POS_ONE;
if (right & FP_CLASS_GT_ZERO_LT_POS_ONE)
result |= FP_CLASS_POS_ZERO | FP_CLASS_GT_ZERO_LT_POS_ONE;
if (right & FP_CLASS_LT_NEG_ONE)
result |= FP_CLASS_LT_ZERO_GT_NEG_ONE | FP_CLASS_NEG_ONE | FP_CLASS_LT_NEG_ONE;
if (right & FP_CLASS_NEG_ONE)
result |= FP_CLASS_LT_ZERO_GT_NEG_ONE;
if (right & FP_CLASS_LT_ZERO_GT_NEG_ONE)
result |= FP_CLASS_NEG_ZERO | FP_CLASS_LT_ZERO_GT_NEG_ONE;
}
if (left & FP_CLASS_NEG_ZERO) {
if (mulz) {
result |= FP_CLASS_POS_ZERO;
} else {
if (right & (FP_CLASS_ANY_INF | FP_CLASS_NAN))
result |= FP_CLASS_NAN;
if (right & (FP_CLASS_ANY_NEG_FINITE | FP_CLASS_NEG_ZERO))
result |= FP_CLASS_POS_ZERO;
if (right & (FP_CLASS_ANY_POS_FINITE | FP_CLASS_POS_ZERO))
result |= FP_CLASS_NEG_ZERO;
}
}
if (left & FP_CLASS_POS_ZERO) {
if (mulz) {
result |= FP_CLASS_POS_ZERO;
} else {
if (right & (FP_CLASS_ANY_INF | FP_CLASS_NAN))
result |= FP_CLASS_NAN;
if (right & (FP_CLASS_ANY_POS_FINITE | FP_CLASS_POS_ZERO))
result |= FP_CLASS_POS_ZERO;
if (right & (FP_CLASS_ANY_NEG_FINITE | FP_CLASS_NEG_ZERO))
result |= FP_CLASS_NEG_ZERO;
}
}
if (src_eq || src_neg_eq) {
/* This case can't create new ones. */
if (!(left & (FP_CLASS_POS_ONE | FP_CLASS_NEG_ONE)))
result &= ~(FP_CLASS_POS_ONE | FP_CLASS_NEG_ONE);
if (src_eq)
result &= ~(FP_CLASS_ANY_NEG | FP_CLASS_NEG_ZERO);
else if (src_neg_eq && mulz)
result &= ~FP_CLASS_ANY_POS;
else if (src_neg_eq)
result &= ~(FP_CLASS_ANY_POS | FP_CLASS_POS_ZERO);
}
if ((left | right) & FP_CLASS_NON_INTEGRAL) {
if (result & (FP_CLASS_LT_NEG_ONE | FP_CLASS_LT_ZERO_GT_NEG_ONE |
FP_CLASS_GT_ZERO_LT_POS_ONE | FP_CLASS_GT_POS_ONE))
result |= FP_CLASS_NON_INTEGRAL;
} else {
result &= ~(FP_CLASS_LT_ZERO_GT_NEG_ONE | FP_CLASS_GT_ZERO_LT_POS_ONE);
}
return result;
}
static fp_class_mask
fadd_fp_class(fp_class_mask left, fp_class_mask right)
{
/* For runtime performance, shortcut the common completely unknown case. */
if (left == FP_CLASS_UNKNOWN && right == FP_CLASS_UNKNOWN)
return FP_CLASS_UNKNOWN;
fp_class_mask result = (left | right) & FP_CLASS_NAN;
/* X + Y is NaN if either operand is NaN or if one operand is +Inf and
* the other is -Inf.
*/
if (left & FP_CLASS_NEG_INF) {
if (right & FP_CLASS_POS_INF)
result |= FP_CLASS_NAN;
if (right & (FP_CLASS_ANY_FINITE | FP_CLASS_NEG_INF))
result |= FP_CLASS_NEG_INF;
}
if (left & FP_CLASS_POS_INF) {
if (right & FP_CLASS_NEG_INF)
result |= FP_CLASS_NAN;
if (right & (FP_CLASS_ANY_FINITE | FP_CLASS_POS_INF))
result |= FP_CLASS_POS_INF;
}
if (left & FP_CLASS_LT_NEG_ONE) {
result |= (right & FP_CLASS_ANY_INF);
if (right & FP_CLASS_LT_NEG_ONE)
result |= FP_CLASS_NEG_INF | FP_CLASS_LT_NEG_ONE;
if (right & (FP_CLASS_NEG_ONE | FP_CLASS_LT_ZERO_GT_NEG_ONE | FP_CLASS_ANY_ZERO))
result |= FP_CLASS_LT_NEG_ONE;
if (right & (FP_CLASS_GT_ZERO_LT_POS_ONE | FP_CLASS_POS_ONE))
result |= FP_CLASS_LT_NEG_ONE | FP_CLASS_NEG_ONE | FP_CLASS_LT_ZERO_GT_NEG_ONE;
if (right & FP_CLASS_GT_POS_ONE)
result |= FP_CLASS_ANY_FINITE;
}
if (left & FP_CLASS_GT_POS_ONE) {
result |= (right & FP_CLASS_ANY_INF);
if (right & FP_CLASS_GT_POS_ONE)
result |= FP_CLASS_POS_INF | FP_CLASS_GT_POS_ONE;
if (right & (FP_CLASS_POS_ONE | FP_CLASS_GT_ZERO_LT_POS_ONE | FP_CLASS_ANY_ZERO))
result |= FP_CLASS_GT_POS_ONE;
if (right & (FP_CLASS_LT_ZERO_GT_NEG_ONE | FP_CLASS_NEG_ONE))
result |= FP_CLASS_GT_POS_ONE | FP_CLASS_POS_ONE | FP_CLASS_GT_ZERO_LT_POS_ONE;
if (right & FP_CLASS_LT_NEG_ONE)
result |= FP_CLASS_ANY_FINITE;
}
if (left & FP_CLASS_NEG_ONE) {
result |= (right & FP_CLASS_ANY_INF);
if (right & (FP_CLASS_LT_NEG_ONE | FP_CLASS_NEG_ONE | FP_CLASS_LT_ZERO_GT_NEG_ONE))
result |= FP_CLASS_LT_NEG_ONE;
if (right & (FP_CLASS_LT_ZERO_GT_NEG_ONE | FP_CLASS_ANY_ZERO | FP_CLASS_GT_ZERO_LT_POS_ONE))
result |= FP_CLASS_NEG_ONE;
if (right & FP_CLASS_GT_ZERO_LT_POS_ONE)
result |= FP_CLASS_LT_ZERO_GT_NEG_ONE;
if (right & FP_CLASS_POS_ONE)
result |= FP_CLASS_POS_ZERO;
if (right & FP_CLASS_GT_POS_ONE)
result |= FP_CLASS_GT_ZERO_LT_POS_ONE | FP_CLASS_POS_ONE | FP_CLASS_GT_POS_ONE;
}
if (left & FP_CLASS_POS_ONE) {
result |= (right & FP_CLASS_ANY_INF);
if (right & (FP_CLASS_GT_POS_ONE | FP_CLASS_POS_ONE | FP_CLASS_GT_ZERO_LT_POS_ONE))
result |= FP_CLASS_GT_POS_ONE;
if (right & (FP_CLASS_GT_ZERO_LT_POS_ONE | FP_CLASS_ANY_ZERO | FP_CLASS_LT_ZERO_GT_NEG_ONE))
result |= FP_CLASS_POS_ONE;
if (right & FP_CLASS_LT_ZERO_GT_NEG_ONE)
result |= FP_CLASS_GT_ZERO_LT_POS_ONE;
if (right & FP_CLASS_NEG_ONE)
result |= FP_CLASS_POS_ZERO;
if (right & FP_CLASS_LT_NEG_ONE)
result |= FP_CLASS_LT_ZERO_GT_NEG_ONE | FP_CLASS_NEG_ONE | FP_CLASS_LT_NEG_ONE;
}
if (left & FP_CLASS_LT_ZERO_GT_NEG_ONE) {
result |= (right & FP_CLASS_ANY_INF);
if (right & FP_CLASS_LT_NEG_ONE)
result |= FP_CLASS_LT_NEG_ONE;
if (right & FP_CLASS_NEG_ONE)
result |= FP_CLASS_LT_NEG_ONE | FP_CLASS_NEG_ONE;
if (right & FP_CLASS_LT_ZERO_GT_NEG_ONE)
result |= FP_CLASS_LT_NEG_ONE | FP_CLASS_NEG_ONE | FP_CLASS_LT_ZERO_GT_NEG_ONE;
if (right & FP_CLASS_ANY_ZERO)
result |= FP_CLASS_LT_ZERO_GT_NEG_ONE;
if (right & FP_CLASS_GT_ZERO_LT_POS_ONE)
result |= FP_CLASS_LT_ZERO_GT_NEG_ONE | FP_CLASS_ANY_ZERO | FP_CLASS_GT_ZERO_LT_POS_ONE;
if (right & FP_CLASS_POS_ONE)
result |= FP_CLASS_GT_ZERO_LT_POS_ONE | FP_CLASS_POS_ONE;
if (right & FP_CLASS_GT_POS_ONE)
result |= FP_CLASS_GT_ZERO_LT_POS_ONE | FP_CLASS_POS_ONE | FP_CLASS_GT_POS_ONE;
}
if (left & FP_CLASS_GT_ZERO_LT_POS_ONE) {
result |= (right & FP_CLASS_ANY_INF);
if (right & FP_CLASS_GT_POS_ONE)
result |= FP_CLASS_GT_POS_ONE;
if (right & FP_CLASS_POS_ONE)
result |= FP_CLASS_GT_POS_ONE | FP_CLASS_POS_ONE;
if (right & FP_CLASS_GT_ZERO_LT_POS_ONE)
result |= FP_CLASS_GT_POS_ONE | FP_CLASS_POS_ONE | FP_CLASS_GT_ZERO_LT_POS_ONE;
if (right & FP_CLASS_ANY_ZERO)
result |= FP_CLASS_GT_ZERO_LT_POS_ONE;
if (right & FP_CLASS_LT_ZERO_GT_NEG_ONE)
result |= FP_CLASS_GT_ZERO_LT_POS_ONE | FP_CLASS_ANY_ZERO | FP_CLASS_LT_ZERO_GT_NEG_ONE;
if (right & FP_CLASS_NEG_ONE)
result |= FP_CLASS_LT_ZERO_GT_NEG_ONE | FP_CLASS_NEG_ONE;
if (right & FP_CLASS_LT_NEG_ONE)
result |= FP_CLASS_LT_ZERO_GT_NEG_ONE | FP_CLASS_NEG_ONE | FP_CLASS_LT_NEG_ONE;
}
if (left & FP_CLASS_NEG_ZERO) {
result |= right;
}
if (left & FP_CLASS_POS_ZERO) {
result |= right & ~FP_CLASS_NEG_ZERO;
if (right & FP_CLASS_NEG_ZERO)
result |= FP_CLASS_POS_ZERO;
}
if ((left | right) & FP_CLASS_NON_INTEGRAL) {
if (result & (FP_CLASS_LT_NEG_ONE | FP_CLASS_LT_ZERO_GT_NEG_ONE |
FP_CLASS_GT_ZERO_LT_POS_ONE | FP_CLASS_GT_POS_ONE))
result |= FP_CLASS_NON_INTEGRAL;
} else {
result &= ~(FP_CLASS_LT_ZERO_GT_NEG_ONE | FP_CLASS_GT_ZERO_LT_POS_ONE);
}
return result;
}
static fp_class_mask
frcp_fp_class(fp_class_mask src)
{
fp_class_mask result = src & FP_CLASS_NAN;
/* Inf/Zero result in Zero/Inf.*/
if (src & FP_CLASS_NEG_INF)
result |= FP_CLASS_NEG_ZERO;
if (src & FP_CLASS_POS_INF)
result |= FP_CLASS_POS_ZERO;
if (src & FP_CLASS_NEG_ZERO)
result |= FP_CLASS_NEG_INF;
if (src & FP_CLASS_POS_ZERO)
result |= FP_CLASS_POS_INF;
/* One results in one. */
if (src & FP_CLASS_NEG_ONE)
result |= FP_CLASS_NEG_ONE;
if (src & FP_CLASS_POS_ONE)
result |= FP_CLASS_POS_ONE;
if (src & FP_CLASS_LT_NEG_ONE)
result |= FP_CLASS_LT_ZERO_GT_NEG_ONE | FP_CLASS_NEG_ZERO | FP_CLASS_NON_INTEGRAL;
if (src & FP_CLASS_GT_POS_ONE)
result |= FP_CLASS_GT_ZERO_LT_POS_ONE | FP_CLASS_POS_ZERO | FP_CLASS_NON_INTEGRAL;
if (src & FP_CLASS_LT_ZERO_GT_NEG_ONE)
result |= FP_CLASS_LT_NEG_ONE | FP_CLASS_NEG_INF | FP_CLASS_NON_INTEGRAL;
if (src & FP_CLASS_GT_ZERO_LT_POS_ONE)
result |= FP_CLASS_GT_POS_ONE | FP_CLASS_POS_INF | FP_CLASS_NON_INTEGRAL;
return result;
}
static fp_class_mask
fsqrt_fp_class(fp_class_mask src)
{
fp_class_mask result = src & (FP_CLASS_NAN | FP_CLASS_ANY_ZERO);
if (src & FP_CLASS_ANY_NEG)
result |= FP_CLASS_NAN;
if (src & FP_CLASS_GT_ZERO_LT_POS_ONE)
result |= FP_CLASS_GT_ZERO_LT_POS_ONE | FP_CLASS_NON_INTEGRAL;
if (src & FP_CLASS_POS_ONE)
result |= FP_CLASS_POS_ONE;
if (src & FP_CLASS_GT_POS_ONE)
result |= FP_CLASS_GT_POS_ONE | FP_CLASS_POS_ONE | FP_CLASS_NON_INTEGRAL;
if (src & FP_CLASS_POS_INF)
result |= FP_CLASS_POS_INF;
return result;
}
static fp_class_mask
fmin_part(fp_class_mask upper_bound, fp_class_mask value)
{
/* Find the highest value in upper_bound, and return all
* smaller or equal values from value.
*/
upper_bound &= FP_CLASS_ANY_NEG | FP_CLASS_ANY_ZERO | FP_CLASS_ANY_POS;
value &= FP_CLASS_ANY_NEG | FP_CLASS_ANY_ZERO | FP_CLASS_ANY_POS;
/* This works even in the case where upper_bound is 0 */
return value & BITFIELD_MASK(util_last_bit(upper_bound));
}
static fp_class_mask
fmin_fp_class(fp_class_mask left, fp_class_mask right)
{
fp_class_mask result = 0;
/* If one source is NaN, we have to include the whole range of the other source. */
if (left & FP_CLASS_NAN)
result |= right;
if (right & FP_CLASS_NAN)
result |= left;
result |= fmin_part(left, right);
result |= fmin_part(right, left);
/* Could probably do better, but meh. */
if ((left | right) & FP_CLASS_NON_INTEGRAL) {
if (result & (FP_CLASS_LT_NEG_ONE | FP_CLASS_LT_ZERO_GT_NEG_ONE |
FP_CLASS_GT_ZERO_LT_POS_ONE | FP_CLASS_GT_POS_ONE))
result |= FP_CLASS_NON_INTEGRAL;
}
return result;
}
static fp_class_mask
handle_sz(const nir_alu_instr *alu, fp_class_mask src)
{
if (nir_alu_instr_is_signed_zero_preserve(alu) || !(src & FP_CLASS_ANY_ZERO))
return src;
return src | FP_CLASS_ANY_ZERO;
}
static fp_class_mask
intrinsic_fp_class(const nir_intrinsic_instr *intrin)
{
switch (intrin->intrinsic) {
case nir_intrinsic_load_typed_buffer_amd: {
const enum pipe_format format = nir_intrinsic_format(intrin);
if (format == PIPE_FORMAT_NONE)
return FP_CLASS_UNKNOWN;
const struct util_format_description *desc = util_format_description(format);
int i = util_format_get_first_non_void_channel(format);
if (i == -1)
return FP_CLASS_UNKNOWN;
bool is_signed = desc->channel[i].type == UTIL_FORMAT_TYPE_SIGNED;
bool is_unsigned = desc->channel[i].type == UTIL_FORMAT_TYPE_UNSIGNED;
bool normalized = desc->channel[i].normalized;
if ((!is_signed && !is_unsigned) || desc->channel[i].pure_integer)
return FP_CLASS_UNKNOWN;
fp_class_mask result = FP_CLASS_POS_ZERO | FP_CLASS_POS_ONE;
result |= is_signed ? FP_CLASS_NEG_ONE : 0;
if (normalized) {
result |= FP_CLASS_GT_ZERO_LT_POS_ONE | FP_CLASS_NON_INTEGRAL;
result |= is_signed ? FP_CLASS_LT_ZERO_GT_NEG_ONE : 0;
} else {
result |= FP_CLASS_GT_POS_ONE;
result |= is_signed ? FP_CLASS_LT_NEG_ONE : 0;
}
return result;
}
case nir_intrinsic_load_front_face_fsign:
return FP_CLASS_POS_ONE | FP_CLASS_NEG_ONE;
default:
return FP_CLASS_UNKNOWN;
}
}
static fp_class_mask
tex_fp_class(nir_tex_instr *tex)
{
/* Not much to analyze, except shadow compare. */
if (!tex->is_shadow)
return FP_CLASS_UNKNOWN;
fp_class_mask result = FP_CLASS_POS_ZERO | FP_CLASS_POS_ONE;
/* Gather returns 0 or 1, other ops can interpolate.
* Cube corners are special even for gathers.
*/
if (tex->op != nir_texop_tg4 || tex->sampler_dim == GLSL_SAMPLER_DIM_CUBE)
result |= FP_CLASS_GT_ZERO_LT_POS_ONE | FP_CLASS_NON_INTEGRAL;
return result;
}
/**
* Analyze an expression to determine the possible fp classes of its result
*/
static void
process_fp_query(struct analysis_state *state, struct analysis_query *aq, uint32_t *result,
const uint32_t *src_res)
{
struct fp_query q = *(struct fp_query *)aq;
const nir_def *def = q.def;
if (nir_def_is_const(def)) {
*result = analyze_fp_constant(nir_def_as_load_const(def));
return;
} else if (nir_def_is_intrinsic(def)) {
*result = intrinsic_fp_class(nir_def_as_intrinsic(def));
return;
} else if (nir_def_is_tex(def)) {
*result = tex_fp_class(nir_def_as_tex(def));
return;
} else if (!nir_def_is_alu(def)) {
*result = FP_CLASS_UNKNOWN;
return;
}
const nir_alu_instr *const alu = nir_def_as_alu(def);
if (!aq->pushed_queries) {
switch (alu->op) {
case nir_op_bcsel:
push_fp_query(state, alu->src[1].src.ssa);
push_fp_query(state, alu->src[2].src.ssa);
return;
case nir_op_mov:
case nir_op_fabs:
case nir_op_fexp2:
case nir_op_flog2:
case nir_op_frcp:
case nir_op_fsqrt:
case nir_op_frsq:
case nir_op_fneg:
case nir_op_fsat:
case nir_op_fsign:
case nir_op_ffloor:
case nir_op_fceil:
case nir_op_ftrunc:
case nir_op_fround_even:
case nir_op_ffract:
case nir_op_fsin:
case nir_op_fcos:
case nir_op_fsin_amd:
case nir_op_fcos_amd:
case nir_op_f2f16:
case nir_op_f2f16_rtz:
case nir_op_f2f16_rtne:
case nir_op_f2f32:
case nir_op_f2f64:
case nir_op_fdot2:
case nir_op_fdot3:
case nir_op_fdot4:
case nir_op_fdot8:
case nir_op_fdot16:
case nir_op_fdot2_replicated:
case nir_op_fdot3_replicated:
case nir_op_fdot4_replicated:
case nir_op_fdot8_replicated:
case nir_op_fdot16_replicated:
push_fp_query(state, alu->src[0].src.ssa);
return;
case nir_op_fadd:
case nir_op_fsub:
case nir_op_fmax:
case nir_op_fmin:
case nir_op_fmul:
case nir_op_fmulz:
case nir_op_fpow:
case nir_op_vec2:
push_fp_query(state, alu->src[0].src.ssa);
push_fp_query(state, alu->src[1].src.ssa);
return;
case nir_op_ffma:
case nir_op_ffmaz:
case nir_op_flrp:
push_fp_query(state, alu->src[0].src.ssa);
push_fp_query(state, alu->src[1].src.ssa);
push_fp_query(state, alu->src[2].src.ssa);
return;
default:
break;
}
}
fp_class_mask r = FP_CLASS_UNKNOWN;
switch (alu->op) {
case nir_op_b2i16:
case nir_op_b2i32:
case nir_op_b2i64:
/* b2i32 will generate either 0x00000000 or 0x00000001. When those bit
* patterns are interpreted as floating point, they are 0.0 and
* 1.401298464324817e-45. The latter is subnormal.
*/
r = FP_CLASS_POS_ZERO | FP_CLASS_GT_ZERO_LT_POS_ONE;
break;
case nir_op_b2f16:
case nir_op_b2f32:
case nir_op_b2f64:
r = FP_CLASS_POS_ZERO | FP_CLASS_POS_ONE;
break;
case nir_op_vec2:
case nir_op_bcsel:
r = src_res[0] | src_res[1];
break;
case nir_op_i2f16:
case nir_op_i2f32:
case nir_op_i2f64:
r &= ~FP_CLASS_NAN;
r &= ~FP_CLASS_NON_INTEGRAL;
r &= ~FP_CLASS_GT_ZERO_LT_POS_ONE;
r &= ~FP_CLASS_LT_ZERO_GT_NEG_ONE;
r &= ~FP_CLASS_NEG_ZERO;
if (alu->def.bit_size > 16 || alu->src[0].src.ssa->bit_size <= 16)
r &= ~FP_CLASS_ANY_INF;
break;
case nir_op_u2f16:
case nir_op_u2f32:
case nir_op_u2f64:
r &= ~FP_CLASS_NAN;
r &= ~FP_CLASS_NON_INTEGRAL;
r &= ~FP_CLASS_GT_ZERO_LT_POS_ONE;
r &= ~FP_CLASS_NEG_ZERO;
r &= ~FP_CLASS_ANY_NEG;
if (alu->def.bit_size > 16 || alu->src[0].src.ssa->bit_size < 16)
r &= ~FP_CLASS_ANY_INF;
break;
case nir_op_f2f16:
case nir_op_f2f16_rtz:
case nir_op_f2f16_rtne:
case nir_op_f2f32:
case nir_op_f2f64: {
r = handle_sz(alu, src_res[0]);
if (alu->src[0].src.ssa->bit_size > alu->def.bit_size) {
bool rtz = alu->op == nir_op_f2f16_rtz;
if (alu->op != nir_op_f2f16_rtne && alu->op != nir_op_f2f16_rtz) {
nir_shader *shader = nir_cf_node_get_function(&alu->instr.block->cf_node)->function->shader;
unsigned execution_mode = shader->info.float_controls_execution_mode;
rtz = nir_is_rounding_mode_rtz(execution_mode, alu->def.bit_size);
}
/* Unless we are rounding towards zero, large values can create Inf. */
if (r & FP_CLASS_LT_NEG_ONE) {
if (!rtz)
r |= FP_CLASS_NEG_INF;
r |= FP_CLASS_NEG_ONE;
}
if (r & FP_CLASS_GT_POS_ONE) {
if (!rtz)
r |= FP_CLASS_POS_INF;
r |= FP_CLASS_POS_ONE;
}
/* Underflow can create new zeros. */
if (r & FP_CLASS_LT_ZERO_GT_NEG_ONE) {
if (!rtz)
r |= FP_CLASS_NEG_ONE;
r |= FP_CLASS_NEG_ZERO;
}
if (r & FP_CLASS_GT_ZERO_LT_POS_ONE) {
if (!rtz)
r |= FP_CLASS_POS_ONE;
r |= FP_CLASS_POS_ZERO;
}
}
break;
}
case nir_op_fneg:
r = fneg_fp_class(src_res[0]);
break;
case nir_op_fabs:
r = src_res[0];
r |= fneg_fp_class(r & (FP_CLASS_ANY_NEG | FP_CLASS_NEG_ZERO));
r &= ~(FP_CLASS_ANY_NEG | FP_CLASS_NEG_ZERO);
break;
case nir_op_fadd: {
r = fadd_fp_class(src_res[0], src_res[1]);
break;
}
case nir_op_fsub: {
r = fadd_fp_class(src_res[0], fneg_fp_class(src_res[1]));
break;
}
nir/range-analysis: Adjust result range of exp2 to account for flush-to-zero Fixes piglit tests (new in piglit!110): - fs-underflow-exp2-compare-zero.shader_test Bugzilla: https://bugs.freedesktop.org/show_bug.cgi?id=111308 Fixes: 405de7ccb6c ("nir/range-analysis: Rudimentary value range analysis pass") Reviewed-by: Caio Marcelo de Oliveira Filho <caio.oliveira@intel.com> Most of the shaders affected are, unsurprisingly, in Unigine Heaven. All Gen6+ platforms had similar results. (Ice Lake shown) total instructions in shared programs: 16278207 -> 16278465 (<.01%) instructions in affected programs: 11374 -> 11632 (2.27%) helped: 0 HURT: 58 HURT stats (abs) min: 2 max: 13 x̄: 4.45 x̃: 4 HURT stats (rel) min: 0.54% max: 4.11% x̄: 2.42% x̃: 2.82% 95% mean confidence interval for instructions value: 3.77 5.13 95% mean confidence interval for instructions %-change: 2.19% 2.64% Instructions are HURT. total cycles in shared programs: 367134284 -> 367135159 (<.01%) cycles in affected programs: 81207 -> 82082 (1.08%) helped: 17 HURT: 36 helped stats (abs) min: 6 max: 356 x̄: 90.35 x̃: 6 helped stats (rel) min: 0.69% max: 21.45% x̄: 5.71% x̃: 0.78% HURT stats (abs) min: 4 max: 235 x̄: 66.97 x̃: 16 HURT stats (rel) min: 0.35% max: 27.58% x̄: 5.34% x̃: 1.09% 95% mean confidence interval for cycles value: -20.36 53.38 95% mean confidence interval for cycles %-change: -1.08% 4.67% Inconclusive result (value mean confidence interval includes 0). No changes on any earlier platforms.
2019-08-07 08:56:22 -07:00
case nir_op_fexp2: {
fp_class_mask src = src_res[0];
r = 0;
nir/range-analysis: Adjust result range of exp2 to account for flush-to-zero Fixes piglit tests (new in piglit!110): - fs-underflow-exp2-compare-zero.shader_test Bugzilla: https://bugs.freedesktop.org/show_bug.cgi?id=111308 Fixes: 405de7ccb6c ("nir/range-analysis: Rudimentary value range analysis pass") Reviewed-by: Caio Marcelo de Oliveira Filho <caio.oliveira@intel.com> Most of the shaders affected are, unsurprisingly, in Unigine Heaven. All Gen6+ platforms had similar results. (Ice Lake shown) total instructions in shared programs: 16278207 -> 16278465 (<.01%) instructions in affected programs: 11374 -> 11632 (2.27%) helped: 0 HURT: 58 HURT stats (abs) min: 2 max: 13 x̄: 4.45 x̃: 4 HURT stats (rel) min: 0.54% max: 4.11% x̄: 2.42% x̃: 2.82% 95% mean confidence interval for instructions value: 3.77 5.13 95% mean confidence interval for instructions %-change: 2.19% 2.64% Instructions are HURT. total cycles in shared programs: 367134284 -> 367135159 (<.01%) cycles in affected programs: 81207 -> 82082 (1.08%) helped: 17 HURT: 36 helped stats (abs) min: 6 max: 356 x̄: 90.35 x̃: 6 helped stats (rel) min: 0.69% max: 21.45% x̄: 5.71% x̃: 0.78% HURT stats (abs) min: 4 max: 235 x̄: 66.97 x̃: 16 HURT stats (rel) min: 0.35% max: 27.58% x̄: 5.34% x̃: 1.09% 95% mean confidence interval for cycles value: -20.36 53.38 95% mean confidence interval for cycles %-change: -1.08% 4.67% Inconclusive result (value mean confidence interval includes 0). No changes on any earlier platforms.
2019-08-07 08:56:22 -07:00
/* If the parameter might be less than zero, the mathematically result
* will be on (0, 1). For sufficiently large magnitude negative
* parameters, the result will flush to zero.
*/
if (src & FP_CLASS_NEG_INF)
r |= FP_CLASS_POS_ZERO;
nir/range-analysis: Adjust result range of exp2 to account for flush-to-zero Fixes piglit tests (new in piglit!110): - fs-underflow-exp2-compare-zero.shader_test Bugzilla: https://bugs.freedesktop.org/show_bug.cgi?id=111308 Fixes: 405de7ccb6c ("nir/range-analysis: Rudimentary value range analysis pass") Reviewed-by: Caio Marcelo de Oliveira Filho <caio.oliveira@intel.com> Most of the shaders affected are, unsurprisingly, in Unigine Heaven. All Gen6+ platforms had similar results. (Ice Lake shown) total instructions in shared programs: 16278207 -> 16278465 (<.01%) instructions in affected programs: 11374 -> 11632 (2.27%) helped: 0 HURT: 58 HURT stats (abs) min: 2 max: 13 x̄: 4.45 x̃: 4 HURT stats (rel) min: 0.54% max: 4.11% x̄: 2.42% x̃: 2.82% 95% mean confidence interval for instructions value: 3.77 5.13 95% mean confidence interval for instructions %-change: 2.19% 2.64% Instructions are HURT. total cycles in shared programs: 367134284 -> 367135159 (<.01%) cycles in affected programs: 81207 -> 82082 (1.08%) helped: 17 HURT: 36 helped stats (abs) min: 6 max: 356 x̄: 90.35 x̃: 6 helped stats (rel) min: 0.69% max: 21.45% x̄: 5.71% x̃: 0.78% HURT stats (abs) min: 4 max: 235 x̄: 66.97 x̃: 16 HURT stats (rel) min: 0.35% max: 27.58% x̄: 5.34% x̃: 1.09% 95% mean confidence interval for cycles value: -20.36 53.38 95% mean confidence interval for cycles %-change: -1.08% 4.67% Inconclusive result (value mean confidence interval includes 0). No changes on any earlier platforms.
2019-08-07 08:56:22 -07:00
if (src & FP_CLASS_LT_NEG_ONE)
r |= FP_CLASS_POS_ZERO | FP_CLASS_GT_ZERO_LT_POS_ONE | FP_CLASS_NON_INTEGRAL;
nir/range-analysis: Adjust result range of exp2 to account for flush-to-zero Fixes piglit tests (new in piglit!110): - fs-underflow-exp2-compare-zero.shader_test Bugzilla: https://bugs.freedesktop.org/show_bug.cgi?id=111308 Fixes: 405de7ccb6c ("nir/range-analysis: Rudimentary value range analysis pass") Reviewed-by: Caio Marcelo de Oliveira Filho <caio.oliveira@intel.com> Most of the shaders affected are, unsurprisingly, in Unigine Heaven. All Gen6+ platforms had similar results. (Ice Lake shown) total instructions in shared programs: 16278207 -> 16278465 (<.01%) instructions in affected programs: 11374 -> 11632 (2.27%) helped: 0 HURT: 58 HURT stats (abs) min: 2 max: 13 x̄: 4.45 x̃: 4 HURT stats (rel) min: 0.54% max: 4.11% x̄: 2.42% x̃: 2.82% 95% mean confidence interval for instructions value: 3.77 5.13 95% mean confidence interval for instructions %-change: 2.19% 2.64% Instructions are HURT. total cycles in shared programs: 367134284 -> 367135159 (<.01%) cycles in affected programs: 81207 -> 82082 (1.08%) helped: 17 HURT: 36 helped stats (abs) min: 6 max: 356 x̄: 90.35 x̃: 6 helped stats (rel) min: 0.69% max: 21.45% x̄: 5.71% x̃: 0.78% HURT stats (abs) min: 4 max: 235 x̄: 66.97 x̃: 16 HURT stats (rel) min: 0.35% max: 27.58% x̄: 5.34% x̃: 1.09% 95% mean confidence interval for cycles value: -20.36 53.38 95% mean confidence interval for cycles %-change: -1.08% 4.67% Inconclusive result (value mean confidence interval includes 0). No changes on any earlier platforms.
2019-08-07 08:56:22 -07:00
if (src & (FP_CLASS_NEG_ONE | FP_CLASS_LT_ZERO_GT_NEG_ONE))
r |= FP_CLASS_GT_ZERO_LT_POS_ONE | FP_CLASS_NON_INTEGRAL;
if (src & (FP_CLASS_LT_ZERO_GT_NEG_ONE | FP_CLASS_ANY_ZERO | FP_CLASS_GT_ZERO_LT_POS_ONE))
r |= FP_CLASS_POS_ONE;
if (src & (FP_CLASS_GT_ZERO_LT_POS_ONE | FP_CLASS_POS_ONE))
r |= FP_CLASS_GT_POS_ONE;
if (src & (FP_CLASS_GT_POS_ONE))
r |= FP_CLASS_GT_POS_ONE | FP_CLASS_POS_INF;
if (src & FP_CLASS_POS_INF)
r |= FP_CLASS_POS_INF;
if (src & FP_CLASS_NON_INTEGRAL)
r |= FP_CLASS_NON_INTEGRAL;
nir/range_analysis: Fix analysis of fmin, fmax, or fsat with NaN source Recall that when either value is NaN, fmax will pick the other value. This means the result range of the fmax will either be the "ideal" result range (calculated above) or the range of the non-NaN value. Previously, something like fmax({gt_zero}, {lt_zero, is_a_number}) would return a range of gt_zero. However, if the "gt_zero" parameter is NaN, the actual result will be the "lt_zero" parameter. This analysis depends on the is_a_number analysis also added in this MR. Assuming this doesn't cause any unforeseen problems, I believe we should wait a bit, then nominate a subset of the series for the stable branches. This fixes the piglit tests tests/spec/glsl-1.30/execution/range_analysis_fmax_of_nan.shader_test tests/spec/glsl-1.30/execution/range_analysis_fmin_of_nan.shader_test from https://gitlab.freedesktop.org/mesa/piglit/-/merge_requests/463. Even with the added fsat fixes, range_analysis_fsat_of_nan.shader_test still fails. There are some other issues there that will be addressed in later commits (in another MR). v2: Add fsat fixes. Suggested by Rhys. Fixes: 405de7ccb6c ("nir/range-analysis: Rudimentary value range analysis pass") Reviewed-by: Rhys Perry <pendingchaos02@gmail.com> Shader-db results: All Intel platforms had similar results. (Tiger Lake shown) total instructions in shared programs: 21049290 -> 21049314 (<.01%) instructions in affected programs: 3175 -> 3199 (0.76%) helped: 0 HURT: 17 HURT stats (abs) min: 1 max: 3 x̄: 1.41 x̃: 1 HURT stats (rel) min: 0.20% max: 1.89% x̄: 0.97% x̃: 0.92% 95% mean confidence interval for instructions value: 1.09 1.73 95% mean confidence interval for instructions %-change: 0.75% 1.19% Instructions are HURT. total cycles in shared programs: 855136176 -> 855136406 (<.01%) cycles in affected programs: 37579 -> 37809 (0.61%) helped: 0 HURT: 17 HURT stats (abs) min: 12 max: 20 x̄: 13.53 x̃: 14 HURT stats (rel) min: 0.17% max: 1.13% x̄: 0.79% x̃: 0.91% 95% mean confidence interval for cycles value: 12.53 14.53 95% mean confidence interval for cycles %-change: 0.63% 0.94% Cycles are HURT. Fossil-db results: Tiger Lake Instructions in all programs: 160901033 -> 160902591 (+0.0%) SENDs in all programs: 6812270 -> 6812270 (+0.0%) Loops in all programs: 38225 -> 38225 (+0.0%) Cycles in all programs: 7430016795 -> 7429003266 (-0.0%) Spills in all programs: 192582 -> 192582 (+0.0%) Fills in all programs: 304539 -> 304539 (+0.0%) Ice Lake Instructions in all programs: 145299102 -> 145301634 (+0.0%) SENDs in all programs: 6863890 -> 6863890 (+0.0%) Loops in all programs: 38219 -> 38219 (+0.0%) Cycles in all programs: 8798390846 -> 8798589772 (+0.0%) Spills in all programs: 216880 -> 216880 (+0.0%) Fills in all programs: 334250 -> 334250 (+0.0%) Skylake Instructions in all programs: 135889478 -> 135892010 (+0.0%) SENDs in all programs: 6802916 -> 6802916 (+0.0%) Loops in all programs: 38216 -> 38216 (+0.0%) Cycles in all programs: 8442624166 -> 8442597324 (-0.0%) Spills in all programs: 194839 -> 194839 (+0.0%) Fills in all programs: 301116 -> 301116 (+0.0%) Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/9108>
2021-01-27 19:42:44 -08:00
if (src & FP_CLASS_NAN)
r |= FP_CLASS_NAN;
break;
}
case nir_op_flog2: {
r = 0;
if (src_res[0] & (FP_CLASS_ANY_NEG | FP_CLASS_NAN))
r |= FP_CLASS_NAN;
if (src_res[0] & FP_CLASS_ANY_ZERO)
r |= FP_CLASS_NEG_INF;
if (src_res[0] & FP_CLASS_GT_ZERO_LT_POS_ONE)
r |= FP_CLASS_ANY_NEG | FP_CLASS_NON_INTEGRAL;
if (src_res[0] & FP_CLASS_POS_ONE)
r |= FP_CLASS_POS_ZERO;
if (src_res[0] & FP_CLASS_GT_POS_ONE)
r |= FP_CLASS_GT_ZERO_LT_POS_ONE | FP_CLASS_POS_ONE | FP_CLASS_GT_POS_ONE | FP_CLASS_NON_INTEGRAL;
if (src_res[0] & FP_CLASS_POS_INF)
r |= FP_CLASS_POS_INF;
break;
}
case nir_op_fmax: {
fp_class_mask left = fneg_fp_class(src_res[0]);
fp_class_mask right = fneg_fp_class(src_res[1]);
r = fneg_fp_class(fmin_fp_class(left, right));
break;
}
case nir_op_fmin:
r = fmin_fp_class(src_res[0], src_res[1]);
break;
case nir_op_fmul:
case nir_op_fmulz: {
bool mulz = alu->op == nir_op_fmulz;
bool src_eq = nir_alu_srcs_equal(alu, alu, 0, 1);
bool src_neg_eq = !nir_src_is_const(alu->src[0].src) && nir_alu_srcs_negative_equal(alu, alu, 0, 1);
r = fmul_fp_class(src_res[0], src_res[1], mulz, src_eq, src_neg_eq);
break;
}
case nir_op_frcp:
r = frcp_fp_class(handle_sz(alu, src_res[0]));
break;
nir/range-analysis: Use types to provide better ranges from bcsel and mov Reviewed-by: Caio Marcelo de Oliveira Filho <caio.oliveira@intel.com> All Gen7+ platforms had similar results. (Ice Lake shown) total instructions in shared programs: 16328255 -> 16315391 (-0.08%) instructions in affected programs: 218318 -> 205454 (-5.89%) helped: 988 HURT: 0 helped stats (abs) min: 1 max: 72 x̄: 13.02 x̃: 10 helped stats (rel) min: 0.33% max: 16.04% x̄: 6.27% x̃: 4.88% 95% mean confidence interval for instructions value: -13.69 -12.35 95% mean confidence interval for instructions %-change: -6.55% -5.99% Instructions are helped. total cycles in shared programs: 363683977 -> 363615417 (-0.02%) cycles in affected programs: 1475193 -> 1406633 (-4.65%) helped: 923 HURT: 36 helped stats (abs) min: 1 max: 624 x̄: 75.78 x̃: 48 helped stats (rel) min: 0.08% max: 13.89% x̄: 5.20% x̃: 5.08% HURT stats (abs) min: 1 max: 179 x̄: 38.58 x̃: 4 HURT stats (rel) min: 0.06% max: 16.56% x̄: 3.33% x̃: 0.29% 95% mean confidence interval for cycles value: -75.88 -67.10 95% mean confidence interval for cycles %-change: -5.10% -4.66% Cycles are helped. Sandy Bridge total instructions in shared programs: 10785779 -> 10785654 (<.01%) instructions in affected programs: 13855 -> 13730 (-0.90%) helped: 67 HURT: 0 helped stats (abs) min: 1 max: 15 x̄: 1.87 x̃: 1 helped stats (rel) min: 0.20% max: 3.45% x̄: 0.97% x̃: 0.78% 95% mean confidence interval for instructions value: -2.47 -1.26 95% mean confidence interval for instructions %-change: -1.13% -0.81% Instructions are helped. total cycles in shared programs: 153704799 -> 153704481 (<.01%) cycles in affected programs: 101509 -> 101191 (-0.31%) helped: 38 HURT: 13 helped stats (abs) min: 1 max: 38 x̄: 12.53 x̃: 16 helped stats (rel) min: 0.07% max: 2.69% x̄: 0.87% x̃: 0.53% HURT stats (abs) min: 1 max: 36 x̄: 12.15 x̃: 7 HURT stats (rel) min: 0.06% max: 2.53% x̄: 0.73% x̃: 0.44% 95% mean confidence interval for cycles value: -10.24 -2.24 95% mean confidence interval for cycles %-change: -0.75% -0.17% Cycles are helped. LOST: 2 GAINED: 0 No shader-db change on Iron Lake or GM45.
2019-08-12 17:28:35 -07:00
case nir_op_mov:
r = src_res[0];
break;
case nir_op_fsat: {
r = src_res[0];
/* max(+0.0, x) */
if (r & (FP_CLASS_ANY_NEG | FP_CLASS_NEG_ZERO | FP_CLASS_NAN)) {
r &= ~(FP_CLASS_ANY_NEG | FP_CLASS_NEG_ZERO | FP_CLASS_NAN);
r |= FP_CLASS_POS_ZERO;
}
/* min(+1.0, x) */
if (r & (FP_CLASS_GT_POS_ONE | FP_CLASS_POS_INF)) {
r &= ~(FP_CLASS_GT_POS_ONE | FP_CLASS_POS_INF);
r |= FP_CLASS_POS_ONE;
}
if (!(r & FP_CLASS_GT_ZERO_LT_POS_ONE))
r &= ~FP_CLASS_NON_INTEGRAL;
break;
}
case nir_op_fsign:
r = 0;
if (src_res[0] & FP_CLASS_ANY_NEG)
r |= FP_CLASS_NEG_ONE;
if (src_res[0] & FP_CLASS_ANY_ZERO)
r |= FP_CLASS_ANY_ZERO;
if (src_res[0] & FP_CLASS_ANY_POS)
r |= FP_CLASS_POS_ONE;
/* fsign is -1, 0, or 1, even for NaN */
if (src_res[0] & FP_CLASS_NAN)
r |= FP_CLASS_NEG_ONE | FP_CLASS_ANY_ZERO | FP_CLASS_POS_ONE;
break;
case nir_op_fsqrt:
r = fsqrt_fp_class(src_res[0]);
break;
case nir_op_frsq:
r = frcp_fp_class(fsqrt_fp_class(handle_sz(alu, src_res[0])));
break;
case nir_op_ffloor: {
/* In IEEE 754, floor(NaN) is NaN, and floor(±Inf) is ±Inf. See
* https://pubs.opengroup.org/onlinepubs/9699919799.2016edition/functions/floor.html
*/
r = src_res[0];
if (r & FP_CLASS_NON_INTEGRAL) {
if (r & FP_CLASS_LT_ZERO_GT_NEG_ONE)
r |= FP_CLASS_NEG_ONE;
if (r & FP_CLASS_GT_ZERO_LT_POS_ONE)
r |= FP_CLASS_POS_ZERO;
if (r & FP_CLASS_GT_POS_ONE)
r |= FP_CLASS_POS_ONE;
r &= ~(FP_CLASS_NON_INTEGRAL | FP_CLASS_LT_ZERO_GT_NEG_ONE | FP_CLASS_GT_ZERO_LT_POS_ONE);
}
break;
}
case nir_op_fceil: {
/* In IEEE 754, ceil(NaN) is NaN, and ceil(±Inf) is ±Inf. See
* https://pubs.opengroup.org/onlinepubs/9699919799.2016edition/functions/ceil.html
*/
r = src_res[0];
if (r & FP_CLASS_NON_INTEGRAL) {
if (r & FP_CLASS_LT_NEG_ONE)
r |= FP_CLASS_NEG_ONE;
if (r & FP_CLASS_LT_ZERO_GT_NEG_ONE)
r |= FP_CLASS_NEG_ZERO;
if (r & FP_CLASS_GT_ZERO_LT_POS_ONE)
r |= FP_CLASS_POS_ONE;
r &= ~(FP_CLASS_NON_INTEGRAL | FP_CLASS_LT_ZERO_GT_NEG_ONE | FP_CLASS_GT_ZERO_LT_POS_ONE);
}
break;
}
case nir_op_ftrunc: {
/* In IEEE 754, trunc(NaN) is NaN, and trunc(±Inf) is ±Inf. See
* https://pubs.opengroup.org/onlinepubs/9699919799.2016edition/functions/trunc.html
*/
r = src_res[0];
if (r & FP_CLASS_NON_INTEGRAL) {
if (r & FP_CLASS_LT_NEG_ONE)
r |= FP_CLASS_NEG_ONE;
if (r & FP_CLASS_LT_ZERO_GT_NEG_ONE)
r |= FP_CLASS_NEG_ZERO;
if (r & FP_CLASS_GT_ZERO_LT_POS_ONE)
r |= FP_CLASS_POS_ZERO;
if (r & FP_CLASS_GT_POS_ONE)
r |= FP_CLASS_POS_ONE;
r &= ~(FP_CLASS_NON_INTEGRAL | FP_CLASS_LT_ZERO_GT_NEG_ONE | FP_CLASS_GT_ZERO_LT_POS_ONE);
}
break;
}
case nir_op_fround_even: {
r = src_res[0];
if (r & FP_CLASS_NON_INTEGRAL) {
if (r & FP_CLASS_LT_NEG_ONE)
r |= FP_CLASS_NEG_ONE;
if (r & FP_CLASS_LT_ZERO_GT_NEG_ONE)
r |= FP_CLASS_NEG_ZERO | FP_CLASS_NEG_ONE;
if (r & FP_CLASS_GT_ZERO_LT_POS_ONE)
r |= FP_CLASS_POS_ZERO | FP_CLASS_POS_ONE;
if (r & FP_CLASS_GT_POS_ONE)
r |= FP_CLASS_POS_ONE;
r &= ~(FP_CLASS_NON_INTEGRAL | FP_CLASS_LT_ZERO_GT_NEG_ONE | FP_CLASS_GT_ZERO_LT_POS_ONE);
}
break;
}
case nir_op_ffract: {
r = 0;
/* fract(±Inf) is NaN. */
if (src_res[0] & (FP_CLASS_ANY_INF | FP_CLASS_NAN))
r |= FP_CLASS_NAN;
/* fract(non_integral) is in (0, 1). */
if (src_res[0] & FP_CLASS_NON_INTEGRAL)
r |= FP_CLASS_GT_ZERO_LT_POS_ONE | FP_CLASS_NON_INTEGRAL;
/* fract(small, negative) can be 1.0. */
if (src_res[0] & FP_CLASS_LT_ZERO_GT_NEG_ONE)
r |= FP_CLASS_POS_ONE;
/* fract(integral) is +0.0. */
if (src_res[0] & (FP_CLASS_LT_NEG_ONE | FP_CLASS_NEG_ONE | FP_CLASS_ANY_ZERO | FP_CLASS_POS_ONE | FP_CLASS_GT_POS_ONE))
r |= FP_CLASS_POS_ZERO;
break;
}
case nir_op_fsin:
case nir_op_fcos:
case nir_op_fsin_amd:
case nir_op_fcos_amd: {
/* [-1, +1], and sin/cos(Inf) is NaN */
r = FP_CLASS_NEG_ONE | FP_CLASS_LT_ZERO_GT_NEG_ONE | FP_CLASS_ANY_ZERO |
FP_CLASS_GT_ZERO_LT_POS_ONE | FP_CLASS_POS_ONE | FP_CLASS_NON_INTEGRAL;
if (src_res[0] & (FP_CLASS_NAN | FP_CLASS_ANY_INF))
r |= FP_CLASS_NAN;
break;
}
nir/range_analysis: Teach range analysis about fdot opcodes This really, really helps on platforms where fabs() isn't free. A great many shaders use a * frsq(fabs(fdot(a, a))) to normalize a vector. Since the result of the fdot must be non-negative, the fabs can be eliminated by an existing algebraic rule. shader-db results: r300 (run on R420 - X800XL) total instructions in shared programs: 1369807 -> 1368550 (-0.09%) instructions in affected programs: 59986 -> 58729 (-2.10%) helped: 609 HURT: 0 total vinst in shared programs: 512899 -> 512861 (<.01%) vinst in affected programs: 1522 -> 1484 (-2.50%) helped: 36 HURT: 0 total sinst in shared programs: 260690 -> 260570 (-0.05%) sinst in affected programs: 1419 -> 1299 (-8.46%) helped: 120 HURT: 0 total consts in shared programs: 957295 -> 957230 (<.01%) consts in affected programs: 849 -> 784 (-7.66%) helped: 65 HURT: 0 LOST: 0 GAINED: 3 The 3 gained shaders are all vertex shaders from XCom: Enemy Unknown. I'm guessing that game is never going to run on my X800XL. :) i915 total instructions in shared programs: 791121 -> 780843 (-1.30%) instructions in affected programs: 220170 -> 209892 (-4.67%) helped: 2085 HURT: 0 total temps in shared programs: 47765 -> 47766 (<.01%) temps in affected programs: 9 -> 10 (11.11%) helped: 0 HURT: 1 total const in shared programs: 93048 -> 92983 (-0.07%) const in affected programs: 784 -> 719 (-8.29%) helped: 65 HURT: 0 LOST: 0 GAINED: 36 Haswell, Ivy Bridge, and Sandy Bridge had similar results. (Haswell shown) total instructions in shared programs: 16702250 -> 16697908 (-0.03%) instructions in affected programs: 119277 -> 114935 (-3.64%) helped: 1065 HURT: 0 helped stats (abs) min: 1 max: 20 x̄: 4.08 x̃: 4 helped stats (rel) min: 0.48% max: 10.17% x̄: 3.66% x̃: 3.94% 95% mean confidence interval for instructions value: -4.26 -3.89 95% mean confidence interval for instructions %-change: -3.76% -3.56% Instructions are helped. total cycles in shared programs: 880772068 -> 880734134 (<.01%) cycles in affected programs: 2134456 -> 2096522 (-1.78%) helped: 941 HURT: 324 helped stats (abs) min: 2 max: 2180 x̄: 123.06 x̃: 44 helped stats (rel) min: 0.04% max: 49.96% x̄: 7.08% x̃: 3.81% HURT stats (abs) min: 2 max: 2098 x̄: 240.33 x̃: 35 HURT stats (rel) min: 0.04% max: 77.07% x̄: 12.34% x̃: 3.00% 95% mean confidence interval for cycles value: -47.93 -12.04 95% mean confidence interval for cycles %-change: -2.87% -1.34% Cycles are helped. No shader-db changes on any other Intel platform. Reviewed-by: Jason Ekstrand <jason.ekstrand@collabora.com> Reviewed-by: Emma Anholt <emma@anholt.net> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/17181>
2022-06-21 16:47:31 -07:00
case nir_op_fdot2:
case nir_op_fdot3:
case nir_op_fdot4:
case nir_op_fdot8:
case nir_op_fdot16:
case nir_op_fdot2_replicated:
case nir_op_fdot3_replicated:
case nir_op_fdot4_replicated:
case nir_op_fdot8_replicated:
case nir_op_fdot16_replicated: {
/* If the two sources are the same SSA value, then the result is either
* NaN or some number >= 0. If one source is the negation of the other,
* the result is either NaN or some number <= 0.
*
* In either of these two cases, if one source is a number, then the
* other must also be a number. Since it should not be possible to get
* Inf-Inf in the dot-product, the result must also be a number.
*/
if (nir_alu_srcs_equal(alu, alu, 0, 1)) {
r = FP_CLASS_ANY_POS | FP_CLASS_POS_ZERO;
nir/range_analysis: Teach range analysis about fdot opcodes This really, really helps on platforms where fabs() isn't free. A great many shaders use a * frsq(fabs(fdot(a, a))) to normalize a vector. Since the result of the fdot must be non-negative, the fabs can be eliminated by an existing algebraic rule. shader-db results: r300 (run on R420 - X800XL) total instructions in shared programs: 1369807 -> 1368550 (-0.09%) instructions in affected programs: 59986 -> 58729 (-2.10%) helped: 609 HURT: 0 total vinst in shared programs: 512899 -> 512861 (<.01%) vinst in affected programs: 1522 -> 1484 (-2.50%) helped: 36 HURT: 0 total sinst in shared programs: 260690 -> 260570 (-0.05%) sinst in affected programs: 1419 -> 1299 (-8.46%) helped: 120 HURT: 0 total consts in shared programs: 957295 -> 957230 (<.01%) consts in affected programs: 849 -> 784 (-7.66%) helped: 65 HURT: 0 LOST: 0 GAINED: 3 The 3 gained shaders are all vertex shaders from XCom: Enemy Unknown. I'm guessing that game is never going to run on my X800XL. :) i915 total instructions in shared programs: 791121 -> 780843 (-1.30%) instructions in affected programs: 220170 -> 209892 (-4.67%) helped: 2085 HURT: 0 total temps in shared programs: 47765 -> 47766 (<.01%) temps in affected programs: 9 -> 10 (11.11%) helped: 0 HURT: 1 total const in shared programs: 93048 -> 92983 (-0.07%) const in affected programs: 784 -> 719 (-8.29%) helped: 65 HURT: 0 LOST: 0 GAINED: 36 Haswell, Ivy Bridge, and Sandy Bridge had similar results. (Haswell shown) total instructions in shared programs: 16702250 -> 16697908 (-0.03%) instructions in affected programs: 119277 -> 114935 (-3.64%) helped: 1065 HURT: 0 helped stats (abs) min: 1 max: 20 x̄: 4.08 x̃: 4 helped stats (rel) min: 0.48% max: 10.17% x̄: 3.66% x̃: 3.94% 95% mean confidence interval for instructions value: -4.26 -3.89 95% mean confidence interval for instructions %-change: -3.76% -3.56% Instructions are helped. total cycles in shared programs: 880772068 -> 880734134 (<.01%) cycles in affected programs: 2134456 -> 2096522 (-1.78%) helped: 941 HURT: 324 helped stats (abs) min: 2 max: 2180 x̄: 123.06 x̃: 44 helped stats (rel) min: 0.04% max: 49.96% x̄: 7.08% x̃: 3.81% HURT stats (abs) min: 2 max: 2098 x̄: 240.33 x̃: 35 HURT stats (rel) min: 0.04% max: 77.07% x̄: 12.34% x̃: 3.00% 95% mean confidence interval for cycles value: -47.93 -12.04 95% mean confidence interval for cycles %-change: -2.87% -1.34% Cycles are helped. No shader-db changes on any other Intel platform. Reviewed-by: Jason Ekstrand <jason.ekstrand@collabora.com> Reviewed-by: Emma Anholt <emma@anholt.net> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/17181>
2022-06-21 16:47:31 -07:00
} else if (nir_alu_srcs_negative_equal(alu, alu, 0, 1)) {
r = FP_CLASS_ANY_NEG | FP_CLASS_NEG_ZERO;
nir/range_analysis: Teach range analysis about fdot opcodes This really, really helps on platforms where fabs() isn't free. A great many shaders use a * frsq(fabs(fdot(a, a))) to normalize a vector. Since the result of the fdot must be non-negative, the fabs can be eliminated by an existing algebraic rule. shader-db results: r300 (run on R420 - X800XL) total instructions in shared programs: 1369807 -> 1368550 (-0.09%) instructions in affected programs: 59986 -> 58729 (-2.10%) helped: 609 HURT: 0 total vinst in shared programs: 512899 -> 512861 (<.01%) vinst in affected programs: 1522 -> 1484 (-2.50%) helped: 36 HURT: 0 total sinst in shared programs: 260690 -> 260570 (-0.05%) sinst in affected programs: 1419 -> 1299 (-8.46%) helped: 120 HURT: 0 total consts in shared programs: 957295 -> 957230 (<.01%) consts in affected programs: 849 -> 784 (-7.66%) helped: 65 HURT: 0 LOST: 0 GAINED: 3 The 3 gained shaders are all vertex shaders from XCom: Enemy Unknown. I'm guessing that game is never going to run on my X800XL. :) i915 total instructions in shared programs: 791121 -> 780843 (-1.30%) instructions in affected programs: 220170 -> 209892 (-4.67%) helped: 2085 HURT: 0 total temps in shared programs: 47765 -> 47766 (<.01%) temps in affected programs: 9 -> 10 (11.11%) helped: 0 HURT: 1 total const in shared programs: 93048 -> 92983 (-0.07%) const in affected programs: 784 -> 719 (-8.29%) helped: 65 HURT: 0 LOST: 0 GAINED: 36 Haswell, Ivy Bridge, and Sandy Bridge had similar results. (Haswell shown) total instructions in shared programs: 16702250 -> 16697908 (-0.03%) instructions in affected programs: 119277 -> 114935 (-3.64%) helped: 1065 HURT: 0 helped stats (abs) min: 1 max: 20 x̄: 4.08 x̃: 4 helped stats (rel) min: 0.48% max: 10.17% x̄: 3.66% x̃: 3.94% 95% mean confidence interval for instructions value: -4.26 -3.89 95% mean confidence interval for instructions %-change: -3.76% -3.56% Instructions are helped. total cycles in shared programs: 880772068 -> 880734134 (<.01%) cycles in affected programs: 2134456 -> 2096522 (-1.78%) helped: 941 HURT: 324 helped stats (abs) min: 2 max: 2180 x̄: 123.06 x̃: 44 helped stats (rel) min: 0.04% max: 49.96% x̄: 7.08% x̃: 3.81% HURT stats (abs) min: 2 max: 2098 x̄: 240.33 x̃: 35 HURT stats (rel) min: 0.04% max: 77.07% x̄: 12.34% x̃: 3.00% 95% mean confidence interval for cycles value: -47.93 -12.04 95% mean confidence interval for cycles %-change: -2.87% -1.34% Cycles are helped. No shader-db changes on any other Intel platform. Reviewed-by: Jason Ekstrand <jason.ekstrand@collabora.com> Reviewed-by: Emma Anholt <emma@anholt.net> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/17181>
2022-06-21 16:47:31 -07:00
} else {
r = FP_CLASS_UNKNOWN;
nir/range_analysis: Teach range analysis about fdot opcodes This really, really helps on platforms where fabs() isn't free. A great many shaders use a * frsq(fabs(fdot(a, a))) to normalize a vector. Since the result of the fdot must be non-negative, the fabs can be eliminated by an existing algebraic rule. shader-db results: r300 (run on R420 - X800XL) total instructions in shared programs: 1369807 -> 1368550 (-0.09%) instructions in affected programs: 59986 -> 58729 (-2.10%) helped: 609 HURT: 0 total vinst in shared programs: 512899 -> 512861 (<.01%) vinst in affected programs: 1522 -> 1484 (-2.50%) helped: 36 HURT: 0 total sinst in shared programs: 260690 -> 260570 (-0.05%) sinst in affected programs: 1419 -> 1299 (-8.46%) helped: 120 HURT: 0 total consts in shared programs: 957295 -> 957230 (<.01%) consts in affected programs: 849 -> 784 (-7.66%) helped: 65 HURT: 0 LOST: 0 GAINED: 3 The 3 gained shaders are all vertex shaders from XCom: Enemy Unknown. I'm guessing that game is never going to run on my X800XL. :) i915 total instructions in shared programs: 791121 -> 780843 (-1.30%) instructions in affected programs: 220170 -> 209892 (-4.67%) helped: 2085 HURT: 0 total temps in shared programs: 47765 -> 47766 (<.01%) temps in affected programs: 9 -> 10 (11.11%) helped: 0 HURT: 1 total const in shared programs: 93048 -> 92983 (-0.07%) const in affected programs: 784 -> 719 (-8.29%) helped: 65 HURT: 0 LOST: 0 GAINED: 36 Haswell, Ivy Bridge, and Sandy Bridge had similar results. (Haswell shown) total instructions in shared programs: 16702250 -> 16697908 (-0.03%) instructions in affected programs: 119277 -> 114935 (-3.64%) helped: 1065 HURT: 0 helped stats (abs) min: 1 max: 20 x̄: 4.08 x̃: 4 helped stats (rel) min: 0.48% max: 10.17% x̄: 3.66% x̃: 3.94% 95% mean confidence interval for instructions value: -4.26 -3.89 95% mean confidence interval for instructions %-change: -3.76% -3.56% Instructions are helped. total cycles in shared programs: 880772068 -> 880734134 (<.01%) cycles in affected programs: 2134456 -> 2096522 (-1.78%) helped: 941 HURT: 324 helped stats (abs) min: 2 max: 2180 x̄: 123.06 x̃: 44 helped stats (rel) min: 0.04% max: 49.96% x̄: 7.08% x̃: 3.81% HURT stats (abs) min: 2 max: 2098 x̄: 240.33 x̃: 35 HURT stats (rel) min: 0.04% max: 77.07% x̄: 12.34% x̃: 3.00% 95% mean confidence interval for cycles value: -47.93 -12.04 95% mean confidence interval for cycles %-change: -2.87% -1.34% Cycles are helped. No shader-db changes on any other Intel platform. Reviewed-by: Jason Ekstrand <jason.ekstrand@collabora.com> Reviewed-by: Emma Anholt <emma@anholt.net> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/17181>
2022-06-21 16:47:31 -07:00
}
if (src_res[0] & FP_CLASS_NAN)
r |= FP_CLASS_NAN;
if (src_res[0] & FP_CLASS_NON_INTEGRAL)
r |= FP_CLASS_NON_INTEGRAL;
nir/range_analysis: Teach range analysis about fdot opcodes This really, really helps on platforms where fabs() isn't free. A great many shaders use a * frsq(fabs(fdot(a, a))) to normalize a vector. Since the result of the fdot must be non-negative, the fabs can be eliminated by an existing algebraic rule. shader-db results: r300 (run on R420 - X800XL) total instructions in shared programs: 1369807 -> 1368550 (-0.09%) instructions in affected programs: 59986 -> 58729 (-2.10%) helped: 609 HURT: 0 total vinst in shared programs: 512899 -> 512861 (<.01%) vinst in affected programs: 1522 -> 1484 (-2.50%) helped: 36 HURT: 0 total sinst in shared programs: 260690 -> 260570 (-0.05%) sinst in affected programs: 1419 -> 1299 (-8.46%) helped: 120 HURT: 0 total consts in shared programs: 957295 -> 957230 (<.01%) consts in affected programs: 849 -> 784 (-7.66%) helped: 65 HURT: 0 LOST: 0 GAINED: 3 The 3 gained shaders are all vertex shaders from XCom: Enemy Unknown. I'm guessing that game is never going to run on my X800XL. :) i915 total instructions in shared programs: 791121 -> 780843 (-1.30%) instructions in affected programs: 220170 -> 209892 (-4.67%) helped: 2085 HURT: 0 total temps in shared programs: 47765 -> 47766 (<.01%) temps in affected programs: 9 -> 10 (11.11%) helped: 0 HURT: 1 total const in shared programs: 93048 -> 92983 (-0.07%) const in affected programs: 784 -> 719 (-8.29%) helped: 65 HURT: 0 LOST: 0 GAINED: 36 Haswell, Ivy Bridge, and Sandy Bridge had similar results. (Haswell shown) total instructions in shared programs: 16702250 -> 16697908 (-0.03%) instructions in affected programs: 119277 -> 114935 (-3.64%) helped: 1065 HURT: 0 helped stats (abs) min: 1 max: 20 x̄: 4.08 x̃: 4 helped stats (rel) min: 0.48% max: 10.17% x̄: 3.66% x̃: 3.94% 95% mean confidence interval for instructions value: -4.26 -3.89 95% mean confidence interval for instructions %-change: -3.76% -3.56% Instructions are helped. total cycles in shared programs: 880772068 -> 880734134 (<.01%) cycles in affected programs: 2134456 -> 2096522 (-1.78%) helped: 941 HURT: 324 helped stats (abs) min: 2 max: 2180 x̄: 123.06 x̃: 44 helped stats (rel) min: 0.04% max: 49.96% x̄: 7.08% x̃: 3.81% HURT stats (abs) min: 2 max: 2098 x̄: 240.33 x̃: 35 HURT stats (rel) min: 0.04% max: 77.07% x̄: 12.34% x̃: 3.00% 95% mean confidence interval for cycles value: -47.93 -12.04 95% mean confidence interval for cycles %-change: -2.87% -1.34% Cycles are helped. No shader-db changes on any other Intel platform. Reviewed-by: Jason Ekstrand <jason.ekstrand@collabora.com> Reviewed-by: Emma Anholt <emma@anholt.net> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/17181>
2022-06-21 16:47:31 -07:00
break;
}
nir/range-analysis: Range tracking for fpow One shader from Metro Last Light and the rest from Rochard. In the Rochard cases, something like: min(1.0, max(pow(saturate(x), y), z)) was transformed to saturate(max(pow(saturate(x), y), z)) because the result of the pow must be >= 0. The Metro Last Light case was similar. An instance of min(pow(abs(x), y), 1.0) became saturate(pow(abs(x), y)) v2: Fix some comments. Suggested by Caio. v3: Fix setting is_intgral when the exponent might be negative. See also Mesa MR !1778. Reviewed-by: Caio Marcelo de Oliveira Filho <caio.oliveira@intel.com> All Intel platforms had similar results. (Ice Lake shown) total instructions in shared programs: 16280670 -> 16280659 (<.01%) instructions in affected programs: 1130 -> 1119 (-0.97%) helped: 11 HURT: 0 helped stats (abs) min: 1 max: 1 x̄: 1.00 x̃: 1 helped stats (rel) min: 0.72% max: 1.43% x̄: 1.03% x̃: 0.97% 95% mean confidence interval for instructions value: -1.00 -1.00 95% mean confidence interval for instructions %-change: -1.19% -0.86% Instructions are helped. total cycles in shared programs: 367168430 -> 367168270 (<.01%) cycles in affected programs: 10281 -> 10121 (-1.56%) helped: 10 HURT: 1 helped stats (abs) min: 16 max: 18 x̄: 17.00 x̃: 17 helped stats (rel) min: 1.31% max: 2.43% x̄: 1.79% x̃: 1.70% HURT stats (abs) min: 10 max: 10 x̄: 10.00 x̃: 10 HURT stats (rel) min: 3.10% max: 3.10% x̄: 3.10% x̃: 3.10% 95% mean confidence interval for cycles value: -20.06 -9.04 95% mean confidence interval for cycles %-change: -2.36% -0.32% Cycles are helped.
2019-08-09 12:48:27 -07:00
case nir_op_fpow: {
/* This is a basic port of the old range analysis, the opcode is very
* underdefined. But improvements are likely possible.
* Due to flush-to-zero semanatics of floating-point numbers with very
nir/range-analysis: Range tracking for fpow One shader from Metro Last Light and the rest from Rochard. In the Rochard cases, something like: min(1.0, max(pow(saturate(x), y), z)) was transformed to saturate(max(pow(saturate(x), y), z)) because the result of the pow must be >= 0. The Metro Last Light case was similar. An instance of min(pow(abs(x), y), 1.0) became saturate(pow(abs(x), y)) v2: Fix some comments. Suggested by Caio. v3: Fix setting is_intgral when the exponent might be negative. See also Mesa MR !1778. Reviewed-by: Caio Marcelo de Oliveira Filho <caio.oliveira@intel.com> All Intel platforms had similar results. (Ice Lake shown) total instructions in shared programs: 16280670 -> 16280659 (<.01%) instructions in affected programs: 1130 -> 1119 (-0.97%) helped: 11 HURT: 0 helped stats (abs) min: 1 max: 1 x̄: 1.00 x̃: 1 helped stats (rel) min: 0.72% max: 1.43% x̄: 1.03% x̃: 0.97% 95% mean confidence interval for instructions value: -1.00 -1.00 95% mean confidence interval for instructions %-change: -1.19% -0.86% Instructions are helped. total cycles in shared programs: 367168430 -> 367168270 (<.01%) cycles in affected programs: 10281 -> 10121 (-1.56%) helped: 10 HURT: 1 helped stats (abs) min: 16 max: 18 x̄: 17.00 x̃: 17 helped stats (rel) min: 1.31% max: 2.43% x̄: 1.79% x̃: 1.70% HURT stats (abs) min: 10 max: 10 x̄: 10.00 x̃: 10 HURT stats (rel) min: 3.10% max: 3.10% x̄: 3.10% x̃: 3.10% 95% mean confidence interval for cycles value: -20.06 -9.04 95% mean confidence interval for cycles %-change: -2.36% -0.32% Cycles are helped.
2019-08-09 12:48:27 -07:00
* small mangnitudes, we can never really be sure a result will be
* non-zero.
*
* NIR uses pow() and powf() to constant evaluate nir_op_fpow. The man
* page for that function says:
*
* If y is 0, the result is 1.0 (even if x is a NaN).
*
* gt_zero: pow(*, eq_zero)
* | pow(eq_zero, lt_zero) # 0^-y = +inf
* | pow(eq_zero, le_zero) # 0^-y = +inf or 0^0 = 1.0
* ;
*
* eq_zero: pow(eq_zero, gt_zero)
* ;
*
* ge_zero: pow(gt_zero, gt_zero)
* | pow(gt_zero, ge_zero)
* | pow(gt_zero, lt_zero)
* | pow(gt_zero, le_zero)
* | pow(gt_zero, ne_zero)
* | pow(gt_zero, unknown)
* | pow(ge_zero, gt_zero)
* | pow(ge_zero, ge_zero)
* | pow(ge_zero, lt_zero)
* | pow(ge_zero, le_zero)
* | pow(ge_zero, ne_zero)
* | pow(ge_zero, unknown)
* | pow(eq_zero, ge_zero) # 0^0 = 1.0 or 0^+y = 0.0
* | pow(eq_zero, ne_zero) # 0^-y = +inf or 0^+y = 0.0
* | pow(eq_zero, unknown) # union of all other y cases
* ;
*
* All other cases are unknown.
*
* We could do better if the right operand is a constant, integral
* value.
*/
fp_class_mask left = src_res[0];
fp_class_mask right = src_res[1];
if (!(right & (FP_CLASS_ANY_NEG | FP_CLASS_ANY_POS))) {
r = FP_CLASS_ANY_POS;
} else if (left & (FP_CLASS_ANY_NEG | FP_CLASS_NEG_ZERO)) {
r = FP_CLASS_UNKNOWN;
} else {
r = FP_CLASS_ANY_POS | FP_CLASS_ANY_ZERO;
if ((right & (FP_CLASS_ANY_NEG | FP_CLASS_NON_INTEGRAL)) || (left & FP_CLASS_NON_INTEGRAL))
r |= FP_CLASS_NON_INTEGRAL;
}
/* Various cases can result in NaN, so assume the worst. */
r |= FP_CLASS_NAN;
nir/range-analysis: Range tracking for fpow One shader from Metro Last Light and the rest from Rochard. In the Rochard cases, something like: min(1.0, max(pow(saturate(x), y), z)) was transformed to saturate(max(pow(saturate(x), y), z)) because the result of the pow must be >= 0. The Metro Last Light case was similar. An instance of min(pow(abs(x), y), 1.0) became saturate(pow(abs(x), y)) v2: Fix some comments. Suggested by Caio. v3: Fix setting is_intgral when the exponent might be negative. See also Mesa MR !1778. Reviewed-by: Caio Marcelo de Oliveira Filho <caio.oliveira@intel.com> All Intel platforms had similar results. (Ice Lake shown) total instructions in shared programs: 16280670 -> 16280659 (<.01%) instructions in affected programs: 1130 -> 1119 (-0.97%) helped: 11 HURT: 0 helped stats (abs) min: 1 max: 1 x̄: 1.00 x̃: 1 helped stats (rel) min: 0.72% max: 1.43% x̄: 1.03% x̃: 0.97% 95% mean confidence interval for instructions value: -1.00 -1.00 95% mean confidence interval for instructions %-change: -1.19% -0.86% Instructions are helped. total cycles in shared programs: 367168430 -> 367168270 (<.01%) cycles in affected programs: 10281 -> 10121 (-1.56%) helped: 10 HURT: 1 helped stats (abs) min: 16 max: 18 x̄: 17.00 x̃: 17 helped stats (rel) min: 1.31% max: 2.43% x̄: 1.79% x̃: 1.70% HURT stats (abs) min: 10 max: 10 x̄: 10.00 x̃: 10 HURT stats (rel) min: 3.10% max: 3.10% x̄: 3.10% x̃: 3.10% 95% mean confidence interval for cycles value: -20.06 -9.04 95% mean confidence interval for cycles %-change: -2.36% -0.32% Cycles are helped.
2019-08-09 12:48:27 -07:00
break;
}
case nir_op_ffma:
case nir_op_ffmaz: {
bool mulz = alu->op == nir_op_ffmaz;
bool src_eq = nir_alu_srcs_equal(alu, alu, 0, 1);
bool src_neg_eq = !nir_src_is_const(alu->src[0].src) && nir_alu_srcs_negative_equal(alu, alu, 0, 1);
fp_class_mask r_mul = fmul_fp_class(src_res[0], src_res[1], mulz, src_eq, src_neg_eq);
r = fadd_fp_class(r_mul, src_res[2]);
/* fma(a, b, +0.0) can be -0.0 if a * b underflows.
* When fused, the underflow is not flushed before the addition.
*/
bool mul_underflow = (((src_res[0] & FP_CLASS_LT_ZERO_GT_NEG_ONE) && (src_res[1] & FP_CLASS_GT_ZERO_LT_POS_ONE)) ||
((src_res[1] & FP_CLASS_LT_ZERO_GT_NEG_ONE) && (src_res[0] & FP_CLASS_GT_ZERO_LT_POS_ONE)));
if (!src_eq && mul_underflow && (src_res[2] & FP_CLASS_POS_ZERO))
r |= FP_CLASS_NEG_ZERO;
break;
}
case nir_op_flrp: {
/* Decompose the flrp to first + third * (second + -first) */
fp_class_mask inner_fadd_class =
fadd_fp_class(src_res[1], fneg_fp_class(src_res[0]));
fp_class_mask fmul_class =
fmul_fp_class(src_res[2], inner_fadd_class, false, false, false);
r = fadd_fp_class(src_res[0], fmul_class);
/* Various cases can result in NaN, so assume the worst. */
r |= FP_CLASS_NAN;
break;
}
default:
r = FP_CLASS_UNKNOWN;
break;
}
if (nir_alu_type_get_base_type(nir_op_infos[alu->op].output_type) == nir_type_float)
r = handle_sz(alu, r);
assert((r & FP_CLASS_UNKNOWN) == r);
assert((r & ~FP_CLASS_NON_INTEGRAL) != 0);
assert(!(r & FP_CLASS_NON_INTEGRAL) || (r & (FP_CLASS_LT_NEG_ONE | FP_CLASS_LT_ZERO_GT_NEG_ONE |
FP_CLASS_GT_POS_ONE | FP_CLASS_GT_ZERO_LT_POS_ONE)));
*result = r;
}
fp_class_mask
nir_analyze_fp_class(nir_fp_analysis_state *fp_state, const nir_def *def)
{
struct fp_query query_alloc[64];
uint32_t result_alloc[64];
struct analysis_state state;
state.range_ht = fp_state;
util_dynarray_init_from_stack(&state.query_stack, query_alloc, sizeof(query_alloc));
util_dynarray_init_from_stack(&state.result_stack, result_alloc, sizeof(result_alloc));
state.query_size = sizeof(struct fp_query);
state.get_key = &get_fp_key;
state.lookup = &fp_lookup;
state.insert = &fp_insert;
state.process_query = &process_fp_query;
push_fp_query(&state, def);
return perform_analysis(&state);
}
nir_fp_analysis_state
nir_create_fp_analysis_state(nir_function_impl *impl)
{
nir_fp_analysis_state state;
state.impl = impl;
/* Over-allocate the bitset, so that we can keep using the allocated table memory
* even when new SSA values are added. */
state.size = BITSET_BYTES(impl->ssa_alloc + impl->ssa_alloc / 4u);
state.max = -1;
state.bitset = calloc(state.size, 1);
util_sparse_array_init(&state.arr, 2, 256);
return state;
}
void
nir_invalidate_fp_analysis_state(nir_fp_analysis_state *state)
{
if (BITSET_BYTES(state->impl->ssa_alloc) > state->size) {
state->size = BITSET_BYTES(state->impl->ssa_alloc + state->impl->ssa_alloc / 4u);
free(state->bitset);
state->bitset = calloc(state->size, 1);
} else if (state->max >= 0) {
memset(state->bitset, 0, BITSET_BYTES(state->max + 1));
}
state->max = -1;
}
void
nir_free_fp_analysis_state(nir_fp_analysis_state *state)
{
util_sparse_array_finish(&state->arr);
free(state->bitset);
}
static uint32_t
bitmask(uint32_t size)
{
return size >= 32 ? 0xffffffffu : ((uint32_t)1 << size) - 1u;
}
static uint64_t
mul_clamp(uint32_t a, uint32_t b)
{
if (a != 0 && (a * b) / a != b)
return (uint64_t)UINT32_MAX + 1;
else
return a * b;
}
/* recursively gather at most "buf_size" phi/bcsel sources */
static unsigned
search_phi_bcsel(nir_scalar scalar, nir_scalar *buf, unsigned buf_size, struct set *visited)
{
if (_mesa_set_search(visited, scalar.def))
return 0;
_mesa_set_add(visited, scalar.def);
if (nir_def_instr_type(scalar.def) == nir_instr_type_phi) {
nir_phi_instr *phi = nir_def_as_phi(scalar.def);
unsigned num_sources_left = exec_list_length(&phi->srcs);
if (buf_size >= num_sources_left) {
unsigned total_added = 0;
nir_foreach_phi_src(src, phi) {
num_sources_left--;
unsigned added = search_phi_bcsel(nir_get_scalar(src->src.ssa, scalar.comp),
buf + total_added, buf_size - num_sources_left, visited);
assert(added <= buf_size);
buf_size -= added;
total_added += added;
}
return total_added;
}
}
if (nir_scalar_is_alu(scalar)) {
nir_op op = nir_scalar_alu_op(scalar);
if ((op == nir_op_bcsel || op == nir_op_b32csel) && buf_size >= 2) {
nir_scalar src1 = nir_scalar_chase_alu_src(scalar, 1);
nir_scalar src2 = nir_scalar_chase_alu_src(scalar, 2);
unsigned added = search_phi_bcsel(src1, buf, buf_size - 1, visited);
buf_size -= added;
added += search_phi_bcsel(src2, buf + added, buf_size, visited);
return added;
}
}
buf[0] = scalar;
return 1;
}
static uint32_t
get_max_workgroup_invocations(nir_shader *nir)
{
if (!nir->options || !nir->options->max_workgroup_invocations)
return UINT16_MAX;
return nir->options->max_workgroup_invocations;
}
static uint32_t
get_max_workgroup_count(nir_shader *nir, unsigned dim)
{
/* max_workgroup_count represents the maximum compute shader / kernel
* dispatchable work size. On most hardware, this is essentially
* unbounded. On some hardware max_workgroup_count[1] and
* max_workgroup_count[2] may be smaller.
*/
if (!nir->options || !nir->options->max_workgroup_count[dim])
return UINT32_MAX;
return nir->options->max_workgroup_count[dim];
}
struct scalar_query {
struct analysis_query head;
nir_scalar scalar;
};
static void
push_scalar_query(struct analysis_state *state, nir_scalar scalar)
{
struct scalar_query *pushed_q = push_analysis_query(state, sizeof(struct scalar_query));
pushed_q->scalar = scalar;
}
static uint32_t
get_scalar_key(struct analysis_query *q)
{
nir_scalar scalar = ((struct scalar_query *)q)->scalar;
/* keys can't be 0, so we have to add 1 to the index */
unsigned shift_amount = ffs(NIR_MAX_VEC_COMPONENTS) - 1;
return nir_scalar_is_const(scalar)
? UINT32_MAX
: ((scalar.def->index + 1) << shift_amount) | scalar.comp;
}
static bool
scalar_lookup(void *table, uint32_t key, uint32_t *value)
{
struct hash_table *ht = table;
struct hash_entry *he = _mesa_hash_table_search(ht, (void *)(uintptr_t)key);
if (he)
*value = (uintptr_t)he->data;
return he != NULL;
}
static void
scalar_insert(void *table, uint32_t key, uint32_t value)
{
struct hash_table *ht = table;
_mesa_hash_table_insert(ht, (void *)(uintptr_t)key, (void *)(uintptr_t)value);
}
static void
get_intrinsic_uub(struct analysis_state *state, struct scalar_query q, uint32_t *result,
const uint32_t *src)
{
nir_shader *shader = state->shader;
nir_intrinsic_instr *intrin = nir_def_as_intrinsic(q.scalar.def);
switch (intrin->intrinsic) {
case nir_intrinsic_load_local_invocation_index:
/* The local invocation index is used under the hood by RADV for
* some non-compute-like shaders (eg. LS and NGG). These technically
* run in workgroups on the HW, even though this fact is not exposed
* by the API.
* They can safely use the same code path here as variable sized
* compute-like shader stages.
*/
if (!mesa_shader_stage_uses_workgroup(shader->info.stage) ||
shader->info.workgroup_size_variable) {
*result = get_max_workgroup_invocations(shader) - 1;
} else {
*result = (shader->info.workgroup_size[0] *
shader->info.workgroup_size[1] *
shader->info.workgroup_size[2]) -
1u;
}
break;
case nir_intrinsic_load_local_invocation_id:
if (shader->info.workgroup_size_variable)
*result = get_max_workgroup_invocations(shader) - 1u;
else
*result = shader->info.workgroup_size[q.scalar.comp] - 1u;
break;
case nir_intrinsic_load_workgroup_id:
*result = get_max_workgroup_count(shader, q.scalar.comp) - 1u;
break;
case nir_intrinsic_load_num_workgroups:
*result = get_max_workgroup_count(shader, q.scalar.comp);
break;
case nir_intrinsic_load_global_invocation_id:
if (shader->info.workgroup_size_variable) {
*result = mul_clamp(get_max_workgroup_invocations(shader),
get_max_workgroup_count(shader, q.scalar.comp)) -
1u;
} else {
*result = (shader->info.workgroup_size[q.scalar.comp] *
get_max_workgroup_count(shader, q.scalar.comp)) -
1u;
}
break;
case nir_intrinsic_load_invocation_id:
if (shader->info.stage == MESA_SHADER_TESS_CTRL)
*result = shader->info.tess.tcs_vertices_out
? (shader->info.tess.tcs_vertices_out - 1)
: 511; /* Generous maximum output patch size of 512 */
break;
case nir_intrinsic_load_subgroup_invocation:
case nir_intrinsic_first_invocation:
*result = shader->info.max_subgroup_size - 1;
break;
case nir_intrinsic_mbcnt_amd: {
if (!q.head.pushed_queries) {
push_scalar_query(state, nir_get_scalar(intrin->src[1].ssa, 0));
return;
} else {
uint32_t src0 = shader->info.max_subgroup_size - 1;
uint32_t src1 = src[0];
if (src0 + src1 >= src0) /* check overflow */
*result = src0 + src1;
}
break;
}
case nir_intrinsic_load_subgroup_size:
if (shader->info.api_subgroup_size)
*result = shader->info.api_subgroup_size;
else
*result = shader->info.max_subgroup_size;
break;
case nir_intrinsic_load_subgroup_id:
case nir_intrinsic_load_num_subgroups: {
uint32_t workgroup_size = get_max_workgroup_invocations(shader);
if (mesa_shader_stage_uses_workgroup(shader->info.stage) &&
!shader->info.workgroup_size_variable) {
workgroup_size = shader->info.workgroup_size[0] *
shader->info.workgroup_size[1] *
shader->info.workgroup_size[2];
}
*result = DIV_ROUND_UP(workgroup_size, shader->info.min_subgroup_size);
if (intrin->intrinsic == nir_intrinsic_load_subgroup_id)
(*result)--;
break;
}
case nir_intrinsic_reduce:
case nir_intrinsic_inclusive_scan:
case nir_intrinsic_exclusive_scan: {
nir_op op = nir_intrinsic_reduction_op(intrin);
switch (op) {
case nir_op_umin:
case nir_op_umax:
case nir_op_imax:
case nir_op_imin:
case nir_op_iand:
case nir_op_ior:
case nir_op_ixor:
case nir_op_iadd:
if (!q.head.pushed_queries) {
push_scalar_query(state, nir_get_scalar(intrin->src[0].ssa, q.scalar.comp));
return;
}
break;
default:
return;
}
unsigned bit_size = q.scalar.def->bit_size;
bool exclusive = intrin->intrinsic == nir_intrinsic_exclusive_scan;
switch (op) {
case nir_op_umin:
case nir_op_umax:
case nir_op_imax:
case nir_op_imin:
case nir_op_iand:
*result = src[0];
break;
case nir_op_ior:
case nir_op_ixor:
*result = bitmask(util_last_bit64(src[0]));
break;
case nir_op_iadd:
*result = MIN2(*result, (uint64_t)src[0] * (shader->info.max_subgroup_size - exclusive));
break;
default:
UNREACHABLE("unhandled op");
}
if (exclusive) {
uint32_t identity = nir_const_value_as_uint(nir_alu_binop_identity(op, bit_size), bit_size);
*result = MAX2(*result, identity);
}
break;
}
case nir_intrinsic_read_first_invocation:
case nir_intrinsic_read_invocation:
case nir_intrinsic_shuffle:
case nir_intrinsic_shuffle_xor:
case nir_intrinsic_shuffle_up:
case nir_intrinsic_shuffle_down:
case nir_intrinsic_quad_broadcast:
case nir_intrinsic_quad_swap_horizontal:
case nir_intrinsic_quad_swap_vertical:
case nir_intrinsic_quad_swap_diagonal:
case nir_intrinsic_quad_swizzle_amd:
case nir_intrinsic_masked_swizzle_amd:
if (!q.head.pushed_queries) {
push_scalar_query(state, nir_get_scalar(intrin->src[0].ssa, q.scalar.comp));
return;
} else {
*result = src[0];
}
break;
case nir_intrinsic_write_invocation_amd:
if (!q.head.pushed_queries) {
push_scalar_query(state, nir_get_scalar(intrin->src[0].ssa, q.scalar.comp));
push_scalar_query(state, nir_get_scalar(intrin->src[1].ssa, q.scalar.comp));
return;
} else {
*result = MAX2(src[0], src[1]);
}
break;
case nir_intrinsic_load_tess_rel_patch_id_amd:
case nir_intrinsic_load_tcs_num_patches_amd:
/* Very generous maximum: TCS/TES executed by largest possible workgroup */
*result = get_max_workgroup_invocations(shader) / MAX2(shader->info.tess.tcs_vertices_out, 1u);
break;
case nir_intrinsic_load_typed_buffer_amd: {
const enum pipe_format format = nir_intrinsic_format(intrin);
if (format == PIPE_FORMAT_NONE)
break;
const struct util_format_description *desc = util_format_description(format);
if (desc->channel[q.scalar.comp].type != UTIL_FORMAT_TYPE_UNSIGNED)
break;
if (desc->channel[q.scalar.comp].normalized) {
*result = fui(1.0);
break;
}
const uint32_t chan_max = u_uintN_max(desc->channel[q.scalar.comp].size);
*result = desc->channel[q.scalar.comp].pure_integer ? chan_max : fui(chan_max);
break;
}
case nir_intrinsic_load_ttmp_register_amd:
case nir_intrinsic_load_scalar_arg_amd:
case nir_intrinsic_load_vector_arg_amd: {
uint32_t upper_bound = nir_intrinsic_arg_upper_bound_u32_amd(intrin);
if (upper_bound)
*result = upper_bound;
break;
}
case nir_intrinsic_image_samples:
if (state->shader->options->max_samples > 0)
*result = state->shader->options->max_samples;
break;
default:
break;
}
}
static void
get_alu_uub(struct analysis_state *state, struct scalar_query q, uint32_t *result, const uint32_t *src)
{
nir_op op = nir_scalar_alu_op(q.scalar);
/* Early exit for unsupported ALU opcodes. */
switch (op) {
case nir_op_umin:
case nir_op_imin:
case nir_op_imax:
case nir_op_umax:
case nir_op_iand:
case nir_op_ior:
case nir_op_ixor:
case nir_op_ishl:
case nir_op_imul:
case nir_op_ushr:
case nir_op_ishr:
case nir_op_iadd:
case nir_op_umod:
case nir_op_udiv:
case nir_op_bcsel:
case nir_op_b32csel:
case nir_op_ubfe:
nir/range_analysis: Handle bfi and bitfield_select in get_alu_uub I noticed some things related to this while implementing support for bitfield_select / BFN in BRW. shader-db: Lunar Lake total instructions in shared programs: 17183140 -> 17183128 (<.01%) instructions in affected programs: 3830 -> 3818 (-0.31%) helped: 6 / HURT: 0 total cycles in shared programs: 889936934 -> 889936056 (<.01%) cycles in affected programs: 253758 -> 252880 (-0.35%) helped: 4 / HURT: 2 No shader-db changes on any other Intel platform. fossil-db: Lunar Lake Totals: Instrs: 233285343 -> 233284796 (-0.00%); split: -0.00%, +0.00% Cycle count: 32756777978 -> 32756399804 (-0.00%); split: -0.00%, +0.00% Max live registers: 71738646 -> 71738626 (-0.00%) Non SSA regs after NIR: 67837900 -> 67837902 (+0.00%) Totals from 177 (0.02% of 790723) affected shaders: Instrs: 389849 -> 389302 (-0.14%); split: -0.14%, +0.00% Cycle count: 356341872 -> 355963698 (-0.11%); split: -0.11%, +0.01% Max live registers: 39364 -> 39344 (-0.05%) Non SSA regs after NIR: 70453 -> 70455 (+0.00%) Meteor Lake, DG2, and Ice Lake had similar results. (Meteor Lake shown) Totals: Instrs: 264095611 -> 264095358 (-0.00%) Cycle count: 26555705299 -> 26554303407 (-0.01%); split: -0.01%, +0.00% Fill count: 613233 -> 613231 (-0.00%) Totals from 123 (0.01% of 905547) affected shaders: Instrs: 334830 -> 334577 (-0.08%) Cycle count: 326531667 -> 325129775 (-0.43%); split: -0.65%, +0.22% Fill count: 4145 -> 4143 (-0.05%) Tiger Lake and Skylake had similar results. (Tiger Lake shown) Totals: Instrs: 269733849 -> 269733590 (-0.00%) Cycle count: 25240548036 -> 25241435039 (+0.00%); split: -0.00%, +0.01% Totals from 123 (0.01% of 903812) affected shaders: Instrs: 338617 -> 338358 (-0.08%) Cycle count: 326605644 -> 327492647 (+0.27%); split: -0.13%, +0.40% Reviewed-by: Matt Turner <mattst88@gmail.com> Reviewed-by: Georg Lehmann <dadschoorse@gmail.com> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/37186>
2025-06-16 11:38:53 -07:00
case nir_op_bfi:
case nir_op_bfm:
nir/range_analysis: Handle bfi and bitfield_select in get_alu_uub I noticed some things related to this while implementing support for bitfield_select / BFN in BRW. shader-db: Lunar Lake total instructions in shared programs: 17183140 -> 17183128 (<.01%) instructions in affected programs: 3830 -> 3818 (-0.31%) helped: 6 / HURT: 0 total cycles in shared programs: 889936934 -> 889936056 (<.01%) cycles in affected programs: 253758 -> 252880 (-0.35%) helped: 4 / HURT: 2 No shader-db changes on any other Intel platform. fossil-db: Lunar Lake Totals: Instrs: 233285343 -> 233284796 (-0.00%); split: -0.00%, +0.00% Cycle count: 32756777978 -> 32756399804 (-0.00%); split: -0.00%, +0.00% Max live registers: 71738646 -> 71738626 (-0.00%) Non SSA regs after NIR: 67837900 -> 67837902 (+0.00%) Totals from 177 (0.02% of 790723) affected shaders: Instrs: 389849 -> 389302 (-0.14%); split: -0.14%, +0.00% Cycle count: 356341872 -> 355963698 (-0.11%); split: -0.11%, +0.01% Max live registers: 39364 -> 39344 (-0.05%) Non SSA regs after NIR: 70453 -> 70455 (+0.00%) Meteor Lake, DG2, and Ice Lake had similar results. (Meteor Lake shown) Totals: Instrs: 264095611 -> 264095358 (-0.00%) Cycle count: 26555705299 -> 26554303407 (-0.01%); split: -0.01%, +0.00% Fill count: 613233 -> 613231 (-0.00%) Totals from 123 (0.01% of 905547) affected shaders: Instrs: 334830 -> 334577 (-0.08%) Cycle count: 326531667 -> 325129775 (-0.43%); split: -0.65%, +0.22% Fill count: 4145 -> 4143 (-0.05%) Tiger Lake and Skylake had similar results. (Tiger Lake shown) Totals: Instrs: 269733849 -> 269733590 (-0.00%) Cycle count: 25240548036 -> 25241435039 (+0.00%); split: -0.00%, +0.01% Totals from 123 (0.01% of 903812) affected shaders: Instrs: 338617 -> 338358 (-0.08%) Cycle count: 326605644 -> 327492647 (+0.27%); split: -0.13%, +0.40% Reviewed-by: Matt Turner <mattst88@gmail.com> Reviewed-by: Georg Lehmann <dadschoorse@gmail.com> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/37186>
2025-06-16 11:38:53 -07:00
case nir_op_bitfield_select:
case nir_op_extract_u8:
case nir_op_extract_i8:
case nir_op_extract_u16:
case nir_op_extract_i16:
case nir_op_b2i8:
case nir_op_b2i16:
case nir_op_b2i32:
break;
case nir_op_u2u1:
case nir_op_u2u8:
case nir_op_u2u16:
case nir_op_u2u32:
if (nir_scalar_chase_alu_src(q.scalar, 0).def->bit_size > 32) {
/* If src is >32 bits, return max */
return;
}
break;
case nir_op_fsat:
case nir_op_fmul:
case nir_op_fmulz:
case nir_op_f2u32:
case nir_op_f2i32:
if (nir_scalar_chase_alu_src(q.scalar, 0).def->bit_size != 32) {
/* Only 32bit floats support for now, return max */
return;
}
break;
case nir_op_bit_count:
if (nir_scalar_chase_alu_src(q.scalar, 0).def->bit_size > 32) {
*result = nir_scalar_chase_alu_src(q.scalar, 0).def->bit_size;
return;
}
break;
default:
return;
}
if (!q.head.pushed_queries) {
for (unsigned i = 0; i < nir_op_infos[op].num_inputs; i++)
push_scalar_query(state, nir_scalar_chase_alu_src(q.scalar, i));
return;
}
uint32_t max = bitmask(q.scalar.def->bit_size);
switch (op) {
case nir_op_umin:
*result = src[0] < src[1] ? src[0] : src[1];
break;
case nir_op_imin:
case nir_op_imax:
case nir_op_umax:
*result = src[0] > src[1] ? src[0] : src[1];
break;
case nir_op_iand: {
nir_scalar src0_scalar = nir_scalar_chase_alu_src(q.scalar, 0);
nir_scalar src1_scalar = nir_scalar_chase_alu_src(q.scalar, 1);
if (nir_scalar_is_const(src0_scalar))
*result = bitmask(util_last_bit64(src[1])) & nir_scalar_as_uint(src0_scalar);
else if (nir_scalar_is_const(src1_scalar))
*result = bitmask(util_last_bit64(src[0])) & nir_scalar_as_uint(src1_scalar);
else
*result = bitmask(util_last_bit64(src[0])) & bitmask(util_last_bit64(src[1]));
break;
}
case nir_op_ior:
case nir_op_ixor: {
nir_scalar src0_scalar = nir_scalar_chase_alu_src(q.scalar, 0);
nir_scalar src1_scalar = nir_scalar_chase_alu_src(q.scalar, 1);
if (nir_scalar_is_const(src0_scalar))
*result = bitmask(util_last_bit64(src[1])) | nir_scalar_as_uint(src0_scalar);
else if (nir_scalar_is_const(src1_scalar))
*result = bitmask(util_last_bit64(src[0])) | nir_scalar_as_uint(src1_scalar);
else
*result = bitmask(util_last_bit64(src[0])) | bitmask(util_last_bit64(src[1]));
break;
}
case nir_op_ishl: {
uint32_t src1 = MIN2(src[1], q.scalar.def->bit_size - 1u);
if (util_last_bit64(src[0]) + src1 <= q.scalar.def->bit_size) /* check overflow */
*result = src[0] << src1;
*result = MIN2(*result, max);
nir_scalar src1_scalar = nir_scalar_chase_alu_src(q.scalar, 1);
if (nir_scalar_is_const(src1_scalar)) {
uint32_t const_val = 1u << (nir_scalar_as_uint(src1_scalar) & (q.scalar.def->bit_size - 1u));
*result = MIN2(*result, max / const_val * const_val);
}
break;
}
case nir_op_imul: {
if (src[0] == 0 || (src[0] * src[1]) / src[0] == src[1]) /* check overflow */
*result = src[0] * src[1];
*result = MIN2(*result, max);
nir_scalar src0_scalar = nir_scalar_chase_alu_src(q.scalar, 0);
nir_scalar src1_scalar = nir_scalar_chase_alu_src(q.scalar, 1);
if (nir_scalar_is_const(src0_scalar)) {
uint32_t const_val = nir_scalar_as_uint(src0_scalar);
*result = const_val ? MIN2(*result, max / const_val * const_val) : 0;
} else if (nir_scalar_is_const(src1_scalar)) {
uint32_t const_val = nir_scalar_as_uint(src1_scalar);
*result = const_val ? MIN2(*result, max / const_val * const_val) : 0;
}
break;
}
case nir_op_ushr: {
nir_scalar src1_scalar = nir_scalar_chase_alu_src(q.scalar, 1);
uint32_t mask = q.scalar.def->bit_size - 1u;
if (nir_scalar_is_const(src1_scalar))
*result = src[0] >> (nir_scalar_as_uint(src1_scalar) & mask);
else
*result = src[0];
break;
}
case nir_op_ishr: {
nir_scalar src1_scalar = nir_scalar_chase_alu_src(q.scalar, 1);
uint32_t mask = q.scalar.def->bit_size - 1u;
if (src[0] <= 2147483647 && nir_scalar_is_const(src1_scalar))
*result = src[0] >> (nir_scalar_as_uint(src1_scalar) & mask);
else
*result = src[0];
break;
}
case nir_op_iadd:
if (src[0] + src[1] >= src[0]) /* check overflow */
*result = src[0] + src[1];
*result = MIN2(*result, max);
break;
case nir_op_umod:
*result = src[1] ? src[1] - 1 : 0;
break;
case nir_op_udiv: {
nir_scalar src1_scalar = nir_scalar_chase_alu_src(q.scalar, 1);
if (nir_scalar_is_const(src1_scalar))
*result = nir_scalar_as_uint(src1_scalar)
? src[0] / nir_scalar_as_uint(src1_scalar)
: 0;
else
*result = src[0];
break;
}
case nir_op_bcsel:
case nir_op_b32csel:
*result = src[1] > src[2] ? src[1] : src[2];
break;
case nir_op_ubfe:
*result = bitmask(MIN2(src[2], q.scalar.def->bit_size));
break;
case nir_op_bfm: {
nir_scalar src1_scalar = nir_scalar_chase_alu_src(q.scalar, 1);
if (nir_scalar_is_const(src1_scalar)) {
uint32_t src0 = MIN2(src[0], 31);
uint32_t src1 = nir_scalar_as_uint(src1_scalar) & 0x1fu;
*result = bitmask(src0) << src1;
} else {
uint32_t src0 = MIN2(src[0], 31);
uint32_t src1 = MIN2(src[1], 31);
*result = bitmask(MIN2(src0 + src1, 32));
}
break;
}
nir/range_analysis: Handle bfi and bitfield_select in get_alu_uub I noticed some things related to this while implementing support for bitfield_select / BFN in BRW. shader-db: Lunar Lake total instructions in shared programs: 17183140 -> 17183128 (<.01%) instructions in affected programs: 3830 -> 3818 (-0.31%) helped: 6 / HURT: 0 total cycles in shared programs: 889936934 -> 889936056 (<.01%) cycles in affected programs: 253758 -> 252880 (-0.35%) helped: 4 / HURT: 2 No shader-db changes on any other Intel platform. fossil-db: Lunar Lake Totals: Instrs: 233285343 -> 233284796 (-0.00%); split: -0.00%, +0.00% Cycle count: 32756777978 -> 32756399804 (-0.00%); split: -0.00%, +0.00% Max live registers: 71738646 -> 71738626 (-0.00%) Non SSA regs after NIR: 67837900 -> 67837902 (+0.00%) Totals from 177 (0.02% of 790723) affected shaders: Instrs: 389849 -> 389302 (-0.14%); split: -0.14%, +0.00% Cycle count: 356341872 -> 355963698 (-0.11%); split: -0.11%, +0.01% Max live registers: 39364 -> 39344 (-0.05%) Non SSA regs after NIR: 70453 -> 70455 (+0.00%) Meteor Lake, DG2, and Ice Lake had similar results. (Meteor Lake shown) Totals: Instrs: 264095611 -> 264095358 (-0.00%) Cycle count: 26555705299 -> 26554303407 (-0.01%); split: -0.01%, +0.00% Fill count: 613233 -> 613231 (-0.00%) Totals from 123 (0.01% of 905547) affected shaders: Instrs: 334830 -> 334577 (-0.08%) Cycle count: 326531667 -> 325129775 (-0.43%); split: -0.65%, +0.22% Fill count: 4145 -> 4143 (-0.05%) Tiger Lake and Skylake had similar results. (Tiger Lake shown) Totals: Instrs: 269733849 -> 269733590 (-0.00%) Cycle count: 25240548036 -> 25241435039 (+0.00%); split: -0.00%, +0.01% Totals from 123 (0.01% of 903812) affected shaders: Instrs: 338617 -> 338358 (-0.08%) Cycle count: 326605644 -> 327492647 (+0.27%); split: -0.13%, +0.40% Reviewed-by: Matt Turner <mattst88@gmail.com> Reviewed-by: Georg Lehmann <dadschoorse@gmail.com> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/37186>
2025-06-16 11:38:53 -07:00
case nir_op_bfi: {
nir_scalar src0_scalar = nir_scalar_chase_alu_src(q.scalar, 0);
const uint64_t s1 = bitmask(util_last_bit64(src[1]));
const uint64_t s2 = bitmask(util_last_bit64(src[2]));
if (nir_scalar_is_const(src0_scalar)) {
const uint64_t s0 = nir_scalar_as_uint(src0_scalar);
/* This case should be eliminated by opt_algebraic. */
if (s0 == 0) {
*result = s2;
} else {
const int x = ffsll(s0) - 1;
*result = (s0 & (s1 << x)) | (~s0 & s2);
}
} else {
const uint64_t s0 = bitmask(util_last_bit64(src[0]));
/* Due to the unpredictable shift, the true maximum value of (s0 &
* (s1 << x)) cannot be known. However, it cannot be larger than
* s0.
*
* inot doesn't work in get_alu_uub. It is known that (~s0 & s2)
* cannot be larger than s2, so just use s2 as a loose upper bound.
*/
*result = s0 | s2;
}
break;
}
case nir_op_bitfield_select: {
nir_scalar src0_scalar = nir_scalar_chase_alu_src(q.scalar, 0);
const uint64_t s1 = bitmask(util_last_bit64(src[1]));
const uint64_t s2 = bitmask(util_last_bit64(src[2]));
if (nir_scalar_is_const(src0_scalar)) {
const uint64_t s0 = nir_scalar_as_uint(src0_scalar);
*result = (s0 & s1) | (~s0 & s2);
} else {
const uint64_t s0 = bitmask(util_last_bit64(src[0]));
/* inot doesn't work in get_alu_uub. It is known that (~s0 & s2)
* cannot be larger than s2, so just use s2 as a loose upper bound.
*/
*result = (s0 & s1) | s2;
}
break;
}
/* limited floating-point support for f2u32(fmul(load_input(), <constant>)) */
case nir_op_f2i32:
case nir_op_f2u32:
/* infinity/NaN starts at 0x7f800000u, negative numbers at 0x80000000 */
if (src[0] < 0x7f800000u) {
float val;
memcpy(&val, &src[0], 4);
*result = (uint32_t)val;
}
break;
case nir_op_fmul:
case nir_op_fmulz:
/* infinity/NaN starts at 0x7f800000u, negative numbers at 0x80000000 */
if (src[0] < 0x7f800000u && src[1] < 0x7f800000u) {
float src0_f, src1_f;
memcpy(&src0_f, &src[0], 4);
memcpy(&src1_f, &src[1], 4);
/* not a proper rounding-up multiplication, but should be good enough */
float max_f = ceilf(src0_f) * ceilf(src1_f);
memcpy(result, &max_f, 4);
}
break;
case nir_op_fsat:
*result = 0x3f800000u;
break;
case nir_op_u2u1:
case nir_op_u2u8:
case nir_op_u2u16:
case nir_op_u2u32:
*result = MIN2(src[0], max);
break;
case nir_op_b2i8:
case nir_op_b2i16:
case nir_op_b2i32:
*result = 1;
break;
case nir_op_msad_4x8:
*result = MIN2((uint64_t)src[2] + 4 * 255, UINT32_MAX);
break;
case nir_op_extract_u8:
*result = MIN2(src[0], UINT8_MAX);
break;
case nir_op_extract_i8:
*result = (src[0] >= 0x80) ? max : MIN2(src[0], INT8_MAX);
break;
case nir_op_extract_u16:
*result = MIN2(src[0], UINT16_MAX);
break;
case nir_op_extract_i16:
*result = (src[0] >= 0x8000) ? max : MIN2(src[0], INT16_MAX);
break;
case nir_op_bit_count:
*result = util_last_bit64(src[0]);
break;
default:
break;
}
}
static void
get_tex_uub(struct analysis_state *state, struct scalar_query q, uint32_t *result, const uint32_t *src)
{
nir_tex_instr *tex = nir_scalar_as_tex(q.scalar);
if (tex->op == nir_texop_texture_samples && state->shader->options->max_samples > 0)
*result = state->shader->options->max_samples;
}
static void
get_phi_uub(struct analysis_state *state, struct scalar_query q, uint32_t *result, const uint32_t *src)
{
nir_phi_instr *phi = nir_def_as_phi(q.scalar.def);
if (exec_list_is_empty(&phi->srcs))
return;
if (q.head.pushed_queries) {
*result = src[0];
for (unsigned i = 1; i < q.head.pushed_queries; i++)
*result = MAX2(*result, src[i]);
return;
}
nir_cf_node *prev = nir_cf_node_prev(&phi->instr.block->cf_node);
if (!prev || prev->type == nir_cf_node_block) {
/* Resolve cycles by inserting max into range_ht. */
uint32_t max = bitmask(q.scalar.def->bit_size);
scalar_insert(state->range_ht, get_scalar_key(&q.head), max);
struct set *visited = _mesa_pointer_set_create(NULL);
nir_scalar *defs = alloca(sizeof(nir_scalar) * 64);
unsigned def_count = search_phi_bcsel(q.scalar, defs, 64, visited);
_mesa_set_destroy(visited, NULL);
for (unsigned i = 0; i < def_count; i++)
push_scalar_query(state, defs[i]);
} else {
nir_foreach_phi_src(src, phi)
push_scalar_query(state, nir_get_scalar(src->src.ssa, q.scalar.comp));
}
}
static void
process_uub_query(struct analysis_state *state, struct analysis_query *aq, uint32_t *result,
const uint32_t *src)
{
struct scalar_query q = *(struct scalar_query *)aq;
*result = bitmask(q.scalar.def->bit_size);
if (nir_scalar_is_const(q.scalar))
*result = nir_scalar_as_uint(q.scalar);
else if (nir_scalar_is_intrinsic(q.scalar))
get_intrinsic_uub(state, q, result, src);
else if (nir_scalar_is_alu(q.scalar))
get_alu_uub(state, q, result, src);
else if (nir_def_instr_type(q.scalar.def) == nir_instr_type_tex)
get_tex_uub(state, q, result, src);
else if (nir_def_instr_type(q.scalar.def) == nir_instr_type_phi)
get_phi_uub(state, q, result, src);
}
uint32_t
nir_unsigned_upper_bound(nir_shader *shader, struct hash_table *range_ht,
nir_scalar scalar)
{
struct scalar_query query_alloc[16];
uint32_t result_alloc[16];
struct analysis_state state;
state.shader = shader;
state.range_ht = range_ht;
util_dynarray_init_from_stack(&state.query_stack, query_alloc, sizeof(query_alloc));
util_dynarray_init_from_stack(&state.result_stack, result_alloc, sizeof(result_alloc));
state.query_size = sizeof(struct scalar_query);
state.get_key = &get_scalar_key;
state.lookup = &scalar_lookup,
state.insert = &scalar_insert,
state.process_query = &process_uub_query;
push_scalar_query(&state, scalar);
_mesa_hash_table_set_deleted_key(range_ht, (void *)(uintptr_t)UINT32_MAX);
return perform_analysis(&state);
}
bool
nir_addition_might_overflow(nir_shader *shader, struct hash_table *range_ht,
nir_scalar ssa, unsigned const_val)
{
uint32_t ub = nir_unsigned_upper_bound(shader, range_ht, ssa);
return const_val + ub < const_val;
}
static uint64_t
ssa_def_bits_used(const nir_def *def, int recur)
{
uint64_t bits_used = 0;
uint64_t all_bits = BITFIELD64_MASK(def->bit_size);
/* Querying the bits used from a vector is too hard of a question to
* answer. Return the conservative answer that all bits are used. To
* handle this, the function would need to be extended to be a query of a
* single component of the vector. That would also necessary to fully
* handle the 'num_components > 1' inside the loop below.
*
* FINISHME: This restriction will eventually need to be restricted to be
* useful for hardware that uses u16vec2 as the native 16-bit integer type.
*/
if (def->num_components > 1)
return all_bits;
/* Limit recursion */
if (recur-- <= 0)
return all_bits;
nir_foreach_use(src, def) {
switch (nir_src_parent_instr(src)->type) {
case nir_instr_type_alu: {
nir_alu_instr *use_alu = nir_instr_as_alu(nir_src_parent_instr(src));
unsigned src_idx = container_of(src, nir_alu_src, src) - use_alu->src;
/* If a user of the value produces a vector result, return the
* conservative answer that all bits are used. It is possible to
* answer this query by looping over the components used. For example,
*
* vec4 32 ssa_5 = load_const(0x0000f000, 0x00000f00, 0x000000f0, 0x0000000f)
* ...
* vec4 32 ssa_8 = iand ssa_7.xxxx, ssa_5
*
* could conceivably return 0x0000ffff when queyring the bits used of
* ssa_7. This is unlikely to be worth the effort because the
* question can eventually answered after the shader has been
* scalarized.
*/
if (use_alu->def.num_components > 1)
return all_bits;
switch (use_alu->op) {
case nir_op_u2u8:
case nir_op_i2i8:
case nir_op_u2u16:
case nir_op_i2i16:
case nir_op_u2u32:
case nir_op_i2i32:
case nir_op_u2u64:
case nir_op_i2i64: {
uint64_t def_bits_used = ssa_def_bits_used(&use_alu->def, recur);
/* If one of the sign-extended bits is used, set the last src bit
* as used.
*/
if ((use_alu->op == nir_op_i2i8 || use_alu->op == nir_op_i2i16 ||
use_alu->op == nir_op_i2i32 || use_alu->op == nir_op_i2i64) &&
def_bits_used & ~all_bits)
def_bits_used |= BITFIELD64_BIT(def->bit_size - 1);
bits_used |= def_bits_used & all_bits;
break;
}
case nir_op_extract_u8:
case nir_op_extract_i8:
case nir_op_extract_u16:
case nir_op_extract_i16:
if (src_idx == 0 && nir_src_is_const(use_alu->src[1].src)) {
unsigned chunk = nir_alu_src_as_uint(use_alu->src[1]);
uint64_t defs_bits_used = ssa_def_bits_used(&use_alu->def, recur);
unsigned field_bits = use_alu->op == nir_op_extract_u8 ||
use_alu->op == nir_op_extract_i8 ? 8 : 16;
uint64_t field_bitmask = BITFIELD64_MASK(field_bits);
/* If one of the sign-extended bits is used, set the last src bit
* as used.
*/
if ((use_alu->op == nir_op_extract_i8 ||
use_alu->op == nir_op_extract_i16) &&
defs_bits_used & ~field_bitmask)
defs_bits_used |= BITFIELD64_BIT(field_bits - 1);
bits_used |= (field_bitmask & defs_bits_used) <<
(chunk * field_bits);
break;
} else {
return all_bits;
}
case nir_op_ishl:
case nir_op_ishr:
case nir_op_ushr:
if (src_idx == 0 && nir_src_is_const(use_alu->src[1].src)) {
unsigned bit_size = def->bit_size;
unsigned shift = nir_alu_src_as_uint(use_alu->src[1]) & (bit_size - 1);
uint64_t def_bits_used = ssa_def_bits_used(&use_alu->def, recur);
/* If one of the sign-extended bits is used, set the "last src
* bit before shifting" as used.
*/
if (use_alu->op == nir_op_ishr &&
def_bits_used & ~(all_bits >> shift))
def_bits_used |= BITFIELD64_BIT(bit_size - 1 - shift);
/* Reverse the shift to get the bits before shifting. */
if (use_alu->op == nir_op_ushr || use_alu->op == nir_op_ishr)
bits_used |= (def_bits_used << shift) & all_bits;
else
bits_used |= def_bits_used >> shift;
break;
} else if (src_idx == 1) {
bits_used |= use_alu->def.bit_size - 1;
break;
} else {
return all_bits;
}
case nir_op_iand:
case nir_op_ior:
assert(src_idx < 2);
if (nir_src_is_const(use_alu->src[1 - src_idx].src)) {
uint64_t other_src = nir_alu_src_as_uint(use_alu->src[1 - src_idx]);
if (use_alu->op == nir_op_iand)
bits_used |= ssa_def_bits_used(&use_alu->def, recur) & other_src;
else
bits_used |= ssa_def_bits_used(&use_alu->def, recur) & ~other_src;
break;
} else {
return all_bits;
}
case nir_op_ibfe:
case nir_op_ubfe:
if (src_idx == 0 && nir_src_is_const(use_alu->src[1].src)) {
uint64_t def_bits_used = ssa_def_bits_used(&use_alu->def, recur);
unsigned bit_size = use_alu->def.bit_size;
unsigned offset = nir_alu_src_as_uint(use_alu->src[1]) & (bit_size - 1);
unsigned bits = nir_src_is_const(use_alu->src[2].src) ?
nir_alu_src_as_uint(use_alu->src[2]) & (bit_size - 1) :
/* Worst case if bits is not constant. */
(bit_size - offset);
uint64_t field_bitmask = BITFIELD64_MASK(bits);
/* If one of the sign-extended bits is used, set the last src
* bit as used.
* If bits is not constant, all bits can be the last one.
*/
if (use_alu->op == nir_op_ibfe &&
(def_bits_used >> offset) & ~field_bitmask) {
if (nir_alu_src_as_uint(use_alu->src[2]))
def_bits_used |= BITFIELD64_BIT(bits - 1);
else
def_bits_used |= field_bitmask;
}
bits_used |= (field_bitmask & def_bits_used) << offset;
break;
} else if (src_idx == 1 || src_idx == 2) {
bits_used |= use_alu->src[0].src.ssa->bit_size - 1;
break;
} else {
return all_bits;
}
case nir_op_imul24:
case nir_op_umul24:
bits_used |= all_bits & 0xffffff;
break;
case nir_op_mov:
bits_used |= ssa_def_bits_used(&use_alu->def, recur);
break;
case nir_op_bcsel:
if (src_idx == 0)
bits_used |= 0x1;
else
bits_used |= ssa_def_bits_used(&use_alu->def, recur);
break;
default:
/* We don't know what this op does */
return all_bits;
}
break;
}
case nir_instr_type_intrinsic: {
nir_intrinsic_instr *use_intrin =
nir_instr_as_intrinsic(nir_src_parent_instr(src));
unsigned src_idx = src - use_intrin->src;
switch (use_intrin->intrinsic) {
case nir_intrinsic_read_invocation:
case nir_intrinsic_shuffle:
case nir_intrinsic_shuffle_up:
case nir_intrinsic_shuffle_down:
case nir_intrinsic_shuffle_xor:
case nir_intrinsic_quad_broadcast:
case nir_intrinsic_quad_swap_horizontal:
case nir_intrinsic_quad_swap_vertical:
case nir_intrinsic_quad_swap_diagonal:
if (src_idx == 0) {
bits_used |= ssa_def_bits_used(&use_intrin->def, recur);
} else {
if (use_intrin->intrinsic == nir_intrinsic_quad_broadcast) {
bits_used |= 3;
} else {
/* Subgroups larger than 128 are not a thing */
bits_used |= 127;
}
}
break;
case nir_intrinsic_reduce:
case nir_intrinsic_inclusive_scan:
case nir_intrinsic_exclusive_scan:
assert(src_idx == 0);
switch (nir_intrinsic_reduction_op(use_intrin)) {
case nir_op_iadd:
case nir_op_imul:
case nir_op_ior:
case nir_op_iand:
case nir_op_ixor:
bits_used |= ssa_def_bits_used(&use_intrin->def, recur);
break;
default:
return all_bits;
}
break;
default:
/* We don't know what this op does */
return all_bits;
}
break;
}
case nir_instr_type_phi: {
nir_phi_instr *use_phi = nir_instr_as_phi(nir_src_parent_instr(src));
bits_used |= ssa_def_bits_used(&use_phi->def, recur);
break;
}
default:
return all_bits;
}
/* If we've somehow shown that all our bits are used, we're done */
assert((bits_used & ~all_bits) == 0);
if (bits_used == all_bits)
return all_bits;
}
return bits_used;
}
uint64_t
nir_def_bits_used(const nir_def *def)
{
return ssa_def_bits_used(def, 2);
}
static void
get_alu_num_lsb(struct analysis_state *state, struct scalar_query q, uint32_t *result, const uint32_t *src)
{
nir_op op = nir_scalar_alu_op(q.scalar);
switch (op) {
case nir_op_ior:
case nir_op_ixor:
case nir_op_iadd:
case nir_op_iand:
case nir_op_imul:
if (!q.head.pushed_queries) {
push_scalar_query(state, nir_scalar_chase_alu_src(q.scalar, 0));
push_scalar_query(state, nir_scalar_chase_alu_src(q.scalar, 1));
return;
}
break;
case nir_op_ishl:
if (!q.head.pushed_queries) {
push_scalar_query(state, nir_scalar_chase_alu_src(q.scalar, 0));
return;
}
break;
case nir_op_ishr:
case nir_op_ushr:
if (!q.head.pushed_queries) {
if (nir_scalar_is_const(nir_scalar_chase_alu_src(q.scalar, 1)))
push_scalar_query(state, nir_scalar_chase_alu_src(q.scalar, 0));
return;
}
break;
case nir_op_bcsel:
if (!q.head.pushed_queries) {
push_scalar_query(state, nir_scalar_chase_alu_src(q.scalar, 1));
push_scalar_query(state, nir_scalar_chase_alu_src(q.scalar, 2));
return;
}
break;
default:
return;
}
switch (op) {
case nir_op_ior:
case nir_op_ixor:
case nir_op_iadd: {
*result = MIN2(src[0], src[1]);
break;
}
case nir_op_iand: {
*result = MAX2(src[0], src[1]);
break;
}
case nir_op_imul: {
*result = MIN2(src[0] + src[1], q.scalar.def->bit_size);
break;
}
case nir_op_ishl: {
nir_scalar src1 = nir_scalar_chase_alu_src(q.scalar, 1);
uint32_t mask = q.scalar.def->bit_size - 1;
unsigned amount = nir_scalar_is_const(src1) ? nir_scalar_as_uint(src1) & mask : 0;
*result = MIN2(src[0] + amount, q.scalar.def->bit_size);
break;
}
case nir_op_ishr:
case nir_op_ushr: {
nir_scalar src1 = nir_scalar_chase_alu_src(q.scalar, 1);
unsigned amount = nir_scalar_as_uint(src1) & (q.scalar.def->bit_size - 1);
*result = amount > src[0] ? 0 : src[0] - amount;
break;
}
case nir_op_bcsel: {
*result = MIN2(src[0], src[1]);
break;
}
default:
UNREACHABLE("Unknown opcode");
}
}
static void
process_num_lsb_query(struct analysis_state *state, struct analysis_query *aq, uint32_t *result,
const uint32_t *src)
{
struct scalar_query q = *(struct scalar_query *)aq;
*result = 0;
if (nir_scalar_is_const(q.scalar)) {
uint64_t val = nir_scalar_as_uint(q.scalar);
*result = val ? ffsll(val) - 1 : q.scalar.def->bit_size;
} else if (nir_scalar_is_alu(q.scalar)) {
get_alu_num_lsb(state, q, result, src);
}
}
unsigned
nir_def_num_lsb_zero(struct hash_table *numlsb_ht, nir_scalar def)
{
struct scalar_query query_alloc[16];
uint32_t result_alloc[16];
struct analysis_state state;
state.shader = NULL;
state.range_ht = numlsb_ht;
util_dynarray_init_from_stack(&state.query_stack, query_alloc, sizeof(query_alloc));
util_dynarray_init_from_stack(&state.result_stack, result_alloc, sizeof(result_alloc));
state.query_size = sizeof(struct scalar_query);
state.get_key = &get_scalar_key;
state.lookup = &scalar_lookup,
state.insert = &scalar_insert,
state.process_query = &process_num_lsb_query;
push_scalar_query(&state, def);
_mesa_hash_table_set_deleted_key(numlsb_ht, (void *)(uintptr_t)UINT32_MAX);
return perform_analysis(&state);
}