[clang] [Clang] Default OpenMP and C/C++ compilations to uniform blocks (PR #221279)
Joseph Huber via cfe-commits
cfe-commits at lists.llvm.org
Sat Sep 5 15:40:50 PDT 2026
https://github.com/jhuber6 updated https://github.com/llvm/llvm-project/pull/221279
>From 5169ab996b6cac4b36846878b0493ce9ce89561d Mon Sep 17 00:00:00 2001
From: Joseph Huber <huberjn at outlook.com>
Date: Fri, 4 Sep 2026 11:27:24 -0500
Subject: [PATCH 1/5] [Clang] Default OpenMP and C/C++ compilations to uniform
blocks
Summary:
This is an OpenCL feature, and the more generic one so it's still the
default. We opt-in on targets where we know we don't need to support the
OpenCL model. This can be overridden with `-f[no-]offload-uniform`
---
clang/include/clang/Basic/LangOptions.def | 2 +-
clang/include/clang/Basic/LangOptions.h | 3 +
clang/include/clang/Options/Options.td | 5 +-
clang/lib/Basic/LangOptions.cpp | 9 +++
clang/lib/Frontend/CompilerInvocation.cpp | 3 +
clang/test/CodeGen/amdgpu-abi-version.c | 56 +++++--------
clang/test/CodeGen/amdgpu-address-spaces.cpp | 2 +-
.../amdgpu-builtin-is-invocable-subarch.c | 4 +-
.../CodeGen/amdgpu-builtin-is-invocable.c | 6 +-
.../amdgpu-builtin-processor-is-subarch.c | 6 +-
.../CodeGen/amdgpu-builtin-processor-is.c | 6 +-
.../CodeGen/amdgpu-offload-uniform-block.c | 17 ++++
clang/test/CodeGen/nvptx_attributes.c | 2 +-
.../CodeGenCXX/dynamic-cast-address-space.cpp | 8 +-
clang/test/CodeGenSYCL/function-attrs.cpp | 4 +-
clang/test/Headers/gpuintrin.c | 80 ++++++-------------
clang/test/OpenMP/amdgcn-attributes.cpp | 4 +-
.../amdgcn_target_global_constructor.cpp | 8 +-
clang/test/OpenMP/amdgcn_weak_alias.c | 2 +-
clang/test/OpenMP/amdgcn_weak_alias.cpp | 2 +-
20 files changed, 107 insertions(+), 122 deletions(-)
create mode 100644 clang/test/CodeGen/amdgpu-offload-uniform-block.c
diff --git a/clang/include/clang/Basic/LangOptions.def b/clang/include/clang/Basic/LangOptions.def
index d7637f2dfd507..6016f0003ca28 100644
--- a/clang/include/clang/Basic/LangOptions.def
+++ b/clang/include/clang/Basic/LangOptions.def
@@ -272,7 +272,7 @@ LANGOPT(SYCLIsHost , 1, 0, NotCompatible, "SYCL host compilation")
ENUM_LANGOPT(SYCLVersion , SYCLMajorVersion, 2, SYCL_None, NotCompatible, "Version of the SYCL standard used")
LANGOPT(HIPUseNewLaunchAPI, 1, 0, NotCompatible, "Use new kernel launching API for HIP")
-LANGOPT(OffloadUniformBlock, 1, 0, NotCompatible, "Assume that kernels are launched with uniform block sizes (default true for CUDA/HIP and false otherwise)")
+LANGOPT(OffloadUniformBlock, 1, 0, NotCompatible, "Assume that kernels are launched with uniform block sizes (default true for CUDA/HIP, OpenMP device, SYCL device, AMDGPU/NVPTX C/C++, and OpenCL 1.2)")
LANGOPT(HIPStdPar, 1, 0, NotCompatible, "Enable Standard Parallel Algorithm Acceleration for HIP (experimental)")
LANGOPT(HIPStdParInterposeAlloc, 1, 0, NotCompatible, "Replace allocations / deallocations with HIP RT calls when Standard Parallel Algorithm Acceleration for HIP is enabled (Experimental)")
diff --git a/clang/include/clang/Basic/LangOptions.h b/clang/include/clang/Basic/LangOptions.h
index 7539e000d03f9..2b664818d0a48 100644
--- a/clang/include/clang/Basic/LangOptions.h
+++ b/clang/include/clang/Basic/LangOptions.h
@@ -823,6 +823,9 @@ class LangOptions : public LangOptionsBase {
return OpenMPIsTargetDevice || CUDAIsDevice || SYCLIsDevice;
}
+ /// Languages that launch full work-groups opt in.
+ bool defaultOffloadUniformBlock(const llvm::Triple &T) const;
+
/// Return the OpenMP version.
llvm::omp::Version getOpenMPVersion() const {
return llvm::omp::Version(OpenMP);
diff --git a/clang/include/clang/Options/Options.td b/clang/include/clang/Options/Options.td
index fe8e8fdb75de8..605b791f053dc 100644
--- a/clang/include/clang/Options/Options.td
+++ b/clang/include/clang/Options/Options.td
@@ -1147,10 +1147,11 @@ def b : JoinedOrSeparate<["-"], "b">, Flags<[LinkerInput]>,
Group<Link_Group>;
defm offload_uniform_block : BoolFOption<"offload-uniform-block",
- LangOpts<"OffloadUniformBlock">, Default<"LangOpts->CUDA || (LangOpts->OpenCL && LangOpts->OpenCLVersion <= 120)">,
+ LangOpts<"OffloadUniformBlock">,
+ Default<"LangOpts->defaultOffloadUniformBlock(T)">,
PosFlag<SetTrue, [], [ClangOption, CC1Option], "Assume">,
NegFlag<SetFalse, [], [ClangOption, CC1Option], "Don't assume">,
- BothFlags<[], [ClangOption], " that kernels are launched with uniform block sizes (default true for CUDA/HIP and false otherwise)">>;
+ BothFlags<[], [ClangOption], " that kernels are launched with uniform block sizes (default true for CUDA/HIP, OpenMP device, SYCL device, AMDGPU/NVPTX C/C++, and OpenCL 1.2)">>;
def fcomplex_arithmetic_EQ : Joined<["-"], "fcomplex-arithmetic=">, Group<f_Group>,
Visibility<[ClangOption, CC1Option, FlangOption, FC1Option]>,
diff --git a/clang/lib/Basic/LangOptions.cpp b/clang/lib/Basic/LangOptions.cpp
index a2c069c575e61..3842d307c7e48 100644
--- a/clang/lib/Basic/LangOptions.cpp
+++ b/clang/lib/Basic/LangOptions.cpp
@@ -74,6 +74,15 @@ unsigned LangOptions::getOpenCLCompatibleVersion() const {
llvm_unreachable("Unknown OpenCL version");
}
+bool LangOptions::defaultOffloadUniformBlock(const llvm::Triple &T) const {
+ if (OpenCL)
+ return getOpenCLCompatibleVersion() <= 120;
+ if (CUDA || OpenMPIsTargetDevice || SYCLIsDevice)
+ return true;
+ // Direct C/C++ for AMDGPU/NVPTX.
+ return !HLSL && (T.isAMDGPU() || T.isNVPTX());
+}
+
void LangOptions::remapPathPrefix(SmallVectorImpl<char> &Path) const {
for (const auto &Entry : MacroPrefixMap)
if (llvm::sys::path::replace_path_prefix(Path, Entry.first, Entry.second))
diff --git a/clang/lib/Frontend/CompilerInvocation.cpp b/clang/lib/Frontend/CompilerInvocation.cpp
index ea8368908879a..2089426582f1b 100644
--- a/clang/lib/Frontend/CompilerInvocation.cpp
+++ b/clang/lib/Frontend/CompilerInvocation.cpp
@@ -4349,6 +4349,9 @@ bool CompilerInvocation::ParseLangArgs(LangOptions &Opts, ArgList &Args,
Opts.OpenMP && !Args.hasArg(options::OPT_fnoopenmp_use_tls);
Opts.OpenMPIsTargetDevice =
Opts.OpenMP && Args.hasArg(options::OPT_fopenmp_is_target_device);
+ if (!Args.getLastArg(OPT_foffload_uniform_block,
+ OPT_fno_offload_uniform_block))
+ Opts.OffloadUniformBlock = Opts.defaultOffloadUniformBlock(T);
Opts.OpenMPIRBuilder =
Opts.OpenMP && Args.hasArg(options::OPT_fopenmp_enable_irbuilder);
bool IsTargetSpecified =
diff --git a/clang/test/CodeGen/amdgpu-abi-version.c b/clang/test/CodeGen/amdgpu-abi-version.c
index 00efbb82e346c..80862211190d6 100644
--- a/clang/test/CodeGen/amdgpu-abi-version.c
+++ b/clang/test/CodeGen/amdgpu-abi-version.c
@@ -11,58 +11,40 @@
// LLVM-NEXT: [[TMP0:%.*]] = load i32, ptr addrspace(4) @__oclc_ABI_version, align 4
// LLVM-NEXT: [[TMP1:%.*]] = icmp sge i32 [[TMP0]], 500
// LLVM-NEXT: [[TMP2:%.*]] = call align 8 dereferenceable(256) ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr()
-// LLVM-NEXT: [[TMP3:%.*]] = getelementptr i8, ptr addrspace(4) [[TMP2]], i64 0
-// LLVM-NEXT: [[TMP4:%.*]] = load i32, ptr addrspace(4) [[TMP3]], align 4, !invariant.load [[META1:![0-9]+]], !noundef [[META1]]
-// LLVM-NEXT: [[TMP5:%.*]] = call i32 @llvm.amdgcn.workgroup.id.x()
-// LLVM-NEXT: [[TMP6:%.*]] = icmp ult i32 [[TMP5]], [[TMP4]]
-// LLVM-NEXT: [[TMP7:%.*]] = select i1 [[TMP6]], i32 12, i32 18
-// LLVM-NEXT: [[TMP8:%.*]] = getelementptr inbounds i8, ptr addrspace(4) [[TMP2]], i32 [[TMP7]]
-// LLVM-NEXT: [[TMP9:%.*]] = load i16, ptr addrspace(4) [[TMP8]], align 2, !range [[RNG2:![0-9]+]], !invariant.load [[META1]], !noundef [[META1]]
-// LLVM-NEXT: [[TMP10:%.*]] = zext i16 [[TMP9]] to i32
-// LLVM-NEXT: [[TMP11:%.*]] = call ptr addrspace(4) @llvm.amdgcn.dispatch.ptr()
-// LLVM-NEXT: [[TMP12:%.*]] = getelementptr inbounds i8, ptr addrspace(4) [[TMP11]], i64 4
-// LLVM-NEXT: [[TMP13:%.*]] = load i16, ptr addrspace(4) [[TMP12]], align 2, !range [[RNG2]], !invariant.load [[META1]], !noundef [[META1]]
-// LLVM-NEXT: [[TMP14:%.*]] = zext i16 [[TMP13]] to i32
-// LLVM-NEXT: [[TMP15:%.*]] = call i32 @llvm.amdgcn.workgroup.id.x()
-// LLVM-NEXT: [[TMP16:%.*]] = getelementptr inbounds i8, ptr addrspace(4) [[TMP11]], i64 12
-// LLVM-NEXT: [[TMP17:%.*]] = load i32, ptr addrspace(4) [[TMP16]], align 4, !range [[RNG3:![0-9]+]], !invariant.load [[META1]]
-// LLVM-NEXT: [[TMP18:%.*]] = mul i32 [[TMP15]], [[TMP14]]
-// LLVM-NEXT: [[TMP19:%.*]] = sub i32 [[TMP17]], [[TMP18]]
-// LLVM-NEXT: [[TMP20:%.*]] = icmp ult i32 [[TMP19]], [[TMP14]]
-// LLVM-NEXT: [[TMP21:%.*]] = select i1 [[TMP20]], i32 [[TMP19]], i32 [[TMP14]]
-// LLVM-NEXT: [[TMP22:%.*]] = select i1 [[TMP1]], i32 [[TMP10]], i32 [[TMP21]]
-// LLVM-NEXT: ret i32 [[TMP22]]
+// LLVM-NEXT: [[TMP3:%.*]] = getelementptr inbounds i8, ptr addrspace(4) [[TMP2]], i64 12
+// LLVM-NEXT: [[TMP4:%.*]] = load i16, ptr addrspace(4) [[TMP3]], align 2, !range [[RNG1:![0-9]+]], !invariant.load [[META2:![0-9]+]], !noundef [[META2]]
+// LLVM-NEXT: [[TMP5:%.*]] = zext i16 [[TMP4]] to i32
+// LLVM-NEXT: [[TMP6:%.*]] = call ptr addrspace(4) @llvm.amdgcn.dispatch.ptr()
+// LLVM-NEXT: [[TMP7:%.*]] = getelementptr inbounds i8, ptr addrspace(4) [[TMP6]], i64 4
+// LLVM-NEXT: [[TMP8:%.*]] = load i16, ptr addrspace(4) [[TMP7]], align 2, !range [[RNG1]], !invariant.load [[META2]], !noundef [[META2]]
+// LLVM-NEXT: [[TMP9:%.*]] = zext i16 [[TMP8]] to i32
+// LLVM-NEXT: [[TMP10:%.*]] = select i1 [[TMP1]], i32 [[TMP5]], i32 [[TMP9]]
+// LLVM-NEXT: ret i32 [[TMP10]]
//
// LLVMENV-LABEL: define dso_local i32 @foo(
// LLVMENV-SAME: ) #[[ATTR0:[0-9]+]] {
// LLVMENV-NEXT: [[ENTRY:.*:]]
// LLVMENV-NEXT: [[TMP0:%.*]] = call align 8 dereferenceable(256) ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr()
-// LLVMENV-NEXT: [[TMP1:%.*]] = getelementptr i8, ptr addrspace(4) [[TMP0]], i64 0
-// LLVMENV-NEXT: [[TMP2:%.*]] = load i32, ptr addrspace(4) [[TMP1]], align 4, !invariant.load [[META1:![0-9]+]], !noundef [[META1]]
-// LLVMENV-NEXT: [[TMP3:%.*]] = call i32 @llvm.amdgcn.workgroup.id.x()
-// LLVMENV-NEXT: [[TMP4:%.*]] = icmp ult i32 [[TMP3]], [[TMP2]]
-// LLVMENV-NEXT: [[TMP5:%.*]] = select i1 [[TMP4]], i32 12, i32 18
-// LLVMENV-NEXT: [[TMP6:%.*]] = getelementptr inbounds i8, ptr addrspace(4) [[TMP0]], i32 [[TMP5]]
-// LLVMENV-NEXT: [[TMP7:%.*]] = load i16, ptr addrspace(4) [[TMP6]], align 2, !range [[RNG2:![0-9]+]], !invariant.load [[META1]], !noundef [[META1]]
-// LLVMENV-NEXT: [[TMP8:%.*]] = zext i16 [[TMP7]] to i32
-// LLVMENV-NEXT: ret i32 [[TMP8]]
+// LLVMENV-NEXT: [[TMP1:%.*]] = getelementptr inbounds i8, ptr addrspace(4) [[TMP0]], i64 12
+// LLVMENV-NEXT: [[TMP2:%.*]] = load i16, ptr addrspace(4) [[TMP1]], align 2, !range [[RNG1:![0-9]+]], !invariant.load [[META2:![0-9]+]], !noundef [[META2]]
+// LLVMENV-NEXT: [[TMP3:%.*]] = zext i16 [[TMP2]] to i32
+// LLVMENV-NEXT: ret i32 [[TMP3]]
//
int foo() { return __builtin_amdgcn_workgroup_size_x(); }
//.
-// LLVM: attributes #[[ATTR0]] = { convergent noinline nounwind optnone "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
+// LLVM: attributes #[[ATTR0]] = { convergent noinline nounwind optnone "no-trapping-math"="true" "stack-protector-buffer-size"="8" "uniform-work-group-size" }
// LLVM: attributes #[[ATTR1:[0-9]+]] = { nocallback nofree nosync nounwind speculatable willreturn memory(none) }
//.
-// LLVMENV: attributes #[[ATTR0]] = { convergent noinline nounwind optnone "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
+// LLVMENV: attributes #[[ATTR0]] = { convergent noinline nounwind optnone "no-trapping-math"="true" "stack-protector-buffer-size"="8" "uniform-work-group-size" }
// LLVMENV: attributes #[[ATTR1:[0-9]+]] = { nocallback nofree nosync nounwind speculatable willreturn memory(none) }
//.
// LLVM: [[META0:![0-9]+]] = !{!"{{.*}}clang version {{.*}}"}
-// LLVM: [[META1]] = !{}
-// LLVM: [[RNG2]] = !{i16 1, i16 1025}
-// LLVM: [[RNG3]] = !{i32 1, i32 0}
+// LLVM: [[RNG1]] = !{i16 1, i16 1025}
+// LLVM: [[META2]] = !{}
//.
// LLVMENV: [[META0:![0-9]+]] = !{!"{{.*}}clang version {{.*}}"}
-// LLVMENV: [[META1]] = !{}
-// LLVMENV: [[RNG2]] = !{i16 1, i16 1025}
+// LLVMENV: [[RNG1]] = !{i16 1, i16 1025}
+// LLVMENV: [[META2]] = !{}
//.
//// NOTE: These prefixes are unused and the list is autogenerated. Do not add tests below this line:
// CHECK: {{.*}}
diff --git a/clang/test/CodeGen/amdgpu-address-spaces.cpp b/clang/test/CodeGen/amdgpu-address-spaces.cpp
index ce86b287d113e..ea57b35314cde 100644
--- a/clang/test/CodeGen/amdgpu-address-spaces.cpp
+++ b/clang/test/CodeGen/amdgpu-address-spaces.cpp
@@ -60,7 +60,7 @@ extern "C" [[clang::amdgpu_kernel]] void foo() {
bbb = 0;
}
//.
-// CHECK: attributes #[[ATTR0]] = { convergent mustprogress noinline nounwind optnone "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
+// CHECK: attributes #[[ATTR0]] = { convergent mustprogress noinline nounwind optnone "no-trapping-math"="true" "stack-protector-buffer-size"="8" "uniform-work-group-size" }
//.
// CHECK: [[META0:![0-9]+]] = !{i32 1, !"amdhsa_code_object_version", i32 600}
// CHECK: [[META1:![0-9]+]] = !{!"{{.*}}clang version {{.*}}"}
diff --git a/clang/test/CodeGen/amdgpu-builtin-is-invocable-subarch.c b/clang/test/CodeGen/amdgpu-builtin-is-invocable-subarch.c
index 5539286404d94..801965b196b5f 100644
--- a/clang/test/CodeGen/amdgpu-builtin-is-invocable-subarch.c
+++ b/clang/test/CodeGen/amdgpu-builtin-is-invocable-subarch.c
@@ -26,9 +26,9 @@ void foo() {
return __builtin_trap();
}
//.
-// AMDGPU900: attributes #[[ATTR0]] = { convergent noinline nounwind optnone "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
+// AMDGPU900: attributes #[[ATTR0]] = { convergent noinline nounwind optnone "no-trapping-math"="true" "stack-protector-buffer-size"="8" "uniform-work-group-size" }
//.
-// AMDGPU1010: attributes #[[ATTR0]] = { convergent noinline nounwind optnone "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
+// AMDGPU1010: attributes #[[ATTR0]] = { convergent noinline nounwind optnone "no-trapping-math"="true" "stack-protector-buffer-size"="8" "uniform-work-group-size" }
// AMDGPU1010: attributes #[[ATTR1:[0-9]+]] = { cold noreturn nounwind memory(inaccessiblemem: write) }
// AMDGPU1010: attributes #[[ATTR2]] = { noreturn nounwind }
//.
diff --git a/clang/test/CodeGen/amdgpu-builtin-is-invocable.c b/clang/test/CodeGen/amdgpu-builtin-is-invocable.c
index 0c4cb720c45a3..52b9051e58ab5 100644
--- a/clang/test/CodeGen/amdgpu-builtin-is-invocable.c
+++ b/clang/test/CodeGen/amdgpu-builtin-is-invocable.c
@@ -53,13 +53,13 @@ void foo() {
return __builtin_trap();
}
//.
-// AMDGCN-GFX900: attributes #[[ATTR0]] = { convergent noinline nounwind optnone "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
+// AMDGCN-GFX900: attributes #[[ATTR0]] = { convergent noinline nounwind optnone "no-trapping-math"="true" "stack-protector-buffer-size"="8" "uniform-work-group-size" }
//.
-// AMDGCN-GFX1010: attributes #[[ATTR0]] = { convergent noinline nounwind optnone "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
+// AMDGCN-GFX1010: attributes #[[ATTR0]] = { convergent noinline nounwind optnone "no-trapping-math"="true" "stack-protector-buffer-size"="8" "uniform-work-group-size" }
// AMDGCN-GFX1010: attributes #[[ATTR1:[0-9]+]] = { cold noreturn nounwind memory(inaccessiblemem: write) }
// AMDGCN-GFX1010: attributes #[[ATTR2]] = { noreturn nounwind }
//.
-// AMDGCNSPIRV: attributes #[[ATTR0]] = { noinline nounwind optnone "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
+// AMDGCNSPIRV: attributes #[[ATTR0]] = { noinline nounwind optnone "no-trapping-math"="true" "stack-protector-buffer-size"="8" "uniform-work-group-size" }
// AMDGCNSPIRV: attributes #[[ATTR1:[0-9]+]] = { nounwind }
// AMDGCNSPIRV: attributes #[[ATTR2:[0-9]+]] = { cold noreturn nounwind memory(inaccessiblemem: write) }
// AMDGCNSPIRV: attributes #[[ATTR3]] = { noreturn nounwind }
diff --git a/clang/test/CodeGen/amdgpu-builtin-processor-is-subarch.c b/clang/test/CodeGen/amdgpu-builtin-processor-is-subarch.c
index 0a2871b32ca02..720ae1c5a62ed 100644
--- a/clang/test/CodeGen/amdgpu-builtin-processor-is-subarch.c
+++ b/clang/test/CodeGen/amdgpu-builtin-processor-is-subarch.c
@@ -33,13 +33,13 @@ void gfx900_name() {
}
//.
-// AMDGPU900: attributes #[[ATTR0]] = { convergent noinline nounwind optnone "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
+// AMDGPU900: attributes #[[ATTR0]] = { convergent noinline nounwind optnone "no-trapping-math"="true" "stack-protector-buffer-size"="8" "uniform-work-group-size" }
// AMDGPU900: attributes #[[ATTR1:[0-9]+]] = { cold noreturn nounwind memory(inaccessiblemem: write) }
// AMDGPU900: attributes #[[ATTR2]] = { noreturn nounwind }
//.
-// AMDGPU9_NOCPU: attributes #[[ATTR0]] = { convergent noinline nounwind optnone "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
+// AMDGPU9_NOCPU: attributes #[[ATTR0]] = { convergent noinline nounwind optnone "no-trapping-math"="true" "stack-protector-buffer-size"="8" "uniform-work-group-size" }
//.
-// AMDGPU900_CPU: attributes #[[ATTR0]] = { convergent noinline nounwind optnone "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-cpu"="gfx900" }
+// AMDGPU900_CPU: attributes #[[ATTR0]] = { convergent noinline nounwind optnone "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-cpu"="gfx900" "uniform-work-group-size" }
// AMDGPU900_CPU: attributes #[[ATTR1:[0-9]+]] = { cold noreturn nounwind memory(inaccessiblemem: write) }
// AMDGPU900_CPU: attributes #[[ATTR2]] = { noreturn nounwind }
//.
diff --git a/clang/test/CodeGen/amdgpu-builtin-processor-is.c b/clang/test/CodeGen/amdgpu-builtin-processor-is.c
index bf02b2da77849..36df7847529d0 100644
--- a/clang/test/CodeGen/amdgpu-builtin-processor-is.c
+++ b/clang/test/CodeGen/amdgpu-builtin-processor-is.c
@@ -62,13 +62,13 @@ void foo() {
return __builtin_trap();
}
//.
-// AMDGCN-GFX900: attributes #[[ATTR0]] = { convergent noinline nounwind optnone "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
+// AMDGCN-GFX900: attributes #[[ATTR0]] = { convergent noinline nounwind optnone "no-trapping-math"="true" "stack-protector-buffer-size"="8" "uniform-work-group-size" }
// AMDGCN-GFX900: attributes #[[ATTR1:[0-9]+]] = { cold noreturn nounwind memory(inaccessiblemem: write) }
// AMDGCN-GFX900: attributes #[[ATTR2]] = { noreturn nounwind }
//.
-// AMDGCN-GFX1010: attributes #[[ATTR0]] = { convergent noinline nounwind optnone "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
+// AMDGCN-GFX1010: attributes #[[ATTR0]] = { convergent noinline nounwind optnone "no-trapping-math"="true" "stack-protector-buffer-size"="8" "uniform-work-group-size" }
//.
-// AMDGCNSPIRV: attributes #[[ATTR0]] = { noinline nounwind optnone "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
+// AMDGCNSPIRV: attributes #[[ATTR0]] = { noinline nounwind optnone "no-trapping-math"="true" "stack-protector-buffer-size"="8" "uniform-work-group-size" }
// AMDGCNSPIRV: attributes #[[ATTR1:[0-9]+]] = { nounwind }
// AMDGCNSPIRV: attributes #[[ATTR2:[0-9]+]] = { cold noreturn nounwind memory(inaccessiblemem: write) }
// AMDGCNSPIRV: attributes #[[ATTR3]] = { noreturn nounwind }
diff --git a/clang/test/CodeGen/amdgpu-offload-uniform-block.c b/clang/test/CodeGen/amdgpu-offload-uniform-block.c
new file mode 100644
index 0000000000000..dae42a7db5901
--- /dev/null
+++ b/clang/test/CodeGen/amdgpu-offload-uniform-block.c
@@ -0,0 +1,17 @@
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -emit-llvm %s -o - \
+// RUN: | FileCheck %s --check-prefixes=CHECK,UNIFORM
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fno-offload-uniform-block \
+// RUN: -emit-llvm %s -o - | FileCheck %s --check-prefixes=CHECK,REMAINDER
+
+#ifdef __AMDGPU__
+int foo(void) { return __builtin_amdgcn_workgroup_size_x(); }
+#else
+int foo(void) { return 0; }
+#endif
+
+// CHECK-LABEL: define{{.*}} i32 @foo(
+// UNIFORM: getelementptr inbounds i8, ptr addrspace(4) {{.*}}, i64 12
+// UNIFORM-NOT: select i1
+// UNIFORM: "uniform-work-group-size"
+// REMAINDER: select i1 {{.*}}, i32 12, i32 18
+// REMAINDER-NOT: "uniform-work-group-size"
diff --git a/clang/test/CodeGen/nvptx_attributes.c b/clang/test/CodeGen/nvptx_attributes.c
index fbba84ff05d5a..72f16be197a11 100644
--- a/clang/test/CodeGen/nvptx_attributes.c
+++ b/clang/test/CodeGen/nvptx_attributes.c
@@ -16,7 +16,7 @@ __attribute__((nvptx_kernel)) void foo(int *ret) {
}
//.
-// CHECK: attributes #[[ATTR0]] = { convergent noinline nounwind optnone "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-cpu"="sm_61" }
+// CHECK: attributes #[[ATTR0]] = { convergent noinline nounwind optnone "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-cpu"="sm_61" "uniform-work-group-size" }
//.
// CHECK: [[META0:![0-9]+]] = !{!"{{.*}}clang version {{.*}}"}
//.
diff --git a/clang/test/CodeGenCXX/dynamic-cast-address-space.cpp b/clang/test/CodeGenCXX/dynamic-cast-address-space.cpp
index a398a622606ba..6dffaaa09d4ee 100644
--- a/clang/test/CodeGenCXX/dynamic-cast-address-space.cpp
+++ b/clang/test/CodeGenCXX/dynamic-cast-address-space.cpp
@@ -101,15 +101,15 @@ const B& f(A *a) {
//.
-// CHECK: attributes #[[ATTR0]] = { convergent mustprogress noinline optnone "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
+// CHECK: attributes #[[ATTR0]] = { convergent mustprogress noinline optnone "no-trapping-math"="true" "stack-protector-buffer-size"="8" "uniform-work-group-size" }
// CHECK: attributes #[[ATTR1:[0-9]+]] = { nounwind willreturn memory(read) }
-// CHECK: attributes #[[ATTR2:[0-9]+]] = { convergent "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
+// CHECK: attributes #[[ATTR2:[0-9]+]] = { convergent "no-trapping-math"="true" "stack-protector-buffer-size"="8" "uniform-work-group-size" }
// CHECK: attributes #[[ATTR3]] = { nounwind }
// CHECK: attributes #[[ATTR4]] = { noreturn }
//.
-// WITH-NONZERO-DEFAULT-AS: attributes #[[ATTR0]] = { mustprogress noinline optnone "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
+// WITH-NONZERO-DEFAULT-AS: attributes #[[ATTR0]] = { mustprogress noinline optnone "no-trapping-math"="true" "stack-protector-buffer-size"="8" "uniform-work-group-size" }
// WITH-NONZERO-DEFAULT-AS: attributes #[[ATTR1:[0-9]+]] = { nounwind willreturn memory(read) }
-// WITH-NONZERO-DEFAULT-AS: attributes #[[ATTR2:[0-9]+]] = { "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
+// WITH-NONZERO-DEFAULT-AS: attributes #[[ATTR2:[0-9]+]] = { "no-trapping-math"="true" "stack-protector-buffer-size"="8" "uniform-work-group-size" }
// WITH-NONZERO-DEFAULT-AS: attributes #[[ATTR3]] = { nounwind }
// WITH-NONZERO-DEFAULT-AS: attributes #[[ATTR4]] = { noreturn }
//.
diff --git a/clang/test/CodeGenSYCL/function-attrs.cpp b/clang/test/CodeGenSYCL/function-attrs.cpp
index 60d3cf10055ec..5a664889c77e7 100644
--- a/clang/test/CodeGenSYCL/function-attrs.cpp
+++ b/clang/test/CodeGenSYCL/function-attrs.cpp
@@ -39,6 +39,6 @@ int main() {
return 0;
}
//.
-// CHECK: attributes #[[ATTR0]] = { convergent mustprogress noinline norecurse nounwind optnone "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
-// CHECK: attributes #[[ATTR1]] = { convergent nounwind }
+// CHECK: attributes #[[ATTR0]] = { convergent mustprogress noinline norecurse nounwind optnone "no-trapping-math"="true" "stack-protector-buffer-size"="8" "uniform-work-group-size" }
+// CHECK: attributes #[[ATTR1]] = { convergent nounwind "uniform-work-group-size" }
//.
diff --git a/clang/test/Headers/gpuintrin.c b/clang/test/Headers/gpuintrin.c
index 7c0141230689e..632b5011f4142 100644
--- a/clang/test/Headers/gpuintrin.c
+++ b/clang/test/Headers/gpuintrin.c
@@ -88,15 +88,10 @@ __gpu_kernel void foo() {
// AMDGPU-NEXT: [[TMP1:%.*]] = getelementptr i8, ptr addrspace(4) [[TMP0]], i32 12
// AMDGPU-NEXT: [[TMP2:%.*]] = load i32, ptr addrspace(4) [[TMP1]], align 4, !range [[RNG2:![0-9]+]], !invariant.load [[META3:![0-9]+]]
// AMDGPU-NEXT: [[TMP3:%.*]] = call align 8 dereferenceable(256) ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr()
-// AMDGPU-NEXT: [[TMP4:%.*]] = getelementptr i8, ptr addrspace(4) [[TMP3]], i64 0
-// AMDGPU-NEXT: [[TMP5:%.*]] = load i32, ptr addrspace(4) [[TMP4]], align 4, !invariant.load [[META3]], !noundef [[META3]]
-// AMDGPU-NEXT: [[TMP6:%.*]] = call i32 @llvm.amdgcn.workgroup.id.x()
-// AMDGPU-NEXT: [[TMP7:%.*]] = icmp ult i32 [[TMP6]], [[TMP5]]
-// AMDGPU-NEXT: [[TMP8:%.*]] = select i1 [[TMP7]], i32 12, i32 18
-// AMDGPU-NEXT: [[TMP9:%.*]] = getelementptr inbounds i8, ptr addrspace(4) [[TMP3]], i32 [[TMP8]]
-// AMDGPU-NEXT: [[TMP10:%.*]] = load i16, ptr addrspace(4) [[TMP9]], align 2, !range [[RNG4:![0-9]+]], !invariant.load [[META3]], !noundef [[META3]]
-// AMDGPU-NEXT: [[TMP11:%.*]] = zext i16 [[TMP10]] to i32
-// AMDGPU-NEXT: [[DIV:%.*]] = udiv i32 [[TMP2]], [[TMP11]]
+// AMDGPU-NEXT: [[TMP4:%.*]] = getelementptr inbounds i8, ptr addrspace(4) [[TMP3]], i64 12
+// AMDGPU-NEXT: [[TMP5:%.*]] = load i16, ptr addrspace(4) [[TMP4]], align 2, !range [[RNG4:![0-9]+]], !invariant.load [[META3]], !noundef [[META3]]
+// AMDGPU-NEXT: [[TMP6:%.*]] = zext i16 [[TMP5]] to i32
+// AMDGPU-NEXT: [[DIV:%.*]] = udiv i32 [[TMP2]], [[TMP6]]
// AMDGPU-NEXT: ret i32 [[DIV]]
//
//
@@ -107,15 +102,10 @@ __gpu_kernel void foo() {
// AMDGPU-NEXT: [[TMP1:%.*]] = getelementptr i8, ptr addrspace(4) [[TMP0]], i32 16
// AMDGPU-NEXT: [[TMP2:%.*]] = load i32, ptr addrspace(4) [[TMP1]], align 4, !range [[RNG2]], !invariant.load [[META3]]
// AMDGPU-NEXT: [[TMP3:%.*]] = call align 8 dereferenceable(256) ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr()
-// AMDGPU-NEXT: [[TMP4:%.*]] = getelementptr i8, ptr addrspace(4) [[TMP3]], i64 4
-// AMDGPU-NEXT: [[TMP5:%.*]] = load i32, ptr addrspace(4) [[TMP4]], align 4, !invariant.load [[META3]], !noundef [[META3]]
-// AMDGPU-NEXT: [[TMP6:%.*]] = call i32 @llvm.amdgcn.workgroup.id.y()
-// AMDGPU-NEXT: [[TMP7:%.*]] = icmp ult i32 [[TMP6]], [[TMP5]]
-// AMDGPU-NEXT: [[TMP8:%.*]] = select i1 [[TMP7]], i32 14, i32 20
-// AMDGPU-NEXT: [[TMP9:%.*]] = getelementptr inbounds i8, ptr addrspace(4) [[TMP3]], i32 [[TMP8]]
-// AMDGPU-NEXT: [[TMP10:%.*]] = load i16, ptr addrspace(4) [[TMP9]], align 2, !range [[RNG4]], !invariant.load [[META3]], !noundef [[META3]]
-// AMDGPU-NEXT: [[TMP11:%.*]] = zext i16 [[TMP10]] to i32
-// AMDGPU-NEXT: [[DIV:%.*]] = udiv i32 [[TMP2]], [[TMP11]]
+// AMDGPU-NEXT: [[TMP4:%.*]] = getelementptr inbounds i8, ptr addrspace(4) [[TMP3]], i64 14
+// AMDGPU-NEXT: [[TMP5:%.*]] = load i16, ptr addrspace(4) [[TMP4]], align 2, !range [[RNG4]], !invariant.load [[META3]], !noundef [[META3]]
+// AMDGPU-NEXT: [[TMP6:%.*]] = zext i16 [[TMP5]] to i32
+// AMDGPU-NEXT: [[DIV:%.*]] = udiv i32 [[TMP2]], [[TMP6]]
// AMDGPU-NEXT: ret i32 [[DIV]]
//
//
@@ -126,15 +116,10 @@ __gpu_kernel void foo() {
// AMDGPU-NEXT: [[TMP1:%.*]] = getelementptr i8, ptr addrspace(4) [[TMP0]], i32 20
// AMDGPU-NEXT: [[TMP2:%.*]] = load i32, ptr addrspace(4) [[TMP1]], align 4, !range [[RNG2]], !invariant.load [[META3]]
// AMDGPU-NEXT: [[TMP3:%.*]] = call align 8 dereferenceable(256) ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr()
-// AMDGPU-NEXT: [[TMP4:%.*]] = getelementptr i8, ptr addrspace(4) [[TMP3]], i64 8
-// AMDGPU-NEXT: [[TMP5:%.*]] = load i32, ptr addrspace(4) [[TMP4]], align 4, !invariant.load [[META3]], !noundef [[META3]]
-// AMDGPU-NEXT: [[TMP6:%.*]] = call i32 @llvm.amdgcn.workgroup.id.z()
-// AMDGPU-NEXT: [[TMP7:%.*]] = icmp ult i32 [[TMP6]], [[TMP5]]
-// AMDGPU-NEXT: [[TMP8:%.*]] = select i1 [[TMP7]], i32 16, i32 22
-// AMDGPU-NEXT: [[TMP9:%.*]] = getelementptr inbounds i8, ptr addrspace(4) [[TMP3]], i32 [[TMP8]]
-// AMDGPU-NEXT: [[TMP10:%.*]] = load i16, ptr addrspace(4) [[TMP9]], align 2, !range [[RNG4]], !invariant.load [[META3]], !noundef [[META3]]
-// AMDGPU-NEXT: [[TMP11:%.*]] = zext i16 [[TMP10]] to i32
-// AMDGPU-NEXT: [[DIV:%.*]] = udiv i32 [[TMP2]], [[TMP11]]
+// AMDGPU-NEXT: [[TMP4:%.*]] = getelementptr inbounds i8, ptr addrspace(4) [[TMP3]], i64 16
+// AMDGPU-NEXT: [[TMP5:%.*]] = load i16, ptr addrspace(4) [[TMP4]], align 2, !range [[RNG4]], !invariant.load [[META3]], !noundef [[META3]]
+// AMDGPU-NEXT: [[TMP6:%.*]] = zext i16 [[TMP5]] to i32
+// AMDGPU-NEXT: [[DIV:%.*]] = udiv i32 [[TMP2]], [[TMP6]]
// AMDGPU-NEXT: ret i32 [[DIV]]
//
//
@@ -229,45 +214,30 @@ __gpu_kernel void foo() {
// AMDGPU-SAME: ) #[[ATTR0]] {
// AMDGPU-NEXT: [[ENTRY:.*:]]
// AMDGPU-NEXT: [[TMP0:%.*]] = call align 8 dereferenceable(256) ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr()
-// AMDGPU-NEXT: [[TMP1:%.*]] = getelementptr i8, ptr addrspace(4) [[TMP0]], i64 0
-// AMDGPU-NEXT: [[TMP2:%.*]] = load i32, ptr addrspace(4) [[TMP1]], align 4, !invariant.load [[META3]], !noundef [[META3]]
-// AMDGPU-NEXT: [[TMP3:%.*]] = call i32 @llvm.amdgcn.workgroup.id.x()
-// AMDGPU-NEXT: [[TMP4:%.*]] = icmp ult i32 [[TMP3]], [[TMP2]]
-// AMDGPU-NEXT: [[TMP5:%.*]] = select i1 [[TMP4]], i32 12, i32 18
-// AMDGPU-NEXT: [[TMP6:%.*]] = getelementptr inbounds i8, ptr addrspace(4) [[TMP0]], i32 [[TMP5]]
-// AMDGPU-NEXT: [[TMP7:%.*]] = load i16, ptr addrspace(4) [[TMP6]], align 2, !range [[RNG4]], !invariant.load [[META3]], !noundef [[META3]]
-// AMDGPU-NEXT: [[TMP8:%.*]] = zext i16 [[TMP7]] to i32
-// AMDGPU-NEXT: ret i32 [[TMP8]]
+// AMDGPU-NEXT: [[TMP1:%.*]] = getelementptr inbounds i8, ptr addrspace(4) [[TMP0]], i64 12
+// AMDGPU-NEXT: [[TMP2:%.*]] = load i16, ptr addrspace(4) [[TMP1]], align 2, !range [[RNG4]], !invariant.load [[META3]], !noundef [[META3]]
+// AMDGPU-NEXT: [[TMP3:%.*]] = zext i16 [[TMP2]] to i32
+// AMDGPU-NEXT: ret i32 [[TMP3]]
//
//
// AMDGPU-LABEL: define internal i32 @__gpu_num_threads_y(
// AMDGPU-SAME: ) #[[ATTR0]] {
// AMDGPU-NEXT: [[ENTRY:.*:]]
// AMDGPU-NEXT: [[TMP0:%.*]] = call align 8 dereferenceable(256) ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr()
-// AMDGPU-NEXT: [[TMP1:%.*]] = getelementptr i8, ptr addrspace(4) [[TMP0]], i64 4
-// AMDGPU-NEXT: [[TMP2:%.*]] = load i32, ptr addrspace(4) [[TMP1]], align 4, !invariant.load [[META3]], !noundef [[META3]]
-// AMDGPU-NEXT: [[TMP3:%.*]] = call i32 @llvm.amdgcn.workgroup.id.y()
-// AMDGPU-NEXT: [[TMP4:%.*]] = icmp ult i32 [[TMP3]], [[TMP2]]
-// AMDGPU-NEXT: [[TMP5:%.*]] = select i1 [[TMP4]], i32 14, i32 20
-// AMDGPU-NEXT: [[TMP6:%.*]] = getelementptr inbounds i8, ptr addrspace(4) [[TMP0]], i32 [[TMP5]]
-// AMDGPU-NEXT: [[TMP7:%.*]] = load i16, ptr addrspace(4) [[TMP6]], align 2, !range [[RNG4]], !invariant.load [[META3]], !noundef [[META3]]
-// AMDGPU-NEXT: [[TMP8:%.*]] = zext i16 [[TMP7]] to i32
-// AMDGPU-NEXT: ret i32 [[TMP8]]
+// AMDGPU-NEXT: [[TMP1:%.*]] = getelementptr inbounds i8, ptr addrspace(4) [[TMP0]], i64 14
+// AMDGPU-NEXT: [[TMP2:%.*]] = load i16, ptr addrspace(4) [[TMP1]], align 2, !range [[RNG4]], !invariant.load [[META3]], !noundef [[META3]]
+// AMDGPU-NEXT: [[TMP3:%.*]] = zext i16 [[TMP2]] to i32
+// AMDGPU-NEXT: ret i32 [[TMP3]]
//
//
// AMDGPU-LABEL: define internal i32 @__gpu_num_threads_z(
// AMDGPU-SAME: ) #[[ATTR0]] {
// AMDGPU-NEXT: [[ENTRY:.*:]]
// AMDGPU-NEXT: [[TMP0:%.*]] = call align 8 dereferenceable(256) ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr()
-// AMDGPU-NEXT: [[TMP1:%.*]] = getelementptr i8, ptr addrspace(4) [[TMP0]], i64 8
-// AMDGPU-NEXT: [[TMP2:%.*]] = load i32, ptr addrspace(4) [[TMP1]], align 4, !invariant.load [[META3]], !noundef [[META3]]
-// AMDGPU-NEXT: [[TMP3:%.*]] = call i32 @llvm.amdgcn.workgroup.id.z()
-// AMDGPU-NEXT: [[TMP4:%.*]] = icmp ult i32 [[TMP3]], [[TMP2]]
-// AMDGPU-NEXT: [[TMP5:%.*]] = select i1 [[TMP4]], i32 16, i32 22
-// AMDGPU-NEXT: [[TMP6:%.*]] = getelementptr inbounds i8, ptr addrspace(4) [[TMP0]], i32 [[TMP5]]
-// AMDGPU-NEXT: [[TMP7:%.*]] = load i16, ptr addrspace(4) [[TMP6]], align 2, !range [[RNG4]], !invariant.load [[META3]], !noundef [[META3]]
-// AMDGPU-NEXT: [[TMP8:%.*]] = zext i16 [[TMP7]] to i32
-// AMDGPU-NEXT: ret i32 [[TMP8]]
+// AMDGPU-NEXT: [[TMP1:%.*]] = getelementptr inbounds i8, ptr addrspace(4) [[TMP0]], i64 16
+// AMDGPU-NEXT: [[TMP2:%.*]] = load i16, ptr addrspace(4) [[TMP1]], align 2, !range [[RNG4]], !invariant.load [[META3]], !noundef [[META3]]
+// AMDGPU-NEXT: [[TMP3:%.*]] = zext i16 [[TMP2]] to i32
+// AMDGPU-NEXT: ret i32 [[TMP3]]
//
//
// AMDGPU-LABEL: define internal i32 @__gpu_num_threads(
@@ -1691,7 +1661,7 @@ __gpu_kernel void foo() {
// SPIRV-LABEL: define internal spir_func void @__gpu_exit(
// SPIRV-SAME: ) #[[ATTR1:[0-9]+]] {
// SPIRV-NEXT: [[ENTRY:.*:]]
-// SPIRV-NEXT: call void @llvm.trap()
+// SPIRV-NEXT: call void @llvm.trap() #[[ATTR8:[0-9]+]]
// SPIRV-NEXT: unreachable
//
//.
diff --git a/clang/test/OpenMP/amdgcn-attributes.cpp b/clang/test/OpenMP/amdgcn-attributes.cpp
index 789c8d6d72218..15588c0befc63 100644
--- a/clang/test/OpenMP/amdgcn-attributes.cpp
+++ b/clang/test/OpenMP/amdgcn-attributes.cpp
@@ -34,5 +34,5 @@ int callable(int x) {
// DEFAULT: attributes #0 = { convergent mustprogress noinline norecurse nounwind optnone "amdgpu-flat-work-group-size"="1,42" "kernel" "no-trapping-math"="true" "omp_target_thread_limit"="42" "stack-protector-buffer-size"="8" "uniform-work-group-size" }
// NOIEEE: attributes #0 = { convergent mustprogress noinline norecurse nounwind optnone "amdgpu-flat-work-group-size"="1,42" "amdgpu-ieee"="false" "kernel" "no-trapping-math"="true" "omp_target_thread_limit"="42" "stack-protector-buffer-size"="8" "uniform-work-group-size" }
-// DEFAULT: attributes #2 = { convergent mustprogress noinline nounwind optnone "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
-// NOIEEE: attributes #2 = { convergent mustprogress noinline nounwind optnone "amdgpu-ieee"="false" "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
+// DEFAULT: attributes #2 = { convergent mustprogress noinline nounwind optnone "no-trapping-math"="true" "stack-protector-buffer-size"="8" "uniform-work-group-size" }
+// NOIEEE: attributes #2 = { convergent mustprogress noinline nounwind optnone "amdgpu-ieee"="false" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "uniform-work-group-size" }
diff --git a/clang/test/OpenMP/amdgcn_target_global_constructor.cpp b/clang/test/OpenMP/amdgcn_target_global_constructor.cpp
index 7f3dc4c6e63fd..0e6161c060a93 100644
--- a/clang/test/OpenMP/amdgcn_target_global_constructor.cpp
+++ b/clang/test/OpenMP/amdgcn_target_global_constructor.cpp
@@ -97,10 +97,10 @@ S A;
//
//.
// CHECK: attributes #[[ATTR0]] = { convergent noinline nounwind "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
-// CHECK: attributes #[[ATTR1]] = { convergent mustprogress noinline nounwind optnone "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
-// CHECK: attributes #[[ATTR2:[0-9]+]] = { convergent "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
-// CHECK: attributes #[[ATTR3]] = { convergent }
-// CHECK: attributes #[[ATTR4]] = { convergent nounwind }
+// CHECK: attributes #[[ATTR1]] = { convergent mustprogress noinline nounwind optnone "no-trapping-math"="true" "stack-protector-buffer-size"="8" "uniform-work-group-size" }
+// CHECK: attributes #[[ATTR2:[0-9]+]] = { convergent "no-trapping-math"="true" "stack-protector-buffer-size"="8" "uniform-work-group-size" }
+// CHECK: attributes #[[ATTR3]] = { convergent "uniform-work-group-size" }
+// CHECK: attributes #[[ATTR4]] = { convergent nounwind "uniform-work-group-size" }
//.
// CHECK: [[META0:![0-9]+]] = !{i32 1, !"A", i32 0, i32 0}
// CHECK: [[META1:![0-9]+]] = !{i32 1, !"amdhsa_code_object_version", i32 600}
diff --git a/clang/test/OpenMP/amdgcn_weak_alias.c b/clang/test/OpenMP/amdgcn_weak_alias.c
index 3ad72a41176bc..1bd2468558559 100644
--- a/clang/test/OpenMP/amdgcn_weak_alias.c
+++ b/clang/test/OpenMP/amdgcn_weak_alias.c
@@ -98,7 +98,7 @@ extern int __attribute__((alias("__Three_var"))) Three_var_;
//.
// HOST: attributes #[[ATTR0]] = { noinline nounwind optnone "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-features"="+cx8,+mmx,+sse,+sse2,+x87" }
//.
-// DEVICE: attributes #[[ATTR0]] = { convergent noinline nounwind optnone "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
+// DEVICE: attributes #[[ATTR0]] = { convergent noinline nounwind optnone "no-trapping-math"="true" "stack-protector-buffer-size"="8" "uniform-work-group-size" }
//.
// HOST: [[META0:![0-9]+]] = !{i32 1, !"__Two_var", i32 0, i32 0}
// HOST: [[META1:![0-9]+]] = !{i32 1, !"__Three_var", i32 0, i32 1}
diff --git a/clang/test/OpenMP/amdgcn_weak_alias.cpp b/clang/test/OpenMP/amdgcn_weak_alias.cpp
index f7ce91e8231c2..140dd58f1cd7f 100644
--- a/clang/test/OpenMP/amdgcn_weak_alias.cpp
+++ b/clang/test/OpenMP/amdgcn_weak_alias.cpp
@@ -115,7 +115,7 @@ int Four_(void) __attribute__((alias("_Z6__Fourv")));
//.
// HOST: attributes #[[ATTR0]] = { mustprogress noinline nounwind optnone "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-features"="+cx8,+mmx,+sse,+sse2,+x87" }
//.
-// DEVICE: attributes #[[ATTR0]] = { convergent mustprogress noinline nounwind optnone "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
+// DEVICE: attributes #[[ATTR0]] = { convergent mustprogress noinline nounwind optnone "no-trapping-math"="true" "stack-protector-buffer-size"="8" "uniform-work-group-size" }
//.
// HOST: [[META0:![0-9]+]] = !{i32 7, !"openmp", i32 51}
// HOST: [[META1:![0-9]+]] = !{!"{{.*}}clang version {{.*}}"}
>From 3a9c31b3b3f7bf6989df61da8e9f0f551f765757 Mon Sep 17 00:00:00 2001
From: Joseph Huber <huberjn at outlook.com>
Date: Fri, 4 Sep 2026 11:52:16 -0500
Subject: [PATCH 2/5] attempt to use language
---
clang/include/clang/Basic/LangOptions.def | 2 +-
clang/include/clang/Basic/LangOptions.h | 5 +++--
clang/include/clang/Options/Options.td | 4 ++--
clang/lib/Basic/LangOptions.cpp | 7 ++-----
clang/lib/CodeGen/CGCall.cpp | 11 +++++------
clang/lib/Frontend/CompilerInvocation.cpp | 3 ---
clang/test/CodeGen/amdgpu-offload-uniform-block.c | 4 ++++
7 files changed, 17 insertions(+), 19 deletions(-)
diff --git a/clang/include/clang/Basic/LangOptions.def b/clang/include/clang/Basic/LangOptions.def
index 6016f0003ca28..b0c7aba3d147a 100644
--- a/clang/include/clang/Basic/LangOptions.def
+++ b/clang/include/clang/Basic/LangOptions.def
@@ -272,7 +272,7 @@ LANGOPT(SYCLIsHost , 1, 0, NotCompatible, "SYCL host compilation")
ENUM_LANGOPT(SYCLVersion , SYCLMajorVersion, 2, SYCL_None, NotCompatible, "Version of the SYCL standard used")
LANGOPT(HIPUseNewLaunchAPI, 1, 0, NotCompatible, "Use new kernel launching API for HIP")
-LANGOPT(OffloadUniformBlock, 1, 0, NotCompatible, "Assume that kernels are launched with uniform block sizes (default true for CUDA/HIP, OpenMP device, SYCL device, AMDGPU/NVPTX C/C++, and OpenCL 1.2)")
+LANGOPT(OffloadUniformBlock, 1, 0, NotCompatible, "Assume that kernels are launched with uniform block sizes (default true except OpenCL 2.0+)")
LANGOPT(HIPStdPar, 1, 0, NotCompatible, "Enable Standard Parallel Algorithm Acceleration for HIP (experimental)")
LANGOPT(HIPStdParInterposeAlloc, 1, 0, NotCompatible, "Replace allocations / deallocations with HIP RT calls when Standard Parallel Algorithm Acceleration for HIP is enabled (Experimental)")
diff --git a/clang/include/clang/Basic/LangOptions.h b/clang/include/clang/Basic/LangOptions.h
index 2b664818d0a48..5886d3d51bdbc 100644
--- a/clang/include/clang/Basic/LangOptions.h
+++ b/clang/include/clang/Basic/LangOptions.h
@@ -823,8 +823,9 @@ class LangOptions : public LangOptionsBase {
return OpenMPIsTargetDevice || CUDAIsDevice || SYCLIsDevice;
}
- /// Languages that launch full work-groups opt in.
- bool defaultOffloadUniformBlock(const llvm::Triple &T) const;
+ /// Languages that launch full work-groups opt in; leftover groups are
+ /// OpenCL 2.0+.
+ bool defaultOffloadUniformBlock() const;
/// Return the OpenMP version.
llvm::omp::Version getOpenMPVersion() const {
diff --git a/clang/include/clang/Options/Options.td b/clang/include/clang/Options/Options.td
index 605b791f053dc..a575abd600a33 100644
--- a/clang/include/clang/Options/Options.td
+++ b/clang/include/clang/Options/Options.td
@@ -1148,10 +1148,10 @@ def b : JoinedOrSeparate<["-"], "b">, Flags<[LinkerInput]>,
defm offload_uniform_block : BoolFOption<"offload-uniform-block",
LangOpts<"OffloadUniformBlock">,
- Default<"LangOpts->defaultOffloadUniformBlock(T)">,
+ Default<"LangOpts->defaultOffloadUniformBlock()">,
PosFlag<SetTrue, [], [ClangOption, CC1Option], "Assume">,
NegFlag<SetFalse, [], [ClangOption, CC1Option], "Don't assume">,
- BothFlags<[], [ClangOption], " that kernels are launched with uniform block sizes (default true for CUDA/HIP, OpenMP device, SYCL device, AMDGPU/NVPTX C/C++, and OpenCL 1.2)">>;
+ BothFlags<[], [ClangOption], " that kernels are launched with uniform block sizes (default true except OpenCL 2.0+)">>;
def fcomplex_arithmetic_EQ : Joined<["-"], "fcomplex-arithmetic=">, Group<f_Group>,
Visibility<[ClangOption, CC1Option, FlangOption, FC1Option]>,
diff --git a/clang/lib/Basic/LangOptions.cpp b/clang/lib/Basic/LangOptions.cpp
index 3842d307c7e48..8295365eef969 100644
--- a/clang/lib/Basic/LangOptions.cpp
+++ b/clang/lib/Basic/LangOptions.cpp
@@ -74,13 +74,10 @@ unsigned LangOptions::getOpenCLCompatibleVersion() const {
llvm_unreachable("Unknown OpenCL version");
}
-bool LangOptions::defaultOffloadUniformBlock(const llvm::Triple &T) const {
+bool LangOptions::defaultOffloadUniformBlock() const {
if (OpenCL)
return getOpenCLCompatibleVersion() <= 120;
- if (CUDA || OpenMPIsTargetDevice || SYCLIsDevice)
- return true;
- // Direct C/C++ for AMDGPU/NVPTX.
- return !HLSL && (T.isAMDGPU() || T.isNVPTX());
+ return !HLSL;
}
void LangOptions::remapPathPrefix(SmallVectorImpl<char> &Path) const {
diff --git a/clang/lib/CodeGen/CGCall.cpp b/clang/lib/CodeGen/CGCall.cpp
index 1221829871b9f..51d3e379a1452 100644
--- a/clang/lib/CodeGen/CGCall.cpp
+++ b/clang/lib/CodeGen/CGCall.cpp
@@ -2865,12 +2865,11 @@ void CodeGenModule::ConstructAttributeList(StringRef Name,
NumElemsParam);
}
- // OpenCL v2.0 Work groups may be whether uniform or not.
- // '-cl-uniform-work-group-size' compile option gets a hint
- // to the compiler that the global work-size be a multiple of
- // the work-group size specified to clEnqueueNDRangeKernel
- // (i.e. work groups are uniform).
- if (getLangOpts().OffloadUniformBlock)
+ // Leftover work-groups are an OpenCL 2.0+ language feature. Host C/C++
+ // inherit OffloadUniformBlock but must not receive the GPU ABI attribute.
+ if (getLangOpts().OffloadUniformBlock &&
+ (getLangOpts().CUDA || getLangOpts().OpenCL ||
+ getLangOpts().isTargetDevice() || getTarget().getTriple().isGPU()))
FuncAttrs.addAttribute("uniform-work-group-size");
if (TargetDecl->hasAttr<ArmLocallyStreamingAttr>())
diff --git a/clang/lib/Frontend/CompilerInvocation.cpp b/clang/lib/Frontend/CompilerInvocation.cpp
index 2089426582f1b..ea8368908879a 100644
--- a/clang/lib/Frontend/CompilerInvocation.cpp
+++ b/clang/lib/Frontend/CompilerInvocation.cpp
@@ -4349,9 +4349,6 @@ bool CompilerInvocation::ParseLangArgs(LangOptions &Opts, ArgList &Args,
Opts.OpenMP && !Args.hasArg(options::OPT_fnoopenmp_use_tls);
Opts.OpenMPIsTargetDevice =
Opts.OpenMP && Args.hasArg(options::OPT_fopenmp_is_target_device);
- if (!Args.getLastArg(OPT_foffload_uniform_block,
- OPT_fno_offload_uniform_block))
- Opts.OffloadUniformBlock = Opts.defaultOffloadUniformBlock(T);
Opts.OpenMPIRBuilder =
Opts.OpenMP && Args.hasArg(options::OPT_fopenmp_enable_irbuilder);
bool IsTargetSpecified =
diff --git a/clang/test/CodeGen/amdgpu-offload-uniform-block.c b/clang/test/CodeGen/amdgpu-offload-uniform-block.c
index dae42a7db5901..08f1f4c1c6107 100644
--- a/clang/test/CodeGen/amdgpu-offload-uniform-block.c
+++ b/clang/test/CodeGen/amdgpu-offload-uniform-block.c
@@ -2,6 +2,8 @@
// RUN: | FileCheck %s --check-prefixes=CHECK,UNIFORM
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fno-offload-uniform-block \
// RUN: -emit-llvm %s -o - | FileCheck %s --check-prefixes=CHECK,REMAINDER
+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -emit-llvm %s -o - \
+// RUN: | FileCheck %s --check-prefix=HOST
#ifdef __AMDGPU__
int foo(void) { return __builtin_amdgcn_workgroup_size_x(); }
@@ -15,3 +17,5 @@ int foo(void) { return 0; }
// UNIFORM: "uniform-work-group-size"
// REMAINDER: select i1 {{.*}}, i32 12, i32 18
// REMAINDER-NOT: "uniform-work-group-size"
+// HOST-LABEL: define{{.*}} i32 @foo(
+// HOST-NOT: "uniform-work-group-size"
>From b7e662efbecf744b99e44dd1838faa9368d5b300 Mon Sep 17 00:00:00 2001
From: Joseph Huber <huberjn at outlook.com>
Date: Fri, 4 Sep 2026 15:53:58 -0500
Subject: [PATCH 3/5] Try to simplify description
---
clang/include/clang/Basic/LangOptions.def | 2 +-
clang/include/clang/Basic/LangOptions.h | 4 ++--
clang/include/clang/Options/Options.td | 2 +-
3 files changed, 4 insertions(+), 4 deletions(-)
diff --git a/clang/include/clang/Basic/LangOptions.def b/clang/include/clang/Basic/LangOptions.def
index b0c7aba3d147a..d1f4a81a98655 100644
--- a/clang/include/clang/Basic/LangOptions.def
+++ b/clang/include/clang/Basic/LangOptions.def
@@ -272,7 +272,7 @@ LANGOPT(SYCLIsHost , 1, 0, NotCompatible, "SYCL host compilation")
ENUM_LANGOPT(SYCLVersion , SYCLMajorVersion, 2, SYCL_None, NotCompatible, "Version of the SYCL standard used")
LANGOPT(HIPUseNewLaunchAPI, 1, 0, NotCompatible, "Use new kernel launching API for HIP")
-LANGOPT(OffloadUniformBlock, 1, 0, NotCompatible, "Assume that kernels are launched with uniform block sizes (default true except OpenCL 2.0+)")
+LANGOPT(OffloadUniformBlock, 1, 0, NotCompatible, "Assume that kernels are launched with uniform block sizes")
LANGOPT(HIPStdPar, 1, 0, NotCompatible, "Enable Standard Parallel Algorithm Acceleration for HIP (experimental)")
LANGOPT(HIPStdParInterposeAlloc, 1, 0, NotCompatible, "Replace allocations / deallocations with HIP RT calls when Standard Parallel Algorithm Acceleration for HIP is enabled (Experimental)")
diff --git a/clang/include/clang/Basic/LangOptions.h b/clang/include/clang/Basic/LangOptions.h
index 5886d3d51bdbc..bca51263049ea 100644
--- a/clang/include/clang/Basic/LangOptions.h
+++ b/clang/include/clang/Basic/LangOptions.h
@@ -823,8 +823,8 @@ class LangOptions : public LangOptionsBase {
return OpenMPIsTargetDevice || CUDAIsDevice || SYCLIsDevice;
}
- /// Languages that launch full work-groups opt in; leftover groups are
- /// OpenCL 2.0+.
+ /// Default for OffloadUniformBlock: false for OpenCL 2.0+ and HLSL, true
+ /// otherwise. OpenCL may opt in with -cl-uniform-work-group-size.
bool defaultOffloadUniformBlock() const;
/// Return the OpenMP version.
diff --git a/clang/include/clang/Options/Options.td b/clang/include/clang/Options/Options.td
index a575abd600a33..72590ab683007 100644
--- a/clang/include/clang/Options/Options.td
+++ b/clang/include/clang/Options/Options.td
@@ -1151,7 +1151,7 @@ defm offload_uniform_block : BoolFOption<"offload-uniform-block",
Default<"LangOpts->defaultOffloadUniformBlock()">,
PosFlag<SetTrue, [], [ClangOption, CC1Option], "Assume">,
NegFlag<SetFalse, [], [ClangOption, CC1Option], "Don't assume">,
- BothFlags<[], [ClangOption], " that kernels are launched with uniform block sizes (default true except OpenCL 2.0+)">>;
+ BothFlags<[], [ClangOption], " that kernels are launched with uniform block sizes">>;
def fcomplex_arithmetic_EQ : Joined<["-"], "fcomplex-arithmetic=">, Group<f_Group>,
Visibility<[ClangOption, CC1Option, FlangOption, FC1Option]>,
>From 1ac05f0c8eb9ab1130567d6ba54ac75caa834daa Mon Sep 17 00:00:00 2001
From: Joseph Huber <huberjn at outlook.com>
Date: Fri, 4 Sep 2026 20:22:16 -0500
Subject: [PATCH 4/5] test update
---
clang/test/CodeGen/link-builtin-bitcode.c | 6 +-
.../test/CodeGenHIP/sret-lifetime-markers.cpp | 2 +-
.../CodeGenSYCL/kernel-caller-entry-point.cpp | 6 +-
clang/test/CodeGenSYCL/sycl-module-id.cpp | 7 +-
clang/test/OpenMP/nvptx_parallel_codegen.cpp | 20 +--
...parallel_for_simd_firstprivate_codegen.cpp | 154 +++++++++---------
...bute_parallel_for_simd_private_codegen.cpp | 82 +++++-----
7 files changed, 139 insertions(+), 138 deletions(-)
diff --git a/clang/test/CodeGen/link-builtin-bitcode.c b/clang/test/CodeGen/link-builtin-bitcode.c
index 2cae7c027a4f3..35ad3af939a98 100644
--- a/clang/test/CodeGen/link-builtin-bitcode.c
+++ b/clang/test/CodeGen/link-builtin-bitcode.c
@@ -44,6 +44,6 @@ int bar() { return no_attr() + attr_in_target() + attr_not_in_target() + attr_in
// CHECK-SAME: () #[[ATTR_INCOMPATIBLE:[0-9]+]] {
// CHECK: attributes #[[ATTR_BAR]] = { {{.*}} "no-trapping-math"="true" {{.*}} }
-// CHECK: attributes #[[ATTR_COMPATIBLE]] = { {{.*}} "target-features"="+16-bit-insts,+atomic-buffer-global-pk-add-f16-insts,+atomic-fadd-rtn-insts,+atomic-fmin-fmax-global-f64,+ci-insts,+cube-insts,+cvt-pknorm-vop2-insts,+dl-insts,+dot1-insts,+dot10-insts,+dot2-insts,+dot3-insts,+dot4-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dpp,+flat-buffer-global-fadd-f64-inst,+flat-global-insts,+gfx8-insts,+gfx9-insts,+gfx90a-insts,+gws,+image-insts,+lerp-inst,+mai-insts,+mqsad-insts,+mqsad-pk-insts,+msad-insts,+qsad-insts,+s-memrealtime,+s-memtime-inst,+sad-insts,+vmem-to-lds-load-insts,+wavefrontsize64" }
-// CHECK: attributes #[[ATTR_EXTEND]] = { {{.*}} "target-features"="+16-bit-insts,+atomic-buffer-global-pk-add-f16-insts,+atomic-fadd-rtn-insts,+atomic-fmin-fmax-global-f64,+ci-insts,+cube-insts,+cvt-pknorm-vop2-insts,+dl-insts,+dot1-insts,+dot10-insts,+dot2-insts,+dot3-insts,+dot4-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dot8-insts,+dpp,+flat-buffer-global-fadd-f64-inst,+flat-global-insts,+gfx8-insts,+gfx9-insts,+gfx90a-insts,+gws,+image-insts,+lerp-inst,+mai-insts,+mqsad-insts,+mqsad-pk-insts,+msad-insts,+qsad-insts,+s-memrealtime,+s-memtime-inst,+sad-insts,+vmem-to-lds-load-insts,+wavefrontsize64" }
-// CHECK: attributes #[[ATTR_INCOMPATIBLE]] = { {{.*}} "target-features"="+16-bit-insts,+atomic-buffer-global-pk-add-f16-insts,+atomic-fadd-rtn-insts,+atomic-fmin-fmax-global-f64,+ci-insts,+cube-insts,+cvt-pknorm-vop2-insts,+dl-insts,+dot1-insts,+dot10-insts,+dot2-insts,+dot3-insts,+dot4-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dpp,+flat-buffer-global-fadd-f64-inst,+flat-global-insts,+gfx8-insts,+gfx90a-insts,+gws,+image-insts,+lerp-inst,+mai-insts,+mqsad-insts,+mqsad-pk-insts,+msad-insts,+qsad-insts,+s-memrealtime,+s-memtime-inst,+sad-insts,+vmem-to-lds-load-insts,+wavefrontsize64,-gfx9-insts" }
+// CHECK: attributes #[[ATTR_COMPATIBLE]] = { {{.*}} "target-features"="+16-bit-insts,+atomic-buffer-global-pk-add-f16-insts,+atomic-fadd-rtn-insts,+atomic-fmin-fmax-global-f64,+ci-insts,+cube-insts,+cvt-pknorm-vop2-insts,+dl-insts,+dot1-insts,+dot10-insts,+dot2-insts,+dot3-insts,+dot4-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dpp,+flat-buffer-global-fadd-f64-inst,+flat-global-insts,+gfx8-insts,+gfx9-insts,+gfx90a-insts,+gws,+image-insts,+lerp-inst,+mai-insts,+mqsad-insts,+mqsad-pk-insts,+msad-insts,+qsad-insts,+s-memrealtime,+s-memtime-inst,+sad-insts,+vmem-to-lds-load-insts,+wavefrontsize64" "uniform-work-group-size" }
+// CHECK: attributes #[[ATTR_EXTEND]] = { {{.*}} "target-features"="+16-bit-insts,+atomic-buffer-global-pk-add-f16-insts,+atomic-fadd-rtn-insts,+atomic-fmin-fmax-global-f64,+ci-insts,+cube-insts,+cvt-pknorm-vop2-insts,+dl-insts,+dot1-insts,+dot10-insts,+dot2-insts,+dot3-insts,+dot4-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dot8-insts,+dpp,+flat-buffer-global-fadd-f64-inst,+flat-global-insts,+gfx8-insts,+gfx9-insts,+gfx90a-insts,+gws,+image-insts,+lerp-inst,+mai-insts,+mqsad-insts,+mqsad-pk-insts,+msad-insts,+qsad-insts,+s-memrealtime,+s-memtime-inst,+sad-insts,+vmem-to-lds-load-insts,+wavefrontsize64" "uniform-work-group-size" }
+// CHECK: attributes #[[ATTR_INCOMPATIBLE]] = { {{.*}} "target-features"="+16-bit-insts,+atomic-buffer-global-pk-add-f16-insts,+atomic-fadd-rtn-insts,+atomic-fmin-fmax-global-f64,+ci-insts,+cube-insts,+cvt-pknorm-vop2-insts,+dl-insts,+dot1-insts,+dot10-insts,+dot2-insts,+dot3-insts,+dot4-insts,+dot5-insts,+dot6-insts,+dot7-insts,+dpp,+flat-buffer-global-fadd-f64-inst,+flat-global-insts,+gfx8-insts,+gfx90a-insts,+gws,+image-insts,+lerp-inst,+mai-insts,+mqsad-insts,+mqsad-pk-insts,+msad-insts,+qsad-insts,+s-memrealtime,+s-memtime-inst,+sad-insts,+vmem-to-lds-load-insts,+wavefrontsize64,-gfx9-insts" "uniform-work-group-size" }
diff --git a/clang/test/CodeGenHIP/sret-lifetime-markers.cpp b/clang/test/CodeGenHIP/sret-lifetime-markers.cpp
index 4056f24e7d514..5beaf65d152f4 100644
--- a/clang/test/CodeGenHIP/sret-lifetime-markers.cpp
+++ b/clang/test/CodeGenHIP/sret-lifetime-markers.cpp
@@ -55,7 +55,7 @@ template void add_functor<4>(SuperScalar<4>*, int, SuperScalar<4>&);
// SPIRV-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [32 x i8], ptr addrspace(4) [[DATA]], i64 [[IDXPROM]]
// SPIRV-NEXT: call addrspace(4) void @llvm.lifetime.start.p0(ptr nonnull [[TMP]]) #[[ATTR3:[0-9]+]]
// SPIRV-NEXT: [[TMP_ASCAST:%.*]] = addrspacecast ptr [[TMP]] to ptr addrspace(4)
-// SPIRV-NEXT: call spir_func addrspace(4) void @_Z16atomic_fetch_addILi4EE11SuperScalarIXT_EEPS1_RKS1_(ptr addrspace(4) dead_on_unwind writable sret([[STRUCT_SUPERSCALAR]]) align 8 [[TMP_ASCAST]], ptr addrspace(4) noundef [[ARRAYIDX]], ptr addrspace(4) noundef align 8 dereferenceable(32) [[UPDATE]]) #[[ATTR3]]
+// SPIRV-NEXT: call spir_func addrspace(4) void @_Z16atomic_fetch_addILi4EE11SuperScalarIXT_EEPS1_RKS1_(ptr addrspace(4) dead_on_unwind writable sret([[STRUCT_SUPERSCALAR]]) align 8 [[TMP_ASCAST]], ptr addrspace(4) noundef [[ARRAYIDX]], ptr addrspace(4) noundef align 8 dereferenceable(32) [[UPDATE]]) #[[ATTR4:[0-9]+]]
// SPIRV-NEXT: call addrspace(4) void @llvm.lifetime.end.p0(ptr nonnull [[TMP]]) #[[ATTR3]]
// SPIRV-NEXT: ret void
//
diff --git a/clang/test/CodeGenSYCL/kernel-caller-entry-point.cpp b/clang/test/CodeGenSYCL/kernel-caller-entry-point.cpp
index 0185af123aa8c..e0f3bd5eba52e 100644
--- a/clang/test/CodeGenSYCL/kernel-caller-entry-point.cpp
+++ b/clang/test/CodeGenSYCL/kernel-caller-entry-point.cpp
@@ -703,10 +703,10 @@ int main() {
// CHECK-SPIR-NEXT: }
// CHECK-AMDGCN: #[[AMDGCN_ATTR0]] = { convergent mustprogress noinline norecurse nounwind optnone "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"=
-// CHECK-AMDGCN: #[[AMDGCN_ATTR1]] = { convergent nounwind }
+// CHECK-AMDGCN: #[[AMDGCN_ATTR1]] = { convergent nounwind "uniform-work-group-size" }
//
// CHECK-NVPTX: #[[NVPTX_ATTR0]] = { convergent mustprogress noinline norecurse nounwind optnone "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"=
-// CHECK-NVPTX: #[[NVPTX_ATTR1]] = { convergent nounwind }
+// CHECK-NVPTX: #[[NVPTX_ATTR1]] = { convergent nounwind "uniform-work-group-size" }
//
// CHECK-SPIR: #[[SPIR_ATTR0]] = { convergent mustprogress noinline norecurse nounwind optnone "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"=
-// CHECK-SPIR: #[[SPIR_ATTR1]] = { convergent nounwind }
+// CHECK-SPIR: #[[SPIR_ATTR1]] = { convergent nounwind "uniform-work-group-size" }
diff --git a/clang/test/CodeGenSYCL/sycl-module-id.cpp b/clang/test/CodeGenSYCL/sycl-module-id.cpp
index d9beded116a65..d4d2fd8c1617c 100644
--- a/clang/test/CodeGenSYCL/sycl-module-id.cpp
+++ b/clang/test/CodeGenSYCL/sycl-module-id.cpp
@@ -22,6 +22,7 @@ void use() { kernel_single_task<KN>(K{}); }
// Both the kernel and the sycl_external function must carry sycl-module-id
// with a value equal to the module identifier (the source file path).
// CHECK: source_filename = "[[MODID:.*]]"
-// CHECK-DAG: define {{.*}}spir_kernel {{.*}}@{{.*}}2KN{{.*}} #[[SATTR:[0-9]+]]
-// CHECK-DAG: define {{.*}}spir_func {{.*}}@_Z3exti{{.*}} #[[SATTR]]
-// CHECK: attributes #[[SATTR]] = { {{.*}}"sycl-module-id"="[[MODID]]"
+// CHECK-DAG: define {{.*}}spir_kernel {{.*}}@{{.*}}2KN{{.*}} #[[KATTR:[0-9]+]]
+// CHECK-DAG: define {{.*}}spir_func {{.*}}@_Z3exti{{.*}} #[[SATTR:[0-9]+]]
+// CHECK-DAG: attributes #[[KATTR]] = { {{.*}}"sycl-module-id"="[[MODID]]"
+// CHECK-DAG: attributes #[[SATTR]] = { {{.*}}"sycl-module-id"="[[MODID]]"
diff --git a/clang/test/OpenMP/nvptx_parallel_codegen.cpp b/clang/test/OpenMP/nvptx_parallel_codegen.cpp
index f57a010b38015..ecb4027bf7984 100644
--- a/clang/test/OpenMP/nvptx_parallel_codegen.cpp
+++ b/clang/test/OpenMP/nvptx_parallel_codegen.cpp
@@ -382,7 +382,7 @@ int bar(int n){
//
//
// CHECK2-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l26_omp_outlined_wrapper
-// CHECK2-SAME: (i16 noundef zeroext [[TMP0:%.*]], i32 noundef [[TMP1:%.*]]) #[[ATTR1]] {
+// CHECK2-SAME: (i16 noundef zeroext [[TMP0:%.*]], i32 noundef [[TMP1:%.*]]) #[[ATTR2:[0-9]+]] {
// CHECK2-NEXT: entry:
// CHECK2-NEXT: [[DOTADDR:%.*]] = alloca i16, align 2
// CHECK2-NEXT: [[DOTADDR1:%.*]] = alloca i32, align 4
@@ -392,7 +392,7 @@ int bar(int n){
// CHECK2-NEXT: store i32 [[TMP1]], ptr [[DOTADDR1]], align 4
// CHECK2-NEXT: store i32 0, ptr [[DOTZERO_ADDR]], align 4
// CHECK2-NEXT: call void @__kmpc_get_shared_variables(ptr [[GLOBAL_ARGS]])
-// CHECK2-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l26_omp_outlined(ptr [[DOTADDR1]], ptr [[DOTZERO_ADDR]]) #[[ATTR2:[0-9]+]]
+// CHECK2-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l26_omp_outlined(ptr [[DOTADDR1]], ptr [[DOTZERO_ADDR]]) #[[ATTR3:[0-9]+]]
// CHECK2-NEXT: ret void
//
//
@@ -409,7 +409,7 @@ int bar(int n){
//
//
// CHECK2-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l26_omp_outlined1_wrapper
-// CHECK2-SAME: (i16 noundef zeroext [[TMP0:%.*]], i32 noundef [[TMP1:%.*]]) #[[ATTR1]] {
+// CHECK2-SAME: (i16 noundef zeroext [[TMP0:%.*]], i32 noundef [[TMP1:%.*]]) #[[ATTR2]] {
// CHECK2-NEXT: entry:
// CHECK2-NEXT: [[DOTADDR:%.*]] = alloca i16, align 2
// CHECK2-NEXT: [[DOTADDR1:%.*]] = alloca i32, align 4
@@ -419,7 +419,7 @@ int bar(int n){
// CHECK2-NEXT: store i32 [[TMP1]], ptr [[DOTADDR1]], align 4
// CHECK2-NEXT: store i32 0, ptr [[DOTZERO_ADDR]], align 4
// CHECK2-NEXT: call void @__kmpc_get_shared_variables(ptr [[GLOBAL_ARGS]])
-// CHECK2-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l26_omp_outlined1(ptr [[DOTADDR1]], ptr [[DOTZERO_ADDR]]) #[[ATTR2]]
+// CHECK2-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l26_omp_outlined1(ptr [[DOTADDR1]], ptr [[DOTZERO_ADDR]]) #[[ATTR3]]
// CHECK2-NEXT: ret void
//
//
@@ -436,7 +436,7 @@ int bar(int n){
//
//
// CHECK2-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l26_omp_outlined2_wrapper
-// CHECK2-SAME: (i16 noundef zeroext [[TMP0:%.*]], i32 noundef [[TMP1:%.*]]) #[[ATTR1]] {
+// CHECK2-SAME: (i16 noundef zeroext [[TMP0:%.*]], i32 noundef [[TMP1:%.*]]) #[[ATTR2]] {
// CHECK2-NEXT: entry:
// CHECK2-NEXT: [[DOTADDR:%.*]] = alloca i16, align 2
// CHECK2-NEXT: [[DOTADDR1:%.*]] = alloca i32, align 4
@@ -446,7 +446,7 @@ int bar(int n){
// CHECK2-NEXT: store i32 [[TMP1]], ptr [[DOTADDR1]], align 4
// CHECK2-NEXT: store i32 0, ptr [[DOTZERO_ADDR]], align 4
// CHECK2-NEXT: call void @__kmpc_get_shared_variables(ptr [[GLOBAL_ARGS]])
-// CHECK2-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l26_omp_outlined2(ptr [[DOTADDR1]], ptr [[DOTZERO_ADDR]]) #[[ATTR2]]
+// CHECK2-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l26_omp_outlined2(ptr [[DOTADDR1]], ptr [[DOTZERO_ADDR]]) #[[ATTR3]]
// CHECK2-NEXT: ret void
//
//
@@ -508,7 +508,7 @@ int bar(int n){
//
//
// CHECK2-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l43_omp_outlined_wrapper
-// CHECK2-SAME: (i16 noundef zeroext [[TMP0:%.*]], i32 noundef [[TMP1:%.*]]) #[[ATTR1]] {
+// CHECK2-SAME: (i16 noundef zeroext [[TMP0:%.*]], i32 noundef [[TMP1:%.*]]) #[[ATTR2]] {
// CHECK2-NEXT: entry:
// CHECK2-NEXT: [[DOTADDR:%.*]] = alloca i16, align 2
// CHECK2-NEXT: [[DOTADDR1:%.*]] = alloca i32, align 4
@@ -518,7 +518,7 @@ int bar(int n){
// CHECK2-NEXT: store i32 [[TMP1]], ptr [[DOTADDR1]], align 4
// CHECK2-NEXT: store i32 0, ptr [[DOTZERO_ADDR]], align 4
// CHECK2-NEXT: call void @__kmpc_get_shared_variables(ptr [[GLOBAL_ARGS]])
-// CHECK2-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l43_omp_outlined(ptr [[DOTADDR1]], ptr [[DOTZERO_ADDR]]) #[[ATTR2]]
+// CHECK2-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l43_omp_outlined(ptr [[DOTADDR1]], ptr [[DOTZERO_ADDR]]) #[[ATTR3]]
// CHECK2-NEXT: ret void
//
//
@@ -594,7 +594,7 @@ int bar(int n){
//
//
// CHECK2-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l55_omp_outlined_wrapper
-// CHECK2-SAME: (i16 noundef zeroext [[TMP0:%.*]], i32 noundef [[TMP1:%.*]]) #[[ATTR1]] {
+// CHECK2-SAME: (i16 noundef zeroext [[TMP0:%.*]], i32 noundef [[TMP1:%.*]]) #[[ATTR2]] {
// CHECK2-NEXT: entry:
// CHECK2-NEXT: [[DOTADDR:%.*]] = alloca i16, align 2
// CHECK2-NEXT: [[DOTADDR1:%.*]] = alloca i32, align 4
@@ -607,6 +607,6 @@ int bar(int n){
// CHECK2-NEXT: [[TMP2:%.*]] = load ptr, ptr [[GLOBAL_ARGS]], align 4
// CHECK2-NEXT: [[TMP3:%.*]] = getelementptr inbounds ptr, ptr [[TMP2]], i32 0
// CHECK2-NEXT: [[TMP4:%.*]] = load ptr, ptr [[TMP3]], align 4
-// CHECK2-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l55_omp_outlined(ptr [[DOTADDR1]], ptr [[DOTZERO_ADDR]], ptr [[TMP4]]) #[[ATTR2]]
+// CHECK2-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z9ftemplateIiET_i_l55_omp_outlined(ptr [[DOTADDR1]], ptr [[DOTZERO_ADDR]], ptr [[TMP4]]) #[[ATTR3]]
// CHECK2-NEXT: ret void
//
diff --git a/clang/test/OpenMP/target_teams_distribute_parallel_for_simd_firstprivate_codegen.cpp b/clang/test/OpenMP/target_teams_distribute_parallel_for_simd_firstprivate_codegen.cpp
index 5e7f3689dd05d..87752764636f4 100644
--- a/clang/test/OpenMP/target_teams_distribute_parallel_for_simd_firstprivate_codegen.cpp
+++ b/clang/test/OpenMP/target_teams_distribute_parallel_for_simd_firstprivate_codegen.cpp
@@ -3604,17 +3604,17 @@ int main() {
// CHECK13: omp.arraycpy.body:
// CHECK13-NEXT: [[OMP_ARRAYCPY_SRCELEMENTPAST:%.*]] = phi ptr [ [[TMP1]], [[ENTRY:%.*]] ], [ [[OMP_ARRAYCPY_SRC_ELEMENT:%.*]], [[OMP_ARRAYCPY_BODY]] ]
// CHECK13-NEXT: [[OMP_ARRAYCPY_DESTELEMENTPAST:%.*]] = phi ptr [ [[ARRAY_BEGIN]], [[ENTRY]] ], [ [[OMP_ARRAYCPY_DEST_ELEMENT:%.*]], [[OMP_ARRAYCPY_BODY]] ]
-// CHECK13-NEXT: call void @_ZN2StC1Ev(ptr noundef nonnull align 4 dereferenceable(8) [[AGG_TMP]])
-// CHECK13-NEXT: call void @_ZN1SIfEC1ERKS0_2St(ptr noundef nonnull align 4 dereferenceable(4) [[OMP_ARRAYCPY_DESTELEMENTPAST]], ptr noundef nonnull align 4 dereferenceable(4) [[OMP_ARRAYCPY_SRCELEMENTPAST]], ptr nofreeobj noundef align 4 dereferenceable(8) [[AGG_TMP]])
-// CHECK13-NEXT: call void @_ZN2StD1Ev(ptr noundef nonnull align 4 dead_on_return(8) dereferenceable(8) [[AGG_TMP]]) #[[ATTR3:[0-9]+]]
+// CHECK13-NEXT: call void @_ZN2StC1Ev(ptr noundef nonnull align 4 dereferenceable(8) [[AGG_TMP]]) #[[ATTR4:[0-9]+]]
+// CHECK13-NEXT: call void @_ZN1SIfEC1ERKS0_2St(ptr noundef nonnull align 4 dereferenceable(4) [[OMP_ARRAYCPY_DESTELEMENTPAST]], ptr noundef nonnull align 4 dereferenceable(4) [[OMP_ARRAYCPY_SRCELEMENTPAST]], ptr nofreeobj noundef align 4 dereferenceable(8) [[AGG_TMP]]) #[[ATTR4]]
+// CHECK13-NEXT: call void @_ZN2StD1Ev(ptr noundef nonnull align 4 dead_on_return(8) dereferenceable(8) [[AGG_TMP]]) #[[ATTR5:[0-9]+]]
// CHECK13-NEXT: [[OMP_ARRAYCPY_DEST_ELEMENT]] = getelementptr [[STRUCT_S]], ptr [[OMP_ARRAYCPY_DESTELEMENTPAST]], i32 1
// CHECK13-NEXT: [[OMP_ARRAYCPY_SRC_ELEMENT]] = getelementptr [[STRUCT_S]], ptr [[OMP_ARRAYCPY_SRCELEMENTPAST]], i32 1
// CHECK13-NEXT: [[OMP_ARRAYCPY_DONE:%.*]] = icmp eq ptr [[OMP_ARRAYCPY_DEST_ELEMENT]], [[TMP3]]
// CHECK13-NEXT: br i1 [[OMP_ARRAYCPY_DONE]], label [[OMP_ARRAYCPY_DONE3]], label [[OMP_ARRAYCPY_BODY]]
// CHECK13: omp.arraycpy.done3:
-// CHECK13-NEXT: call void @_ZN2StC1Ev(ptr noundef nonnull align 4 dereferenceable(8) [[AGG_TMP5]])
-// CHECK13-NEXT: call void @_ZN1SIfEC1ERKS0_2St(ptr noundef nonnull align 4 dereferenceable(4) [[VAR4]], ptr noundef nonnull align 4 dereferenceable(4) [[TMP2]], ptr nofreeobj noundef align 4 dereferenceable(8) [[AGG_TMP5]])
-// CHECK13-NEXT: call void @_ZN2StD1Ev(ptr noundef nonnull align 4 dead_on_return(8) dereferenceable(8) [[AGG_TMP5]]) #[[ATTR3]]
+// CHECK13-NEXT: call void @_ZN2StC1Ev(ptr noundef nonnull align 4 dereferenceable(8) [[AGG_TMP5]]) #[[ATTR4]]
+// CHECK13-NEXT: call void @_ZN1SIfEC1ERKS0_2St(ptr noundef nonnull align 4 dereferenceable(4) [[VAR4]], ptr noundef nonnull align 4 dereferenceable(4) [[TMP2]], ptr nofreeobj noundef align 4 dereferenceable(8) [[AGG_TMP5]]) #[[ATTR4]]
+// CHECK13-NEXT: call void @_ZN2StD1Ev(ptr noundef nonnull align 4 dead_on_return(8) dereferenceable(8) [[AGG_TMP5]]) #[[ATTR5]]
// CHECK13-NEXT: [[TMP4:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8
// CHECK13-NEXT: [[TMP5:%.*]] = load i32, ptr [[TMP4]], align 4
// CHECK13-NEXT: call void @__kmpc_for_static_init_4(ptr @[[GLOB1:[0-9]+]], i32 [[TMP5]], i32 92, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_COMB_LB]], ptr [[DOTOMP_COMB_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)
@@ -3671,14 +3671,14 @@ int main() {
// CHECK13-NEXT: store i32 2, ptr [[I]], align 4
// CHECK13-NEXT: br label [[DOTOMP_FINAL_DONE]]
// CHECK13: .omp.final.done:
-// CHECK13-NEXT: call void @_ZN1SIfED1Ev(ptr noundef nonnull align 4 dead_on_return(4) dereferenceable(4) [[VAR4]]) #[[ATTR3]]
+// CHECK13-NEXT: call void @_ZN1SIfED1Ev(ptr noundef nonnull align 4 dead_on_return(4) dereferenceable(4) [[VAR4]]) #[[ATTR5]]
// CHECK13-NEXT: [[ARRAY_BEGIN7:%.*]] = getelementptr inbounds [2 x [[STRUCT_S]]], ptr [[S_ARR2]], i32 0, i32 0
// CHECK13-NEXT: [[TMP25:%.*]] = getelementptr inbounds [[STRUCT_S]], ptr [[ARRAY_BEGIN7]], i64 2
// CHECK13-NEXT: br label [[ARRAYDESTROY_BODY:%.*]]
// CHECK13: arraydestroy.body:
// CHECK13-NEXT: [[ARRAYDESTROY_ELEMENTPAST:%.*]] = phi ptr [ [[TMP25]], [[DOTOMP_FINAL_DONE]] ], [ [[ARRAYDESTROY_ELEMENT:%.*]], [[ARRAYDESTROY_BODY]] ]
// CHECK13-NEXT: [[ARRAYDESTROY_ELEMENT]] = getelementptr inbounds [[STRUCT_S]], ptr [[ARRAYDESTROY_ELEMENTPAST]], i64 -1
-// CHECK13-NEXT: call void @_ZN1SIfED1Ev(ptr noundef nonnull align 4 dead_on_return(4) dereferenceable(4) [[ARRAYDESTROY_ELEMENT]]) #[[ATTR3]]
+// CHECK13-NEXT: call void @_ZN1SIfED1Ev(ptr noundef nonnull align 4 dead_on_return(4) dereferenceable(4) [[ARRAYDESTROY_ELEMENT]]) #[[ATTR5]]
// CHECK13-NEXT: [[ARRAYDESTROY_DONE:%.*]] = icmp eq ptr [[ARRAYDESTROY_ELEMENT]], [[ARRAY_BEGIN7]]
// CHECK13-NEXT: br i1 [[ARRAYDESTROY_DONE]], label [[ARRAYDESTROY_DONE8:%.*]], label [[ARRAYDESTROY_BODY]]
// CHECK13: arraydestroy.done8:
@@ -3691,7 +3691,7 @@ int main() {
// CHECK13-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 8
// CHECK13-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 8
// CHECK13-NEXT: [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8
-// CHECK13-NEXT: call void @_ZN2StC2Ev(ptr noundef nonnull align 4 dereferenceable(8) [[THIS1]])
+// CHECK13-NEXT: call void @_ZN2StC2Ev(ptr noundef nonnull align 4 dereferenceable(8) [[THIS1]]) #[[ATTR4]]
// CHECK13-NEXT: ret void
//
//
@@ -3706,7 +3706,7 @@ int main() {
// CHECK13-NEXT: store ptr [[T]], ptr [[T_INDIRECT_ADDR]], align 8
// CHECK13-NEXT: [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8
// CHECK13-NEXT: [[TMP0:%.*]] = load ptr, ptr [[S_ADDR]], align 8
-// CHECK13-NEXT: call void @_ZN1SIfEC2ERKS0_2St(ptr noundef nonnull align 4 dereferenceable(4) [[THIS1]], ptr noundef nonnull align 4 dereferenceable(4) [[TMP0]], ptr nofreeobj noundef align 4 dereferenceable(8) [[T]])
+// CHECK13-NEXT: call void @_ZN1SIfEC2ERKS0_2St(ptr noundef nonnull align 4 dereferenceable(4) [[THIS1]], ptr noundef nonnull align 4 dereferenceable(4) [[TMP0]], ptr nofreeobj noundef align 4 dereferenceable(8) [[T]]) #[[ATTR4]]
// CHECK13-NEXT: ret void
//
//
@@ -3716,7 +3716,7 @@ int main() {
// CHECK13-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 8
// CHECK13-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 8
// CHECK13-NEXT: [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8
-// CHECK13-NEXT: call void @_ZN2StD2Ev(ptr noundef nonnull align 4 dead_on_return(8) dereferenceable(8) [[THIS1]]) #[[ATTR3]]
+// CHECK13-NEXT: call void @_ZN2StD2Ev(ptr noundef nonnull align 4 dead_on_return(8) dereferenceable(8) [[THIS1]]) #[[ATTR5]]
// CHECK13-NEXT: ret void
//
//
@@ -3774,17 +3774,17 @@ int main() {
// CHECK13: omp.arraycpy.body:
// CHECK13-NEXT: [[OMP_ARRAYCPY_SRCELEMENTPAST:%.*]] = phi ptr [ [[TMP1]], [[ENTRY:%.*]] ], [ [[OMP_ARRAYCPY_SRC_ELEMENT:%.*]], [[OMP_ARRAYCPY_BODY]] ]
// CHECK13-NEXT: [[OMP_ARRAYCPY_DESTELEMENTPAST:%.*]] = phi ptr [ [[ARRAY_BEGIN]], [[ENTRY]] ], [ [[OMP_ARRAYCPY_DEST_ELEMENT:%.*]], [[OMP_ARRAYCPY_BODY]] ]
-// CHECK13-NEXT: call void @_ZN2StC1Ev(ptr noundef nonnull align 4 dereferenceable(8) [[AGG_TMP]])
-// CHECK13-NEXT: call void @_ZN1SIfEC1ERKS0_2St(ptr noundef nonnull align 4 dereferenceable(4) [[OMP_ARRAYCPY_DESTELEMENTPAST]], ptr noundef nonnull align 4 dereferenceable(4) [[OMP_ARRAYCPY_SRCELEMENTPAST]], ptr nofreeobj noundef align 4 dereferenceable(8) [[AGG_TMP]])
-// CHECK13-NEXT: call void @_ZN2StD1Ev(ptr noundef nonnull align 4 dead_on_return(8) dereferenceable(8) [[AGG_TMP]]) #[[ATTR3]]
+// CHECK13-NEXT: call void @_ZN2StC1Ev(ptr noundef nonnull align 4 dereferenceable(8) [[AGG_TMP]]) #[[ATTR4]]
+// CHECK13-NEXT: call void @_ZN1SIfEC1ERKS0_2St(ptr noundef nonnull align 4 dereferenceable(4) [[OMP_ARRAYCPY_DESTELEMENTPAST]], ptr noundef nonnull align 4 dereferenceable(4) [[OMP_ARRAYCPY_SRCELEMENTPAST]], ptr nofreeobj noundef align 4 dereferenceable(8) [[AGG_TMP]]) #[[ATTR4]]
+// CHECK13-NEXT: call void @_ZN2StD1Ev(ptr noundef nonnull align 4 dead_on_return(8) dereferenceable(8) [[AGG_TMP]]) #[[ATTR5]]
// CHECK13-NEXT: [[OMP_ARRAYCPY_DEST_ELEMENT]] = getelementptr [[STRUCT_S]], ptr [[OMP_ARRAYCPY_DESTELEMENTPAST]], i32 1
// CHECK13-NEXT: [[OMP_ARRAYCPY_SRC_ELEMENT]] = getelementptr [[STRUCT_S]], ptr [[OMP_ARRAYCPY_SRCELEMENTPAST]], i32 1
// CHECK13-NEXT: [[OMP_ARRAYCPY_DONE:%.*]] = icmp eq ptr [[OMP_ARRAYCPY_DEST_ELEMENT]], [[TMP5]]
// CHECK13-NEXT: br i1 [[OMP_ARRAYCPY_DONE]], label [[OMP_ARRAYCPY_DONE4]], label [[OMP_ARRAYCPY_BODY]]
// CHECK13: omp.arraycpy.done4:
-// CHECK13-NEXT: call void @_ZN2StC1Ev(ptr noundef nonnull align 4 dereferenceable(8) [[AGG_TMP6]])
-// CHECK13-NEXT: call void @_ZN1SIfEC1ERKS0_2St(ptr noundef nonnull align 4 dereferenceable(4) [[VAR5]], ptr noundef nonnull align 4 dereferenceable(4) [[TMP2]], ptr nofreeobj noundef align 4 dereferenceable(8) [[AGG_TMP6]])
-// CHECK13-NEXT: call void @_ZN2StD1Ev(ptr noundef nonnull align 4 dead_on_return(8) dereferenceable(8) [[AGG_TMP6]]) #[[ATTR3]]
+// CHECK13-NEXT: call void @_ZN2StC1Ev(ptr noundef nonnull align 4 dereferenceable(8) [[AGG_TMP6]]) #[[ATTR4]]
+// CHECK13-NEXT: call void @_ZN1SIfEC1ERKS0_2St(ptr noundef nonnull align 4 dereferenceable(4) [[VAR5]], ptr noundef nonnull align 4 dereferenceable(4) [[TMP2]], ptr nofreeobj noundef align 4 dereferenceable(8) [[AGG_TMP6]]) #[[ATTR4]]
+// CHECK13-NEXT: call void @_ZN2StD1Ev(ptr noundef nonnull align 4 dead_on_return(8) dereferenceable(8) [[AGG_TMP6]]) #[[ATTR5]]
// CHECK13-NEXT: [[TMP6:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8
// CHECK13-NEXT: [[TMP7:%.*]] = load i32, ptr [[TMP6]], align 4
// CHECK13-NEXT: call void @__kmpc_for_static_init_4(ptr @[[GLOB2:[0-9]+]], i32 [[TMP7]], i32 34, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)
@@ -3848,14 +3848,14 @@ int main() {
// CHECK13-NEXT: store i32 2, ptr [[I]], align 4
// CHECK13-NEXT: br label [[DOTOMP_FINAL_DONE]]
// CHECK13: .omp.final.done:
-// CHECK13-NEXT: call void @_ZN1SIfED1Ev(ptr noundef nonnull align 4 dead_on_return(4) dereferenceable(4) [[VAR5]]) #[[ATTR3]]
+// CHECK13-NEXT: call void @_ZN1SIfED1Ev(ptr noundef nonnull align 4 dead_on_return(4) dereferenceable(4) [[VAR5]]) #[[ATTR5]]
// CHECK13-NEXT: [[ARRAY_BEGIN12:%.*]] = getelementptr inbounds [2 x [[STRUCT_S]]], ptr [[S_ARR3]], i32 0, i32 0
// CHECK13-NEXT: [[TMP24:%.*]] = getelementptr inbounds [[STRUCT_S]], ptr [[ARRAY_BEGIN12]], i64 2
// CHECK13-NEXT: br label [[ARRAYDESTROY_BODY:%.*]]
// CHECK13: arraydestroy.body:
// CHECK13-NEXT: [[ARRAYDESTROY_ELEMENTPAST:%.*]] = phi ptr [ [[TMP24]], [[DOTOMP_FINAL_DONE]] ], [ [[ARRAYDESTROY_ELEMENT:%.*]], [[ARRAYDESTROY_BODY]] ]
// CHECK13-NEXT: [[ARRAYDESTROY_ELEMENT]] = getelementptr inbounds [[STRUCT_S]], ptr [[ARRAYDESTROY_ELEMENTPAST]], i64 -1
-// CHECK13-NEXT: call void @_ZN1SIfED1Ev(ptr noundef nonnull align 4 dead_on_return(4) dereferenceable(4) [[ARRAYDESTROY_ELEMENT]]) #[[ATTR3]]
+// CHECK13-NEXT: call void @_ZN1SIfED1Ev(ptr noundef nonnull align 4 dead_on_return(4) dereferenceable(4) [[ARRAYDESTROY_ELEMENT]]) #[[ATTR5]]
// CHECK13-NEXT: [[ARRAYDESTROY_DONE:%.*]] = icmp eq ptr [[ARRAYDESTROY_ELEMENT]], [[ARRAY_BEGIN12]]
// CHECK13-NEXT: br i1 [[ARRAYDESTROY_DONE]], label [[ARRAYDESTROY_DONE13:%.*]], label [[ARRAYDESTROY_BODY]]
// CHECK13: arraydestroy.done13:
@@ -3868,7 +3868,7 @@ int main() {
// CHECK13-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 8
// CHECK13-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 8
// CHECK13-NEXT: [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8
-// CHECK13-NEXT: call void @_ZN1SIfED2Ev(ptr noundef nonnull align 4 dead_on_return(4) dereferenceable(4) [[THIS1]]) #[[ATTR3]]
+// CHECK13-NEXT: call void @_ZN1SIfED2Ev(ptr noundef nonnull align 4 dead_on_return(4) dereferenceable(4) [[THIS1]]) #[[ATTR5]]
// CHECK13-NEXT: ret void
//
//
@@ -3954,18 +3954,18 @@ int main() {
// CHECK13: omp.arraycpy.body:
// CHECK13-NEXT: [[OMP_ARRAYCPY_SRCELEMENTPAST:%.*]] = phi ptr [ [[TMP1]], [[ENTRY:%.*]] ], [ [[OMP_ARRAYCPY_SRC_ELEMENT:%.*]], [[OMP_ARRAYCPY_BODY]] ]
// CHECK13-NEXT: [[OMP_ARRAYCPY_DESTELEMENTPAST:%.*]] = phi ptr [ [[ARRAY_BEGIN]], [[ENTRY]] ], [ [[OMP_ARRAYCPY_DEST_ELEMENT:%.*]], [[OMP_ARRAYCPY_BODY]] ]
-// CHECK13-NEXT: call void @_ZN2StC1Ev(ptr noundef nonnull align 4 dereferenceable(8) [[AGG_TMP]])
-// CHECK13-NEXT: call void @_ZN1SIiEC1ERKS0_2St(ptr noundef nonnull align 4 dereferenceable(4) [[OMP_ARRAYCPY_DESTELEMENTPAST]], ptr noundef nonnull align 4 dereferenceable(4) [[OMP_ARRAYCPY_SRCELEMENTPAST]], ptr nofreeobj noundef align 4 dereferenceable(8) [[AGG_TMP]])
-// CHECK13-NEXT: call void @_ZN2StD1Ev(ptr noundef nonnull align 4 dead_on_return(8) dereferenceable(8) [[AGG_TMP]]) #[[ATTR3]]
+// CHECK13-NEXT: call void @_ZN2StC1Ev(ptr noundef nonnull align 4 dereferenceable(8) [[AGG_TMP]]) #[[ATTR4]]
+// CHECK13-NEXT: call void @_ZN1SIiEC1ERKS0_2St(ptr noundef nonnull align 4 dereferenceable(4) [[OMP_ARRAYCPY_DESTELEMENTPAST]], ptr noundef nonnull align 4 dereferenceable(4) [[OMP_ARRAYCPY_SRCELEMENTPAST]], ptr nofreeobj noundef align 4 dereferenceable(8) [[AGG_TMP]]) #[[ATTR4]]
+// CHECK13-NEXT: call void @_ZN2StD1Ev(ptr noundef nonnull align 4 dead_on_return(8) dereferenceable(8) [[AGG_TMP]]) #[[ATTR5]]
// CHECK13-NEXT: [[OMP_ARRAYCPY_DEST_ELEMENT]] = getelementptr [[STRUCT_S_0]], ptr [[OMP_ARRAYCPY_DESTELEMENTPAST]], i32 1
// CHECK13-NEXT: [[OMP_ARRAYCPY_SRC_ELEMENT]] = getelementptr [[STRUCT_S_0]], ptr [[OMP_ARRAYCPY_SRCELEMENTPAST]], i32 1
// CHECK13-NEXT: [[OMP_ARRAYCPY_DONE:%.*]] = icmp eq ptr [[OMP_ARRAYCPY_DEST_ELEMENT]], [[TMP3]]
// CHECK13-NEXT: br i1 [[OMP_ARRAYCPY_DONE]], label [[OMP_ARRAYCPY_DONE4]], label [[OMP_ARRAYCPY_BODY]]
// CHECK13: omp.arraycpy.done4:
// CHECK13-NEXT: [[TMP4:%.*]] = load ptr, ptr [[TMP]], align 8, !nonnull [[META7]], !align [[META8]]
-// CHECK13-NEXT: call void @_ZN2StC1Ev(ptr noundef nonnull align 4 dereferenceable(8) [[AGG_TMP6]])
-// CHECK13-NEXT: call void @_ZN1SIiEC1ERKS0_2St(ptr noundef nonnull align 4 dereferenceable(4) [[VAR5]], ptr noundef nonnull align 4 dereferenceable(4) [[TMP4]], ptr nofreeobj noundef align 4 dereferenceable(8) [[AGG_TMP6]])
-// CHECK13-NEXT: call void @_ZN2StD1Ev(ptr noundef nonnull align 4 dead_on_return(8) dereferenceable(8) [[AGG_TMP6]]) #[[ATTR3]]
+// CHECK13-NEXT: call void @_ZN2StC1Ev(ptr noundef nonnull align 4 dereferenceable(8) [[AGG_TMP6]]) #[[ATTR4]]
+// CHECK13-NEXT: call void @_ZN1SIiEC1ERKS0_2St(ptr noundef nonnull align 4 dereferenceable(4) [[VAR5]], ptr noundef nonnull align 4 dereferenceable(4) [[TMP4]], ptr nofreeobj noundef align 4 dereferenceable(8) [[AGG_TMP6]]) #[[ATTR4]]
+// CHECK13-NEXT: call void @_ZN2StD1Ev(ptr noundef nonnull align 4 dead_on_return(8) dereferenceable(8) [[AGG_TMP6]]) #[[ATTR5]]
// CHECK13-NEXT: store ptr [[VAR5]], ptr [[_TMP7]], align 8
// CHECK13-NEXT: [[TMP5:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8
// CHECK13-NEXT: [[TMP6:%.*]] = load i32, ptr [[TMP5]], align 4
@@ -4021,14 +4021,14 @@ int main() {
// CHECK13-NEXT: store i32 2, ptr [[I]], align 4
// CHECK13-NEXT: br label [[DOTOMP_FINAL_DONE]]
// CHECK13: .omp.final.done:
-// CHECK13-NEXT: call void @_ZN1SIiED1Ev(ptr noundef nonnull align 4 dead_on_return(4) dereferenceable(4) [[VAR5]]) #[[ATTR3]]
+// CHECK13-NEXT: call void @_ZN1SIiED1Ev(ptr noundef nonnull align 4 dead_on_return(4) dereferenceable(4) [[VAR5]]) #[[ATTR5]]
// CHECK13-NEXT: [[ARRAY_BEGIN9:%.*]] = getelementptr inbounds [2 x [[STRUCT_S_0]]], ptr [[S_ARR3]], i32 0, i32 0
// CHECK13-NEXT: [[TMP25:%.*]] = getelementptr inbounds [[STRUCT_S_0]], ptr [[ARRAY_BEGIN9]], i64 2
// CHECK13-NEXT: br label [[ARRAYDESTROY_BODY:%.*]]
// CHECK13: arraydestroy.body:
// CHECK13-NEXT: [[ARRAYDESTROY_ELEMENTPAST:%.*]] = phi ptr [ [[TMP25]], [[DOTOMP_FINAL_DONE]] ], [ [[ARRAYDESTROY_ELEMENT:%.*]], [[ARRAYDESTROY_BODY]] ]
// CHECK13-NEXT: [[ARRAYDESTROY_ELEMENT]] = getelementptr inbounds [[STRUCT_S_0]], ptr [[ARRAYDESTROY_ELEMENTPAST]], i64 -1
-// CHECK13-NEXT: call void @_ZN1SIiED1Ev(ptr noundef nonnull align 4 dead_on_return(4) dereferenceable(4) [[ARRAYDESTROY_ELEMENT]]) #[[ATTR3]]
+// CHECK13-NEXT: call void @_ZN1SIiED1Ev(ptr noundef nonnull align 4 dead_on_return(4) dereferenceable(4) [[ARRAYDESTROY_ELEMENT]]) #[[ATTR5]]
// CHECK13-NEXT: [[ARRAYDESTROY_DONE:%.*]] = icmp eq ptr [[ARRAYDESTROY_ELEMENT]], [[ARRAY_BEGIN9]]
// CHECK13-NEXT: br i1 [[ARRAYDESTROY_DONE]], label [[ARRAYDESTROY_DONE10:%.*]], label [[ARRAYDESTROY_BODY]]
// CHECK13: arraydestroy.done10:
@@ -4046,7 +4046,7 @@ int main() {
// CHECK13-NEXT: store ptr [[T]], ptr [[T_INDIRECT_ADDR]], align 8
// CHECK13-NEXT: [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8
// CHECK13-NEXT: [[TMP0:%.*]] = load ptr, ptr [[S_ADDR]], align 8
-// CHECK13-NEXT: call void @_ZN1SIiEC2ERKS0_2St(ptr noundef nonnull align 4 dereferenceable(4) [[THIS1]], ptr noundef nonnull align 4 dereferenceable(4) [[TMP0]], ptr nofreeobj noundef align 4 dereferenceable(8) [[T]])
+// CHECK13-NEXT: call void @_ZN1SIiEC2ERKS0_2St(ptr noundef nonnull align 4 dereferenceable(4) [[THIS1]], ptr noundef nonnull align 4 dereferenceable(4) [[TMP0]], ptr nofreeobj noundef align 4 dereferenceable(8) [[T]]) #[[ATTR4]]
// CHECK13-NEXT: ret void
//
//
@@ -4105,18 +4105,18 @@ int main() {
// CHECK13: omp.arraycpy.body:
// CHECK13-NEXT: [[OMP_ARRAYCPY_SRCELEMENTPAST:%.*]] = phi ptr [ [[TMP1]], [[ENTRY:%.*]] ], [ [[OMP_ARRAYCPY_SRC_ELEMENT:%.*]], [[OMP_ARRAYCPY_BODY]] ]
// CHECK13-NEXT: [[OMP_ARRAYCPY_DESTELEMENTPAST:%.*]] = phi ptr [ [[ARRAY_BEGIN]], [[ENTRY]] ], [ [[OMP_ARRAYCPY_DEST_ELEMENT:%.*]], [[OMP_ARRAYCPY_BODY]] ]
-// CHECK13-NEXT: call void @_ZN2StC1Ev(ptr noundef nonnull align 4 dereferenceable(8) [[AGG_TMP]])
-// CHECK13-NEXT: call void @_ZN1SIiEC1ERKS0_2St(ptr noundef nonnull align 4 dereferenceable(4) [[OMP_ARRAYCPY_DESTELEMENTPAST]], ptr noundef nonnull align 4 dereferenceable(4) [[OMP_ARRAYCPY_SRCELEMENTPAST]], ptr nofreeobj noundef align 4 dereferenceable(8) [[AGG_TMP]])
-// CHECK13-NEXT: call void @_ZN2StD1Ev(ptr noundef nonnull align 4 dead_on_return(8) dereferenceable(8) [[AGG_TMP]]) #[[ATTR3]]
+// CHECK13-NEXT: call void @_ZN2StC1Ev(ptr noundef nonnull align 4 dereferenceable(8) [[AGG_TMP]]) #[[ATTR4]]
+// CHECK13-NEXT: call void @_ZN1SIiEC1ERKS0_2St(ptr noundef nonnull align 4 dereferenceable(4) [[OMP_ARRAYCPY_DESTELEMENTPAST]], ptr noundef nonnull align 4 dereferenceable(4) [[OMP_ARRAYCPY_SRCELEMENTPAST]], ptr nofreeobj noundef align 4 dereferenceable(8) [[AGG_TMP]]) #[[ATTR4]]
+// CHECK13-NEXT: call void @_ZN2StD1Ev(ptr noundef nonnull align 4 dead_on_return(8) dereferenceable(8) [[AGG_TMP]]) #[[ATTR5]]
// CHECK13-NEXT: [[OMP_ARRAYCPY_DEST_ELEMENT]] = getelementptr [[STRUCT_S_0]], ptr [[OMP_ARRAYCPY_DESTELEMENTPAST]], i32 1
// CHECK13-NEXT: [[OMP_ARRAYCPY_SRC_ELEMENT]] = getelementptr [[STRUCT_S_0]], ptr [[OMP_ARRAYCPY_SRCELEMENTPAST]], i32 1
// CHECK13-NEXT: [[OMP_ARRAYCPY_DONE:%.*]] = icmp eq ptr [[OMP_ARRAYCPY_DEST_ELEMENT]], [[TMP5]]
// CHECK13-NEXT: br i1 [[OMP_ARRAYCPY_DONE]], label [[OMP_ARRAYCPY_DONE5]], label [[OMP_ARRAYCPY_BODY]]
// CHECK13: omp.arraycpy.done5:
// CHECK13-NEXT: [[TMP6:%.*]] = load ptr, ptr [[TMP]], align 8, !nonnull [[META7]], !align [[META8]]
-// CHECK13-NEXT: call void @_ZN2StC1Ev(ptr noundef nonnull align 4 dereferenceable(8) [[AGG_TMP7]])
-// CHECK13-NEXT: call void @_ZN1SIiEC1ERKS0_2St(ptr noundef nonnull align 4 dereferenceable(4) [[VAR6]], ptr noundef nonnull align 4 dereferenceable(4) [[TMP6]], ptr nofreeobj noundef align 4 dereferenceable(8) [[AGG_TMP7]])
-// CHECK13-NEXT: call void @_ZN2StD1Ev(ptr noundef nonnull align 4 dead_on_return(8) dereferenceable(8) [[AGG_TMP7]]) #[[ATTR3]]
+// CHECK13-NEXT: call void @_ZN2StC1Ev(ptr noundef nonnull align 4 dereferenceable(8) [[AGG_TMP7]]) #[[ATTR4]]
+// CHECK13-NEXT: call void @_ZN1SIiEC1ERKS0_2St(ptr noundef nonnull align 4 dereferenceable(4) [[VAR6]], ptr noundef nonnull align 4 dereferenceable(4) [[TMP6]], ptr nofreeobj noundef align 4 dereferenceable(8) [[AGG_TMP7]]) #[[ATTR4]]
+// CHECK13-NEXT: call void @_ZN2StD1Ev(ptr noundef nonnull align 4 dead_on_return(8) dereferenceable(8) [[AGG_TMP7]]) #[[ATTR5]]
// CHECK13-NEXT: store ptr [[VAR6]], ptr [[_TMP8]], align 8
// CHECK13-NEXT: [[TMP7:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8
// CHECK13-NEXT: [[TMP8:%.*]] = load i32, ptr [[TMP7]], align 4
@@ -4178,14 +4178,14 @@ int main() {
// CHECK13-NEXT: store i32 2, ptr [[I]], align 4
// CHECK13-NEXT: br label [[DOTOMP_FINAL_DONE]]
// CHECK13: .omp.final.done:
-// CHECK13-NEXT: call void @_ZN1SIiED1Ev(ptr noundef nonnull align 4 dead_on_return(4) dereferenceable(4) [[VAR6]]) #[[ATTR3]]
+// CHECK13-NEXT: call void @_ZN1SIiED1Ev(ptr noundef nonnull align 4 dead_on_return(4) dereferenceable(4) [[VAR6]]) #[[ATTR5]]
// CHECK13-NEXT: [[ARRAY_BEGIN13:%.*]] = getelementptr inbounds [2 x [[STRUCT_S_0]]], ptr [[S_ARR4]], i32 0, i32 0
// CHECK13-NEXT: [[TMP24:%.*]] = getelementptr inbounds [[STRUCT_S_0]], ptr [[ARRAY_BEGIN13]], i64 2
// CHECK13-NEXT: br label [[ARRAYDESTROY_BODY:%.*]]
// CHECK13: arraydestroy.body:
// CHECK13-NEXT: [[ARRAYDESTROY_ELEMENTPAST:%.*]] = phi ptr [ [[TMP24]], [[DOTOMP_FINAL_DONE]] ], [ [[ARRAYDESTROY_ELEMENT:%.*]], [[ARRAYDESTROY_BODY]] ]
// CHECK13-NEXT: [[ARRAYDESTROY_ELEMENT]] = getelementptr inbounds [[STRUCT_S_0]], ptr [[ARRAYDESTROY_ELEMENTPAST]], i64 -1
-// CHECK13-NEXT: call void @_ZN1SIiED1Ev(ptr noundef nonnull align 4 dead_on_return(4) dereferenceable(4) [[ARRAYDESTROY_ELEMENT]]) #[[ATTR3]]
+// CHECK13-NEXT: call void @_ZN1SIiED1Ev(ptr noundef nonnull align 4 dead_on_return(4) dereferenceable(4) [[ARRAYDESTROY_ELEMENT]]) #[[ATTR5]]
// CHECK13-NEXT: [[ARRAYDESTROY_DONE:%.*]] = icmp eq ptr [[ARRAYDESTROY_ELEMENT]], [[ARRAY_BEGIN13]]
// CHECK13-NEXT: br i1 [[ARRAYDESTROY_DONE]], label [[ARRAYDESTROY_DONE14:%.*]], label [[ARRAYDESTROY_BODY]]
// CHECK13: arraydestroy.done14:
@@ -4198,7 +4198,7 @@ int main() {
// CHECK13-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 8
// CHECK13-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 8
// CHECK13-NEXT: [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8
-// CHECK13-NEXT: call void @_ZN1SIiED2Ev(ptr noundef nonnull align 4 dead_on_return(4) dereferenceable(4) [[THIS1]]) #[[ATTR3]]
+// CHECK13-NEXT: call void @_ZN1SIiED2Ev(ptr noundef nonnull align 4 dead_on_return(4) dereferenceable(4) [[THIS1]]) #[[ATTR5]]
// CHECK13-NEXT: ret void
//
//
@@ -4373,17 +4373,17 @@ int main() {
// CHECK15: omp.arraycpy.body:
// CHECK15-NEXT: [[OMP_ARRAYCPY_SRCELEMENTPAST:%.*]] = phi ptr [ [[TMP1]], [[ENTRY:%.*]] ], [ [[OMP_ARRAYCPY_SRC_ELEMENT:%.*]], [[OMP_ARRAYCPY_BODY]] ]
// CHECK15-NEXT: [[OMP_ARRAYCPY_DESTELEMENTPAST:%.*]] = phi ptr [ [[ARRAY_BEGIN]], [[ENTRY]] ], [ [[OMP_ARRAYCPY_DEST_ELEMENT:%.*]], [[OMP_ARRAYCPY_BODY]] ]
-// CHECK15-NEXT: call void @_ZN2StC1Ev(ptr noundef nonnull align 4 dereferenceable(8) [[AGG_TMP]])
-// CHECK15-NEXT: call void @_ZN1SIfEC1ERKS0_2St(ptr noundef nonnull align 4 dereferenceable(4) [[OMP_ARRAYCPY_DESTELEMENTPAST]], ptr noundef nonnull align 4 dereferenceable(4) [[OMP_ARRAYCPY_SRCELEMENTPAST]], ptr nofreeobj noundef align 4 dereferenceable(8) [[AGG_TMP]])
-// CHECK15-NEXT: call void @_ZN2StD1Ev(ptr noundef nonnull align 4 dead_on_return(8) dereferenceable(8) [[AGG_TMP]]) #[[ATTR3:[0-9]+]]
+// CHECK15-NEXT: call void @_ZN2StC1Ev(ptr noundef nonnull align 4 dereferenceable(8) [[AGG_TMP]]) #[[ATTR4:[0-9]+]]
+// CHECK15-NEXT: call void @_ZN1SIfEC1ERKS0_2St(ptr noundef nonnull align 4 dereferenceable(4) [[OMP_ARRAYCPY_DESTELEMENTPAST]], ptr noundef nonnull align 4 dereferenceable(4) [[OMP_ARRAYCPY_SRCELEMENTPAST]], ptr nofreeobj noundef align 4 dereferenceable(8) [[AGG_TMP]]) #[[ATTR4]]
+// CHECK15-NEXT: call void @_ZN2StD1Ev(ptr noundef nonnull align 4 dead_on_return(8) dereferenceable(8) [[AGG_TMP]]) #[[ATTR5:[0-9]+]]
// CHECK15-NEXT: [[OMP_ARRAYCPY_DEST_ELEMENT]] = getelementptr [[STRUCT_S]], ptr [[OMP_ARRAYCPY_DESTELEMENTPAST]], i32 1
// CHECK15-NEXT: [[OMP_ARRAYCPY_SRC_ELEMENT]] = getelementptr [[STRUCT_S]], ptr [[OMP_ARRAYCPY_SRCELEMENTPAST]], i32 1
// CHECK15-NEXT: [[OMP_ARRAYCPY_DONE:%.*]] = icmp eq ptr [[OMP_ARRAYCPY_DEST_ELEMENT]], [[TMP3]]
// CHECK15-NEXT: br i1 [[OMP_ARRAYCPY_DONE]], label [[OMP_ARRAYCPY_DONE3]], label [[OMP_ARRAYCPY_BODY]]
// CHECK15: omp.arraycpy.done3:
-// CHECK15-NEXT: call void @_ZN2StC1Ev(ptr noundef nonnull align 4 dereferenceable(8) [[AGG_TMP5]])
-// CHECK15-NEXT: call void @_ZN1SIfEC1ERKS0_2St(ptr noundef nonnull align 4 dereferenceable(4) [[VAR4]], ptr noundef nonnull align 4 dereferenceable(4) [[TMP2]], ptr nofreeobj noundef align 4 dereferenceable(8) [[AGG_TMP5]])
-// CHECK15-NEXT: call void @_ZN2StD1Ev(ptr noundef nonnull align 4 dead_on_return(8) dereferenceable(8) [[AGG_TMP5]]) #[[ATTR3]]
+// CHECK15-NEXT: call void @_ZN2StC1Ev(ptr noundef nonnull align 4 dereferenceable(8) [[AGG_TMP5]]) #[[ATTR4]]
+// CHECK15-NEXT: call void @_ZN1SIfEC1ERKS0_2St(ptr noundef nonnull align 4 dereferenceable(4) [[VAR4]], ptr noundef nonnull align 4 dereferenceable(4) [[TMP2]], ptr nofreeobj noundef align 4 dereferenceable(8) [[AGG_TMP5]]) #[[ATTR4]]
+// CHECK15-NEXT: call void @_ZN2StD1Ev(ptr noundef nonnull align 4 dead_on_return(8) dereferenceable(8) [[AGG_TMP5]]) #[[ATTR5]]
// CHECK15-NEXT: [[TMP4:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 4
// CHECK15-NEXT: [[TMP5:%.*]] = load i32, ptr [[TMP4]], align 4
// CHECK15-NEXT: call void @__kmpc_for_static_init_4(ptr @[[GLOB1:[0-9]+]], i32 [[TMP5]], i32 92, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_COMB_LB]], ptr [[DOTOMP_COMB_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)
@@ -4438,14 +4438,14 @@ int main() {
// CHECK15-NEXT: store i32 2, ptr [[I]], align 4
// CHECK15-NEXT: br label [[DOTOMP_FINAL_DONE]]
// CHECK15: .omp.final.done:
-// CHECK15-NEXT: call void @_ZN1SIfED1Ev(ptr noundef nonnull align 4 dead_on_return(4) dereferenceable(4) [[VAR4]]) #[[ATTR3]]
+// CHECK15-NEXT: call void @_ZN1SIfED1Ev(ptr noundef nonnull align 4 dead_on_return(4) dereferenceable(4) [[VAR4]]) #[[ATTR5]]
// CHECK15-NEXT: [[ARRAY_BEGIN7:%.*]] = getelementptr inbounds [2 x [[STRUCT_S]]], ptr [[S_ARR2]], i32 0, i32 0
// CHECK15-NEXT: [[TMP23:%.*]] = getelementptr inbounds [[STRUCT_S]], ptr [[ARRAY_BEGIN7]], i32 2
// CHECK15-NEXT: br label [[ARRAYDESTROY_BODY:%.*]]
// CHECK15: arraydestroy.body:
// CHECK15-NEXT: [[ARRAYDESTROY_ELEMENTPAST:%.*]] = phi ptr [ [[TMP23]], [[DOTOMP_FINAL_DONE]] ], [ [[ARRAYDESTROY_ELEMENT:%.*]], [[ARRAYDESTROY_BODY]] ]
// CHECK15-NEXT: [[ARRAYDESTROY_ELEMENT]] = getelementptr inbounds [[STRUCT_S]], ptr [[ARRAYDESTROY_ELEMENTPAST]], i32 -1
-// CHECK15-NEXT: call void @_ZN1SIfED1Ev(ptr noundef nonnull align 4 dead_on_return(4) dereferenceable(4) [[ARRAYDESTROY_ELEMENT]]) #[[ATTR3]]
+// CHECK15-NEXT: call void @_ZN1SIfED1Ev(ptr noundef nonnull align 4 dead_on_return(4) dereferenceable(4) [[ARRAYDESTROY_ELEMENT]]) #[[ATTR5]]
// CHECK15-NEXT: [[ARRAYDESTROY_DONE:%.*]] = icmp eq ptr [[ARRAYDESTROY_ELEMENT]], [[ARRAY_BEGIN7]]
// CHECK15-NEXT: br i1 [[ARRAYDESTROY_DONE]], label [[ARRAYDESTROY_DONE8:%.*]], label [[ARRAYDESTROY_BODY]]
// CHECK15: arraydestroy.done8:
@@ -4458,7 +4458,7 @@ int main() {
// CHECK15-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 4
// CHECK15-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 4
// CHECK15-NEXT: [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 4
-// CHECK15-NEXT: call void @_ZN2StC2Ev(ptr noundef nonnull align 4 dereferenceable(8) [[THIS1]])
+// CHECK15-NEXT: call void @_ZN2StC2Ev(ptr noundef nonnull align 4 dereferenceable(8) [[THIS1]]) #[[ATTR4]]
// CHECK15-NEXT: ret void
//
//
@@ -4473,7 +4473,7 @@ int main() {
// CHECK15-NEXT: store ptr [[T]], ptr [[T_INDIRECT_ADDR]], align 4
// CHECK15-NEXT: [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 4
// CHECK15-NEXT: [[TMP0:%.*]] = load ptr, ptr [[S_ADDR]], align 4
-// CHECK15-NEXT: call void @_ZN1SIfEC2ERKS0_2St(ptr noundef nonnull align 4 dereferenceable(4) [[THIS1]], ptr noundef nonnull align 4 dereferenceable(4) [[TMP0]], ptr nofreeobj noundef align 4 dereferenceable(8) [[T]])
+// CHECK15-NEXT: call void @_ZN1SIfEC2ERKS0_2St(ptr noundef nonnull align 4 dereferenceable(4) [[THIS1]], ptr noundef nonnull align 4 dereferenceable(4) [[TMP0]], ptr nofreeobj noundef align 4 dereferenceable(8) [[T]]) #[[ATTR4]]
// CHECK15-NEXT: ret void
//
//
@@ -4483,7 +4483,7 @@ int main() {
// CHECK15-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 4
// CHECK15-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 4
// CHECK15-NEXT: [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 4
-// CHECK15-NEXT: call void @_ZN2StD2Ev(ptr noundef nonnull align 4 dead_on_return(8) dereferenceable(8) [[THIS1]]) #[[ATTR3]]
+// CHECK15-NEXT: call void @_ZN2StD2Ev(ptr noundef nonnull align 4 dead_on_return(8) dereferenceable(8) [[THIS1]]) #[[ATTR5]]
// CHECK15-NEXT: ret void
//
//
@@ -4539,17 +4539,17 @@ int main() {
// CHECK15: omp.arraycpy.body:
// CHECK15-NEXT: [[OMP_ARRAYCPY_SRCELEMENTPAST:%.*]] = phi ptr [ [[TMP1]], [[ENTRY:%.*]] ], [ [[OMP_ARRAYCPY_SRC_ELEMENT:%.*]], [[OMP_ARRAYCPY_BODY]] ]
// CHECK15-NEXT: [[OMP_ARRAYCPY_DESTELEMENTPAST:%.*]] = phi ptr [ [[ARRAY_BEGIN]], [[ENTRY]] ], [ [[OMP_ARRAYCPY_DEST_ELEMENT:%.*]], [[OMP_ARRAYCPY_BODY]] ]
-// CHECK15-NEXT: call void @_ZN2StC1Ev(ptr noundef nonnull align 4 dereferenceable(8) [[AGG_TMP]])
-// CHECK15-NEXT: call void @_ZN1SIfEC1ERKS0_2St(ptr noundef nonnull align 4 dereferenceable(4) [[OMP_ARRAYCPY_DESTELEMENTPAST]], ptr noundef nonnull align 4 dereferenceable(4) [[OMP_ARRAYCPY_SRCELEMENTPAST]], ptr nofreeobj noundef align 4 dereferenceable(8) [[AGG_TMP]])
-// CHECK15-NEXT: call void @_ZN2StD1Ev(ptr noundef nonnull align 4 dead_on_return(8) dereferenceable(8) [[AGG_TMP]]) #[[ATTR3]]
+// CHECK15-NEXT: call void @_ZN2StC1Ev(ptr noundef nonnull align 4 dereferenceable(8) [[AGG_TMP]]) #[[ATTR4]]
+// CHECK15-NEXT: call void @_ZN1SIfEC1ERKS0_2St(ptr noundef nonnull align 4 dereferenceable(4) [[OMP_ARRAYCPY_DESTELEMENTPAST]], ptr noundef nonnull align 4 dereferenceable(4) [[OMP_ARRAYCPY_SRCELEMENTPAST]], ptr nofreeobj noundef align 4 dereferenceable(8) [[AGG_TMP]]) #[[ATTR4]]
+// CHECK15-NEXT: call void @_ZN2StD1Ev(ptr noundef nonnull align 4 dead_on_return(8) dereferenceable(8) [[AGG_TMP]]) #[[ATTR5]]
// CHECK15-NEXT: [[OMP_ARRAYCPY_DEST_ELEMENT]] = getelementptr [[STRUCT_S]], ptr [[OMP_ARRAYCPY_DESTELEMENTPAST]], i32 1
// CHECK15-NEXT: [[OMP_ARRAYCPY_SRC_ELEMENT]] = getelementptr [[STRUCT_S]], ptr [[OMP_ARRAYCPY_SRCELEMENTPAST]], i32 1
// CHECK15-NEXT: [[OMP_ARRAYCPY_DONE:%.*]] = icmp eq ptr [[OMP_ARRAYCPY_DEST_ELEMENT]], [[TMP5]]
// CHECK15-NEXT: br i1 [[OMP_ARRAYCPY_DONE]], label [[OMP_ARRAYCPY_DONE3]], label [[OMP_ARRAYCPY_BODY]]
// CHECK15: omp.arraycpy.done3:
-// CHECK15-NEXT: call void @_ZN2StC1Ev(ptr noundef nonnull align 4 dereferenceable(8) [[AGG_TMP5]])
-// CHECK15-NEXT: call void @_ZN1SIfEC1ERKS0_2St(ptr noundef nonnull align 4 dereferenceable(4) [[VAR4]], ptr noundef nonnull align 4 dereferenceable(4) [[TMP2]], ptr nofreeobj noundef align 4 dereferenceable(8) [[AGG_TMP5]])
-// CHECK15-NEXT: call void @_ZN2StD1Ev(ptr noundef nonnull align 4 dead_on_return(8) dereferenceable(8) [[AGG_TMP5]]) #[[ATTR3]]
+// CHECK15-NEXT: call void @_ZN2StC1Ev(ptr noundef nonnull align 4 dereferenceable(8) [[AGG_TMP5]]) #[[ATTR4]]
+// CHECK15-NEXT: call void @_ZN1SIfEC1ERKS0_2St(ptr noundef nonnull align 4 dereferenceable(4) [[VAR4]], ptr noundef nonnull align 4 dereferenceable(4) [[TMP2]], ptr nofreeobj noundef align 4 dereferenceable(8) [[AGG_TMP5]]) #[[ATTR4]]
+// CHECK15-NEXT: call void @_ZN2StD1Ev(ptr noundef nonnull align 4 dead_on_return(8) dereferenceable(8) [[AGG_TMP5]]) #[[ATTR5]]
// CHECK15-NEXT: [[TMP6:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 4
// CHECK15-NEXT: [[TMP7:%.*]] = load i32, ptr [[TMP6]], align 4
// CHECK15-NEXT: call void @__kmpc_for_static_init_4(ptr @[[GLOB2:[0-9]+]], i32 [[TMP7]], i32 34, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)
@@ -4611,14 +4611,14 @@ int main() {
// CHECK15-NEXT: store i32 2, ptr [[I]], align 4
// CHECK15-NEXT: br label [[DOTOMP_FINAL_DONE]]
// CHECK15: .omp.final.done:
-// CHECK15-NEXT: call void @_ZN1SIfED1Ev(ptr noundef nonnull align 4 dead_on_return(4) dereferenceable(4) [[VAR4]]) #[[ATTR3]]
+// CHECK15-NEXT: call void @_ZN1SIfED1Ev(ptr noundef nonnull align 4 dead_on_return(4) dereferenceable(4) [[VAR4]]) #[[ATTR5]]
// CHECK15-NEXT: [[ARRAY_BEGIN10:%.*]] = getelementptr inbounds [2 x [[STRUCT_S]]], ptr [[S_ARR2]], i32 0, i32 0
// CHECK15-NEXT: [[TMP24:%.*]] = getelementptr inbounds [[STRUCT_S]], ptr [[ARRAY_BEGIN10]], i32 2
// CHECK15-NEXT: br label [[ARRAYDESTROY_BODY:%.*]]
// CHECK15: arraydestroy.body:
// CHECK15-NEXT: [[ARRAYDESTROY_ELEMENTPAST:%.*]] = phi ptr [ [[TMP24]], [[DOTOMP_FINAL_DONE]] ], [ [[ARRAYDESTROY_ELEMENT:%.*]], [[ARRAYDESTROY_BODY]] ]
// CHECK15-NEXT: [[ARRAYDESTROY_ELEMENT]] = getelementptr inbounds [[STRUCT_S]], ptr [[ARRAYDESTROY_ELEMENTPAST]], i32 -1
-// CHECK15-NEXT: call void @_ZN1SIfED1Ev(ptr noundef nonnull align 4 dead_on_return(4) dereferenceable(4) [[ARRAYDESTROY_ELEMENT]]) #[[ATTR3]]
+// CHECK15-NEXT: call void @_ZN1SIfED1Ev(ptr noundef nonnull align 4 dead_on_return(4) dereferenceable(4) [[ARRAYDESTROY_ELEMENT]]) #[[ATTR5]]
// CHECK15-NEXT: [[ARRAYDESTROY_DONE:%.*]] = icmp eq ptr [[ARRAYDESTROY_ELEMENT]], [[ARRAY_BEGIN10]]
// CHECK15-NEXT: br i1 [[ARRAYDESTROY_DONE]], label [[ARRAYDESTROY_DONE11:%.*]], label [[ARRAYDESTROY_BODY]]
// CHECK15: arraydestroy.done11:
@@ -4631,7 +4631,7 @@ int main() {
// CHECK15-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 4
// CHECK15-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 4
// CHECK15-NEXT: [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 4
-// CHECK15-NEXT: call void @_ZN1SIfED2Ev(ptr noundef nonnull align 4 dead_on_return(4) dereferenceable(4) [[THIS1]]) #[[ATTR3]]
+// CHECK15-NEXT: call void @_ZN1SIfED2Ev(ptr noundef nonnull align 4 dead_on_return(4) dereferenceable(4) [[THIS1]]) #[[ATTR5]]
// CHECK15-NEXT: ret void
//
//
@@ -4717,18 +4717,18 @@ int main() {
// CHECK15: omp.arraycpy.body:
// CHECK15-NEXT: [[OMP_ARRAYCPY_SRCELEMENTPAST:%.*]] = phi ptr [ [[TMP1]], [[ENTRY:%.*]] ], [ [[OMP_ARRAYCPY_SRC_ELEMENT:%.*]], [[OMP_ARRAYCPY_BODY]] ]
// CHECK15-NEXT: [[OMP_ARRAYCPY_DESTELEMENTPAST:%.*]] = phi ptr [ [[ARRAY_BEGIN]], [[ENTRY]] ], [ [[OMP_ARRAYCPY_DEST_ELEMENT:%.*]], [[OMP_ARRAYCPY_BODY]] ]
-// CHECK15-NEXT: call void @_ZN2StC1Ev(ptr noundef nonnull align 4 dereferenceable(8) [[AGG_TMP]])
-// CHECK15-NEXT: call void @_ZN1SIiEC1ERKS0_2St(ptr noundef nonnull align 4 dereferenceable(4) [[OMP_ARRAYCPY_DESTELEMENTPAST]], ptr noundef nonnull align 4 dereferenceable(4) [[OMP_ARRAYCPY_SRCELEMENTPAST]], ptr nofreeobj noundef align 4 dereferenceable(8) [[AGG_TMP]])
-// CHECK15-NEXT: call void @_ZN2StD1Ev(ptr noundef nonnull align 4 dead_on_return(8) dereferenceable(8) [[AGG_TMP]]) #[[ATTR3]]
+// CHECK15-NEXT: call void @_ZN2StC1Ev(ptr noundef nonnull align 4 dereferenceable(8) [[AGG_TMP]]) #[[ATTR4]]
+// CHECK15-NEXT: call void @_ZN1SIiEC1ERKS0_2St(ptr noundef nonnull align 4 dereferenceable(4) [[OMP_ARRAYCPY_DESTELEMENTPAST]], ptr noundef nonnull align 4 dereferenceable(4) [[OMP_ARRAYCPY_SRCELEMENTPAST]], ptr nofreeobj noundef align 4 dereferenceable(8) [[AGG_TMP]]) #[[ATTR4]]
+// CHECK15-NEXT: call void @_ZN2StD1Ev(ptr noundef nonnull align 4 dead_on_return(8) dereferenceable(8) [[AGG_TMP]]) #[[ATTR5]]
// CHECK15-NEXT: [[OMP_ARRAYCPY_DEST_ELEMENT]] = getelementptr [[STRUCT_S_0]], ptr [[OMP_ARRAYCPY_DESTELEMENTPAST]], i32 1
// CHECK15-NEXT: [[OMP_ARRAYCPY_SRC_ELEMENT]] = getelementptr [[STRUCT_S_0]], ptr [[OMP_ARRAYCPY_SRCELEMENTPAST]], i32 1
// CHECK15-NEXT: [[OMP_ARRAYCPY_DONE:%.*]] = icmp eq ptr [[OMP_ARRAYCPY_DEST_ELEMENT]], [[TMP3]]
// CHECK15-NEXT: br i1 [[OMP_ARRAYCPY_DONE]], label [[OMP_ARRAYCPY_DONE4]], label [[OMP_ARRAYCPY_BODY]]
// CHECK15: omp.arraycpy.done4:
// CHECK15-NEXT: [[TMP4:%.*]] = load ptr, ptr [[TMP]], align 4, !nonnull [[META8]], !align [[META9]]
-// CHECK15-NEXT: call void @_ZN2StC1Ev(ptr noundef nonnull align 4 dereferenceable(8) [[AGG_TMP6]])
-// CHECK15-NEXT: call void @_ZN1SIiEC1ERKS0_2St(ptr noundef nonnull align 4 dereferenceable(4) [[VAR5]], ptr noundef nonnull align 4 dereferenceable(4) [[TMP4]], ptr nofreeobj noundef align 4 dereferenceable(8) [[AGG_TMP6]])
-// CHECK15-NEXT: call void @_ZN2StD1Ev(ptr noundef nonnull align 4 dead_on_return(8) dereferenceable(8) [[AGG_TMP6]]) #[[ATTR3]]
+// CHECK15-NEXT: call void @_ZN2StC1Ev(ptr noundef nonnull align 4 dereferenceable(8) [[AGG_TMP6]]) #[[ATTR4]]
+// CHECK15-NEXT: call void @_ZN1SIiEC1ERKS0_2St(ptr noundef nonnull align 4 dereferenceable(4) [[VAR5]], ptr noundef nonnull align 4 dereferenceable(4) [[TMP4]], ptr nofreeobj noundef align 4 dereferenceable(8) [[AGG_TMP6]]) #[[ATTR4]]
+// CHECK15-NEXT: call void @_ZN2StD1Ev(ptr noundef nonnull align 4 dead_on_return(8) dereferenceable(8) [[AGG_TMP6]]) #[[ATTR5]]
// CHECK15-NEXT: store ptr [[VAR5]], ptr [[_TMP7]], align 4
// CHECK15-NEXT: [[TMP5:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 4
// CHECK15-NEXT: [[TMP6:%.*]] = load i32, ptr [[TMP5]], align 4
@@ -4782,14 +4782,14 @@ int main() {
// CHECK15-NEXT: store i32 2, ptr [[I]], align 4
// CHECK15-NEXT: br label [[DOTOMP_FINAL_DONE]]
// CHECK15: .omp.final.done:
-// CHECK15-NEXT: call void @_ZN1SIiED1Ev(ptr noundef nonnull align 4 dead_on_return(4) dereferenceable(4) [[VAR5]]) #[[ATTR3]]
+// CHECK15-NEXT: call void @_ZN1SIiED1Ev(ptr noundef nonnull align 4 dead_on_return(4) dereferenceable(4) [[VAR5]]) #[[ATTR5]]
// CHECK15-NEXT: [[ARRAY_BEGIN9:%.*]] = getelementptr inbounds [2 x [[STRUCT_S_0]]], ptr [[S_ARR3]], i32 0, i32 0
// CHECK15-NEXT: [[TMP23:%.*]] = getelementptr inbounds [[STRUCT_S_0]], ptr [[ARRAY_BEGIN9]], i32 2
// CHECK15-NEXT: br label [[ARRAYDESTROY_BODY:%.*]]
// CHECK15: arraydestroy.body:
// CHECK15-NEXT: [[ARRAYDESTROY_ELEMENTPAST:%.*]] = phi ptr [ [[TMP23]], [[DOTOMP_FINAL_DONE]] ], [ [[ARRAYDESTROY_ELEMENT:%.*]], [[ARRAYDESTROY_BODY]] ]
// CHECK15-NEXT: [[ARRAYDESTROY_ELEMENT]] = getelementptr inbounds [[STRUCT_S_0]], ptr [[ARRAYDESTROY_ELEMENTPAST]], i32 -1
-// CHECK15-NEXT: call void @_ZN1SIiED1Ev(ptr noundef nonnull align 4 dead_on_return(4) dereferenceable(4) [[ARRAYDESTROY_ELEMENT]]) #[[ATTR3]]
+// CHECK15-NEXT: call void @_ZN1SIiED1Ev(ptr noundef nonnull align 4 dead_on_return(4) dereferenceable(4) [[ARRAYDESTROY_ELEMENT]]) #[[ATTR5]]
// CHECK15-NEXT: [[ARRAYDESTROY_DONE:%.*]] = icmp eq ptr [[ARRAYDESTROY_ELEMENT]], [[ARRAY_BEGIN9]]
// CHECK15-NEXT: br i1 [[ARRAYDESTROY_DONE]], label [[ARRAYDESTROY_DONE10:%.*]], label [[ARRAYDESTROY_BODY]]
// CHECK15: arraydestroy.done10:
@@ -4807,7 +4807,7 @@ int main() {
// CHECK15-NEXT: store ptr [[T]], ptr [[T_INDIRECT_ADDR]], align 4
// CHECK15-NEXT: [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 4
// CHECK15-NEXT: [[TMP0:%.*]] = load ptr, ptr [[S_ADDR]], align 4
-// CHECK15-NEXT: call void @_ZN1SIiEC2ERKS0_2St(ptr noundef nonnull align 4 dereferenceable(4) [[THIS1]], ptr noundef nonnull align 4 dereferenceable(4) [[TMP0]], ptr nofreeobj noundef align 4 dereferenceable(8) [[T]])
+// CHECK15-NEXT: call void @_ZN1SIiEC2ERKS0_2St(ptr noundef nonnull align 4 dereferenceable(4) [[THIS1]], ptr noundef nonnull align 4 dereferenceable(4) [[TMP0]], ptr nofreeobj noundef align 4 dereferenceable(8) [[T]]) #[[ATTR4]]
// CHECK15-NEXT: ret void
//
//
@@ -4864,18 +4864,18 @@ int main() {
// CHECK15: omp.arraycpy.body:
// CHECK15-NEXT: [[OMP_ARRAYCPY_SRCELEMENTPAST:%.*]] = phi ptr [ [[TMP1]], [[ENTRY:%.*]] ], [ [[OMP_ARRAYCPY_SRC_ELEMENT:%.*]], [[OMP_ARRAYCPY_BODY]] ]
// CHECK15-NEXT: [[OMP_ARRAYCPY_DESTELEMENTPAST:%.*]] = phi ptr [ [[ARRAY_BEGIN]], [[ENTRY]] ], [ [[OMP_ARRAYCPY_DEST_ELEMENT:%.*]], [[OMP_ARRAYCPY_BODY]] ]
-// CHECK15-NEXT: call void @_ZN2StC1Ev(ptr noundef nonnull align 4 dereferenceable(8) [[AGG_TMP]])
-// CHECK15-NEXT: call void @_ZN1SIiEC1ERKS0_2St(ptr noundef nonnull align 4 dereferenceable(4) [[OMP_ARRAYCPY_DESTELEMENTPAST]], ptr noundef nonnull align 4 dereferenceable(4) [[OMP_ARRAYCPY_SRCELEMENTPAST]], ptr nofreeobj noundef align 4 dereferenceable(8) [[AGG_TMP]])
-// CHECK15-NEXT: call void @_ZN2StD1Ev(ptr noundef nonnull align 4 dead_on_return(8) dereferenceable(8) [[AGG_TMP]]) #[[ATTR3]]
+// CHECK15-NEXT: call void @_ZN2StC1Ev(ptr noundef nonnull align 4 dereferenceable(8) [[AGG_TMP]]) #[[ATTR4]]
+// CHECK15-NEXT: call void @_ZN1SIiEC1ERKS0_2St(ptr noundef nonnull align 4 dereferenceable(4) [[OMP_ARRAYCPY_DESTELEMENTPAST]], ptr noundef nonnull align 4 dereferenceable(4) [[OMP_ARRAYCPY_SRCELEMENTPAST]], ptr nofreeobj noundef align 4 dereferenceable(8) [[AGG_TMP]]) #[[ATTR4]]
+// CHECK15-NEXT: call void @_ZN2StD1Ev(ptr noundef nonnull align 4 dead_on_return(8) dereferenceable(8) [[AGG_TMP]]) #[[ATTR5]]
// CHECK15-NEXT: [[OMP_ARRAYCPY_DEST_ELEMENT]] = getelementptr [[STRUCT_S_0]], ptr [[OMP_ARRAYCPY_DESTELEMENTPAST]], i32 1
// CHECK15-NEXT: [[OMP_ARRAYCPY_SRC_ELEMENT]] = getelementptr [[STRUCT_S_0]], ptr [[OMP_ARRAYCPY_SRCELEMENTPAST]], i32 1
// CHECK15-NEXT: [[OMP_ARRAYCPY_DONE:%.*]] = icmp eq ptr [[OMP_ARRAYCPY_DEST_ELEMENT]], [[TMP5]]
// CHECK15-NEXT: br i1 [[OMP_ARRAYCPY_DONE]], label [[OMP_ARRAYCPY_DONE4]], label [[OMP_ARRAYCPY_BODY]]
// CHECK15: omp.arraycpy.done4:
// CHECK15-NEXT: [[TMP6:%.*]] = load ptr, ptr [[TMP]], align 4, !nonnull [[META8]], !align [[META9]]
-// CHECK15-NEXT: call void @_ZN2StC1Ev(ptr noundef nonnull align 4 dereferenceable(8) [[AGG_TMP6]])
-// CHECK15-NEXT: call void @_ZN1SIiEC1ERKS0_2St(ptr noundef nonnull align 4 dereferenceable(4) [[VAR5]], ptr noundef nonnull align 4 dereferenceable(4) [[TMP6]], ptr nofreeobj noundef align 4 dereferenceable(8) [[AGG_TMP6]])
-// CHECK15-NEXT: call void @_ZN2StD1Ev(ptr noundef nonnull align 4 dead_on_return(8) dereferenceable(8) [[AGG_TMP6]]) #[[ATTR3]]
+// CHECK15-NEXT: call void @_ZN2StC1Ev(ptr noundef nonnull align 4 dereferenceable(8) [[AGG_TMP6]]) #[[ATTR4]]
+// CHECK15-NEXT: call void @_ZN1SIiEC1ERKS0_2St(ptr noundef nonnull align 4 dereferenceable(4) [[VAR5]], ptr noundef nonnull align 4 dereferenceable(4) [[TMP6]], ptr nofreeobj noundef align 4 dereferenceable(8) [[AGG_TMP6]]) #[[ATTR4]]
+// CHECK15-NEXT: call void @_ZN2StD1Ev(ptr noundef nonnull align 4 dead_on_return(8) dereferenceable(8) [[AGG_TMP6]]) #[[ATTR5]]
// CHECK15-NEXT: store ptr [[VAR5]], ptr [[_TMP7]], align 4
// CHECK15-NEXT: [[TMP7:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 4
// CHECK15-NEXT: [[TMP8:%.*]] = load i32, ptr [[TMP7]], align 4
@@ -4935,14 +4935,14 @@ int main() {
// CHECK15-NEXT: store i32 2, ptr [[I]], align 4
// CHECK15-NEXT: br label [[DOTOMP_FINAL_DONE]]
// CHECK15: .omp.final.done:
-// CHECK15-NEXT: call void @_ZN1SIiED1Ev(ptr noundef nonnull align 4 dead_on_return(4) dereferenceable(4) [[VAR5]]) #[[ATTR3]]
+// CHECK15-NEXT: call void @_ZN1SIiED1Ev(ptr noundef nonnull align 4 dead_on_return(4) dereferenceable(4) [[VAR5]]) #[[ATTR5]]
// CHECK15-NEXT: [[ARRAY_BEGIN11:%.*]] = getelementptr inbounds [2 x [[STRUCT_S_0]]], ptr [[S_ARR3]], i32 0, i32 0
// CHECK15-NEXT: [[TMP24:%.*]] = getelementptr inbounds [[STRUCT_S_0]], ptr [[ARRAY_BEGIN11]], i32 2
// CHECK15-NEXT: br label [[ARRAYDESTROY_BODY:%.*]]
// CHECK15: arraydestroy.body:
// CHECK15-NEXT: [[ARRAYDESTROY_ELEMENTPAST:%.*]] = phi ptr [ [[TMP24]], [[DOTOMP_FINAL_DONE]] ], [ [[ARRAYDESTROY_ELEMENT:%.*]], [[ARRAYDESTROY_BODY]] ]
// CHECK15-NEXT: [[ARRAYDESTROY_ELEMENT]] = getelementptr inbounds [[STRUCT_S_0]], ptr [[ARRAYDESTROY_ELEMENTPAST]], i32 -1
-// CHECK15-NEXT: call void @_ZN1SIiED1Ev(ptr noundef nonnull align 4 dead_on_return(4) dereferenceable(4) [[ARRAYDESTROY_ELEMENT]]) #[[ATTR3]]
+// CHECK15-NEXT: call void @_ZN1SIiED1Ev(ptr noundef nonnull align 4 dead_on_return(4) dereferenceable(4) [[ARRAYDESTROY_ELEMENT]]) #[[ATTR5]]
// CHECK15-NEXT: [[ARRAYDESTROY_DONE:%.*]] = icmp eq ptr [[ARRAYDESTROY_ELEMENT]], [[ARRAY_BEGIN11]]
// CHECK15-NEXT: br i1 [[ARRAYDESTROY_DONE]], label [[ARRAYDESTROY_DONE12:%.*]], label [[ARRAYDESTROY_BODY]]
// CHECK15: arraydestroy.done12:
@@ -4955,7 +4955,7 @@ int main() {
// CHECK15-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 4
// CHECK15-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 4
// CHECK15-NEXT: [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 4
-// CHECK15-NEXT: call void @_ZN1SIiED2Ev(ptr noundef nonnull align 4 dead_on_return(4) dereferenceable(4) [[THIS1]]) #[[ATTR3]]
+// CHECK15-NEXT: call void @_ZN1SIiED2Ev(ptr noundef nonnull align 4 dead_on_return(4) dereferenceable(4) [[THIS1]]) #[[ATTR5]]
// CHECK15-NEXT: ret void
//
//
@@ -5250,7 +5250,7 @@ int main() {
// CHECK17-NEXT: store ptr [[TMP13]], ptr [[TMP12]], align 8, !llvm.access.group [[ACC_GRP11]]
// CHECK17-NEXT: [[TMP14:%.*]] = getelementptr inbounds nuw [[CLASS_ANON]], ptr [[REF_TMP]], i32 0, i32 2
// CHECK17-NEXT: store ptr [[SIVAR_ADDR]], ptr [[TMP14]], align 8, !llvm.access.group [[ACC_GRP11]]
-// CHECK17-NEXT: call void @"_ZZZ4mainENK3$_0clEvENKUlvE_clEv"(ptr noundef nonnull align 8 dereferenceable(24) [[REF_TMP]]), !llvm.access.group [[ACC_GRP11]]
+// CHECK17-NEXT: call void @"_ZZZ4mainENK3$_0clEvENKUlvE_clEv"(ptr noundef nonnull align 8 dereferenceable(24) [[REF_TMP]]) #[[ATTR3:[0-9]+]], !llvm.access.group [[ACC_GRP11]]
// CHECK17-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]
// CHECK17: omp.body.continue:
// CHECK17-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]
diff --git a/clang/test/OpenMP/target_teams_distribute_parallel_for_simd_private_codegen.cpp b/clang/test/OpenMP/target_teams_distribute_parallel_for_simd_private_codegen.cpp
index 6d9d437979044..0509529d6f517 100644
--- a/clang/test/OpenMP/target_teams_distribute_parallel_for_simd_private_codegen.cpp
+++ b/clang/test/OpenMP/target_teams_distribute_parallel_for_simd_private_codegen.cpp
@@ -2977,12 +2977,12 @@ int main() {
// CHECK13-NEXT: br label [[ARRAYCTOR_LOOP:%.*]]
// CHECK13: arrayctor.loop:
// CHECK13-NEXT: [[ARRAYCTOR_CUR:%.*]] = phi ptr [ [[ARRAY_BEGIN]], [[ENTRY:%.*]] ], [ [[ARRAYCTOR_NEXT:%.*]], [[ARRAYCTOR_LOOP]] ]
-// CHECK13-NEXT: call void @_ZN1SIfEC1Ev(ptr noundef nonnull align 4 dereferenceable(4) [[ARRAYCTOR_CUR]])
+// CHECK13-NEXT: call void @_ZN1SIfEC1Ev(ptr noundef nonnull align 4 dereferenceable(4) [[ARRAYCTOR_CUR]]) #[[ATTR4:[0-9]+]]
// CHECK13-NEXT: [[ARRAYCTOR_NEXT]] = getelementptr inbounds [[STRUCT_S]], ptr [[ARRAYCTOR_CUR]], i64 1
// CHECK13-NEXT: [[ARRAYCTOR_DONE:%.*]] = icmp eq ptr [[ARRAYCTOR_NEXT]], [[ARRAYCTOR_END]]
// CHECK13-NEXT: br i1 [[ARRAYCTOR_DONE]], label [[ARRAYCTOR_CONT:%.*]], label [[ARRAYCTOR_LOOP]]
// CHECK13: arrayctor.cont:
-// CHECK13-NEXT: call void @_ZN1SIfEC1Ev(ptr noundef nonnull align 4 dereferenceable(4) [[VAR]])
+// CHECK13-NEXT: call void @_ZN1SIfEC1Ev(ptr noundef nonnull align 4 dereferenceable(4) [[VAR]]) #[[ATTR4]]
// CHECK13-NEXT: [[TMP0:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8
// CHECK13-NEXT: [[TMP1:%.*]] = load i32, ptr [[TMP0]], align 4
// CHECK13-NEXT: call void @__kmpc_for_static_init_4(ptr @[[GLOB1:[0-9]+]], i32 [[TMP1]], i32 92, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_COMB_LB]], ptr [[DOTOMP_COMB_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)
@@ -3033,14 +3033,14 @@ int main() {
// CHECK13-NEXT: store i32 2, ptr [[I]], align 4
// CHECK13-NEXT: br label [[DOTOMP_FINAL_DONE]]
// CHECK13: .omp.final.done:
-// CHECK13-NEXT: call void @_ZN1SIfED1Ev(ptr noundef nonnull align 4 dead_on_return(4) dereferenceable(4) [[VAR]]) #[[ATTR2:[0-9]+]]
+// CHECK13-NEXT: call void @_ZN1SIfED1Ev(ptr noundef nonnull align 4 dead_on_return(4) dereferenceable(4) [[VAR]]) #[[ATTR5:[0-9]+]]
// CHECK13-NEXT: [[ARRAY_BEGIN2:%.*]] = getelementptr inbounds [2 x [[STRUCT_S]]], ptr [[S_ARR]], i32 0, i32 0
// CHECK13-NEXT: [[TMP17:%.*]] = getelementptr inbounds [[STRUCT_S]], ptr [[ARRAY_BEGIN2]], i64 2
// CHECK13-NEXT: br label [[ARRAYDESTROY_BODY:%.*]]
// CHECK13: arraydestroy.body:
// CHECK13-NEXT: [[ARRAYDESTROY_ELEMENTPAST:%.*]] = phi ptr [ [[TMP17]], [[DOTOMP_FINAL_DONE]] ], [ [[ARRAYDESTROY_ELEMENT:%.*]], [[ARRAYDESTROY_BODY]] ]
// CHECK13-NEXT: [[ARRAYDESTROY_ELEMENT]] = getelementptr inbounds [[STRUCT_S]], ptr [[ARRAYDESTROY_ELEMENTPAST]], i64 -1
-// CHECK13-NEXT: call void @_ZN1SIfED1Ev(ptr noundef nonnull align 4 dead_on_return(4) dereferenceable(4) [[ARRAYDESTROY_ELEMENT]]) #[[ATTR2]]
+// CHECK13-NEXT: call void @_ZN1SIfED1Ev(ptr noundef nonnull align 4 dead_on_return(4) dereferenceable(4) [[ARRAYDESTROY_ELEMENT]]) #[[ATTR5]]
// CHECK13-NEXT: [[ARRAYDESTROY_DONE:%.*]] = icmp eq ptr [[ARRAYDESTROY_ELEMENT]], [[ARRAY_BEGIN2]]
// CHECK13-NEXT: br i1 [[ARRAYDESTROY_DONE]], label [[ARRAYDESTROY_DONE3:%.*]], label [[ARRAYDESTROY_BODY]]
// CHECK13: arraydestroy.done3:
@@ -3053,7 +3053,7 @@ int main() {
// CHECK13-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 8
// CHECK13-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 8
// CHECK13-NEXT: [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8
-// CHECK13-NEXT: call void @_ZN1SIfEC2Ev(ptr noundef nonnull align 4 dereferenceable(4) [[THIS1]])
+// CHECK13-NEXT: call void @_ZN1SIfEC2Ev(ptr noundef nonnull align 4 dereferenceable(4) [[THIS1]]) #[[ATTR4]]
// CHECK13-NEXT: ret void
//
//
@@ -3095,12 +3095,12 @@ int main() {
// CHECK13-NEXT: br label [[ARRAYCTOR_LOOP:%.*]]
// CHECK13: arrayctor.loop:
// CHECK13-NEXT: [[ARRAYCTOR_CUR:%.*]] = phi ptr [ [[ARRAY_BEGIN]], [[ENTRY:%.*]] ], [ [[ARRAYCTOR_NEXT:%.*]], [[ARRAYCTOR_LOOP]] ]
-// CHECK13-NEXT: call void @_ZN1SIfEC1Ev(ptr noundef nonnull align 4 dereferenceable(4) [[ARRAYCTOR_CUR]])
+// CHECK13-NEXT: call void @_ZN1SIfEC1Ev(ptr noundef nonnull align 4 dereferenceable(4) [[ARRAYCTOR_CUR]]) #[[ATTR4]]
// CHECK13-NEXT: [[ARRAYCTOR_NEXT]] = getelementptr inbounds [[STRUCT_S]], ptr [[ARRAYCTOR_CUR]], i64 1
// CHECK13-NEXT: [[ARRAYCTOR_DONE:%.*]] = icmp eq ptr [[ARRAYCTOR_NEXT]], [[ARRAYCTOR_END]]
// CHECK13-NEXT: br i1 [[ARRAYCTOR_DONE]], label [[ARRAYCTOR_CONT:%.*]], label [[ARRAYCTOR_LOOP]]
// CHECK13: arrayctor.cont:
-// CHECK13-NEXT: call void @_ZN1SIfEC1Ev(ptr noundef nonnull align 4 dereferenceable(4) [[VAR]])
+// CHECK13-NEXT: call void @_ZN1SIfEC1Ev(ptr noundef nonnull align 4 dereferenceable(4) [[VAR]]) #[[ATTR4]]
// CHECK13-NEXT: [[TMP2:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8
// CHECK13-NEXT: [[TMP3:%.*]] = load i32, ptr [[TMP2]], align 4
// CHECK13-NEXT: call void @__kmpc_for_static_init_4(ptr @[[GLOB2:[0-9]+]], i32 [[TMP3]], i32 34, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)
@@ -3164,14 +3164,14 @@ int main() {
// CHECK13-NEXT: store i32 2, ptr [[I]], align 4
// CHECK13-NEXT: br label [[DOTOMP_FINAL_DONE]]
// CHECK13: .omp.final.done:
-// CHECK13-NEXT: call void @_ZN1SIfED1Ev(ptr noundef nonnull align 4 dead_on_return(4) dereferenceable(4) [[VAR]]) #[[ATTR2]]
+// CHECK13-NEXT: call void @_ZN1SIfED1Ev(ptr noundef nonnull align 4 dead_on_return(4) dereferenceable(4) [[VAR]]) #[[ATTR5]]
// CHECK13-NEXT: [[ARRAY_BEGIN7:%.*]] = getelementptr inbounds [2 x [[STRUCT_S]]], ptr [[S_ARR]], i32 0, i32 0
// CHECK13-NEXT: [[TMP20:%.*]] = getelementptr inbounds [[STRUCT_S]], ptr [[ARRAY_BEGIN7]], i64 2
// CHECK13-NEXT: br label [[ARRAYDESTROY_BODY:%.*]]
// CHECK13: arraydestroy.body:
// CHECK13-NEXT: [[ARRAYDESTROY_ELEMENTPAST:%.*]] = phi ptr [ [[TMP20]], [[DOTOMP_FINAL_DONE]] ], [ [[ARRAYDESTROY_ELEMENT:%.*]], [[ARRAYDESTROY_BODY]] ]
// CHECK13-NEXT: [[ARRAYDESTROY_ELEMENT]] = getelementptr inbounds [[STRUCT_S]], ptr [[ARRAYDESTROY_ELEMENTPAST]], i64 -1
-// CHECK13-NEXT: call void @_ZN1SIfED1Ev(ptr noundef nonnull align 4 dead_on_return(4) dereferenceable(4) [[ARRAYDESTROY_ELEMENT]]) #[[ATTR2]]
+// CHECK13-NEXT: call void @_ZN1SIfED1Ev(ptr noundef nonnull align 4 dead_on_return(4) dereferenceable(4) [[ARRAYDESTROY_ELEMENT]]) #[[ATTR5]]
// CHECK13-NEXT: [[ARRAYDESTROY_DONE:%.*]] = icmp eq ptr [[ARRAYDESTROY_ELEMENT]], [[ARRAY_BEGIN7]]
// CHECK13-NEXT: br i1 [[ARRAYDESTROY_DONE]], label [[ARRAYDESTROY_DONE8:%.*]], label [[ARRAYDESTROY_BODY]]
// CHECK13: arraydestroy.done8:
@@ -3184,7 +3184,7 @@ int main() {
// CHECK13-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 8
// CHECK13-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 8
// CHECK13-NEXT: [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8
-// CHECK13-NEXT: call void @_ZN1SIfED2Ev(ptr noundef nonnull align 4 dead_on_return(4) dereferenceable(4) [[THIS1]]) #[[ATTR2]]
+// CHECK13-NEXT: call void @_ZN1SIfED2Ev(ptr noundef nonnull align 4 dead_on_return(4) dereferenceable(4) [[THIS1]]) #[[ATTR5]]
// CHECK13-NEXT: ret void
//
//
@@ -3233,12 +3233,12 @@ int main() {
// CHECK13-NEXT: br label [[ARRAYCTOR_LOOP:%.*]]
// CHECK13: arrayctor.loop:
// CHECK13-NEXT: [[ARRAYCTOR_CUR:%.*]] = phi ptr [ [[ARRAY_BEGIN]], [[ENTRY:%.*]] ], [ [[ARRAYCTOR_NEXT:%.*]], [[ARRAYCTOR_LOOP]] ]
-// CHECK13-NEXT: call void @_ZN1SIiEC1Ev(ptr noundef nonnull align 4 dereferenceable(4) [[ARRAYCTOR_CUR]])
+// CHECK13-NEXT: call void @_ZN1SIiEC1Ev(ptr noundef nonnull align 4 dereferenceable(4) [[ARRAYCTOR_CUR]]) #[[ATTR4]]
// CHECK13-NEXT: [[ARRAYCTOR_NEXT]] = getelementptr inbounds [[STRUCT_S_0]], ptr [[ARRAYCTOR_CUR]], i64 1
// CHECK13-NEXT: [[ARRAYCTOR_DONE:%.*]] = icmp eq ptr [[ARRAYCTOR_NEXT]], [[ARRAYCTOR_END]]
// CHECK13-NEXT: br i1 [[ARRAYCTOR_DONE]], label [[ARRAYCTOR_CONT:%.*]], label [[ARRAYCTOR_LOOP]]
// CHECK13: arrayctor.cont:
-// CHECK13-NEXT: call void @_ZN1SIiEC1Ev(ptr noundef nonnull align 4 dereferenceable(4) [[VAR]])
+// CHECK13-NEXT: call void @_ZN1SIiEC1Ev(ptr noundef nonnull align 4 dereferenceable(4) [[VAR]]) #[[ATTR4]]
// CHECK13-NEXT: store ptr [[VAR]], ptr [[_TMP2]], align 8
// CHECK13-NEXT: [[TMP0:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8
// CHECK13-NEXT: [[TMP1:%.*]] = load i32, ptr [[TMP0]], align 4
@@ -3290,14 +3290,14 @@ int main() {
// CHECK13-NEXT: store i32 2, ptr [[I]], align 4
// CHECK13-NEXT: br label [[DOTOMP_FINAL_DONE]]
// CHECK13: .omp.final.done:
-// CHECK13-NEXT: call void @_ZN1SIiED1Ev(ptr noundef nonnull align 4 dead_on_return(4) dereferenceable(4) [[VAR]]) #[[ATTR2]]
+// CHECK13-NEXT: call void @_ZN1SIiED1Ev(ptr noundef nonnull align 4 dead_on_return(4) dereferenceable(4) [[VAR]]) #[[ATTR5]]
// CHECK13-NEXT: [[ARRAY_BEGIN4:%.*]] = getelementptr inbounds [2 x [[STRUCT_S_0]]], ptr [[S_ARR]], i32 0, i32 0
// CHECK13-NEXT: [[TMP17:%.*]] = getelementptr inbounds [[STRUCT_S_0]], ptr [[ARRAY_BEGIN4]], i64 2
// CHECK13-NEXT: br label [[ARRAYDESTROY_BODY:%.*]]
// CHECK13: arraydestroy.body:
// CHECK13-NEXT: [[ARRAYDESTROY_ELEMENTPAST:%.*]] = phi ptr [ [[TMP17]], [[DOTOMP_FINAL_DONE]] ], [ [[ARRAYDESTROY_ELEMENT:%.*]], [[ARRAYDESTROY_BODY]] ]
// CHECK13-NEXT: [[ARRAYDESTROY_ELEMENT]] = getelementptr inbounds [[STRUCT_S_0]], ptr [[ARRAYDESTROY_ELEMENTPAST]], i64 -1
-// CHECK13-NEXT: call void @_ZN1SIiED1Ev(ptr noundef nonnull align 4 dead_on_return(4) dereferenceable(4) [[ARRAYDESTROY_ELEMENT]]) #[[ATTR2]]
+// CHECK13-NEXT: call void @_ZN1SIiED1Ev(ptr noundef nonnull align 4 dead_on_return(4) dereferenceable(4) [[ARRAYDESTROY_ELEMENT]]) #[[ATTR5]]
// CHECK13-NEXT: [[ARRAYDESTROY_DONE:%.*]] = icmp eq ptr [[ARRAYDESTROY_ELEMENT]], [[ARRAY_BEGIN4]]
// CHECK13-NEXT: br i1 [[ARRAYDESTROY_DONE]], label [[ARRAYDESTROY_DONE5:%.*]], label [[ARRAYDESTROY_BODY]]
// CHECK13: arraydestroy.done5:
@@ -3310,7 +3310,7 @@ int main() {
// CHECK13-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 8
// CHECK13-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 8
// CHECK13-NEXT: [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8
-// CHECK13-NEXT: call void @_ZN1SIiEC2Ev(ptr noundef nonnull align 4 dereferenceable(4) [[THIS1]])
+// CHECK13-NEXT: call void @_ZN1SIiEC2Ev(ptr noundef nonnull align 4 dereferenceable(4) [[THIS1]]) #[[ATTR4]]
// CHECK13-NEXT: ret void
//
//
@@ -3354,12 +3354,12 @@ int main() {
// CHECK13-NEXT: br label [[ARRAYCTOR_LOOP:%.*]]
// CHECK13: arrayctor.loop:
// CHECK13-NEXT: [[ARRAYCTOR_CUR:%.*]] = phi ptr [ [[ARRAY_BEGIN]], [[ENTRY:%.*]] ], [ [[ARRAYCTOR_NEXT:%.*]], [[ARRAYCTOR_LOOP]] ]
-// CHECK13-NEXT: call void @_ZN1SIiEC1Ev(ptr noundef nonnull align 4 dereferenceable(4) [[ARRAYCTOR_CUR]])
+// CHECK13-NEXT: call void @_ZN1SIiEC1Ev(ptr noundef nonnull align 4 dereferenceable(4) [[ARRAYCTOR_CUR]]) #[[ATTR4]]
// CHECK13-NEXT: [[ARRAYCTOR_NEXT]] = getelementptr inbounds [[STRUCT_S_0]], ptr [[ARRAYCTOR_CUR]], i64 1
// CHECK13-NEXT: [[ARRAYCTOR_DONE:%.*]] = icmp eq ptr [[ARRAYCTOR_NEXT]], [[ARRAYCTOR_END]]
// CHECK13-NEXT: br i1 [[ARRAYCTOR_DONE]], label [[ARRAYCTOR_CONT:%.*]], label [[ARRAYCTOR_LOOP]]
// CHECK13: arrayctor.cont:
-// CHECK13-NEXT: call void @_ZN1SIiEC1Ev(ptr noundef nonnull align 4 dereferenceable(4) [[VAR]])
+// CHECK13-NEXT: call void @_ZN1SIiEC1Ev(ptr noundef nonnull align 4 dereferenceable(4) [[VAR]]) #[[ATTR4]]
// CHECK13-NEXT: store ptr [[VAR]], ptr [[_TMP3]], align 8
// CHECK13-NEXT: [[TMP2:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8
// CHECK13-NEXT: [[TMP3:%.*]] = load i32, ptr [[TMP2]], align 4
@@ -3421,14 +3421,14 @@ int main() {
// CHECK13-NEXT: store i32 2, ptr [[I]], align 4
// CHECK13-NEXT: br label [[DOTOMP_FINAL_DONE]]
// CHECK13: .omp.final.done:
-// CHECK13-NEXT: call void @_ZN1SIiED1Ev(ptr noundef nonnull align 4 dead_on_return(4) dereferenceable(4) [[VAR]]) #[[ATTR2]]
+// CHECK13-NEXT: call void @_ZN1SIiED1Ev(ptr noundef nonnull align 4 dead_on_return(4) dereferenceable(4) [[VAR]]) #[[ATTR5]]
// CHECK13-NEXT: [[ARRAY_BEGIN8:%.*]] = getelementptr inbounds [2 x [[STRUCT_S_0]]], ptr [[S_ARR]], i32 0, i32 0
// CHECK13-NEXT: [[TMP19:%.*]] = getelementptr inbounds [[STRUCT_S_0]], ptr [[ARRAY_BEGIN8]], i64 2
// CHECK13-NEXT: br label [[ARRAYDESTROY_BODY:%.*]]
// CHECK13: arraydestroy.body:
// CHECK13-NEXT: [[ARRAYDESTROY_ELEMENTPAST:%.*]] = phi ptr [ [[TMP19]], [[DOTOMP_FINAL_DONE]] ], [ [[ARRAYDESTROY_ELEMENT:%.*]], [[ARRAYDESTROY_BODY]] ]
// CHECK13-NEXT: [[ARRAYDESTROY_ELEMENT]] = getelementptr inbounds [[STRUCT_S_0]], ptr [[ARRAYDESTROY_ELEMENTPAST]], i64 -1
-// CHECK13-NEXT: call void @_ZN1SIiED1Ev(ptr noundef nonnull align 4 dead_on_return(4) dereferenceable(4) [[ARRAYDESTROY_ELEMENT]]) #[[ATTR2]]
+// CHECK13-NEXT: call void @_ZN1SIiED1Ev(ptr noundef nonnull align 4 dead_on_return(4) dereferenceable(4) [[ARRAYDESTROY_ELEMENT]]) #[[ATTR5]]
// CHECK13-NEXT: [[ARRAYDESTROY_DONE:%.*]] = icmp eq ptr [[ARRAYDESTROY_ELEMENT]], [[ARRAY_BEGIN8]]
// CHECK13-NEXT: br i1 [[ARRAYDESTROY_DONE]], label [[ARRAYDESTROY_DONE9:%.*]], label [[ARRAYDESTROY_BODY]]
// CHECK13: arraydestroy.done9:
@@ -3441,7 +3441,7 @@ int main() {
// CHECK13-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 8
// CHECK13-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 8
// CHECK13-NEXT: [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 8
-// CHECK13-NEXT: call void @_ZN1SIiED2Ev(ptr noundef nonnull align 4 dead_on_return(4) dereferenceable(4) [[THIS1]]) #[[ATTR2]]
+// CHECK13-NEXT: call void @_ZN1SIiED2Ev(ptr noundef nonnull align 4 dead_on_return(4) dereferenceable(4) [[THIS1]]) #[[ATTR5]]
// CHECK13-NEXT: ret void
//
//
@@ -3531,12 +3531,12 @@ int main() {
// CHECK15-NEXT: br label [[ARRAYCTOR_LOOP:%.*]]
// CHECK15: arrayctor.loop:
// CHECK15-NEXT: [[ARRAYCTOR_CUR:%.*]] = phi ptr [ [[ARRAY_BEGIN]], [[ENTRY:%.*]] ], [ [[ARRAYCTOR_NEXT:%.*]], [[ARRAYCTOR_LOOP]] ]
-// CHECK15-NEXT: call void @_ZN1SIfEC1Ev(ptr noundef nonnull align 4 dereferenceable(4) [[ARRAYCTOR_CUR]])
+// CHECK15-NEXT: call void @_ZN1SIfEC1Ev(ptr noundef nonnull align 4 dereferenceable(4) [[ARRAYCTOR_CUR]]) #[[ATTR4:[0-9]+]]
// CHECK15-NEXT: [[ARRAYCTOR_NEXT]] = getelementptr inbounds [[STRUCT_S]], ptr [[ARRAYCTOR_CUR]], i32 1
// CHECK15-NEXT: [[ARRAYCTOR_DONE:%.*]] = icmp eq ptr [[ARRAYCTOR_NEXT]], [[ARRAYCTOR_END]]
// CHECK15-NEXT: br i1 [[ARRAYCTOR_DONE]], label [[ARRAYCTOR_CONT:%.*]], label [[ARRAYCTOR_LOOP]]
// CHECK15: arrayctor.cont:
-// CHECK15-NEXT: call void @_ZN1SIfEC1Ev(ptr noundef nonnull align 4 dereferenceable(4) [[VAR]])
+// CHECK15-NEXT: call void @_ZN1SIfEC1Ev(ptr noundef nonnull align 4 dereferenceable(4) [[VAR]]) #[[ATTR4]]
// CHECK15-NEXT: [[TMP0:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 4
// CHECK15-NEXT: [[TMP1:%.*]] = load i32, ptr [[TMP0]], align 4
// CHECK15-NEXT: call void @__kmpc_for_static_init_4(ptr @[[GLOB1:[0-9]+]], i32 [[TMP1]], i32 92, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_COMB_LB]], ptr [[DOTOMP_COMB_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)
@@ -3585,14 +3585,14 @@ int main() {
// CHECK15-NEXT: store i32 2, ptr [[I]], align 4
// CHECK15-NEXT: br label [[DOTOMP_FINAL_DONE]]
// CHECK15: .omp.final.done:
-// CHECK15-NEXT: call void @_ZN1SIfED1Ev(ptr noundef nonnull align 4 dead_on_return(4) dereferenceable(4) [[VAR]]) #[[ATTR2:[0-9]+]]
+// CHECK15-NEXT: call void @_ZN1SIfED1Ev(ptr noundef nonnull align 4 dead_on_return(4) dereferenceable(4) [[VAR]]) #[[ATTR5:[0-9]+]]
// CHECK15-NEXT: [[ARRAY_BEGIN2:%.*]] = getelementptr inbounds [2 x [[STRUCT_S]]], ptr [[S_ARR]], i32 0, i32 0
// CHECK15-NEXT: [[TMP15:%.*]] = getelementptr inbounds [[STRUCT_S]], ptr [[ARRAY_BEGIN2]], i32 2
// CHECK15-NEXT: br label [[ARRAYDESTROY_BODY:%.*]]
// CHECK15: arraydestroy.body:
// CHECK15-NEXT: [[ARRAYDESTROY_ELEMENTPAST:%.*]] = phi ptr [ [[TMP15]], [[DOTOMP_FINAL_DONE]] ], [ [[ARRAYDESTROY_ELEMENT:%.*]], [[ARRAYDESTROY_BODY]] ]
// CHECK15-NEXT: [[ARRAYDESTROY_ELEMENT]] = getelementptr inbounds [[STRUCT_S]], ptr [[ARRAYDESTROY_ELEMENTPAST]], i32 -1
-// CHECK15-NEXT: call void @_ZN1SIfED1Ev(ptr noundef nonnull align 4 dead_on_return(4) dereferenceable(4) [[ARRAYDESTROY_ELEMENT]]) #[[ATTR2]]
+// CHECK15-NEXT: call void @_ZN1SIfED1Ev(ptr noundef nonnull align 4 dead_on_return(4) dereferenceable(4) [[ARRAYDESTROY_ELEMENT]]) #[[ATTR5]]
// CHECK15-NEXT: [[ARRAYDESTROY_DONE:%.*]] = icmp eq ptr [[ARRAYDESTROY_ELEMENT]], [[ARRAY_BEGIN2]]
// CHECK15-NEXT: br i1 [[ARRAYDESTROY_DONE]], label [[ARRAYDESTROY_DONE3:%.*]], label [[ARRAYDESTROY_BODY]]
// CHECK15: arraydestroy.done3:
@@ -3605,7 +3605,7 @@ int main() {
// CHECK15-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 4
// CHECK15-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 4
// CHECK15-NEXT: [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 4
-// CHECK15-NEXT: call void @_ZN1SIfEC2Ev(ptr noundef nonnull align 4 dereferenceable(4) [[THIS1]])
+// CHECK15-NEXT: call void @_ZN1SIfEC2Ev(ptr noundef nonnull align 4 dereferenceable(4) [[THIS1]]) #[[ATTR4]]
// CHECK15-NEXT: ret void
//
//
@@ -3645,12 +3645,12 @@ int main() {
// CHECK15-NEXT: br label [[ARRAYCTOR_LOOP:%.*]]
// CHECK15: arrayctor.loop:
// CHECK15-NEXT: [[ARRAYCTOR_CUR:%.*]] = phi ptr [ [[ARRAY_BEGIN]], [[ENTRY:%.*]] ], [ [[ARRAYCTOR_NEXT:%.*]], [[ARRAYCTOR_LOOP]] ]
-// CHECK15-NEXT: call void @_ZN1SIfEC1Ev(ptr noundef nonnull align 4 dereferenceable(4) [[ARRAYCTOR_CUR]])
+// CHECK15-NEXT: call void @_ZN1SIfEC1Ev(ptr noundef nonnull align 4 dereferenceable(4) [[ARRAYCTOR_CUR]]) #[[ATTR4]]
// CHECK15-NEXT: [[ARRAYCTOR_NEXT]] = getelementptr inbounds [[STRUCT_S]], ptr [[ARRAYCTOR_CUR]], i32 1
// CHECK15-NEXT: [[ARRAYCTOR_DONE:%.*]] = icmp eq ptr [[ARRAYCTOR_NEXT]], [[ARRAYCTOR_END]]
// CHECK15-NEXT: br i1 [[ARRAYCTOR_DONE]], label [[ARRAYCTOR_CONT:%.*]], label [[ARRAYCTOR_LOOP]]
// CHECK15: arrayctor.cont:
-// CHECK15-NEXT: call void @_ZN1SIfEC1Ev(ptr noundef nonnull align 4 dereferenceable(4) [[VAR]])
+// CHECK15-NEXT: call void @_ZN1SIfEC1Ev(ptr noundef nonnull align 4 dereferenceable(4) [[VAR]]) #[[ATTR4]]
// CHECK15-NEXT: [[TMP2:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 4
// CHECK15-NEXT: [[TMP3:%.*]] = load i32, ptr [[TMP2]], align 4
// CHECK15-NEXT: call void @__kmpc_for_static_init_4(ptr @[[GLOB2:[0-9]+]], i32 [[TMP3]], i32 34, ptr [[DOTOMP_IS_LAST]], ptr [[DOTOMP_LB]], ptr [[DOTOMP_UB]], ptr [[DOTOMP_STRIDE]], i32 1, i32 1)
@@ -3712,14 +3712,14 @@ int main() {
// CHECK15-NEXT: store i32 2, ptr [[I]], align 4
// CHECK15-NEXT: br label [[DOTOMP_FINAL_DONE]]
// CHECK15: .omp.final.done:
-// CHECK15-NEXT: call void @_ZN1SIfED1Ev(ptr noundef nonnull align 4 dead_on_return(4) dereferenceable(4) [[VAR]]) #[[ATTR2]]
+// CHECK15-NEXT: call void @_ZN1SIfED1Ev(ptr noundef nonnull align 4 dead_on_return(4) dereferenceable(4) [[VAR]]) #[[ATTR5]]
// CHECK15-NEXT: [[ARRAY_BEGIN5:%.*]] = getelementptr inbounds [2 x [[STRUCT_S]]], ptr [[S_ARR]], i32 0, i32 0
// CHECK15-NEXT: [[TMP20:%.*]] = getelementptr inbounds [[STRUCT_S]], ptr [[ARRAY_BEGIN5]], i32 2
// CHECK15-NEXT: br label [[ARRAYDESTROY_BODY:%.*]]
// CHECK15: arraydestroy.body:
// CHECK15-NEXT: [[ARRAYDESTROY_ELEMENTPAST:%.*]] = phi ptr [ [[TMP20]], [[DOTOMP_FINAL_DONE]] ], [ [[ARRAYDESTROY_ELEMENT:%.*]], [[ARRAYDESTROY_BODY]] ]
// CHECK15-NEXT: [[ARRAYDESTROY_ELEMENT]] = getelementptr inbounds [[STRUCT_S]], ptr [[ARRAYDESTROY_ELEMENTPAST]], i32 -1
-// CHECK15-NEXT: call void @_ZN1SIfED1Ev(ptr noundef nonnull align 4 dead_on_return(4) dereferenceable(4) [[ARRAYDESTROY_ELEMENT]]) #[[ATTR2]]
+// CHECK15-NEXT: call void @_ZN1SIfED1Ev(ptr noundef nonnull align 4 dead_on_return(4) dereferenceable(4) [[ARRAYDESTROY_ELEMENT]]) #[[ATTR5]]
// CHECK15-NEXT: [[ARRAYDESTROY_DONE:%.*]] = icmp eq ptr [[ARRAYDESTROY_ELEMENT]], [[ARRAY_BEGIN5]]
// CHECK15-NEXT: br i1 [[ARRAYDESTROY_DONE]], label [[ARRAYDESTROY_DONE6:%.*]], label [[ARRAYDESTROY_BODY]]
// CHECK15: arraydestroy.done6:
@@ -3732,7 +3732,7 @@ int main() {
// CHECK15-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 4
// CHECK15-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 4
// CHECK15-NEXT: [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 4
-// CHECK15-NEXT: call void @_ZN1SIfED2Ev(ptr noundef nonnull align 4 dead_on_return(4) dereferenceable(4) [[THIS1]]) #[[ATTR2]]
+// CHECK15-NEXT: call void @_ZN1SIfED2Ev(ptr noundef nonnull align 4 dead_on_return(4) dereferenceable(4) [[THIS1]]) #[[ATTR5]]
// CHECK15-NEXT: ret void
//
//
@@ -3781,12 +3781,12 @@ int main() {
// CHECK15-NEXT: br label [[ARRAYCTOR_LOOP:%.*]]
// CHECK15: arrayctor.loop:
// CHECK15-NEXT: [[ARRAYCTOR_CUR:%.*]] = phi ptr [ [[ARRAY_BEGIN]], [[ENTRY:%.*]] ], [ [[ARRAYCTOR_NEXT:%.*]], [[ARRAYCTOR_LOOP]] ]
-// CHECK15-NEXT: call void @_ZN1SIiEC1Ev(ptr noundef nonnull align 4 dereferenceable(4) [[ARRAYCTOR_CUR]])
+// CHECK15-NEXT: call void @_ZN1SIiEC1Ev(ptr noundef nonnull align 4 dereferenceable(4) [[ARRAYCTOR_CUR]]) #[[ATTR4]]
// CHECK15-NEXT: [[ARRAYCTOR_NEXT]] = getelementptr inbounds [[STRUCT_S_0]], ptr [[ARRAYCTOR_CUR]], i32 1
// CHECK15-NEXT: [[ARRAYCTOR_DONE:%.*]] = icmp eq ptr [[ARRAYCTOR_NEXT]], [[ARRAYCTOR_END]]
// CHECK15-NEXT: br i1 [[ARRAYCTOR_DONE]], label [[ARRAYCTOR_CONT:%.*]], label [[ARRAYCTOR_LOOP]]
// CHECK15: arrayctor.cont:
-// CHECK15-NEXT: call void @_ZN1SIiEC1Ev(ptr noundef nonnull align 4 dereferenceable(4) [[VAR]])
+// CHECK15-NEXT: call void @_ZN1SIiEC1Ev(ptr noundef nonnull align 4 dereferenceable(4) [[VAR]]) #[[ATTR4]]
// CHECK15-NEXT: store ptr [[VAR]], ptr [[_TMP2]], align 4
// CHECK15-NEXT: [[TMP0:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 4
// CHECK15-NEXT: [[TMP1:%.*]] = load i32, ptr [[TMP0]], align 4
@@ -3836,14 +3836,14 @@ int main() {
// CHECK15-NEXT: store i32 2, ptr [[I]], align 4
// CHECK15-NEXT: br label [[DOTOMP_FINAL_DONE]]
// CHECK15: .omp.final.done:
-// CHECK15-NEXT: call void @_ZN1SIiED1Ev(ptr noundef nonnull align 4 dead_on_return(4) dereferenceable(4) [[VAR]]) #[[ATTR2]]
+// CHECK15-NEXT: call void @_ZN1SIiED1Ev(ptr noundef nonnull align 4 dead_on_return(4) dereferenceable(4) [[VAR]]) #[[ATTR5]]
// CHECK15-NEXT: [[ARRAY_BEGIN4:%.*]] = getelementptr inbounds [2 x [[STRUCT_S_0]]], ptr [[S_ARR]], i32 0, i32 0
// CHECK15-NEXT: [[TMP15:%.*]] = getelementptr inbounds [[STRUCT_S_0]], ptr [[ARRAY_BEGIN4]], i32 2
// CHECK15-NEXT: br label [[ARRAYDESTROY_BODY:%.*]]
// CHECK15: arraydestroy.body:
// CHECK15-NEXT: [[ARRAYDESTROY_ELEMENTPAST:%.*]] = phi ptr [ [[TMP15]], [[DOTOMP_FINAL_DONE]] ], [ [[ARRAYDESTROY_ELEMENT:%.*]], [[ARRAYDESTROY_BODY]] ]
// CHECK15-NEXT: [[ARRAYDESTROY_ELEMENT]] = getelementptr inbounds [[STRUCT_S_0]], ptr [[ARRAYDESTROY_ELEMENTPAST]], i32 -1
-// CHECK15-NEXT: call void @_ZN1SIiED1Ev(ptr noundef nonnull align 4 dead_on_return(4) dereferenceable(4) [[ARRAYDESTROY_ELEMENT]]) #[[ATTR2]]
+// CHECK15-NEXT: call void @_ZN1SIiED1Ev(ptr noundef nonnull align 4 dead_on_return(4) dereferenceable(4) [[ARRAYDESTROY_ELEMENT]]) #[[ATTR5]]
// CHECK15-NEXT: [[ARRAYDESTROY_DONE:%.*]] = icmp eq ptr [[ARRAYDESTROY_ELEMENT]], [[ARRAY_BEGIN4]]
// CHECK15-NEXT: br i1 [[ARRAYDESTROY_DONE]], label [[ARRAYDESTROY_DONE5:%.*]], label [[ARRAYDESTROY_BODY]]
// CHECK15: arraydestroy.done5:
@@ -3856,7 +3856,7 @@ int main() {
// CHECK15-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 4
// CHECK15-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 4
// CHECK15-NEXT: [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 4
-// CHECK15-NEXT: call void @_ZN1SIiEC2Ev(ptr noundef nonnull align 4 dereferenceable(4) [[THIS1]])
+// CHECK15-NEXT: call void @_ZN1SIiEC2Ev(ptr noundef nonnull align 4 dereferenceable(4) [[THIS1]]) #[[ATTR4]]
// CHECK15-NEXT: ret void
//
//
@@ -3898,12 +3898,12 @@ int main() {
// CHECK15-NEXT: br label [[ARRAYCTOR_LOOP:%.*]]
// CHECK15: arrayctor.loop:
// CHECK15-NEXT: [[ARRAYCTOR_CUR:%.*]] = phi ptr [ [[ARRAY_BEGIN]], [[ENTRY:%.*]] ], [ [[ARRAYCTOR_NEXT:%.*]], [[ARRAYCTOR_LOOP]] ]
-// CHECK15-NEXT: call void @_ZN1SIiEC1Ev(ptr noundef nonnull align 4 dereferenceable(4) [[ARRAYCTOR_CUR]])
+// CHECK15-NEXT: call void @_ZN1SIiEC1Ev(ptr noundef nonnull align 4 dereferenceable(4) [[ARRAYCTOR_CUR]]) #[[ATTR4]]
// CHECK15-NEXT: [[ARRAYCTOR_NEXT]] = getelementptr inbounds [[STRUCT_S_0]], ptr [[ARRAYCTOR_CUR]], i32 1
// CHECK15-NEXT: [[ARRAYCTOR_DONE:%.*]] = icmp eq ptr [[ARRAYCTOR_NEXT]], [[ARRAYCTOR_END]]
// CHECK15-NEXT: br i1 [[ARRAYCTOR_DONE]], label [[ARRAYCTOR_CONT:%.*]], label [[ARRAYCTOR_LOOP]]
// CHECK15: arrayctor.cont:
-// CHECK15-NEXT: call void @_ZN1SIiEC1Ev(ptr noundef nonnull align 4 dereferenceable(4) [[VAR]])
+// CHECK15-NEXT: call void @_ZN1SIiEC1Ev(ptr noundef nonnull align 4 dereferenceable(4) [[VAR]]) #[[ATTR4]]
// CHECK15-NEXT: store ptr [[VAR]], ptr [[_TMP2]], align 4
// CHECK15-NEXT: [[TMP2:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 4
// CHECK15-NEXT: [[TMP3:%.*]] = load i32, ptr [[TMP2]], align 4
@@ -3963,14 +3963,14 @@ int main() {
// CHECK15-NEXT: store i32 2, ptr [[I]], align 4
// CHECK15-NEXT: br label [[DOTOMP_FINAL_DONE]]
// CHECK15: .omp.final.done:
-// CHECK15-NEXT: call void @_ZN1SIiED1Ev(ptr noundef nonnull align 4 dead_on_return(4) dereferenceable(4) [[VAR]]) #[[ATTR2]]
+// CHECK15-NEXT: call void @_ZN1SIiED1Ev(ptr noundef nonnull align 4 dead_on_return(4) dereferenceable(4) [[VAR]]) #[[ATTR5]]
// CHECK15-NEXT: [[ARRAY_BEGIN6:%.*]] = getelementptr inbounds [2 x [[STRUCT_S_0]]], ptr [[S_ARR]], i32 0, i32 0
// CHECK15-NEXT: [[TMP19:%.*]] = getelementptr inbounds [[STRUCT_S_0]], ptr [[ARRAY_BEGIN6]], i32 2
// CHECK15-NEXT: br label [[ARRAYDESTROY_BODY:%.*]]
// CHECK15: arraydestroy.body:
// CHECK15-NEXT: [[ARRAYDESTROY_ELEMENTPAST:%.*]] = phi ptr [ [[TMP19]], [[DOTOMP_FINAL_DONE]] ], [ [[ARRAYDESTROY_ELEMENT:%.*]], [[ARRAYDESTROY_BODY]] ]
// CHECK15-NEXT: [[ARRAYDESTROY_ELEMENT]] = getelementptr inbounds [[STRUCT_S_0]], ptr [[ARRAYDESTROY_ELEMENTPAST]], i32 -1
-// CHECK15-NEXT: call void @_ZN1SIiED1Ev(ptr noundef nonnull align 4 dead_on_return(4) dereferenceable(4) [[ARRAYDESTROY_ELEMENT]]) #[[ATTR2]]
+// CHECK15-NEXT: call void @_ZN1SIiED1Ev(ptr noundef nonnull align 4 dead_on_return(4) dereferenceable(4) [[ARRAYDESTROY_ELEMENT]]) #[[ATTR5]]
// CHECK15-NEXT: [[ARRAYDESTROY_DONE:%.*]] = icmp eq ptr [[ARRAYDESTROY_ELEMENT]], [[ARRAY_BEGIN6]]
// CHECK15-NEXT: br i1 [[ARRAYDESTROY_DONE]], label [[ARRAYDESTROY_DONE7:%.*]], label [[ARRAYDESTROY_BODY]]
// CHECK15: arraydestroy.done7:
@@ -3983,7 +3983,7 @@ int main() {
// CHECK15-NEXT: [[THIS_ADDR:%.*]] = alloca ptr, align 4
// CHECK15-NEXT: store ptr [[THIS]], ptr [[THIS_ADDR]], align 4
// CHECK15-NEXT: [[THIS1:%.*]] = load ptr, ptr [[THIS_ADDR]], align 4
-// CHECK15-NEXT: call void @_ZN1SIiED2Ev(ptr noundef nonnull align 4 dead_on_return(4) dereferenceable(4) [[THIS1]]) #[[ATTR2]]
+// CHECK15-NEXT: call void @_ZN1SIiED2Ev(ptr noundef nonnull align 4 dead_on_return(4) dereferenceable(4) [[THIS1]]) #[[ATTR5]]
// CHECK15-NEXT: ret void
//
//
@@ -4193,7 +4193,7 @@ int main() {
// CHECK17-NEXT: store ptr [[TMP13]], ptr [[TMP12]], align 8, !llvm.access.group [[ACC_GRP9]]
// CHECK17-NEXT: [[TMP14:%.*]] = getelementptr inbounds nuw [[CLASS_ANON]], ptr [[REF_TMP]], i32 0, i32 2
// CHECK17-NEXT: store ptr [[SIVAR]], ptr [[TMP14]], align 8, !llvm.access.group [[ACC_GRP9]]
-// CHECK17-NEXT: call void @"_ZZZ4mainENK3$_0clEvENKUlvE_clEv"(ptr noundef nonnull align 8 dereferenceable(24) [[REF_TMP]]), !llvm.access.group [[ACC_GRP9]]
+// CHECK17-NEXT: call void @"_ZZZ4mainENK3$_0clEvENKUlvE_clEv"(ptr noundef nonnull align 8 dereferenceable(24) [[REF_TMP]]) #[[ATTR3:[0-9]+]], !llvm.access.group [[ACC_GRP9]]
// CHECK17-NEXT: br label [[OMP_BODY_CONTINUE:%.*]]
// CHECK17: omp.body.continue:
// CHECK17-NEXT: br label [[OMP_INNER_FOR_INC:%.*]]
>From 7fbbf2e1d714ec6f41e33e9ac79efcb54883ce2e Mon Sep 17 00:00:00 2001
From: Joseph Huber <huberjn at outlook.com>
Date: Sat, 5 Sep 2026 17:17:18 -0500
Subject: [PATCH 5/5] use target hook isntead of a triple
---
clang/include/clang/Options/Options.td | 3 ++-
clang/lib/CodeGen/CGCall.cpp | 7 +++----
clang/lib/CodeGen/TargetInfo.h | 3 +++
clang/lib/CodeGen/Targets/AMDGPU.cpp | 1 +
clang/lib/CodeGen/Targets/NVPTX.cpp | 1 +
clang/lib/CodeGen/Targets/SPIR.cpp | 4 ++++
6 files changed, 14 insertions(+), 5 deletions(-)
diff --git a/clang/include/clang/Options/Options.td b/clang/include/clang/Options/Options.td
index 72590ab683007..3d4b05add57e7 100644
--- a/clang/include/clang/Options/Options.td
+++ b/clang/include/clang/Options/Options.td
@@ -1151,7 +1151,8 @@ defm offload_uniform_block : BoolFOption<"offload-uniform-block",
Default<"LangOpts->defaultOffloadUniformBlock()">,
PosFlag<SetTrue, [], [ClangOption, CC1Option], "Assume">,
NegFlag<SetFalse, [], [ClangOption, CC1Option], "Don't assume">,
- BothFlags<[], [ClangOption], " that kernels are launched with uniform block sizes">>;
+ BothFlags<[], [ClangOption], " that kernels are launched with uniform block sizes "
+ "(OpenCL 2.0 and later instead default to non-uniform, see -cl-uniform-work-group-size)">>;
def fcomplex_arithmetic_EQ : Joined<["-"], "fcomplex-arithmetic=">, Group<f_Group>,
Visibility<[ClangOption, CC1Option, FlangOption, FC1Option]>,
diff --git a/clang/lib/CodeGen/CGCall.cpp b/clang/lib/CodeGen/CGCall.cpp
index 51d3e379a1452..8ba01b05e5466 100644
--- a/clang/lib/CodeGen/CGCall.cpp
+++ b/clang/lib/CodeGen/CGCall.cpp
@@ -2865,11 +2865,10 @@ void CodeGenModule::ConstructAttributeList(StringRef Name,
NumElemsParam);
}
- // Leftover work-groups are an OpenCL 2.0+ language feature. Host C/C++
- // inherit OffloadUniformBlock but must not receive the GPU ABI attribute.
if (getLangOpts().OffloadUniformBlock &&
- (getLangOpts().CUDA || getLangOpts().OpenCL ||
- getLangOpts().isTargetDevice() || getTarget().getTriple().isGPU()))
+ (getLangOpts().OpenCL || getLangOpts().CUDA ||
+ getLangOpts().isTargetDevice() ||
+ getTargetCodeGenInfo().hasUniformWorkGroupLaunch()))
FuncAttrs.addAttribute("uniform-work-group-size");
if (TargetDecl->hasAttr<ArmLocallyStreamingAttr>())
diff --git a/clang/lib/CodeGen/TargetInfo.h b/clang/lib/CodeGen/TargetInfo.h
index fae32068b1570..d9a63bc232101 100644
--- a/clang/lib/CodeGen/TargetInfo.h
+++ b/clang/lib/CodeGen/TargetInfo.h
@@ -334,6 +334,9 @@ class TargetCodeGenInfo {
/// Get LLVM calling convention for device kernels.
virtual unsigned getDeviceKernelCallingConv() const;
+ /// \return true if the target's launch model delivers complete work-groups.
+ virtual bool hasUniformWorkGroupLaunch() const { return false; }
+
/// Get target specific null pointer.
/// \param T is the LLVM type of the null pointer.
/// \param QT is the clang QualType of the null pointer.
diff --git a/clang/lib/CodeGen/Targets/AMDGPU.cpp b/clang/lib/CodeGen/Targets/AMDGPU.cpp
index 07e2eac39305d..0ceb477eea4b7 100644
--- a/clang/lib/CodeGen/Targets/AMDGPU.cpp
+++ b/clang/lib/CodeGen/Targets/AMDGPU.cpp
@@ -302,6 +302,7 @@ class AMDGPUTargetCodeGenInfo : public TargetCodeGenInfo {
void setTargetAttributes(const Decl *D, llvm::GlobalValue *GV,
CodeGen::CodeGenModule &M) const override;
unsigned getDeviceKernelCallingConv() const override;
+ bool hasUniformWorkGroupLaunch() const override { return true; }
llvm::Constant *getNullPointer(const CodeGen::CodeGenModule &CGM,
llvm::PointerType *T, QualType QT) const override;
diff --git a/clang/lib/CodeGen/Targets/NVPTX.cpp b/clang/lib/CodeGen/Targets/NVPTX.cpp
index 00fe271c85b0b..aac4dd82a2e61 100644
--- a/clang/lib/CodeGen/Targets/NVPTX.cpp
+++ b/clang/lib/CodeGen/Targets/NVPTX.cpp
@@ -51,6 +51,7 @@ class NVPTXTargetCodeGenInfo : public TargetCodeGenInfo {
void setTargetAttributes(const Decl *D, llvm::GlobalValue *GV,
CodeGen::CodeGenModule &M) const override;
bool shouldEmitStaticExternCAliases() const override;
+ bool hasUniformWorkGroupLaunch() const override { return true; }
StringRef getLLVMSyncScopeStr(const LangOptions &LangOpts, SyncScope Scope,
llvm::AtomicOrdering Ordering) const override;
diff --git a/clang/lib/CodeGen/Targets/SPIR.cpp b/clang/lib/CodeGen/Targets/SPIR.cpp
index 7f7f1a2a1fe8c..67b6538789a14 100644
--- a/clang/lib/CodeGen/Targets/SPIR.cpp
+++ b/clang/lib/CodeGen/Targets/SPIR.cpp
@@ -140,6 +140,10 @@ class SPIRVTargetCodeGenInfo : public CommonSPIRTargetCodeGenInfo {
return getABIInfo().getTarget().getTriple().getVendor() !=
llvm::Triple::AMD;
}
+ bool hasUniformWorkGroupLaunch() const override {
+ return getABIInfo().getTarget().getTriple().getVendor() ==
+ llvm::Triple::AMD;
+ }
LangAS getSRetAddrSpace(const CXXRecordDecl *RD) const override;
};
More information about the cfe-commits
mailing list