[GitHub] [incubator-tvm] masahi commented on a change in pull request #4550: [Perf] Add CublasLt extern support for better Igemm performance

[GitBox]

masahi commented on a change in pull request #4550: [Perf] Add CublasLt extern 
support for better Igemm performance
URL: https://github.com/apache/incubator-tvm/pull/4550#discussion_r361422741
 
 

 ##########
 File path: src/runtime/contrib/cublas/cublas.cc
 ##########
 @@ -342,12 +434,33 @@ TVM_REGISTER_GLOBAL("tvm.contrib.cublas.matmul")
     }
 });
 
-TVM_REGISTER_GLOBAL("tvm.contrib.cublas.batch_matmul")
+#if CUDART_VERSION >= 10010
+TVM_REGISTER_GLOBAL("tvm.contrib.cublaslt.matmul")
 .set_body([](TVMArgs args, TVMRetValue* ret) {
     DLTensor* A = args[0];
-    DLTensor* C = args[2];
 
+    CuBlasThreadEntry* entry_ptr = CuBlasThreadEntry::ThreadLocal();
+
+    TryEnableTensorCore(entry_ptr->handle);
+
+    int version;
+    CHECK_CUBLAS_ERROR(cublasGetVersion(entry_ptr->handle, &version));
+    if (TypeMatch(A->dtype, kDLInt, 8) && version >= 10100) {
 
 Review comment:
   version >= 10100 check is no longer necessary by #if guard. I would remove 
this if block and replace it with
   ```
   CHECK(TypeMatch(A->dtype, kDLInt, 8), "Expects dtype to be int8");
   cublasLtHandle_t ltHandle;
   CHECK_CUBLAS_ERROR(cublasLtCreate(&ltHandle));
   CallLtIgemm(args, ret, ltHandle);
   CHECK_CUBLAS_ERROR(cublasLtDestroy(ltHandle));
   ```

----------------------------------------------------------------
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
 
For queries about this service, please contact Infrastructure at:
us...@infra.apache.org


With regards,
Apache Git Services