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

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

git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@370808 91177308-0d34-0410-b5e6-96231b3b80d8

diff --git a/lib/CodeGen/TargetInfo.cpp b/lib/CodeGen/TargetInfo.cpp
index 231a20c68f5a6bc7359f345fc9ac7c7ce995eaa9..760895a493ff2d5ec7fc206368ef4785ca18ca7e 100644 (file)
--- a/lib/CodeGen/TargetInfo.cpp
+++ b/lib/CodeGen/TargetInfo.cpp
@@ -7915,8 +7915,9 @@ void AMDGPUTargetCodeGenInfo::setTargetAttributes(
 
   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::setTargetAttributes(
       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");
   }

diff --git a/test/CodeGenCUDA/kernel-amdgcn.cu b/test/CodeGenCUDA/kernel-amdgcn.cu
index ffa6c9549f07338fd5ef477803c0c82eb5335b42..135d3030480c6ae95c947401f37fe680b75614a9 100644 (file)
--- a/test/CodeGenCUDA/kernel-amdgcn.cu
+++ b/test/CodeGenCUDA/kernel-amdgcn.cu
@@ -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"