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

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>

diff --git a/src/amd/llvm/ac_llvm_build.c b/src/amd/llvm/ac_llvm_build.c
index 04ca1c5..0a79b6b 100644 (file)
--- a/src/amd/llvm/ac_llvm_build.c
+++ b/src/amd/llvm/ac_llvm_build.c
@@ -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);
 }

diff --git a/src/amd/llvm/ac_nir_to_llvm.c b/src/amd/llvm/ac_nir_to_llvm.c
index 4d014fd..252f543 100644 (file)
--- a/src/amd/llvm/ac_nir_to_llvm.c
+++ b/src/amd/llvm/ac_nir_to_llvm.c
@@ -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);

diff --git a/src/amd/vulkan/radv_nir_to_llvm.c b/src/amd/vulkan/radv_nir_to_llvm.c
index e2f9361..c6fd844 100644 (file)
--- a/src/amd/vulkan/radv_nir_to_llvm.c
+++ b/src/amd/vulkan/radv_nir_to_llvm.c
@@ -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], "");
          }
       }