[llvm] 3590945 - [AMDGPU] Add attribute to AMDGPU ctor / dtor to indicate single threadedness

[Joseph Huber via llvm-commits]

Author: Joseph Huber
Date: 2023-05-24T07:24:17-05:00
New Revision: 3590945a11c79a539738dcecd5d7706b6449d0c4

URL: https://github.com/llvm/llvm-project/commit/3590945a11c79a539738dcecd5d7706b6449d0c4
DIFF: https://github.com/llvm/llvm-project/commit/3590945a11c79a539738dcecd5d7706b6449d0c4.diff

LOG: [AMDGPU] Add attribute to AMDGPU ctor / dtor to indicate single threadedness

We only expect these ctor / dtor functions to be called with a single
thread. Add the appropriate attributes to indicate this to the backend.

Reviewed By: arsenm

Differential Revision: https://reviews.llvm.org/D151153

Added: 
    

Modified: 
    llvm/lib/Target/AMDGPU/AMDGPUCtorDtorLowering.cpp
    llvm/test/CodeGen/AMDGPU/lower-ctor-dtor-constexpr-alias.ll
    llvm/test/CodeGen/AMDGPU/lower-ctor-dtor.ll
    llvm/test/CodeGen/AMDGPU/lower-multiple-ctor-dtor.ll

Removed: 
    


################################################################################
diff  --git a/llvm/lib/Target/AMDGPU/AMDGPUCtorDtorLowering.cpp b/llvm/lib/Target/AMDGPU/AMDGPUCtorDtorLowering.cpp
index 05786d64d5e7c..a13447586bd4b 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPUCtorDtorLowering.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPUCtorDtorLowering.cpp
@@ -38,6 +38,7 @@ static Function *createInitOrFiniKernelFunction(Module &M, bool IsCtor) {
       FunctionType::get(Type::getVoidTy(M.getContext()), false),
       GlobalValue::WeakODRLinkage, 0, InitOrFiniKernelName, &M);
   InitOrFiniKernel->setCallingConv(CallingConv::AMDGPU_KERNEL);
+  InitOrFiniKernel->addFnAttr("amdgpu-flat-work-group-size", "1,1");
   if (IsCtor)
     InitOrFiniKernel->addFnAttr("device-init");
   else
@@ -58,6 +59,7 @@ static Function *createInitOrFiniKernelFunction(Module &M, bool IsCtor) {
 // void call_init_array_callbacks() {
 //   for (auto start = __init_array_start; start != __init_array_end; ++start)
 //     reinterpret_cast<InitCallback *>(*start)();
+// }
 static void createInitOrFiniCalls(Function &F, bool IsCtor) {
   Module &M = *F.getParent();
   LLVMContext &C = M.getContext();

diff  --git a/llvm/test/CodeGen/AMDGPU/lower-ctor-dtor-constexpr-alias.ll b/llvm/test/CodeGen/AMDGPU/lower-ctor-dtor-constexpr-alias.ll
index 35849d318f05b..344ee62b44065 100644
--- a/llvm/test/CodeGen/AMDGPU/lower-ctor-dtor-constexpr-alias.ll
+++ b/llvm/test/CodeGen/AMDGPU/lower-ctor-dtor-constexpr-alias.ll
@@ -64,5 +64,5 @@ define void @bar() addrspace(1) {
 ; CHECK:       while.end:
 ; CHECK-NEXT:    ret void
 
-; CHECK: attributes #[[ATTR0:[0-9]+]] = { "device-init" }
-; CHECK: attributes #[[ATTR1:[0-9]+]] = { "device-fini" }
+; CHECK: attributes #[[ATTR0:[0-9]+]] = { "amdgpu-flat-work-group-size"="1,1" "device-init" }
+; CHECK: attributes #[[ATTR1:[0-9]+]] = { "amdgpu-flat-work-group-size"="1,1" "device-fini" }

diff  --git a/llvm/test/CodeGen/AMDGPU/lower-ctor-dtor.ll b/llvm/test/CodeGen/AMDGPU/lower-ctor-dtor.ll
index a8eec5cbf1969..5c8e56dd93933 100644
--- a/llvm/test/CodeGen/AMDGPU/lower-ctor-dtor.ll
+++ b/llvm/test/CodeGen/AMDGPU/lower-ctor-dtor.ll
@@ -6,6 +6,7 @@
 ; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx700 -filetype=obj -o - < %s | llvm-readelf -s - 2>&1 | FileCheck %s -check-prefix=VISIBILITY
 ; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx700 -filetype=obj -o - < %s | llvm-readelf -S - 2>&1 | FileCheck %s -check-prefix=SECTION
 ; RUN: llc -mtriple=amdgcn-amd-amdhsa -amdgpu-lower-global-ctor-dtor=0 -mcpu=gfx700 -filetype=obj -o - < %s | llvm-readelf -s - 2>&1 | FileCheck %s -check-prefix=DISABLED
+; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx700 -filetype=obj -o - < %s | llvm-readelf --notes - 2>&1 | FileCheck %s -check-prefix=METADATA
 
 @llvm.global_ctors = appending addrspace(1) global [1 x { i32, ptr, ptr }] [{ i32, ptr, ptr } { i32 1, ptr @foo, ptr null }]
 @llvm.global_dtors = appending addrspace(1) global [1 x { i32, ptr, ptr }] [{ i32, ptr, ptr } { i32 1, ptr @bar, ptr null }]
@@ -54,6 +55,15 @@
 ; DISABLED-NOT: OBJECT GLOBAL DEFAULT {{.*}} amdgcn.device.init.kd
 ; DISABLED-NOT: FUNC   GLOBAL PROTECTED {{.*}} amdgcn.device.fini
 ; DISABLED-NOT: OBJECT   GLOBAL DEFAULT {{.*}} amdgcn.device.fini.kd
+; METADATA:  amdhsa.kernels:
+; METADATA:    .kind:           init
+; METADATA:    .max_flat_workgroup_size: 1
+; METADATA:    .name:           amdgcn.device.init
+; METADATA:    .symbol:         amdgcn.device.init.kd
+; METADATA:    .kind:           fini
+; METADATA:    .max_flat_workgroup_size: 1
+; METADATA:    .name:           amdgcn.device.fini
+; METADATA:    .symbol:         amdgcn.device.fini.kd
 
 define internal void @foo() {
   ret void
@@ -63,5 +73,5 @@ define internal void @bar() {
   ret void
 }
 
-; CHECK: attributes #0 = { "device-init" }
-; CHECK: attributes #1 = { "device-fini" }
+; CHECK: attributes #0 = { "amdgpu-flat-work-group-size"="1,1" "device-init" }
+; CHECK: attributes #1 = { "amdgpu-flat-work-group-size"="1,1" "device-fini" }

diff  --git a/llvm/test/CodeGen/AMDGPU/lower-multiple-ctor-dtor.ll b/llvm/test/CodeGen/AMDGPU/lower-multiple-ctor-dtor.ll
index a5286ec45dee1..57362b71cf665 100644
--- a/llvm/test/CodeGen/AMDGPU/lower-multiple-ctor-dtor.ll
+++ b/llvm/test/CodeGen/AMDGPU/lower-multiple-ctor-dtor.ll
@@ -57,5 +57,5 @@ define internal void @bar.5() {
   ret void
 }
 
-; CHECK: attributes #0 = { "device-init" }
-; CHECK: attributes #1 = { "device-fini" }
+; CHECK: attributes #0 = { "amdgpu-flat-work-group-size"="1,1" "device-init" }
+; CHECK: attributes #1 = { "amdgpu-flat-work-group-size"="1,1" "device-fini" }