[BUG] Weird precision problem in specific tiling description

[vincentccc]

Describe the bug
A clear and concise description of what the bug is.
I got a stochastic result in conv2d kernel with the tiling
cutlass::arch::Sm80,
cutlass::gemm::GemmShape<128, 96, 64>,
cutlass::gemm::GemmShape<32, 96, 64 >,
cutlass::gemm::GemmShape<16, 8, 32>,
and stage3, the workload shows below:

Conv2d Workload input: s8 (1, 256,256,128) weight: s8(128,3,3,128) padding: (1,1,1,1) strides(1,1)

some fact may help:

1. I ran the kernel several times, some of the results are not same, and the mismatch items index are random, but the channel dim of the index always be 0-63
2. I found when I set the workload to be small, the error will not shows up, like (1, 128, 128, 64) (64,3,3,64)
3. other schedule works well, eg: 128x64x64 64x64x64

It is very hard for me to debug what is going on cause I cannot reproduce the error stably.
Is there any hints or method about how to fix/debug this bug or the approximate bug code location.

Steps/Code to reproduce bug
Follow this guide http://matthewrocklin.com/blog/work/2018/02/28/minimal-bug-reports to craft a minimal bug report. This helps us reproduce the issue you're having and resolve the issue more quickly.

Expected behavior
A clear and concise description of what you expected to happen.

Environment details (please complete the following information):

• Environment location: [Bare-metal, Docker, Cloud(specify cloud provider)]

Additional context
Add any other context about the problem here.

[hwu36]

sorry. i think it is just that non power of two tile size is not valid in cutlass 2.x conv.

[vincentccc]

sorry. i think it is just that non power of two tile size is not valid in cutlass 2.x conv.

Thanks for your reply, is that means this tiling will work on cutlass 3.x?
BTW, I am still curious about why it not works, is there any hard code or strong prior?

[thakkarV]

We do support non Po2 tile shapes in 3.X api yes

[hwu36]

BTW, I am still curious about why it not works, is there any hard code or strong prior?

in 2.x, shared memory layout is not non power of 2 friendly.