OpenHarmony/third_party_mesa3d
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

ac/llvm: don't set "readnone" on non-memory intrinsics

Browse Source 
It's illegal and LLVM always knows which intrinsics don't read memory.
This started failing IR validation with LLVM 16.

Reviewed-by: Pierre-Eric Pelloux-Prayer <pierre-eric.pelloux-prayer@amd.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/20146>
This commit is contained in:
Marek Olšák
2022-12-04 06:00:15 -05:00
committed by Marge Bot
parent 50bece9322
commit d3a03d24fc
3 changed files with 99 additions and 142 deletions
Download Patch File Download Diff File Expand all files Collapse all files

173
src/amd/llvm/ac_llvm_build.c
View File

@@ -510,9 +510,7 @@ LLVMValueRef ac_build_ballot(struct ac_llvm_context *ctx, LLVMValueRef value)
args[0] = ac_to_integer(ctx, args[0]);
return ac_build_intrinsic(
ctx, name, ctx->iN_wavemask, args, 3,
AC_FUNC_ATTR_READNONE);
return ac_build_intrinsic(ctx, name, ctx->iN_wavemask, args, 3, 0);
}
LLVMValueRef ac_get_i1_sgpr_mask(struct ac_llvm_context *ctx, LLVMValueRef value)
@@ -530,9 +528,7 @@ LLVMValueRef ac_get_i1_sgpr_mask(struct ac_llvm_context *ctx, LLVMValueRef value
LLVMConstInt(ctx->i32, LLVMIntNE, 0),
};
return ac_build_intrinsic(
ctx, name, ctx->iN_wavemask, args, 3,
AC_FUNC_ATTR_READNONE);
return ac_build_intrinsic(ctx, name, ctx->iN_wavemask, args, 3, 0);
}
LLVMValueRef ac_build_vote_all(struct ac_llvm_context *ctx, LLVMValueRef value)
@@ -694,7 +690,7 @@ LLVMValueRef ac_build_round(struct ac_llvm_context *ctx, LLVMValueRef value)
else
name = "llvm.rint.f64";
return ac_build_intrinsic(ctx, name, LLVMTypeOf(value), &value, 1, AC_FUNC_ATTR_READNONE);
return ac_build_intrinsic(ctx, name, LLVMTypeOf(value), &value, 1, 0);
}
LLVMValueRef ac_build_fdiv(struct ac_llvm_context *ctx, LLVMValueRef num, LLVMValueRef den)
@@ -714,7 +710,7 @@ LLVMValueRef ac_build_fdiv(struct ac_llvm_context *ctx, LLVMValueRef num, LLVMVa
name = "llvm.amdgcn.rcp.f64";
LLVMValueRef rcp =
ac_build_intrinsic(ctx, name, LLVMTypeOf(den), &den, 1, AC_FUNC_ATTR_READNONE);
ac_build_intrinsic(ctx, name, LLVMTypeOf(den), &den, 1, 0);
return LLVMBuildFMul(ctx->builder, num, rcp, "");
}
@@ -783,10 +779,10 @@ static void build_cube_intrinsic(struct ac_llvm_context *ctx, LLVMValueRef in[3]
{
LLVMTypeRef f32 = ctx->f32;
out->stc[1] = ac_build_intrinsic(ctx, "llvm.amdgcn.cubetc", f32, in, 3, AC_FUNC_ATTR_READNONE);
out->stc[0] = ac_build_intrinsic(ctx, "llvm.amdgcn.cubesc", f32, in, 3, AC_FUNC_ATTR_READNONE);
out->ma = ac_build_intrinsic(ctx, "llvm.amdgcn.cubema", f32, in, 3, AC_FUNC_ATTR_READNONE);
out->id = ac_build_intrinsic(ctx, "llvm.amdgcn.cubeid", f32, in, 3, AC_FUNC_ATTR_READNONE);
out->stc[1] = ac_build_intrinsic(ctx, "llvm.amdgcn.cubetc", f32, in, 3, 0);
out->stc[0] = ac_build_intrinsic(ctx, "llvm.amdgcn.cubesc", f32, in, 3, 0);
out->ma = ac_build_intrinsic(ctx, "llvm.amdgcn.cubema", f32, in, 3, 0);
out->id = ac_build_intrinsic(ctx, "llvm.amdgcn.cubeid", f32, in, 3, 0);
}
/**
@@ -839,7 +835,7 @@ static void build_cube_select(struct ac_llvm_context *ctx,
/* Select ma */
tmp = LLVMBuildSelect(builder, is_ma_z, coords[2],
LLVMBuildSelect(builder, is_ma_y, coords[1], coords[0], ""), "");
tmp = ac_build_intrinsic(ctx, "llvm.fabs.f32", ctx->f32, &tmp, 1, AC_FUNC_ATTR_READNONE);
tmp = ac_build_intrinsic(ctx, "llvm.fabs.f32", ctx->f32, &tmp, 1, 0);
*out_ma = LLVMBuildFMul(builder, tmp, LLVMConstReal(f32, 2.0), "");
}
@@ -886,7 +882,7 @@ void ac_prepare_cube_coords(struct ac_llvm_context *ctx, bool is_deriv, bool is_
build_cube_intrinsic(ctx, coords_arg, &selcoords);
invma =
ac_build_intrinsic(ctx, "llvm.fabs.f32", ctx->f32, &selcoords.ma, 1, AC_FUNC_ATTR_READNONE);
ac_build_intrinsic(ctx, "llvm.fabs.f32", ctx->f32, &selcoords.ma, 1, 0);
invma = ac_build_fdiv(ctx, LLVMConstReal(ctx->f32, 1.0), invma);
for (int i = 0; i < 2; ++i)
@@ -967,21 +963,21 @@ LLVMValueRef ac_build_fs_interp(struct ac_llvm_context *ctx, LLVMValueRef llvm_c
args[2] = params;
p = ac_build_intrinsic(ctx, "llvm.amdgcn.lds.param.load",
ctx->f32, args, 3, AC_FUNC_ATTR_READNONE);
ctx->f32, args, 3, 0);
args[0] = p;
args[1] = i;
args[2] = p;
p10 = ac_build_intrinsic(ctx, "llvm.amdgcn.interp.inreg.p10",
ctx->f32, args, 3, AC_FUNC_ATTR_READNONE);
ctx->f32, args, 3, 0);
args[0] = p;
args[1] = j;
args[2] = p10;
return ac_build_intrinsic(ctx, "llvm.amdgcn.interp.inreg.p2",
ctx->f32, args, 3, AC_FUNC_ATTR_READNONE);
ctx->f32, args, 3, 0);
} else {
LLVMValueRef p1;
@@ -992,7 +988,7 @@ LLVMValueRef ac_build_fs_interp(struct ac_llvm_context *ctx, LLVMValueRef llvm_c
args[3] = params;
p1 = ac_build_intrinsic(ctx, "llvm.amdgcn.interp.p1",
ctx->f32, args, 4, AC_FUNC_ATTR_READNONE);
ctx->f32, args, 4, 0);
args[0] = p1;
args[1] = j;
@@ -1001,7 +997,7 @@ LLVMValueRef ac_build_fs_interp(struct ac_llvm_context *ctx, LLVMValueRef llvm_c
args[4] = params;
return ac_build_intrinsic(ctx, "llvm.amdgcn.interp.p2",
ctx->f32, args, 5, AC_FUNC_ATTR_READNONE);
ctx->f32, args, 5, 0);
}
}
@@ -1020,7 +1016,7 @@ LLVMValueRef ac_build_fs_interp_f16(struct ac_llvm_context *ctx, LLVMValueRef ll
args[2] = params;
p = ac_build_intrinsic(ctx, "llvm.amdgcn.lds.param.load",
ctx->f32, args, 3, AC_FUNC_ATTR_READNONE);
ctx->f32, args, 3, 0);
args[0] = p;
args[1] = i;
@@ -1028,7 +1024,7 @@ LLVMValueRef ac_build_fs_interp_f16(struct ac_llvm_context *ctx, LLVMValueRef ll
args[3] = high_16bits ? ctx->i1true : ctx->i1false;
p10 = ac_build_intrinsic(ctx, "llvm.amdgcn.interp.inreg.p10.f16",
ctx->f32, args, 4, AC_FUNC_ATTR_READNONE);
ctx->f32, args, 4, 0);
args[0] = p;
args[1] = j;
@@ -1036,7 +1032,7 @@ LLVMValueRef ac_build_fs_interp_f16(struct ac_llvm_context *ctx, LLVMValueRef ll
args[3] = high_16bits ? ctx->i1true : ctx->i1false;
return ac_build_intrinsic(ctx, "llvm.amdgcn.interp.inreg.p2.f16",
ctx->f16, args, 4, AC_FUNC_ATTR_READNONE);
ctx->f16, args, 4, 0);
} else {
LLVMValueRef p1;
@@ -1048,7 +1044,7 @@ LLVMValueRef ac_build_fs_interp_f16(struct ac_llvm_context *ctx, LLVMValueRef ll
args[4] = params;
p1 = ac_build_intrinsic(ctx, "llvm.amdgcn.interp.p1.f16", ctx->f32, args, 5,
AC_FUNC_ATTR_READNONE);
0);
args[0] = p1;
args[1] = j;
@@ -1058,7 +1054,7 @@ LLVMValueRef ac_build_fs_interp_f16(struct ac_llvm_context *ctx, LLVMValueRef ll
args[5] = params;
return ac_build_intrinsic(ctx, "llvm.amdgcn.interp.p2.f16", ctx->f16, args, 6,
AC_FUNC_ATTR_READNONE);
0);
}
}
@@ -1076,17 +1072,16 @@ LLVMValueRef ac_build_fs_interp_mov(struct ac_llvm_context *ctx, LLVMValueRef pa
args[2] = params;
p = ac_build_intrinsic(ctx, "llvm.amdgcn.lds.param.load",
ctx->f32, args, 3, AC_FUNC_ATTR_READNONE);
ctx->f32, args, 3, 0);
p = ac_build_quad_swizzle(ctx, p, 0, 0, 0 ,0);
return ac_build_intrinsic(ctx, "llvm.amdgcn.wqm.f32", ctx->f32, &p, 1, AC_FUNC_ATTR_READNONE);
return ac_build_intrinsic(ctx, "llvm.amdgcn.wqm.f32", ctx->f32, &p, 1, 0);
} else {
args[0] = parameter;
args[1] = llvm_chan;
args[2] = attr_number;
args[3] = params;
return ac_build_intrinsic(ctx, "llvm.amdgcn.interp.mov", ctx->f32, args, 4,
AC_FUNC_ATTR_READNONE);
return ac_build_intrinsic(ctx, "llvm.amdgcn.interp.mov", ctx->f32, args, 4, 0);
}
}
@@ -1502,7 +1497,7 @@ static LLVMValueRef ac_ufN_to_float(struct ac_llvm_context *ctx, LLVMValueRef sr
mantissa, ctx->i1true, /* result can be undef when arg is 0 */
};
LLVMValueRef ctlz =
ac_build_intrinsic(ctx, "llvm.ctlz.i32", ctx->i32, params, 2, AC_FUNC_ATTR_READNONE);
ac_build_intrinsic(ctx, "llvm.ctlz.i32", ctx->i32, params, 2, 0);
/* Shift such that the leading 1 ends up as the LSB of the exponent field. */
tmp = LLVMBuildSub(ctx->builder, ctlz, LLVMConstInt(ctx->i32, 8, false), "");
@@ -1872,7 +1867,7 @@ void ac_build_sendmsg(struct ac_llvm_context *ctx, uint32_t msg, LLVMValueRef wa
LLVMValueRef ac_build_imsb(struct ac_llvm_context *ctx, LLVMValueRef arg, LLVMTypeRef dst_type)
{
LLVMValueRef msb =
ac_build_intrinsic(ctx, "llvm.amdgcn.sffbh.i32", dst_type, &arg, 1, AC_FUNC_ATTR_READNONE);
ac_build_intrinsic(ctx, "llvm.amdgcn.sffbh.i32", dst_type, &arg, 1, 0);
/* The HW returns the last bit index from MSB, but NIR/TGSI wants
* the index from LSB. Invert it by doing "31 - msb". */
@@ -1931,7 +1926,7 @@ LLVMValueRef ac_build_umsb(struct ac_llvm_context *ctx, LLVMValueRef arg, LLVMTy
ctx->i1true,
};
LLVMValueRef msb = ac_build_intrinsic(ctx, intrin_name, type, params, 2, AC_FUNC_ATTR_READNONE);
LLVMValueRef msb = ac_build_intrinsic(ctx, intrin_name, type, params, 2, 0);
if (!rev) {
/* The HW returns the last bit index from MSB, but TGSI/NIR wants
@@ -1957,7 +1952,7 @@ LLVMValueRef ac_build_fmin(struct ac_llvm_context *ctx, LLVMValueRef a, LLVMValu
ac_build_type_name_for_intr(LLVMTypeOf(a), type, sizeof(type));
snprintf(name, sizeof(name), "llvm.minnum.%s", type);
LLVMValueRef args[2] = {a, b};
return ac_build_intrinsic(ctx, name, LLVMTypeOf(a), args, 2, AC_FUNC_ATTR_READNONE);
return ac_build_intrinsic(ctx, name, LLVMTypeOf(a), args, 2, 0);
}
LLVMValueRef ac_build_fmax(struct ac_llvm_context *ctx, LLVMValueRef a, LLVMValueRef b)
@@ -1967,7 +1962,7 @@ LLVMValueRef ac_build_fmax(struct ac_llvm_context *ctx, LLVMValueRef a, LLVMValu
ac_build_type_name_for_intr(LLVMTypeOf(a), type, sizeof(type));
snprintf(name, sizeof(name), "llvm.maxnum.%s", type);
LLVMValueRef args[2] = {a, b};
return ac_build_intrinsic(ctx, name, LLVMTypeOf(a), args, 2, AC_FUNC_ATTR_READNONE);
return ac_build_intrinsic(ctx, name, LLVMTypeOf(a), args, 2, 0);
}
LLVMValueRef ac_build_imin(struct ac_llvm_context *ctx, LLVMValueRef a, LLVMValueRef b)
@@ -2373,21 +2368,18 @@ LLVMValueRef ac_build_image_get_sample_count(struct ac_llvm_context *ctx, LLVMVa
LLVMValueRef ac_build_cvt_pkrtz_f16(struct ac_llvm_context *ctx, LLVMValueRef args[2])
{
return ac_build_intrinsic(ctx, "llvm.amdgcn.cvt.pkrtz", ctx->v2f16, args, 2,
AC_FUNC_ATTR_READNONE);
return ac_build_intrinsic(ctx, "llvm.amdgcn.cvt.pkrtz", ctx->v2f16, args, 2, 0);
}
LLVMValueRef ac_build_cvt_pknorm_i16(struct ac_llvm_context *ctx, LLVMValueRef args[2])
{
LLVMValueRef res = ac_build_intrinsic(ctx, "llvm.amdgcn.cvt.pknorm.i16", ctx->v2i16, args, 2,
AC_FUNC_ATTR_READNONE);
LLVMValueRef res = ac_build_intrinsic(ctx, "llvm.amdgcn.cvt.pknorm.i16", ctx->v2i16, args, 2, 0);
return LLVMBuildBitCast(ctx->builder, res, ctx->i32, "");
}
LLVMValueRef ac_build_cvt_pknorm_u16(struct ac_llvm_context *ctx, LLVMValueRef args[2])
{
LLVMValueRef res = ac_build_intrinsic(ctx, "llvm.amdgcn.cvt.pknorm.u16", ctx->v2i16, args, 2,
AC_FUNC_ATTR_READNONE);
LLVMValueRef res = ac_build_intrinsic(ctx, "llvm.amdgcn.cvt.pknorm.u16", ctx->v2i16, args, 2, 0);
return LLVMBuildBitCast(ctx->builder, res, ctx->i32, "");
}
@@ -2438,7 +2430,7 @@ LLVMValueRef ac_build_cvt_pk_i16(struct ac_llvm_context *ctx, LLVMValueRef args[
}
LLVMValueRef res =
ac_build_intrinsic(ctx, "llvm.amdgcn.cvt.pk.i16", ctx->v2i16, args, 2, AC_FUNC_ATTR_READNONE);
ac_build_intrinsic(ctx, "llvm.amdgcn.cvt.pk.i16", ctx->v2i16, args, 2, 0);
return LLVMBuildBitCast(ctx->builder, res, ctx->i32, "");
}
@@ -2460,13 +2452,13 @@ LLVMValueRef ac_build_cvt_pk_u16(struct ac_llvm_context *ctx, LLVMValueRef args[
}
LLVMValueRef res =
ac_build_intrinsic(ctx, "llvm.amdgcn.cvt.pk.u16", ctx->v2i16, args, 2, AC_FUNC_ATTR_READNONE);
ac_build_intrinsic(ctx, "llvm.amdgcn.cvt.pk.u16", ctx->v2i16, args, 2, 0);
return LLVMBuildBitCast(ctx->builder, res, ctx->i32, "");
}
LLVMValueRef ac_build_wqm_vote(struct ac_llvm_context *ctx, LLVMValueRef i1)
{
return ac_build_intrinsic(ctx, "llvm.amdgcn.wqm.vote", ctx->i1, &i1, 1, AC_FUNC_ATTR_READNONE);
return ac_build_intrinsic(ctx, "llvm.amdgcn.wqm.vote", ctx->i1, &i1, 1, 0);
}
void ac_build_kill_if_false(struct ac_llvm_context *ctx, LLVMValueRef i1)
@@ -2484,7 +2476,7 @@ LLVMValueRef ac_build_bfe(struct ac_llvm_context *ctx, LLVMValueRef input, LLVMV
};
return ac_build_intrinsic(ctx, is_signed ? "llvm.amdgcn.sbfe.i32" : "llvm.amdgcn.ubfe.i32",
ctx->i32, args, 3, AC_FUNC_ATTR_READNONE);
ctx->i32, args, 3, 0);
}
LLVMValueRef ac_build_imad(struct ac_llvm_context *ctx, LLVMValueRef s0, LLVMValueRef s1,
@@ -2497,10 +2489,8 @@ LLVMValueRef ac_build_fmad(struct ac_llvm_context *ctx, LLVMValueRef s0, LLVMVal
LLVMValueRef s2)
{
/* FMA is better on GFX10, because it has FMA units instead of MUL-ADD units. */
if (ctx->gfx_level >= GFX10) {
return ac_build_intrinsic(ctx, "llvm.fma.f32", ctx->f32, (LLVMValueRef[]){s0, s1, s2}, 3,
AC_FUNC_ATTR_READNONE);
}
if (ctx->gfx_level >= GFX10)
return ac_build_intrinsic(ctx, "llvm.fma.f32", ctx->f32, (LLVMValueRef[]){s0, s1, s2}, 3, 0);
return LLVMBuildFAdd(ctx->builder, LLVMBuildFMul(ctx->builder, s0, s1, ""), s2, "");
}
@@ -2582,8 +2572,7 @@ LLVMValueRef ac_build_fsat(struct ac_llvm_context *ctx, LLVMValueRef src,
src,
};
result = ac_build_intrinsic(ctx, intr, type, params, 3,
AC_FUNC_ATTR_READNONE);
result = ac_build_intrinsic(ctx, intr, type, params, 3, 0);
}
if (ctx->gfx_level < GFX9 && bitsize == 32) {
@@ -2613,7 +2602,7 @@ LLVMValueRef ac_build_fract(struct ac_llvm_context *ctx, LLVMValueRef src0, unsi
LLVMValueRef params[] = {
src0,
};
return ac_build_intrinsic(ctx, intr, type, params, 1, AC_FUNC_ATTR_READNONE);
return ac_build_intrinsic(ctx, intr, type, params, 1, 0);
}
LLVMValueRef ac_const_uint_vec(struct ac_llvm_context *ctx, LLVMTypeRef type, uint64_t value)
@@ -2699,30 +2688,22 @@ LLVMValueRef ac_build_bit_count(struct ac_llvm_context *ctx, LLVMValueRef src0)
switch (bitsize) {
case 128:
result = ac_build_intrinsic(ctx, "llvm.ctpop.i128", ctx->i128, (LLVMValueRef[]){src0}, 1,
AC_FUNC_ATTR_READNONE);
result = ac_build_intrinsic(ctx, "llvm.ctpop.i128", ctx->i128, (LLVMValueRef[]){src0}, 1, 0);
result = LLVMBuildTrunc(ctx->builder, result, ctx->i32, "");
break;
case 64:
result = ac_build_intrinsic(ctx, "llvm.ctpop.i64", ctx->i64, (LLVMValueRef[]){src0}, 1,
AC_FUNC_ATTR_READNONE);
result = ac_build_intrinsic(ctx, "llvm.ctpop.i64", ctx->i64, (LLVMValueRef[]){src0}, 1, 0);
result = LLVMBuildTrunc(ctx->builder, result, ctx->i32, "");
break;
case 32:
result = ac_build_intrinsic(ctx, "llvm.ctpop.i32", ctx->i32, (LLVMValueRef[]){src0}, 1,
AC_FUNC_ATTR_READNONE);
result = ac_build_intrinsic(ctx, "llvm.ctpop.i32", ctx->i32, (LLVMValueRef[]){src0}, 1, 0);
break;
case 16:
result = ac_build_intrinsic(ctx, "llvm.ctpop.i16", ctx->i16, (LLVMValueRef[]){src0}, 1,
AC_FUNC_ATTR_READNONE);
result = ac_build_intrinsic(ctx, "llvm.ctpop.i16", ctx->i16, (LLVMValueRef[]){src0}, 1, 0);
result = LLVMBuildZExt(ctx->builder, result, ctx->i32, "");
break;
case 8:
result = ac_build_intrinsic(ctx, "llvm.ctpop.i8", ctx->i8, (LLVMValueRef[]){src0}, 1,
AC_FUNC_ATTR_READNONE);
result = ac_build_intrinsic(ctx, "llvm.ctpop.i8", ctx->i8, (LLVMValueRef[]){src0}, 1, 0);
result = LLVMBuildZExt(ctx->builder, result, ctx->i32, "");
break;
default:
@@ -2742,25 +2723,18 @@ LLVMValueRef ac_build_bitfield_reverse(struct ac_llvm_context *ctx, LLVMValueRef
switch (bitsize) {
case 64:
result = ac_build_intrinsic(ctx, "llvm.bitreverse.i64", ctx->i64, (LLVMValueRef[]){src0}, 1,
AC_FUNC_ATTR_READNONE);
result = ac_build_intrinsic(ctx, "llvm.bitreverse.i64", ctx->i64, (LLVMValueRef[]){src0}, 1, 0);
result = LLVMBuildTrunc(ctx->builder, result, ctx->i32, "");
break;
case 32:
result = ac_build_intrinsic(ctx, "llvm.bitreverse.i32", ctx->i32, (LLVMValueRef[]){src0}, 1,
AC_FUNC_ATTR_READNONE);
result = ac_build_intrinsic(ctx, "llvm.bitreverse.i32", ctx->i32, (LLVMValueRef[]){src0}, 1, 0);
break;
case 16:
result = ac_build_intrinsic(ctx, "llvm.bitreverse.i16", ctx->i16, (LLVMValueRef[]){src0}, 1,
AC_FUNC_ATTR_READNONE);
result = ac_build_intrinsic(ctx, "llvm.bitreverse.i16", ctx->i16, (LLVMValueRef[]){src0}, 1, 0);
result = LLVMBuildZExt(ctx->builder, result, ctx->i32, "");
break;
case 8:
result = ac_build_intrinsic(ctx, "llvm.bitreverse.i8", ctx->i8, (LLVMValueRef[]){src0}, 1,
AC_FUNC_ATTR_READNONE);
result = ac_build_intrinsic(ctx, "llvm.bitreverse.i8", ctx->i8, (LLVMValueRef[]){src0}, 1, 0);
result = LLVMBuildZExt(ctx->builder, result, ctx->i32, "");
break;
default:
@@ -2784,7 +2758,7 @@ LLVMValueRef ac_build_sudot_4x8(struct ac_llvm_context *ctx, LLVMValueRef s0, LL
src[4] = s2;
src[5] = LLVMConstInt(ctx->i1, clamp, false);
return ac_build_intrinsic(ctx, name, ctx->i32, src, 6, AC_FUNC_ATTR_READNONE);
return ac_build_intrinsic(ctx, name, ctx->i32, src, 6, 0);
}
void ac_init_exec_full_mask(struct ac_llvm_context *ctx)
@@ -2863,7 +2837,7 @@ LLVMValueRef ac_find_lsb(struct ac_llvm_context *ctx, LLVMTypeRef dst_type, LLVM
ctx->i1true,
};
LLVMValueRef lsb = ac_build_intrinsic(ctx, intrin_name, type, params, 2, AC_FUNC_ATTR_READNONE);
LLVMValueRef lsb = ac_build_intrinsic(ctx, intrin_name, type, params, 2, 0);
if (src0_bitsize == 64) {
lsb = LLVMBuildTrunc(ctx->builder, lsb, ctx->i32, "");
@@ -3222,8 +3196,7 @@ static LLVMValueRef _ac_build_readlane(struct ac_llvm_context *ctx, LLVMValueRef
result =
ac_build_intrinsic(ctx, lane == NULL ? "llvm.amdgcn.readfirstlane" : "llvm.amdgcn.readlane",
ctx->i32, (LLVMValueRef[]){src, lane}, lane == NULL ? 1 : 2,
AC_FUNC_ATTR_READNONE);
ctx->i32, (LLVMValueRef[]){src, lane}, lane == NULL ? 1 : 2, 0);
return LLVMBuildTrunc(ctx->builder, result, type, "");
}
@@ -3286,8 +3259,7 @@ LLVMValueRef ac_build_writelane(struct ac_llvm_context *ctx, LLVMValueRef src, L
LLVMValueRef lane)
{
return ac_build_intrinsic(ctx, "llvm.amdgcn.writelane", ctx->i32,
(LLVMValueRef[]){value, lane, src}, 3,
AC_FUNC_ATTR_READNONE);
(LLVMValueRef[]){value, lane, src}, 3, 0);
}
LLVMValueRef ac_build_mbcnt_add(struct ac_llvm_context *ctx, LLVMValueRef mask, LLVMValueRef add_src)
@@ -3297,15 +3269,15 @@ LLVMValueRef ac_build_mbcnt_add(struct ac_llvm_context *ctx, LLVMValueRef mask,
if (ctx->wave_size == 32) {
val = ac_build_intrinsic(ctx, "llvm.amdgcn.mbcnt.lo", ctx->i32,
(LLVMValueRef[]){mask, add}, 2, AC_FUNC_ATTR_READNONE);
(LLVMValueRef[]){mask, add}, 2, 0);
} else {
LLVMValueRef mask_vec = LLVMBuildBitCast(ctx->builder, mask, ctx->v2i32, "");
LLVMValueRef mask_lo = LLVMBuildExtractElement(ctx->builder, mask_vec, ctx->i32_0, "");
LLVMValueRef mask_hi = LLVMBuildExtractElement(ctx->builder, mask_vec, ctx->i32_1, "");
val = ac_build_intrinsic(ctx, "llvm.amdgcn.mbcnt.lo", ctx->i32,
(LLVMValueRef[]){mask_lo, add}, 2, AC_FUNC_ATTR_READNONE);
(LLVMValueRef[]){mask_lo, add}, 2, 0);
val = ac_build_intrinsic(ctx, "llvm.amdgcn.mbcnt.hi", ctx->i32, (LLVMValueRef[]){mask_hi, val},
2, AC_FUNC_ATTR_READNONE);
2, 0);
}
if (add == ctx->i32_0)
@@ -3371,7 +3343,7 @@ static LLVMValueRef _ac_build_dpp(struct ac_llvm_context *ctx, LLVMValueRef old,
(LLVMValueRef[]){old, src, LLVMConstInt(ctx->i32, dpp_ctrl, 0),
LLVMConstInt(ctx->i32, row_mask, 0), LLVMConstInt(ctx->i32, bank_mask, 0),
LLVMConstInt(ctx->i1, bound_ctrl, 0)},
6, AC_FUNC_ATTR_READNONE);
6, 0);
return LLVMBuildTrunc(ctx->builder, res, type, "");
}
@@ -3424,7 +3396,7 @@ static LLVMValueRef _ac_build_permlane16(struct ac_llvm_context *ctx, LLVMValueR
result =
ac_build_intrinsic(ctx, exchange_rows ? "llvm.amdgcn.permlanex16" : "llvm.amdgcn.permlane16",
ctx->i32, args, 6, AC_FUNC_ATTR_READNONE);
ctx->i32, args, 6, 0);
return LLVMBuildTrunc(ctx->builder, result, type, "");
}
@@ -3469,7 +3441,7 @@ static LLVMValueRef _ac_build_ds_swizzle(struct ac_llvm_context *ctx, LLVMValueR
ret = ac_build_intrinsic(ctx, "llvm.amdgcn.ds.swizzle", ctx->i32,
(LLVMValueRef[]){src, LLVMConstInt(ctx->i32, mask, 0)}, 2,
AC_FUNC_ATTR_READNONE);
0);
return LLVMBuildTrunc(ctx->builder, ret, src_type, "");
}
@@ -3511,8 +3483,7 @@ static LLVMValueRef ac_build_wwm(struct ac_llvm_context *ctx, LLVMValueRef src)
ac_build_type_name_for_intr(LLVMTypeOf(src), type, sizeof(type));
snprintf(name, sizeof(name), "llvm.amdgcn.wwm.%s", type);
ret = ac_build_intrinsic(ctx, name, LLVMTypeOf(src), (LLVMValueRef[]){src}, 1,
AC_FUNC_ATTR_READNONE);
ret = ac_build_intrinsic(ctx, name, LLVMTypeOf(src), (LLVMValueRef[]){src}, 1, 0);
if (bitsize < 32)
ret = LLVMBuildTrunc(ctx->builder, ret, ac_to_integer_type(ctx, src_type), "");
@@ -3537,8 +3508,7 @@ static LLVMValueRef ac_build_set_inactive(struct ac_llvm_context *ctx, LLVMValue
ac_build_type_name_for_intr(LLVMTypeOf(src), type, sizeof(type));
snprintf(name, sizeof(name), "llvm.amdgcn.set.inactive.%s", type);
LLVMValueRef ret =
ac_build_intrinsic(ctx, name, LLVMTypeOf(src), (LLVMValueRef[]){src, inactive}, 2,
AC_FUNC_ATTR_READNONE);
ac_build_intrinsic(ctx, name, LLVMTypeOf(src), (LLVMValueRef[]){src, inactive}, 2, 0);
if (bitsize < 32)
ret = LLVMBuildTrunc(ctx->builder, ret, src_type, "");
@@ -3701,8 +3671,7 @@ static LLVMValueRef ac_build_alu_op(struct ac_llvm_context *ctx, LLVMValueRef lh
case nir_op_fmin:
return ac_build_intrinsic(
ctx, _64bit ? "llvm.minnum.f64" : _32bit ? "llvm.minnum.f32" : "llvm.minnum.f16",
_64bit ? ctx->f64 : _32bit ? ctx->f32 : ctx->f16, (LLVMValueRef[]){lhs, rhs}, 2,
AC_FUNC_ATTR_READNONE);
_64bit ? ctx->f64 : _32bit ? ctx->f32 : ctx->f16, (LLVMValueRef[]){lhs, rhs}, 2, 0);
case nir_op_imax:
return LLVMBuildSelect(ctx->builder, LLVMBuildICmp(ctx->builder, LLVMIntSGT, lhs, rhs, ""),
lhs, rhs, "");
@@ -3712,8 +3681,7 @@ static LLVMValueRef ac_build_alu_op(struct ac_llvm_context *ctx, LLVMValueRef lh
case nir_op_fmax:
return ac_build_intrinsic(
ctx, _64bit ? "llvm.maxnum.f64" : _32bit ? "llvm.maxnum.f32" : "llvm.maxnum.f16",
_64bit ? ctx->f64 : _32bit ? ctx->f32 : ctx->f16, (LLVMValueRef[]){lhs, rhs}, 2,
AC_FUNC_ATTR_READNONE);
_64bit ? ctx->f64 : _32bit ? ctx->f32 : ctx->f16, (LLVMValueRef[]){lhs, rhs}, 2, 0);
case nir_op_iand:
return LLVMBuildAnd(ctx->builder, lhs, rhs, "");
case nir_op_ior:
@@ -4256,8 +4224,7 @@ LLVMValueRef ac_build_shuffle(struct ac_llvm_context *ctx, LLVMValueRef src, LLV
src = LLVMBuildZExt(ctx->builder, src, ctx->i32, "");
result =
ac_build_intrinsic(ctx, "llvm.amdgcn.ds.bpermute", ctx->i32, (LLVMValueRef[]){index, src}, 2,
AC_FUNC_ATTR_READNONE);
ac_build_intrinsic(ctx, "llvm.amdgcn.ds.bpermute", ctx->i32, (LLVMValueRef[]){index, src}, 2, 0);
return LLVMBuildTrunc(ctx->builder, result, type, "");
}
@@ -4280,7 +4247,7 @@ LLVMValueRef ac_build_frexp_exp(struct ac_llvm_context *ctx, LLVMValueRef src0,
LLVMValueRef params[] = {
src0,
};
return ac_build_intrinsic(ctx, intr, type, params, 1, AC_FUNC_ATTR_READNONE);
return ac_build_intrinsic(ctx, intr, type, params, 1, 0);
}
LLVMValueRef ac_build_frexp_mant(struct ac_llvm_context *ctx, LLVMValueRef src0, unsigned bitsize)
{
@@ -4301,7 +4268,7 @@ LLVMValueRef ac_build_frexp_mant(struct ac_llvm_context *ctx, LLVMValueRef src0,
LLVMValueRef params[] = {
src0,
};
return ac_build_intrinsic(ctx, intr, type, params, 1, AC_FUNC_ATTR_READNONE);
return ac_build_intrinsic(ctx, intr, type, params, 1, 0);
}
LLVMValueRef ac_build_canonicalize(struct ac_llvm_context *ctx, LLVMValueRef src0, unsigned bitsize)
@@ -4323,7 +4290,7 @@ LLVMValueRef ac_build_canonicalize(struct ac_llvm_context *ctx, LLVMValueRef src
LLVMValueRef params[] = {
src0,
};
return ac_build_intrinsic(ctx, intr, type, params, 1, AC_FUNC_ATTR_READNONE);
return ac_build_intrinsic(ctx, intr, type, params, 1, 0);
}
/*
@@ -4351,8 +4318,7 @@ LLVMValueRef ac_build_load_helper_invocation(struct ac_llvm_context *ctx)
if (LLVM_VERSION_MAJOR >= 13) {
result = ac_build_intrinsic(ctx, "llvm.amdgcn.live.mask", ctx->i1, NULL, 0, 0);
} else {
result = ac_build_intrinsic(ctx, "llvm.amdgcn.ps.live", ctx->i1, NULL, 0,
AC_FUNC_ATTR_READNONE);
result = ac_build_intrinsic(ctx, "llvm.amdgcn.ps.live", ctx->i1, NULL, 0, 0);
}
return LLVMBuildNot(ctx->builder, result, "");
}
@@ -4367,7 +4333,7 @@ LLVMValueRef ac_build_is_helper_invocation(struct ac_llvm_context *ctx)
/* !(exact && postponed) */
LLVMValueRef exact =
ac_build_intrinsic(ctx, "llvm.amdgcn.ps.live", ctx->i1, NULL, 0, AC_FUNC_ATTR_READNONE);
ac_build_intrinsic(ctx, "llvm.amdgcn.ps.live", ctx->i1, NULL, 0, 0);
LLVMValueRef postponed = LLVMBuildLoad2(ctx->builder, ctx->i1, ctx->postponed_kill, "");
return LLVMBuildNot(ctx->builder, LLVMBuildAnd(ctx->builder, exact, postponed, ""), "");
@@ -4691,6 +4657,5 @@ LLVMValueRef ac_build_is_inf_or_nan(struct ac_llvm_context *ctx, LLVMValueRef a)
a,
LLVMConstInt(ctx->i32, S_NAN | Q_NAN | N_INFINITY | P_INFINITY, 0),
};
return ac_build_intrinsic(ctx, "llvm.amdgcn.class.f32", ctx->i1, args, 2,
AC_FUNC_ATTR_READNONE);
return ac_build_intrinsic(ctx, "llvm.amdgcn.class.f32", ctx->i1, args, 2, 0);
}

63
src/amd/llvm/ac_nir_to_llvm.c
View File

@@ -155,7 +155,7 @@ static LLVMValueRef emit_intrin_1f_param(struct ac_llvm_context *ctx, const char
ac_build_type_name_for_intr(LLVMTypeOf(params[0]), type, sizeof(type));
ASSERTED const int length = snprintf(name, sizeof(name), "%s.%s", intrin, type);
assert(length < sizeof(name));
return ac_build_intrinsic(ctx, name, result_type, params, 1, AC_FUNC_ATTR_READNONE);
return ac_build_intrinsic(ctx, name, result_type, params, 1, 0);
}
static LLVMValueRef emit_intrin_1f_param_scalar(struct ac_llvm_context *ctx, const char *intrin,
@@ -179,7 +179,7 @@ static LLVMValueRef emit_intrin_1f_param_scalar(struct ac_llvm_context *ctx, con
assert(length < sizeof(name));
ret = LLVMBuildInsertElement(
ctx->builder, ret,
ac_build_intrinsic(ctx, name, elem_type, params, 1, AC_FUNC_ATTR_READNONE),
ac_build_intrinsic(ctx, name, elem_type, params, 1, 0),
LLVMConstInt(ctx->i32, i, 0), "");
}
return ret;
@@ -198,7 +198,7 @@ static LLVMValueRef emit_intrin_2f_param(struct ac_llvm_context *ctx, const char
ac_build_type_name_for_intr(LLVMTypeOf(params[0]), type, sizeof(type));
ASSERTED const int length = snprintf(name, sizeof(name), "%s.%s", intrin, type);
assert(length < sizeof(name));
return ac_build_intrinsic(ctx, name, result_type, params, 2, AC_FUNC_ATTR_READNONE);
return ac_build_intrinsic(ctx, name, result_type, params, 2, 0);
}
static LLVMValueRef emit_intrin_3f_param(struct ac_llvm_context *ctx, const char *intrin,
@@ -215,7 +215,7 @@ static LLVMValueRef emit_intrin_3f_param(struct ac_llvm_context *ctx, const char
ac_build_type_name_for_intr(LLVMTypeOf(params[0]), type, sizeof(type));
ASSERTED const int length = snprintf(name, sizeof(name), "%s.%s", intrin, type);
assert(length < sizeof(name));
return ac_build_intrinsic(ctx, name, result_type, params, 3, AC_FUNC_ATTR_READNONE);
return ac_build_intrinsic(ctx, name, result_type, params, 3, 0);
}
static LLVMValueRef emit_bcsel(struct ac_llvm_context *ctx, LLVMValueRef src0, LLVMValueRef src1,
@@ -250,7 +250,7 @@ static LLVMValueRef emit_uint_carry(struct ac_llvm_context *ctx, const char *int
LLVMValueRef params[] = {src0, src1};
ret_type = LLVMStructTypeInContext(ctx->context, types, 2, true);
res = ac_build_intrinsic(ctx, intrin, ret_type, params, 2, AC_FUNC_ATTR_READNONE);
res = ac_build_intrinsic(ctx, intrin, ret_type, params, 2, 0);
res = LLVMBuildExtractValue(ctx->builder, res, 1, "");
res = LLVMBuildZExt(ctx->builder, res, ctx->i32, "");
@@ -326,7 +326,7 @@ static LLVMValueRef emit_f2f16(struct ac_llvm_context *ctx, LLVMValueRef src0)
args[0] = result;
args[1] = LLVMConstInt(ctx->i32, N_SUBNORMAL | P_SUBNORMAL, false);
cond =
ac_build_intrinsic(ctx, "llvm.amdgcn.class.f16", ctx->i1, args, 2, AC_FUNC_ATTR_READNONE);
ac_build_intrinsic(ctx, "llvm.amdgcn.class.f16", ctx->i1, args, 2, 0);
}
/* need to convert back up to f32 */
@@ -627,7 +627,7 @@ static bool visit_alu(struct ac_nir_context *ctx, const nir_alu_instr *instr)
ac_build_type_name_for_intr(def_type, type, sizeof(type));
snprintf(name, sizeof(name), "llvm.%cadd.sat.%s",
instr->op == nir_op_uadd_sat ? 'u' : 's', type);
result = ac_build_intrinsic(&ctx->ac, name, def_type, src, 2, AC_FUNC_ATTR_READNONE);
result = ac_build_intrinsic(&ctx->ac, name, def_type, src, 2, 0);
break;
}
case nir_op_usub_sat:
@@ -636,7 +636,7 @@ static bool visit_alu(struct ac_nir_context *ctx, const nir_alu_instr *instr)
ac_build_type_name_for_intr(def_type, type, sizeof(type));
snprintf(name, sizeof(name), "llvm.%csub.sat.%s",
instr->op == nir_op_usub_sat ? 'u' : 's', type);
result = ac_build_intrinsic(&ctx->ac, name, def_type, src, 2, AC_FUNC_ATTR_READNONE);
result = ac_build_intrinsic(&ctx->ac, name, def_type, src, 2, 0);
break;
}
case nir_op_fadd:
@@ -690,7 +690,7 @@ static bool visit_alu(struct ac_nir_context *ctx, const nir_alu_instr *instr)
src[0] = ac_to_float(&ctx->ac, src[0]);
src[1] = ac_to_float(&ctx->ac, src[1]);
result = ac_build_intrinsic(&ctx->ac, "llvm.amdgcn.fmul.legacy", ctx->ac.f32,
src, 2, AC_FUNC_ATTR_READNONE);
src, 2, 0);
break;
case nir_op_frcp:
/* For doubles, we need precise division to pass GLCTS. */
@@ -889,7 +889,7 @@ static bool visit_alu(struct ac_nir_context *ctx, const nir_alu_instr *instr)
ac_to_float_type(&ctx->ac, def_type), src[0]);
result = ac_build_intrinsic(&ctx->ac, "llvm.amdgcn.fmul.legacy", ctx->ac.f32,
(LLVMValueRef[]){result, ac_to_float(&ctx->ac, src[1])},
2, AC_FUNC_ATTR_READNONE);
2, 0);
result = emit_intrin_1f_param(&ctx->ac, "llvm.exp2",
ac_to_float_type(&ctx->ac, def_type), result);
break;
@@ -926,19 +926,19 @@ static bool visit_alu(struct ac_nir_context *ctx, const nir_alu_instr *instr)
src[1] = ac_to_float(&ctx->ac, src[1]);
src[2] = ac_to_float(&ctx->ac, src[2]);
result = ac_build_intrinsic(&ctx->ac, "llvm.amdgcn.fma.legacy", ctx->ac.f32,
src, 3, AC_FUNC_ATTR_READNONE);
src, 3, 0);
break;
case nir_op_ldexp:
src[0] = ac_to_float(&ctx->ac, src[0]);
if (ac_get_elem_bits(&ctx->ac, def_type) == 32)
result = ac_build_intrinsic(&ctx->ac, "llvm.amdgcn.ldexp.f32", ctx->ac.f32, src, 2,
AC_FUNC_ATTR_READNONE);
0);
else if (ac_get_elem_bits(&ctx->ac, def_type) == 16)
result = ac_build_intrinsic(&ctx->ac, "llvm.amdgcn.ldexp.f16", ctx->ac.f16, src, 2,
AC_FUNC_ATTR_READNONE);
0);
else
result = ac_build_intrinsic(&ctx->ac, "llvm.amdgcn.ldexp.f64", ctx->ac.f64, src, 2,
AC_FUNC_ATTR_READNONE);
0);
break;
case nir_op_bfm:
result = emit_bfm(&ctx->ac, src[0], src[1]);
@@ -1082,14 +1082,14 @@ static bool visit_alu(struct ac_nir_context *ctx, const nir_alu_instr *instr)
break;
case nir_op_ifind_msb_rev:
result = ac_build_intrinsic(&ctx->ac, "llvm.amdgcn.sffbh.i32", ctx->ac.i32, &src[0], 1,
AC_FUNC_ATTR_READNONE);
0);
break;
case nir_op_uclz: {
LLVMValueRef params[2] = {
src[0],
ctx->ac.i1false,
};
result = ac_build_intrinsic(&ctx->ac, "llvm.ctlz.i32", ctx->ac.i32, params, 2, AC_FUNC_ATTR_READNONE);
result = ac_build_intrinsic(&ctx->ac, "llvm.ctlz.i32", ctx->ac.i32, params, 2, 0);
break;
}
case nir_op_uadd_carry:
@@ -1262,11 +1262,10 @@ static bool visit_alu(struct ac_nir_context *ctx, const nir_alu_instr *instr)
for (unsigned chan = 0; chan < 3; chan++)
in[chan] = ac_llvm_extract_elem(&ctx->ac, src[0], chan);
results[0] = ac_build_intrinsic(&ctx->ac, "llvm.amdgcn.cubesc", ctx->ac.f32, in, 3,
AC_FUNC_ATTR_READNONE);
0);
results[1] = ac_build_intrinsic(&ctx->ac, "llvm.amdgcn.cubetc", ctx->ac.f32, in, 3,
AC_FUNC_ATTR_READNONE);
LLVMValueRef ma = ac_build_intrinsic(&ctx->ac, "llvm.amdgcn.cubema", ctx->ac.f32, in, 3,
AC_FUNC_ATTR_READNONE);
0);
LLVMValueRef ma = ac_build_intrinsic(&ctx->ac, "llvm.amdgcn.cubema", ctx->ac.f32, in, 3, 0);
results[0] = ac_build_fdiv(&ctx->ac, results[0], ma);
results[1] = ac_build_fdiv(&ctx->ac, results[1], ma);
LLVMValueRef offset = LLVMConstReal(ctx->ac.f32, 0.5);
@@ -1281,8 +1280,7 @@ static bool visit_alu(struct ac_nir_context *ctx, const nir_alu_instr *instr)
LLVMValueRef in[3];
for (unsigned chan = 0; chan < 3; chan++)
in[chan] = ac_llvm_extract_elem(&ctx->ac, src[0], chan);
result = ac_build_intrinsic(&ctx->ac, "llvm.amdgcn.cubeid", ctx->ac.f32, in, 3,
AC_FUNC_ATTR_READNONE);
result = ac_build_intrinsic(&ctx->ac, "llvm.amdgcn.cubeid", ctx->ac.f32, in, 3, 0);
break;
}
@@ -1319,7 +1317,7 @@ static bool visit_alu(struct ac_nir_context *ctx, const nir_alu_instr *instr)
} else {
const char *name = "llvm.amdgcn.sdot4";
src[3] = LLVMConstInt(ctx->ac.i1, instr->op == nir_op_sdot_4x8_iadd_sat, false);
result = ac_build_intrinsic(&ctx->ac, name, def_type, src, 4, AC_FUNC_ATTR_READNONE);
result = ac_build_intrinsic(&ctx->ac, name, def_type, src, 4, 0);
}
break;
}
@@ -1333,7 +1331,7 @@ static bool visit_alu(struct ac_nir_context *ctx, const nir_alu_instr *instr)
case nir_op_udot_4x8_uadd_sat: {
const char *name = "llvm.amdgcn.udot4";
src[3] = LLVMConstInt(ctx->ac.i1, instr->op == nir_op_udot_4x8_uadd_sat, false);
result = ac_build_intrinsic(&ctx->ac, name, def_type, src, 4, AC_FUNC_ATTR_READNONE);
result = ac_build_intrinsic(&ctx->ac, name, def_type, src, 4, 0);
break;
}
@@ -1348,14 +1346,13 @@ static bool visit_alu(struct ac_nir_context *ctx, const nir_alu_instr *instr)
src[1] = LLVMBuildBitCast(ctx->ac.builder, src[1], ctx->ac.v2i16, "");
src[3] = LLVMConstInt(ctx->ac.i1, instr->op == nir_op_sdot_2x16_iadd_sat ||
instr->op == nir_op_udot_2x16_uadd_sat, false);
result = ac_build_intrinsic(&ctx->ac, name, def_type, src, 4, AC_FUNC_ATTR_READNONE);
result = ac_build_intrinsic(&ctx->ac, name, def_type, src, 4, 0);
break;
}
case nir_op_sad_u8x4:
result = ac_build_intrinsic(&ctx->ac, "llvm.amdgcn.sad.u8", ctx->ac.i32,
(LLVMValueRef[]){src[0], src[1], src[2]}, 3,
AC_FUNC_ATTR_READNONE);
(LLVMValueRef[]){src[0], src[1], src[2]}, 3, 0);
break;
default:
@@ -3080,8 +3077,7 @@ static LLVMValueRef visit_first_invocation(struct ac_nir_context *ctx)
/* The second argument is whether cttz(0) should be defined, but we do not care. */
LLVMValueRef args[] = {active_set, ctx->ac.i1false};
LLVMValueRef result = ac_build_intrinsic(&ctx->ac, intr, ctx->ac.iN_wavemask, args, 2,
AC_FUNC_ATTR_READNONE);
LLVMValueRef result = ac_build_intrinsic(&ctx->ac, intr, ctx->ac.iN_wavemask, args, 2, 0);
return LLVMBuildTrunc(ctx->ac.builder, result, ctx->ac.i32, "");
}
@@ -4072,8 +4068,7 @@ static bool visit_intrinsic(struct ac_nir_context *ctx, nir_intrinsic_instr *ins
src = LLVMBuildZExt(ctx->ac.builder, src, ctx->ac.i32, "");
result = ac_build_intrinsic(&ctx->ac, "llvm.amdgcn.readlane", ctx->ac.i32,
(LLVMValueRef[]){src, index_val}, 2,
AC_FUNC_ATTR_READNONE);
(LLVMValueRef[]){src, index_val}, 2, 0);
result = LLVMBuildTrunc(ctx->ac.builder, result, type, "");
@@ -4344,8 +4339,7 @@ static bool visit_intrinsic(struct ac_nir_context *ctx, nir_intrinsic_instr *ins
result = ac_build_intrinsic(&ctx->ac, "llvm.amdgcn.perm", ctx->ac.i32,
(LLVMValueRef[]){get_src(ctx, instr->src[0]),
get_src(ctx, instr->src[1]),
get_src(ctx, instr->src[2])},
3, AC_FUNC_ATTR_READNONE);
get_src(ctx, instr->src[2])}, 3, 0);
break;
case nir_intrinsic_lane_permute_16_amd:
result = ac_build_intrinsic(&ctx->ac, "llvm.amdgcn.permlane16", ctx->ac.i32,
@@ -4354,8 +4348,7 @@ static bool visit_intrinsic(struct ac_nir_context *ctx, nir_intrinsic_instr *ins
get_src(ctx, instr->src[1]),
get_src(ctx, instr->src[2]),
ctx->ac.i1false,
ctx->ac.i1false},
6, AC_FUNC_ATTR_READNONE);
ctx->ac.i1false}, 6, 0);
break;
case nir_intrinsic_load_force_vrs_rates_amd:
result = ac_get_arg(&ctx->ac, ctx->args->force_vrs_rates);

5
src/amd/vulkan/radv_nir_to_llvm.c
View File

@@ -175,8 +175,7 @@ create_function(struct radv_shader_context *ctx, gl_shader_stage stage, bool has
ctx->max_workgroup_size, ctx->options);
ctx->ring_offsets = ac_build_intrinsic(&ctx->ac, "llvm.amdgcn.implicit.buffer.ptr",
LLVMPointerType(ctx->ac.i8, AC_ADDR_SPACE_CONST), NULL, 0,
AC_FUNC_ATTR_READNONE);
LLVMPointerType(ctx->ac.i8, AC_ADDR_SPACE_CONST), NULL, 0, 0);
ctx->ring_offsets = LLVMBuildBitCast(ctx->ac.builder, ctx->ring_offsets,
ac_array_in_const_addr_space(ctx->ac.v4i32), "");
@@ -657,7 +656,7 @@ si_llvm_init_export_args(struct radv_shader_context *ctx, LLVMValueRef *values,
LLVMValueRef class_args[2] = {values[i],
LLVMConstInt(ctx->ac.i32, S_NAN | Q_NAN, false)};
LLVMValueRef isnan = ac_build_intrinsic(&ctx->ac, "llvm.amdgcn.class.f32", ctx->ac.i1,
class_args, 2, AC_FUNC_ATTR_READNONE);
class_args, 2, 0);
values[i] = LLVMBuildSelect(ctx->ac.builder, isnan, ctx->ac.f32_0, values[i], "");
}
}