From c403a478076a16172d9b50e16c288b0d360f42ce Mon Sep 17 00:00:00 2001
From: Changpeng Fang
Date: Tue, 2 Apr 2024 13:15:21 -0700
Subject: [PATCH 0001/1683] AMDGPU: Update the description of cache policy for
buffer intrinsics, NFC (#87364)
Explicitly added gfx940, which has SC and NT bits. There must be better
ways of documenting the cache policy. But this is what I can do now.
Fixes: SWDEV-449810
---
llvm/include/llvm/IR/IntrinsicsAMDGPU.td | 518 ++++++++++-------------
1 file changed, 233 insertions(+), 285 deletions(-)
diff --git a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
index 3de20bb44e0c..6bbc13f1de86 100644
--- a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
+++ b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
@@ -851,16 +851,19 @@ class AMDGPUImageDimIntrinsic sdnodeprops> : Intrinsic<
P_.RetTypes, // vdata(VGPR) -- for load/atomic-with-return
!listconcat(
- !foreach(arg, P_.DataArgs, arg.Type), // vdata(VGPR) -- for store/atomic
- !if(P_.IsAtomic, [], [llvm_i32_ty]), // dmask(imm)
- P_.AddrTypes, // vaddr(VGPR)
- [llvm_v8i32_ty], // rsrc(SGPR)
- !if(P_.IsSample, [llvm_v4i32_ty, // samp(SGPR)
- llvm_i1_ty], []), // unorm(imm)
- [llvm_i32_ty, // texfailctrl(imm; bit 0 = tfe, bit 1 = lwe)
- llvm_i32_ty]), // cachepolicy(imm; bit 0 = glc, bit 1 = slc, bit 2 = dlc;
- // gfx12+ imm: bits [0-2] = th, bits [3-4] = scope)
-
+ !foreach(arg, P_.DataArgs, arg.Type), // vdata(VGPR) -- for store/atomic
+ !if(P_.IsAtomic, [], [llvm_i32_ty]), // dmask(imm)
+ P_.AddrTypes, // vaddr(VGPR)
+ [llvm_v8i32_ty], // rsrc(SGPR)
+ !if(P_.IsSample, [llvm_v4i32_ty, // samp(SGPR)
+ llvm_i1_ty], []), // unorm(imm)
+ [llvm_i32_ty, // texfailctrl(imm; bit 0 = tfe, bit 1 = lwe)
+ llvm_i32_ty]), // auxiliary/cachepolicy(imm):
+ // bit 0 = glc, bit 1 = slc,
+ // bit 2 = dlc (gfx10/gfx11),
+ // bit 4 = scc (gfx90a)
+ // gfx940: bit 0 = sc0, bit 1 = nt, bit 4 = sc1
+ // gfx12+: bits [0-2] = th, bits [3-4] = scope
!listconcat(props, [IntrNoCallback, IntrNoFree, IntrWillReturn],
!if(P_.IsAtomic, [], [ImmArg.DmaskArgIndex>>]),
!if(P_.IsSample, [ImmArg.UnormArgIndex>>], []),
@@ -1085,11 +1088,15 @@ def int_amdgcn_buffer_load : AMDGPUBufferLoad;
// the offset argument is uniform.
def int_amdgcn_s_buffer_load : DefaultAttrsIntrinsic <
[llvm_any_ty],
- [llvm_v4i32_ty, // rsrc(SGPR)
- llvm_i32_ty, // byte offset
- llvm_i32_ty], // cachepolicy(imm; bit 0 = glc, bit 1 = slc, bit 2 = dlc;
- // gfx12+ imm: bits [0-2] = th, bits [3-4] = scope)
- // Note: volatile bit is **not** permitted here.
+ [llvm_v4i32_ty, // rsrc(SGPR)
+ llvm_i32_ty, // byte offset
+ llvm_i32_ty], // auxiliary/cachepolicy(imm):
+ // bit 0 = glc, bit 1 = slc, bit 2 = dlc (gfx10/gfx11),
+ // bit 3 = swz, bit 4 = scc (gfx90a)
+ // gfx940: bit 0 = sc0, bit 1 = nt, bit 3 = swz, bit 4 = sc1
+ // gfx12+: bits [0-2] = th, bits [3-4] = scope,
+ // bit 6 = swz
+ // Note: volatile bit is **not** permitted here.
[IntrNoMem, ImmArg>]>,
AMDGPURsrcIntrinsic<0>;
@@ -1123,19 +1130,16 @@ def int_amdgcn_buffer_store : AMDGPUBufferStore;
// operation is volatile.
class AMDGPURawBufferLoad : DefaultAttrsIntrinsic <
[data_ty],
- [llvm_v4i32_ty, // rsrc(SGPR)
- llvm_i32_ty, // offset(VGPR/imm, included in bounds checking and swizzling)
- llvm_i32_ty, // soffset(SGPR/imm, excluded from bounds checking and swizzling)
- llvm_i32_ty], // auxiliary data (imm, cachepolicy (bit 0 = glc,
- // bit 1 = slc,
- // bit 2 = dlc on gfx10/gfx11),
- // swizzled buffer (bit 3 = swz),
- // gfx12+:
- // cachepolicy (bits [0-2] = th,
- // bits [3-4] = scope)
- // swizzled buffer (bit 6 = swz),
- // all:
- // volatile op (bit 31, stripped at lowering))
+ [llvm_v4i32_ty, // rsrc(SGPR)
+ llvm_i32_ty, // offset(VGPR/imm, included in bounds checking and swizzling)
+ llvm_i32_ty, // soffset(SGPR/imm, excluded from bounds checking and swizzling)
+ llvm_i32_ty], // auxiliary/cachepolicy(imm):
+ // bit 0 = glc, bit 1 = slc, bit 2 = dlc (gfx10/gfx11),
+ // bit 3 = swz, bit 4 = scc (gfx90a)
+ // gfx940: bit 0 = sc0, bit 1 = nt, bit 3 = swz, bit 4 = sc1
+ // gfx12+: bits [0-2] = th, bits [3-4] = scope,
+ // bit 6 = swz
+ // all: volatile op (bit 31, stripped at lowering)
[IntrReadMem, ImmArg>], "", [SDNPMemOperand]>,
AMDGPURsrcIntrinsic<0>;
def int_amdgcn_raw_buffer_load_format : AMDGPURawBufferLoad;
@@ -1143,20 +1147,16 @@ def int_amdgcn_raw_buffer_load : AMDGPURawBufferLoad;
class AMDGPURawPtrBufferLoad : DefaultAttrsIntrinsic <
[data_ty],
- [AMDGPUBufferRsrcTy, // rsrc(SGPR)
- llvm_i32_ty, // offset(VGPR/imm, included in bounds checking and swizzling)
- llvm_i32_ty, // soffset(SGPR/imm, excluded from bounds checking and swizzling)
- llvm_i32_ty], // auxiliary data (imm, cachepolicy (bit 0 = glc,
- // bit 1 = slc,
- // bit 2 = dlc on gfx10/gfx11),
- // swizzled buffer (bit 3 = swz),
- // gfx12+:
- // cachepolicy (bits [0-2] = th,
- // bits [3-4] = scope)
- // swizzled buffer (bit 6 = swz),
- // all:
- // volatile op (bit 31, stripped at lowering))
-
+ [AMDGPUBufferRsrcTy, // rsrc(SGPR)
+ llvm_i32_ty, // offset(VGPR/imm, included in bounds checking and swizzling)
+ llvm_i32_ty, // soffset(SGPR/imm, excluded from bounds checking and swizzling)
+ llvm_i32_ty], // auxiliary/cachepolicy(imm):
+ // bit 0 = glc, bit 1 = slc, bit 2 = dlc (gfx10/gfx11),
+ // bit 3 = swz, bit 4 = scc (gfx90a)
+ // gfx940: bit 0 = sc0, bit 1 = nt, bit 3 = swz, bit 4 = sc1
+ // gfx12+: bits [0-2] = th, bits [3-4] = scope,
+ // bit 6 = swz
+ // all: volatile op (bit 31, stripped at lowering)
[IntrArgMemOnly, IntrReadMem, ReadOnly>, NoCapture>,
ImmArg>], "", [SDNPMemOperand]>,
AMDGPURsrcIntrinsic<0>;
@@ -1165,20 +1165,17 @@ def int_amdgcn_raw_ptr_buffer_load : AMDGPURawPtrBufferLoad;
class AMDGPUStructBufferLoad : DefaultAttrsIntrinsic <
[data_ty],
- [llvm_v4i32_ty, // rsrc(SGPR)
- llvm_i32_ty, // vindex(VGPR)
- llvm_i32_ty, // offset(VGPR/imm, included in bounds checking and swizzling)
- llvm_i32_ty, // soffset(SGPR/imm, excluded from bounds checking and swizzling)
- llvm_i32_ty], // auxiliary data (imm, cachepolicy (bit 0 = glc,
- // bit 1 = slc,
- // bit 2 = dlc on gfx10/gfx11),
- // swizzled buffer (bit 3 = swz),
- // gfx12+:
- // cachepolicy (bits [0-2] = th,
- // bits [3-4] = scope)
- // swizzled buffer (bit 6 = swz),
- // all:
- // volatile op (bit 31, stripped at lowering))
+ [llvm_v4i32_ty, // rsrc(SGPR)
+ llvm_i32_ty, // vindex(VGPR)
+ llvm_i32_ty, // offset(VGPR/imm, included in bounds checking and swizzling)
+ llvm_i32_ty, // soffset(SGPR/imm, excluded from bounds checking and swizzling)
+ llvm_i32_ty], // auxiliary/cachepolicy(imm):
+ // bit 0 = glc, bit 1 = slc, bit 2 = dlc (gfx10/gfx11),
+ // bit 3 = swz, bit 4 = scc (gfx90a)
+ // gfx940: bit 0 = sc0, bit 1 = nt, bit 3 = swz, bit 4 = sc1
+ // gfx12+: bits [0-2] = th, bits [3-4] = scope,
+ // bit 6 = swz
+ // all: volatile op (bit 31, stripped at lowering)
[IntrReadMem, ImmArg>], "", [SDNPMemOperand]>,
AMDGPURsrcIntrinsic<0>;
def int_amdgcn_struct_buffer_load_format : AMDGPUStructBufferLoad;
@@ -1186,20 +1183,17 @@ def int_amdgcn_struct_buffer_load : AMDGPUStructBufferLoad;
class AMDGPUStructPtrBufferLoad : DefaultAttrsIntrinsic <
[data_ty],
- [AMDGPUBufferRsrcTy, // rsrc(SGPR)
- llvm_i32_ty, // vindex(VGPR)
- llvm_i32_ty, // offset(VGPR/imm, included in bounds checking and swizzling)
- llvm_i32_ty, // soffset(SGPR/imm, excluded from bounds checking and swizzling)
- llvm_i32_ty], // auxiliary data (imm, cachepolicy (bit 0 = glc,
- // bit 1 = slc,
- // bit 2 = dlc on gfx10/gfx11),
- // swizzled buffer (bit 3 = swz),
- // gfx12+:
- // cachepolicy (bits [0-2] = th,
- // bits [3-4] = scope)
- // swizzled buffer (bit 6 = swz),
- // all:
- // volatile op (bit 31, stripped at lowering))
+ [AMDGPUBufferRsrcTy, // rsrc(SGPR)
+ llvm_i32_ty, // vindex(VGPR)
+ llvm_i32_ty, // offset(VGPR/imm, included in bounds checking and swizzling)
+ llvm_i32_ty, // soffset(SGPR/imm, excluded from bounds checking and swizzling)
+ llvm_i32_ty], // auxiliary/cachepolicy(imm):
+ // bit 0 = glc, bit 1 = slc, bit 2 = dlc (gfx10/gfx11),
+ // bit 3 = swz, bit 4 = scc (gfx90a)
+ // gfx940: bit 0 = sc0, bit 1 = nt, bit 3 = swz, bit 4 = sc1
+ // gfx12+: bits [0-2] = th, bits [3-4] = scope,
+ // bit 6 = swz
+ // all: volatile op (bit 31, stripped at lowering)
[IntrArgMemOnly, IntrReadMem, ReadOnly>, NoCapture>,
ImmArg>], "", [SDNPMemOperand]>,
AMDGPURsrcIntrinsic<0>;
@@ -1208,20 +1202,17 @@ def int_amdgcn_struct_ptr_buffer_load : AMDGPUStructPtrBufferLoad;
class AMDGPURawBufferStore : DefaultAttrsIntrinsic <
[],
- [data_ty, // vdata(VGPR)
- llvm_v4i32_ty, // rsrc(SGPR)
- llvm_i32_ty, // offset(VGPR/imm, included in bounds checking and swizzling)
- llvm_i32_ty, // soffset(SGPR/imm, excluded from bounds checking and swizzling)
- llvm_i32_ty], // auxiliary data (imm, cachepolicy (bit 0 = glc,
- // bit 1 = slc,
- // bit 2 = dlc on gfx10/gfx11),
- // swizzled buffer (bit 3 = swz),
- // gfx12+:
- // cachepolicy (bits [0-2] = th,
- // bits [3-4] = scope)
- // swizzled buffer (bit 6 = swz),
- // all:
- // volatile op (bit 31, stripped at lowering))
+ [data_ty, // vdata(VGPR)
+ llvm_v4i32_ty, // rsrc(SGPR)
+ llvm_i32_ty, // offset(VGPR/imm, included in bounds checking and swizzling)
+ llvm_i32_ty, // soffset(SGPR/imm, excluded from bounds checking and swizzling)
+ llvm_i32_ty], // auxiliary/cachepolicy(imm):
+ // bit 0 = glc, bit 1 = slc, bit 2 = dlc (gfx10/gfx11),
+ // bit 3 = swz, bit 4 = scc (gfx90a)
+ // gfx940: bit 0 = sc0, bit 1 = nt, bit 3 = swz, bit 4 = sc1
+ // gfx12+: bits [0-2] = th, bits [3-4] = scope,
+ // bit 6 = swz
+ // all: volatile op (bit 31, stripped at lowering)
[IntrWriteMem, ImmArg>], "", [SDNPMemOperand]>,
AMDGPURsrcIntrinsic<1>;
def int_amdgcn_raw_buffer_store_format : AMDGPURawBufferStore;
@@ -1229,20 +1220,17 @@ def int_amdgcn_raw_buffer_store : AMDGPURawBufferStore;
class AMDGPURawPtrBufferStore : DefaultAttrsIntrinsic <
[],
- [data_ty, // vdata(VGPR)
- AMDGPUBufferRsrcTy, // rsrc(SGPR)
- llvm_i32_ty, // offset(VGPR/imm, included in bounds checking and swizzling)
- llvm_i32_ty, // soffset(SGPR/imm, excluded from bounds checking and swizzling)
- llvm_i32_ty], // auxiliary data (imm, cachepolicy (bit 0 = glc,
- // bit 1 = slc,
- // bit 2 = dlc on gfx10/gfx11),
- // swizzled buffer (bit 3 = swz),
- // gfx12+:
- // cachepolicy (bits [0-2] = th,
- // bits [3-4] = scope)
- // swizzled buffer (bit 6 = swz),
- // all:
- // volatile op (bit 31, stripped at lowering))
+ [data_ty, // vdata(VGPR)
+ AMDGPUBufferRsrcTy, // rsrc(SGPR)
+ llvm_i32_ty, // offset(VGPR/imm, included in bounds checking and swizzling)
+ llvm_i32_ty, // soffset(SGPR/imm, excluded from bounds checking and swizzling)
+ llvm_i32_ty], // auxiliary/cachepolicy(imm):
+ // bit 0 = glc, bit 1 = slc, bit 2 = dlc (gfx10/gfx11),
+ // bit 3 = swz, bit 4 = scc (gfx90a)
+ // gfx940: bit 0 = sc0, bit 1 = nt, bit 3 = swz, bit 4 = sc1
+ // gfx12+: bits [0-2] = th, bits [3-4] = scope,
+ // bit 6 = swz
+ // all: volatile op (bit 31, stripped at lowering)
[IntrArgMemOnly, IntrWriteMem, WriteOnly>, NoCapture>,
ImmArg>], "", [SDNPMemOperand]>,
AMDGPURsrcIntrinsic<1>;
@@ -1251,21 +1239,18 @@ def int_amdgcn_raw_ptr_buffer_store : AMDGPURawPtrBufferStore;
class AMDGPUStructBufferStore : DefaultAttrsIntrinsic <
[],
- [data_ty, // vdata(VGPR)
- llvm_v4i32_ty, // rsrc(SGPR)
- llvm_i32_ty, // vindex(VGPR)
- llvm_i32_ty, // offset(VGPR/imm, included in bounds checking and swizzling)
- llvm_i32_ty, // soffset(SGPR/imm, excluded from bounds checking and swizzling)
- llvm_i32_ty], // auxiliary data (imm, cachepolicy (bit 0 = glc,
- // bit 1 = slc,
- // bit 2 = dlc on gfx10/gfx11),
- // swizzled buffer (bit 3 = swz),
- // gfx12+:
- // cachepolicy (bits [0-2] = th,
- // bits [3-4] = scope)
- // swizzled buffer (bit 6 = swz),
- // all:
- // volatile op (bit 31, stripped at lowering))
+ [data_ty, // vdata(VGPR)
+ llvm_v4i32_ty, // rsrc(SGPR)
+ llvm_i32_ty, // vindex(VGPR)
+ llvm_i32_ty, // offset(VGPR/imm, included in bounds checking and swizzling)
+ llvm_i32_ty, // soffset(SGPR/imm, excluded from bounds checking and swizzling)
+ llvm_i32_ty], // auxiliary/cachepolicy(imm):
+ // bit 0 = glc, bit 1 = slc, bit 2 = dlc (gfx10/gfx11),
+ // bit 3 = swz, bit 4 = scc (gfx90a)
+ // gfx940: bit 0 = sc0, bit 1 = nt, bit 3 = swz, bit 4 = sc1
+ // gfx12+: bits [0-2] = th, bits [3-4] = scope,
+ // bit 6 = swz
+ // all: volatile op (bit 31, stripped at lowering)
[IntrWriteMem, ImmArg>], "", [SDNPMemOperand]>,
AMDGPURsrcIntrinsic<1>;
def int_amdgcn_struct_buffer_store_format : AMDGPUStructBufferStore;
@@ -1273,21 +1258,18 @@ def int_amdgcn_struct_buffer_store : AMDGPUStructBufferStore;
class AMDGPUStructPtrBufferStore : DefaultAttrsIntrinsic <
[],
- [data_ty, // vdata(VGPR)
- AMDGPUBufferRsrcTy, // rsrc(SGPR)
- llvm_i32_ty, // vindex(VGPR)
- llvm_i32_ty, // offset(VGPR/imm, included in bounds checking and swizzling)
- llvm_i32_ty, // soffset(SGPR/imm, excluded from bounds checking and swizzling)
- llvm_i32_ty], // auxiliary data (imm, cachepolicy (bit 0 = glc,
- // bit 1 = slc,
- // bit 2 = dlc on gfx10/gfx11),
- // swizzled buffer (bit 3 = swz),
- // gfx12+:
- // cachepolicy (bits [0-2] = th,
- // bits [3-4] = scope)
- // swizzled buffer (bit 6 = swz),
- // all:
- // volatile op (bit 31, stripped at lowering))
+ [data_ty, // vdata(VGPR)
+ AMDGPUBufferRsrcTy, // rsrc(SGPR)
+ llvm_i32_ty, // vindex(VGPR)
+ llvm_i32_ty, // offset(VGPR/imm, included in bounds checking and swizzling)
+ llvm_i32_ty, // soffset(SGPR/imm, excluded from bounds checking and swizzling)
+ llvm_i32_ty], // auxiliary/cachepolicy(imm):
+ // bit 0 = glc, bit 1 = slc, bit 2 = dlc (gfx10/gfx11),
+ // bit 3 = swz, bit 4 = scc (gfx90a)
+ // gfx940: bit 0 = sc0, bit 1 = nt, bit 3 = swz, bit 4 = sc1
+ // gfx12+: bits [0-2] = th, bits [3-4] = scope,
+ // bit 6 = swz
+ // all: volatile op (bit 31, stripped at lowering)
[IntrArgMemOnly, IntrWriteMem, WriteOnly>, NoCapture>,
ImmArg>], "", [SDNPMemOperand]>,
AMDGPURsrcIntrinsic<1>;
@@ -1540,33 +1522,29 @@ def int_amdgcn_raw_tbuffer_load : DefaultAttrsIntrinsic <
llvm_i32_ty, // offset(VGPR/imm, included in bounds checking and swizzling)
llvm_i32_ty, // soffset(SGPR/imm, excluded from bounds checking and swizzling)
llvm_i32_ty, // format(imm; bits 3..0 = dfmt, bits 6..4 = nfmt)
- llvm_i32_ty], // auxiliary data (imm, cachepolicy (bit 0 = glc,
- // bit 1 = slc,
- // bit 2 = dlc on gfx10/gfx11),
- // swizzled buffer (bit 3 = swz))
- // gfx12+:
- // cachepolicy (bits [0-2] = th,
- // bits [3-4] = scope)
- // swizzled buffer (bit 6 = swz)
+ llvm_i32_ty], // auxiliary/cachepolicy(imm):
+ // bit 0 = glc, bit 1 = slc, bit 2 = dlc (gfx10/gfx11),
+ // bit 3 = swz, bit 4 = scc (gfx90a)
+ // gfx940: bit 0 = sc0, bit 1 = nt, bit 3 = swz, bit 4 = sc1
+ // gfx12+: bits [0-2] = th, bits [3-4] = scope,
+ // bit 6 = swz
[IntrReadMem,
ImmArg>, ImmArg>], "", [SDNPMemOperand]>,
AMDGPURsrcIntrinsic<0>;
def int_amdgcn_raw_ptr_tbuffer_load : DefaultAttrsIntrinsic <
- [llvm_any_ty], // overloaded for types f32/i32, v2f32/v2i32, v4f32/v4i32
+ [llvm_any_ty], // overloaded for types f32/i32, v2f32/v2i32, v4f32/v4i32
[AMDGPUBufferRsrcTy, // rsrc(SGPR)
- llvm_i32_ty, // offset(VGPR/imm, included in bounds` checking and swizzling)
- llvm_i32_ty, // soffset(SGPR/imm, excluded from bounds checking and swizzling)
- llvm_i32_ty, // format(imm; bits 3..0 = dfmt, bits 6..4 = nfmt)
- llvm_i32_ty], // auxiliary data (imm, cachepolicy (bit 0 = glc,
- // bit 1 = slc,
- // bit 2 = dlc on gfx10/gfx11),
- // swizzled buffer (bit 3 = swz),
- // gfx12+:
- // cachepolicy (bits [0-2] = th,
- // bits [3-4] = scope)
- // swizzled buffer (bit 6 = swz)
- // volatile op (bit 31, stripped at lowering))
+ llvm_i32_ty, // offset(VGPR/imm, included in bounds` checking and swizzling)
+ llvm_i32_ty, // soffset(SGPR/imm, excluded from bounds checking and swizzling)
+ llvm_i32_ty, // format(imm; bits 3..0 = dfmt, bits 6..4 = nfmt)
+ llvm_i32_ty], // auxiliary/cachepolicy(imm):
+ // bit 0 = glc, bit 1 = slc, bit 2 = dlc (gfx10/gfx11),
+ // bit 3 = swz, bit 4 = scc (gfx90a)
+ // gfx940: bit 0 = sc0, bit 1 = nt, bit 3 = swz, bit 4 = sc1
+ // gfx12+: bits [0-2] = th, bits [3-4] = scope,
+ // bit 6 = swz
+ // all: volatile op (bit 31, stripped at lowering)
[IntrArgMemOnly, IntrReadMem, ReadOnly>, NoCapture>,
ImmArg>, ImmArg>], "", [SDNPMemOperand]>,
AMDGPURsrcIntrinsic<0>;
@@ -1578,16 +1556,13 @@ def int_amdgcn_raw_tbuffer_store : DefaultAttrsIntrinsic <
llvm_i32_ty, // offset(VGPR/imm, included in bounds checking and swizzling)
llvm_i32_ty, // soffset(SGPR/imm, excluded from bounds checking and swizzling)
llvm_i32_ty, // format(imm; bits 3..0 = dfmt, bits 6..4 = nfmt)
- llvm_i32_ty], // auxiliary data (imm, cachepolicy (bit 0 = glc,
- // bit 1 = slc,
- // bit 2 = dlc on gfx10/gfx11),
- // swizzled buffer (bit 3 = swz),
- // gfx12+:
- // cachepolicy (bits [0-2] = th,
- // bits [3-4] = scope)
- // swizzled buffer (bit 6 = swz),
- // all:
- // volatile op (bit 31, stripped at lowering))
+ llvm_i32_ty], // auxiliary/cachepolicy(imm):
+ // bit 0 = glc, bit 1 = slc, bit 2 = dlc (gfx10/gfx11),
+ // bit 3 = swz, bit 4 = scc (gfx90a)
+ // gfx940: bit 0 = sc0, bit 1 = nt, bit 3 = swz, bit 4 = sc1
+ // gfx12+: bits [0-2] = th, bits [3-4] = scope,
+ // bit 6 = swz
+ // all: volatile op (bit 31, stripped at lowering)
[IntrWriteMem,
ImmArg>, ImmArg>], "", [SDNPMemOperand]>,
AMDGPURsrcIntrinsic<1>;
@@ -1599,16 +1574,13 @@ def int_amdgcn_raw_ptr_tbuffer_store : DefaultAttrsIntrinsic <
llvm_i32_ty, // offset(VGPR/imm, included in bounds checking and swizzling)
llvm_i32_ty, // soffset(SGPR/imm, excluded from bounds checking and swizzling)
llvm_i32_ty, // format(imm; bits 3..0 = dfmt, bits 6..4 = nfmt)
- llvm_i32_ty], // auxiliary data (imm, cachepolicy (bit 0 = glc,
- // bit 1 = slc,
- // bit 2 = dlc on gfx10/gfx11),
- // swizzled buffer (bit 3 = swz),
- // gfx12+:
- // cachepolicy (bits [0-2] = th,
- // bits [3-4] = scope)
- // swizzled buffer (bit 6 = swz),
- // all:
- // volatile op (bit 31, stripped at lowering))
+ llvm_i32_ty], // auxiliary/cachepolicy(imm):
+ // bit 0 = glc, bit 1 = slc, bit 2 = dlc (gfx10/gfx11),
+ // bit 3 = swz, bit 4 = scc (gfx90a)
+ // gfx940: bit 0 = sc0, bit 1 = nt, bit 3 = swz, bit 4 = sc1
+ // gfx12+: bits [0-2] = th, bits [3-4] = scope,
+ // bit 6 = swz
+ // all: volatile op (bit 31, stripped at lowering)
[IntrArgMemOnly, IntrWriteMem, WriteOnly>, NoCapture>,
ImmArg>, ImmArg>], "", [SDNPMemOperand]>,
AMDGPURsrcIntrinsic<1>;
@@ -1620,59 +1592,50 @@ def int_amdgcn_struct_tbuffer_load : DefaultAttrsIntrinsic <
llvm_i32_ty, // offset(VGPR/imm, included in bounds checking and swizzling)
llvm_i32_ty, // soffset(SGPR/imm, excluded from bounds checking and swizzling)
llvm_i32_ty, // format(imm; bits 3..0 = dfmt, bits 6..4 = nfmt)
- llvm_i32_ty], // auxiliary data (imm, cachepolicy (bit 0 = glc,
- // bit 1 = slc,
- // bit 2 = dlc on gfx10/gfx11),
- // swizzled buffer (bit 3 = swz),
- // gfx12+:
- // cachepolicy (bits [0-2] = th,
- // bits [3-4] = scope)
- // swizzled buffer (bit 6 = swz),
- // all:
- // volatile op (bit 31, stripped at lowering))
+ llvm_i32_ty], // auxiliary/cachepolicy(imm):
+ // bit 0 = glc, bit 1 = slc, bit 2 = dlc (gfx10/gfx11),
+ // bit 3 = swz, bit 4 = scc (gfx90a)
+ // gfx940: bit 0 = sc0, bit 1 = nt, bit 3 = swz, bit 4 = sc1
+ // gfx12+: bits [0-2] = th, bits [3-4] = scope,
+ // bit 6 = swz
+ // all: volatile op (bit 31, stripped at lowering)
[IntrReadMem,
ImmArg>, ImmArg>], "", [SDNPMemOperand]>,
AMDGPURsrcIntrinsic<0>;
def int_amdgcn_struct_ptr_tbuffer_load : DefaultAttrsIntrinsic <
- [llvm_any_ty], // overloaded for types f32/i32, v2f32/v2i32, v4f32/v4i32
+ [llvm_any_ty], // overloaded for types f32/i32, v2f32/v2i32, v4f32/v4i32
[AMDGPUBufferRsrcTy, // rsrc(SGPR)
- llvm_i32_ty, // vindex(VGPR)
- llvm_i32_ty, // offset(VGPR/imm, included in bounds checking and swizzling)
- llvm_i32_ty, // soffset(SGPR/imm, excluded from bounds checking and swizzling)
- llvm_i32_ty, // format(imm; bits 3..0 = dfmt, bits 6..4 = nfmt)
- llvm_i32_ty], // auxiliary data (imm, cachepolicy (bit 0 = glc,
- // bit 1 = slc,
- // bit 2 = dlc on gfx10/gfx11),
- // swizzled buffer (bit 3 = swz),
- // gfx12+:
- // cachepolicy (bits [0-2] = th,
- // bits [3-4] = scope)
- // swizzled buffer (bit 6 = swz),
- // all:
- // volatile op (bit 31, stripped at lowering))
+ llvm_i32_ty, // vindex(VGPR)
+ llvm_i32_ty, // offset(VGPR/imm, included in bounds checking and swizzling)
+ llvm_i32_ty, // soffset(SGPR/imm, excluded from bounds checking and swizzling)
+ llvm_i32_ty, // format(imm; bits 3..0 = dfmt, bits 6..4 = nfmt)
+ llvm_i32_ty], // auxiliary/cachepolicy(imm):
+ // bit 0 = glc, bit 1 = slc, bit 2 = dlc (gfx10/gfx11),
+ // bit 3 = swz, bit 4 = scc (gfx90a)
+ // gfx940: bit 0 = sc0, bit 1 = nt, bit 3 = swz, bit 4 = sc1
+ // gfx12+: bits [0-2] = th, bits [3-4] = scope,
+ // bit 6 = swz
+ // all: volatile op (bit 31, stripped at lowering)
[IntrArgMemOnly, IntrReadMem, ReadOnly>, NoCapture>,
ImmArg>, ImmArg>], "", [SDNPMemOperand]>,
AMDGPURsrcIntrinsic<0>;
def int_amdgcn_struct_ptr_tbuffer_store : DefaultAttrsIntrinsic <
[],
- [llvm_any_ty, // vdata(VGPR), overloaded for types f32/i32, v2f32/v2i32, v4f32/v4i32
+ [llvm_any_ty, // vdata(VGPR), overloaded for types f32/i32, v2f32/v2i32, v4f32/v4i32
AMDGPUBufferRsrcTy, // rsrc(SGPR)
- llvm_i32_ty, // vindex(VGPR)
- llvm_i32_ty, // offset(VGPR/imm, included in bounds checking and swizzling)
- llvm_i32_ty, // soffset(SGPR/imm, excluded from bounds checking and swizzling)
- llvm_i32_ty, // format(imm; bits 3..0 = dfmt, bits 6..4 = nfmt)
- llvm_i32_ty], // auxiliary data (imm, cachepolicy (bit 0 = glc,
- // bit 1 = slc,
- // bit 2 = dlc on gfx10/gfx11),
- // swizzled buffer (bit 3 = swz),
- // gfx12+:
- // cachepolicy (bits [0-2] = th,
- // bits [3-4] = scope)
- // swizzled buffer (bit 6 = swz),
- // all:
- // volatile op (bit 31, stripped at lowering))
+ llvm_i32_ty, // vindex(VGPR)
+ llvm_i32_ty, // offset(VGPR/imm, included in bounds checking and swizzling)
+ llvm_i32_ty, // soffset(SGPR/imm, excluded from bounds checking and swizzling)
+ llvm_i32_ty, // format(imm; bits 3..0 = dfmt, bits 6..4 = nfmt)
+ llvm_i32_ty], // auxiliary/cachepolicy(imm):
+ // bit 0 = glc, bit 1 = slc, bit 2 = dlc (gfx10/gfx11),
+ // bit 3 = swz, bit 4 = scc (gfx90a)
+ // gfx940: bit 0 = sc0, bit 1 = nt, bit 3 = swz, bit 4 = sc1
+ // gfx12+: bits [0-2] = th, bits [3-4] = scope,
+ // bit 6 = swz
+ // all: volatile op (bit 31, stripped at lowering)
[IntrArgMemOnly, IntrWriteMem, WriteOnly>, NoCapture>,
ImmArg>, ImmArg>], "", [SDNPMemOperand]>,
AMDGPURsrcIntrinsic<1>;
@@ -1685,16 +1648,13 @@ def int_amdgcn_struct_tbuffer_store : DefaultAttrsIntrinsic <
llvm_i32_ty, // offset(VGPR/imm, included in bounds checking and swizzling)
llvm_i32_ty, // soffset(SGPR/imm, excluded from bounds checking and swizzling)
llvm_i32_ty, // format(imm; bits 3..0 = dfmt, bits 6..4 = nfmt)
- llvm_i32_ty], // auxiliary data (imm, cachepolicy (bit 0 = glc,
- // bit 1 = slc,
- // bit 2 = dlc on gfx10/gfx11),
- // swizzled buffer (bit 3 = swz),
- // gfx12+:
- // cachepolicy (bits [0-2] = th,
- // bits [3-4] = scope)
- // swizzled buffer (bit 6 = swz),
- // all:
- // volatile op (bit 31, stripped at lowering))
+ llvm_i32_ty], // auxiliary/cachepolicy(imm):
+ // bit 0 = glc, bit 1 = slc, bit 2 = dlc (gfx10/gfx11),
+ // bit 3 = swz, bit 4 = scc (gfx90a)
+ // gfx940: bit 0 = sc0, bit 1 = nt, bit 3 = swz, bit 4 = sc1
+ // gfx12+: bits [0-2] = th, bits [3-4] = scope,
+ // bit 6 = swz
+ // all: volatile op (bit 31, stripped at lowering)
[IntrWriteMem,
ImmArg>, ImmArg>], "", [SDNPMemOperand]>,
AMDGPURsrcIntrinsic<1>;
@@ -1746,44 +1706,38 @@ def int_amdgcn_buffer_atomic_fadd : AMDGPUBufferAtomicFP;
class AMDGPURawBufferLoadLDS : Intrinsic <
[],
- [llvm_v4i32_ty, // rsrc(SGPR)
- LLVMQualPointerType<3>, // LDS base offset
- llvm_i32_ty, // Data byte size: 1/2/4
- llvm_i32_ty, // voffset(VGPR, included in bounds checking and swizzling)
- llvm_i32_ty, // soffset(SGPR/imm, excluded from bounds checking and swizzling)
- llvm_i32_ty, // imm offset(imm, included in bounds checking and swizzling)
- llvm_i32_ty], // auxiliary data (imm, cachepolicy (bit 0 = glc,
- // bit 1 = slc,
- // bit 2 = dlc on gfx10/gfx11))
- // swizzled buffer (bit 3 = swz),
- // gfx12+:
- // cachepolicy (bits [0-2] = th,
- // bits [3-4] = scope)
- // swizzled buffer (bit 6 = swz),
- // all:
- // volatile op (bit 31, stripped at lowering))
+ [llvm_v4i32_ty, // rsrc(SGPR)
+ LLVMQualPointerType<3>, // LDS base offset
+ llvm_i32_ty, // Data byte size: 1/2/4
+ llvm_i32_ty, // voffset(VGPR, included in bounds checking and swizzling)
+ llvm_i32_ty, // soffset(SGPR/imm, excluded from bounds checking and swizzling)
+ llvm_i32_ty, // imm offset(imm, included in bounds checking and swizzling)
+ llvm_i32_ty], // auxiliary/cachepolicy(imm):
+ // bit 0 = glc, bit 1 = slc, bit 2 = dlc (gfx10/gfx11),
+ // bit 3 = swz, bit 4 = scc (gfx90a)
+ // gfx940: bit 0 = sc0, bit 1 = nt, bit 3 = swz, bit 4 = sc1
+ // gfx12+: bits [0-2] = th, bits [3-4] = scope,
+ // bit 6 = swz
+ // all: volatile op (bit 31, stripped at lowering)
[IntrWillReturn, NoCapture>, ImmArg>, ImmArg>,
ImmArg>, IntrNoCallback, IntrNoFree], "", [SDNPMemOperand]>, AMDGPURsrcIntrinsic<0>;
def int_amdgcn_raw_buffer_load_lds : AMDGPURawBufferLoadLDS;
class AMDGPURawPtrBufferLoadLDS : Intrinsic <
[],
- [AMDGPUBufferRsrcTy, // rsrc(SGPR)
- LLVMQualPointerType<3>, // LDS base offset
- llvm_i32_ty, // Data byte size: 1/2/4
- llvm_i32_ty, // voffset(VGPR, included in bounds checking and swizzling)
- llvm_i32_ty, // soffset(SGPR/imm, excluded from bounds checking and swizzling)
- llvm_i32_ty, // imm offset(imm, included in bounds checking and swizzling)
- llvm_i32_ty], // auxiliary data (imm, cachepolicy (bit 0 = glc,
- // bit 1 = slc,
- // bit 2 = dlc on gfx10/gfx11))
- // swizzled buffer (bit 3 = swz),
- // gfx12+:
- // cachepolicy (bits [0-2] = th,
- // bits [3-4] = scope)
- // swizzled buffer (bit 6 = swz),
- // all:
- // volatile op (bit 31, stripped at lowering))
+ [AMDGPUBufferRsrcTy, // rsrc(SGPR)
+ LLVMQualPointerType<3>, // LDS base offset
+ llvm_i32_ty, // Data byte size: 1/2/4
+ llvm_i32_ty, // voffset(VGPR, included in bounds checking and swizzling)
+ llvm_i32_ty, // soffset(SGPR/imm, excluded from bounds checking and swizzling)
+ llvm_i32_ty, // imm offset(imm, included in bounds checking and swizzling)
+ llvm_i32_ty], // auxiliary/cachepolicy(imm):
+ // bit 0 = glc, bit 1 = slc, bit 2 = dlc (gfx10/gfx11),
+ // bit 3 = swz, bit 4 = scc (gfx90a)
+ // gfx940: bit 0 = sc0, bit 1 = nt, bit 3 = swz, bit 4 = sc1
+ // gfx12+: bits [0-2] = th, bits [3-4] = scope,
+ // bit 6 = swz
+ // all: volatile op (bit 31, stripped at lowering)
[IntrWillReturn, IntrArgMemOnly,
ReadOnly>, NoCapture>,
WriteOnly>, NoCapture>,
@@ -1793,46 +1747,40 @@ def int_amdgcn_raw_ptr_buffer_load_lds : AMDGPURawPtrBufferLoadLDS;
class AMDGPUStructBufferLoadLDS : Intrinsic <
[],
- [llvm_v4i32_ty, // rsrc(SGPR)
- LLVMQualPointerType<3>, // LDS base offset
- llvm_i32_ty, // Data byte size: 1/2/4
- llvm_i32_ty, // vindex(VGPR)
- llvm_i32_ty, // voffset(VGPR, included in bounds checking and swizzling)
- llvm_i32_ty, // soffset(SGPR/imm, excluded from bounds checking and swizzling)
- llvm_i32_ty, // imm offset(imm, included in bounds checking and swizzling)
- llvm_i32_ty], // auxiliary data (imm, cachepolicy (bit 0 = glc,
- // bit 1 = slc,
- // bit 2 = dlc on gfx10/gfx11))
- // swizzled buffer (bit 3 = swz),
- // gfx12+:
- // cachepolicy (bits [0-2] = th,
- // bits [3-4] = scope)
- // swizzled buffer (bit 6 = swz),
- // all:
- // volatile op (bit 31, stripped at lowering))
+ [llvm_v4i32_ty, // rsrc(SGPR)
+ LLVMQualPointerType<3>, // LDS base offset
+ llvm_i32_ty, // Data byte size: 1/2/4
+ llvm_i32_ty, // vindex(VGPR)
+ llvm_i32_ty, // voffset(VGPR, included in bounds checking and swizzling)
+ llvm_i32_ty, // soffset(SGPR/imm, excluded from bounds checking and swizzling)
+ llvm_i32_ty, // imm offset(imm, included in bounds checking and swizzling)
+ llvm_i32_ty], // auxiliary/cachepolicy(imm):
+ // bit 0 = glc, bit 1 = slc, bit 2 = dlc (gfx10/gfx11),
+ // bit 3 = swz, bit 4 = scc (gfx90a)
+ // gfx940: bit 0 = sc0, bit 1 = nt, bit 3 = swz, bit 4 = sc1
+ // gfx12+: bits [0-2] = th, bits [3-4] = scope,
+ // bit 6 = swz
+ // all: volatile op (bit 31, stripped at lowering)
[IntrWillReturn, NoCapture>, ImmArg>, ImmArg>,
ImmArg>, IntrNoCallback, IntrNoFree], "", [SDNPMemOperand]>, AMDGPURsrcIntrinsic<0>;
def int_amdgcn_struct_buffer_load_lds : AMDGPUStructBufferLoadLDS;
class AMDGPUStructPtrBufferLoadLDS : Intrinsic <
[],
- [AMDGPUBufferRsrcTy, // rsrc(SGPR)
- LLVMQualPointerType<3> , // LDS base offset
- llvm_i32_ty, // Data byte size: 1/2/4
- llvm_i32_ty, // vindex(VGPR)
- llvm_i32_ty, // voffset(VGPR, included in bounds checking and swizzling)
- llvm_i32_ty, // soffset(SGPR/imm, excluded from bounds checking and swizzling)
- llvm_i32_ty, // imm offset(imm, included in bounds checking and swizzling)
- llvm_i32_ty], // auxiliary data (imm, cachepolicy (bit 0 = glc,
- // bit 1 = slc,
- // bit 2 = dlc on gfx10/gfx11))
- // swizzled buffer (bit 3 = swz),
- // gfx12+:
- // cachepolicy (bits [0-2] = th,
- // bits [3-4] = scope)
- // swizzled buffer (bit 6 = swz),
- // all:
- // volatile op (bit 31, stripped at lowering))
+ [AMDGPUBufferRsrcTy, // rsrc(SGPR)
+ LLVMQualPointerType<3>, // LDS base offset
+ llvm_i32_ty, // Data byte size: 1/2/4
+ llvm_i32_ty, // vindex(VGPR)
+ llvm_i32_ty, // voffset(VGPR, included in bounds checking and swizzling)
+ llvm_i32_ty, // soffset(SGPR/imm, excluded from bounds checking and swizzling)
+ llvm_i32_ty, // imm offset(imm, included in bounds checking and swizzling)
+ llvm_i32_ty], // auxiliary/cachepolicy(imm):
+ // bit 0 = glc, bit 1 = slc, bit 2 = dlc (gfx10/gfx11),
+ // bit 3 = swz, bit 4 = scc (gfx90a)
+ // gfx940: bit 0 = sc0, bit 1 = nt, bit 3 = swz, bit 4 = sc1
+ // gfx12+: bits [0-2] = th, bits [3-4] = scope,
+ // bit 6 = swz
+ // all: volatile op (bit 31, stripped at lowering)
[IntrWillReturn, IntrArgMemOnly,
ReadOnly>, NoCapture>,
WriteOnly>, NoCapture>,
--
GitLab
From 6335de4a23e4b71020861b872f51bafb79fbe419 Mon Sep 17 00:00:00 2001
From: Joseph Huber
Date: Tue, 2 Apr 2024 15:23:11 -0500
Subject: [PATCH 0002/1683] [libc] Disable '_exit' on the GPU build
Summary:
There are other dependencies to enable `unistd.h` on the GPU which
prevented the header from being generated. This is a POSIX extension and
isn't part of the core `libc`, so we can just disable this for now to
get the bots gree.
---
libc/config/gpu/entrypoints.txt | 3 ---
libc/config/gpu/headers.txt | 1 -
2 files changed, 4 deletions(-)
diff --git a/libc/config/gpu/entrypoints.txt b/libc/config/gpu/entrypoints.txt
index 827b5937e636..4fb87cb9f5a3 100644
--- a/libc/config/gpu/entrypoints.txt
+++ b/libc/config/gpu/entrypoints.txt
@@ -211,9 +211,6 @@ set(TARGET_LIBC_ENTRYPOINTS
# gpu/rpc.h entrypoints
libc.src.gpu.rpc_host_call
-
- # unistd.h entrypoints
- libc.src.unistd._exit
)
set(TARGET_LIBM_ENTRYPOINTS
diff --git a/libc/config/gpu/headers.txt b/libc/config/gpu/headers.txt
index 37c063a7ef6f..dd16938da8a4 100644
--- a/libc/config/gpu/headers.txt
+++ b/libc/config/gpu/headers.txt
@@ -12,7 +12,6 @@ set(TARGET_PUBLIC_HEADERS
libc.include.errno
libc.include.stdlib
libc.include.stdio
- libc.include.unistd
# Header for RPC extensions
libc.include.gpu_rpc
--
GitLab
From 7f2a41b643af1348631dee27dc1321046370a758 Mon Sep 17 00:00:00 2001
From: Matt Arsenault
Date: Mon, 1 Apr 2024 16:35:28 +0300
Subject: [PATCH 0003/1683] MachineScheduler: Simplify usage of TargetInstrInfo
---
llvm/lib/CodeGen/MachineScheduler.cpp | 16 ++++------------
1 file changed, 4 insertions(+), 12 deletions(-)
diff --git a/llvm/lib/CodeGen/MachineScheduler.cpp b/llvm/lib/CodeGen/MachineScheduler.cpp
index cbeb02a72552..203bdffc49ae 100644
--- a/llvm/lib/CodeGen/MachineScheduler.cpp
+++ b/llvm/lib/CodeGen/MachineScheduler.cpp
@@ -3263,14 +3263,10 @@ void GenericScheduler::initialize(ScheduleDAGMI *dag) {
// are disabled, then these HazardRecs will be disabled.
const InstrItineraryData *Itin = SchedModel->getInstrItineraries();
if (!Top.HazardRec) {
- Top.HazardRec =
- DAG->MF.getSubtarget().getInstrInfo()->CreateTargetMIHazardRecognizer(
- Itin, DAG);
+ Top.HazardRec = DAG->TII->CreateTargetMIHazardRecognizer(Itin, DAG);
}
if (!Bot.HazardRec) {
- Bot.HazardRec =
- DAG->MF.getSubtarget().getInstrInfo()->CreateTargetMIHazardRecognizer(
- Itin, DAG);
+ Bot.HazardRec = DAG->TII->CreateTargetMIHazardRecognizer(Itin, DAG);
}
TopCand.SU = nullptr;
BotCand.SU = nullptr;
@@ -3877,14 +3873,10 @@ void PostGenericScheduler::initialize(ScheduleDAGMI *Dag) {
// or are disabled, then these HazardRecs will be disabled.
const InstrItineraryData *Itin = SchedModel->getInstrItineraries();
if (!Top.HazardRec) {
- Top.HazardRec =
- DAG->MF.getSubtarget().getInstrInfo()->CreateTargetMIHazardRecognizer(
- Itin, DAG);
+ Top.HazardRec = DAG->TII->CreateTargetMIHazardRecognizer(Itin, DAG);
}
if (!Bot.HazardRec) {
- Bot.HazardRec =
- DAG->MF.getSubtarget().getInstrInfo()->CreateTargetMIHazardRecognizer(
- Itin, DAG);
+ Bot.HazardRec = DAG->TII->CreateTargetMIHazardRecognizer(Itin, DAG);
}
}
--
GitLab
From 6b2a4668b99ef47f567709bcf6f57bc60c35eea9 Mon Sep 17 00:00:00 2001
From: Aaron Ballman
Date: Tue, 2 Apr 2024 16:25:22 -0400
Subject: [PATCH 0004/1683] [C11] Claim conformance to WG14 N1482
This paper added a pile of macros for atomic initialization which we
started supporting in Clang 4.0.
---
clang/test/C/C11/n1482.c | 59 ++++++++++++++++++++++++++++++++++++++++
clang/www/c_status.html | 2 +-
2 files changed, 60 insertions(+), 1 deletion(-)
create mode 100644 clang/test/C/C11/n1482.c
diff --git a/clang/test/C/C11/n1482.c b/clang/test/C/C11/n1482.c
new file mode 100644
index 000000000000..493c7d0a18f1
--- /dev/null
+++ b/clang/test/C/C11/n1482.c
@@ -0,0 +1,59 @@
+// RUN: %clang_cc1 -verify -ffreestanding -std=c11 %s
+// expected-no-diagnostics
+
+/* WG14 N1482: Clang 4
+ * Explicit initializers for atomics
+ *
+ * NB: We can only test the compile time behavior from the paper, not the
+ * runtime behavior.
+ */
+
+#include
+
+#ifndef ATOMIC_BOOL_LOCK_FREE
+#error "Missing ATOMIC_BOOL_LOCK_FREE"
+#endif
+
+#ifndef ATOMIC_CHAR_LOCK_FREE
+#error "Missing ATOMIC_CHAR_LOCK_FREE"
+#endif
+
+#ifndef ATOMIC_CHAR16_T_LOCK_FREE
+#error "Missing ATOMIC_CHAR16_T_LOCK_FREE"
+#endif
+
+#ifndef ATOMIC_CHAR32_T_LOCK_FREE
+#error "Missing ATOMIC_CHAR32_T_LOCK_FREE"
+#endif
+
+#ifndef ATOMIC_WCHAR_T_LOCK_FREE
+#error "Missing ATOMIC_WCHAR_T_LOCK_FREE"
+#endif
+
+#ifndef ATOMIC_SHORT_LOCK_FREE
+#error "Missing ATOMIC_SHORT_LOCK_FREE"
+#endif
+
+#ifndef ATOMIC_INT_LOCK_FREE
+#error "Missing ATOMIC_INT_LOCK_FREE"
+#endif
+
+#ifndef ATOMIC_LONG_LOCK_FREE
+#error "Missing ATOMIC_LONG_LOCK_FREE"
+#endif
+
+#ifndef ATOMIC_LLONG_LOCK_FREE
+#error "Missing ATOMIC_LLONG_LOCK_FREE"
+#endif
+
+#ifndef ATOMIC_POINTER_LOCK_FREE
+#error "Missing ATOMIC_POINTER_LOCK_FREE"
+#endif
+
+#ifndef ATOMIC_VAR_INIT
+#error "Missing ATOMIC_VAR_INIT"
+#endif
+
+#ifndef atomic_init
+#error "Missing atomic_init"
+#endif
diff --git a/clang/www/c_status.html b/clang/www/c_status.html
index 62100611219e..370d99b55b13 100644
--- a/clang/www/c_status.html
+++ b/clang/www/c_status.html
@@ -548,7 +548,7 @@ conformance.
| Explicit initializers for atomics |
N1482 |
- Unknown |
+ Clang 4 |
| Atomics proposal (minus ternary op) |
--
GitLab
From 817c6915aeadbed100efee9e70eb42d9368d5ee6 Mon Sep 17 00:00:00 2001
From: Tom Stellard
Date: Tue, 2 Apr 2024 13:36:08 -0700
Subject: [PATCH 0005/1683] [workflows] Use issue-write workflow for warning
about private emails (#87399)
This allows us to use the less privileged pull_request event for the
validate_email job, since this workflow no longer writes a comment
directly.
---
.github/workflows/email-check.yaml | 32 ++++++++++++++++--------------
.github/workflows/issue-write.yml | 4 +++-
2 files changed, 20 insertions(+), 16 deletions(-)
diff --git a/.github/workflows/email-check.yaml b/.github/workflows/email-check.yaml
index ac53b5e527b0..8f32d020975f 100644
--- a/.github/workflows/email-check.yaml
+++ b/.github/workflows/email-check.yaml
@@ -1,7 +1,7 @@
name: "Check for private emails used in PRs"
on:
- pull_request_target:
+ pull_request:
types:
- opened
@@ -10,8 +10,6 @@ permissions:
jobs:
validate_email:
- permissions:
- pull-requests: write
runs-on: ubuntu-latest
if: github.repository == 'llvm/llvm-project'
steps:
@@ -25,20 +23,24 @@ jobs:
run: |
git log -1
echo "EMAIL=$(git show -s --format='%ae' HEAD~0)" >> $GITHUB_OUTPUT
+ # Create empty comment file
+ echo "[]" > comments
- name: Validate author email
if: ${{ endsWith(steps.author.outputs.EMAIL, 'noreply.github.com') }}
- uses: actions/github-script@v6
env:
- EMAIL: ${{ steps.author.outputs.EMAIL }}
+ COMMENT: >-
+ ⚠️ We detected that you are using a GitHub private e-mail address to contribute to the repo.
+ Please turn off [Keep my email addresses private](https://github.com/settings/emails) setting in your account.
+ See [LLVM Discourse](https://discourse.llvm.org/t/hidden-emails-on-github-should-we-do-something-about-it) for more information.
+ run: |
+ cat << EOF > comments
+ [{"body" : "$COMMENT"}]
+ EOF
+
+ - uses: actions/upload-artifact@26f96dfa697d77e81fd5907df203aa23a56210a8 #v4.3.0
+ if: always()
with:
- script: |
- const { EMAIL } = process.env
- await github.rest.issues.createComment({
- issue_number: context.issue.number,
- owner: context.repo.owner,
- repo: context.repo.repo,
- body: `⚠️ We detected that you are using a GitHub private e-mail address to contribute to the repo.
- Please turn off [Keep my email addresses private](https://github.com/settings/emails) setting in your account.
- See [LLVM Discourse](https://discourse.llvm.org/t/hidden-emails-on-github-should-we-do-something-about-it) for more information.
- `})
+ name: workflow-args
+ path: |
+ comments
diff --git a/.github/workflows/issue-write.yml b/.github/workflows/issue-write.yml
index 4a564a5076ba..e003be006c4e 100644
--- a/.github/workflows/issue-write.yml
+++ b/.github/workflows/issue-write.yml
@@ -2,7 +2,9 @@ name: Comment on an issue
on:
workflow_run:
- workflows: ["Check code formatting"]
+ workflows:
+ - "Check code formatting"
+ - "Check for private emails used in PRs"
types:
- completed
--
GitLab
From d8db13edd60b195e6692deb04d63e633319e486a Mon Sep 17 00:00:00 2001
From: Michael Maitland
Date: Tue, 2 Apr 2024 16:43:57 -0400
Subject: [PATCH 0006/1683] [RISCV][NFC] Delete some unused pseudo multiclasses
(#87401)
We only use the `RM` equivalents now.
---
.../Target/RISCV/RISCVInstrInfoVPseudos.td | 26 -------------------
1 file changed, 26 deletions(-)
diff --git a/llvm/lib/Target/RISCV/RISCVInstrInfoVPseudos.td b/llvm/lib/Target/RISCV/RISCVInstrInfoVPseudos.td
index 9424c461b34f..af20b11514ca 100644
--- a/llvm/lib/Target/RISCV/RISCVInstrInfoVPseudos.td
+++ b/llvm/lib/Target/RISCV/RISCVInstrInfoVPseudos.td
@@ -2910,22 +2910,6 @@ multiclass VPseudoVMAX_VV_VF {
}
}
-multiclass VPseudoVALU_VV_VF {
- foreach m = MxListF in {
- defm "" : VPseudoBinaryFV_VV,
- SchedBinary<"WriteVFALUV", "ReadVFALUV", "ReadVFALUV", m.MX,
- forceMergeOpRead=true>;
- }
-
- foreach f = FPList in {
- foreach m = f.MxList in {
- defm "" : VPseudoBinaryV_VF,
- SchedBinary<"WriteVFALUF", "ReadVFALUV", "ReadVFALUF", m.MX,
- forceMergeOpRead=true>;
- }
- }
-}
-
multiclass VPseudoVALU_VV_VF_RM {
foreach m = MxListF in {
defm "" : VPseudoBinaryFV_VV_RM,
@@ -2942,16 +2926,6 @@ multiclass VPseudoVALU_VV_VF_RM {
}
}
-multiclass VPseudoVALU_VF {
- foreach f = FPList in {
- foreach m = f.MxList in {
- defm "" : VPseudoBinaryV_VF,
- SchedBinary<"WriteVFALUF", "ReadVFALUV", "ReadVFALUF", m.MX,
- forceMergeOpRead=true>;
- }
- }
-}
-
multiclass VPseudoVALU_VF_RM {
foreach f = FPList in {
foreach m = f.MxList in {
--
GitLab
From 6261c53c6f9c82bfca5df2ed5917876d748ed079 Mon Sep 17 00:00:00 2001
From: Florian Hahn
Date: Tue, 2 Apr 2024 21:48:35 +0100
Subject: [PATCH 0007/1683] [VPlan] Make sure OR VPInstructions are treated as
disjoint ops.
Make sure that VPInstructions with OR opcodes are properly registered as
disjoint ops.
Fixes https://github.com/llvm/llvm-project/issues/87378.
---
.../Vectorize/LoopVectorizationPlanner.h | 5 +-
llvm/lib/Transforms/Vectorize/VPlan.h | 22 ++-
...ruction-or-drop-poison-generating-flags.ll | 132 ++++++++++++++++++
3 files changed, 155 insertions(+), 4 deletions(-)
create mode 100644 llvm/test/Transforms/LoopVectorize/pr87378-vpinstruction-or-drop-poison-generating-flags.ll
diff --git a/llvm/lib/Transforms/Vectorize/LoopVectorizationPlanner.h b/llvm/lib/Transforms/Vectorize/LoopVectorizationPlanner.h
index 5d03b66b0ce3..ebca2d855a46 100644
--- a/llvm/lib/Transforms/Vectorize/LoopVectorizationPlanner.h
+++ b/llvm/lib/Transforms/Vectorize/LoopVectorizationPlanner.h
@@ -175,7 +175,10 @@ public:
VPValue *createOr(VPValue *LHS, VPValue *RHS, DebugLoc DL = {},
const Twine &Name = "") {
- return createInstruction(Instruction::BinaryOps::Or, {LHS, RHS}, DL, Name);
+
+ return tryInsertInstruction(new VPInstruction(
+ Instruction::BinaryOps::Or, {LHS, RHS},
+ VPRecipeWithIRFlags::DisjointFlagsTy(false), DL, Name));
}
VPValue *createSelect(VPValue *Cond, VPValue *TrueVal, VPValue *FalseVal,
diff --git a/llvm/lib/Transforms/Vectorize/VPlan.h b/llvm/lib/Transforms/Vectorize/VPlan.h
index 3baca43f7276..707a826ecdc2 100644
--- a/llvm/lib/Transforms/Vectorize/VPlan.h
+++ b/llvm/lib/Transforms/Vectorize/VPlan.h
@@ -913,6 +913,11 @@ public:
WrapFlagsTy(bool HasNUW, bool HasNSW) : HasNUW(HasNUW), HasNSW(HasNSW) {}
};
+ struct DisjointFlagsTy {
+ char IsDisjoint : 1;
+ DisjointFlagsTy(bool IsDisjoint) : IsDisjoint(IsDisjoint) {}
+ };
+
protected:
struct GEPFlagsTy {
char IsInBounds : 1;
@@ -920,9 +925,6 @@ protected:
};
private:
- struct DisjointFlagsTy {
- char IsDisjoint : 1;
- };
struct ExactFlagsTy {
char IsExact : 1;
};
@@ -1016,6 +1018,12 @@ public:
: VPSingleDefRecipe(SC, Operands, DL), OpType(OperationType::FPMathOp),
FMFs(FMFs) {}
+ template
+ VPRecipeWithIRFlags(const unsigned char SC, IterT Operands,
+ DisjointFlagsTy DisjointFlags, DebugLoc DL = {})
+ : VPSingleDefRecipe(SC, Operands, DL), OpType(OperationType::DisjointOp),
+ DisjointFlags(DisjointFlags) {}
+
protected:
template
VPRecipeWithIRFlags(const unsigned char SC, IterT Operands,
@@ -1221,6 +1229,14 @@ public:
: VPRecipeWithIRFlags(VPDef::VPInstructionSC, Operands, WrapFlags, DL),
Opcode(Opcode), Name(Name.str()) {}
+ VPInstruction(unsigned Opcode, std::initializer_list Operands,
+ DisjointFlagsTy DisjointFlag, DebugLoc DL = {},
+ const Twine &Name = "")
+ : VPRecipeWithIRFlags(VPDef::VPInstructionSC, Operands, DisjointFlag, DL),
+ Opcode(Opcode), Name(Name.str()) {
+ assert(Opcode == Instruction::Or && "only OR opcodes can be disjoint");
+ }
+
VPInstruction(unsigned Opcode, std::initializer_list Operands,
FastMathFlags FMFs, DebugLoc DL = {}, const Twine &Name = "");
diff --git a/llvm/test/Transforms/LoopVectorize/pr87378-vpinstruction-or-drop-poison-generating-flags.ll b/llvm/test/Transforms/LoopVectorize/pr87378-vpinstruction-or-drop-poison-generating-flags.ll
new file mode 100644
index 000000000000..4e38630209b2
--- /dev/null
+++ b/llvm/test/Transforms/LoopVectorize/pr87378-vpinstruction-or-drop-poison-generating-flags.ll
@@ -0,0 +1,132 @@
+; NOTE: Assertions have been autogenerated by utils/update_test_checks.py UTC_ARGS: --version 4
+; RUN: opt -p loop-vectorize -mattr="+v" -S %s | FileCheck %s
+
+target datalayout = "e-m:e-p:64:64-i64:64-i128:128-n32:64-S128"
+target triple = "riscv64-unknown-linux-gnu"
+
+; Test case for https://github.com/llvm/llvm-project/issues/87378.
+define void @pr87378_vpinstruction_or_drop_poison_generating_flags(ptr %arg, i64 %a, i64 %b, i64 %c) {
+; CHECK-LABEL: define void @pr87378_vpinstruction_or_drop_poison_generating_flags(
+; CHECK-SAME: ptr [[ARG:%.*]], i64 [[A:%.*]], i64 [[B:%.*]], i64 [[C:%.*]]) #[[ATTR0:[0-9]+]] {
+; CHECK-NEXT: entry:
+; CHECK-NEXT: [[TMP0:%.*]] = call i64 @llvm.vscale.i64()
+; CHECK-NEXT: [[TMP1:%.*]] = mul i64 [[TMP0]], 8
+; CHECK-NEXT: [[MIN_ITERS_CHECK:%.*]] = icmp ult i64 1001, [[TMP1]]
+; CHECK-NEXT: br i1 [[MIN_ITERS_CHECK]], label [[SCALAR_PH:%.*]], label [[VECTOR_PH:%.*]]
+; CHECK: vector.ph:
+; CHECK-NEXT: [[TMP2:%.*]] = call i64 @llvm.vscale.i64()
+; CHECK-NEXT: [[TMP3:%.*]] = mul i64 [[TMP2]], 8
+; CHECK-NEXT: [[N_MOD_VF:%.*]] = urem i64 1001, [[TMP3]]
+; CHECK-NEXT: [[N_VEC:%.*]] = sub i64 1001, [[N_MOD_VF]]
+; CHECK-NEXT: [[TMP4:%.*]] = call i64 @llvm.vscale.i64()
+; CHECK-NEXT: [[TMP5:%.*]] = mul i64 [[TMP4]], 8
+; CHECK-NEXT: [[TMP6:%.*]] = call @llvm.experimental.stepvector.nxv8i64()
+; CHECK-NEXT: [[TMP7:%.*]] = add [[TMP6]], zeroinitializer
+; CHECK-NEXT: [[TMP8:%.*]] = mul [[TMP7]], shufflevector ( insertelement ( poison, i64 1, i64 0), poison, zeroinitializer)
+; CHECK-NEXT: [[INDUCTION:%.*]] = add zeroinitializer, [[TMP8]]
+; CHECK-NEXT: [[TMP9:%.*]] = call i64 @llvm.vscale.i64()
+; CHECK-NEXT: [[TMP10:%.*]] = mul i64 [[TMP9]], 8
+; CHECK-NEXT: [[TMP11:%.*]] = mul i64 1, [[TMP10]]
+; CHECK-NEXT: [[DOTSPLATINSERT:%.*]] = insertelement poison, i64 [[TMP11]], i64 0
+; CHECK-NEXT: [[DOTSPLAT:%.*]] = shufflevector [[DOTSPLATINSERT]], poison, zeroinitializer
+; CHECK-NEXT: [[BROADCAST_SPLATINSERT:%.*]] = insertelement poison, i64 [[A]], i64 0
+; CHECK-NEXT: [[BROADCAST_SPLAT:%.*]] = shufflevector [[BROADCAST_SPLATINSERT]], poison, zeroinitializer
+; CHECK-NEXT: [[BROADCAST_SPLATINSERT1:%.*]] = insertelement poison, i64 [[B]], i64 0
+; CHECK-NEXT: [[BROADCAST_SPLAT2:%.*]] = shufflevector [[BROADCAST_SPLATINSERT1]], poison, zeroinitializer
+; CHECK-NEXT: [[BROADCAST_SPLATINSERT3:%.*]] = insertelement poison, i64 [[C]], i64 0
+; CHECK-NEXT: [[BROADCAST_SPLAT4:%.*]] = shufflevector [[BROADCAST_SPLATINSERT3]], poison, zeroinitializer
+; CHECK-NEXT: br label [[VECTOR_BODY:%.*]]
+; CHECK: vector.body:
+; CHECK-NEXT: [[INDEX:%.*]] = phi i64 [ 0, [[VECTOR_PH]] ], [ [[INDEX_NEXT:%.*]], [[VECTOR_BODY]] ]
+; CHECK-NEXT: [[VEC_IND:%.*]] = phi [ [[INDUCTION]], [[VECTOR_PH]] ], [ [[VEC_IND_NEXT:%.*]], [[VECTOR_BODY]] ]
+; CHECK-NEXT: [[TMP12:%.*]] = add i64 [[INDEX]], 0
+; CHECK-NEXT: [[BROADCAST_SPLATINSERT5:%.*]] = insertelement poison, i64 [[TMP12]], i64 0
+; CHECK-NEXT: [[BROADCAST_SPLAT6:%.*]] = shufflevector [[BROADCAST_SPLATINSERT5]], poison, zeroinitializer
+; CHECK-NEXT: [[TMP13:%.*]] = icmp ule [[VEC_IND]], [[BROADCAST_SPLAT]]
+; CHECK-NEXT: [[TMP14:%.*]] = icmp ule [[VEC_IND]], [[BROADCAST_SPLAT2]]
+; CHECK-NEXT: [[TMP15:%.*]] = select [[TMP13]], [[TMP14]], zeroinitializer
+; CHECK-NEXT: [[TMP16:%.*]] = xor [[TMP13]], shufflevector ( insertelement ( poison, i1 true, i64 0), poison, zeroinitializer)
+; CHECK-NEXT: [[TMP17:%.*]] = or [[TMP15]], [[TMP16]]
+; CHECK-NEXT: [[TMP18:%.*]] = icmp ule [[VEC_IND]], [[BROADCAST_SPLAT4]]
+; CHECK-NEXT: [[TMP19:%.*]] = select [[TMP17]], [[TMP18]], zeroinitializer
+; CHECK-NEXT: [[TMP20:%.*]] = xor [[TMP14]], shufflevector ( insertelement ( poison, i1 true, i64 0), poison, zeroinitializer)
+; CHECK-NEXT: [[TMP21:%.*]] = select [[TMP13]], [[TMP20]], zeroinitializer
+; CHECK-NEXT: [[TMP22:%.*]] = or [[TMP19]], [[TMP21]]
+; CHECK-NEXT: [[PREDPHI:%.*]] = select [[TMP19]], [[BROADCAST_SPLAT6]], shufflevector ( insertelement ( poison, i64 poison, i64 0), poison, zeroinitializer)
+; CHECK-NEXT: [[TMP23:%.*]] = extractelement [[PREDPHI]], i32 0
+; CHECK-NEXT: [[TMP24:%.*]] = getelementptr i16, ptr [[ARG]], i64 [[TMP23]]
+; CHECK-NEXT: [[TMP25:%.*]] = getelementptr i16, ptr [[TMP24]], i32 0
+; CHECK-NEXT: call void @llvm.masked.store.nxv8i16.p0( zeroinitializer, ptr [[TMP25]], i32 2, [[TMP22]])
+; CHECK-NEXT: [[INDEX_NEXT]] = add nuw i64 [[INDEX]], [[TMP5]]
+; CHECK-NEXT: [[VEC_IND_NEXT]] = add [[VEC_IND]], [[DOTSPLAT]]
+; CHECK-NEXT: [[TMP26:%.*]] = icmp eq i64 [[INDEX_NEXT]], [[N_VEC]]
+; CHECK-NEXT: br i1 [[TMP26]], label [[MIDDLE_BLOCK:%.*]], label [[VECTOR_BODY]], !llvm.loop [[LOOP0:![0-9]+]]
+; CHECK: middle.block:
+; CHECK-NEXT: [[CMP_N:%.*]] = icmp eq i64 1001, [[N_VEC]]
+; CHECK-NEXT: br i1 [[CMP_N]], label [[EXIT:%.*]], label [[SCALAR_PH]]
+; CHECK: scalar.ph:
+; CHECK-NEXT: [[BC_RESUME_VAL:%.*]] = phi i64 [ [[N_VEC]], [[MIDDLE_BLOCK]] ], [ 0, [[ENTRY:%.*]] ]
+; CHECK-NEXT: br label [[LOOP_HEADER:%.*]]
+; CHECK: loop.header:
+; CHECK-NEXT: [[IV:%.*]] = phi i64 [ [[BC_RESUME_VAL]], [[SCALAR_PH]] ], [ [[IV_NEXT:%.*]], [[LOOP_LATCH:%.*]] ]
+; CHECK-NEXT: [[C_1:%.*]] = icmp ule i64 [[IV]], [[A]]
+; CHECK-NEXT: br i1 [[C_1]], label [[THEN_1:%.*]], label [[ELSE_1:%.*]]
+; CHECK: then.1:
+; CHECK-NEXT: [[C_2:%.*]] = icmp ule i64 [[IV]], [[B]]
+; CHECK-NEXT: br i1 [[C_2]], label [[ELSE_1]], label [[MERGE:%.*]]
+; CHECK: else.1:
+; CHECK-NEXT: [[C_3:%.*]] = icmp ule i64 [[IV]], [[C]]
+; CHECK-NEXT: br i1 [[C_3]], label [[THEN_2:%.*]], label [[LOOP_LATCH]]
+; CHECK: then.2:
+; CHECK-NEXT: br label [[MERGE]]
+; CHECK: merge:
+; CHECK-NEXT: [[IDX:%.*]] = phi i64 [ poison, [[THEN_1]] ], [ [[IV]], [[THEN_2]] ]
+; CHECK-NEXT: [[GETELEMENTPTR:%.*]] = getelementptr i16, ptr [[ARG]], i64 [[IDX]]
+; CHECK-NEXT: store i16 0, ptr [[GETELEMENTPTR]], align 2
+; CHECK-NEXT: br label [[LOOP_LATCH]]
+; CHECK: loop.latch:
+; CHECK-NEXT: [[IV_NEXT]] = add i64 [[IV]], 1
+; CHECK-NEXT: [[ICMP:%.*]] = icmp eq i64 [[IV]], 1000
+; CHECK-NEXT: br i1 [[ICMP]], label [[EXIT]], label [[LOOP_HEADER]], !llvm.loop [[LOOP3:![0-9]+]]
+; CHECK: exit:
+; CHECK-NEXT: ret void
+;
+entry:
+ br label %loop.header
+
+loop.header:
+ %iv = phi i64 [ 0, %entry ], [ %iv.next, %loop.latch ]
+ %c.1 = icmp ule i64 %iv, %a
+ br i1 %c.1, label %then.1, label %else.1
+
+then.1:
+ %c.2 = icmp ule i64 %iv, %b
+ br i1 %c.2, label %else.1, label %merge
+
+else.1:
+ %c.3 = icmp ule i64 %iv, %c
+ br i1 %c.3, label %then.2, label %loop.latch
+
+then.2:
+ br label %merge
+
+merge:
+ %idx = phi i64 [ poison, %then.1 ], [ %iv, %then.2 ]
+ %getelementptr = getelementptr i16, ptr %arg, i64 %idx
+ store i16 0, ptr %getelementptr, align 2
+ br label %loop.latch
+
+loop.latch:
+ %iv.next = add i64 %iv, 1
+ %icmp = icmp eq i64 %iv, 1000
+ br i1 %icmp, label %exit, label %loop.header
+
+exit:
+ ret void
+}
+;.
+; CHECK: [[LOOP0]] = distinct !{[[LOOP0]], [[META1:![0-9]+]], [[META2:![0-9]+]]}
+; CHECK: [[META1]] = !{!"llvm.loop.isvectorized", i32 1}
+; CHECK: [[META2]] = !{!"llvm.loop.unroll.runtime.disable"}
+; CHECK: [[LOOP3]] = distinct !{[[LOOP3]], [[META2]], [[META1]]}
+;.
--
GitLab
From c3e3d59fab8ae8161810c861d78c7b5fcabb1a2e Mon Sep 17 00:00:00 2001
From: Han-Chung Wang
Date: Tue, 2 Apr 2024 13:49:28 -0700
Subject: [PATCH 0008/1683] [mlir][tensor] Fix tensor::PackOp fold() handling
of padding value (#87296)
We can't just check if it is a splat constant or not. We should also
check if the value match.
---
mlir/lib/Dialect/Tensor/IR/TensorOps.cpp | 16 +++++++----
mlir/test/Dialect/Tensor/canonicalize.mlir | 33 ++++++++++++++++++++++
2 files changed, 44 insertions(+), 5 deletions(-)
diff --git a/mlir/lib/Dialect/Tensor/IR/TensorOps.cpp b/mlir/lib/Dialect/Tensor/IR/TensorOps.cpp
index 38a9ad60bb79..0ce40e813712 100644
--- a/mlir/lib/Dialect/Tensor/IR/TensorOps.cpp
+++ b/mlir/lib/Dialect/Tensor/IR/TensorOps.cpp
@@ -1068,10 +1068,13 @@ void EmptyOp::getCanonicalizationPatterns(RewritePatternSet &results,
/// Try to remove a tensor operation if it would only reshape a constant.
/// Removes the op and replaces the constant with a new constant of the result
-/// shape.
-static OpFoldResult reshapeConstantSource(DenseElementsAttr source,
- TensorType result) {
- if (source && source.isSplat() && result.hasStaticShape())
+/// shape. When an optional cst attribute is passed, it is reshaped only if the
+/// splat value matches the value in the attribute.
+static OpFoldResult
+reshapeConstantSource(DenseElementsAttr source, TensorType result,
+ std::optional cst = std::nullopt) {
+ if (source && source.isSplat() && result.hasStaticShape() &&
+ (!cst.has_value() || source.getSplatValue() == cst.value()))
return source.resizeSplat(result);
return {};
@@ -4143,9 +4146,12 @@ bool PackOp::isLikePad() {
}
OpFoldResult PackOp::fold(FoldAdaptor adaptor) {
+ std::optional paddingValue;
+ if (auto pad = adaptor.getPaddingValue())
+ paddingValue = pad;
if (OpFoldResult reshapedSource = reshapeConstantSource(
llvm::dyn_cast_if_present(adaptor.getSource()),
- getResult().getType()))
+ getDestType(), paddingValue))
return reshapedSource;
return {};
}
diff --git a/mlir/test/Dialect/Tensor/canonicalize.mlir b/mlir/test/Dialect/Tensor/canonicalize.mlir
index 9ab54fe9c133..ac365c9d297e 100644
--- a/mlir/test/Dialect/Tensor/canonicalize.mlir
+++ b/mlir/test/Dialect/Tensor/canonicalize.mlir
@@ -830,6 +830,39 @@ func.func @fold_pack_constant_splat(%dest : tensor<8x16x8x32xf32>) -> tensor<8x1
// -----
+// CHECK-LABEL: func @fold_padding_value_pack_constant_splat
+// CHECK-NOT: tensor.pack
+// CHECK: arith.constant dense<1.000000e-01> : tensor<8x16x8x32xf32>
+func.func @fold_padding_value_pack_constant_splat(%dest : tensor<8x16x8x32xf32>) -> tensor<8x16x8x32xf32> {
+ %pad = arith.constant 1.000000e-01 : f32
+ %cst = arith.constant dense<1.000000e-01> : tensor<63x127xf32>
+ %0 = tensor.pack %cst
+ padding_value(%pad : f32)
+ outer_dims_perm = [1, 0] inner_dims_pos = [0, 1]
+ inner_tiles = [8, 32] into %dest : tensor<63x127xf32> -> tensor<8x16x8x32xf32>
+ return %0 : tensor<8x16x8x32xf32>
+}
+
+
+// -----
+
+// CHECK-LABEL: func @nofold_padding_value_pack_constant_splat
+// CHECK: arith.constant dense<1.000000e-01> : tensor<63x127xf32>
+// CHECK: tensor.pack
+func.func @nofold_padding_value_pack_constant_splat(%dest : tensor<8x16x8x32xf32>) -> tensor<8x16x8x32xf32> {
+ %pad = arith.constant 0.0 : f32
+ %cst = arith.constant dense<1.000000e-01> : tensor<63x127xf32>
+ %0 = tensor.pack %cst
+ padding_value(%pad : f32)
+ outer_dims_perm = [1, 0]
+ inner_dims_pos = [0, 1]
+ inner_tiles = [8, 32]
+ into %dest : tensor<63x127xf32> -> tensor<8x16x8x32xf32>
+ return %0 : tensor<8x16x8x32xf32>
+}
+
+// -----
+
func.func @fold_padding_value_pack(%arg0: tensor<1200x500000xf32>) -> tensor<31250x1200x16x1xf32> {
%cst = arith.constant 0.000000e+00 : f32
%0 = tensor.empty() : tensor<31250x1200x16x1xf32>
--
GitLab
From f119a4ffb885ed588c46de1d51f4185572142ca2 Mon Sep 17 00:00:00 2001
From: Chris Bieneman
Date: Tue, 2 Apr 2024 15:52:24 -0500
Subject: [PATCH 0009/1683] [HLSL] Fix broken spir-v test
This fixes a test broken in 3d469c0e7c3072f0dad0f5e9bd0c74dffaf83cd3.
fast-forwarded.
../clang/test/CodeGenHLSL/builtins/wave_get_lane_index_subcall.hlsl
---
.../test/CodeGenHLSL/builtins/wave_get_lane_index_subcall.hlsl | 2 +-
1 file changed, 1 insertion(+), 1 deletion(-)
diff --git a/clang/test/CodeGenHLSL/builtins/wave_get_lane_index_subcall.hlsl b/clang/test/CodeGenHLSL/builtins/wave_get_lane_index_subcall.hlsl
index 379c8f118f52..6ea80d692cd2 100644
--- a/clang/test/CodeGenHLSL/builtins/wave_get_lane_index_subcall.hlsl
+++ b/clang/test/CodeGenHLSL/builtins/wave_get_lane_index_subcall.hlsl
@@ -12,7 +12,7 @@ uint test_1() {
// CHECK: define spir_func noundef i32 @_Z6test_2v() [[A0]] {
// CHECK: %[[C2:[0-9]+]] = call token @llvm.experimental.convergence.entry()
-// CHECK: call spir_func noundef i32 @_Z6test_1v() [ "convergencectrl"(token %[[C2]]) ]
+// CHECK: call spir_func noundef i32 @_Z6test_1v() {{#[0-9]+}} [ "convergencectrl"(token %[[C2]]) ]
uint test_2() {
return test_1();
}
--
GitLab
From 4ea24946e356be31446fc30ca3d11cc5783ba2a6 Mon Sep 17 00:00:00 2001
From: Jonathan Peyton
Date: Tue, 2 Apr 2024 15:56:50 -0500
Subject: [PATCH 0010/1683] [OpenMP] Fix nested parallel with tasking (#87309)
When a nested parallel region ends, the runtime calls __kmp_join_call().
During this call, the primary thread of the nested parallel region will
reset its tid (retval of omp_get_thread_num()) to what it was in the
outer parallel region. A data race occurs with the current code when
another worker thread from the nested inner parallel region tries to
steal tasks from the primary thread's task deque. The worker thread
reads the tid value directly from the primary thread's data structure
and may read the wrong value.
This change just uses the calculated victim_tid from execute_tasks()
directly in the steal_task() routine rather than reading tid from the
data structure.
Fixes: #87307
---
openmp/runtime/src/kmp_tasking.cpp | 15 ++++----
openmp/runtime/test/tasking/issue-87307.c | 43 +++++++++++++++++++++++
2 files changed, 52 insertions(+), 6 deletions(-)
create mode 100644 openmp/runtime/test/tasking/issue-87307.c
diff --git a/openmp/runtime/src/kmp_tasking.cpp b/openmp/runtime/src/kmp_tasking.cpp
index 155e17ba7ec8..6303bb0d63f0 100644
--- a/openmp/runtime/src/kmp_tasking.cpp
+++ b/openmp/runtime/src/kmp_tasking.cpp
@@ -3219,7 +3219,7 @@ static kmp_task_t *__kmp_remove_my_task(kmp_info_t *thread, kmp_int32 gtid,
// __kmp_steal_task: remove a task from another thread's deque
// Assume that calling thread has already checked existence of
// task_team thread_data before calling this routine.
-static kmp_task_t *__kmp_steal_task(kmp_info_t *victim_thr, kmp_int32 gtid,
+static kmp_task_t *__kmp_steal_task(kmp_int32 victim_tid, kmp_int32 gtid,
kmp_task_team_t *task_team,
std::atomic *unfinished_threads,
int *thread_finished,
@@ -3229,15 +3229,18 @@ static kmp_task_t *__kmp_steal_task(kmp_info_t *victim_thr, kmp_int32 gtid,
kmp_taskdata_t *current;
kmp_thread_data_t *victim_td, *threads_data;
kmp_int32 target;
- kmp_int32 victim_tid;
+ kmp_info_t *victim_thr;
KMP_DEBUG_ASSERT(__kmp_tasking_mode != tskm_immediate_exec);
threads_data = task_team->tt.tt_threads_data;
KMP_DEBUG_ASSERT(threads_data != NULL); // Caller should check this condition
+ KMP_DEBUG_ASSERT(victim_tid >= 0);
+ KMP_DEBUG_ASSERT(victim_tid < task_team->tt.tt_nproc);
- victim_tid = victim_thr->th.th_info.ds.ds_tid;
victim_td = &threads_data[victim_tid];
+ victim_thr = victim_td->td.td_thr;
+ (void)victim_thr; // Use in TRACE messages which aren't always enabled.
KA_TRACE(10, ("__kmp_steal_task(enter): T#%d try to steal from T#%d: "
"task_team=%p ntasks=%d head=%u tail=%u\n",
@@ -3452,9 +3455,9 @@ static inline int __kmp_execute_tasks_template(
if (!asleep) {
// We have a victim to try to steal from
- task = __kmp_steal_task(other_thread, gtid, task_team,
- unfinished_threads, thread_finished,
- is_constrained);
+ task =
+ __kmp_steal_task(victim_tid, gtid, task_team, unfinished_threads,
+ thread_finished, is_constrained);
}
if (task != NULL) { // set last stolen to victim
if (threads_data[tid].td.td_deque_last_stolen != victim_tid) {
diff --git a/openmp/runtime/test/tasking/issue-87307.c b/openmp/runtime/test/tasking/issue-87307.c
new file mode 100644
index 000000000000..f889ae20329e
--- /dev/null
+++ b/openmp/runtime/test/tasking/issue-87307.c
@@ -0,0 +1,43 @@
+// RUN: %libomp-compile-and-run
+#include
+#include
+#include
+
+int a;
+
+void inc_a() {
+#pragma omp task
+ {
+#pragma omp atomic
+ a++;
+ }
+}
+
+int main() {
+ int n;
+ int nth_outer;
+ omp_set_max_active_levels(2);
+ omp_set_dynamic(0);
+
+ for (n = 0; n < 200; ++n) {
+ a = 0;
+#pragma omp parallel num_threads(8)
+ {
+ if (omp_get_thread_num() == 0)
+ nth_outer = omp_get_num_threads();
+#pragma omp parallel num_threads(2)
+ {
+ int i;
+#pragma omp master
+ for (i = 0; i < 50; ++i)
+ inc_a();
+ }
+ }
+ if (a != nth_outer * 50) {
+ fprintf(stderr, "error: a (%d) != %d\n", a, nth_outer * 50);
+ return EXIT_FAILURE;
+ }
+ }
+
+ return EXIT_SUCCESS;
+}
--
GitLab
From 89271b46761749503dffe94c60b9cbe0bda80284 Mon Sep 17 00:00:00 2001
From: Florian Hahn
Date: Tue, 2 Apr 2024 22:02:25 +0100
Subject: [PATCH 0011/1683] [LV] Add test depending on target to RISCV
subdirectory.
---
.../pr87378-vpinstruction-or-drop-poison-generating-flags.ll | 0
1 file changed, 0 insertions(+), 0 deletions(-)
rename llvm/test/Transforms/LoopVectorize/{ => RISCV}/pr87378-vpinstruction-or-drop-poison-generating-flags.ll (100%)
diff --git a/llvm/test/Transforms/LoopVectorize/pr87378-vpinstruction-or-drop-poison-generating-flags.ll b/llvm/test/Transforms/LoopVectorize/RISCV/pr87378-vpinstruction-or-drop-poison-generating-flags.ll
similarity index 100%
rename from llvm/test/Transforms/LoopVectorize/pr87378-vpinstruction-or-drop-poison-generating-flags.ll
rename to llvm/test/Transforms/LoopVectorize/RISCV/pr87378-vpinstruction-or-drop-poison-generating-flags.ll
--
GitLab
From 7327014b49ba0e4c8227edaed569d21d3cc1ec74 Mon Sep 17 00:00:00 2001
From: Joseph Huber
Date: Tue, 2 Apr 2024 16:25:18 -0500
Subject: [PATCH 0012/1683] [libc] Implement temporary `printf` on the GPU
(#85331)
Summary:
This patch adds a temporary implementation that uses a struct-based
interface in lieu of varargs support. Once varargs support exists we
will move this implementation to the "real" printf implementation.
Conceptually, this patch has the client copy over its format string and
arguments to the server. The server will then scan the format string
searching for any specifiers that are actually a string. If it is a
string then we will send the pointer back to the server to tell it to
copy it back. This copied value will then replace the pointer when the
final formatting is done.
This will require a built-in extension to the varargs support to get
access to the underlying struct. The varargs used on the GPU will simply
be a struct wrapped in a varargs ABI.
---
libc/config/gpu/entrypoints.txt | 1 +
libc/include/llvm-libc-types/rpc_opcodes_t.h | 3 +
libc/spec/gpu_ext.td | 8 +
libc/src/__support/arg_list.h | 38 +++++
libc/src/gpu/CMakeLists.txt | 12 ++
libc/src/gpu/rpc_fprintf.cpp | 71 +++++++++
libc/src/gpu/rpc_fprintf.h | 22 +++
.../test/integration/src/stdio/CMakeLists.txt | 3 +
.../integration/src/stdio/gpu/CMakeLists.txt | 21 +++
.../test/integration/src/stdio/gpu/printf.cpp | 88 +++++++++++
libc/utils/gpu/server/CMakeLists.txt | 10 +-
libc/utils/gpu/server/rpc_server.cpp | 148 ++++++++++++++++++
12 files changed, 424 insertions(+), 1 deletion(-)
create mode 100644 libc/src/gpu/rpc_fprintf.cpp
create mode 100644 libc/src/gpu/rpc_fprintf.h
create mode 100644 libc/test/integration/src/stdio/gpu/CMakeLists.txt
create mode 100644 libc/test/integration/src/stdio/gpu/printf.cpp
diff --git a/libc/config/gpu/entrypoints.txt b/libc/config/gpu/entrypoints.txt
index 4fb87cb9f5a3..b678350e9fcb 100644
--- a/libc/config/gpu/entrypoints.txt
+++ b/libc/config/gpu/entrypoints.txt
@@ -211,6 +211,7 @@ set(TARGET_LIBC_ENTRYPOINTS
# gpu/rpc.h entrypoints
libc.src.gpu.rpc_host_call
+ libc.src.gpu.rpc_fprintf
)
set(TARGET_LIBM_ENTRYPOINTS
diff --git a/libc/include/llvm-libc-types/rpc_opcodes_t.h b/libc/include/llvm-libc-types/rpc_opcodes_t.h
index 919ea039c18e..faed7b5f5ff4 100644
--- a/libc/include/llvm-libc-types/rpc_opcodes_t.h
+++ b/libc/include/llvm-libc-types/rpc_opcodes_t.h
@@ -31,6 +31,9 @@ typedef enum {
RPC_FTELL,
RPC_FFLUSH,
RPC_UNGETC,
+ RPC_PRINTF_TO_STDOUT,
+ RPC_PRINTF_TO_STDERR,
+ RPC_PRINTF_TO_STREAM,
RPC_LAST = 0xFFFF,
} rpc_opcode_t;
diff --git a/libc/spec/gpu_ext.td b/libc/spec/gpu_ext.td
index dce81ff77862..5400e0afa756 100644
--- a/libc/spec/gpu_ext.td
+++ b/libc/spec/gpu_ext.td
@@ -10,6 +10,14 @@ def GPUExtensions : StandardSpec<"GPUExtensions"> {
RetValSpec,
[ArgSpec, ArgSpec, ArgSpec]
>,
+ FunctionSpec<
+ "rpc_fprintf",
+ RetValSpec,
+ [ArgSpec,
+ ArgSpec,
+ ArgSpec,
+ ArgSpec]
+ >,
]
>;
let Headers = [
diff --git a/libc/src/__support/arg_list.h b/libc/src/__support/arg_list.h
index 9de17651142f..0965e12afd56 100644
--- a/libc/src/__support/arg_list.h
+++ b/libc/src/__support/arg_list.h
@@ -13,6 +13,7 @@
#include
#include
+#include
namespace LIBC_NAMESPACE {
namespace internal {
@@ -60,6 +61,43 @@ public:
size_t read_count() const { return arg_counter; }
};
+// Used for the GPU implementation of `printf`. This models a variadic list as a
+// simple array of pointers that are built manually by the implementation.
+class StructArgList {
+ void *ptr;
+ void *end;
+
+public:
+ LIBC_INLINE StructArgList(void *ptr, size_t size)
+ : ptr(ptr), end(reinterpret_cast(ptr) + size) {}
+ LIBC_INLINE StructArgList(const StructArgList &other) {
+ ptr = other.ptr;
+ end = other.end;
+ }
+ LIBC_INLINE StructArgList() = default;
+ LIBC_INLINE ~StructArgList() = default;
+
+ LIBC_INLINE StructArgList &operator=(const StructArgList &rhs) {
+ ptr = rhs.ptr;
+ return *this;
+ }
+
+ LIBC_INLINE void *get_ptr() const { return ptr; }
+
+ template LIBC_INLINE T next_var() {
+ ptr = reinterpret_cast(
+ ((reinterpret_cast(ptr) + alignof(T) - 1) / alignof(T)) *
+ alignof(T));
+
+ if (ptr >= end)
+ return T(-1);
+
+ T val = *reinterpret_cast(ptr);
+ ptr = reinterpret_cast(ptr) + sizeof(T);
+ return val;
+ }
+};
+
} // namespace internal
} // namespace LIBC_NAMESPACE
diff --git a/libc/src/gpu/CMakeLists.txt b/libc/src/gpu/CMakeLists.txt
index e20228516b51..4508abea7a88 100644
--- a/libc/src/gpu/CMakeLists.txt
+++ b/libc/src/gpu/CMakeLists.txt
@@ -8,3 +8,15 @@ add_entrypoint_object(
libc.src.__support.RPC.rpc_client
libc.src.__support.GPU.utils
)
+
+add_entrypoint_object(
+ rpc_fprintf
+ SRCS
+ rpc_fprintf.cpp
+ HDRS
+ rpc_fprintf.h
+ DEPENDS
+ libc.src.stdio.gpu.gpu_file
+ libc.src.__support.RPC.rpc_client
+ libc.src.__support.GPU.utils
+)
diff --git a/libc/src/gpu/rpc_fprintf.cpp b/libc/src/gpu/rpc_fprintf.cpp
new file mode 100644
index 000000000000..7b0e60b59baf
--- /dev/null
+++ b/libc/src/gpu/rpc_fprintf.cpp
@@ -0,0 +1,71 @@
+//===-- GPU implementation of fprintf -------------------------------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#include "rpc_fprintf.h"
+
+#include "src/__support/CPP/string_view.h"
+#include "src/__support/GPU/utils.h"
+#include "src/__support/RPC/rpc_client.h"
+#include "src/__support/common.h"
+#include "src/stdio/gpu/file.h"
+
+namespace LIBC_NAMESPACE {
+
+template
+int fprintf_impl(::FILE *__restrict file, const char *__restrict format,
+ size_t format_size, void *args, size_t args_size) {
+ uint64_t mask = gpu::get_lane_mask();
+ rpc::Client::Port port = rpc::client.open();
+
+ if constexpr (opcode == RPC_PRINTF_TO_STREAM) {
+ port.send([&](rpc::Buffer *buffer) {
+ buffer->data[0] = reinterpret_cast(file);
+ });
+ }
+
+ port.send_n(format, format_size);
+ port.send_n(args, args_size);
+
+ uint32_t ret = 0;
+ for (;;) {
+ const char *str = nullptr;
+ port.recv([&](rpc::Buffer *buffer) {
+ ret = static_cast(buffer->data[0]);
+ str = reinterpret_cast(buffer->data[1]);
+ });
+ // If any lanes have a string argument it needs to be copied back.
+ if (!gpu::ballot(mask, str))
+ break;
+
+ uint64_t size = str ? internal::string_length(str) + 1 : 0;
+ port.send_n(str, size);
+ }
+
+ port.close();
+ return ret;
+}
+
+// TODO: This is a stand-in function that uses a struct pointer and size in
+// place of varargs. Once varargs support is added we will use that to
+// implement the real version.
+LLVM_LIBC_FUNCTION(int, rpc_fprintf,
+ (::FILE *__restrict stream, const char *__restrict format,
+ void *args, size_t size)) {
+ cpp::string_view str(format);
+ if (stream == stdout)
+ return fprintf_impl(stream, format, str.size() + 1,
+ args, size);
+ else if (stream == stderr)
+ return fprintf_impl(stream, format, str.size() + 1,
+ args, size);
+ else
+ return fprintf_impl(stream, format, str.size() + 1,
+ args, size);
+}
+
+} // namespace LIBC_NAMESPACE
diff --git a/libc/src/gpu/rpc_fprintf.h b/libc/src/gpu/rpc_fprintf.h
new file mode 100644
index 000000000000..053f7b4f818a
--- /dev/null
+++ b/libc/src/gpu/rpc_fprintf.h
@@ -0,0 +1,22 @@
+//===-- Implementation header for RPC functions -----------------*- C++ -*-===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef LLVM_LIBC_SRC_GPU_RPC_HOST_CALL_H
+#define LLVM_LIBC_SRC_GPU_RPC_HOST_CALL_H
+
+#include
+#include
+
+namespace LIBC_NAMESPACE {
+
+int rpc_fprintf(::FILE *__restrict stream, const char *__restrict format,
+ void *argc, size_t size);
+
+} // namespace LIBC_NAMESPACE
+
+#endif // LLVM_LIBC_SRC_GPU_RPC_HOST_CALL_H
diff --git a/libc/test/integration/src/stdio/CMakeLists.txt b/libc/test/integration/src/stdio/CMakeLists.txt
index 61caa2e5fa17..51c5ee25a6b2 100644
--- a/libc/test/integration/src/stdio/CMakeLists.txt
+++ b/libc/test/integration/src/stdio/CMakeLists.txt
@@ -1,3 +1,6 @@
+if(EXISTS ${CMAKE_CURRENT_SOURCE_DIR}/${LIBC_TARGET_OS})
+ add_subdirectory(${LIBC_TARGET_OS})
+endif()
add_custom_target(stdio-integration-tests)
add_dependencies(libc-integration-tests stdio-integration-tests)
diff --git a/libc/test/integration/src/stdio/gpu/CMakeLists.txt b/libc/test/integration/src/stdio/gpu/CMakeLists.txt
new file mode 100644
index 000000000000..6327c45e1ea5
--- /dev/null
+++ b/libc/test/integration/src/stdio/gpu/CMakeLists.txt
@@ -0,0 +1,21 @@
+add_custom_target(stdio-gpu-integration-tests)
+add_dependencies(libc-integration-tests stdio-gpu-integration-tests)
+
+# Create an output directory for any temporary test files.
+file(MAKE_DIRECTORY ${CMAKE_CURRENT_BINARY_DIR}/testdata)
+
+# These tests are not for correctness testing, but are instead a convenient way
+# to generate hermetic binaries for comparitive binary size testing.
+add_integration_test(
+ printf_test
+ SUITE
+ stdio-gpu-integration-tests
+ SRCS
+ printf.cpp
+ DEPENDS
+ libc.src.gpu.rpc_fprintf
+ libc.src.stdio.fopen
+ LOADER_ARGS
+ --threads 32
+ --blocks 4
+)
diff --git a/libc/test/integration/src/stdio/gpu/printf.cpp b/libc/test/integration/src/stdio/gpu/printf.cpp
new file mode 100644
index 000000000000..97ad4ace1dca
--- /dev/null
+++ b/libc/test/integration/src/stdio/gpu/printf.cpp
@@ -0,0 +1,88 @@
+//===-- RPC test to check args to printf ----------------------------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#include "test/IntegrationTest/test.h"
+
+#include "src/__support/GPU/utils.h"
+#include "src/gpu/rpc_fprintf.h"
+#include "src/stdio/fopen.h"
+
+using namespace LIBC_NAMESPACE;
+
+FILE *file = LIBC_NAMESPACE::fopen("testdata/test_data.txt", "w");
+
+TEST_MAIN(int argc, char **argv, char **envp) {
+ ASSERT_TRUE(file && "failed to open file");
+ // Check basic printing.
+ int written = 0;
+ written = LIBC_NAMESPACE::rpc_fprintf(file, "A simple string\n", nullptr, 0);
+ ASSERT_EQ(written, 16);
+
+ const char *str = "A simple string\n";
+ written = LIBC_NAMESPACE::rpc_fprintf(file, "%s", &str, sizeof(void *));
+ ASSERT_EQ(written, 16);
+
+ // Check printing a different value with each thread.
+ uint64_t thread_id = gpu::get_thread_id();
+ written = LIBC_NAMESPACE::rpc_fprintf(file, "%8ld\n", &thread_id,
+ sizeof(thread_id));
+ ASSERT_EQ(written, 9);
+
+ struct {
+ uint32_t x = 1;
+ char c = 'c';
+ double f = 1.0;
+ } args1;
+ written =
+ LIBC_NAMESPACE::rpc_fprintf(file, "%d%c%.1f\n", &args1, sizeof(args1));
+ ASSERT_EQ(written, 6);
+
+ struct {
+ uint32_t x = 1;
+ const char *str = "A simple string\n";
+ } args2;
+ written =
+ LIBC_NAMESPACE::rpc_fprintf(file, "%032b%s\n", &args2, sizeof(args2));
+ ASSERT_EQ(written, 49);
+
+ // Check that the server correctly handles divergent numbers of arguments.
+ const char *format = gpu::get_thread_id() % 2 ? "%s" : "%20ld\n";
+ written = LIBC_NAMESPACE::rpc_fprintf(file, format, &str, sizeof(void *));
+ ASSERT_EQ(written, gpu::get_thread_id() % 2 ? 16 : 21);
+
+ format = gpu::get_thread_id() % 2 ? "%s" : str;
+ written = LIBC_NAMESPACE::rpc_fprintf(file, format, &str, sizeof(void *));
+ ASSERT_EQ(written, 16);
+
+ // Check that we handle null arguments correctly.
+ struct {
+ void *null = nullptr;
+ } args3;
+ written = LIBC_NAMESPACE::rpc_fprintf(file, "%p", &args3, sizeof(args3));
+ ASSERT_EQ(written, 9);
+
+#ifndef LIBC_COPT_PRINTF_NO_NULLPTR_CHECKS
+ written = LIBC_NAMESPACE::rpc_fprintf(file, "%s", &args3, sizeof(args3));
+ ASSERT_EQ(written, 6);
+#endif // LIBC_COPT_PRINTF_NO_NULLPTR_CHECKS
+
+ // Check for extremely abused variable width arguments
+ struct {
+ uint32_t x = 1;
+ uint32_t y = 2;
+ double f = 1.0;
+ } args4;
+ written = LIBC_NAMESPACE::rpc_fprintf(file, "%**d", &args4, sizeof(args4));
+ ASSERT_EQ(written, 4);
+ written = LIBC_NAMESPACE::rpc_fprintf(file, "%**d%6d", &args4, sizeof(args4));
+ ASSERT_EQ(written, 10);
+ written = LIBC_NAMESPACE::rpc_fprintf(file, "%**.**f", &args4, sizeof(args4));
+ ASSERT_EQ(written, 7);
+
+ return 0;
+}
diff --git a/libc/utils/gpu/server/CMakeLists.txt b/libc/utils/gpu/server/CMakeLists.txt
index 6fca72cfae95..94347ef39497 100644
--- a/libc/utils/gpu/server/CMakeLists.txt
+++ b/libc/utils/gpu/server/CMakeLists.txt
@@ -1,4 +1,8 @@
-add_library(llvmlibc_rpc_server STATIC rpc_server.cpp)
+add_library(llvmlibc_rpc_server STATIC
+ ${LIBC_SOURCE_DIR}/src/stdio/printf_core/writer.cpp
+ ${LIBC_SOURCE_DIR}/src/stdio/printf_core/converter.cpp
+ rpc_server.cpp
+)
# Include the RPC implemenation from libc.
target_include_directories(llvmlibc_rpc_server PRIVATE ${LIBC_SOURCE_DIR})
@@ -9,6 +13,10 @@ target_include_directories(llvmlibc_rpc_server PUBLIC ${CMAKE_CURRENT_SOURCE_DIR
target_compile_options(llvmlibc_rpc_server PUBLIC
$<$:-Wno-attributes>)
target_compile_definitions(llvmlibc_rpc_server PUBLIC
+ LIBC_COPT_USE_C_ASSERT
+ LIBC_COPT_ARRAY_ARG_LIST
+ LIBC_COPT_PRINTF_DISABLE_WRITE_INT
+ LIBC_COPT_PRINTF_DISABLE_INDEX_MODE
LIBC_NAMESPACE=${LIBC_NAMESPACE})
# Install the server and associated header.
diff --git a/libc/utils/gpu/server/rpc_server.cpp b/libc/utils/gpu/server/rpc_server.cpp
index fd306642fdcc..095f3fa13ffa 100644
--- a/libc/utils/gpu/server/rpc_server.cpp
+++ b/libc/utils/gpu/server/rpc_server.cpp
@@ -14,7 +14,13 @@
#include "llvmlibc_rpc_server.h"
#include "src/__support/RPC/rpc.h"
+#include "src/__support/arg_list.h"
+#include "src/stdio/printf_core/converter.h"
+#include "src/stdio/printf_core/parser.h"
+#include "src/stdio/printf_core/writer.h"
+
#include "src/stdio/gpu/file.h"
+#include
#include
#include
#include
@@ -25,6 +31,7 @@
#include
using namespace LIBC_NAMESPACE;
+using namespace LIBC_NAMESPACE::printf_core;
static_assert(sizeof(rpc_buffer_t) == sizeof(rpc::Buffer),
"Buffer size mismatch");
@@ -32,6 +39,141 @@ static_assert(sizeof(rpc_buffer_t) == sizeof(rpc::Buffer),
static_assert(RPC_MAXIMUM_PORT_COUNT == rpc::MAX_PORT_COUNT,
"Incorrect maximum port count");
+template void handle_printf(rpc::Server::Port &port) {
+ FILE *files[lane_size] = {nullptr};
+ // Get the appropriate output stream to use.
+ if (port.get_opcode() == RPC_PRINTF_TO_STREAM)
+ port.recv([&](rpc::Buffer *buffer, uint32_t id) {
+ files[id] = reinterpret_cast(buffer->data[0]);
+ });
+ else if (port.get_opcode() == RPC_PRINTF_TO_STDOUT)
+ std::fill(files, files + lane_size, stdout);
+ else
+ std::fill(files, files + lane_size, stderr);
+
+ uint64_t format_sizes[lane_size] = {0};
+ void *format[lane_size] = {nullptr};
+
+ uint64_t args_sizes[lane_size] = {0};
+ void *args[lane_size] = {nullptr};
+
+ // Recieve the format string and arguments from the client.
+ port.recv_n(format, format_sizes,
+ [&](uint64_t size) { return new char[size]; });
+ port.recv_n(args, args_sizes, [&](uint64_t size) { return new char[size]; });
+
+ // Identify any arguments that are actually pointers to strings on the client.
+ // Additionally we want to determine how much buffer space we need to print.
+ std::vector strs_to_copy[lane_size];
+ int buffer_size[lane_size] = {0};
+ for (uint32_t lane = 0; lane < lane_size; ++lane) {
+ if (!format[lane])
+ continue;
+
+ WriteBuffer wb(nullptr, 0);
+ Writer writer(&wb);
+
+ internal::StructArgList printf_args(args[lane], args_sizes[lane]);
+ Parser parser(
+ reinterpret_cast(format[lane]), printf_args);
+
+ for (FormatSection cur_section = parser.get_next_section();
+ !cur_section.raw_string.empty();
+ cur_section = parser.get_next_section()) {
+ if (cur_section.has_conv && cur_section.conv_name == 's' &&
+ cur_section.conv_val_ptr) {
+ strs_to_copy[lane].emplace_back(cur_section.conv_val_ptr);
+ } else if (cur_section.has_conv) {
+ // Ignore conversion errors for the first pass.
+ convert(&writer, cur_section);
+ } else {
+ writer.write(cur_section.raw_string);
+ }
+ }
+ buffer_size[lane] = writer.get_chars_written();
+ }
+
+ // Recieve any strings from the client and push them into a buffer.
+ std::vector copied_strs[lane_size];
+ while (std::any_of(std::begin(strs_to_copy), std::end(strs_to_copy),
+ [](const auto &v) { return !v.empty() && v.back(); })) {
+ port.send([&](rpc::Buffer *buffer, uint32_t id) {
+ void *ptr = !strs_to_copy[id].empty() ? strs_to_copy[id].back() : nullptr;
+ buffer->data[1] = reinterpret_cast(ptr);
+ if (!strs_to_copy[id].empty())
+ strs_to_copy[id].pop_back();
+ });
+ uint64_t str_sizes[lane_size] = {0};
+ void *strs[lane_size] = {nullptr};
+ port.recv_n(strs, str_sizes, [](uint64_t size) { return new char[size]; });
+ for (uint32_t lane = 0; lane < lane_size; ++lane) {
+ if (!strs[lane])
+ continue;
+
+ copied_strs[lane].emplace_back(strs[lane]);
+ buffer_size[lane] += str_sizes[lane];
+ }
+ }
+
+ // Perform the final formatting and printing using the LLVM C library printf.
+ int results[lane_size] = {0};
+ std::vector to_be_deleted;
+ for (uint32_t lane = 0; lane < lane_size; ++lane) {
+ if (!format[lane])
+ continue;
+
+ std::unique_ptr buffer(new char[buffer_size[lane]]);
+ WriteBuffer wb(buffer.get(), buffer_size[lane]);
+ Writer writer(&wb);
+
+ internal::StructArgList printf_args(args[lane], args_sizes[lane]);
+ Parser parser(
+ reinterpret_cast(format[lane]), printf_args);
+
+ // Parse and print the format string using the arguments we copied from
+ // the client.
+ int ret = 0;
+ for (FormatSection cur_section = parser.get_next_section();
+ !cur_section.raw_string.empty();
+ cur_section = parser.get_next_section()) {
+ // If this argument was a string we use the memory buffer we copied from
+ // the client by replacing the raw pointer with the copied one.
+ if (cur_section.has_conv && cur_section.conv_name == 's') {
+ if (!copied_strs[lane].empty()) {
+ cur_section.conv_val_ptr = copied_strs[lane].back();
+ to_be_deleted.push_back(copied_strs[lane].back());
+ copied_strs[lane].pop_back();
+ } else {
+ cur_section.conv_val_ptr = nullptr;
+ }
+ }
+ if (cur_section.has_conv) {
+ ret = convert(&writer, cur_section);
+ if (ret == -1)
+ break;
+ } else {
+ writer.write(cur_section.raw_string);
+ }
+ }
+
+ results[lane] =
+ fwrite(buffer.get(), 1, writer.get_chars_written(), files[lane]);
+ if (results[lane] != writer.get_chars_written() || ret == -1)
+ results[lane] = -1;
+ }
+
+ // Send the final return value and signal completion by setting the string
+ // argument to null.
+ port.send([&](rpc::Buffer *buffer, uint32_t id) {
+ buffer->data[0] = static_cast(results[id]);
+ buffer->data[1] = reinterpret_cast(nullptr);
+ delete[] reinterpret_cast(format[id]);
+ delete[] reinterpret_cast(args[id]);
+ });
+ for (void *ptr : to_be_deleted)
+ delete[] reinterpret_cast(ptr);
+}
+
template
rpc_status_t handle_server_impl(
rpc::Server &server,
@@ -195,6 +337,12 @@ rpc_status_t handle_server_impl(
});
break;
}
+ case RPC_PRINTF_TO_STREAM:
+ case RPC_PRINTF_TO_STDOUT:
+ case RPC_PRINTF_TO_STDERR: {
+ handle_printf(*port);
+ break;
+ }
case RPC_NOOP: {
port->recv([](rpc::Buffer *) {});
break;
--
GitLab
From 5029949952f4dc745dcb7799c7449a02fe8309c3 Mon Sep 17 00:00:00 2001
From: Daniil Kovalev
Date: Wed, 3 Apr 2024 00:28:47 +0300
Subject: [PATCH 0013/1683] [PAC][llvm-readobj][AArch64][ELF] Support
`GNU_PROPERTY_AARCH64_FEATURE_PAUTH` (#85231)
This adds support for `GNU_PROPERTY_AARCH64_FEATURE_PAUTH` feature (as
defined in https://github.com/ARM-software/abi-aa/pull/240) handling in
llvm-readobj and llvm-readelf. The following constants for supported
platforms are also introduced:
- `AARCH64_PAUTH_PLATFORM_INVALID = 0x0`
- `AARCH64_PAUTH_PLATFORM_BAREMETAL = 0x1`
- `AARCH64_PAUTH_PLATFORM_LLVM_LINUX = 0x10000002`
For the llvm_linux platform, output of the tools contains descriptions
of PAuth features which are enabled/disabled depending on the version
value. Version value bits correspond to the following `LangOptions`
defined in #85232:
- bit 0: `PointerAuthIntrinsics`;
- bit 1: `PointerAuthCalls`;
- bit 2: `PointerAuthReturns`;
- bit 3: `PointerAuthAuthTraps`;
- bit 4: `PointerAuthVTPtrAddressDiscrimination`;
- bit 5: `PointerAuthVTPtrTypeDiscrimination`;
- bit 6: `PointerAuthInitFini`.
Support for `.note.AARCH64-PAUTH-ABI-tag` is dropped since it's deleted
from the spec in ARM-software/abi-aa#250.
---
llvm/include/llvm/BinaryFormat/ELF.h | 26 +-
.../ELF/AArch64/aarch64-feature-pauth.s | 305 ++++++++++++------
.../ELF/AArch64/aarch64-note-gnu-property.s | 2 +
llvm/tools/llvm-readobj/ELFDumper.cpp | 127 ++++----
4 files changed, 303 insertions(+), 157 deletions(-)
diff --git a/llvm/include/llvm/BinaryFormat/ELF.h b/llvm/include/llvm/BinaryFormat/ELF.h
index 877f3f7862c8..ed267c100b76 100644
--- a/llvm/include/llvm/BinaryFormat/ELF.h
+++ b/llvm/include/llvm/BinaryFormat/ELF.h
@@ -1712,11 +1712,6 @@ enum {
NT_ANDROID_TYPE_MEMTAG = 4,
};
-// ARM note types.
-enum {
- NT_ARM_TYPE_PAUTH_ABI_TAG = 1,
-};
-
// Memory tagging values used in NT_ANDROID_TYPE_MEMTAG notes.
enum {
// Enumeration to determine the tagging mode. In Android-land, 'SYNC' means
@@ -1740,6 +1735,7 @@ enum : unsigned {
GNU_PROPERTY_STACK_SIZE = 1,
GNU_PROPERTY_NO_COPY_ON_PROTECTED = 2,
GNU_PROPERTY_AARCH64_FEATURE_1_AND = 0xc0000000,
+ GNU_PROPERTY_AARCH64_FEATURE_PAUTH = 0xc0000001,
GNU_PROPERTY_X86_FEATURE_1_AND = 0xc0000002,
GNU_PROPERTY_X86_UINT32_OR_LO = 0xc0008000,
@@ -1758,6 +1754,26 @@ enum : unsigned {
GNU_PROPERTY_AARCH64_FEATURE_1_GCS = 1 << 2,
};
+// aarch64 PAuth platforms.
+enum : unsigned {
+ AARCH64_PAUTH_PLATFORM_INVALID = 0x0,
+ AARCH64_PAUTH_PLATFORM_BAREMETAL = 0x1,
+ AARCH64_PAUTH_PLATFORM_LLVM_LINUX = 0x10000002,
+};
+
+// Bit positions of version flags for AARCH64_PAUTH_PLATFORM_LLVM_LINUX.
+enum : unsigned {
+ AARCH64_PAUTH_PLATFORM_LLVM_LINUX_VERSION_INTRINSICS = 0,
+ AARCH64_PAUTH_PLATFORM_LLVM_LINUX_VERSION_CALLS = 1,
+ AARCH64_PAUTH_PLATFORM_LLVM_LINUX_VERSION_RETURNS = 2,
+ AARCH64_PAUTH_PLATFORM_LLVM_LINUX_VERSION_AUTHTRAPS = 3,
+ AARCH64_PAUTH_PLATFORM_LLVM_LINUX_VERSION_VPTRADDRDISCR = 4,
+ AARCH64_PAUTH_PLATFORM_LLVM_LINUX_VERSION_VPTRTYPEDISCR = 5,
+ AARCH64_PAUTH_PLATFORM_LLVM_LINUX_VERSION_INITFINI = 6,
+ AARCH64_PAUTH_PLATFORM_LLVM_LINUX_VERSION_LAST =
+ AARCH64_PAUTH_PLATFORM_LLVM_LINUX_VERSION_INITFINI,
+};
+
// x86 processor feature bits.
enum : unsigned {
GNU_PROPERTY_X86_FEATURE_1_IBT = 1 << 0,
diff --git a/llvm/test/tools/llvm-readobj/ELF/AArch64/aarch64-feature-pauth.s b/llvm/test/tools/llvm-readobj/ELF/AArch64/aarch64-feature-pauth.s
index f28d92eae857..24918e883d9b 100644
--- a/llvm/test/tools/llvm-readobj/ELF/AArch64/aarch64-feature-pauth.s
+++ b/llvm/test/tools/llvm-readobj/ELF/AArch64/aarch64-feature-pauth.s
@@ -1,98 +1,211 @@
# RUN: rm -rf %t && split-file %s %t && cd %t
-# RUN: llvm-mc -filetype=obj -triple=aarch64-linux-gnu abi-tag.s -o tag.o
-# RUN: llvm-mc -filetype=obj -triple=aarch64-linux-gnu abi-tag-short.s -o tag-short.o
-# RUN: llvm-mc -filetype=obj -triple=aarch64-linux-gnu abi-tag-long.s -o tag-long.o
-
-# RUN: llvm-readelf --notes tag.o | FileCheck --check-prefix NORMAL %s
-# RUN: llvm-readelf --notes tag-short.o | FileCheck --check-prefix SHORT %s
-# RUN: llvm-readelf --notes tag-long.o | FileCheck --check-prefix LONG %s
-
-# NORMAL: AArch64 PAuth ABI tag: platform 0x2a, version 0x1
-# SHORT: AArch64 PAuth ABI tag:
-# LONG: AArch64 PAuth ABI tag: platform 0x2a, version 0x1, additional info 0xEFCDAB8967452301
-
-# RUN: llvm-readobj --notes tag.o | FileCheck --check-prefix LLVM-NORMAL %s
-# RUN: llvm-readobj --notes tag-short.o | FileCheck --check-prefix LLVM-SHORT %s
-# RUN: llvm-readobj --notes tag-long.o | FileCheck --check-prefix LLVM-LONG %s
-
-# LLVM-SHORT: Notes [
-# LLVM-SHORT-NEXT: NoteSection {
-# LLVM-SHORT-NEXT: Name: .note.AARCH64-PAUTH-ABI-tag
-# LLVM-SHORT-NEXT: Offset: 0x40
-# LLVM-SHORT-NEXT: Size: 0x1C
-# LLVM-SHORT-NEXT: Note {
-# LLVM-SHORT-NEXT: Owner: ARM
-# LLVM-SHORT-NEXT: Data size: 0xC
-# LLVM-SHORT-NEXT: Type: NT_ARM_TYPE_PAUTH_ABI_TAG
-# LLVM-SHORT-NEXT: Description data (
-# LLVM-SHORT-NEXT: 0000: 2A000000 00000000 01000000
-# LLVM-SHORT-NEXT: )
-# LLVM-SHORT-NEXT: }
-# LLVM-SHORT-NEXT: }
-# LLVM-SHORT-NEXT: ]
-
-# LLVM-NORMAL: Notes [
-# LLVM-NORMAL-NEXT: NoteSection {
-# LLVM-NORMAL-NEXT: Name: .note.AARCH64-PAUTH-ABI-tag
-# LLVM-NORMAL-NEXT: Offset: 0x40
-# LLVM-NORMAL-NEXT: Size: 0x20
-# LLVM-NORMAL-NEXT: Note {
-# LLVM-NORMAL-NEXT: Owner: ARM
-# LLVM-NORMAL-NEXT: Data size: 0x10
-# LLVM-NORMAL-NEXT: Type: NT_ARM_TYPE_PAUTH_ABI_TAG
-# LLVM-NORMAL-NEXT: Platform: 42
-# LLVM-NORMAL-NEXT: Version: 1
-# LLVM-NORMAL-NEXT: }
-# LLVM-NORMAL-NEXT: }
-# LLVM-NORMAL-NEXT: ]
-
-# LLVM-LONG: Notes [
-# LLVM-LONG-NEXT: NoteSection {
-# LLVM-LONG-NEXT: Name: .note.AARCH64-PAUTH-ABI-tag
-# LLVM-LONG-NEXT: Offset: 0x40
-# LLVM-LONG-NEXT: Size: 0x28
-# LLVM-LONG-NEXT: Note {
-# LLVM-LONG-NEXT: Owner: ARM
-# LLVM-LONG-NEXT: Data size: 0x18
-# LLVM-LONG-NEXT: Type: NT_ARM_TYPE_PAUTH_ABI_TAG
-# LLVM-LONG-NEXT: Platform: 42
-# LLVM-LONG-NEXT: Version: 1
-# LLVM-LONG-NEXT: Additional info: EFCDAB8967452301
-# LLVM-LONG-NEXT: }
-# LLVM-LONG-NEXT: }
-# LLVM-LONG-NEXT: ]
-
-#--- abi-tag.s
-
-.section ".note.AARCH64-PAUTH-ABI-tag", "a"
-.long 4
-.long 16
-.long 1
-.asciz "ARM"
-
-.quad 42 // platform
-.quad 1 // version
-
-#--- abi-tag-short.s
-
-.section ".note.AARCH64-PAUTH-ABI-tag", "a"
-.long 4
-.long 12
-.long 1
-.asciz "ARM"
-
-.quad 42
-.word 1
-
-#--- abi-tag-long.s
-
-.section ".note.AARCH64-PAUTH-ABI-tag", "a"
-.long 4
-.long 24
-.long 1
-.asciz "ARM"
-
-.quad 42 // platform
-.quad 1 // version
-.quad 0x0123456789ABCDEF // extra data
+#--- gnu-42-1.s
+
+.section ".note.gnu.property", "a"
+ .long 4 // Name length is always 4 ("GNU")
+ .long end - begin // Data length
+ .long 5 // Type: NT_GNU_PROPERTY_TYPE_0
+ .asciz "GNU" // Name
+ .p2align 3
+begin:
+ # PAuth ABI property note
+ .long 0xc0000001 // Type: GNU_PROPERTY_AARCH64_FEATURE_PAUTH
+ .long 16 // Data size
+ .quad 42 // PAuth ABI platform
+ .quad 1 // PAuth ABI version
+ .p2align 3 // Align to 8 byte for 64 bit
+end:
+
+# RUN: llvm-mc -filetype=obj -triple aarch64-linux-gnu gnu-42-1.s -o gnu-42-1.o
+# RUN: llvm-readelf --notes gnu-42-1.o | \
+# RUN: FileCheck --check-prefix=ELF -DPLATFORM="0x2a (unknown)" -DVERSION=0x1 %s
+# RUN: llvm-readobj --notes gnu-42-1.o | \
+# RUN: FileCheck --check-prefix=OBJ -DPLATFORM="0x2a (unknown)" -DVERSION=0x1 %s
+
+# ELF: Displaying notes found in: .note.gnu.property
+# ELF-NEXT: Owner Data size Description
+# ELF-NEXT: GNU 0x00000018 NT_GNU_PROPERTY_TYPE_0 (property note)
+# ELF-NEXT: AArch64 PAuth ABI core info: platform [[PLATFORM]], version [[VERSION]]
+
+# OBJ: Notes [
+# OBJ-NEXT: NoteSection {
+# OBJ-NEXT: Name: .note.gnu.property
+# OBJ-NEXT: Offset: 0x40
+# OBJ-NEXT: Size: 0x28
+# OBJ-NEXT: Note {
+# OBJ-NEXT: Owner: GNU
+# OBJ-NEXT: Data size: 0x18
+# OBJ-NEXT: Type: NT_GNU_PROPERTY_TYPE_0 (property note)
+# OBJ-NEXT: Property [
+# OBJ-NEXT: AArch64 PAuth ABI core info: platform [[PLATFORM]], version [[VERSION]]
+# OBJ-NEXT: ]
+# OBJ-NEXT: }
+# OBJ-NEXT: }
+# OBJ-NEXT: ]
+
+#--- gnu-0-0.s
+
+.section ".note.gnu.property", "a"
+ .long 4 // Name length is always 4 ("GNU")
+ .long end - begin // Data length
+ .long 5 // Type: NT_GNU_PROPERTY_TYPE_0
+ .asciz "GNU" // Name
+ .p2align 3
+begin:
+ # PAuth ABI property note
+ .long 0xc0000001 // Type: GNU_PROPERTY_AARCH64_FEATURE_PAUTH
+ .long 16 // Data size
+ .quad 0 // PAuth ABI platform
+ .quad 0 // PAuth ABI version
+ .p2align 3 // Align to 8 byte for 64 bit
+end:
+
+# RUN: llvm-mc -filetype=obj -triple aarch64-linux-gnu gnu-0-0.s -o gnu-0-0.o
+# RUN: llvm-readelf --notes gnu-0-0.o | \
+# RUN: FileCheck --check-prefix=ELF -DPLATFORM="0x0 (invalid)" -DVERSION=0x0 %s
+# RUN: llvm-readobj --notes gnu-0-0.o | \
+# RUN: FileCheck --check-prefix=OBJ -DPLATFORM="0x0 (invalid)" -DVERSION=0x0 %s
+
+#--- gnu-1-0.s
+
+.section ".note.gnu.property", "a"
+ .long 4 // Name length is always 4 ("GNU")
+ .long end - begin // Data length
+ .long 5 // Type: NT_GNU_PROPERTY_TYPE_0
+ .asciz "GNU" // Name
+ .p2align 3
+begin:
+ # PAuth ABI property note
+ .long 0xc0000001 // Type: GNU_PROPERTY_AARCH64_FEATURE_PAUTH
+ .long 16 // Data size
+ .quad 1 // PAuth ABI platform
+ .quad 0 // PAuth ABI version
+ .p2align 3 // Align to 8 byte for 64 bit
+end:
+
+# RUN: llvm-mc -filetype=obj -triple aarch64-linux-gnu gnu-1-0.s -o gnu-1-0.o
+# RUN: llvm-readelf --notes gnu-1-0.o | \
+# RUN: FileCheck --check-prefix=ELF -DPLATFORM="0x1 (baremetal)" -DVERSION=0x0 %s
+# RUN: llvm-readobj --notes gnu-1-0.o | \
+# RUN: FileCheck --check-prefix=OBJ -DPLATFORM="0x1 (baremetal)" -DVERSION=0x0 %s
+
+#--- gnu-0x10000002-85.s
+
+.section ".note.gnu.property", "a"
+ .long 4 // Name length is always 4 ("GNU")
+ .long end - begin // Data length
+ .long 5 // Type: NT_GNU_PROPERTY_TYPE_0
+ .asciz "GNU" // Name
+ .p2align 3
+begin:
+ # PAuth ABI property note
+ .long 0xc0000001 // Type: GNU_PROPERTY_AARCH64_FEATURE_PAUTH
+ .long 16 // Data size
+ .quad 0x10000002 // PAuth ABI platform
+ .quad 85 // PAuth ABI version
+ .p2align 3 // Align to 8 byte for 64 bit
+end:
+
+# RUN: llvm-mc -filetype=obj -triple aarch64-linux-gnu gnu-0x10000002-85.s -o gnu-0x10000002-85.o
+# RUN: llvm-readelf --notes gnu-0x10000002-85.o | \
+# RUN: FileCheck --check-prefix=ELF -DPLATFORM="0x10000002 (llvm_linux)" \
+# RUN: -DVERSION="0x55 (PointerAuthIntrinsics, !PointerAuthCalls, PointerAuthReturns, !PointerAuthAuthTraps, PointerAuthVTPtrAddressDiscrimination, !PointerAuthVTPtrTypeDiscrimination, PointerAuthInitFini)" %s
+# RUN: llvm-readobj --notes gnu-0x10000002-85.o | \
+# RUN: FileCheck --check-prefix=OBJ -DPLATFORM="0x10000002 (llvm_linux)" \
+# RUN: -DVERSION="0x55 (PointerAuthIntrinsics, !PointerAuthCalls, PointerAuthReturns, !PointerAuthAuthTraps, PointerAuthVTPtrAddressDiscrimination, !PointerAuthVTPtrTypeDiscrimination, PointerAuthInitFini)" %s
+
+#--- gnu-0x10000002-128.s
+
+.section ".note.gnu.property", "a"
+ .long 4 // Name length is always 4 ("GNU")
+ .long end - begin // Data length
+ .long 5 // Type: NT_GNU_PROPERTY_TYPE_0
+ .asciz "GNU" // Name
+ .p2align 3
+begin:
+ # PAuth ABI property note
+ .long 0xc0000001 // Type: GNU_PROPERTY_AARCH64_FEATURE_PAUTH
+ .long 16 // Data size
+ .quad 0x10000002 // PAuth ABI platform
+ .quad 128 // PAuth ABI version
+ .p2align 3 // Align to 8 byte for 64 bit
+end:
+
+# RUN: llvm-mc -filetype=obj -triple aarch64-linux-gnu gnu-0x10000002-128.s -o gnu-0x10000002-128.o
+# RUN: llvm-readelf --notes gnu-0x10000002-128.o | \
+# RUN: FileCheck --check-prefix=ELF -DPLATFORM="0x10000002 (llvm_linux)" -DVERSION="0x80 (unknown)" %s
+# RUN: llvm-readobj --notes gnu-0x10000002-128.o | \
+# RUN: FileCheck --check-prefix=OBJ -DPLATFORM="0x10000002 (llvm_linux)" -DVERSION="0x80 (unknown)" %s
+
+#--- gnu-short.s
+
+.section ".note.gnu.property", "a"
+ .long 4 // Name length is always 4 ("GNU")
+ .long end - begin // Data length
+ .long 5 // Type: NT_GNU_PROPERTY_TYPE_0
+ .asciz "GNU" // Name
+ .p2align 3
+begin:
+ # PAuth ABI property note
+ .long 0xc0000001 // Type: GNU_PROPERTY_AARCH64_FEATURE_PAUTH
+ .long 12 // Data size
+ .quad 42 // PAuth ABI platform
+ .word 1 // PAuth ABI version
+ .p2align 3 // Align to 8 byte for 64 bit
+end:
+
+# RUN: llvm-mc -filetype=obj -triple aarch64-linux-gnu gnu-short.s -o gnu-short.o
+# RUN: llvm-readelf --notes gnu-short.o | \
+# RUN: FileCheck --check-prefix=ELF-ERR -DSIZE=28 -DDATASIZE=18 \
+# RUN: -DERR="" %s
+# RUN: llvm-readobj --notes gnu-short.o | \
+# RUN: FileCheck --check-prefix=OBJ-ERR -DSIZE=28 -DDATASIZE=18 \
+# RUN: -DERR="" %s
+
+# ELF-ERR: Displaying notes found in: .note.gnu.property
+# ELF-ERR-NEXT: Owner Data size Description
+# ELF-ERR-NEXT: GNU 0x000000[[DATASIZE]] NT_GNU_PROPERTY_TYPE_0 (property note)
+# ELF-ERR-NEXT: AArch64 PAuth ABI core info: [[ERR]]
+
+# OBJ-ERR: Notes [
+# OBJ-ERR-NEXT: NoteSection {
+# OBJ-ERR-NEXT: Name: .note.gnu.property
+# OBJ-ERR-NEXT: Offset: 0x40
+# OBJ-ERR-NEXT: Size: 0x[[SIZE]]
+# OBJ-ERR-NEXT: Note {
+# OBJ-ERR-NEXT: Owner: GNU
+# OBJ-ERR-NEXT: Data size: 0x[[DATASIZE]]
+# OBJ-ERR-NEXT: Type: NT_GNU_PROPERTY_TYPE_0 (property note)
+# OBJ-ERR-NEXT: Property [
+# OBJ-ERR-NEXT: AArch64 PAuth ABI core info: [[ERR]]
+# OBJ-ERR-NEXT: ]
+# OBJ-ERR-NEXT: }
+# OBJ-ERR-NEXT: }
+# OBJ-ERR-NEXT: ]
+
+#--- gnu-long.s
+
+.section ".note.gnu.property", "a"
+ .long 4 // Name length is always 4 ("GNU")
+ .long end - begin // Data length
+ .long 5 // Type: NT_GNU_PROPERTY_TYPE_0
+ .asciz "GNU" // Name
+ .p2align 3
+begin:
+ # PAuth ABI property note
+ .long 0xc0000001 // Type: GNU_PROPERTY_AARCH64_FEATURE_PAUTH
+ .long 24 // Data size
+ .quad 42 // PAuth ABI platform
+ .quad 1 // PAuth ABI version
+ .quad 0x0123456789ABCDEF
+ .p2align 3 // Align to 8 byte for 64 bit
+end:
+
+# RUN: llvm-mc -filetype=obj -triple aarch64-linux-gnu gnu-long.s -o gnu-long.o
+# RUN: llvm-readelf --notes gnu-long.o | \
+# RUN: FileCheck --check-prefix=ELF-ERR -DSIZE=30 -DDATASIZE=20 \
+# RUN: -DERR="" %s
+# RUN: llvm-readobj --notes gnu-long.o | \
+# RUN: FileCheck --check-prefix=OBJ-ERR -DSIZE=30 -DDATASIZE=20 \
+# RUN: -DERR="" %s
diff --git a/llvm/test/tools/llvm-readobj/ELF/AArch64/aarch64-note-gnu-property.s b/llvm/test/tools/llvm-readobj/ELF/AArch64/aarch64-note-gnu-property.s
index 377e6f93448c..b517f0b38155 100644
--- a/llvm/test/tools/llvm-readobj/ELF/AArch64/aarch64-note-gnu-property.s
+++ b/llvm/test/tools/llvm-readobj/ELF/AArch64/aarch64-note-gnu-property.s
@@ -1,3 +1,5 @@
+// See tests for GNU_PROPERTY_AARCH64_FEATURE_PAUTH in aarch64-feature-pauth.s
+
// RUN: llvm-mc -filetype=obj -triple aarch64-linux-gnu %s -o %t
// RUN: llvm-readelf --notes %t | FileCheck %s --check-prefix=GNU
// RUN: llvm-readobj --notes %t | FileCheck %s --check-prefix=LLVM
diff --git a/llvm/tools/llvm-readobj/ELFDumper.cpp b/llvm/tools/llvm-readobj/ELFDumper.cpp
index 4b406ef12aec..29a032557642 100644
--- a/llvm/tools/llvm-readobj/ELFDumper.cpp
+++ b/llvm/tools/llvm-readobj/ELFDumper.cpp
@@ -61,6 +61,7 @@
#include "llvm/Support/SystemZ/zOSSupport.h"
#include "llvm/Support/raw_ostream.h"
#include
+#include
#include
#include
#include
@@ -5105,6 +5106,73 @@ template void GNUELFDumper::printAddrsig() {
}
}
+template
+static bool printAArch64PAuthABICoreInfo(raw_ostream &OS, uint32_t DataSize,
+ ArrayRef Desc) {
+ OS << " AArch64 PAuth ABI core info: ";
+ // DataSize - size without padding, Desc.size() - size with padding
+ if (DataSize != 16) {
+ OS << format("", DataSize);
+ return false;
+ }
+
+ uint64_t Platform =
+ support::endian::read64(Desc.data() + 0);
+ uint64_t Version = support::endian::read64(Desc.data() + 8);
+
+ const char *PlatformDesc = [Platform]() {
+ switch (Platform) {
+ case AARCH64_PAUTH_PLATFORM_INVALID:
+ return "invalid";
+ case AARCH64_PAUTH_PLATFORM_BAREMETAL:
+ return "baremetal";
+ case AARCH64_PAUTH_PLATFORM_LLVM_LINUX:
+ return "llvm_linux";
+ default:
+ return "unknown";
+ }
+ }();
+
+ std::string VersionDesc = [Platform, Version]() -> std::string {
+ if (Platform != AARCH64_PAUTH_PLATFORM_LLVM_LINUX)
+ return "";
+ if (Version >= (1 << (AARCH64_PAUTH_PLATFORM_LLVM_LINUX_VERSION_LAST + 1)))
+ return "unknown";
+
+ std::array
+ Flags;
+ Flags[AARCH64_PAUTH_PLATFORM_LLVM_LINUX_VERSION_INTRINSICS] = "Intrinsics";
+ Flags[AARCH64_PAUTH_PLATFORM_LLVM_LINUX_VERSION_CALLS] = "Calls";
+ Flags[AARCH64_PAUTH_PLATFORM_LLVM_LINUX_VERSION_RETURNS] = "Returns";
+ Flags[AARCH64_PAUTH_PLATFORM_LLVM_LINUX_VERSION_AUTHTRAPS] = "AuthTraps";
+ Flags[AARCH64_PAUTH_PLATFORM_LLVM_LINUX_VERSION_VPTRADDRDISCR] =
+ "VTPtrAddressDiscrimination";
+ Flags[AARCH64_PAUTH_PLATFORM_LLVM_LINUX_VERSION_VPTRTYPEDISCR] =
+ "VTPtrTypeDiscrimination";
+ Flags[AARCH64_PAUTH_PLATFORM_LLVM_LINUX_VERSION_INITFINI] = "InitFini";
+
+ static_assert(AARCH64_PAUTH_PLATFORM_LLVM_LINUX_VERSION_INITFINI ==
+ AARCH64_PAUTH_PLATFORM_LLVM_LINUX_VERSION_LAST,
+ "Update when new enum items are defined");
+
+ std::string Desc;
+ for (uint32_t I = 0; I < Flags.size(); ++I) {
+ if (!(Version & (1 << I)))
+ Desc += '!';
+ Desc += Twine("PointerAuth" + Flags[I] + ", ").str();
+ }
+ Desc.resize(Desc.size() - 2); // Trim last ", "
+ return Desc;
+ }();
+
+ OS << format("platform 0x%x (%s), version 0x%x", Platform, PlatformDesc,
+ Version);
+ if (!VersionDesc.empty())
+ OS << format(" (%s)", VersionDesc.c_str());
+
+ return true;
+}
+
template
static std::string getGNUProperty(uint32_t Type, uint32_t DataSize,
ArrayRef Data) {
@@ -5162,6 +5230,9 @@ static std::string getGNUProperty(uint32_t Type, uint32_t DataSize,
if (PrData)
OS << format("", PrData);
return OS.str();
+ case GNU_PROPERTY_AARCH64_FEATURE_PAUTH:
+ printAArch64PAuthABICoreInfo(OS, DataSize, Data);
+ return OS.str();
case GNU_PROPERTY_X86_FEATURE_2_NEEDED:
case GNU_PROPERTY_X86_FEATURE_2_USED:
OS << "x86 feature "
@@ -5363,29 +5434,6 @@ static bool printAndroidNote(raw_ostream &OS, uint32_t NoteType,
return true;
}
-template
-static bool printAArch64Note(raw_ostream &OS, uint32_t NoteType,
- ArrayRef Desc) {
- if (NoteType != NT_ARM_TYPE_PAUTH_ABI_TAG)
- return false;
-
- OS << " AArch64 PAuth ABI tag: ";
- if (Desc.size() < 16) {
- OS << format("", Desc.size());
- return false;
- }
-
- uint64_t Platform = endian::read64(Desc.data() + 0);
- uint64_t Version = endian::read64(Desc.data() + 8);
- OS << format("platform 0x%" PRIx64 ", version 0x%" PRIx64, Platform, Version);
-
- if (Desc.size() > 16)
- OS << ", additional info 0x"
- << toHex(ArrayRef(Desc.data() + 16, Desc.size() - 16));
-
- return true;
-}
-
template
void GNUELFDumper::printMemtag(
const ArrayRef> DynamicEntries,
@@ -5783,10 +5831,6 @@ const NoteType AndroidNoteTypes[] = {
"NT_ANDROID_TYPE_MEMTAG (Android memory tagging information)"},
};
-const NoteType ARMNoteTypes[] = {
- {ELF::NT_ARM_TYPE_PAUTH_ABI_TAG, "NT_ARM_TYPE_PAUTH_ABI_TAG"},
-};
-
const NoteType CoreNoteTypes[] = {
{ELF::NT_PRSTATUS, "NT_PRSTATUS (prstatus structure)"},
{ELF::NT_FPREGSET, "NT_FPREGSET (floating point registers)"},
@@ -5905,8 +5949,6 @@ StringRef getNoteTypeName(const typename ELFT::Note &Note, unsigned ELFType) {
return FindNote(LLVMOMPOFFLOADNoteTypes);
if (Name == "Android")
return FindNote(AndroidNoteTypes);
- if (Name == "ARM")
- return FindNote(ARMNoteTypes);
if (ELFType == ELF::ET_CORE)
return FindNote(CoreNoteTypes);
@@ -6062,9 +6104,6 @@ template void GNUELFDumper::printNotes() {
} else if (Name == "Android") {
if (printAndroidNote(OS, Type, Descriptor))
return Error::success();
- } else if (Name == "ARM") {
- if (printAArch64Note(OS, Type, Descriptor))
- return Error::success();
}
if (!Descriptor.empty()) {
OS << " description data:";
@@ -7702,27 +7741,6 @@ static bool printAndroidNoteLLVMStyle(uint32_t NoteType, ArrayRef Desc,
return true;
}
-template
-static bool printAarch64NoteLLVMStyle(uint32_t NoteType, ArrayRef Desc,
- ScopedPrinter &W) {
- if (NoteType != NT_ARM_TYPE_PAUTH_ABI_TAG)
- return false;
-
- if (Desc.size() < 16)
- return false;
-
- uint64_t platform = endian::read64(Desc.data() + 0);
- uint64_t version = endian::read64(Desc.data() + 8);
- W.printNumber("Platform", platform);
- W.printNumber("Version", version);
-
- if (Desc.size() > 16)
- W.printString("Additional info",
- toHex(ArrayRef(Desc.data() + 16, Desc.size() - 16)));
-
- return true;
-}
-
template
void LLVMELFDumper::printMemtag(
const ArrayRef> DynamicEntries,
@@ -7859,9 +7877,6 @@ template void LLVMELFDumper::printNotes() {
} else if (Name == "Android") {
if (printAndroidNoteLLVMStyle(Type, Descriptor, W))
return Error::success();
- } else if (Name == "ARM") {
- if (printAarch64NoteLLVMStyle(Type, Descriptor, W))
- return Error::success();
}
if (!Descriptor.empty()) {
W.printBinaryBlock("Description data", Descriptor);
--
GitLab
From 2cf8118e3aa60f406ec41e88bdd4304f39744e89 Mon Sep 17 00:00:00 2001
From: Joseph Huber
Date: Tue, 2 Apr 2024 16:36:36 -0500
Subject: [PATCH 0014/1683] [Libomptarget] Add RPC-based `printf`
implementation for OpenMP (#85638)
Summary:
This patch adds an implementation of `printf` that's provided by the GPU
C library runtime. This `pritnf` currently implemented using the same
wrapper handling that OpenMP sets up. This will be removed once we have
proper varargs support.
This `printf` differs from the one CUDA offers in that it is synchronous
and uses a finite size. Additionally we support pretty much every format
specifier except the `%n` option.
Depends on https://github.com/llvm/llvm-project/pull/85331
---
openmp/libomptarget/DeviceRTL/CMakeLists.txt | 5 +++
openmp/libomptarget/DeviceRTL/src/LibC.cpp | 13 +++++++
openmp/libomptarget/test/libc/printf.c | 36 ++++++++++++++++++++
3 files changed, 54 insertions(+)
create mode 100644 openmp/libomptarget/test/libc/printf.c
diff --git a/openmp/libomptarget/DeviceRTL/CMakeLists.txt b/openmp/libomptarget/DeviceRTL/CMakeLists.txt
index 2509f1276cce..2e7f28df24d6 100644
--- a/openmp/libomptarget/DeviceRTL/CMakeLists.txt
+++ b/openmp/libomptarget/DeviceRTL/CMakeLists.txt
@@ -122,6 +122,11 @@ set(clang_opt_flags -O3 -mllvm -openmp-opt-disable -DSHARED_SCRATCHPAD_SIZE=512
set(link_opt_flags -O3 -openmp-opt-disable -attributor-enable=module -vectorize-slp=false )
set(link_export_flag -passes=internalize -internalize-public-api-file=${source_directory}/exports)
+# If the user built with the GPU C library enabled we will use that instead.
+if(${LIBOMPTARGET_GPU_LIBC_SUPPORT})
+ list(APPEND clang_opt_flags -DOMPTARGET_HAS_LIBC)
+endif()
+
# Prepend -I to each list element
set (LIBOMPTARGET_LLVM_INCLUDE_DIRS_DEVICERTL "${LIBOMPTARGET_LLVM_INCLUDE_DIRS}")
list(TRANSFORM LIBOMPTARGET_LLVM_INCLUDE_DIRS_DEVICERTL PREPEND "-I")
diff --git a/openmp/libomptarget/DeviceRTL/src/LibC.cpp b/openmp/libomptarget/DeviceRTL/src/LibC.cpp
index af675b97256f..33fec8135ef0 100644
--- a/openmp/libomptarget/DeviceRTL/src/LibC.cpp
+++ b/openmp/libomptarget/DeviceRTL/src/LibC.cpp
@@ -53,10 +53,23 @@ void memset(void *dst, int C, size_t count) {
dstc[I] = C;
}
+// If the user built with the GPU C library enabled we will assume that we can
+// call it.
+#ifdef OMPTARGET_HAS_LIBC
+
+// TODO: Remove this handling once we have varargs support.
+extern struct FILE *stdout;
+int32_t rpc_fprintf(FILE *, const char *, void *, uint64_t);
+
+int32_t __llvm_omp_vprintf(const char *Format, void *Arguments, uint32_t Size) {
+ return rpc_fprintf(stdout, Format, Arguments, Size);
+}
+#else
/// printf() calls are rewritten by CGGPUBuiltin to __llvm_omp_vprintf
int32_t __llvm_omp_vprintf(const char *Format, void *Arguments, uint32_t Size) {
return impl::omp_vprintf(Format, Arguments, Size);
}
+#endif
}
#pragma omp end declare target
diff --git a/openmp/libomptarget/test/libc/printf.c b/openmp/libomptarget/test/libc/printf.c
new file mode 100644
index 000000000000..64cdd805d355
--- /dev/null
+++ b/openmp/libomptarget/test/libc/printf.c
@@ -0,0 +1,36 @@
+// RUN: %libomptarget-compile-run-and-check-generic
+
+// REQUIRES: libc
+
+#include
+
+int main() {
+ // CHECK: PASS
+#pragma omp target
+ { printf("PASS\n"); }
+
+ // CHECK: PASS
+#pragma omp target
+ { printf("%s\n", "PASS"); }
+
+ // CHECK: PASS
+ // CHECK: PASS
+ // CHECK: PASS
+ // CHECK: PASS
+ // CHECK: PASS
+ // CHECK: PASS
+ // CHECK: PASS
+ // CHECK: PASS
+#pragma omp target teams num_teams(4)
+#pragma omp parallel num_threads(2)
+ { printf("PASS\n"); }
+
+ // CHECK: PASS
+ char str[] = {'P', 'A', 'S', 'S', '\0'};
+#pragma omp target map(to : str)
+ { printf("%s\n", str); }
+
+ // CHECK: 11111111111
+#pragma omp target
+ { printf("%s%-.0f%4b%c%ld\n", "1111", 1.0, 0xf, '1', 1lu); }
+}
--
GitLab
From a7f4576ff4e296ff42b16d9d91aadf82b5ea325c Mon Sep 17 00:00:00 2001
From: Owen Pan
Date: Tue, 2 Apr 2024 14:48:14 -0700
Subject: [PATCH 0015/1683] [clang-format] Fix a regression in annotating
TrailingReturnArrow (#86624)
Fixes #86559.
---
clang/lib/Format/TokenAnnotator.cpp | 2 ++
clang/unittests/Format/TokenAnnotatorTest.cpp | 4 ++++
2 files changed, 6 insertions(+)
diff --git a/clang/lib/Format/TokenAnnotator.cpp b/clang/lib/Format/TokenAnnotator.cpp
index a405a348403b..3e9988d50945 100644
--- a/clang/lib/Format/TokenAnnotator.cpp
+++ b/clang/lib/Format/TokenAnnotator.cpp
@@ -3889,6 +3889,8 @@ void TokenAnnotator::calculateFormattingInformation(AnnotatedLine &Line) const {
}
} else if (ClosingParen) {
for (auto *Tok = ClosingParen->Next; Tok; Tok = Tok->Next) {
+ if (Tok->is(TT_CtorInitializerColon))
+ break;
if (Tok->is(tok::arrow)) {
Tok->setType(TT_TrailingReturnArrow);
break;
diff --git a/clang/unittests/Format/TokenAnnotatorTest.cpp b/clang/unittests/Format/TokenAnnotatorTest.cpp
index 2539d3d76ef0..9425647d9998 100644
--- a/clang/unittests/Format/TokenAnnotatorTest.cpp
+++ b/clang/unittests/Format/TokenAnnotatorTest.cpp
@@ -1916,6 +1916,10 @@ TEST_F(TokenAnnotatorTest, UnderstandsTrailingReturnArrow) {
ASSERT_EQ(Tokens.size(), 12u) << Tokens;
EXPECT_TOKEN(Tokens[7], tok::arrow, TT_Unknown);
+ Tokens = annotate("__attribute__((cold)) C() : Base(obj->func()) {}");
+ ASSERT_EQ(Tokens.size(), 21u) << Tokens;
+ EXPECT_TOKEN(Tokens[13], tok::arrow, TT_Unknown);
+
// Mixed
Tokens = annotate("auto f() -> int { auto a = b()->c; }");
ASSERT_EQ(Tokens.size(), 18u) << Tokens;
--
GitLab
From e381586f259568bf244fcd857ce91fc5cb38b959 Mon Sep 17 00:00:00 2001
From: Joseph Huber
Date: Tue, 2 Apr 2024 17:06:03 -0500
Subject: [PATCH 0016/1683] [libc] Work around lack of '__has_builtin' for GPU
server (#87417)
Summary:
The RPC server build for the GPU support needs to be build from the
"projects" phase of the LLVM build. That means it is built with the same
compile that LLVM supports, which currently is GCC 7.4 in most cases.
A previous patch removed the `LIBC_HAS_BUILTIN` indirection we used,
which regressed the case where we used the `libc` source externally. The
files that we need to use here are `converter.cpp` and `writer.cpp`
which currently are compatible with C++17, so there aren't issues with
the code itself. However, older GCC does not have this builtin which
makes the checks fail.
This patch just adds in a simple wrapper that allows it to correctly
ignore everything if using a compiler that doesn't support it.
---
libc/src/__support/macros/config.h | 6 ++++++
1 file changed, 6 insertions(+)
diff --git a/libc/src/__support/macros/config.h b/libc/src/__support/macros/config.h
index 3f200f0d62ba..6390c7992325 100644
--- a/libc/src/__support/macros/config.h
+++ b/libc/src/__support/macros/config.h
@@ -13,6 +13,12 @@
#ifndef LLVM_LIBC_SRC___SUPPORT_MACROS_CONFIG_H
#define LLVM_LIBC_SRC___SUPPORT_MACROS_CONFIG_H
+// Workaround for compilers that do not support builtin detection.
+// FIXME: This is only required for the GPU portion which should be moved.
+#ifndef __has_builtin
+#define __has_builtin(b) 0
+#endif
+
// Compiler feature-detection.
// clang.llvm.org/docs/LanguageExtensions.html#has-feature-and-has-extension
#ifdef __has_feature
--
GitLab
From 68217a52fb9fec8a88623f97a90899b8d27eefb3 Mon Sep 17 00:00:00 2001
From: Haowei
Date: Tue, 2 Apr 2024 15:11:17 -0700
Subject: [PATCH 0017/1683] [Fuchsia] Add SWIG flags to Fuchsia Clang stage2
build (#87421)
This patch adds SWIG cmake flags to the stage2 build in Fuchsia
Clang configuration.
---
clang/cmake/caches/Fuchsia.cmake | 2 ++
1 file changed, 2 insertions(+)
diff --git a/clang/cmake/caches/Fuchsia.cmake b/clang/cmake/caches/Fuchsia.cmake
index df69d7d0dd41..393d97a4cf1a 100644
--- a/clang/cmake/caches/Fuchsia.cmake
+++ b/clang/cmake/caches/Fuchsia.cmake
@@ -71,6 +71,8 @@ set(_FUCHSIA_BOOTSTRAP_PASSTHROUGH
Python3_LIBRARIES
Python3_INCLUDE_DIRS
Python3_RPATH
+ SWIG_DIR
+ SWIG_EXECUTABLE
CMAKE_FIND_PACKAGE_PREFER_CONFIG
CMAKE_SYSROOT
CMAKE_MODULE_LINKER_FLAGS
--
GitLab
From b4adb42151bbfa80be4cf6d076cbe5edf680693e Mon Sep 17 00:00:00 2001
From: Jon Chesterfield
Date: Tue, 2 Apr 2024 22:14:03 +0100
Subject: [PATCH 0018/1683] Use setup_host_tool for clang-ast-dump, fixes 76707
---
clang/lib/Tooling/CMakeLists.txt | 6 ++++--
1 file changed, 4 insertions(+), 2 deletions(-)
diff --git a/clang/lib/Tooling/CMakeLists.txt b/clang/lib/Tooling/CMakeLists.txt
index 91e6cbdcbc44..8b4ab0e21296 100644
--- a/clang/lib/Tooling/CMakeLists.txt
+++ b/clang/lib/Tooling/CMakeLists.txt
@@ -53,14 +53,16 @@ else()
list(APPEND implicitDirs -I ${implicitDir})
endforeach()
+ setup_host_tool(clang-ast-dump CLANG_AST_DUMP clang_ast_dump_exe clang_ast_dump_target)
+
include(GetClangResourceDir)
get_clang_resource_dir(resource_dir PREFIX ${LLVM_BINARY_DIR})
add_custom_command(
COMMENT Generate ASTNodeAPI.json
OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/ASTNodeAPI.json
- DEPENDS clang-ast-dump clang-resource-headers
+ DEPENDS ${clang_ast_dump_target} clang-resource-headers
COMMAND
- $
+ ${clang_ast_dump_exe}
# Skip this in debug mode because parsing AST.h is too slow
--skip-processing=${skip_expensive_processing}
-I ${resource_dir}/include
--
GitLab
From 633bc3bfda71c55bc38d5a3bfdb426bab61ff101 Mon Sep 17 00:00:00 2001
From: Vitaly Buka
Date: Tue, 2 Apr 2024 15:45:03 -0700
Subject: [PATCH 0019/1683] [CodeGen][NFC] Make an opt<> static
---
clang/lib/CodeGen/BackendUtil.cpp | 4 ++--
1 file changed, 2 insertions(+), 2 deletions(-)
diff --git a/clang/lib/CodeGen/BackendUtil.cpp b/clang/lib/CodeGen/BackendUtil.cpp
index 1220c575d1df..c8b2a93ae47a 100644
--- a/clang/lib/CodeGen/BackendUtil.cpp
+++ b/clang/lib/CodeGen/BackendUtil.cpp
@@ -100,8 +100,8 @@ using namespace llvm;
namespace llvm {
extern cl::opt PrintPipelinePasses;
-cl::opt ClRemoveTraps("clang-remove-traps", cl::Optional,
- cl::desc("Insert remove-traps pass."));
+static cl::opt ClRemoveTraps("clang-remove-traps", cl::Optional,
+ cl::desc("Insert remove-traps pass."));
// Experiment to move sanitizers earlier.
static cl::opt ClSanitizeOnOptimizerEarlyEP(
--
GitLab
From 12c7371296e59c22debdd906f632c5e6574e3a44 Mon Sep 17 00:00:00 2001
From: Changpeng Fang
Date: Tue, 2 Apr 2024 15:51:28 -0700
Subject: [PATCH 0020/1683] AMDGPU: Use PseudoInstr instead of Pseudo Mnemonic
for SIMCInstr, NFC (#87420)
Pseudo Mnemonic could be of other uses.
---
llvm/lib/Target/AMDGPU/DSInstructions.td | 4 ++--
1 file changed, 2 insertions(+), 2 deletions(-)
diff --git a/llvm/lib/Target/AMDGPU/DSInstructions.td b/llvm/lib/Target/AMDGPU/DSInstructions.td
index e944dde15990..0773ef7f3234 100644
--- a/llvm/lib/Target/AMDGPU/DSInstructions.td
+++ b/llvm/lib/Target/AMDGPU/DSInstructions.td
@@ -1192,7 +1192,7 @@ def : GCNPat <
class Base_DS_Real_gfx6_gfx7_gfx10_gfx11_gfx12 op, DS_Pseudo ps, int ef,
string opName = ps.Mnemonic,
bit hasGDS = true>
- : DS_Real, SIMCInstr {
+ : DS_Real, SIMCInstr {
let Inst{7-0} = !if(ps.has_offset0, offset0, 0);
let Inst{15-8} = !if(ps.has_offset1, offset1, 0);
@@ -1557,7 +1557,7 @@ defm DS_MAX_SRC2_F64 : DS_Real_gfx6_gfx7_gfx10<0x0d3>;
class DS_Real_vi op, DS_Pseudo ps> :
DS_Real ,
- SIMCInstr {
+ SIMCInstr {
let AssemblerPredicate = isGFX8GFX9;
let DecoderNamespace = "GFX8";
--
GitLab
From 84ae8cb4af9abafe9f45e69744607aadb38d649a Mon Sep 17 00:00:00 2001
From: =?UTF-8?q?Jan=20Kokem=C3=BCller?=
Date: Wed, 3 Apr 2024 01:09:26 +0200
Subject: [PATCH 0021/1683] [libc++] `std::ranges::advance`: avoid unneeded
bounds checks when advancing iterator (#84126)
Currently, the bounds check in `std::ranges::advance(it, n, s)` is done
_before_ `n` is checked. This results in one extra, unneeded bounds
check.
Thus, `std::ranges::advance(it, 1, s)` currently is _not_ simply
equivalent to:
```c++
if (it != s) {
++it;
}
```
This difference in behavior matters when the check involves some
"expensive" logic. For example, the `==` operator of
`std::istreambuf_iterator` may actually have to read the underlying
`streambuf`.
Swapping around the checks in the `while` results in the expected
behavior.
---
libcxx/include/__iterator/advance.h | 4 +-
.../iterator_count_sentinel.pass.cpp | 76 ++++++++++++++-----
libcxx/test/support/test_iterators.h | 22 +++++-
3 files changed, 80 insertions(+), 22 deletions(-)
diff --git a/libcxx/include/__iterator/advance.h b/libcxx/include/__iterator/advance.h
index 7959bdeae326..296db1aaab65 100644
--- a/libcxx/include/__iterator/advance.h
+++ b/libcxx/include/__iterator/advance.h
@@ -170,14 +170,14 @@ public:
} else {
// Otherwise, if `n` is non-negative, while `bool(i != bound_sentinel)` is true, increments `i` but at
// most `n` times.
- while (__i != __bound_sentinel && __n > 0) {
+ while (__n > 0 && __i != __bound_sentinel) {
++__i;
--__n;
}
// Otherwise, while `bool(i != bound_sentinel)` is true, decrements `i` but at most `-n` times.
if constexpr (bidirectional_iterator<_Ip> && same_as<_Ip, _Sp>) {
- while (__i != __bound_sentinel && __n < 0) {
+ while (__n < 0 && __i != __bound_sentinel) {
--__i;
++__n;
}
diff --git a/libcxx/test/std/iterators/iterator.primitives/range.iter.ops/range.iter.ops.advance/iterator_count_sentinel.pass.cpp b/libcxx/test/std/iterators/iterator.primitives/range.iter.ops/range.iter.ops.advance/iterator_count_sentinel.pass.cpp
index a1c156401821..76439ef93a60 100644
--- a/libcxx/test/std/iterators/iterator.primitives/range.iter.ops/range.iter.ops.advance/iterator_count_sentinel.pass.cpp
+++ b/libcxx/test/std/iterators/iterator.primitives/range.iter.ops/range.iter.ops.advance/iterator_count_sentinel.pass.cpp
@@ -21,9 +21,12 @@
#include "../types.h"
template
-constexpr void check_forward(int* first, int* last, std::iter_difference_t n, int* expected) {
+constexpr void
+check_forward(int* first, int* last, std::iter_difference_t n, int* expected, int expected_equals_count = -1) {
using Difference = std::iter_difference_t;
Difference const M = (expected - first); // expected travel distance
+ // `expected_equals_count` is only relevant when `Count` is true.
+ assert(Count || expected_equals_count == -1);
{
It it(first);
@@ -42,6 +45,7 @@ constexpr void check_forward(int* first, int* last, std::iter_difference_t n
// regardless of the iterator category.
assert(it.stride_count() == M);
assert(it.stride_displacement() == M);
+ assert(it.equals_count() == expected_equals_count);
}
}
@@ -74,9 +78,20 @@ constexpr void check_forward_sized_sentinel(int* first, int* last, std::iter_dif
}
}
-template
-constexpr void check_backward(int* first, int* last, std::iter_difference_t n, int* expected) {
- static_assert(std::random_access_iterator, "This test doesn't support non random access iterators");
+struct Expected {
+ int stride_count;
+ int stride_displacement;
+ int equals_count;
+};
+
+template
+constexpr void
+check_backward(int* first, int* last, std::iter_difference_t n, int* expected, Expected expected_counts) {
+ // Check preconditions for `advance` when called with negative `n`
+ // (see [range.iter.op.advance]). In addition, allow `n == 0`.
+ assert(n <= 0);
+ static_assert(std::bidirectional_iterator);
+
using Difference = std::iter_difference_t;
Difference const M = (expected - last); // expected travel distance (which is negative)
@@ -92,9 +107,14 @@ constexpr void check_backward(int* first, int* last, std::iter_difference_t
{
auto it = stride_counting_iterator(It(last));
auto sent = stride_counting_iterator(It(first));
+ static_assert(std::bidirectional_iterator>);
+ static_assert(Count == !std::sized_sentinel_for);
+
(void)std::ranges::advance(it, n, sent);
- assert(it.stride_count() <= 1);
- assert(it.stride_displacement() <= 1);
+
+ assert(it.stride_count() == expected_counts.stride_count);
+ assert(it.stride_displacement() == expected_counts.stride_displacement);
+ assert(it.equals_count() == expected_counts.equals_count);
}
}
@@ -171,13 +191,17 @@ constexpr bool test() {
{
int* expected = n > size ? range + size : range + n;
+ int equals_count = n > size ? size + 1 : n;
+
+ // clang-format off
check_forward>( range, range+size, n, expected);
check_forward>( range, range+size, n, expected);
- check_forward>( range, range+size, n, expected);
- check_forward>(range, range+size, n, expected);
- check_forward>(range, range+size, n, expected);
- check_forward>( range, range+size, n, expected);
- check_forward( range, range+size, n, expected);
+ check_forward>( range, range+size, n, expected, equals_count);
+ check_forward>(range, range+size, n, expected, equals_count);
+ check_forward>(range, range+size, n, expected, equals_count);
+ check_forward>( range, range+size, n, expected, equals_count);
+ check_forward( range, range+size, n, expected, equals_count);
+ // clang-format on
check_forward_sized_sentinel>( range, range+size, n, expected);
check_forward_sized_sentinel>( range, range+size, n, expected);
@@ -188,14 +212,32 @@ constexpr bool test() {
check_forward_sized_sentinel( range, range+size, n, expected);
}
+ // Input and forward iterators are not tested as the backwards case does
+ // not apply for them.
{
- // Note that we can only test ranges::advance with a negative n for iterators that
- // are sized sentinels for themselves, because ranges::advance is UB otherwise.
- // In particular, that excludes bidirectional_iterators since those are not sized sentinels.
int* expected = n > size ? range : range + size - n;
- check_backward>(range, range+size, -n, expected);
- check_backward>( range, range+size, -n, expected);
- check_backward( range, range+size, -n, expected);
+ {
+ Expected expected_counts = {
+ .stride_count = static_cast(range + size - expected),
+ .stride_displacement = -expected_counts.stride_count,
+ .equals_count = n > size ? size + 1 : n,
+ };
+
+ check_backward>(range, range + size, -n, expected, expected_counts);
+ }
+ {
+ Expected expected_counts = {
+ // If `n >= size`, the algorithm can just do `it = std::move(sent);`
+ // instead of doing iterator arithmetic.
+ .stride_count = (n >= size) ? 0 : 1,
+ .stride_displacement = (n >= size) ? 0 : 1,
+ .equals_count = 0,
+ };
+
+ check_backward>(range, range + size, -n, expected, expected_counts);
+ check_backward>(range, range + size, -n, expected, expected_counts);
+ check_backward(range, range + size, -n, expected, expected_counts);
+ }
}
}
}
diff --git a/libcxx/test/support/test_iterators.h b/libcxx/test/support/test_iterators.h
index c92ce375348f..7ffb74990fa4 100644
--- a/libcxx/test/support/test_iterators.h
+++ b/libcxx/test/support/test_iterators.h
@@ -725,11 +725,14 @@ struct common_input_iterator {
# endif // TEST_STD_VER >= 20
// Iterator adaptor that counts the number of times the iterator has had a successor/predecessor
-// operation called. Has two recorders:
+// operation or an equality comparison operation called. Has three recorders:
// * `stride_count`, which records the total number of calls to an op++, op--, op+=, or op-=.
// * `stride_displacement`, which records the displacement of the calls. This means that both
// op++/op+= will increase the displacement counter by 1, and op--/op-= will decrease the
// displacement counter by 1.
+// * `equals_count`, which records the total number of calls to an op== or op!=. If compared
+// against a sentinel object, that sentinel object must call the `record_equality_comparison`
+// function so that the comparison is counted correctly.
template
class stride_counting_iterator {
public:
@@ -754,6 +757,8 @@ public:
constexpr difference_type stride_displacement() const { return stride_displacement_; }
+ constexpr difference_type equals_count() const { return equals_count_; }
+
constexpr decltype(auto) operator*() const { return *It(base_); }
constexpr decltype(auto) operator[](difference_type n) const { return It(base_)[n]; }
@@ -838,10 +843,13 @@ public:
return base(x) - base(y);
}
+ constexpr void record_equality_comparison() const { ++equals_count_; }
+
constexpr bool operator==(stride_counting_iterator const& other) const
requires std::sentinel_for
{
- return It(base_) == It(other.base_);
+ record_equality_comparison();
+ return It(base_) == It(other.base_);
}
friend constexpr bool operator<(stride_counting_iterator const& x, stride_counting_iterator const& y)
@@ -875,6 +883,7 @@ private:
decltype(base(std::declval())) base_;
difference_type stride_count_ = 0;
difference_type stride_displacement_ = 0;
+ mutable difference_type equals_count_ = 0;
};
template
stride_counting_iterator(It) -> stride_counting_iterator