[QST]why 3090 get different result with 4090 or 3060 when call get<0>(tensor)
Nobody has claimed this yet.
- 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
- 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 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