diff --git a/csrc/sliding_tile_attention/st_attn/st_attn_h100.cu b/csrc/sliding_tile_attention/st_attn/st_attn_h100.cu index aaad99678..6cd633ee7 100644 --- a/csrc/sliding_tile_attention/st_attn/st_attn_h100.cu +++ b/csrc/sliding_tile_attention/st_attn/st_attn_h100.cu @@ -526,7 +526,15 @@ sta_forward(torch::Tensor q, torch::Tensor k, torch::Tensor v, torch::Tensor o, mem_size ); fwd_attend_ker<128, false, false, 2, 0, 5, 5, 6, 10><<>>(g); - } else { + } else if (kernel_t_size == 5 && kernel_h_size == 6 && kernel_w_size == 1){ + cudaFuncSetAttribute( + fwd_attend_ker<128, false, false, 2, 3, 0, 5, 6, 10>, + cudaFuncAttributeMaxDynamicSharedMemorySize, + mem_size + ); + fwd_attend_ker<128, false, false, 2, 3, 0, 5, 6, 10><<>>(g); + } + else { // print error std::cout << "Invalid kernel size" << std::endl; //print kernel size