gemini-code-assist[bot] commented on code in PR #18604:
URL: https://github.com/apache/tvm/pull/18604#discussion_r2642573727
##########
src/runtime/file_utils.cc:
##########
@@ -80,13 +85,15 @@ void FunctionInfo::Save(dmlc::Stream* writer) const {
writer->Write(arg_types);
writer->Write(launch_param_tags);
writer->Write(arg_extra_tags);
+ writer->Write(has_programmatic_dependent_launch);
}
bool FunctionInfo::Load(dmlc::Stream* reader) {
if (!reader->Read(&name)) return false;
if (!reader->Read(&arg_types)) return false;
if (!reader->Read(&launch_param_tags)) return false;
if (!reader->Read(&arg_extra_tags)) return false;
+ if (!reader->Read(&has_programmatic_dependent_launch)) return false;
Review Comment:
This unconditional read of `has_programmatic_dependent_launch` breaks binary
backward compatibility. If you try to load a module that was saved with an
older version of TVM (which wouldn't have this field in its binary metadata),
this `reader->Read` call will fail, causing the entire module loading to fail.
This is inconsistent with the JSON loading logic, which correctly handles
this as an optional field. Please update the binary loading logic to handle
this field optionally to maintain backward compatibility.
##########
src/runtime/meta_data.h:
##########
@@ -60,6 +58,8 @@ struct FunctionInfo {
enum class ArgExtraTags : int { kNone = 0, kTensorMap = 1 };
std::vector<ArgExtraTags> arg_extra_tags;
+ bool has_programmatic_dependent_launch;
Review Comment:
The new member `has_programmatic_dependent_launch` is not initialized. When
loading from older JSON metadata that doesn't contain this field, its value
will be indeterminate. To ensure backward compatibility and predictable
behavior, please provide a default value.
```suggestion
bool has_programmatic_dependent_launch = false;
```
##########
src/runtime/cuda/cuda_module.cc:
##########
@@ -200,9 +202,33 @@ class CUDAWrappedFunc {
}
}
CUstream strm = static_cast<CUstream>(TVMFFIEnvGetStream(kDLCUDA,
device_id));
- CUresult result = cuLaunchKernel(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,
void_args, nullptr);
+
+ CUresult result;
+
+ if (has_programmatic_dependent_launch_) {
+ CUlaunchConfig config;
+ CUlaunchAttribute attribute[1];
Review Comment:
For robustness and to avoid any potential issues with uninitialized struct
members, it's a good practice to zero-initialize `config` and `attribute`. This
ensures that any fields not explicitly set (including any padding or future
fields added to the CUDA structs) are properly initialized to zero.
```suggestion
CUlaunchConfig config{};
CUlaunchAttribute attribute[1]{};
```
--
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.
To unsubscribe, e-mail: [email protected]
For queries about this service, please contact Infrastructure at:
[email protected]
---------------------------------------------------------------------
To unsubscribe, e-mail: [email protected]
For additional commands, e-mail: [email protected]
