gentoo-mirror/gentoo
1
0
Fork 0
mirror of https://github.com/gentoo-mirror/gentoo.git synced 2026-08-25 07:08:13 -07:00
Code Issues Projects Releases Wiki Activity
Files
stable
gentoo/dev-libs/rocm-device-libs/files/rocm-device-libs-7.2.0-llvm-22-compat.patch
Sv. Lockal b7390c94c0 dev-libs/rocm-device-libs: add 7.2.0 
Signed-off-by: Sv. Lockal <lockalsash@gmail.com>
Part-of: https://github.com/gentoo/gentoo/pull/45826
Signed-off-by: Sam James <sam@gentoo.org>
2026-03-06 00:44:01 +00:00

234 lines
7.6 KiB
Diff
Raw Permalink Blame History

Fix error: '__builtin_amdgcn_lerp' needs target feature cube-insts.
Copies files from:
https://github.com/ROCm/llvm-project/blame/amd-staging/amd/device-libs/ockl/src/image.cl
https://github.com/ROCm/llvm-project/blame/amd-staging/amd/device-libs/ockl/src/media.cl
--- a/ockl/src/image.cl
+++ b/ockl/src/image.cl
@@ -13,6 +13,7 @@
#define EII() __oclc_ISA_version != 9010
#define RATTR __attribute__((pure))
+#define CRATTR __attribute__((pure, target("cube-insts")))
#define ERATTR __attribute__((pure, target("extended-image-insts")))
#define WATTR
#define GATTR __attribute__((const))
@@ -510,14 +511,14 @@ OCKL_MANGLE_T(image_load,3D)(TSHARP i, int4 c)
return my_image_load_3d_v4f32_i32(c.x, c.y, c.z, LOAD_TSHARP(i));
}
-RATTR float4
-OCKL_MANGLE_T(image_load,CM)(TSHARP i, int2 c, int f)
+CRATTR float4
+OCKL_MANGLE_T(image_load, CM)(TSHARP i, int2 c, int f)
{
return my_image_load_cube_v4f32_i32(c.x, c.y, f, LOAD_TSHARP(i));
}
-RATTR float4
-OCKL_MANGLE_T(image_load,CMa)(TSHARP i, int4 c, int f)
+CRATTR float4
+OCKL_MANGLE_T(image_load, CMa)(TSHARP i, int4 c, int f)
{
f = LS_ARRAY_FACE(c.z, f);
return my_image_load_cube_v4f32_i32(c.x, c.y, f, LOAD_TSHARP(i));
@@ -565,14 +566,14 @@ OCKL_MANGLE_T(image_load_lod,3D)(TSHARP i, int4 c, int l)
return my_image_load_mip_3d_v4f32_i32(c.x, c.y, c.z, l, LOAD_TSHARP(i));
}
-RATTR float4
-OCKL_MANGLE_T(image_load_lod,CM)(TSHARP i, int2 c, int f, int l)
+CRATTR float4
+OCKL_MANGLE_T(image_load_lod, CM)(TSHARP i, int2 c, int f, int l)
{
return my_image_load_mip_cube_v4f32_i32(c.x, c.y, f, l, LOAD_TSHARP(i));
}
-RATTR float4
-OCKL_MANGLE_T(image_load_lod,CMa)(TSHARP i, int4 c, int f, int l)
+CRATTR float4
+OCKL_MANGLE_T(image_load_lod, CMa)(TSHARP i, int4 c, int f, int l)
{
f = LS_ARRAY_FACE(c.z, f);
return my_image_load_mip_cube_v4f32_i32(c.x, c.y, f, l, LOAD_TSHARP(i));
@@ -614,14 +615,14 @@ OCKL_MANGLE_T(image_loadh,3D)(TSHARP i, int4 c)
return my_image_load_3d_v4f16_i32(c.x, c.y, c.z, LOAD_TSHARP(i));
}
-RATTR half4
-OCKL_MANGLE_T(image_loadh,CM)(TSHARP i, int2 c, int f)
+CRATTR half4
+OCKL_MANGLE_T(image_loadh, CM)(TSHARP i, int2 c, int f)
{
return my_image_load_cube_v4f16_i32(c.x, c.y, f, LOAD_TSHARP(i));
}
-RATTR half4
-OCKL_MANGLE_T(image_loadh,CMa)(TSHARP i, int4 c, int f)
+CRATTR half4
+OCKL_MANGLE_T(image_loadh, CMa)(TSHARP i, int4 c, int f)
{
f = LS_ARRAY_FACE(c.z, f);
return my_image_load_cube_v4f16_i32(c.x, c.y, f, LOAD_TSHARP(i));
@@ -657,14 +658,14 @@ OCKL_MANGLE_T(image_loadh_lod,3D)(TSHARP i, int4 c, int l)
return my_image_load_mip_3d_v4f16_i32(c.x, c.y, c.z, l, LOAD_TSHARP(i));
}
-RATTR half4
-OCKL_MANGLE_T(image_loadh_lod,CM)(TSHARP i, int2 c, int f, int l)
+CRATTR half4
+OCKL_MANGLE_T(image_loadh_lod, CM)(TSHARP i, int2 c, int f, int l)
{
return my_image_load_mip_cube_v4f16_i32(c.x, c.y, f, l, LOAD_TSHARP(i));
}
-RATTR half4
-OCKL_MANGLE_T(image_loadh_lod,CMa)(TSHARP i, int4 c, int f, int l)
+CRATTR half4
+OCKL_MANGLE_T(image_loadh_lod, CMa)(TSHARP i, int4 c, int f, int l)
{
f = LS_ARRAY_FACE(c.z, f);
return my_image_load_mip_cube_v4f16_i32(c.x, c.y, f, l, LOAD_TSHARP(i));
@@ -950,8 +951,8 @@ OCKL_MANGLE_T(image_sample,3D)(TSHARP i, SSHARP s, float4 c)
return my_image_sample_3d_v4f32_f32(c.x, c.y, c.z, LOAD_TSHARP(i), LOAD_SSHARP(s));
}
-RATTR float4
-OCKL_MANGLE_T(image_sample,CM)(TSHARP i, SSHARP s, float4 c)
+CRATTR float4
+OCKL_MANGLE_T(image_sample, CM)(TSHARP i, SSHARP s, float4 c)
{
CUBE_PREP(c);
if (EII())
@@ -960,8 +961,8 @@ OCKL_MANGLE_T(image_sample,CM)(TSHARP i, SSHARP s, float4 c)
return my_image_sample_cube_v4f32_f32(c.x, c.y, c.z, LOAD_TSHARP(i), LOAD_SSHARP(s));
}
-RATTR float4
-OCKL_MANGLE_T(image_sample,CMa)(TSHARP i, SSHARP s, float4 c)
+CRATTR float4
+OCKL_MANGLE_T(image_sample, CMa)(TSHARP i, SSHARP s, float4 c)
{
CUBE_PREP(c);
c.z = SAMPLE_ARRAY_FACE(c.w, c.z);
@@ -1068,15 +1069,15 @@ OCKL_MANGLE_T(image_sample_lod,3D)(TSHARP i, SSHARP s, float4 c, float l)
return my_image_sample_l_3d_v4f32_f32(c.x, c.y, c.z, l, LOAD_TSHARP(i), LOAD_SSHARP(s));
}
-RATTR float4
-OCKL_MANGLE_T(image_sample_lod,CM)(TSHARP i, SSHARP s, float4 c, float l)
+CRATTR float4
+OCKL_MANGLE_T(image_sample_lod, CM)(TSHARP i, SSHARP s, float4 c, float l)
{
CUBE_PREP(c);
return my_image_sample_l_cube_v4f32_f32(c.x, c.y, c.z, l, LOAD_TSHARP(i), LOAD_SSHARP(s));
}
-RATTR float4
-OCKL_MANGLE_T(image_sample_lod,CMa)(TSHARP i, SSHARP s, float4 c, float l)
+CRATTR float4
+OCKL_MANGLE_T(image_sample_lod, CMa)(TSHARP i, SSHARP s, float4 c, float l)
{
CUBE_PREP(c);
c.z = SAMPLE_ARRAY_FACE(c.w, c.z);
@@ -1135,8 +1136,8 @@ OCKL_MANGLE_T(image_sampleh,3D)(TSHARP i, SSHARP s, float4 c)
return my_image_sample_3d_v4f16_f32(c.x, c.y, c.z, LOAD_TSHARP(i), LOAD_SSHARP(s));
}
-RATTR half4
-OCKL_MANGLE_T(image_sampleh,CM)(TSHARP i, SSHARP s, float4 c)
+CRATTR half4
+OCKL_MANGLE_T(image_sampleh, CM)(TSHARP i, SSHARP s, float4 c)
{
CUBE_PREP(c);
if (EII())
@@ -1145,8 +1146,8 @@ OCKL_MANGLE_T(image_sampleh,CM)(TSHARP i, SSHARP s, float4 c)
return my_image_sample_cube_v4f16_f32(c.x, c.y, c.z, LOAD_TSHARP(i), LOAD_SSHARP(s));
}
-RATTR half4
-OCKL_MANGLE_T(image_sampleh,CMa)(TSHARP i, SSHARP s, float4 c)
+CRATTR half4
+OCKL_MANGLE_T(image_sampleh, CMa)(TSHARP i, SSHARP s, float4 c)
{
CUBE_PREP(c);
c.z = SAMPLE_ARRAY_FACE(c.w, c.z);
@@ -1225,15 +1226,15 @@ OCKL_MANGLE_T(image_sampleh_lod,3D)(TSHARP i, SSHARP s, float4 c, float l)
return my_image_sample_l_3d_v4f16_f32(c.x, c.y, c.z, l, LOAD_TSHARP(i), LOAD_SSHARP(s));
}
-RATTR half4
-OCKL_MANGLE_T(image_sampleh_lod,CM)(TSHARP i, SSHARP s, float4 c, float l)
+CRATTR half4
+OCKL_MANGLE_T(image_sampleh_lod, CM)(TSHARP i, SSHARP s, float4 c, float l)
{
CUBE_PREP(c);
return my_image_sample_l_cube_v4f16_f32(c.x, c.y, c.z, l, LOAD_TSHARP(i), LOAD_SSHARP(s));
}
-RATTR half4
-OCKL_MANGLE_T(image_sampleh_lod,CMa)(TSHARP i, SSHARP s, float4 c, float l)
+CRATTR half4
+OCKL_MANGLE_T(image_sampleh_lod, CMa)(TSHARP i, SSHARP s, float4 c, float l)
{
CUBE_PREP(c);
c.z = SAMPLE_ARRAY_FACE(c.w, c.z);
--- a/ockl/src/media.cl
+++ b/ockl/src/media.cl
@@ -11,6 +11,9 @@
#pragma OPENCL EXTENSION cl_khr_fp16 : enable
#define CATTR __attribute__((const))
+#define LCATTR __attribute__((const, target("lerp-inst")))
+#define QCATTR __attribute__((const, target("qsad-insts")))
+#define SCATTR __attribute__((const, target("sad-insts")))
#define AS_UCHAR4(X) __builtin_astype(X, uchar4)
CATTR uint
@@ -44,7 +47,7 @@ OCKL_MANGLE_U32(bytealign)(uint a, uint b, uint c)
return __builtin_amdgcn_alignbyte(a, b, c);
}
-CATTR uint
+LCATTR uint
OCKL_MANGLE_U32(lerp)(uint a, uint b, uint c)
{
return __builtin_amdgcn_lerp(a, b, c);
@@ -155,13 +158,13 @@ OCKL_MANGLE_U32(pack)(float4 a)
__builtin_amdgcn_cvt_pk_u8_f32(a.s0, 0, 0))));
}
-CATTR ulong
+QCATTR ulong
OCKL_MANGLE_U64(qsad)(ulong a, uint b, ulong c)
{
return __builtin_amdgcn_qsad_pk_u16_u8(a, b, c);
}
-CATTR uint
+SCATTR uint
OCKL_MANGLE_U32(sad)(uint a, uint b, uint c)
{
return __builtin_amdgcn_sad_u8(a, b, c);
@@ -174,13 +177,13 @@ OCKL_MANGLE_U32(sadd)(uint a, uint b, uint c)
return (a > b ? a : b) - (a < b ? a : b) + c;
}
-CATTR uint
+SCATTR uint
OCKL_MANGLE_U32(sadhi)(uint a, uint b, uint c)
{
return __builtin_amdgcn_sad_hi_u8(a, b, c);
}
-CATTR uint
+SCATTR uint
OCKL_MANGLE_U32(sadw)(uint a, uint b, uint c)
{
return __builtin_amdgcn_sad_u16(a, b, c);
@@ -213,4 +216,3 @@ OCKL_MANGLE_F32(unpack3)(uint a)
uchar4 v = AS_UCHAR4(a);
return (float)v.s3;
}
-
Reference in New Issue View Git Blame Copy Permalink