r370808 - [AMDGPU] Set default flat work group size to (1, 256) for HIP

[Yaxun Liu via cfe-commits]

Author: yaxunl
Date: Tue Sep  3 11:50:24 2019
New Revision: 370808

URL: http://llvm.org/viewvc/llvm-project?rev=370808&view=rev
Log:
[AMDGPU] Set default flat work group size to (1,256) for HIP

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

Modified:
    cfe/trunk/lib/CodeGen/TargetInfo.cpp
    cfe/trunk/test/CodeGenCUDA/kernel-amdgcn.cu

Modified: cfe/trunk/lib/CodeGen/TargetInfo.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/TargetInfo.cpp?rev=370808&r1=370807&r2=370808&view=diff
==============================================================================
--- cfe/trunk/lib/CodeGen/TargetInfo.cpp (original)
+++ cfe/trunk/lib/CodeGen/TargetInfo.cpp Tue Sep  3 11:50:24 2019
@@ -7915,8 +7915,9 @@ void AMDGPUTargetCodeGenInfo::setTargetA
 
   const bool IsOpenCLKernel = M.getLangOpts().OpenCL &&
                               FD->hasAttr<OpenCLKernelAttr>();
-  if ((IsOpenCLKernel ||
-       (M.getLangOpts().HIP && FD->hasAttr<CUDAGlobalAttr>())) &&
+  const bool IsHIPKernel = M.getLangOpts().HIP &&
+                           FD->hasAttr<CUDAGlobalAttr>();
+  if ((IsOpenCLKernel || IsHIPKernel) &&
       (M.getTriple().getOS() == llvm::Triple::AMDHSA))
     F->addFnAttr("amdgpu-implicitarg-num-bytes", "56");
 
@@ -7942,7 +7943,7 @@ void AMDGPUTargetCodeGenInfo::setTargetA
       F->addFnAttr("amdgpu-flat-work-group-size", AttrVal);
     } else
       assert(Max == 0 && "Max must be zero");
-  } else if (IsOpenCLKernel) {
+  } else if (IsOpenCLKernel || IsHIPKernel) {
     // By default, restrict the maximum size to 256.
     F->addFnAttr("amdgpu-flat-work-group-size", "1,256");
   }

Modified: cfe/trunk/test/CodeGenCUDA/kernel-amdgcn.cu
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenCUDA/kernel-amdgcn.cu?rev=370808&r1=370807&r2=370808&view=diff
==============================================================================
--- cfe/trunk/test/CodeGenCUDA/kernel-amdgcn.cu (original)
+++ cfe/trunk/test/CodeGenCUDA/kernel-amdgcn.cu Tue Sep  3 11:50:24 2019
@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 -triple amdgcn -fcuda-is-device -emit-llvm %s -o - | FileCheck %s
+// RUN: %clang_cc1 -triple amdgcn -fcuda-is-device -emit-llvm -x hip %s -o - | FileCheck %s
 #include "Inputs/cuda.h"
 
 // CHECK: define amdgpu_kernel void @_ZN1A6kernelEv
@@ -25,7 +25,7 @@ struct Dummy {
   EmptyKernelPtr Empty() { return EmptyKernel<void>; } 
 };
 
-// CHECK: define amdgpu_kernel void @_Z15template_kernelI1AEvT_
+// CHECK: define amdgpu_kernel void @_Z15template_kernelI1AEvT_{{.*}} #[[ATTR:[0-9][0-9]*]]
 template<class T>
 __global__ void template_kernel(T x) {}
 
@@ -39,3 +39,4 @@ int main() {
   launch((void*)D.Empty());
   return 0;
 }
+// CHECK: attributes #[[ATTR]] = {{.*}}"amdgpu-flat-work-group-size"="1,256"