AMDGPU: Use stricter bounds for workitem builtins

[nhaehnle]

LGTM

[arsenm]

Use 1024 for OpenCL

[t-tye]

include/clang/Basic/TargetInfo.h

1060

Is this specifically tied to OpenCL or is it the target's maximum supported work-group size? There are other languages besides OpenCL that have the notion of work-groups, and would seem better to generalize this.

lib/Basic/Targets/AMDGPU.h

73–74

Probably want to update this comment as all runtimes are moving to support the full capabilities of the hardware which is 1024 for all current AMDGCN targets.

include/clang/Basic/TargetInfo.h

public:
   /// \brief Get const supported OpenCL extensions and optional core features.
   const OpenCLOptions &getSupportedOpenCLOpts() const {
       return getTargetOpts().SupportedOpenCLOptions;
   }
 
   /// \brief Get address space for OpenCL type.
   virtual LangAS getOpenCLTypeAddrSpace(const Type *T) const;
 
+  /// \returns Maximum device supported OpenCL workgroup size.
+  virtual unsigned getOpenCLMaxWorkGroupSize(unsigned Dim) const {
+    return 0;
+  }
+
   /// \returns Target specific vtbl ptr address space.
   virtual unsigned getVtblPtrAddressSpace() const {
     return 0;
   }
 
   /// \returns If a target requires an address within a target specific address
   /// space \p AddressSpace to be converted in order to be used, then return the
   /// corresponding target specific DWARF address space.

lib/Basic/Targets/AMDGPU.h

enum GPUKind {
     GK_GFX9
   } GPU;
 
   bool hasFP64 : 1;
   bool hasFMAF : 1;
   bool hasLDEXPF : 1;
   const AddrSpace AS;
 
+  // The maximum supported group size is 1024, but some runtimes currently only
+  // support 256.
+  unsigned MaxWorkGroupSize = 1024;
+
   static bool hasFullSpeedFMAF32(StringRef GPUName) {
     return parseAMDGCNName(GPUName) >= GK_GFX9;
   }
 
   static bool isAMDGCN(const llvm::Triple &TT) {
     return TT.getArch() == llvm::Triple::amdgcn;
   }
 
default:
       return TargetInfo::getOpenCLTypeAddrSpace(T);
     }
   }
 
   llvm::Optional<LangAS> getConstantAddressSpace() const override {
     return getLangASFromTargetAS(AS.Constant);
   }
 
+  unsigned getOpenCLMaxWorkGroupSize(unsigned Dim) const override {
+    return MaxWorkGroupSize;
+  }
+
   /// \returns Target specific vtbl ptr address space.
   unsigned getVtblPtrAddressSpace() const override { return AS.Constant; }
 
   /// \returns If a target requires an address within a target specific address
   /// space \p AddressSpace to be converted in order to be used, then return the
   /// corresponding target specific DWARF address space.
   ///
   /// \returns Otherwise return None and no conversion will be emitted in the

lib/Basic/Targets/AMDGPU.cpp

AMDGPUTargetInfo::AMDGPUTargetInfo(const llvm::Triple &Triple,
   }
 
   MaxAtomicPromoteWidth = MaxAtomicInlineWidth = 64;
 }
 
 void AMDGPUTargetInfo::adjust(LangOptions &Opts) {
   TargetInfo::adjust(Opts);
   setAddressSpaceMap(Opts.OpenCL || !isAMDGCN(getTriple()));
-}
+
 
 ArrayRef<Builtin::Info> AMDGPUTargetInfo::getTargetBuiltins() const {
   return llvm::makeArrayRef(BuiltinInfo, clang::AMDGPU::LastTSBuiltin -
                                              Builtin::FirstTSBuiltin);
 }
 
 void AMDGPUTargetInfo::getTargetDefines(const LangOptions &Opts,
                                         MacroBuilder &Builder) const {

lib/CodeGen/CGBuiltin.cpp

case AMDGPU::BI__builtin_amdgcn_read_exec_hi: {
     CallInst *CI = cast<CallInst>(
       EmitSpecialRegisterBuiltin(*this, E, Int32Ty, Int32Ty, true, RegName));
     CI->setConvergent();
     return CI;
   }
 
   // amdgcn workitem
   case AMDGPU::BI__builtin_amdgcn_workitem_id_x:
-    return emitRangedBuiltin(*this, Intrinsic::amdgcn_workitem_id_x, 0, 1024);
+    return emitRangedBuiltin(*this, Intrinsic::amdgcn_workitem_id_x, 0,
+      getContext().getTargetInfo().getOpenCLMaxWorkGroupSize(0));
   case AMDGPU::BI__builtin_amdgcn_workitem_id_y:
-    return emitRangedBuiltin(*this, Intrinsic::amdgcn_workitem_id_y, 0, 1024);
+    return emitRangedBuiltin(*this, Intrinsic::amdgcn_workitem_id_y, 0,
+      getContext().getTargetInfo().getOpenCLMaxWorkGroupSize(1));
   case AMDGPU::BI__builtin_amdgcn_workitem_id_z:
-    return emitRangedBuiltin(*this, Intrinsic::amdgcn_workitem_id_z, 0, 1024);
+    return emitRangedBuiltin(*this, Intrinsic::amdgcn_workitem_id_z, 0,
+      getContext().getTargetInfo().getOpenCLMaxWorkGroupSize(2));
 
   // r600 intrinsics
   case AMDGPU::BI__builtin_r600_recipsqrt_ieee:
   case AMDGPU::BI__builtin_r600_recipsqrt_ieeef:
     return emitUnaryBuiltin(*this, E, Intrinsic::r600_recipsqrt_ieee);
   case AMDGPU::BI__builtin_r600_read_tidig_x:
-    return emitRangedBuiltin(*this, Intrinsic::r600_read_tidig_x, 0, 1024);
+    return emitRangedBuiltin(*this, Intrinsic::r600_read_tidig_x, 0,
+      getContext().getTargetInfo().getOpenCLMaxWorkGroupSize(0));
   case AMDGPU::BI__builtin_r600_read_tidig_y:
-    return emitRangedBuiltin(*this, Intrinsic::r600_read_tidig_y, 0, 1024);
+    return emitRangedBuiltin(*this, Intrinsic::r600_read_tidig_y, 0,
+      getContext().getTargetInfo().getOpenCLMaxWorkGroupSize(1));
   case AMDGPU::BI__builtin_r600_read_tidig_z:
-    return emitRangedBuiltin(*this, Intrinsic::r600_read_tidig_z, 0, 1024);
+    return emitRangedBuiltin(*this, Intrinsic::r600_read_tidig_z, 0,
+      getContext().getTargetInfo().getOpenCLMaxWorkGroupSize(2));
   default:
     return nullptr;
   }
 }
 
 /// Handle a SystemZ function in which the final argument is a pointer
 /// to an int that receives the post-instruction CC value.  At the LLVM level
 /// this is represented as a function that returns a {result, cc} pair.

test/CodeGenOpenCL/builtins-amdgcn.cl

 
 // CHECK-LABEL: @test_s_getpc
 // CHECK: call i64 @llvm.amdgcn.s.getpc()
 void test_s_getpc(global ulong* out)
 {
   *out = __builtin_amdgcn_s_getpc();
 }
 
-// CHECK-DAG: [[WI_RANGE]] = !{i32 0, i32 1024}
+// CHECK-DAG: [[WI_RANGE]] = !{i32 0, i32 256}
 // CHECK-DAG: attributes #[[NOUNWIND_READONLY:[0-9]+]] = { nounwind readonly }
 // CHECK-DAG: attributes #[[READ_EXEC_ATTRS]] = { convergent }
 // CHECK-DAG: ![[EXEC]] = !{!"exec"}
 // CHECK-DAG: ![[EXEC_LO]] = !{!"exec_lo"}
 // CHECK-DAG: ![[EXEC_HI]] = !{!"exec_hi"}

test/CodeGenOpenCL/builtins-r600.cl

void test_get_local_id(int d, global int *out)
   switch (d) {
   case 0: *out = __builtin_r600_read_tidig_x(); break;
   case 1: *out = __builtin_r600_read_tidig_y(); break;
   case 2: *out = __builtin_r600_read_tidig_z(); break;
   default: *out = 0;
   }
 }
 
-// CHECK-DAG: [[WI_RANGE]] = !{i32 0, i32 1024}
+// CHECK-DAG: [[WI_RANGE]] = !{i32 0, i32 256}