[clang] [Clang] Default OpenMP and C/C++ compilations to uniform blocks (PR #221279)
Joseph Huber via cfe-commits
cfe-commits at lists.llvm.org
Fri Sep 4 13:57:10 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/3] [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/3] 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/3] 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]>,
More information about the cfe-commits
mailing list