[QST] Does examples/36 available on sm 70?

[umiswing]

Hello! I write the following kernel config to support gather-gemm-scatter fusion on sm 70. But the result is not correct. I debug for a while and can't fix it.
Does cutlass support this fusion on sm 70?

struct cutlass_tensorop_h884gemm_256x128_32x2_nn_align8 {                                                                                                
   using Gemm =                                                                                                                                           
      cutlass::gemm::device::GemmUniversal<                                                                                                                
       cutlass::half_t,                                                                                                                                   
       cutlass::layout::RowMajor,                                                                                                                         
       cutlass::half_t,                                                                                                                                   
       cutlass::layout::RowMajor,                                                                                                                         
       cutlass::half_t,                                                                                                                                   
       cutlass::layout::RowMajor,                                                                                                                         
       cutlass::half_t,                                                                                                                                   
       cutlass::arch::OpClassTensorOp,                                                                                                                    
       cutlass::arch::Sm70,                                                                                                                               
       cutlass::gemm::GemmShape<256, 128, 32>,                                                                                                            
       cutlass::gemm::GemmShape<64, 64, 32>,                                                                                                              
       cutlass::gemm::GemmShape<8, 8, 4>,                                                                                                                                                                                                                                                                 
       cutlass::epilogue::thread::LinearCombination<                                                                                                        
         cutlass::half_t,                                                                                                                                   
         8,                                                                                                                                                 
         cutlass::half_t,                                                                                                                                   
         cutlass::half_t                                                                                                                                    
       >,                                                                                                                                                        
       cutlass::gemm::threadblock::GemmIdentityThreadblockSwizzle<8>,                                                                                     
       2,                                                                                                                                                 
       8,                                                                                                                                                 
       8,                                                                                                                                                 
       cutlass::arch::OpMultiplyAdd,                                                                                                                      
       cutlass::ComplexTransform::kNone,                                                                                                                  
       cutlass::ComplexTransform::kNone,                                                                                                                  
       true, // gather a                                                                                                                                  
       false, // gather b                                                                                                                                 
       true // scatter d                                                                                                                                  
>;                     

cutlass version: v2.11.0
hardware: nvidia v100

[hwu36]

sm70 is not tested and it is a bit too old for us to add new things to it.

Can you tell the problem is from gather a or scatter d?

[umiswing]

sm70 is not tested and it is a bit too old for us to add new things to it.

Can you tell the problem is from gather a or scatter d?

I think it is from scatter d. It works well when I remove scatter d fusion.
Do you have any plan to fix it so that I can keep the scatter d fusion on sm 70?

[mnicely]

@umiswing we are in the middle of Hopper enablement. I'll put this on the backlog for now.
Can you provide a complete reproducer, with a flag that triggers the issue?

[umiswing]

@mnicely I will provide it later.

[hwu36]

Do you have any plan to fix it so that I can keep the scatter d fusion on sm 70?

We don't have bw to do it ourself, but we can help you to fix it if you want.

I think it is from scatter d. It works well when I remove scatter d fusion.

in this line https://github.com/NVIDIA/cutlass/blob/master/include/cutlass/epilogue/threadblock/predicated_tile_iterator.h#L333 , we map the row number to the one wanted in the index file (row_offset + thread_start_row_ -> indices_[row_offset + thread_start_row_]). You can dump here to see what is going wrong. Volta uses a different thread map from turing/ampere. This could be the source of the problem.

[umiswing]

@hwu36 Thank you! I might fix it myself in the future. I will just remove this fusion on sm 70 now. :)