@@ -57,10 +57,6 @@ namespace kernels
5757{
5858namespace cutlass_kernels
5959{
60-
61- namespace oss
62- {
63-
6460template <typename T, typename arch, typename ThreadblockShape, typename WarpShape, int Stages>
6561void genericInt8GemmKernelLauncher (int8_t const * A, int8_t const * B, tk::QuantMode quantOption, float const * alphaCol,
6662 float const * alphaRow, T* C, int m, int n, int k, tkc::CutlassGemmConfig gemmConfig, char * workspace,
@@ -304,7 +300,6 @@ void dispatchGemmToCutlass(int8_t const* A, int8_t const* B, tk::QuantMode quant
304300 break ;
305301 }
306302}
307- } // namespace oss
308303
309304template <typename T>
310305CutlassInt8GemmRunner<T>::CutlassInt8GemmRunner()
@@ -330,18 +325,18 @@ void CutlassInt8GemmRunner<T>::dispatchToArch(int8_t const* A, int8_t const* B,
330325 TLLM_LOG_DEBUG (__PRETTY_FUNCTION__);
331326 if (mSm >= 72 && mSm < 75 )
332327 {
333- oss:: dispatchGemmToCutlass<T, cutlass::arch::Sm72>(A, B, quantOption, alphaCol, alphaRow, C, m, n, k,
334- workspacePtr, workspaceBytes, gemmConfig, stream, occupancy);
328+ dispatchGemmToCutlass<T, cutlass::arch::Sm72>(A, B, quantOption, alphaCol, alphaRow, C, m, n, k, workspacePtr ,
329+ workspaceBytes, gemmConfig, stream, occupancy);
335330 }
336331 else if (mSm >= 75 && mSm < 80 )
337332 {
338- oss:: dispatchGemmToCutlass<T, cutlass::arch::Sm75>(A, B, quantOption, alphaCol, alphaRow, C, m, n, k,
339- workspacePtr, workspaceBytes, gemmConfig, stream, occupancy);
333+ dispatchGemmToCutlass<T, cutlass::arch::Sm75>(A, B, quantOption, alphaCol, alphaRow, C, m, n, k, workspacePtr ,
334+ workspaceBytes, gemmConfig, stream, occupancy);
340335 }
341336 else if (mSm >= 80 && mSm <= 90 || mSm >= 120 )
342337 {
343- oss:: dispatchGemmToCutlass<T, cutlass::arch::Sm80>(A, B, quantOption, alphaCol, alphaRow, C, m, n, k,
344- workspacePtr, workspaceBytes, gemmConfig, stream, occupancy);
338+ dispatchGemmToCutlass<T, cutlass::arch::Sm80>(A, B, quantOption, alphaCol, alphaRow, C, m, n, k, workspacePtr ,
339+ workspaceBytes, gemmConfig, stream, occupancy);
345340 }
346341 else
347342 {
0 commit comments