This is an automated email from the ASF dual-hosted git repository.
tqchen pushed a commit to branch main
in repository https://gitbox.apache.org/repos/asf/tvm.git
The following commit(s) were added to refs/heads/main by this push:
new ab566679f8 [REFACTOR][Runtime] Use CamelCase for launch dimension
methods (#20682)
ab566679f8 is described below
commit ab566679f8fbd0e76c1ed44945bc174fe146e09f
Author: Tianqi Chen <[email protected]>
AuthorDate: Mon Oct 12 19:09:02 2026 -0400
[REFACTOR][Runtime] Use CamelCase for launch dimension methods (#20682)
Use CamelCase for launch workload dimension methods and conventional
naming for the local dimension-tag constant, while retaining snake_case
for plain configuration-member accessors. Update the ROCm, Metal, and
Vulkan runtime callers consistently without changing launch decoding.
---
src/backend/metal/runtime/metal_module.mm | 6 +++---
src/backend/rocm/runtime/rocm_module.cc | 6 +++---
src/backend/vulkan/runtime/vulkan_wrapped_func.cc | 4 ++--
src/runtime/launch_param.h | 14 +++++++-------
4 files changed, 15 insertions(+), 15 deletions(-)
diff --git a/src/backend/metal/runtime/metal_module.mm
b/src/backend/metal/runtime/metal_module.mm
index 713a34944a..9df3cca9b9 100644
--- a/src/backend/metal/runtime/metal_module.mm
+++ b/src/backend/metal/runtime/metal_module.mm
@@ -250,7 +250,7 @@ class MetalWrappedFunc {
scache_[device_id] = m_->GetPipelineState(device_id, func_name_);
}
ThreadWorkLoad wl = launch_param_config_.Extract(args);
- int blockSize = wl.block_dim(0) * wl.block_dim(1) * wl.block_dim(2);
+ int blockSize = wl.BlockDim(0) * wl.BlockDim(1) * wl.BlockDim(2);
auto maxTotalThreadsPerThreadgroup =
scache_[device_id].maxTotalThreadsPerThreadgroup;
TVM_FFI_ICHECK_LE(blockSize, maxTotalThreadsPerThreadgroup);
// Reuse the pending compute encoder to batch dispatches.
@@ -267,8 +267,8 @@ class MetalWrappedFunc {
atIndex:num_buffer_args_];
}
// launch
- MTLSize dimGrid = MTLSizeMake(wl.grid_dim(0), wl.grid_dim(1),
wl.grid_dim(2));
- MTLSize dimBlock = MTLSizeMake(wl.block_dim(0), wl.block_dim(1),
wl.block_dim(2));
+ MTLSize dimGrid = MTLSizeMake(wl.GridDim(0), wl.GridDim(1),
wl.GridDim(2));
+ MTLSize dimBlock = MTLSizeMake(wl.BlockDim(0), wl.BlockDim(1),
wl.BlockDim(2));
[encoder dispatchThreadgroups:dimGrid threadsPerThreadgroup:dimBlock];
};
}
diff --git a/src/backend/rocm/runtime/rocm_module.cc
b/src/backend/rocm/runtime/rocm_module.cc
index 5654e184e2..df68d86e4d 100644
--- a/src/backend/rocm/runtime/rocm_module.cc
+++ b/src/backend/rocm/runtime/rocm_module.cc
@@ -180,9 +180,9 @@ class ROCMWrappedFunc {
void* config[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, packed_args,
HIP_LAUNCH_PARAM_BUFFER_SIZE,
&packed_nbytes, HIP_LAUNCH_PARAM_END};
// HIP supports only extra_args.
- ROCM_DRIVER_CALL(hipModuleLaunchKernel(fcache_[device_id], wl.grid_dim(0),
wl.grid_dim(1),
- wl.grid_dim(2), wl.block_dim(0),
wl.block_dim(1),
- wl.block_dim(2), wl.dyn_shmem_size,
strm, nullptr,
+ ROCM_DRIVER_CALL(hipModuleLaunchKernel(fcache_[device_id], wl.GridDim(0),
wl.GridDim(1),
+ wl.GridDim(2), wl.BlockDim(0),
wl.BlockDim(1),
+ wl.BlockDim(2), wl.dyn_shmem_size,
strm, nullptr,
reinterpret_cast<void**>(&config)));
}
diff --git a/src/backend/vulkan/runtime/vulkan_wrapped_func.cc
b/src/backend/vulkan/runtime/vulkan_wrapped_func.cc
index 1e777afa75..ed100379e1 100644
--- a/src/backend/vulkan/runtime/vulkan_wrapped_func.cc
+++ b/src/backend/vulkan/runtime/vulkan_wrapped_func.cc
@@ -89,7 +89,7 @@ void VulkanWrappedFunc::operator()(ffi::PackedArgs args,
ffi::Any* rv,
pack_args);
}
- vkCmdDispatch(state->cmd_buffer_, wl.grid_dim(0), wl.grid_dim(1),
wl.grid_dim(2));
+ vkCmdDispatch(state->cmd_buffer_, wl.GridDim(0), wl.GridDim(1),
wl.GridDim(2));
VkMemoryBarrier barrier_info;
barrier_info.sType = VK_STRUCTURE_TYPE_MEMORY_BARRIER;
barrier_info.pNext = nullptr;
@@ -156,7 +156,7 @@ void VulkanWrappedFunc::operator()(ffi::PackedArgs args,
ffi::Any* rv,
pack_args_storage.data());
}
- vkCmdDispatch(state->cmd_buffer_, wl.grid_dim(0), wl.grid_dim(1),
wl.grid_dim(2));
+ vkCmdDispatch(state->cmd_buffer_, wl.GridDim(0), wl.GridDim(1),
wl.GridDim(2));
VkMemoryBarrier barrier_info;
barrier_info.sType = VK_STRUCTURE_TYPE_MEMORY_BARRIER;
barrier_info.pNext = nullptr;
diff --git a/src/runtime/launch_param.h b/src/runtime/launch_param.h
index 1d3b51b0fc..41b9cbc0ac 100644
--- a/src/runtime/launch_param.h
+++ b/src/runtime/launch_param.h
@@ -47,22 +47,22 @@ struct ThreadWorkLoad {
* \param i The block dimension.
* \return i-th block dim
*/
- inline size_t block_dim(size_t i) const { return work_size[i + 3]; }
+ inline size_t BlockDim(size_t i) const { return work_size[i + 3]; }
/*!
* \param i The grid dimension.
* \return i-th grid dim
*/
- inline size_t grid_dim(size_t i) const { return work_size[i]; }
+ inline size_t GridDim(size_t i) const { return work_size[i]; }
/*!
* \param i The cluster dimension.
* \return i-th cluster dim
*/
- inline size_t cluster_dim(size_t i) const { return work_size[i + 6]; }
+ inline size_t ClusterDim(size_t i) const { return work_size[i + 6]; }
/*!
* \param i The preferred cluster dimension.
* \return i-th preferred cluster dim
*/
- inline size_t preferred_cluster_dim(size_t i) const { return work_size[i +
9]; }
+ inline size_t PreferredClusterDim(size_t i) const { return work_size[i + 9];
}
};
/*! \brief Launch parameters configuration */
@@ -85,7 +85,7 @@ class LaunchParamConfig {
} else if (tag == launch_param::kUseRequiredBlockDimension) {
use_required_block_dimension_ = true;
} else {
- static constexpr const char* dimension_tags[] = {"blockIdx.x",
+ static constexpr const char* kDimensionTags[] = {"blockIdx.x",
"blockIdx.y",
"blockIdx.z",
"threadIdx.x",
@@ -98,7 +98,7 @@ class LaunchParamConfig {
"preferredClusterCtaIdx.y",
"preferredClusterCtaIdx.z"};
uint32_t index = 0;
- while (index < 12 && tag != dimension_tags[index]) {
+ while (index < 12 && tag != kDimensionTags[index]) {
++index;
}
TVM_FFI_ICHECK_LT(index, 12) << "Unsupported launch parameter tag: "
<< tag;
@@ -137,7 +137,7 @@ class LaunchParamConfig {
// return the work dim
size_t work_dim() const { return work_dim_; }
- bool use_programtic_dependent_launch() const { return
use_programmatic_dependent_launch_; }
+ bool use_programmatic_dependent_launch() const { return
use_programmatic_dependent_launch_; }
bool use_cooperative_launch() const { return use_cooperative_launch_; }
