NVIDIA / NVIDIA/cutlass

Three host reference utility defects: wrong MRE formula, operator precedence in padded-diagonal fill, unseeded per-element RNG in gaussian fill

Open
#3,543 1 comment 0 reactions 0 assignees View on GitHub

Nobody has claimed this yet.

CUTLASS C++
Dominant language
C++
Stars
10.5k
Forks
2.1k
Avg merge
3d 11h
Merged PRs (30d)
7

Description

Description

Three independent defects in the host reference utilities under tools/util/include/cutlass/util/reference/host/.

1. TensorMRE computes |lhs - rhs/(rhs+eps)|, not mean relative error

tensor_compare.h:127:

sum += std::abs(double(lhs_) - double(rhs_) / (double(rhs_) + epsilon));

The division binds only to the second term of the subtraction. The intended mean relative error is |lhs-rhs| / (|rhs| + epsilon). Executed check with a host build of this header: for lhs = 8, rhs = 4 everywhere the function reports MRE = 7.0 instead of 0.5. No in-tree caller today, but it is a public utility whose result is wrong by construction.

2. TensorFillPadDiagonal operator precedence fills the wrong triangle

tensor_fill.h:1003-1011:

if (Layout::kRank == 2 &&
    (fill_mode == cutlass::FillMode::kLower) &&
    (coord[0] >= coord[1]) ||
    ((coord[1] - coord[0]) >= alignment)) {
  view.at(coord) = func();
} else if (... same structure for kUpper ...)

&& binds tighter than ||, so the alignment-pad clause fires in both directions regardless of fill mode: a kLower fill also populates elements far above the diagonal (col - row >= alignment) and vice versa. For a 32x32 matrix with kLower and alignment = 4, 406 elements strictly above the diagonal band receive random values where the fill mode requires zeros. The diagonal itself (row == col) is also filled although padded-diagonal semantics exclude it.

3. Gaussian fill ignores the user seed and constructs an RNG per element

tensor_fill.h:180-224 (same pattern in the complex specialization at :263-267 and quaternion at :344-347):

std::random_device rnd_device;
std::mt19937 bernoulli_rnd(rnd_device());
std::bernoulli_distribution bernoulli_dist(pnz);

is constructed fresh inside operator() for every element when pnz < 1. Two consequences:

  • profiler runs using --dist=gaussian,...,pnzA:<x<1> get different data on every run under the same --seed (the Bernoulli stream comes from random_device, not from the seeded engine that RandomUniformFunc keeps as a member; the Gaussian draw itself uses unseeded std::rand() too),
  • an engine construction per element is wasted even at pnz == 1.
Suggested fixes
  1. Parenthesize: sum += std::abs(double(lhs_) - double(rhs_)) / (std::abs(double(rhs_)) + epsilon);
  2. Parenthesize each mode's conjunction and exclude the diagonal band per mode.
  3. Hoist the Bernoulli engine into the functor member seeded from the same seed as the Gaussian source.

Contributor guide

No contributing guide indexed for this repository

First steps

  1. Read the whole issue, then the project's contributing guide.
  2. Comment on the issue to say you are picking it up — it saves two people doing the same work.
  3. Fork the repository and make your change on a branch.
  4. Open a pull request that references the issue number.

Research direction

Start with tools/util/include/cutlass/util/reference/host/tensor_compare.h:127 and tensor_fill.h:180-224, 263-267, 344-347, and 1003-1011. Build or run the host reference utilities using the reported MRE and 32x32 padded-diagonal cases, then exercise Gaussian filling with a fixed seed. Done means the MRE, fill-mode boundaries, diagonal exclusion, and seeded Gaussian output match the stated semantics without per-element RNG construction.

Written by the indexing model from the issue text.

Assessment

Tech stack
cpp
Domain
testing, tooling
Issue type
Bug
Difficulty
4/5
Estimated time
3-5 days
Activity status
Active
Clarity
Clearly specified
Newbie friendliness
64/100

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.