NVIDIA / NVIDIA/cutlass

[QST]why 3090 get different result with 4090 or 3060 when call get<0>(tensor)

Open
#1,898 4 comments 0 reactions 0 assignees View on GitHub

Nobody has claimed this yet.

? - Needs Triage inactive-30d inactive-90d question
Dominant language
C++
Stars
10.5k
Forks
2.1k
Avg merge
3d 11h
Merged PRs (30d)
7

Description

I encountered a strange phenomenon.

for code below:

  auto l = make_identity_tensor(make_shape(_32{}, 64));

  using T_V = cutlass::AlignedArray<cutlass::half_t, 8>;
  using v_g2r_copy_op = UniversalCopy<T_V>;
  using v_g2r_traits = Copy_Traits<v_g2r_copy_op>;
  using v_g2r_copy_atom = Copy_Atom<v_g2r_traits, T_V>;

  using v_g2r_copy_tile = decltype(make_tiled_copy(
      v_g2r_copy_atom{}, Layout<Shape<_16, _8>, Stride<_8, _1>>{},
      Layout<Shape<_1, _8>>{}));

  v_g2r_copy_tile v_copy_tile;
  auto v_g2r_thr_copy = v_copy_tile.get_thread_slice(10);

  auto thr_v_c = v_g2r_thr_copy.partition_S(l);
  printf("get<0>(thr_v_c(0, 0, 0)) : %d\n", int(get<0>(thr_v_c(0, 0, 0))));
  printf("get<1>(thr_v_c(0, 0, 0)): %d\n", int(get<1>(thr_v_c(0, 0, 0))));
  bool can = get<0>(thr_v_c(0, 0, 0)) < 32 && get<1>(thr_v_c(0, 0, 0)) < 64;
  printf("can: %d\n", can);

when cutlass v3.0.0 , run code in cpu, GPU3060, 3090, 4090, it all get right result:

get<0>(thr_v_c(0, 0, 0)) : 1
get<1>(thr_v_c(0, 0, 0)): 16
can: 1

but when cutlass v3.5.0 , run code in cpu, GPU3060, 3090, 4090, only GPU3090 get wrong result:

get<0>(thr_v_c(0, 0, i)) : 0
get<1>(thr_v_c(0, 0, i)): 0
can: 0

Compilation directives: nvcc --std=c++17 -arch=sm_86 -w -I ./include/ test.cu -o a

Can you reproduce it?

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 by compiling and running the provided test.cu with the stated nvcc command on CUTLASS v3.0.0 and v3.5.0, comparing CPU, RTX 3060, RTX 3090, and RTX 4090 results. Trace make_identity_tensor, partition_S, and get<> to determine whether the RTX 3090 discrepancy is reproducible and identify the expected result for the reported case.

Written by the indexing model from the issue text.

Assessment

Tech stack
cpp
Domain
hpc
Issue type
Bug
Difficulty
4/5
Estimated time
3-5 days
Activity status
Stale
Clarity
Mostly clear
Newbie friendliness
25/100

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.