llvmorg-github-actions[bot] wrote:
<!--LLVM PR SUMMARY COMMENT--> @llvm/pr-subscribers-backend-amdgpu Author: Joseph Huber (jhuber6) <details> <summary>Changes</summary> 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` --- Patch is 35.38 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/221279.diff 20 Files Affected: - (modified) clang/include/clang/Basic/LangOptions.def (+1-1) - (modified) clang/include/clang/Basic/LangOptions.h (+3) - (modified) clang/include/clang/Options/Options.td (+3-2) - (modified) clang/lib/Basic/LangOptions.cpp (+9) - (modified) clang/lib/Frontend/CompilerInvocation.cpp (+3) - (modified) clang/test/CodeGen/amdgpu-abi-version.c (+19-37) - (modified) clang/test/CodeGen/amdgpu-address-spaces.cpp (+1-1) - (modified) clang/test/CodeGen/amdgpu-builtin-is-invocable-subarch.c (+2-2) - (modified) clang/test/CodeGen/amdgpu-builtin-is-invocable.c (+3-3) - (modified) clang/test/CodeGen/amdgpu-builtin-processor-is-subarch.c (+3-3) - (modified) clang/test/CodeGen/amdgpu-builtin-processor-is.c (+3-3) - (added) clang/test/CodeGen/amdgpu-offload-uniform-block.c (+17) - (modified) clang/test/CodeGen/nvptx_attributes.c (+1-1) - (modified) clang/test/CodeGenCXX/dynamic-cast-address-space.cpp (+4-4) - (modified) clang/test/CodeGenSYCL/function-attrs.cpp (+2-2) - (modified) clang/test/Headers/gpuintrin.c (+25-55) - (modified) clang/test/OpenMP/amdgcn-attributes.cpp (+2-2) - (modified) clang/test/OpenMP/amdgcn_target_global_constructor.cpp (+4-4) - (modified) clang/test/OpenMP/amdgcn_weak_alias.c (+1-1) - (modified) clang/test/OpenMP/amdgcn_weak_alias.cpp (+1-1) ``````````diff 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-trap... [truncated] `````````` </details> https://github.com/llvm/llvm-project/pull/221279 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
