[clang] 0ffb12c - [HIP] Mark kernels with uniform-work-group-size=true
Yaxun Liu via cfe-commits
cfe-commits at lists.llvm.org
Fri Mar 13 04:24:51 PDT 2020
Author: Yaxun (Sam) Liu
Date: 2020-03-13T06:56:56-04:00
New Revision: 0ffb12ca67fd813a8ae840399626dd5f8fea3178
URL: https://github.com/llvm/llvm-project/commit/0ffb12ca67fd813a8ae840399626dd5f8fea3178
DIFF: https://github.com/llvm/llvm-project/commit/0ffb12ca67fd813a8ae840399626dd5f8fea3178.diff
LOG: [HIP] Mark kernels with uniform-work-group-size=true
Differential Revision: https://reviews.llvm.org/D76076
Added:
Modified:
clang/lib/CodeGen/TargetInfo.cpp
clang/test/CodeGenCUDA/amdgpu-kernel-attrs.cu
Removed:
################################################################################
diff --git a/clang/lib/CodeGen/TargetInfo.cpp b/clang/lib/CodeGen/TargetInfo.cpp
index 21d5cd08c9fa..fb472d541160 100644
--- a/clang/lib/CodeGen/TargetInfo.cpp
+++ b/clang/lib/CodeGen/TargetInfo.cpp
@@ -8091,6 +8091,10 @@ void AMDGPUTargetCodeGenInfo::setTargetAttributes(
(M.getTriple().getOS() == llvm::Triple::AMDHSA))
F->addFnAttr("amdgpu-implicitarg-num-bytes", "56");
+ if (IsHIPKernel)
+ F->addFnAttr("uniform-work-group-size", "true");
+
+
const auto *FlatWGS = FD->getAttr<AMDGPUFlatWorkGroupSizeAttr>();
if (ReqdWGS || FlatWGS) {
unsigned Min = 0;
diff --git a/clang/test/CodeGenCUDA/amdgpu-kernel-attrs.cu b/clang/test/CodeGenCUDA/amdgpu-kernel-attrs.cu
index ece8685932d2..5415bddffc89 100644
--- a/clang/test/CodeGenCUDA/amdgpu-kernel-attrs.cu
+++ b/clang/test/CodeGenCUDA/amdgpu-kernel-attrs.cu
@@ -39,7 +39,7 @@ __global__ void num_vgpr_64() {
// NAMD-NOT: "amdgpu-num-vgpr"
// NAMD-NOT: "amdgpu-num-sgpr"
-// DEFAULT-DAG: attributes [[FLAT_WORK_GROUP_SIZE_DEFAULT]] = {{.*}}"amdgpu-flat-work-group-size"="1,256"
+// DEFAULT-DAG: attributes [[FLAT_WORK_GROUP_SIZE_DEFAULT]] = {{.*}}"amdgpu-flat-work-group-size"="1,256"{{.*}}"uniform-work-group-size"="true"
// MAX1024-DAG: attributes [[FLAT_WORK_GROUP_SIZE_DEFAULT]] = {{.*}}"amdgpu-flat-work-group-size"="1,1024"
// CHECK-DAG: attributes [[FLAT_WORK_GROUP_SIZE_32_64]] = {{.*}}"amdgpu-flat-work-group-size"="32,64"
// CHECK-DAG: attributes [[WAVES_PER_EU_2]] = {{.*}}"amdgpu-waves-per-eu"="2"
More information about the cfe-commits
mailing list