Commit 0ffb12ca authored by Yaxun (Sam) Liu's avatar Yaxun (Sam) Liu
Browse files

[HIP] Mark kernels with uniform-work-group-size=true 

Differential Revision: https://reviews.llvm.org/D76076
parent 2e77f0cf

clang/lib/CodeGen/TargetInfo.cpp

+4 −0
Original line number Diff line number Diff line
@@ -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;

clang/test/CodeGenCUDA/amdgpu-kernel-attrs.cu

+1 −1
Original line number Diff line number Diff line
@@ -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"