Small profiler defects: device-ID validation off by one, kTensorNCHW mapped to TensorNHWC, sparse-meta host/device type mismatch
Nobody has claimed this yet.
- Dominant language
- C++
- Stars
- 10.5k
- Forks
- 2.1k
- Avg merge
- 3d 11h
- Merged PRs (30d)
- 7
Description
Description
Three small, independent defects in tools/profiler/src/.
1. --devices validation is off by one
options.cu:83:
} else if (device > num_devices) {
throw std::runtime_error("Bad device ID: " + std::to_string(device));
}
Valid IDs are [0, num_devices - 1], so device == num_devices passes validation and fails later inside cudaGetDeviceProperties() with a generic error. The comparison should be >= (and arguably also reject negative input earlier with a clear message).
2. construct_layout maps kTensorNCHW onto TensorNHWC
device_allocation.cu:258-259:
case library::LayoutTypeID::kTensorNCHW:
return construct_layout_<cutlass::layout::TensorNHWC>(bytes, layout_id, extent, stride);
Copy-paste from the adjacent NHWC case (get_packed_layout at :133 handles kTensorNCHW correctly with layout::TensorNCHW). Any allocation or capacity computation routed through construct_layout for an NCHW tensor would use NHWC strides. Currently latent: the conv profilers hardcode NHWC allocations, so no in-tree caller reaches this case.
3. Host and device sparse-metadata initialization switch on different type IDs
device_allocation.cu:1671-1690 (initialize_random_sparsemeta_host) switches on kS16/kS32, while its device twin (:1636-1654) switches on kU16/kU32. The E metadata tensors of 3.x sparse kernels are created from unsigned type IDs (NumericTypeMap<uint16_t/uint32_t> -> kU16/kU32), so with --initialization-provider=host the host-side switch falls through to default: and an all-zero metadata buffer is uploaded instead of random valid 2:4 indices. The run stays self-consistent (no false pass/fail), but the sparse pattern degenerates to all-zeros and no longer exercises the metadata path.
Suggested fixes
device >= num_devices.- Use
cutlass::layout::TensorNCHWin the NCHW case. - Switch on
kU16/kU32in the host twin (or handle both signed and unsigned IDs).
Contributor guide
No contributing guide indexed for this repository
First steps
- Read the whole issue, then the project's contributing guide.
- Comment on the issue to say you are picking it up — it saves two people doing the same work.
- Fork the repository and make your change on a branch.
- Open a pull request that references the issue number.
Research direction
Start with the cited branches in tools/profiler/src/options.cu and device_allocation.cu, then inspect the adjacent layout and host/device sparse-metadata initialization cases. Verify the device-boundary check, NCHW layout selection, and unsigned metadata type handling against their neighboring cases and profiler call paths. Done means all three defects are corrected without changing the stated behavior of the other cases.
Written by the indexing model from the issue text.
Assessment
- Tech stack
- cpp
- Domain
- performance, tooling
- Issue type
- Bug
- Difficulty
- 3/5
- Estimated time
- 1-2 days
- Activity status
- Active
- Clarity
- Clearly specified
- Newbie friendliness
- 76/100