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 <[email protected]> Date: Fri, 4 Sep 2026 11:27:24 -0500 Subject: [PATCH 1/2] [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 6c544c7f154a1523ca05c78276428d84fe9e5619 Mon Sep 17 00:00:00 2001 From: Joseph Huber <[email protected]> Date: Fri, 4 Sep 2026 11:52:16 -0500 Subject: [PATCH 2/2] attempt to use language --- clang/include/clang/Basic/LangOptions.def | 2 +- clang/include/clang/Basic/LangOptions.h | 4 ++-- 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, 16 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..83e3dae12898c 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. - 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" _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
