petrex commented on a change in pull request #4342: Add workgroup size
attribute to AMDGPU functions in codegen
URL: https://github.com/apache/incubator-tvm/pull/4342#discussion_r346909319
##########
File path: src/codegen/llvm/codegen_amdgpu.cc
##########
@@ -36,13 +36,39 @@
namespace tvm {
namespace codegen {
+namespace {
+
+// calls the device api to get the max threads per block
+static inline int DetectROCMmaxThreadsPerBlock() {
+ TVMContext tvm_ctx;
+ tvm_ctx.device_type = kDLROCM;
+ tvm_ctx.device_id = 0;
+ tvm::runtime::DeviceAPI* api = tvm::runtime::DeviceAPI::Get(tvm_ctx, true);
+ if (api != nullptr) {
+ TVMRetValue val;
+ api->GetAttr(tvm_ctx, tvm::runtime::kExist, &val);
+ if (val.operator int() == 1) {
+ tvm::runtime::DeviceAPI::Get(tvm_ctx)->
+ GetAttr(tvm_ctx, tvm::runtime::kMaxThreadsPerBlock, &val);
+ return val.operator int();
+ }
+ }
+ LOG(WARNING) << "Cannot get maximum number of threads for AMD codegen";
+ return 1024;
Review comment:
The optimal value depends on your VGPR/SGPR consumption, and that would vary
among kernels/arch. 256 is a safe bet across multiple workloads/architecture.
Also, TVM AMD backend does not use HIP for kernel construction, LLVM backend
generate the ISA. Are you referring to the Nvidia path when you say annotate
the kernel with `__launch_bounds__()`?
----------------------------------------------------------------
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
For queries about this service, please contact Infrastructure at:
[email protected]
With regards,
Apache Git Services
