New tensile client

library/src/blas3/Tensile/gemm.cpp

@@ -163,7 +163,6 @@ namespace
             hipMemcpy(&beta_h, beta, sizeof(T), hipMemcpyDeviceToHost);
         }
 
-        TensileHostCall<T>           hostCall;
         RocblasContractionProblem<T> problem(ContractionProblemType::GEMM,
                                              trans_a,
                                              trans_b,
@@ -179,7 +178,7 @@ namespace
                                              C,
                                              ld_c);
 
-        return callTensileContraction(&problem, handle->host);
+        return handle->host->runContractionProblem(problem);
 
 #else
 

library/src/blas3/Tensile/gemm_strided_batched.cpp

@@ -197,7 +197,6 @@ namespace
             hipMemcpy(&beta_h, beta, sizeof(T), hipMemcpyDeviceToHost);
         }
 
-        TensileHostCall<T>           hostCall;
         RocblasContractionProblem<T> problem(ContractionProblemType::GEMMStridedBatch,
                                              trans_a,
                                              trans_b,
@@ -217,7 +216,7 @@ namespace
                                              stride_c,
                                              b_c);
 
-        return callTensileContraction(&problem, handle->host);
+        return handle->host->runContractionProblem(problem);
 
 #else
         rocblas_status validArgs = validateArgs(handle,

library/src/handle.cpp

@@ -1,6 +1,9 @@
 /* ************************************************************************
  * Copyright 2016-2019 Advanced Micro Devices, Inc.
  * ************************************************************************ */
+#if BUILD_WITH_TENSILE
+#include "Tensile.h"
+#endif
 #include "handle.h"
 #include <cstdio>
 #include <cstdlib>
@@ -10,10 +13,12 @@
  ******************************************************************************/
 _rocblas_handle::_rocblas_handle()
 {
+#if BUILD_WITH_TENSILE
+    static int dummy = (tensileInitialize(), 0);
 #ifdef USE_TENSILE_HOST
-    host = createTensileHost();
-    if(!host)
-        throw rocblas_status_internal_error;
+    static TensileHost* hostImpl = createTensileHost();
+    host                         = hostImpl;
+#endif
 #endif
 
     // default device is active device
@@ -69,10 +74,6 @@ _rocblas_handle::~_rocblas_handle()
     }
     if(device_memory)
         (hipFree)(device_memory);
-
-#ifdef USE_TENSILE_HOST
-    delete host;
-#endif
 }
 
 /*******************************************************************************

library/src/include/tensile_host.hpp

@@ -12,42 +12,41 @@ enum ContractionProblemType
 };
 
 template <typename T>
-class RocblasContractionProblem
+struct RocblasContractionProblem
 {
-public:
     ContractionProblemType problem_type;
     rocblas_operation      trans_a;
     rocblas_operation      trans_b;
-    unsigned long          m;
-    unsigned long          n;
-    unsigned long          k;
+    size_t                 m;
+    size_t                 n;
+    size_t                 k;
     const T                alpha;
     const T*               A;
-    const unsigned long    ld_a;
-    unsigned long          stride_a;
+    size_t                 ld_a;
+    size_t                 stride_a;
     const T*               B;
-    unsigned long          ld_b;
-    unsigned long          stride_b;
+    size_t                 ld_b;
+    size_t                 stride_b;
     const T                beta;
     T*                     C;
-    unsigned long          ld_c;
-    unsigned long          stride_c;
-    unsigned long          batch_size;
+    size_t                 ld_c;
+    size_t                 stride_c;
+    size_t                 batch_size;
 
     RocblasContractionProblem(ContractionProblemType problem_type,
                               rocblas_operation      trans_a,
                               rocblas_operation      trans_b,
-                              unsigned long          m,
-                              unsigned long          n,
-                              unsigned long          k,
+                              size_t                 m,
+                              size_t                 n,
+                              size_t                 k,
                               const T                alpha,
                               const T*               A,
-                              unsigned long          ld_a,
+                              size_t                 ld_a,
                               const T*               B,
-                              unsigned long          ld_b,
+                              size_t                 ld_b,
                               const T                beta,
                               T*                     C,
-                              unsigned long          ld_c)
+                              size_t                 ld_c)
         : problem_type(problem_type)
         , trans_a(trans_a)
         , trans_b(trans_b)
@@ -72,21 +71,21 @@ class RocblasContractionProblem
     RocblasContractionProblem(ContractionProblemType problem_type,
                               rocblas_operation      trans_a,
                               rocblas_operation      trans_b,
-                              unsigned long          m,
-                              unsigned long          n,
-                              unsigned long          k,
+                              size_t                 m,
+                              size_t                 n,
+                              size_t                 k,
                               const T                alpha,
                               const T*               A,
-                              unsigned long          ld_a,
-                              unsigned long          stride_a,
+                              size_t                 ld_a,
+                              size_t                 stride_a,
                               const T*               B,
-                              unsigned long          ld_b,
-                              unsigned long          stride_b,
+                              size_t                 ld_b,
+                              size_t                 stride_b,
                               const T                beta,
                               T*                     C,
-                              unsigned long          ld_c,
-                              unsigned long          stride_c,
-                              unsigned long          batch_size)
+                              size_t                 ld_c,
+                              size_t                 stride_c,
+                              size_t                 batch_size)
         : problem_type(problem_type)
         , trans_a(trans_a)
         , trans_b(trans_b)
@@ -109,24 +108,17 @@ class RocblasContractionProblem
     }
 };
 
-class TensileHost
+struct TensileHost
 {
-public:
-    virtual void initializeHost(const char*) {}
-};
+    template <typename T>
+    rocblas_status runContractionProblem(const RocblasContractionProblem<T>& problem);
 
-template <typename T>
-class TensileHostCall
-{
-public:
-    rocblas_status runContractionProblem(RocblasContractionProblem<T>* problem, TensileHost* host);
+protected:
+    TensileHost() = default; // Prevent instantiating this class except as base class
 };
 
 TensileHost* createTensileHost();
 
-template <typename T>
-rocblas_status callTensileContraction(RocblasContractionProblem<T>* problem, TensileHost* host);
-
 #endif
 
 #endif // __TENSILE_HOST_HPP__

library/src/rocblas_auxiliary.cpp

@@ -2,9 +2,6 @@
  * Copyright 2016-2019 Advanced Micro Devices, Inc.
  *
  * ************************************************************************ */
-#if BUILD_WITH_TENSILE
-#include "Tensile.h"
-#endif
 #include "handle.h"
 #include "logging.h"
 #include "rocblas-auxiliary.h"
@@ -71,20 +68,10 @@ extern "C" rocblas_status rocblas_create_handle(rocblas_handle* handle)
     // allocate on heap
     try
     {
-#if BUILD_WITH_TENSILE
-        static int dummy = (tensileInitialize(), 0);
-#endif
         *handle = new _rocblas_handle();
 
         if((*handle)->layer_mode & rocblas_layer_mode_log_trace)
             log_trace(*handle, "rocblas_create_handle");
-
-#ifdef USE_TENSILE_HOST
-        const char* lib_path = getenv("ROCBLAS_TENSILE_LIBPATH");
-        if(!lib_path)
-            lib_path = "/opt/rocm/";  // TODO: Set default path
-        (*handle)->host->initializeHost(lib_path);
-#endif
     }
     catch(...)
     {