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

[jhuber6]

Summary

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

[arsenm]

Can you check the device metadata for these in the codegen output

[jhuber6]

Adding check on the elf notes

[arsenm]

Should probably document this in AMDGPUUsage. Is this really true for the asan usage?

[jhuber6]

@ampandey-AMD
Is it true that asan and others only runs these with a single thread?

[arsenm]

llvm/lib/Target/AMDGPU/AMDGPUCtorDtorLowering.cpp

41

We should probably teach UniformityAnalysis about this special case for the work item IDs

[jhuber6]

llvm/lib/Target/AMDGPU/AMDGPUCtorDtorLowering.cpp

41

This special case where the number of threads is one? It would definitely simplify anything to do with convergence.

llvm/lib/Target/AMDGPU/AMDGPUCtorDtorLowering.cpp

if (!IsCtor)
     InitOrFiniKernelName = "amdgcn.device.fini";
   if (M.getFunction(InitOrFiniKernelName))
     return nullptr;
 
   Function *InitOrFiniKernel = Function::createWithDefaultAttr(
       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
     InitOrFiniKernel->addFnAttr("device-fini");
   return InitOrFiniKernel;
 }
 
 // The linker will provide the associated symbols to allow us to traverse the
 // global constructors / destructors in priority order. We create the IR
 // required to call each callback in this section. This is equivalent to the
 // following code.
 //
 // extern "C" void * __init_array_start[];
 // extern "C" void * __init_array_end[];
 //
 // using InitCallback = void();
 //
 // 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();
 
   IRBuilder<> IRB(BasicBlock::Create(C, "entry", &F));
   auto *LoopBB = BasicBlock::Create(C, "while.entry", &F);
   auto *ExitBB = BasicBlock::Create(C, "while.end", &F);
   Type *PtrTy = IRB.getPtrTy(AMDGPUAS::GLOBAL_ADDRESS);

llvm/test/CodeGen/AMDGPU/lower-ctor-dtor-constexpr-alias.ll

 ; CHECK-NEXT:    [[CALLBACK:%.*]] = load ptr, ptr addrspace(1) [[PTR]], align 8
 ; CHECK-NEXT:    call void [[CALLBACK]]()
 ; CHECK-NEXT:    [[NEXT]] = getelementptr ptr addrspace(1), ptr addrspace(1) [[PTR]], i64 1
 ; CHECK-NEXT:    [[END:%.*]] = icmp eq ptr addrspace(1) [[NEXT]], @__fini_array_end
 ; CHECK-NEXT:    br i1 [[END]], label [[WHILE_END]], label [[WHILE_ENTRY]]
 ; 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" }

llvm/test/CodeGen/AMDGPU/lower-ctor-dtor.ll

 ; RUN: opt -S -mtriple=amdgcn-- -amdgpu-lower-ctor-dtor < %s | FileCheck %s
 ; RUN: opt -S -mtriple=amdgcn-- -passes=amdgpu-lower-ctor-dtor < %s | FileCheck %s
 
 ; Make sure we get the same result if we run multiple times
 ; RUN: opt -S -mtriple=amdgcn-- -passes=amdgpu-lower-ctor-dtor,amdgpu-lower-ctor-dtor < %s | FileCheck %s
 ; 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 }]
 
 ; CHECK: @__init_array_start = external addrspace(1) constant [0 x ptr addrspace(1)]
 ; CHECK: @__init_array_end = external addrspace(1) constant [0 x ptr addrspace(1)]
 ; CHECK: @__fini_array_start = external addrspace(1) constant [0 x ptr addrspace(1)]
 ; CHECK: @__fini_array_end = external addrspace(1) constant [0 x ptr addrspace(1)]

 ; VISIBILITY: FUNC   WEAK PROTECTED {{.*}} amdgcn.device.fini
 ; VISIBILITY: OBJECT   WEAK DEFAULT {{.*}} amdgcn.device.fini.kd
 ; SECTION: .init_array.1     INIT_ARRAY      {{.*}} {{.*}} 000008 00  WA  0   0  8
 ; SECTION: .fini_array.1     FINI_ARRAY      {{.*}} {{.*}} 000008 00  WA  0   0  8
 ; DISABLED-NOT: FUNC   GLOBAL PROTECTED {{.*}} amdgcn.device.init
 ; 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
 }
 
 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" }

llvm/test/CodeGen/AMDGPU/lower-multiple-ctor-dtor.ll

 define internal void @foo.5() {
   ret void
 }
 
 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" }