Code clean up and documentation.

155 files changed, 635 insertions, 805 deletions

diff --git a/src/ti/linalg/blasblisacc/src/facade.c b/src/ti/linalg/blasblisacc/src/facade.c
index 547d983..8f896ea 100644
--- a/src/ti/linalg/blasblisacc/src/facade.c
+++ b/src/ti/linalg/blasblisacc/src/facade.c
@@ -24,26 +24,16 @@
  *   CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE)
  *   ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF
  *   THE POSSIBILITY OF SUCH DAMAGE.
- *****************************************************************************/
-  
+ *****************************************************************************/ 
 #include "../../cblas/include/cblas.h"
 #include "../../ticblas/ticblas.h"
 
-#ifdef TI_CBLAS_DEBUG
-#include "stdio.h"  
-
-extern char *pool_mk_mem_L1;
-extern char *pool_kn_mem_L1;
-extern char *pool_mn_mem_L1;
-extern char *pool_mk_mem_L2;
-extern char *pool_kn_mem_L2;
-extern char *pool_mn_mem_L2;
-extern char *pool_mk_mem_L3;
-extern char *pool_kn_mem_L3;
-extern char *pool_mn_mem_L3;
-#endif
-
-extern int bli_l3_mem_config(void *msmc_buf, size_t msmc_buf_size, void *ddr_buf, size_t ddr_buf_size, size_t *l1D_SRAM_size_orig, size_t *l2_SRAM_size_orig);
+/*==============================================================================
+ * This file contains functions of the DSP OpenCL layer of ARM+DSP CBLAS library.
+ *============================================================================*/
+
+extern int bli_l3_mem_config(void *msmc_buf, size_t msmc_buf_size, void *ddr_buf, size_t ddr_buf_size, 
+                             size_t *l1D_SRAM_size_orig, size_t *l2_SRAM_size_orig);
 extern int bli_l3_mem_reconfig(size_t l1D_SRAM_size_orig, size_t l2_SRAM_size_orig);
 
 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_acc.h b/src/ti/linalg/blasblisacc/src/ti_cblas_acc.h
index 77d8a47..907391b 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_acc.h
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_acc.h
@@ -1,4 +1,3 @@
-
 /******************************************************************************
  * Copyright (c) 2013-2015, Texas Instruments Incorporated - http://www.ti.com/
  *   All rights reserved.
@@ -26,10 +25,8 @@
  *   ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF
  *   THE POSSIBILITY OF SUCH DAMAGE.
  *****************************************************************************/
-
-
-#ifndef TI_CBLAS_H
-#define TI_CBLAS_H
+#ifndef TI_CBLAS_ACC_H
+#define TI_CBLAS_ACC_H
 
 #ifdef __cplusplus
 #include <cstdlib>
@@ -125,19 +122,6 @@ extern void ti_cblas_init(void);
 extern err_t bli_finalize();
 extern err_t bli_init();
 
-
-#ifdef _cplusplus
-extern "C" { int ti_cblas_finalize(); }
-else
-        int ti_cblas_finalize(void);
-#endif
-
-#ifdef _cplusplus
-extern "C" { int ti_blis_init(); }
-else
-        int ti_blis_init(void);
-#endif
-
 void ti_cblas_auto_finalize(void);
 
 void ti_cblas_mem_free(void *ptr);
@@ -147,183 +131,17 @@ extern pthread_cond_t   CV;
 extern pthread_mutex_t  MUTEX;
 
 
-#ifdef __cplusplus
-extern Kernel*                  ti_cblas_get_kernel(int idx, const char *fname);
-int                                     ti_cblas_delete_kernel(Kernel* K);
-#if 0
-extern Context                          ti_cblas_ocl_context;
-extern std::vector<Device>      ti_cblas_ocl_devices;
-extern CommandQueue                     ti_cblas_ocl_Q;
-extern Program::Binaries        ti_cblas_ocl_binary;
-extern Program                          ti_cblas_ocl_program;
-extern Kernel*                          ti_cblas_ocl_kernels[];
-#else
-extern Context*                         ti_cblas_ocl_context;
-extern std::vector<Device>*     ti_cblas_ocl_devices;
-extern CommandQueue*            ti_cblas_ocl_Q;
-extern Program::Binaries*       ti_cblas_ocl_binary;
-extern Program*                         ti_cblas_ocl_program;
-#endif
-#else
-extern cl_kernel                        ti_cblas_get_kernel(int idx, const char *fname);
-int                                             ti_cblas_delete_kernel(cl_kernel K);
-extern cl_context                       ti_cblas_ocl_context;
-extern cl_command_queue         ti_cblas_ocl_Q;
-extern cl_program               ti_cblas_ocl_program;
-extern cl_kernel                ti_cblas_ocl_kernels[];
-#endif
+extern Kernel*              ti_cblas_get_kernel(const char *fname);
+int                         ti_cblas_delete_kernel(Kernel* K);
+extern Context*             ti_cblas_ocl_context;
+extern std::vector<Device>* ti_cblas_ocl_devices;
+extern CommandQueue*        ti_cblas_ocl_Q;
+extern Program::Binaries*   ti_cblas_ocl_binary;
+extern Program*             ti_cblas_ocl_program;
 
 extern int ti_cblas_init_done;
-extern int ti_cblas_kernel_valid[];
 extern int ti_cblas_offload;
 
-#define TI_CBLAS_CBLAS_CAXPY_IDX 0
-#define TI_CBLAS_CBLAS_CCOPY_IDX 1
-#define TI_CBLAS_CBLAS_CDOTC_SUB_IDX 2
-#define TI_CBLAS_CBLAS_CDOTU_SUB_IDX 3
-#define TI_CBLAS_CBLAS_CGBMV_IDX 4
-#define TI_CBLAS_CBLAS_CGEMM_IDX 5
-#define TI_CBLAS_CBLAS_CGEMV_IDX 6
-#define TI_CBLAS_CBLAS_CGERC_IDX 7
-#define TI_CBLAS_CBLAS_CGERU_IDX 8
-#define TI_CBLAS_CBLAS_CHBMV_IDX 9
-#define TI_CBLAS_CBLAS_CHEMM_IDX 10
-#define TI_CBLAS_CBLAS_CHEMV_IDX 11
-#define TI_CBLAS_CBLAS_CHER_IDX 12
-#define TI_CBLAS_CBLAS_CHER2_IDX 13
-#define TI_CBLAS_CBLAS_CHER2K_IDX 14
-#define TI_CBLAS_CBLAS_CHERK_IDX 15
-#define TI_CBLAS_CBLAS_CHPMV_IDX 16
-#define TI_CBLAS_CBLAS_CHPR_IDX 17
-#define TI_CBLAS_CBLAS_CHPR2_IDX 18
-#define TI_CBLAS_CBLAS_CROTG_IDX 19
-#define TI_CBLAS_CBLAS_CSCAL_IDX 20
-#define TI_CBLAS_CBLAS_CSSCAL_IDX 21
-#define TI_CBLAS_CBLAS_CSWAP_IDX 22
-#define TI_CBLAS_CBLAS_CSYMM_IDX 23
-#define TI_CBLAS_CBLAS_CSYR2K_IDX 24
-#define TI_CBLAS_CBLAS_CSYRK_IDX 25
-#define TI_CBLAS_CBLAS_CTBMV_IDX 26
-#define TI_CBLAS_CBLAS_CTBSV_IDX 27
-#define TI_CBLAS_CBLAS_CTPMV_IDX 28
-#define TI_CBLAS_CBLAS_CTPSV_IDX 29
-#define TI_CBLAS_CBLAS_CTRMM_IDX 30
-#define TI_CBLAS_CBLAS_CTRMV_IDX 31
-#define TI_CBLAS_CBLAS_CTRSM_IDX 32
-#define TI_CBLAS_CBLAS_CTRSV_IDX 33
-#define TI_CBLAS_CBLAS_DASUM_IDX 34
-#define TI_CBLAS_CBLAS_DAXPY_IDX 35
-#define TI_CBLAS_CBLAS_DCOPY_IDX 36
-#define TI_CBLAS_CBLAS_DDOT_IDX 37
-#define TI_CBLAS_CBLAS_DGBMV_IDX 38
-#define TI_CBLAS_CBLAS_DGEMM_IDX 39
-#define TI_CBLAS_CBLAS_DGEMV_IDX 40
-#define TI_CBLAS_CBLAS_DGER_IDX 41
-#define TI_CBLAS_CBLAS_DNRM2_IDX 42
-#define TI_CBLAS_CBLAS_DROT_IDX 43
-#define TI_CBLAS_CBLAS_DROTG_IDX 44
-#define TI_CBLAS_CBLAS_DROTM_IDX 45
-#define TI_CBLAS_CBLAS_DROTMG_IDX 46
-#define TI_CBLAS_CBLAS_DSBMV_IDX 47
-#define TI_CBLAS_CBLAS_DSCAL_IDX 48
-#define TI_CBLAS_CBLAS_DSDOT_IDX 49
-#define TI_CBLAS_CBLAS_DSPMV_IDX 50
-#define TI_CBLAS_CBLAS_DSPR_IDX 51
-#define TI_CBLAS_CBLAS_DSPR2_IDX 52
-#define TI_CBLAS_CBLAS_DSWAP_IDX 53
-#define TI_CBLAS_CBLAS_DSYMM_IDX 54
-#define TI_CBLAS_CBLAS_DSYMV_IDX 55
-#define TI_CBLAS_CBLAS_DSYR_IDX 56
-#define TI_CBLAS_CBLAS_DSYR2_IDX 57
-#define TI_CBLAS_CBLAS_DSYR2K_IDX 58
-#define TI_CBLAS_CBLAS_DSYRK_IDX 59
-#define TI_CBLAS_CBLAS_DTBMV_IDX 60
-#define TI_CBLAS_CBLAS_DTBSV_IDX 61
-#define TI_CBLAS_CBLAS_DTPMV_IDX 62
-#define TI_CBLAS_CBLAS_DTPSV_IDX 63
-#define TI_CBLAS_CBLAS_DTRMM_IDX 64
-#define TI_CBLAS_CBLAS_DTRMV_IDX 65
-#define TI_CBLAS_CBLAS_DTRSM_IDX 66
-#define TI_CBLAS_CBLAS_DTRSV_IDX 67
-#define TI_CBLAS_CBLAS_DZASUM_IDX 68
-#define TI_CBLAS_CBLAS_DZNRM2_IDX 69
-#define TI_CBLAS_CBLAS_ICAMAX_IDX 70
-#define TI_CBLAS_CBLAS_IDAMAX_IDX 71
-#define TI_CBLAS_CBLAS_ISAMAX_IDX 72
-#define TI_CBLAS_CBLAS_IZAMAX_IDX 73
-#define TI_CBLAS_CBLAS_SASUM_IDX 74
-#define TI_CBLAS_CBLAS_SAXPY_IDX 75
-#define TI_CBLAS_CBLAS_SCASUM_IDX 76
-#define TI_CBLAS_CBLAS_SCNRM2_IDX 77
-#define TI_CBLAS_CBLAS_SCOPY_IDX 78
-#define TI_CBLAS_CBLAS_SDOT_IDX 79
-#define TI_CBLAS_CBLAS_SDSDOT_IDX 80
-#define TI_CBLAS_CBLAS_SGBMV_IDX 81
-#define TI_CBLAS_CBLAS_SGEMM_IDX 82
-#define TI_CBLAS_CBLAS_SGEMV_IDX 83
-#define TI_CBLAS_CBLAS_SGER_IDX 84
-#define TI_CBLAS_CBLAS_SNRM2_IDX 85
-#define TI_CBLAS_CBLAS_SROT_IDX 86
-#define TI_CBLAS_CBLAS_SROTG_IDX 87
-#define TI_CBLAS_CBLAS_SROTM_IDX 88
-#define TI_CBLAS_CBLAS_SROTMG_IDX 89
-#define TI_CBLAS_CBLAS_SSBMV_IDX 90
-#define TI_CBLAS_CBLAS_SSCAL_IDX 91
-#define TI_CBLAS_CBLAS_SSPMV_IDX 92
-#define TI_CBLAS_CBLAS_SSPR_IDX 93
-#define TI_CBLAS_CBLAS_SSPR2_IDX 94
-#define TI_CBLAS_CBLAS_SSWAP_IDX 95
-#define TI_CBLAS_CBLAS_SSYMM_IDX 96
-#define TI_CBLAS_CBLAS_SSYMV_IDX 97
-#define TI_CBLAS_CBLAS_SSYR_IDX 98
-#define TI_CBLAS_CBLAS_SSYR2_IDX 99
-#define TI_CBLAS_CBLAS_SSYR2K_IDX 100
-#define TI_CBLAS_CBLAS_SSYRK_IDX 101
-#define TI_CBLAS_CBLAS_STBMV_IDX 102
-#define TI_CBLAS_CBLAS_STBSV_IDX 103
-#define TI_CBLAS_CBLAS_STPMV_IDX 104
-#define TI_CBLAS_CBLAS_STPSV_IDX 105
-#define TI_CBLAS_CBLAS_STRMM_IDX 106
-#define TI_CBLAS_CBLAS_STRMV_IDX 107
-#define TI_CBLAS_CBLAS_STRSM_IDX 108
-#define TI_CBLAS_CBLAS_STRSV_IDX 109
-#define TI_CBLAS_CBLAS_XERBLA_IDX 110
-#define TI_CBLAS_CBLAS_ZAXPY_IDX 111
-#define TI_CBLAS_CBLAS_ZCOPY_IDX 112
-#define TI_CBLAS_CBLAS_ZDOTC_SUB_IDX 113
-#define TI_CBLAS_CBLAS_ZDOTU_SUB_IDX 114
-#define TI_CBLAS_CBLAS_ZDSCAL_IDX 115
-#define TI_CBLAS_CBLAS_ZGBMV_IDX 116
-#define TI_CBLAS_CBLAS_ZGEMM_IDX 117
-#define TI_CBLAS_CBLAS_ZGEMV_IDX 118
-#define TI_CBLAS_CBLAS_ZGERC_IDX 119
-#define TI_CBLAS_CBLAS_ZGERU_IDX 120
-#define TI_CBLAS_CBLAS_ZHBMV_IDX 121
-#define TI_CBLAS_CBLAS_ZHEMM_IDX 122
-#define TI_CBLAS_CBLAS_ZHEMV_IDX 123
-#define TI_CBLAS_CBLAS_ZHER_IDX 124
-#define TI_CBLAS_CBLAS_ZHER2_IDX 125
-#define TI_CBLAS_CBLAS_ZHER2K_IDX 126
-#define TI_CBLAS_CBLAS_ZHERK_IDX 127
-#define TI_CBLAS_CBLAS_ZHPMV_IDX 128
-#define TI_CBLAS_CBLAS_ZHPR_IDX 129
-#define TI_CBLAS_CBLAS_ZHPR2_IDX 130
-#define TI_CBLAS_CBLAS_ZROTG_IDX 131
-#define TI_CBLAS_CBLAS_ZSCAL_IDX 132
-#define TI_CBLAS_CBLAS_ZSWAP_IDX 133
-#define TI_CBLAS_CBLAS_ZSYMM_IDX 134
-#define TI_CBLAS_CBLAS_ZSYR2K_IDX 135
-#define TI_CBLAS_CBLAS_ZSYRK_IDX 136
-#define TI_CBLAS_CBLAS_ZTBMV_IDX 137
-#define TI_CBLAS_CBLAS_ZTBSV_IDX 138
-#define TI_CBLAS_CBLAS_ZTPMV_IDX 139
-#define TI_CBLAS_CBLAS_ZTPSV_IDX 140
-#define TI_CBLAS_CBLAS_ZTRMM_IDX 141
-#define TI_CBLAS_CBLAS_ZTRMV_IDX 142
-#define TI_CBLAS_CBLAS_ZTRSM_IDX 143
-#define TI_CBLAS_CBLAS_ZTRSV_IDX 144
-#define TI_CBLAS_NUM_KERNELS 145
-
 /* Level 3 kernels offload table */
 /* Number of points in each dimension. ARM processing and DSP processing time
    are measured for each point to determine offload or not. */
@@ -336,8 +154,8 @@ extern int ti_cblas_offload;
 #define TRSM_OFFLOAD_TBL_SIZE (NUM_PNT_EACH_DIM*NUM_PNT_EACH_DIM)
 
 /* compile time defaults */
-#ifndef TI_CBLAS_OFFLOAD
-#define TI_CBLAS_OFFLOAD "002"
+#ifndef TI_CBLAS_OFFLOAD_DEF
+#define TI_CBLAS_OFFLOAD_DEF "002"
 #endif
 
 /* macros used for BLAS/LAPACK buffer size calculations */
@@ -396,9 +214,9 @@ extern int csymm_offload_dsp(const enum CBLAS_ORDER Order, const enum CBLAS_SIDE
 extern int zsymm_offload_dsp(const enum CBLAS_ORDER Order, const enum CBLAS_SIDE Side, 
                              int M, int N);
 extern int chemm_offload_dsp(const enum CBLAS_ORDER Order, const enum CBLAS_SIDE Side, 
-                      int M, int N);
+                             int M, int N);
 extern int zhemm_offload_dsp(const enum CBLAS_ORDER Order, const enum CBLAS_SIDE Side, 
-                      int M, int N);
+                             int M, int N);
 extern int ssyrk_offload_dsp(const enum CBLAS_ORDER Order, int N, int K);
 extern int dsyrk_offload_dsp(const enum CBLAS_ORDER Order, int N, int K);
 extern int csyrk_offload_dsp(const enum CBLAS_ORDER Order, int N, int K);
@@ -412,21 +230,21 @@ extern int zsyr2k_offload_dsp(const enum CBLAS_ORDER Order, int N, int K);
 extern int cher2k_offload_dsp(const enum CBLAS_ORDER Order, int N, int K);
 extern int zher2k_offload_dsp(const enum CBLAS_ORDER Order, int N, int K);
 extern int strmm_offload_dsp(const enum CBLAS_ORDER Order, const enum CBLAS_SIDE Side, 
-                      int M, int N);
+                             int M, int N);
 extern int dtrmm_offload_dsp(const enum CBLAS_ORDER Order, const enum CBLAS_SIDE Side, 
-                      int M, int N);
+                             int M, int N);
 extern int ctrmm_offload_dsp(const enum CBLAS_ORDER Order, const enum CBLAS_SIDE Side, 
-                      int M, int N);
+                             int M, int N);
 extern int ztrmm_offload_dsp(const enum CBLAS_ORDER Order, const enum CBLAS_SIDE Side, 
-                      int M, int N);
+                             int M, int N);
 extern int strsm_offload_dsp(const enum CBLAS_ORDER Order, const enum CBLAS_SIDE Side, 
-                      int M, int N);
+                             int M, int N);
 extern int dtrsm_offload_dsp(const enum CBLAS_ORDER Order, const enum CBLAS_SIDE Side, 
-                      int M, int N);
+                             int M, int N);
 extern int ctrsm_offload_dsp(const enum CBLAS_ORDER Order, const enum CBLAS_SIDE Side, 
-                      int M, int N);
+                             int M, int N);
 extern int ztrsm_offload_dsp(const enum CBLAS_ORDER Order, const enum CBLAS_SIDE Side, 
-                      int M, int N);
+                             int M, int N);
 
 #endif
 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_caxpy.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_caxpy.c
index 233c69b..cfc47c7 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_caxpy.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_caxpy.c
@@ -66,7 +66,7 @@ void cblas_caxpy(const int N, const void *alpha, const void *X, const int incX,
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_CAXPY_IDX, "ocl_cblas_caxpy");
+    __K =  ti_cblas_get_kernel("ocl_cblas_caxpy");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_ccopy.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_ccopy.c
index 2a5fa6b..88373f1 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_ccopy.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_ccopy.c
@@ -66,7 +66,7 @@ void cblas_ccopy(const int N, const void *X, const int incX, void *Y, const int
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_CCOPY_IDX, "ocl_cblas_ccopy");
+    __K =  ti_cblas_get_kernel("ocl_cblas_ccopy");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_cdotc_sub.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_cdotc_sub.c
index bb8fd34..5a24501 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_cdotc_sub.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_cdotc_sub.c
@@ -66,7 +66,7 @@ void cblas_cdotc_sub(const int N, const void *X, const int incX, const void *Y,
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_CDOTC_SUB_IDX, "ocl_cblas_cdotc_sub");
+    __K =  ti_cblas_get_kernel("ocl_cblas_cdotc_sub");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_cdotu_sub.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_cdotu_sub.c
index 2ba9005..95473b1 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_cdotu_sub.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_cdotu_sub.c
@@ -66,7 +66,7 @@ void cblas_cdotu_sub(const int N, const void *X, const int incX, const void *Y,
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_CDOTU_SUB_IDX, "ocl_cblas_cdotu_sub");
+    __K =  ti_cblas_get_kernel("ocl_cblas_cdotu_sub");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_cgbmv.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_cgbmv.c
index 4a04d07..cc9dd91 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_cgbmv.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_cgbmv.c
@@ -66,7 +66,7 @@ void cblas_cgbmv(const enum CBLAS_ORDER order, const enum CBLAS_TRANSPOSE TransA
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_CGBMV_IDX, "ocl_cblas_cgbmv");
+    __K =  ti_cblas_get_kernel("ocl_cblas_cgbmv");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_cgemm.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_cgemm.c
index 9880b6d..38bc67c 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_cgemm.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_cgemm.c
@@ -66,7 +66,7 @@ void cblas_cgemm(const enum CBLAS_ORDER Order, const enum CBLAS_TRANSPOSE TransA
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_CGEMM_IDX, "ocl_cblas_cgemm");
+    __K =  ti_cblas_get_kernel("ocl_cblas_cgemm");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_cgemv.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_cgemv.c
index c4d35c8..18e9c5c 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_cgemv.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_cgemv.c
@@ -66,7 +66,7 @@ void cblas_cgemv(const enum CBLAS_ORDER order, const enum CBLAS_TRANSPOSE TransA
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_CGEMV_IDX, "ocl_cblas_cgemv");
+    __K =  ti_cblas_get_kernel("ocl_cblas_cgemv");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_cgerc.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_cgerc.c
index 3c89112..5f4cfc5 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_cgerc.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_cgerc.c
@@ -66,7 +66,7 @@ void cblas_cgerc(const enum CBLAS_ORDER order, const int M, const int N, const v
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_CGERC_IDX, "ocl_cblas_cgerc");
+    __K =  ti_cblas_get_kernel("ocl_cblas_cgerc");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_cgeru.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_cgeru.c
index 3e951aa..1bcc580 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_cgeru.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_cgeru.c
@@ -66,7 +66,7 @@ void cblas_cgeru(const enum CBLAS_ORDER order, const int M, const int N, const v
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_CGERU_IDX, "ocl_cblas_cgeru");
+    __K =  ti_cblas_get_kernel("ocl_cblas_cgeru");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_chbmv.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_chbmv.c
index cfdb1f1..05c77a8 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_chbmv.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_chbmv.c
@@ -66,7 +66,7 @@ void cblas_chbmv(const enum CBLAS_ORDER order, const enum CBLAS_UPLO Uplo, const
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_CHBMV_IDX, "ocl_cblas_chbmv");
+    __K =  ti_cblas_get_kernel("ocl_cblas_chbmv");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_chemm.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_chemm.c
index 0e8eb77..c02e2db 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_chemm.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_chemm.c
@@ -66,7 +66,7 @@ void cblas_chemm(const enum CBLAS_ORDER Order, const enum CBLAS_SIDE Side, const
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_CHEMM_IDX, "ocl_cblas_chemm");
+    __K =  ti_cblas_get_kernel("ocl_cblas_chemm");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_chemv.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_chemv.c
index 7421539..7e72a78 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_chemv.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_chemv.c
@@ -66,7 +66,7 @@ void cblas_chemv(const enum CBLAS_ORDER order, const enum CBLAS_UPLO Uplo, const
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_CHEMV_IDX, "ocl_cblas_chemv");
+    __K =  ti_cblas_get_kernel("ocl_cblas_chemv");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_cher.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_cher.c
index fb566fd..73075de 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_cher.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_cher.c
@@ -66,7 +66,7 @@ void cblas_cher(const enum CBLAS_ORDER order, const enum CBLAS_UPLO Uplo, const
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_CHER_IDX, "ocl_cblas_cher");
+    __K =  ti_cblas_get_kernel("ocl_cblas_cher");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_cher2.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_cher2.c
index f65b307..3240ed2 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_cher2.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_cher2.c
@@ -66,7 +66,7 @@ void cblas_cher2(const enum CBLAS_ORDER order, const enum CBLAS_UPLO Uplo, const
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_CHER2_IDX, "ocl_cblas_cher2");
+    __K =  ti_cblas_get_kernel("ocl_cblas_cher2");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_cher2k.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_cher2k.c
index 51ed406..f119c3b 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_cher2k.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_cher2k.c
@@ -66,7 +66,7 @@ void cblas_cher2k(const enum CBLAS_ORDER Order, const enum CBLAS_UPLO Uplo, cons
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_CHER2K_IDX, "ocl_cblas_cher2k");
+    __K =  ti_cblas_get_kernel("ocl_cblas_cher2k");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_cherk.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_cherk.c
index 5e45c12..5b3dab3 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_cherk.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_cherk.c
@@ -66,7 +66,7 @@ void cblas_cherk(const enum CBLAS_ORDER Order, const enum CBLAS_UPLO Uplo, const
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_CHERK_IDX, "ocl_cblas_cherk");
+    __K =  ti_cblas_get_kernel("ocl_cblas_cherk");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_chpmv.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_chpmv.c
index 03e14dd..bdbe0cf 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_chpmv.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_chpmv.c
@@ -66,7 +66,7 @@ void cblas_chpmv(const enum CBLAS_ORDER order, const enum CBLAS_UPLO Uplo, const
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_CHPMV_IDX, "ocl_cblas_chpmv");
+    __K =  ti_cblas_get_kernel("ocl_cblas_chpmv");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_chpr.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_chpr.c
index 315d80d..c68c478 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_chpr.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_chpr.c
@@ -66,7 +66,7 @@ void cblas_chpr(const enum CBLAS_ORDER order, const enum CBLAS_UPLO Uplo, const
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_CHPR_IDX, "ocl_cblas_chpr");
+    __K =  ti_cblas_get_kernel("ocl_cblas_chpr");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_chpr2.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_chpr2.c
index 7f0f765..ca0ce5c 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_chpr2.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_chpr2.c
@@ -66,7 +66,7 @@ void cblas_chpr2(const enum CBLAS_ORDER order, const enum CBLAS_UPLO Uplo, const
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_CHPR2_IDX, "ocl_cblas_chpr2");
+    __K =  ti_cblas_get_kernel("ocl_cblas_chpr2");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_crotg.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_crotg.c
index 8eb77ee..492e279 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_crotg.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_crotg.c
@@ -66,7 +66,7 @@ void cblas_crotg( void *a, void *b, float *c, void *s)
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_CROTG_IDX, "ocl_cblas_crotg");
+    __K =  ti_cblas_get_kernel("ocl_cblas_crotg");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_cscal.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_cscal.c
index 7933cba..0d0d659 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_cscal.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_cscal.c
@@ -66,7 +66,7 @@ void cblas_cscal(const int N, const void *alpha, void *X, const int incX)
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_CSCAL_IDX, "ocl_cblas_cscal");
+    __K =  ti_cblas_get_kernel("ocl_cblas_cscal");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_csscal.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_csscal.c
index 161b6d5..f18b209 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_csscal.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_csscal.c
@@ -66,7 +66,7 @@ void cblas_csscal(const int N, const float alpha, void *X, const int incX)
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_CSSCAL_IDX, "ocl_cblas_csscal");
+    __K =  ti_cblas_get_kernel("ocl_cblas_csscal");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_cswap.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_cswap.c
index 1dac9fe..c29f25c 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_cswap.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_cswap.c
@@ -78,7 +78,7 @@ void cblas_cswap(const int N, void *X, const int incX, void *Y, const int incY)
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_CSWAP_IDX, "ocl_cblas_cswap");
+    __K =  ti_cblas_get_kernel("ocl_cblas_cswap");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_csymm.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_csymm.c
index 41558f8..5117b67 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_csymm.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_csymm.c
@@ -66,7 +66,7 @@ void cblas_csymm(const enum CBLAS_ORDER Order, const enum CBLAS_SIDE Side, const
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_CSYMM_IDX, "ocl_cblas_csymm");
+    __K =  ti_cblas_get_kernel("ocl_cblas_csymm");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_csyr2k.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_csyr2k.c
index 3a3611f..0ca1f1d 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_csyr2k.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_csyr2k.c
@@ -66,7 +66,7 @@ void cblas_csyr2k(const enum CBLAS_ORDER Order, const enum CBLAS_UPLO Uplo, cons
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_CSYR2K_IDX, "ocl_cblas_csyr2k");
+    __K =  ti_cblas_get_kernel("ocl_cblas_csyr2k");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_csyrk.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_csyrk.c
index 8a0a4a4..d57db3a 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_csyrk.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_csyrk.c
@@ -66,7 +66,7 @@ void cblas_csyrk(const enum CBLAS_ORDER Order, const enum CBLAS_UPLO Uplo, const
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_CSYRK_IDX, "ocl_cblas_csyrk");
+    __K =  ti_cblas_get_kernel("ocl_cblas_csyrk");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_ctbmv.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_ctbmv.c
index 237b7fe..493897d 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_ctbmv.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_ctbmv.c
@@ -66,7 +66,7 @@ void cblas_ctbmv(const enum CBLAS_ORDER order, const enum CBLAS_UPLO Uplo, const
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_CTBMV_IDX, "ocl_cblas_ctbmv");
+    __K =  ti_cblas_get_kernel("ocl_cblas_ctbmv");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_ctbsv.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_ctbsv.c
index b03c316..d12a66b 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_ctbsv.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_ctbsv.c
@@ -66,7 +66,7 @@ void cblas_ctbsv(const enum CBLAS_ORDER order, const enum CBLAS_UPLO Uplo, const
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_CTBSV_IDX, "ocl_cblas_ctbsv");
+    __K =  ti_cblas_get_kernel("ocl_cblas_ctbsv");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_ctpmv.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_ctpmv.c
index 816fed2..c5ef3a6 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_ctpmv.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_ctpmv.c
@@ -66,7 +66,7 @@ void cblas_ctpmv(const enum CBLAS_ORDER order, const enum CBLAS_UPLO Uplo, const
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_CTPMV_IDX, "ocl_cblas_ctpmv");
+    __K =  ti_cblas_get_kernel("ocl_cblas_ctpmv");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_ctpsv.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_ctpsv.c
index 44d25ed..c554877 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_ctpsv.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_ctpsv.c
@@ -66,7 +66,7 @@ void cblas_ctpsv(const enum CBLAS_ORDER order, const enum CBLAS_UPLO Uplo, const
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_CTPSV_IDX, "ocl_cblas_ctpsv");
+    __K =  ti_cblas_get_kernel("ocl_cblas_ctpsv");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_ctrmm.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_ctrmm.c
index 17713af..6740788 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_ctrmm.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_ctrmm.c
@@ -66,7 +66,7 @@ void cblas_ctrmm(const enum CBLAS_ORDER Order, const enum CBLAS_SIDE Side, const
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_CTRMM_IDX, "ocl_cblas_ctrmm");
+    __K =  ti_cblas_get_kernel("ocl_cblas_ctrmm");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_ctrmv.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_ctrmv.c
index e2461a5..890bbd2 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_ctrmv.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_ctrmv.c
@@ -66,7 +66,7 @@ void cblas_ctrmv(const enum CBLAS_ORDER order, const enum CBLAS_UPLO Uplo, const
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_CTRMV_IDX, "ocl_cblas_ctrmv");
+    __K =  ti_cblas_get_kernel("ocl_cblas_ctrmv");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_ctrsm.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_ctrsm.c
index 7279d08..1a104b8 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_ctrsm.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_ctrsm.c
@@ -66,7 +66,7 @@ void cblas_ctrsm(const enum CBLAS_ORDER Order, const enum CBLAS_SIDE Side, const
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_CTRSM_IDX, "ocl_cblas_ctrsm");
+    __K =  ti_cblas_get_kernel("ocl_cblas_ctrsm");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_ctrsv.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_ctrsv.c
index 2ecf168..1969368 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_ctrsv.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_ctrsv.c
@@ -66,7 +66,7 @@ void cblas_ctrsv(const enum CBLAS_ORDER order, const enum CBLAS_UPLO Uplo, const
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_CTRSV_IDX, "ocl_cblas_ctrsv");
+    __K =  ti_cblas_get_kernel("ocl_cblas_ctrsv");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_dasum.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_dasum.c
index 3425c22..466478a 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_dasum.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_dasum.c
@@ -66,7 +66,7 @@ double cblas_dasum(const int N, const double *X, const int incX)
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_DASUM_IDX, "ocl_cblas_dasum");
+    __K =  ti_cblas_get_kernel("ocl_cblas_dasum");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_daxpy.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_daxpy.c
index 64a6824..62fea45 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_daxpy.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_daxpy.c
@@ -66,7 +66,7 @@ void cblas_daxpy(const int N, const double alpha, const double *X, const int inc
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_DAXPY_IDX, "ocl_cblas_daxpy");
+    __K =  ti_cblas_get_kernel("ocl_cblas_daxpy");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_dcopy.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_dcopy.c
index c145e2d..52ab3c7 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_dcopy.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_dcopy.c
@@ -66,7 +66,7 @@ void cblas_dcopy(const int N, const double *X, const int incX, double *Y, const
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_DCOPY_IDX, "ocl_cblas_dcopy");
+    __K =  ti_cblas_get_kernel("ocl_cblas_dcopy");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_ddot.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_ddot.c
index aa1f769..1717502 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_ddot.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_ddot.c
@@ -66,7 +66,7 @@ double cblas_ddot(const int N, const double *X, const int incX, const double *Y,
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_DDOT_IDX, "ocl_cblas_ddot");
+    __K =  ti_cblas_get_kernel("ocl_cblas_ddot");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_dgbmv.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_dgbmv.c
index 1ce1c80..15f41c9 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_dgbmv.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_dgbmv.c
@@ -66,7 +66,7 @@ void cblas_dgbmv(const enum CBLAS_ORDER order, const enum CBLAS_TRANSPOSE TransA
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_DGBMV_IDX, "ocl_cblas_dgbmv");
+    __K =  ti_cblas_get_kernel("ocl_cblas_dgbmv");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_dgemm.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_dgemm.c
index 92ecb14..58f6abc 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_dgemm.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_dgemm.c
@@ -66,7 +66,7 @@ void cblas_dgemm(const enum CBLAS_ORDER Order, const enum CBLAS_TRANSPOSE TransA
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_DGEMM_IDX, "ocl_cblas_dgemm");
+    __K =  ti_cblas_get_kernel("ocl_cblas_dgemm");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_dgemv.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_dgemv.c
index 2c3a5e2..ba2fefb 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_dgemv.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_dgemv.c
@@ -66,7 +66,7 @@ void cblas_dgemv(const enum CBLAS_ORDER order, const enum CBLAS_TRANSPOSE TransA
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_DGEMV_IDX, "ocl_cblas_dgemv");
+    __K =  ti_cblas_get_kernel("ocl_cblas_dgemv");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_dger.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_dger.c
index 1c911e7..c30dede 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_dger.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_dger.c
@@ -66,7 +66,7 @@ void cblas_dger(const enum CBLAS_ORDER order, const int M, const int N, const do
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_DGER_IDX, "ocl_cblas_dger");
+    __K =  ti_cblas_get_kernel("ocl_cblas_dger");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_dnrm2.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_dnrm2.c
index a3c14c6..3b92435 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_dnrm2.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_dnrm2.c
@@ -66,7 +66,7 @@ double cblas_dnrm2(const int N, const double *X, const int incX)
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_DNRM2_IDX, "ocl_cblas_dnrm2");
+    __K =  ti_cblas_get_kernel("ocl_cblas_dnrm2");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_drot.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_drot.c
index dd86c24..dfc2b82 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_drot.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_drot.c
@@ -66,7 +66,7 @@ void cblas_drot(const int N, double *X, const int incX, double *Y, const int inc
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_DROT_IDX, "ocl_cblas_drot");
+    __K =  ti_cblas_get_kernel("ocl_cblas_drot");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_drotg.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_drotg.c
index 5b3732c..a09995b 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_drotg.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_drotg.c
@@ -66,7 +66,7 @@ void cblas_drotg(double *a, double *b, double *c, double *s)
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_DROTG_IDX, "ocl_cblas_drotg");
+    __K =  ti_cblas_get_kernel("ocl_cblas_drotg");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_drotm.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_drotm.c
index 0ba101a..30fd5d0 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_drotm.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_drotm.c
@@ -66,7 +66,7 @@ void cblas_drotm(const int N, double *X, const int incX, double *Y, const int in
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_DROTM_IDX, "ocl_cblas_drotm");
+    __K =  ti_cblas_get_kernel("ocl_cblas_drotm");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_drotmg.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_drotmg.c
index efbcc94..921ece4 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_drotmg.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_drotmg.c
@@ -66,7 +66,7 @@ void cblas_drotmg(double *d1, double *d2, double *b1, const double b2, double *P
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_DROTMG_IDX, "ocl_cblas_drotmg");
+    __K =  ti_cblas_get_kernel("ocl_cblas_drotmg");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_dsbmv.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_dsbmv.c
index b826766..7532d48 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_dsbmv.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_dsbmv.c
@@ -66,7 +66,7 @@ void cblas_dsbmv(const enum CBLAS_ORDER order, const enum CBLAS_UPLO Uplo, const
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_DSBMV_IDX, "ocl_cblas_dsbmv");
+    __K =  ti_cblas_get_kernel("ocl_cblas_dsbmv");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_dscal.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_dscal.c
index 93d81f0..dd75f67 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_dscal.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_dscal.c
@@ -66,7 +66,7 @@ void cblas_dscal(const int N, const double alpha, double *X, const int incX)
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_DSCAL_IDX, "ocl_cblas_dscal");
+    __K =  ti_cblas_get_kernel("ocl_cblas_dscal");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_dsdot.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_dsdot.c
index 8f02791..602414c 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_dsdot.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_dsdot.c
@@ -66,7 +66,7 @@ double cblas_dsdot(const int N, const float *X, const int incX, const float *Y,
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_DSDOT_IDX, "ocl_cblas_dsdot");
+    __K =  ti_cblas_get_kernel("ocl_cblas_dsdot");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_dspmv.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_dspmv.c
index 13b403b..0c24fee 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_dspmv.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_dspmv.c
@@ -66,7 +66,7 @@ void cblas_dspmv(const enum CBLAS_ORDER order, const enum CBLAS_UPLO Uplo, const
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_DSPMV_IDX, "ocl_cblas_dspmv");
+    __K =  ti_cblas_get_kernel("ocl_cblas_dspmv");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_dspr.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_dspr.c
index c347969..2f75cf9 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_dspr.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_dspr.c
@@ -66,7 +66,7 @@ void cblas_dspr(const enum CBLAS_ORDER order, const enum CBLAS_UPLO Uplo, const
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_DSPR_IDX, "ocl_cblas_dspr");
+    __K =  ti_cblas_get_kernel("ocl_cblas_dspr");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_dspr2.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_dspr2.c
index 812bb3b..38620ae 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_dspr2.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_dspr2.c
@@ -66,7 +66,7 @@ void cblas_dspr2(const enum CBLAS_ORDER order, const enum CBLAS_UPLO Uplo, const
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_DSPR2_IDX, "ocl_cblas_dspr2");
+    __K =  ti_cblas_get_kernel("ocl_cblas_dspr2");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_dswap.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_dswap.c
index d17893d..9cd70b9 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_dswap.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_dswap.c
@@ -78,7 +78,7 @@ void cblas_dswap(const int N, double *X, const int incX, double *Y, const int in
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_DSWAP_IDX, "ocl_cblas_dswap");
+    __K =  ti_cblas_get_kernel("ocl_cblas_dswap");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_dsymm.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_dsymm.c
index d5c28bc..4136f9e 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_dsymm.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_dsymm.c
@@ -66,7 +66,7 @@ void cblas_dsymm(const enum CBLAS_ORDER Order, const enum CBLAS_SIDE Side, const
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_DSYMM_IDX, "ocl_cblas_dsymm");
+    __K =  ti_cblas_get_kernel("ocl_cblas_dsymm");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_dsymv.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_dsymv.c
index 80eb1a6..5a5d1c5 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_dsymv.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_dsymv.c
@@ -66,7 +66,7 @@ void cblas_dsymv(const enum CBLAS_ORDER order, const enum CBLAS_UPLO Uplo, const
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_DSYMV_IDX, "ocl_cblas_dsymv");
+    __K =  ti_cblas_get_kernel("ocl_cblas_dsymv");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_dsyr.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_dsyr.c
index 6b05e31..6df6638 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_dsyr.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_dsyr.c
@@ -66,7 +66,7 @@ void cblas_dsyr(const enum CBLAS_ORDER order, const enum CBLAS_UPLO Uplo, const
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_DSYR_IDX, "ocl_cblas_dsyr");
+    __K =  ti_cblas_get_kernel("ocl_cblas_dsyr");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_dsyr2.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_dsyr2.c
index 86a5a48..384c68f 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_dsyr2.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_dsyr2.c
@@ -66,7 +66,7 @@ void cblas_dsyr2(const enum CBLAS_ORDER order, const enum CBLAS_UPLO Uplo, const
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_DSYR2_IDX, "ocl_cblas_dsyr2");
+    __K =  ti_cblas_get_kernel("ocl_cblas_dsyr2");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_dsyr2k.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_dsyr2k.c
index 4caffb7..3546475 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_dsyr2k.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_dsyr2k.c
@@ -66,7 +66,7 @@ void cblas_dsyr2k(const enum CBLAS_ORDER Order, const enum CBLAS_UPLO Uplo, cons
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_DSYR2K_IDX, "ocl_cblas_dsyr2k");
+    __K =  ti_cblas_get_kernel("ocl_cblas_dsyr2k");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_dsyrk.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_dsyrk.c
index 341ff82..923a66a 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_dsyrk.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_dsyrk.c
@@ -66,7 +66,7 @@ void cblas_dsyrk(const enum CBLAS_ORDER Order, const enum CBLAS_UPLO Uplo, const
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_DSYRK_IDX, "ocl_cblas_dsyrk");
+    __K =  ti_cblas_get_kernel("ocl_cblas_dsyrk");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_dtbmv.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_dtbmv.c
index d95c34f..102b842 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_dtbmv.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_dtbmv.c
@@ -66,7 +66,7 @@ void cblas_dtbmv(const enum CBLAS_ORDER order, const enum CBLAS_UPLO Uplo, const
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_DTBMV_IDX, "ocl_cblas_dtbmv");
+    __K =  ti_cblas_get_kernel("ocl_cblas_dtbmv");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_dtbsv.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_dtbsv.c
index c375c4e..dde82ef 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_dtbsv.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_dtbsv.c
@@ -66,7 +66,7 @@ void cblas_dtbsv(const enum CBLAS_ORDER order, const enum CBLAS_UPLO Uplo, const
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_DTBSV_IDX, "ocl_cblas_dtbsv");
+    __K =  ti_cblas_get_kernel("ocl_cblas_dtbsv");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_dtpmv.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_dtpmv.c
index 105c178..0a5da49 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_dtpmv.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_dtpmv.c
@@ -66,7 +66,7 @@ void cblas_dtpmv(const enum CBLAS_ORDER order, const enum CBLAS_UPLO Uplo, const
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_DTPMV_IDX, "ocl_cblas_dtpmv");
+    __K =  ti_cblas_get_kernel("ocl_cblas_dtpmv");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_dtpsv.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_dtpsv.c
index 32311ab..0f2abb2 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_dtpsv.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_dtpsv.c
@@ -66,7 +66,7 @@ void cblas_dtpsv(const enum CBLAS_ORDER order, const enum CBLAS_UPLO Uplo, const
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_DTPSV_IDX, "ocl_cblas_dtpsv");
+    __K =  ti_cblas_get_kernel("ocl_cblas_dtpsv");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_dtrmm.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_dtrmm.c
index cac7e27..28537a3 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_dtrmm.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_dtrmm.c
@@ -66,7 +66,7 @@ void cblas_dtrmm(const enum CBLAS_ORDER Order, const enum CBLAS_SIDE Side, const
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_DTRMM_IDX, "ocl_cblas_dtrmm");
+    __K =  ti_cblas_get_kernel("ocl_cblas_dtrmm");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_dtrmv.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_dtrmv.c
index d281310..14d2000 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_dtrmv.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_dtrmv.c
@@ -66,7 +66,7 @@ void cblas_dtrmv(const enum CBLAS_ORDER order, const enum CBLAS_UPLO Uplo, const
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_DTRMV_IDX, "ocl_cblas_dtrmv");
+    __K =  ti_cblas_get_kernel("ocl_cblas_dtrmv");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_dtrsm.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_dtrsm.c
index a3bbc06..e53f21f 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_dtrsm.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_dtrsm.c
@@ -66,7 +66,7 @@ void cblas_dtrsm(const enum CBLAS_ORDER Order, const enum CBLAS_SIDE Side, const
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_DTRSM_IDX, "ocl_cblas_dtrsm");
+    __K =  ti_cblas_get_kernel("ocl_cblas_dtrsm");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_dtrsv.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_dtrsv.c
index 4001c3f..4d77335 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_dtrsv.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_dtrsv.c
@@ -66,7 +66,7 @@ void cblas_dtrsv(const enum CBLAS_ORDER order, const enum CBLAS_UPLO Uplo, const
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_DTRSV_IDX, "ocl_cblas_dtrsv");
+    __K =  ti_cblas_get_kernel("ocl_cblas_dtrsv");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_dzasum.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_dzasum.c
index 5b79f68..a1bb3d9 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_dzasum.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_dzasum.c
@@ -66,7 +66,7 @@ double cblas_dzasum(const int N, const void *X, const int incX)
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_DZASUM_IDX, "ocl_cblas_dzasum");
+    __K =  ti_cblas_get_kernel("ocl_cblas_dzasum");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_dznrm2.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_dznrm2.c
index 81d7e53..265d6f9 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_dznrm2.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_dznrm2.c
@@ -66,7 +66,7 @@ double cblas_dznrm2(const int N, const void *X, const int incX)
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_DZNRM2_IDX, "ocl_cblas_dznrm2");
+    __K =  ti_cblas_get_kernel("ocl_cblas_dznrm2");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_icamax.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_icamax.c
index 7cbcfc3..860f4d7 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_icamax.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_icamax.c
@@ -66,7 +66,7 @@ CBLAS_INDEX cblas_icamax(const int N, const void *X, const int incX)
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_ICAMAX_IDX, "ocl_cblas_icamax");
+    __K =  ti_cblas_get_kernel("ocl_cblas_icamax");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_idamax.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_idamax.c
index d393ea2..9aab317 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_idamax.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_idamax.c
@@ -66,7 +66,7 @@ CBLAS_INDEX cblas_idamax(const int N, const double *X, const int incX)
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_IDAMAX_IDX, "ocl_cblas_idamax");
+    __K =  ti_cblas_get_kernel("ocl_cblas_idamax");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_isamax.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_isamax.c
index 35d6aa6..b15a9e6 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_isamax.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_isamax.c
@@ -66,7 +66,7 @@ CBLAS_INDEX cblas_isamax(const int N, const float *X, const int incX)
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_ISAMAX_IDX, "ocl_cblas_isamax");
+    __K =  ti_cblas_get_kernel("ocl_cblas_isamax");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_izamax.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_izamax.c
index f235319..7fc6d66 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_izamax.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_izamax.c
@@ -66,7 +66,7 @@ CBLAS_INDEX cblas_izamax(const int N, const void *X, const int incX)
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_IZAMAX_IDX, "ocl_cblas_izamax");
+    __K =  ti_cblas_get_kernel("ocl_cblas_izamax");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_sasum.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_sasum.c
index 60b0b88..dd46358 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_sasum.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_sasum.c
@@ -66,7 +66,7 @@ float cblas_sasum(const int N, const float *X, const int incX)
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_SASUM_IDX, "ocl_cblas_sasum");
+    __K =  ti_cblas_get_kernel("ocl_cblas_sasum");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_saxpy.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_saxpy.c
index 2caee52..cb61305 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_saxpy.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_saxpy.c
@@ -66,7 +66,7 @@ void cblas_saxpy(const int N, const float alpha, const float *X, const int incX,
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_SAXPY_IDX, "ocl_cblas_saxpy");
+    __K =  ti_cblas_get_kernel("ocl_cblas_saxpy");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_scasum.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_scasum.c
index 67c18ea..3845b05 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_scasum.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_scasum.c
@@ -66,7 +66,7 @@ float cblas_scasum(const int N, const void *X, const int incX)
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_SCASUM_IDX, "ocl_cblas_scasum");
+    __K =  ti_cblas_get_kernel("ocl_cblas_scasum");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_scnrm2.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_scnrm2.c
index a9e2713..733a4e2 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_scnrm2.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_scnrm2.c
@@ -66,7 +66,7 @@ float cblas_scnrm2(const int N, const void *X, const int incX)
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_SCNRM2_IDX, "ocl_cblas_scnrm2");
+    __K =  ti_cblas_get_kernel("ocl_cblas_scnrm2");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_scopy.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_scopy.c
index 2c3a8f3..c3d12fb 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_scopy.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_scopy.c
@@ -66,7 +66,7 @@ void cblas_scopy(const int N, const float *X, const int incX, float *Y, const in
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_SCOPY_IDX, "ocl_cblas_scopy");
+    __K =  ti_cblas_get_kernel("ocl_cblas_scopy");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_sdot.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_sdot.c
index ab7d8d0..636464d 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_sdot.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_sdot.c
@@ -66,7 +66,7 @@ float cblas_sdot(const int N, const float *X, const int incX, const float *Y, co
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_SDOT_IDX, "ocl_cblas_sdot");
+    __K =  ti_cblas_get_kernel("ocl_cblas_sdot");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_sdsdot.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_sdsdot.c
index fd5adb9..ed4822a 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_sdsdot.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_sdsdot.c
@@ -66,7 +66,7 @@ float cblas_sdsdot(const int N, const float alpha, const float *X, const int inc
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_SDSDOT_IDX, "ocl_cblas_sdsdot");
+    __K =  ti_cblas_get_kernel("ocl_cblas_sdsdot");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_sgbmv.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_sgbmv.c
index 32d800a..c7f4895 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_sgbmv.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_sgbmv.c
@@ -66,7 +66,7 @@ void cblas_sgbmv(const enum CBLAS_ORDER order, const enum CBLAS_TRANSPOSE TransA
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_SGBMV_IDX, "ocl_cblas_sgbmv");
+    __K =  ti_cblas_get_kernel("ocl_cblas_sgbmv");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_sgemm.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_sgemm.c
index 9599e84..ecd1885 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_sgemm.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_sgemm.c
@@ -66,7 +66,7 @@ void cblas_sgemm(const enum CBLAS_ORDER Order, const enum CBLAS_TRANSPOSE TransA
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_SGEMM_IDX, "ocl_cblas_sgemm");
+    __K =  ti_cblas_get_kernel("ocl_cblas_sgemm");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_sgemv.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_sgemv.c
index 48a3de6..ad3adaf 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_sgemv.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_sgemv.c
@@ -66,7 +66,7 @@ void cblas_sgemv(const enum CBLAS_ORDER order, const enum CBLAS_TRANSPOSE TransA
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_SGEMV_IDX, "ocl_cblas_sgemv");
+    __K =  ti_cblas_get_kernel("ocl_cblas_sgemv");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_sger.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_sger.c
index 6780fa9..b6b3680 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_sger.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_sger.c
@@ -66,7 +66,7 @@ void cblas_sger(const enum CBLAS_ORDER order, const int M, const int N, const fl
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_SGER_IDX, "ocl_cblas_sger");
+    __K =  ti_cblas_get_kernel("ocl_cblas_sger");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_snrm2.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_snrm2.c
index ba378a3..92c0a3a 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_snrm2.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_snrm2.c
@@ -66,7 +66,7 @@ float cblas_snrm2(const int N, const float *X, const int incX)
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_SNRM2_IDX, "ocl_cblas_snrm2");
+    __K =  ti_cblas_get_kernel("ocl_cblas_snrm2");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_srot.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_srot.c
index a7c6bbc..0a02021 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_srot.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_srot.c
@@ -66,7 +66,7 @@ void cblas_srot(const int N, float *X, const int incX, float *Y, const int incY,
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_SROT_IDX, "ocl_cblas_srot");
+    __K =  ti_cblas_get_kernel("ocl_cblas_srot");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_srotg.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_srotg.c
index f399acc..38a5d5a 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_srotg.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_srotg.c
@@ -66,7 +66,7 @@ void cblas_srotg(float *a, float *b, float *c, float *s)
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_SROTG_IDX, "ocl_cblas_srotg");
+    __K =  ti_cblas_get_kernel("ocl_cblas_srotg");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_srotm.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_srotm.c
index 845a954..c823cd2 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_srotm.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_srotm.c
@@ -66,7 +66,7 @@ void cblas_srotm(const int N, float *X, const int incX, float *Y, const int incY
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_SROTM_IDX, "ocl_cblas_srotm");
+    __K =  ti_cblas_get_kernel("ocl_cblas_srotm");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_srotmg.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_srotmg.c
index 8a03f42..68f02d6 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_srotmg.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_srotmg.c
@@ -66,7 +66,7 @@ void cblas_srotmg(float *d1, float *d2, float *b1, const float b2, float *P)
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_SROTMG_IDX, "ocl_cblas_srotmg");
+    __K =  ti_cblas_get_kernel("ocl_cblas_srotmg");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_ssbmv.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_ssbmv.c
index 24f9925..6ca7edd 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_ssbmv.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_ssbmv.c
@@ -66,7 +66,7 @@ void cblas_ssbmv(const enum CBLAS_ORDER order, const enum CBLAS_UPLO Uplo, const
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_SSBMV_IDX, "ocl_cblas_ssbmv");
+    __K =  ti_cblas_get_kernel("ocl_cblas_ssbmv");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_sscal.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_sscal.c
index df05602..7a84815 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_sscal.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_sscal.c
@@ -66,7 +66,7 @@ void cblas_sscal(const int N, const float alpha, float *X, const int incX)
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_SSCAL_IDX, "ocl_cblas_sscal");
+    __K =  ti_cblas_get_kernel("ocl_cblas_sscal");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_sspmv.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_sspmv.c
index 562ccec..41c8ef5 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_sspmv.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_sspmv.c
@@ -66,7 +66,7 @@ void cblas_sspmv(const enum CBLAS_ORDER order, const enum CBLAS_UPLO Uplo, const
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_SSPMV_IDX, "ocl_cblas_sspmv");
+    __K =  ti_cblas_get_kernel("ocl_cblas_sspmv");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_sspr.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_sspr.c
index b2e9342..bae0371 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_sspr.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_sspr.c
@@ -66,7 +66,7 @@ void cblas_sspr(const enum CBLAS_ORDER order, const enum CBLAS_UPLO Uplo, const
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_SSPR_IDX, "ocl_cblas_sspr");
+    __K =  ti_cblas_get_kernel("ocl_cblas_sspr");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_sspr2.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_sspr2.c
index 7b82e54..2156d5f 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_sspr2.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_sspr2.c
@@ -66,7 +66,7 @@ void cblas_sspr2(const enum CBLAS_ORDER order, const enum CBLAS_UPLO Uplo, const
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_SSPR2_IDX, "ocl_cblas_sspr2");
+    __K =  ti_cblas_get_kernel("ocl_cblas_sspr2");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_sswap.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_sswap.c
index d7590dd..19e5d38 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_sswap.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_sswap.c
@@ -78,7 +78,7 @@ void cblas_sswap(const int N, float *X, const int incX, float *Y, const int incY
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_SSWAP_IDX, "ocl_cblas_sswap");
+    __K =  ti_cblas_get_kernel("ocl_cblas_sswap");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_ssymm.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_ssymm.c
index bc86c8e..d62b2ec 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_ssymm.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_ssymm.c
@@ -66,7 +66,7 @@ void cblas_ssymm(const enum CBLAS_ORDER Order, const enum CBLAS_SIDE Side, const
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_SSYMM_IDX, "ocl_cblas_ssymm");
+    __K =  ti_cblas_get_kernel("ocl_cblas_ssymm");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_ssymv.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_ssymv.c
index bfe09ac..2a8d6cf 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_ssymv.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_ssymv.c
@@ -66,7 +66,7 @@ void cblas_ssymv(const enum CBLAS_ORDER order, const enum CBLAS_UPLO Uplo, const
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_SSYMV_IDX, "ocl_cblas_ssymv");
+    __K =  ti_cblas_get_kernel("ocl_cblas_ssymv");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_ssyr.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_ssyr.c
index e6c30c8..ac1b3ca 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_ssyr.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_ssyr.c
@@ -66,7 +66,7 @@ void cblas_ssyr(const enum CBLAS_ORDER order, const enum CBLAS_UPLO Uplo, const
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_SSYR_IDX, "ocl_cblas_ssyr");
+    __K =  ti_cblas_get_kernel("ocl_cblas_ssyr");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_ssyr2.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_ssyr2.c
index fa56ed3..895d1a5 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_ssyr2.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_ssyr2.c
@@ -66,7 +66,7 @@ void cblas_ssyr2(const enum CBLAS_ORDER order, const enum CBLAS_UPLO Uplo, const
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_SSYR2_IDX, "ocl_cblas_ssyr2");
+    __K =  ti_cblas_get_kernel("ocl_cblas_ssyr2");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_ssyr2k.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_ssyr2k.c
index 40d8503..66ff101 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_ssyr2k.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_ssyr2k.c
@@ -66,7 +66,7 @@ void cblas_ssyr2k(const enum CBLAS_ORDER Order, const enum CBLAS_UPLO Uplo, cons
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_SSYR2K_IDX, "ocl_cblas_ssyr2k");
+    __K =  ti_cblas_get_kernel("ocl_cblas_ssyr2k");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_ssyrk.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_ssyrk.c
index f9c0f31..c407e97 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_ssyrk.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_ssyrk.c
@@ -66,7 +66,7 @@ void cblas_ssyrk(const enum CBLAS_ORDER Order, const enum CBLAS_UPLO Uplo, const
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_SSYRK_IDX, "ocl_cblas_ssyrk");
+    __K =  ti_cblas_get_kernel("ocl_cblas_ssyrk");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_stbmv.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_stbmv.c
index 607394c..50fd50d 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_stbmv.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_stbmv.c
@@ -66,7 +66,7 @@ void cblas_stbmv(const enum CBLAS_ORDER order, const enum CBLAS_UPLO Uplo, const
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_STBMV_IDX, "ocl_cblas_stbmv");
+    __K =  ti_cblas_get_kernel("ocl_cblas_stbmv");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_stbsv.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_stbsv.c
index 96e6cab..c6d92a9 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_stbsv.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_stbsv.c
@@ -66,7 +66,7 @@ void cblas_stbsv(const enum CBLAS_ORDER order, const enum CBLAS_UPLO Uplo, const
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_STBSV_IDX, "ocl_cblas_stbsv");
+    __K =  ti_cblas_get_kernel("ocl_cblas_stbsv");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_stpmv.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_stpmv.c
index b45b877..950497c 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_stpmv.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_stpmv.c
@@ -66,7 +66,7 @@ void cblas_stpmv(const enum CBLAS_ORDER order, const enum CBLAS_UPLO Uplo, const
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_STPMV_IDX, "ocl_cblas_stpmv");
+    __K =  ti_cblas_get_kernel("ocl_cblas_stpmv");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_stpsv.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_stpsv.c
index 0681ed8..0b7aa2d 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_stpsv.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_stpsv.c
@@ -66,7 +66,7 @@ void cblas_stpsv(const enum CBLAS_ORDER order, const enum CBLAS_UPLO Uplo, const
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_STPSV_IDX, "ocl_cblas_stpsv");
+    __K =  ti_cblas_get_kernel("ocl_cblas_stpsv");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_strmm.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_strmm.c
index 3839c8e..e818b22 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_strmm.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_strmm.c
@@ -66,7 +66,7 @@ void cblas_strmm(const enum CBLAS_ORDER Order, const enum CBLAS_SIDE Side, const
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_STRMM_IDX, "ocl_cblas_strmm");
+    __K =  ti_cblas_get_kernel("ocl_cblas_strmm");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_strmv.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_strmv.c
index 40ddbc3..e2b1c0d 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_strmv.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_strmv.c
@@ -66,7 +66,7 @@ void cblas_strmv(const enum CBLAS_ORDER order, const enum CBLAS_UPLO Uplo, const
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_STRMV_IDX, "ocl_cblas_strmv");
+    __K =  ti_cblas_get_kernel("ocl_cblas_strmv");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_strsm.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_strsm.c
index 18e329e..a44e53d 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_strsm.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_strsm.c
@@ -66,7 +66,7 @@ void cblas_strsm(const enum CBLAS_ORDER Order, const enum CBLAS_SIDE Side, const
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_STRSM_IDX, "ocl_cblas_strsm");
+    __K =  ti_cblas_get_kernel("ocl_cblas_strsm");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_strsv.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_strsv.c
index 26b9a3c..42689ee 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_strsv.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_strsv.c
@@ -66,7 +66,7 @@ void cblas_strsv(const enum CBLAS_ORDER order, const enum CBLAS_UPLO Uplo, const
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_STRSV_IDX, "ocl_cblas_strsv");
+    __K =  ti_cblas_get_kernel("ocl_cblas_strsv");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_xerbla.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_xerbla.c
index 23efb90..b24f239 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_xerbla.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_xerbla.c
@@ -66,7 +66,7 @@ void cblas_xerbla(int p, const char *rout, const char *form, ...)
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_XERBLA_IDX, "ocl_cblas_xerbla");
+    __K =  ti_cblas_get_kernel("ocl_cblas_xerbla");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_zaxpy.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_zaxpy.c
index fdbd907..f17a820 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_zaxpy.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_zaxpy.c
@@ -66,7 +66,7 @@ void cblas_zaxpy(const int N, const void *alpha, const void *X, const int incX,
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_ZAXPY_IDX, "ocl_cblas_zaxpy");
+    __K =  ti_cblas_get_kernel("ocl_cblas_zaxpy");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_zcopy.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_zcopy.c
index da878b5..a070a59 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_zcopy.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_zcopy.c
@@ -66,7 +66,7 @@ void cblas_zcopy(const int N, const void *X, const int incX, void *Y, const int
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_ZCOPY_IDX, "ocl_cblas_zcopy");
+    __K =  ti_cblas_get_kernel("ocl_cblas_zcopy");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_zdotc_sub.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_zdotc_sub.c
index 4a97d0b..aa1459c 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_zdotc_sub.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_zdotc_sub.c
@@ -66,7 +66,7 @@ void cblas_zdotc_sub(const int N, const void *X, const int incX, const void *Y,
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_ZDOTC_SUB_IDX, "ocl_cblas_zdotc_sub");
+    __K =  ti_cblas_get_kernel("ocl_cblas_zdotc_sub");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_zdotu_sub.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_zdotu_sub.c
index 7d7f3e1..2f22cb1 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_zdotu_sub.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_zdotu_sub.c
@@ -66,7 +66,7 @@ void cblas_zdotu_sub(const int N, const void *X, const int incX, const void *Y,
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_ZDOTU_SUB_IDX, "ocl_cblas_zdotu_sub");
+    __K =  ti_cblas_get_kernel("ocl_cblas_zdotu_sub");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_zdscal.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_zdscal.c
index 7294884..db77056 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_zdscal.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_zdscal.c
@@ -66,7 +66,7 @@ void cblas_zdscal(const int N, const double alpha, void *X, const int incX)
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_ZDSCAL_IDX, "ocl_cblas_zdscal");
+    __K =  ti_cblas_get_kernel("ocl_cblas_zdscal");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_zgbmv.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_zgbmv.c
index 3eb8471..09ebd59 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_zgbmv.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_zgbmv.c
@@ -66,7 +66,7 @@ void cblas_zgbmv(const enum CBLAS_ORDER order, const enum CBLAS_TRANSPOSE TransA
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_ZGBMV_IDX, "ocl_cblas_zgbmv");
+    __K =  ti_cblas_get_kernel("ocl_cblas_zgbmv");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_zgemm.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_zgemm.c
index ac457bd..d4b3fcc 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_zgemm.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_zgemm.c
@@ -66,7 +66,7 @@ void cblas_zgemm(const enum CBLAS_ORDER Order, const enum CBLAS_TRANSPOSE TransA
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_ZGEMM_IDX, "ocl_cblas_zgemm");
+    __K =  ti_cblas_get_kernel("ocl_cblas_zgemm");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_zgemv.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_zgemv.c
index b4767f5..02591ee 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_zgemv.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_zgemv.c
@@ -66,7 +66,7 @@ void cblas_zgemv(const enum CBLAS_ORDER order, const enum CBLAS_TRANSPOSE TransA
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_ZGEMV_IDX, "ocl_cblas_zgemv");
+    __K =  ti_cblas_get_kernel("ocl_cblas_zgemv");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_zgerc.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_zgerc.c
index 58eb984..8153ca5 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_zgerc.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_zgerc.c
@@ -66,7 +66,7 @@ void cblas_zgerc(const enum CBLAS_ORDER order, const int M, const int N, const v
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_ZGERC_IDX, "ocl_cblas_zgerc");
+    __K =  ti_cblas_get_kernel("ocl_cblas_zgerc");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_zgeru.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_zgeru.c
index 23d4c07..7e28de2 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_zgeru.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_zgeru.c
@@ -66,7 +66,7 @@ void cblas_zgeru(const enum CBLAS_ORDER order, const int M, const int N, const v
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_ZGERU_IDX, "ocl_cblas_zgeru");
+    __K =  ti_cblas_get_kernel("ocl_cblas_zgeru");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_zhbmv.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_zhbmv.c
index 67779e4..31a34de 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_zhbmv.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_zhbmv.c
@@ -66,7 +66,7 @@ void cblas_zhbmv(const enum CBLAS_ORDER order, const enum CBLAS_UPLO Uplo, const
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_ZHBMV_IDX, "ocl_cblas_zhbmv");
+    __K =  ti_cblas_get_kernel("ocl_cblas_zhbmv");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_zhemm.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_zhemm.c
index 7e94b80..2492139 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_zhemm.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_zhemm.c
@@ -66,7 +66,7 @@ void cblas_zhemm(const enum CBLAS_ORDER Order, const enum CBLAS_SIDE Side, const
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_ZHEMM_IDX, "ocl_cblas_zhemm");
+    __K =  ti_cblas_get_kernel("ocl_cblas_zhemm");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_zhemv.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_zhemv.c
index c6a4cbe..a63a2f9 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_zhemv.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_zhemv.c
@@ -66,7 +66,7 @@ void cblas_zhemv(const enum CBLAS_ORDER order, const enum CBLAS_UPLO Uplo, const
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_ZHEMV_IDX, "ocl_cblas_zhemv");
+    __K =  ti_cblas_get_kernel("ocl_cblas_zhemv");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_zher.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_zher.c
index 3fdd79e..79f2599 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_zher.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_zher.c
@@ -66,7 +66,7 @@ void cblas_zher(const enum CBLAS_ORDER order, const enum CBLAS_UPLO Uplo, const
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_ZHER_IDX, "ocl_cblas_zher");
+    __K =  ti_cblas_get_kernel("ocl_cblas_zher");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_zher2.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_zher2.c
index 0cdf1ab..aefd5d6 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_zher2.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_zher2.c
@@ -66,7 +66,7 @@ void cblas_zher2(const enum CBLAS_ORDER order, const enum CBLAS_UPLO Uplo, const
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_ZHER2_IDX, "ocl_cblas_zher2");
+    __K =  ti_cblas_get_kernel("ocl_cblas_zher2");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_zher2k.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_zher2k.c
index e15a8ea..55aca4e 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_zher2k.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_zher2k.c
@@ -66,7 +66,7 @@ void cblas_zher2k(const enum CBLAS_ORDER Order, const enum CBLAS_UPLO Uplo, cons
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_ZHER2K_IDX, "ocl_cblas_zher2k");
+    __K =  ti_cblas_get_kernel("ocl_cblas_zher2k");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_zherk.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_zherk.c
index 5eb89b4..f3f8664 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_zherk.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_zherk.c
@@ -66,7 +66,7 @@ void cblas_zherk(const enum CBLAS_ORDER Order, const enum CBLAS_UPLO Uplo, const
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_ZHERK_IDX, "ocl_cblas_zherk");
+    __K =  ti_cblas_get_kernel("ocl_cblas_zherk");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_zhpmv.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_zhpmv.c
index bc9cdf9..d28b866 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_zhpmv.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_zhpmv.c
@@ -66,7 +66,7 @@ void cblas_zhpmv(const enum CBLAS_ORDER order, const enum CBLAS_UPLO Uplo, const
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_ZHPMV_IDX, "ocl_cblas_zhpmv");
+    __K =  ti_cblas_get_kernel("ocl_cblas_zhpmv");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_zhpr.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_zhpr.c
index 1b25b51..d02acb3 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_zhpr.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_zhpr.c
@@ -66,7 +66,7 @@ void cblas_zhpr(const enum CBLAS_ORDER order, const enum CBLAS_UPLO Uplo, const
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_ZHPR_IDX, "ocl_cblas_zhpr");
+    __K =  ti_cblas_get_kernel("ocl_cblas_zhpr");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_zhpr2.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_zhpr2.c
index a1aeaae..95e2e76 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_zhpr2.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_zhpr2.c
@@ -66,7 +66,7 @@ void cblas_zhpr2(const enum CBLAS_ORDER order, const enum CBLAS_UPLO Uplo, const
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_ZHPR2_IDX, "ocl_cblas_zhpr2");
+    __K =  ti_cblas_get_kernel("ocl_cblas_zhpr2");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_zrotg.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_zrotg.c
index f1e437c..29e2707 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_zrotg.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_zrotg.c
@@ -66,7 +66,7 @@ void cblas_zrotg( void *a, void *b, double *c, void *s)
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_ZROTG_IDX, "ocl_cblas_zrotg");
+    __K =  ti_cblas_get_kernel("ocl_cblas_zrotg");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_zscal.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_zscal.c
index 0169b07..b9912f7 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_zscal.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_zscal.c
@@ -66,7 +66,7 @@ void cblas_zscal(const int N, const void *alpha, void *X, const int incX)
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_ZSCAL_IDX, "ocl_cblas_zscal");
+    __K =  ti_cblas_get_kernel("ocl_cblas_zscal");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_zswap.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_zswap.c
index 70a29e9..f96688e 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_zswap.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_zswap.c
@@ -78,7 +78,7 @@ void cblas_zswap(const int N, void *X, const int incX, void *Y, const int incY)
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_ZSWAP_IDX, "ocl_cblas_zswap");
+    __K =  ti_cblas_get_kernel("ocl_cblas_zswap");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_zsymm.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_zsymm.c
index 9175e0e..4239501 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_zsymm.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_zsymm.c
@@ -66,7 +66,7 @@ void cblas_zsymm(const enum CBLAS_ORDER Order, const enum CBLAS_SIDE Side, const
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_ZSYMM_IDX, "ocl_cblas_zsymm");
+    __K =  ti_cblas_get_kernel("ocl_cblas_zsymm");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_zsyr2k.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_zsyr2k.c
index 53b7044..1e54fdf 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_zsyr2k.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_zsyr2k.c
@@ -66,7 +66,7 @@ void cblas_zsyr2k(const enum CBLAS_ORDER Order, const enum CBLAS_UPLO Uplo, cons
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_ZSYR2K_IDX, "ocl_cblas_zsyr2k");
+    __K =  ti_cblas_get_kernel("ocl_cblas_zsyr2k");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_zsyrk.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_zsyrk.c
index f705eb6..e75b9bd 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_zsyrk.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_zsyrk.c
@@ -66,7 +66,7 @@ void cblas_zsyrk(const enum CBLAS_ORDER Order, const enum CBLAS_UPLO Uplo, const
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_ZSYRK_IDX, "ocl_cblas_zsyrk");
+    __K =  ti_cblas_get_kernel("ocl_cblas_zsyrk");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_ztbmv.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_ztbmv.c
index c026942..063f939 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_ztbmv.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_ztbmv.c
@@ -66,7 +66,7 @@ void cblas_ztbmv(const enum CBLAS_ORDER order, const enum CBLAS_UPLO Uplo, const
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_ZTBMV_IDX, "ocl_cblas_ztbmv");
+    __K =  ti_cblas_get_kernel("ocl_cblas_ztbmv");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_ztbsv.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_ztbsv.c
index 754eca7..f9644c4 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_ztbsv.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_ztbsv.c
@@ -66,7 +66,7 @@ void cblas_ztbsv(const enum CBLAS_ORDER order, const enum CBLAS_UPLO Uplo, const
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_ZTBSV_IDX, "ocl_cblas_ztbsv");
+    __K =  ti_cblas_get_kernel("ocl_cblas_ztbsv");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_ztpmv.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_ztpmv.c
index 6fd2dd0..5361390 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_ztpmv.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_ztpmv.c
@@ -66,7 +66,7 @@ void cblas_ztpmv(const enum CBLAS_ORDER order, const enum CBLAS_UPLO Uplo, const
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_ZTPMV_IDX, "ocl_cblas_ztpmv");
+    __K =  ti_cblas_get_kernel("ocl_cblas_ztpmv");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_ztpsv.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_ztpsv.c
index dcf6100..c68987d 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_ztpsv.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_ztpsv.c
@@ -66,7 +66,7 @@ void cblas_ztpsv(const enum CBLAS_ORDER order, const enum CBLAS_UPLO Uplo, const
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_ZTPSV_IDX, "ocl_cblas_ztpsv");
+    __K =  ti_cblas_get_kernel("ocl_cblas_ztpsv");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_ztrmm.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_ztrmm.c
index c010d89..51a4a53 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_ztrmm.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_ztrmm.c
@@ -66,7 +66,7 @@ void cblas_ztrmm(const enum CBLAS_ORDER Order, const enum CBLAS_SIDE Side, const
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_ZTRMM_IDX, "ocl_cblas_ztrmm");
+    __K =  ti_cblas_get_kernel("ocl_cblas_ztrmm");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_ztrmv.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_ztrmv.c
index 1cbe9e0..244a5f9 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_ztrmv.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_ztrmv.c
@@ -66,7 +66,7 @@ void cblas_ztrmv(const enum CBLAS_ORDER order, const enum CBLAS_UPLO Uplo, const
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_ZTRMV_IDX, "ocl_cblas_ztrmv");
+    __K =  ti_cblas_get_kernel("ocl_cblas_ztrmv");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_ztrsm.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_ztrsm.c
index 4cca728..d3583e8 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_ztrsm.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_ztrsm.c
@@ -66,7 +66,7 @@ void cblas_ztrsm(const enum CBLAS_ORDER Order, const enum CBLAS_SIDE Side, const
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_ZTRSM_IDX, "ocl_cblas_ztrsm");
+    __K =  ti_cblas_get_kernel("ocl_cblas_ztrsm");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_ztrsv.c b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_ztrsv.c
index 044c64e..f7b533c 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_ztrsv.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_cblas_ztrsv.c
@@ -66,7 +66,7 @@ void cblas_ztrsv(const enum CBLAS_ORDER order, const enum CBLAS_UPLO Uplo, const
 #else
         cl_kernel __K;
 #endif
-    __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_ZTRSV_IDX, "ocl_cblas_ztrsv");
+    __K =  ti_cblas_get_kernel("ocl_cblas_ztrsv");
 
 #ifdef __cplusplus
         try 
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_initfini.c b/src/ti/linalg/blasblisacc/src/ti_cblas_initfini.c
index 65963d3..4ac5ddf 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_initfini.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_initfini.c
@@ -24,8 +24,7 @@
  *   CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE)
  *   ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF
  *   THE POSSIBILITY OF SUCH DAMAGE.
- *****************************************************************************/
-  
+ *****************************************************************************/ 
 #include "ti_cblas_acc.h"
 #include "../../ticblas/ticblas.h"
 #include <pthread.h>
@@ -34,21 +33,28 @@
 #include "ti_cblas_kernel.dsp_h"
 #endif
 
+/*==============================================================================
+ * This file contains functions of the ARM wrapper of ARM+DSP CBLAS library.
+ * It has the initialization and finalization routines. 
+ *
+ * The standard CBLAS API for each BLAS function can be found in file 
+ * ti_cblas_cblas_<func_name>.c, such as ti_cblas_cblas_dgemm.c for DGEMM. 
+ *============================================================================*/
+
 #define TI_CBLAS_INITFINI_SUCCESS   0
 #define TI_CBLAS_INITFINI_OCL_ERR   1
 #define TI_CBLAS_INITFINI_BLI_ERR   2
 
 /* Global variables */
-Context*                ti_cblas_ocl_context = NULL;
-std::vector<Device>*    ti_cblas_ocl_devices = NULL;
-CommandQueue*           ti_cblas_ocl_Q = NULL;
-Program::Binaries*      ti_cblas_ocl_binary = NULL;
-Program*                        ti_cblas_ocl_program = NULL;
+Context*             ti_cblas_ocl_context = NULL;
+std::vector<Device>* ti_cblas_ocl_devices = NULL;
+CommandQueue*        ti_cblas_ocl_Q       = NULL;
+Program::Binaries*   ti_cblas_ocl_binary  = NULL;
+Program*             ti_cblas_ocl_program = NULL;
 
 int ti_cblas_init_done = 0;       /* flag to check if init is complete */
 int ti_cblas_disable_debug = 0;   /* runtime toggle to disable debug */
 int ti_cblas_offload = TI_CBLAS_OFFLOAD_SIZE;  
-int ti_cblas_kernel_valid[TI_CBLAS_NUM_KERNELS];
 int TI_CBLAS_L1_OFFLOAD = TI_CBLAS_OFFLOAD_NONE;
 int TI_CBLAS_L2_OFFLOAD = TI_CBLAS_OFFLOAD_NONE;
 int TI_CBLAS_L3_OFFLOAD = TI_CBLAS_OFFLOAD_NONE;
@@ -56,172 +62,75 @@ int TI_CBLAS_L3_OFFLOAD = TI_CBLAS_OFFLOAD_NONE;
 pthread_cond_t   CV;
 pthread_mutex_t  MUTEX;
 
+/*============================================================================
+ * Function purpose: report error encoutered in ARM wrapper code.
+ *============================================================================*/
 void ti_cblas_error(const char* msg, int code)
 {
-  fprintf(stderr, "ERROR: (%s,%d)\n", msg, code);
+  fprintf(stderr, "TI CBLAS wrapper ERROR: (%s,%d)\n", msg, code);
 }
 
-extern "C"
+/*============================================================================
+ * Function purpose: initialize BLIS on both ARM and DSP
+ *============================================================================*/
 int ti_blis_init(void)
 {
-        int r_val = TI_CBLAS_INITFINI_SUCCESS;
+    int r_val = TI_CBLAS_INITFINI_SUCCESS;
     
-        TI_CBLAS_DEBUG_PRINT("Initializing BLIS on ARM...\n");   
     /* Initialize BLIS on ARM */
-        bli_init();    
-        TI_CBLAS_DEBUG_PRINT("BLIS initialized on ARM.\n");
-
-    /* Initialize BLIS on DSP */
-        TI_CBLAS_DEBUG_PRINT("Initializing BLIS on DSP...\n");   
-        Event e;
-        Kernel* __K;
-
-        __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_CGEMM_IDX, "ocl_bli_init");
-        try
-        {
-                TI_CBLAS_DEBUG_PRINT("Initializing BLIS on DSP...\n");
-        
-                int err_code;           
-                Buffer buf_err(*ti_cblas_ocl_context, CL_MEM_READ_WRITE|CL_MEM_USE_HOST_PTR, sizeof(int), &err_code);
-                __K->setArg(0, buf_err);
-
-                ti_cblas_ocl_Q->enqueueTask(*__K, 0, &e);
-                e.wait();
-
-        if(err_code != TICBLAS_SUCCESS) {
-          TI_CBLAS_DEBUG_PRINT("Error in offloaded ocl_bli_init with error code %d!\n", err_code);
-          r_val = TI_CBLAS_INITFINI_BLI_ERR;
-        }
-        
-                ti_cblas_delete_kernel(__K);
-                TI_CBLAS_DEBUG_PRINT("BLIS DSP initialization finished.\n");        
-        }
-
-        catch (Error err)
-        {
-                ti_cblas_delete_kernel(__K);
-                ti_cblas_error(err.what(),err.err());
-                r_val = TI_CBLAS_INITFINI_OCL_ERR;
-        }
-
-    return r_val;
-}
-
-extern "C"
-int ti_blis_finalize(void)
-{
-        int r_val = TI_CBLAS_INITFINI_SUCCESS;
-    
-        TI_CBLAS_DEBUG_PRINT("Finalizing BLIS on ARM...\n");
-    /* Finalize BLIS on ARM */
-        bli_finalize();
-        TI_CBLAS_DEBUG_PRINT("BLIS finalized on ARM.\n");
+    TI_CBLAS_DEBUG_PRINT("Initializing BLIS on ARM...\n");   
+    bli_init();    
+    TI_CBLAS_DEBUG_PRINT("BLIS initialized on ARM.\n");
 
-        Event e;
-        Kernel* __K;
+    /* Initialize BLIS on DSP by offloading bli_init() on DSP */
+    TI_CBLAS_DEBUG_PRINT("Initializing BLIS on DSP...\n");   
+    Event e;
+    Kernel* __K;
 
-        __K =  ti_cblas_get_kernel(TI_CBLAS_CBLAS_CGEMM_IDX, "ocl_bli_finalize");
+    __K =  ti_cblas_get_kernel("ocl_bli_init");
+    try
+    {
+        int err_code;        
+        Buffer buf_err(*ti_cblas_ocl_context, CL_MEM_READ_WRITE|CL_MEM_USE_HOST_PTR, sizeof(int), &err_code);
+        __K->setArg(0, buf_err);
 
-    /* Finalize BLIS on DSP */
-        TI_CBLAS_DEBUG_PRINT("Finalizing BLIS on DSP...\n");
-    int err_code;               
-        Buffer buf_err(*ti_cblas_ocl_context, CL_MEM_READ_WRITE|CL_MEM_USE_HOST_PTR, sizeof(int), &err_code);
-        __K->setArg(0, buf_err);
-    
-        try
-        {
-                ti_cblas_ocl_Q->enqueueTask(*__K, 0, &e);
-                e.wait();
+        ti_cblas_ocl_Q->enqueueTask(*__K, 0, &e);
+        e.wait();
 
         if(err_code != TICBLAS_SUCCESS) {
-          TI_CBLAS_DEBUG_PRINT("Error in offloaded ocl_bli_finalize with error code %d!\n", err_code);
+          TI_CBLAS_DEBUG_PRINT("Error in offloaded ocl_bli_init with error code %d!\n", err_code);
           r_val = TI_CBLAS_INITFINI_BLI_ERR;
         }
         
-                ti_cblas_delete_kernel(__K);
-        }
-
-        catch (Error err)
-        {
-                ti_cblas_error(err.what(),err.err());
-                r_val = TI_CBLAS_INITFINI_OCL_ERR;
-        }
-
-        return r_val;
-}
-
-#ifdef __cplusplus
-extern "C"
-#endif
-int ti_cblas_finalize(void)
-{
-        /* If ti_cblas_init_done is equal to 0,
-         * then we know that ti_cblas_init was not called,
-         * and so we can return early.
-         */
-        if(ti_cblas_init_done == 0) {
-                return TI_CBLAS_INITFINI_SUCCESS;
+        ti_cblas_delete_kernel(__K);
+        TI_CBLAS_DEBUG_PRINT("BLIS DSP initialization finished.\n");        
     }
 
-        int r_val = ti_blis_finalize();
-        /*Using same name as ti_cblas_init critical region. See notes in bli_init*/
-#pragma omp critical (ti_cblas_init_critical)
-        {
-                // Destroy Pthread
-                pthread_mutex_destroy(&MUTEX);
-                pthread_cond_destroy (&CV);
-
-                //destroy Command queue, program, devices and context.
-                if(ti_cblas_ocl_Q != NULL)
-                {
-                        delete(ti_cblas_ocl_Q);
-                        ti_cblas_ocl_Q = NULL;
-                }
-                if(ti_cblas_ocl_program != NULL)
-                {
-                        delete(ti_cblas_ocl_program);
-                        ti_cblas_ocl_program = NULL;
-                }
-                if(ti_cblas_ocl_binary != NULL)
-                {
-                        delete(ti_cblas_ocl_binary);
-                        ti_cblas_ocl_binary = NULL;
-                }
-                if(ti_cblas_ocl_devices != NULL)
-                {
-                        delete(ti_cblas_ocl_devices);
-                        ti_cblas_ocl_devices = NULL;
-                }
-                if(ti_cblas_ocl_context != NULL)
-                {
-                        delete(ti_cblas_ocl_context);
-                        ti_cblas_ocl_context = NULL;
-                }
-        }
-    
-        return r_val;
-}
-
-void ti_cblas_auto_finalize(void)
-{
-    int r_val;
-
-    r_val = ti_cblas_finalize();
-    if (r_val != TI_CBLAS_INITFINI_SUCCESS)
+    catch (Error err)
     {
-        fprintf(stderr, "Error: ti_cblas_finalize failed with error code %d!\n", r_val);
+        ti_cblas_delete_kernel(__K);
+        ti_cblas_error(err.what(),err.err());
+        r_val = TI_CBLAS_INITFINI_OCL_ERR;
     }
-}
 
-/* This function is invoked exactly once on startup */
-/* Its purpose is to parse the environment variables and do OpenCL init */
+    return r_val;
+} // ti_blis_init 
+
+
+/*============================================================================
+ * Function purpose: initialize and prepare for CBLAS calls:
+ *    - parse the environment variables 
+ *    - initialize OpenCL 
+ *    - initialize BLIS
+ *
+ * Note: this function is invoked exactly once on startup, when any CBLAS function
+ *       is called the first time. 
+ *============================================================================*/
 void ti_cblas_init(void)
 {
 #pragma omp critical (ti_cblas_init_critical)
   {
     /* Add code for interception */
-    if (!ti_cblas_init_done) 
-    {
 #ifdef TI_CBLAS_DEBUG
       char *no_debug_env = getenv("TI_CBLAS_NO_DEBUG");
       if (no_debug_env) {
@@ -238,8 +147,8 @@ void ti_cblas_init(void)
       /* check environment variables */
       const char *offload_env = getenv("TI_CBLAS_OFFLOAD");
       if (!offload_env)  { 
-        TI_CBLAS_DEBUG_PRINT("Using build time default for offload: TI_CBLAS_OFFLOAD=%s\n", TI_CBLAS_OFFLOAD);
-        offload_env = TI_CBLAS_OFFLOAD;
+        TI_CBLAS_DEBUG_PRINT("Using build time default for offload: TI_CBLAS_OFFLOAD=%s\n", TI_CBLAS_OFFLOAD_DEF);
+        offload_env = TI_CBLAS_OFFLOAD_DEF;
       }
       else {
         TI_CBLAS_DEBUG_PRINT("Using runtime override for offloads: TI_CBLAS_OFFLOAD=%s\n", offload_env);
@@ -280,166 +189,224 @@ void ti_cblas_init(void)
 #else
       const char binary[] = "./ti_cblas_kernel.out";
       unsigned int bin_length;
-#ifdef __cplusplus
       bin_length = ocl_read_binary(binary, (char*&)bin); 
-#else
-      FILE *fp = fopen(binary, "r");
-      if (!fp) {
-        TI_CBLAS_ERROR_EXIT("Could not open OpenCL pre-compiled binary %s for reading\n", binary);
-      }
-      struct stat fileinfo;
-      stat(binary, &fileinfo);
-      bin_length = fileinfo.st_size;
-      bin = (char *)malloc(bin_length);
-      if (!bin) {
-        TI_CBLAS_ERROR_EXIT("Could not malloc of size %d for reading OpenCL binary\n", bin_length);
-      }
-      if (fread((char *)bin, bin_length, 1, fp) != 1) {
-        TI_CBLAS_ERROR_EXIT("Could not read %d bytes of OpenCL binary\n", bin_length);
-      }
-      fclose(fp);
-#endif /* cplusplus */
 #endif /* FAT_BINARY */
     
       /* OpenCL init */
       TI_CBLAS_DEBUG_PRINT("Initializing OpenCL\n");
-#ifdef __cplusplus
-          ti_cblas_ocl_context = new Context(CL_DEVICE_TYPE_ACCELERATOR);
-          ti_cblas_ocl_devices = new std::vector<Device> (ti_cblas_ocl_context->getInfo<CL_CONTEXT_DEVICES>());
-          ti_cblas_ocl_binary = new Program::Binaries(1, std::make_pair(bin, bin_length));
-          ti_cblas_ocl_program = new Program(*ti_cblas_ocl_context, *ti_cblas_ocl_devices, *ti_cblas_ocl_binary);
-          ti_cblas_ocl_program->build(*ti_cblas_ocl_devices);
-          ti_cblas_ocl_Q = new CommandQueue(*ti_cblas_ocl_context, ti_cblas_ocl_devices[0][0], CL_QUEUE_PROFILING_ENABLE);
-#else
-      cl_int err;
-      cl_device_id device;
-      /* Create an in-order command queue by default*/
-      int queue_flags = 0;
-#ifdef TI_CBLAS_PROFILE
-      queue_flags |= CL_QUEUE_PROFILING_ENABLE;
-#endif
-    
-      ti_cblas_ocl_context = clCreateContextFromType(0,CL_DEVICE_TYPE_ACCELERATOR,0,0,&err);
-      TI_CBLAS_OCL_CHKERROR("clCreateContextFromType",err);
-      err = clGetDeviceIDs(0,CL_DEVICE_TYPE_ACCELERATOR,1,&device,0);
-      TI_CBLAS_OCL_CHKERROR("clGetDeviceIDs",err);
-      ti_cblas_ocl_Q = clCreateCommandQueue(ti_cblas_ocl_context, device, queue_flags, &err);
-      TI_CBLAS_OCL_CHKERROR("clCreateCommandQueue",err);
-      ti_cblas_ocl_program = clCreateProgramWithBinary(ti_cblas_ocl_context, 1, &device, &bin_length, &bin, NULL, &err);
-      TI_CBLAS_OCL_CHKERROR("clCreateProgramWithBinary",err);
-      const char *compile_options = "";
-      err = clBuildProgram(ti_cblas_ocl_program, 1, &device, compile_options, 0, 0);
-      TI_CBLAS_OCL_CHKERROR("clBuildProgram",err);
-    
-#endif
+      ti_cblas_ocl_context = new Context(CL_DEVICE_TYPE_ACCELERATOR);
+      ti_cblas_ocl_devices = new std::vector<Device> (ti_cblas_ocl_context->getInfo<CL_CONTEXT_DEVICES>());
+      ti_cblas_ocl_binary = new Program::Binaries(1, std::make_pair(bin, bin_length));
+      ti_cblas_ocl_program = new Program(*ti_cblas_ocl_context, *ti_cblas_ocl_devices, *ti_cblas_ocl_binary);
+      ti_cblas_ocl_program->build(*ti_cblas_ocl_devices);
+      ti_cblas_ocl_Q = new CommandQueue(*ti_cblas_ocl_context, ti_cblas_ocl_devices[0][0], CL_QUEUE_PROFILING_ENABLE);
     
 #ifndef TI_CBLAS_FAT_BINARY
-#ifdef __cplusplus
       delete [] bin;
-#else
-      free((char*)bin);
-#endif
 #endif /* FAT_BINARY */
+      
       TI_CBLAS_DEBUG_PRINT("OpenCL initialized\n");
 
+      /* Initializing pthreads */
       TI_CBLAS_DEBUG_PRINT("Initializing Pthreads\n");
-          /* Initializing pthreads */
-          pthread_cond_init (&CV, 0);
-          pthread_mutex_init(&MUTEX, 0);
-          TI_CBLAS_DEBUG_PRINT("Pthreads initialized\n");
-          TI_CBLAS_DEBUG_PRINT("Initializing BLIS\n");
-          if(ti_blis_init() == TI_CBLAS_INITFINI_SUCCESS) {
+      pthread_cond_init (&CV, 0);
+      pthread_mutex_init(&MUTEX, 0);
+      TI_CBLAS_DEBUG_PRINT("Pthreads initialized\n");
+      TI_CBLAS_DEBUG_PRINT("Initializing BLIS\n");
+      if(ti_blis_init() == TI_CBLAS_INITFINI_SUCCESS) {
         TI_CBLAS_DEBUG_PRINT("BLIS initialized\n");\
       }
       else {
-            TI_CBLAS_DEBUG_PRINT("BLIS NOT initialized!\n");\
+        TI_CBLAS_DEBUG_PRINT("BLIS NOT initialized!\n");\
       }
-        
-          atexit(ti_cblas_auto_finalize);
+      
+      /* Register auto finalization to be called when program exits */
+      atexit(ti_cblas_auto_finalize);
 
-          TI_CBLAS_PROFILE_REPORT("  Initialization took %8.2f us\n", (float) clock_diff);
-          ti_cblas_init_done = 1;
-          TI_CBLAS_DEBUG_PRINT("ti_cblas_init: Finished OpenCL initialization\n");
-        } //end of !ti_cblas_init_done
+      TI_CBLAS_PROFILE_REPORT("Initialization took %8.2f us\n", (float) clock_diff);
+      ti_cblas_init_done = 1;
+      TI_CBLAS_DEBUG_PRINT("ti_cblas_init: Finished initialization\n");
       
   } // End of critical section
 
   return;
-}
+} //ti_cblas_init
 
+/*============================================================================
+ * Function purpose: finalize BLIS on both ARM and DSP
+ *============================================================================*/
+int ti_blis_finalize(void)
+{
+    int r_val = TI_CBLAS_INITFINI_SUCCESS;
+    
+    /* Finalize BLIS on ARM */
+    TI_CBLAS_DEBUG_PRINT("Finalizing BLIS on ARM...\n");
+    bli_finalize();
+    TI_CBLAS_DEBUG_PRINT("BLIS finalized on ARM.\n");
 
+    /* Finalize BLIS on DSP */
+    Event e;
+    Kernel* __K;
 
-void ti_cblas_mem_free(void *ptr)
+    __K =  ti_cblas_get_kernel("ocl_bli_finalize");
+
+    TI_CBLAS_DEBUG_PRINT("Finalizing BLIS on DSP...\n");
+    int err_code;        
+    Buffer buf_err(*ti_cblas_ocl_context, CL_MEM_READ_WRITE|CL_MEM_USE_HOST_PTR, sizeof(int), &err_code);
+    __K->setArg(0, buf_err);
+    
+    try
+    {
+        ti_cblas_ocl_Q->enqueueTask(*__K, 0, &e);
+        e.wait();
+
+        if(err_code != TICBLAS_SUCCESS) {
+          TI_CBLAS_DEBUG_PRINT("Error in offloaded ocl_bli_finalize with error code %d!\n", err_code);
+          r_val = TI_CBLAS_INITFINI_BLI_ERR;
+        }
+        
+        ti_cblas_delete_kernel(__K);
+    }
+
+    catch (Error err)
+    {
+        ti_cblas_error(err.what(),err.err());
+        r_val = TI_CBLAS_INITFINI_OCL_ERR;
+    }
+
+    return r_val;
+} // ti_blis_finalize 
+
+/*============================================================================
+ * Function purpose: finalize after all CBLAS calls:
+ *    - finalize BLIS
+ *    - delete OpenCL context
+ *
+ * Note: this function is invoked exactly once on program exit.
+ *============================================================================*/
+int ti_cblas_finalize(void)
 {
-        pthread_mutex_lock(&MUTEX);
-        __free_msmc(ptr);
-        pthread_cond_broadcast(&CV);
-        pthread_mutex_unlock(&MUTEX);
+    /* If ti_cblas_init_done is equal to 0,
+     * then we know that ti_cblas_init was not called,
+     * and so we can return early.
+     */
+    if(ti_cblas_init_done == 0) {
+        return TI_CBLAS_INITFINI_SUCCESS;
+    }
 
-}
+    int r_val = ti_blis_finalize();
+    /*Using same name as ti_cblas_init critical region. See notes in bli_init*/
+#pragma omp critical (ti_cblas_init_critical)
+    {
+        // Destroy Pthread
+        pthread_mutex_destroy(&MUTEX);
+        pthread_cond_destroy (&CV);
+
+        //destroy Command queue, program, devices and context.
+        if(ti_cblas_ocl_Q != NULL)
+        {
+            delete(ti_cblas_ocl_Q);
+            ti_cblas_ocl_Q = NULL;
+        }
+        if(ti_cblas_ocl_program != NULL)
+        {
+            delete(ti_cblas_ocl_program);
+            ti_cblas_ocl_program = NULL;
+        }
+        if(ti_cblas_ocl_binary != NULL)
+        {
+            delete(ti_cblas_ocl_binary);
+            ti_cblas_ocl_binary = NULL;
+        }
+        if(ti_cblas_ocl_devices != NULL)
+        {
+            delete(ti_cblas_ocl_devices);
+            ti_cblas_ocl_devices = NULL;
+        }
+        if(ti_cblas_ocl_context != NULL)
+        {
+            delete(ti_cblas_ocl_context);
+            ti_cblas_ocl_context = NULL;
+        }
+    }
+    
+    return r_val;
+} // ti_cblas_finalize
 
-void *ti_cblas_mem_alloc(size_t size)
+
+/*============================================================================
+ * Function purpose: auto-finalize on program exit.
+ *============================================================================*/
+void ti_cblas_auto_finalize(void)
 {
-        void *ptr;
-        pthread_mutex_lock(&MUTEX);
-        /*-------------------------------------------------------------------------
+    int r_val;
 
-         * Loop in case of false signal after broadcast.
+    r_val = ti_cblas_finalize();
+    if (r_val != TI_CBLAS_INITFINI_SUCCESS)
+    {
+        fprintf(stderr, "Error: ti_cblas_finalize failed with error code %d!\n", r_val);
+    }
+} //ti_cblas_auto_finalize
 
-         *------------------------------------------------------------------------*/
-        while ((ptr = __malloc_msmc(size)) == 0)
 
-        pthread_cond_wait(&CV, &MUTEX);
-        pthread_mutex_unlock(&MUTEX);
-        return ptr;
+/*============================================================================
+ * Function purpose: free previously allocated MSMC memory
+ *============================================================================*/
+void ti_cblas_mem_free(void *ptr)
+{
+    pthread_mutex_lock(&MUTEX);
+    __free_msmc(ptr);
+    pthread_cond_broadcast(&CV);
+    pthread_mutex_unlock(&MUTEX);
 
 }
 
+/*============================================================================
+ * Function purpose: allocate MSMC memory
+ *============================================================================*/
+void *ti_cblas_mem_alloc(size_t size)
+{
+    void *ptr;
+    pthread_mutex_lock(&MUTEX);
+    
+    /*-------------------------------------------------------------------------
+     * Loop in case of false signal after broadcast.
+     *------------------------------------------------------------------------*/
+    while ((ptr = __malloc_msmc(size)) == 0)
+    {
+        pthread_cond_wait(&CV, &MUTEX);
+    }
 
-/* Returns a handle to the kernel for the specified
- * function with index 'idx'. Initializes the handle if it's
- * not been used before, otherwise returns earlier handle
- */
-#ifdef __cplusplus
-Kernel*
-#else
-cl_kernel
-#endif
-ti_cblas_get_kernel(int idx, const char *fname)
+    pthread_mutex_unlock(&MUTEX);
+    
+    return ptr;
+} //ti_cblas_mem_alloc
+
+
+/*============================================================================
+ * Function purpose: create an OpenCL kernel
+ *============================================================================*/
+Kernel *ti_cblas_get_kernel(const char *fname)
 {
-#ifdef __cplusplus
-        Kernel* __K;
-#else
-        cl_kernel __K;
-#endif
+    Kernel* __K;
     
-#ifdef __cplusplus
     __K = new Kernel(*ti_cblas_ocl_program, fname);
-#else
-    cl_int err;
-    __K = clCreateKernel(ti_cblas_ocl_program,fname,&err);
-    TI_CBLAS_OCL_CHKERROR("clCreateKernel",err);
-#endif
 
     return __K;
 }
 
-#ifdef __cplusplus
+/*============================================================================
+ * Function purpose: delete an OpenCL kernel
+ *============================================================================*/
 int ti_cblas_delete_kernel(Kernel* K)
-#else
-int ti_cblas_delete_kernel(cl_kernel K)
-#endif
 {
-#ifdef __cplusplus
-        if(K != NULL)
-        {
-                delete(K);
-                K=NULL;
-        }
-#else
-        clReleaseKernel(K);
-#endif
-        return 0;
+    if(K != NULL)
+    {
+        delete(K);
+        K=NULL;
+    }
+    
+    return 0;
 }
 
-
+/* Nothing after this line */
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_kernel.cl b/src/ti/linalg/blasblisacc/src/ti_cblas_kernel.cl
index 8fa16d0..b4368eb 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_kernel.cl
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_kernel.cl
@@ -25,6 +25,9 @@
  *   ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF
  *   THE POSSIBILITY OF SUCH DAMAGE.
  *****************************************************************************/
+/*==============================================================================
+ * This file contains the OpenCL kernels of ARM+DSP CBLAS library.
+ *============================================================================*/
 
 #define CBLAS_H
 #define CBLAS_INDEX size_t  /* this may vary between platforms */
@@ -34,16 +37,16 @@ enum CBLAS_UPLO {CblasUpper=121, CblasLower=122};
 enum CBLAS_DIAG {CblasNonUnit=131, CblasUnit=132};
 enum CBLAS_SIDE {CblasLeft=141, CblasRight=142};
 
-
-int ti_bli_init_dsp(void);
+int tiCblasNew(void);
 kernel void ocl_bli_init(global int *err_code)
 {
-  *err_code = ti_bli_init_dsp(); 
+  *err_code = tiCblasNew();
 }
-int ti_bli_finalize_dsp(void);
+
+int tiCblasDelete(void);
 kernel void ocl_bli_finalize(global int *err_code)
 {
-  *err_code = ti_bli_finalize_dsp(); 
+  *err_code = tiCblasDelete();
 }
 
 void cblas_caxpy_facade(const int N, global const void *alpha, global const void *X, const int incX, global void *Y, const int incY);
diff --git a/src/ti/linalg/blasblisacc/src/ti_cblas_mem_config.c b/src/ti/linalg/blasblisacc/src/ti_cblas_mem_config.c
index a506ece..54d6e27 100644
--- a/src/ti/linalg/blasblisacc/src/ti_cblas_mem_config.c
+++ b/src/ti/linalg/blasblisacc/src/ti_cblas_mem_config.c
@@ -33,6 +33,13 @@
 #include "../../ticblas/ticblas.h"
 #include <ti/libarch/libarch.h>
 
+/*==============================================================================
+ * This file contains functions of the DSP OpenCL layer of ARM+DSP CBLAS library.
+ *============================================================================*/
+
+#define TICBLAS_ERROR_MEMCFG   (-1)  /* L1D/L2 config error.           */
+#define TICBLAS_ERROR_MEMRECFG (-2)  /* /L2 reconfig error.            */
+
 extern void bli_init();
 extern void bli_finalize();
 
@@ -41,23 +48,39 @@ int malloc_size;
 extern lib_memdscr_t * blas_memdscr_tab[4];
 #endif
 
-int bli_l3_mem_config(void *msmc_buf, size_t msmc_buf_size, void *ddr_buf, size_t ddr_buf_size, size_t *l1D_SRAM_size_orig, size_t *l2_SRAM_size_orig)
+/*==============================================================================
+ * This function configures L1D and L2 and initializes heap for BLIS.
+ *
+ * Input:
+ *   msmc_buf              - base address of MSMC/L3 buffer
+ *   msmc_buf_size         - size of MSMC/L3 buffer
+ *   ddr_buf               - base address of DDR buffer
+ *   ddr_buf_size          - size of DDR buffer
+ * Output:
+ *   l1D_SRAM_size_orig    - original L1D SRAM size
+ *   l2_SRAM_size_orig     - original L2 SRAM size
+ *============================================================================*/
+int bli_l3_mem_config(void *msmc_buf, size_t msmc_buf_size, 
+                      void *ddr_buf, size_t ddr_buf_size, 
+                      size_t *l1D_SRAM_size_orig, size_t *l2_SRAM_size_orig)
 {
     size_t smem_size_vfast, smem_size_fast, smem_size_med, smem_size_slow;
     void *l1d_SRAM_ptr, *l2_SRAM_ptr;
     int l1d_cfg_err, l2_cfg_err, blas_ret_err_code;
 
 #ifdef TI_CBLAS_DEBUG 
-        malloc_size = 0;
-        printf("Memory buffers passed to bli_l3_mem_config are: MSMC base 0x%x, size %d, DDR base 0x%x, size%d.\n", (unsigned int)msmc_buf, msmc_buf_size, (unsigned int) ddr_buf, ddr_buf_size);
-        printf("Before calling BLIS, malloc_size is %d.\n", malloc_size);
+    malloc_size = 0;
+    printf("Memory buffers passed to bli_l3_mem_config are: MSMC base 0x%x, size %d, DDR base 0x%x, size%d.\n", 
+           (unsigned int)msmc_buf, msmc_buf_size, (unsigned int) ddr_buf, ddr_buf_size);
+    printf("Before calling BLIS, malloc_size is %d.\n", malloc_size);
 #endif
-        
-    /* First, verify the provided/available memory meet requirements */
+    
+    /* First, verify the provided and available memory meet requirements */
     tiCblasGetSizes(&smem_size_vfast, &smem_size_fast, &smem_size_med, &smem_size_slow);
     
 #ifdef TI_CBLAS_DEBUG 
-        printf("Very fast mem size is %d, fast mem size is %d, medium mem size is %d, slow mem size is %d.\n", smem_size_vfast, smem_size_fast, smem_size_med, smem_size_slow);
+    printf("Very fast mem size is %d, fast mem size is %d, medium mem size is %d, slow mem size is %d.\n", 
+           smem_size_vfast, smem_size_fast, smem_size_med, smem_size_slow);
     printf("Total L1D size is: %d\n", lib_get_L1D_total_size());
     printf("Total L2 size is: %d\n", lib_get_L2_total_size());
 #endif
@@ -83,7 +106,6 @@ int bli_l3_mem_config(void *msmc_buf, size_t msmc_buf_size, void *ddr_buf, size_
 #endif
  
     if(*l1D_SRAM_size_orig < smem_size_vfast) {    /* configure L1D if needs more SRAM */
-      /*printf("Configuring L1D SRAM on all cores.\n");*/
       #pragma omp parallel                         
       {
         l1d_cfg_err = lib_L1D_config_SRAM(smem_size_vfast);  
@@ -98,13 +120,6 @@ int bli_l3_mem_config(void *msmc_buf, size_t msmc_buf_size, void *ddr_buf, size_
       }
     }  
  
-#ifdef TI_CBLAS_DEBUG 
-    #pragma omp parallel                         
-    {
-      int core_id = lib_get_coreID();
-    }
-#endif
- 
     /* Configure L2 if necessary */
     *l2_SRAM_size_orig  = lib_get_L2_SRAM_size();   /* get current L2 SRAM size */
     l2_cfg_err = LIB_CACHE_SUCCESS;
@@ -118,7 +133,6 @@ int bli_l3_mem_config(void *msmc_buf, size_t msmc_buf_size, void *ddr_buf, size_
       #pragma omp parallel
       { 
         l2_cfg_err  = lib_L2_config_SRAM(smem_size_fast);    
-      }
 #ifdef TI_CBLAS_DEBUG 
         if(l2_cfg_err != LIB_CACHE_SUCCESS) {
           printf("Error in configuring L2 on core %d!\n", lib_get_coreID());
@@ -127,6 +141,7 @@ int bli_l3_mem_config(void *msmc_buf, size_t msmc_buf_size, void *ddr_buf, size_
           printf("On core %d, new L2 SRAM size is %d.\n", lib_get_coreID(), lib_get_L2_SRAM_size());  
         }
 #endif
+      }
     }
     
     if(l1d_cfg_err || l2_cfg_err) {
@@ -137,8 +152,8 @@ int bli_l3_mem_config(void *msmc_buf, size_t msmc_buf_size, void *ddr_buf, size_
     }
 
 #ifdef TI_CBLAS_DEBUG 
- printf("New L1D SRAM size is: %d\n", lib_get_L1D_SRAM_size());
- printf("New L2 SRAM size is: %d\n", lib_get_L2_SRAM_size());
+    printf("New L1D SRAM size is: %d\n", lib_get_L1D_SRAM_size());
+    printf("New L2 SRAM size is: %d\n", lib_get_L2_SRAM_size());
 #endif
  
     /* get L1D and L2 SRAM base address */
@@ -158,15 +173,16 @@ int bli_l3_mem_config(void *msmc_buf, size_t msmc_buf_size, void *ddr_buf, size_
                                     ddr_buf,       ddr_buf_size);
 
 #ifdef TI_CBLAS_DEBUG 
-        if(blas_ret_err_code == TICBLAS_SUCCESS) {
-          printf("Before calling BLIS, memory descriptor base is 0x%x, used is %d.\n", blas_memdscr_tab[3]->base, blas_memdscr_tab[3]->used);
-        }
-        else {
-          printf("BLAS init error with code %d.\n   ", blas_ret_err_code);
-        }
+    if(blas_ret_err_code == TICBLAS_SUCCESS) {
+      printf("Before calling BLIS, memory descriptor base is 0x%x, used is %d.\n", 
+             blas_memdscr_tab[3]->base, blas_memdscr_tab[3]->used);
+    }
+    else {
+      printf("BLAS init error with code %d.\n   ", blas_ret_err_code);
+    }
 #endif
-        
-        return(blas_ret_err_code);
+    
+    return(blas_ret_err_code);
 } /* bli_l3_mem_config */
 
 /*==============================================================================
@@ -177,11 +193,13 @@ int bli_l3_mem_reconfig(size_t l1D_SRAM_size_orig, size_t l2_SRAM_size_orig)
     int l1d_cfg_err, l2_cfg_err;
 
 #ifdef TI_CBLAS_DEBUG 
-        printf("After calling BLIS, malloc_size is %d.\n", malloc_size);
-        printf("After calling BLIS, used_size in memory descriptor is %d.\n", blas_memdscr_tab[3]->used);       
-        printf("Configuring L1D SRAM and L2 SRAM back to %d and %d.\n", l1D_SRAM_size_orig, l2_SRAM_size_orig); 
+    printf("After calling BLIS, malloc_size is %d.\n", malloc_size);
+    printf("After calling BLIS, used_size in memory descriptor is %d.\n", 
+           blas_memdscr_tab[3]->used);
+    printf("Configuring L1D SRAM and L2 SRAM back to %d and %d.\n", 
+           l1D_SRAM_size_orig, l2_SRAM_size_orig);
 #endif
-        
+    
     /* configure L1D back if necessary */    
     l1d_cfg_err = LIB_CACHE_SUCCESS;
     if(l1D_SRAM_size_orig!=lib_get_L1D_SRAM_size()) {
@@ -214,20 +232,4 @@ int bli_l3_mem_reconfig(size_t l1D_SRAM_size_orig, size_t l2_SRAM_size_orig)
     return(TICBLAS_SUCCESS);            
 } /* bli_l3_mem_reconfig */
 
-/*==============================================================================
- * This function initializes BLIS before first CBLAS call is made. 
- *============================================================================*/
-int ti_bli_init_dsp(void)
-{ 
-  return tiCblasNew();
-}
-
-/*==============================================================================
- * This function frees all memories allocated by ti_bli_init_dsp. 
- *============================================================================*/
-int ti_bli_finalize_dsp(void)
-{
-  return tiCblasDelete();
-}
-
 /* Nothing after this line */
diff --git a/src/ti/linalg/blasblisacc/src/wrap_gen/oclgen.pl b/src/ti/linalg/blasblisacc/src/wrap_gen/oclgen.pl
index 77a49a4..7c5fe27 100755
--- a/src/ti/linalg/blasblisacc/src/wrap_gen/oclgen.pl
+++ b/src/ti/linalg/blasblisacc/src/wrap_gen/oclgen.pl
@@ -395,7 +395,7 @@ Kernel*
 #else
 cl_kernel
 #endif
-${namespace}_get_kernel(int idx, const char *fname)
+${namespace}_get_kernel(const char *fname)
 {
    if (!${namespace}_kernel_valid[idx]) {
 #ifdef __cplusplus
@@ -2204,7 +2204,7 @@ ARM_FROM_PROTO
 #else
         cl_kernel __K;
 #endif
-    __K =  ${namespace}_get_kernel($trampdef, \"ocl_$trampname\");
+    __K =  ${namespace}_get_kernel(\"ocl_$trampname\");
 
 #ifdef __cplusplus
         try 
@@ -2771,7 +2771,7 @@ $::dep
 extern void ${namespace}_error(const char* msg, int code);
 extern void ${namespace}_init(void);
 #ifdef __cplusplus
-extern Kernel*                  ${namespace}_get_kernel(int idx, const char *fname);
+extern Kernel*                  ${namespace}_get_kernel(const char *fname);
 extern Context                  ${namespace}_ocl_context;
 extern std::vector<Device>      ${namespace}_ocl_devices;
 extern CommandQueue             ${namespace}_ocl_Q;
@@ -2779,7 +2779,7 @@ extern Program::Binaries	${namespace}_ocl_binary;
 extern Program                  ${namespace}_ocl_program;
 extern Kernel*                  ${namespace}_ocl_kernels[];
 #else
-extern cl_kernel                ${namespace}_get_kernel(int idx, const char *fname);
+extern cl_kernel                ${namespace}_get_kernelconst char *fname);
 extern cl_context               ${namespace}_ocl_context;
 extern cl_command_queue         ${namespace}_ocl_Q;
 extern cl_program               ${namespace}_ocl_program;
diff --git a/src/ti/linalg/blis/frame/base/bli_mem.c b/src/ti/linalg/blis/frame/base/bli_mem.c
index 4303217..275ff4f 100644
--- a/src/ti/linalg/blis/frame/base/bli_mem.c
+++ b/src/ti/linalg/blis/frame/base/bli_mem.c
@@ -72,37 +72,37 @@ static char   pool_mn_mem[ BLIS_MN_POOL_SIZE ];
 
 //L1
 static void*  pool_mk_blk_ptrs_L1[ BLIS_NUM_MR_X_KC_BLOCKS_L1 ];
-extern char   *pool_mk_mem_L1;
+char   *pool_mk_mem_L1;
 
 static void*  pool_kn_blk_ptrs_L1[ BLIS_NUM_KC_X_NR_BLOCKS_L1 ];
-extern char   *pool_kn_mem_L1;
+char   *pool_kn_mem_L1;
 
 static void*  pool_mn_blk_ptrs_L1[ BLIS_NUM_MR_X_NR_BLOCKS_L1 ];
-extern char   *pool_mn_mem_L1;
+char   *pool_mn_mem_L1;
 
 //
 //L2 Pools
 //
 static void*  pool_mk_blk_ptrs_L2[ BLIS_NUM_MC_X_KC_BLOCKS_L2 ];
-extern char   *pool_mk_mem_L2;
+char   *pool_mk_mem_L2;
 
 static void*  pool_kn_blk_ptrs_L2[ BLIS_NUM_KC_X_NC_BLOCKS_L2 ];
-extern char   *pool_kn_mem_L2;
+char   *pool_kn_mem_L2;
 
 static void*  pool_mn_blk_ptrs_L2[ BLIS_NUM_MC_X_NR_BLOCKS_L2 ];
-extern char   *pool_mn_mem_L2;
+char   *pool_mn_mem_L2;
 
 //
 //L3 Pools
 //
 static void*  pool_mk_blk_ptrs_L3[ BLIS_NUM_MC_X_KC_BLOCKS_L3 ];
-extern char   *pool_mk_mem_L3;
+char   *pool_mk_mem_L3;
 
 static void*  pool_kn_blk_ptrs_L3[ BLIS_NUM_KC_X_NC_BLOCKS_L3 ];
-extern char   *pool_kn_mem_L3;
+char   *pool_kn_mem_L3;
 
 static void*  pool_mn_blk_ptrs_L3[ BLIS_NUM_MC_X_NC_BLOCKS_L3 ];
-extern char   *pool_mn_mem_L3;
+char   *pool_mn_mem_L3;
 
 #else
 static pool_t pools[3];
@@ -119,9 +119,6 @@ static char   pool_mn_mem[ BLIS_MN_POOL_SIZE ];
 
 
 
-
-
-
 void bli_mem_acquire_m( siz_t     req_size,
                         packbuf_t buf_type,
                         mem_t*    mem )
@@ -348,7 +345,90 @@ void bli_mem_acquire_v( siz_t  req_size,
                            mem );
 }
 
+#ifdef BLIS_ENABLE_C66X_MEM_POOLS
+/*==============================================================================
+ * Functions bli_get_mem_sizes() and bli_scratch_mem_alloc() are only used for
+ * C66x implementation.
+ *
+ * BLIS memories are allocated in 3 different ways:
+ * 1. persistent memory: allocated by lib_pmem_salloc() which calls malloc() for now.
+ * 2. very fast, fast, and medium speed scratch memory: allocated by 
+ *    bli_scratch_mem_alloc(). 
+ *       - This may be done just once at boot time if the provided memory address 
+ *         stay the same. This may be the case for DSP-only application where the
+ *         memories for BLIS are statically allocated.
+ *       - This can also be done before every level 3 BLAS function call. This 
+ *         may be the case when memories for BLIS are dynamically allocated in 
+ *         ARM+DSP applications. 
+ *       - Heaps need to be initialized by tiCblas API before allocation. 
+ * 3. slow scratch memory: allocated by lib_smem_salloc() inside BLIS computation.
+ *    The heap (slow scratch) needs to be initialized the same as above and reset 
+ *    whenever a level 3 BLIS function is called.
+ *
+ *============================================================================*/
+#define getNextMultiple(x, y) ( ( ((x)+(y)-1)/(y) )* (y) )
+
+#define BLAS_MEM_SIZE_VFAST (  getNextMultiple(BLIS_MK_POOL_SIZE_L1, BLIS_CACHE_LINE_SIZE) \
+                             + getNextMultiple(BLIS_KN_POOL_SIZE_L1, BLIS_CACHE_LINE_SIZE) \
+                             + getNextMultiple(BLIS_MN_POOL_SIZE_L1, BLIS_CACHE_LINE_SIZE) )
+#define BLAS_MEM_SIZE_FAST  (  getNextMultiple(BLIS_MK_POOL_SIZE_L2, BLIS_CACHE_LINE_SIZE) \
+                             + getNextMultiple(BLIS_KN_POOL_SIZE_L2, BLIS_CACHE_LINE_SIZE) \
+                             + getNextMultiple(BLIS_MN_POOL_SIZE_L2, BLIS_CACHE_LINE_SIZE) )
+#define BLAS_MEM_SIZE_MEDIUM ( getNextMultiple(BLIS_MK_POOL_SIZE_L3, BLIS_CACHE_LINE_SIZE) \
+                             + getNextMultiple(BLIS_KN_POOL_SIZE_L3, BLIS_CACHE_LINE_SIZE) \
+                             + getNextMultiple(BLIS_MN_POOL_SIZE_L3, BLIS_CACHE_LINE_SIZE) )
+#define BLAS_MEM_SIZE_SLOW    (4804)
+
+/*==============================================================================
+ * Purpose: returns the sizes of the non-persistent memories used by BLIS. 
+ *============================================================================*/
+void bli_get_mem_sizes(size_t *smem_size_vfast,  size_t *smem_size_fast, 
+                       size_t *smem_size_medium, size_t *smem_size_slow)
+{
+    *smem_size_vfast  = BLAS_MEM_SIZE_VFAST;  /* very fast scratch memory     */
+    *smem_size_fast   = BLAS_MEM_SIZE_FAST;   /* fast scratch memory          */
+    *smem_size_medium = BLAS_MEM_SIZE_MEDIUM; /* medium speed scratch memory  */ 
+    *smem_size_slow   = BLAS_MEM_SIZE_SLOW;   /* slow scratch memory          */
+}
 
+/*==============================================================================
+ * Purpose: allocates very fast, fast, and medium speed scratch memories from
+ *          initialized heaps. 
+ * Note: slow memoris are allocated during BLIS compuation.
+ *============================================================================*/
+int bli_scratch_mem_alloc()
+{
+    lib_memdscr_t **blas_mem_handle = blasGetMemHandle();
+    
+    pool_mk_mem_L1 = lib_smem_valloc(blas_mem_handle, BLIS_MK_POOL_SIZE_L1, BLIS_CACHE_LINE_SIZE);
+    pool_kn_mem_L1 = lib_smem_valloc(blas_mem_handle, BLIS_KN_POOL_SIZE_L1, BLIS_CACHE_LINE_SIZE);
+    pool_mn_mem_L1 = lib_smem_valloc(blas_mem_handle, BLIS_MN_POOL_SIZE_L1, BLIS_CACHE_LINE_SIZE);
+
+    pool_mk_mem_L2 = lib_smem_falloc(blas_mem_handle, BLIS_MK_POOL_SIZE_L2, BLIS_CACHE_LINE_SIZE);
+    pool_kn_mem_L2 = lib_smem_falloc(blas_mem_handle, BLIS_KN_POOL_SIZE_L2, BLIS_CACHE_LINE_SIZE);
+    pool_mn_mem_L2 = lib_smem_falloc(blas_mem_handle, BLIS_MN_POOL_SIZE_L2, BLIS_CACHE_LINE_SIZE);
+    
+    pool_mk_mem_L3 = lib_smem_malloc(blas_mem_handle, BLIS_MK_POOL_SIZE_L3, BLIS_CACHE_LINE_SIZE);
+    pool_kn_mem_L3 = lib_smem_malloc(blas_mem_handle, BLIS_KN_POOL_SIZE_L3, BLIS_CACHE_LINE_SIZE);
+    pool_mn_mem_L3 = lib_smem_malloc(blas_mem_handle, BLIS_MN_POOL_SIZE_L3, BLIS_CACHE_LINE_SIZE);    
+
+    if(  (pool_mk_mem_L1 == NULL)
+       ||(pool_kn_mem_L1 == NULL) 
+       ||(pool_mn_mem_L1 == NULL)          
+       ||(pool_mk_mem_L2 == NULL) 
+       ||(pool_kn_mem_L2 == NULL) 
+       ||(pool_mn_mem_L2 == NULL)          
+       ||(pool_mk_mem_L3 == NULL) 
+       ||(pool_kn_mem_L3 == NULL) 
+       ||(pool_mn_mem_L3 == NULL) ) {
+      return(BLI_MEM_ALLOC_ERROR);    
+    } 
+    else {
+      bli_mem_init();
+      return(BLI_MEM_ALLOC_SUCCESS);
+    }    
+}
+#endif /* BLIS_ENABLE_C66X_MEM_POOLS */
 
 void bli_mem_init()
 {
diff --git a/src/ti/linalg/blis/frame/base/bli_mem.h b/src/ti/linalg/blis/frame/base/bli_mem.h
index b015f1d..1ea3e16 100644
--- a/src/ti/linalg/blis/frame/base/bli_mem.h
+++ b/src/ti/linalg/blis/frame/base/bli_mem.h
@@ -51,6 +51,13 @@ void bli_mem_init_pool( char*   pool_mem,
                         void**  block_ptrs,
                         pool_t* pool,
                         membuf_t buf_type);
+void bli_get_mem_sizes(size_t *smem_size_vfast,  size_t *smem_size_fast, 
+                       size_t *smem_size_medium, size_t *smem_size_slow);
+int bli_scratch_mem_alloc();
+
+#define BLI_MEM_ALLOC_SUCCESS 0
+#define BLI_MEM_ALLOC_ERROR   1
+
 #else
 void bli_mem_init_pool( char*   pool_mem,
                         siz_t   block_size,
diff --git a/src/ti/linalg/ticblas/src/ticblas.c b/src/ti/linalg/ticblas/src/ticblas.c
index 3e45b5f..d7795d5 100644
--- a/src/ti/linalg/ticblas/src/ticblas.c
+++ b/src/ti/linalg/ticblas/src/ticblas.c
@@ -29,18 +29,11 @@
 #include "../ticblas.h"

 #include "blis.h"

 

-#define getNextMultiple(x, y) ( ( ((x)+(y)-1)/(y) )* (y) )

-

-#define BLAS_MEM_SIZE_VFAST (  getNextMultiple(BLIS_MK_POOL_SIZE_L1, BLIS_CACHE_LINE_SIZE) \

-                             + getNextMultiple(BLIS_KN_POOL_SIZE_L1, BLIS_CACHE_LINE_SIZE) \

-                             + getNextMultiple(BLIS_MN_POOL_SIZE_L1, BLIS_CACHE_LINE_SIZE) )

-#define BLAS_MEM_SIZE_FAST  (  getNextMultiple(BLIS_MK_POOL_SIZE_L2, BLIS_CACHE_LINE_SIZE) \

-                             + getNextMultiple(BLIS_KN_POOL_SIZE_L2, BLIS_CACHE_LINE_SIZE) \

-                             + getNextMultiple(BLIS_MN_POOL_SIZE_L2, BLIS_CACHE_LINE_SIZE) )

-#define BLAS_MEM_SIZE_MEDIUM ( getNextMultiple(BLIS_MK_POOL_SIZE_L3, BLIS_CACHE_LINE_SIZE) \

-                             + getNextMultiple(BLIS_KN_POOL_SIZE_L3, BLIS_CACHE_LINE_SIZE) \

-                             + getNextMultiple(BLIS_MN_POOL_SIZE_L3, BLIS_CACHE_LINE_SIZE) )

-#define BLAS_MEM_SIZE_SLOW    (4804)

+/*=============================================================================

+ * This file contains the CBLAS API Extension for TI-DSP. This extension is used

+ * in the OpenCL DSP layer for ARM+DSP CBLAS library. It may also be used directly

+ * by the user for DSP-only applications. 

+ *===========================================================================*/

 

 /* Define memory descriptors for memory management */

 lib_memdscr_t blas_mem_vfast;

@@ -56,23 +49,8 @@ lib_memdscr_t * blas_memdscr_tab[LIB_MEMTYPE_N] = {
     &blas_mem_slow

 };

 

-// note these pointers must be filled if used functions

-char *pool_mk_mem_L1;

-char *pool_kn_mem_L1;

-char *pool_mn_mem_L1;

-

-char *pool_mk_mem_L2;

-char *pool_kn_mem_L2;

-char *pool_mn_mem_L2;

-

-char *pool_mk_mem_L3;

-char *pool_kn_mem_L3;

-char *pool_mn_mem_L3;

-

-extern void bli_mem_init();

-

 /*==============================================================================

- * This function returns the address of the memory descriptor array

+ * This internal function returns the address of the memory descriptor array.

  *============================================================================*/

 void * blasGetMemHandle()

 {

@@ -80,94 +58,61 @@ void * blasGetMemHandle()
 } /* blasGetMemHandle */

 

 /*==============================================================================

- *   It returns the size requirement of each of the 4 memory types defined in 

- *   the library framework. 

+ * External API. Refer to ticblas.h for detailed documentation.

  *============================================================================*/

 void tiCblasGetSizes(size_t *smem_size_vfast,  size_t *smem_size_fast, 

                      size_t *smem_size_medium, size_t *smem_size_slow)

-{

-    *smem_size_vfast  = BLAS_MEM_SIZE_VFAST;  /* very fast scratch memory     */

-    *smem_size_fast   = BLAS_MEM_SIZE_FAST;   /* fast scratch memory          */

-    *smem_size_medium = BLAS_MEM_SIZE_MEDIUM; /* medium speed scratch memory  */ 

-    *smem_size_slow   = BLAS_MEM_SIZE_SLOW;   /* slow scratch memory          */

-/*

-    printf("BLIS_MK_POOL_SIZE_L1 is %d.\n", BLIS_MK_POOL_SIZE_L1);

-    printf("BLIS_KN_POOL_SIZE_L1 is %d.\n", BLIS_KN_POOL_SIZE_L1);

-    printf("BLIS_MN_POOL_SIZE_L1 is %d.\n", BLIS_MN_POOL_SIZE_L1);

-    printf("BLIS_MK_POOL_SIZE_L2 is %d.\n", BLIS_MK_POOL_SIZE_L2);

-    printf("BLIS_KN_POOL_SIZE_L2 is %d.\n", BLIS_KN_POOL_SIZE_L2);

-    printf("BLIS_MN_POOL_SIZE_L2 is %d.\n", BLIS_MN_POOL_SIZE_L2);

-    printf("BLIS_MK_POOL_SIZE_L3 is %d.\n", BLIS_MK_POOL_SIZE_L3);

-    printf("BLIS_KN_POOL_SIZE_L3 is %d.\n", BLIS_KN_POOL_SIZE_L3);

-    printf("BLIS_MN_POOL_SIZE_L3 is %d.\n", BLIS_MN_POOL_SIZE_L3);

-*/

+{   

+    /* get memory requirement information from BLIS */

+    bli_get_mem_sizes(smem_size_vfast, smem_size_fast, smem_size_medium, smem_size_slow);

 } /* tiCblasGetSizes */

 

 /*==============================================================================

- *   It performs necessary initialization through library framework API in order

- *   to do memory allocations.  

+ * External API. Refer to ticblas.h for detailed documentation.

  *============================================================================*/

 int tiCblasInit(void * mem_vfast_base,  size_t mem_vfast_size,

                 void * mem_fast_base,   size_t mem_fast_size,

                 void * mem_medium_base, size_t mem_medium_size,

                 void * mem_slow_base,   size_t mem_slow_size)

 {

+  size_t mem_vfast_size_req, mem_fast_size_req, mem_medium_size_req, mem_slow_size_req;

   lib_memdscr_t **blas_mem_handle = blasGetMemHandle();

 

+  /* Get the memory size requirements by BLIS */

+  bli_get_mem_sizes(&mem_vfast_size_req, &mem_fast_size_req, 

+                    &mem_medium_size_req, &mem_slow_size_req);

+    

   /* Verify supplied memories meet requirements */    

-  if(  ((mem_vfast_base  == NULL) || (mem_vfast_size  < BLAS_MEM_SIZE_VFAST)  )

-     ||((mem_fast_base   == NULL) || (mem_fast_size   < BLAS_MEM_SIZE_FAST)   )

-     ||((mem_medium_base == NULL) || (mem_medium_size < BLAS_MEM_SIZE_MEDIUM) )

-     ||((mem_slow_base   == NULL) || (mem_slow_size   < BLAS_MEM_SIZE_SLOW) )

+  if(  ( (mem_vfast_base  == NULL) || (mem_vfast_size  < mem_vfast_size_req)  )

+     ||( (mem_fast_base   == NULL) || (mem_fast_size   < mem_fast_size_req)   )

+     ||( (mem_medium_base == NULL) || (mem_medium_size < mem_medium_size_req) )

+     ||( (mem_slow_base   == NULL) || (mem_slow_size   < mem_slow_size_req)   )

     ) {

     return(TICBLAS_ERROR_NOMEM);

   }

   else {

+    /* Initialize all 4 types of scratch heaps */

     lib_smem_vinit(blas_mem_handle, mem_vfast_base,  mem_vfast_size);

     lib_smem_finit(blas_mem_handle, mem_fast_base,   mem_fast_size);

     lib_smem_minit(blas_mem_handle, mem_medium_base, mem_medium_size);

     lib_smem_sinit(blas_mem_handle, mem_slow_base,   mem_slow_size);      

- 

-    pool_mk_mem_L1 = lib_smem_valloc(blas_mem_handle, BLIS_MK_POOL_SIZE_L1, BLIS_CACHE_LINE_SIZE);

-    pool_kn_mem_L1 = lib_smem_valloc(blas_mem_handle, BLIS_KN_POOL_SIZE_L1, BLIS_CACHE_LINE_SIZE);

-    pool_mn_mem_L1 = lib_smem_valloc(blas_mem_handle, BLIS_MN_POOL_SIZE_L1, BLIS_CACHE_LINE_SIZE);

-

-    pool_mk_mem_L2 = lib_smem_falloc(blas_mem_handle, BLIS_MK_POOL_SIZE_L2, BLIS_CACHE_LINE_SIZE);

-    pool_kn_mem_L2 = lib_smem_falloc(blas_mem_handle, BLIS_KN_POOL_SIZE_L2, BLIS_CACHE_LINE_SIZE);

-    pool_mn_mem_L2 = lib_smem_falloc(blas_mem_handle, BLIS_MN_POOL_SIZE_L2, BLIS_CACHE_LINE_SIZE);

     

-    pool_mk_mem_L3 = lib_smem_malloc(blas_mem_handle, BLIS_MK_POOL_SIZE_L3, BLIS_CACHE_LINE_SIZE);

-    pool_kn_mem_L3 = lib_smem_malloc(blas_mem_handle, BLIS_KN_POOL_SIZE_L3, BLIS_CACHE_LINE_SIZE);

-    pool_mn_mem_L3 = lib_smem_malloc(blas_mem_handle, BLIS_MN_POOL_SIZE_L3, BLIS_CACHE_LINE_SIZE);

-/*

-    printf("BLIS_MK_POOL_SIZE_L1 is %d, pool_mk_mem_L1 is 0x%x.\n", BLIS_MK_POOL_SIZE_L1, (unsigned int)pool_mk_mem_L1);

-    printf("BLIS_KN_POOL_SIZE_L1 is %d, pool_kn_mem_L1 is 0x%x.\n", BLIS_KN_POOL_SIZE_L1, (unsigned int)pool_kn_mem_L1);

-    printf("BLIS_MN_POOL_SIZE_L1 is %d, pool_mn_mem_L1 is 0x%x.\n", BLIS_MN_POOL_SIZE_L1, (unsigned int)pool_mn_mem_L1);

-    printf("BLIS_MK_POOL_SIZE_L2 is %d, pool_mk_mem_L2 is 0x%x.\n", BLIS_MK_POOL_SIZE_L2, (unsigned int)pool_mk_mem_L2);

-    printf("BLIS_KN_POOL_SIZE_L2 is %d, pool_kn_mem_L2 is 0x%x.\n", BLIS_KN_POOL_SIZE_L2, (unsigned int)pool_kn_mem_L2);

-    printf("BLIS_MN_POOL_SIZE_L2 is %d, pool_mn_mem_L2 is 0x%x.\n", BLIS_MN_POOL_SIZE_L2, (unsigned int)pool_mn_mem_L2);

-    printf("BLIS_MK_POOL_SIZE_L3 is %d, pool_mk_mem_L3 is 0x%x.\n", BLIS_MK_POOL_SIZE_L3, (unsigned int)pool_mk_mem_L3);

-    printf("BLIS_KN_POOL_SIZE_L3 is %d, pool_kn_mem_L3 is 0x%x.\n", BLIS_KN_POOL_SIZE_L3, (unsigned int)pool_kn_mem_L3);

-    printf("BLIS_MN_POOL_SIZE_L3 is %d, pool_mn_mem_L3 is 0x%x.\n", BLIS_MN_POOL_SIZE_L3, (unsigned int)pool_mn_mem_L3);

-*/  

-    if(  (pool_mk_mem_L1 == NULL)

-       ||(pool_kn_mem_L1 == NULL) 

-       ||(pool_mn_mem_L1 == NULL)          

-       ||(pool_mk_mem_L2 == NULL) 

-       ||(pool_kn_mem_L2 == NULL) 

-       ||(pool_mn_mem_L2 == NULL)          

-       ||(pool_mk_mem_L3 == NULL) 

-       ||(pool_kn_mem_L3 == NULL) 

-       ||(pool_mn_mem_L3 == NULL) ) {

-      return(TICBLAS_ERROR_MEMINIT);    

-    } 

+    /* Make a BLIS call to allocate scratch memory from very fast heap, 

+     * fast heap and medium speed heap for BLIS. Slow memoris are allocated during 

+     * BLIS compuation. Refer to blis/frame/base/bli_mem.c for detailed information. 

+     */

+    if(bli_scratch_mem_alloc() == BLI_MEM_ALLOC_ERROR) {

+      return(TICBLAS_ERROR_MEMALLOC);

+    }

     else {

-      bli_mem_init();

       return(TICBLAS_SUCCESS);

     }

   }

 } /* tiCblasInit */

 

+/*==============================================================================

+ * External API. Refer to ticblas.h for detailed documentation.

+ *============================================================================*/

 int tiCblasNew()

 {

   if(bli_init() == BLIS_SUCCESS) {

@@ -178,6 +123,9 @@ int tiCblasNew()
   }

 }

 

+/*==============================================================================

+ * External API. Refer to ticblas.h for detailed documentation.

+ *============================================================================*/

 int tiCblasDelete()

 {

   if(bli_finalize() == BLIS_SUCCESS) {

diff --git a/src/ti/linalg/ticblas/ticblas.h b/src/ti/linalg/ticblas/ticblas.h
index 18b7a12..813da18 100644
--- a/src/ti/linalg/ticblas/ticblas.h
+++ b/src/ti/linalg/ticblas/ticblas.h
@@ -31,6 +31,13 @@
 #include <stddef.h> 

 

 /** @defgroup ti_cblas_api CBLAS API Extension for TI-DSP 

+ *  @brief This extension contains the initialization and finalization APIs 

+ *         for proper usage of CBLAS optimized for TI's C66x DSP.

+ *            - for ARM+DSP applications with CBLAS API on ARM (host), users

+ *              do not need to use this extension. 

+ *            - for DSP-only applications, uers will need to use this extension

+ *              to set up before making standard CBLAS calls in their programs

+ *              running on the DSP, and to tear down before the program exits. 

  *  @{

  */

 /** @} */

@@ -42,11 +49,9 @@
 /*@{*/

 #define TICBLAS_SUCCESS        (0)  /**< Success. No error.                      */

 #define TICBLAS_ERROR_NOMEM    (1)  /**< Failure. Provided memory is not enough. */

-#define TICBLAS_ERROR_MEMINIT  (2)  /**< Failure. Memory init error.             */

-#define TICBLAS_ERROR_MEMCFG   (3)  /**< Failure. L1D/L2 config error.           */

-#define TICBLAS_ERROR_NEW      (4)  /**< Failure. tiCblasnew error.              */

-#define TICBLAS_ERROR_DELETE   (5)  /**< Failure. tiCblasDelete error.           */

-#define TICBLAS_ERROR_MEMRECFG (6)  /**< Failure. L1D/L2 reconfig error.         */

+#define TICBLAS_ERROR_MEMALLOC (2)  /**< Failure. Memory allocation error.       */

+#define TICBLAS_ERROR_NEW      (3)  /**< Failure. CBLAS creation error.          */

+#define TICBLAS_ERROR_DELETE   (4)  /**< Failure. CBLAS deletion error.          */

 /*@}*/

 /** @} */

 

@@ -67,7 +72,7 @@
  *           based on speed:

  *           - very fast memory, e.g. L1D;

  *           - fast memory, e.g. L2;

- *           - medium memory, e.g. L3/MSMC;

+ *           - medium speed memory, e.g. L3/MSMC;

  *           - slow memory, e.g. DDR.

  *

  *  @param[out]  smem_size_vfast    size of very fast shared memory 

@@ -81,25 +86,34 @@ void tiCblasGetSizes(size_t *smem_size_vfast,  size_t *smem_size_fast,
 

 /**

  *  @ingroup ti_cblas_api

- *  @brief   Function tiCblasNew() creates an instance for CBLAS.

+ *  @brief   Function tiCblasNew() creates and initializes global structures 

+ *           for CBLAS.

  *

  *  @remarks  tiCblasNew() MUST be called before tiCblasInit().

  *

- *  @retval     TICBLAS_SUCCESS  @copydoc TICBLAS_SUCCESS

- *  @retval     TICBLAS_ERROR    @copydoc TICBLAS_ERROR

+ *  @retval     TICBLAS_SUCCESS      @copydoc TICBLAS_SUCCESS

+ *  @retval     TICBLAS_ERROR_NEW    @copydoc TICBLAS_ERROR_NEW

  */

 int tiCblasNew();

 

 /**

  *  @ingroup ti_cblas_api

- *  @brief   Function tiCblasInit() performs heap initialization for CBLAS 

- *           to do dynamic memory allocations.

- *  

+ *  @brief   Function tiCblasInit() performs heap initialization for TI CBLAS 

+ *           to do dynamic memory allocations. 

+ *

  *  @remarks  Users must allocate memories according to the requirements 

- *            given by tiCblasGetSizes() and pass the base and size to this

+ *            given by tiCblasGetSizes() and pass the bases and sizes to this

  *            function. 

  *

- *  @remarks  tiCblasInit() must NOT be called before tiCblasNew(). 

+ *  @remarks  tiCblasNew() MUST be called before tiCblasInit() can be called.

+ *

+ *  @remarks  The provided memory can be shared with other modules and thus its

+ *            content do not need to be preserved from call to call. If the 

+ *            memory bases do not change from call to call, this function only 

+ *            needs to be called just once at boot time. However, if it cannot

+ *            be guaranteed that memory bases stay the same, this function 

+ *            needs to be called every time a level 3 CBLAS function is called

+ *            on DSP. 

  *

  *  @param[in]  mem_vfast_base    base of very fast shared memory 

  *  @param[in]  mem_vfast_size    size of very fast shared memory 

@@ -110,8 +124,9 @@ int tiCblasNew();
  *  @param[in]  mem_slow_base     base of slow shared memory 

  *  @param[in]  mem_slow_size     size of slow shared memory  

  *

- *  @retval     TICBLAS_SUCCESS  @copydoc TICBLAS_SUCCESS

- *  @retval     TICBLAS_ERROR    @copydoc TICBLAS_ERROR

+ *  @retval     TICBLAS_SUCCESS         @copydoc TICBLAS_SUCCESS

+ *  @retval     TICBLAS_ERROR_NOMEM     @copydoc TICBLAS_ERROR_NOMEM

+ *  @retval     TICBLAS_ERROR_MEMALLOC  @copydoc TICBLAS_ERROR_MEMALLOC

  */

 int tiCblasInit(void * mem_vfast_base,  size_t mem_vfast_size,

                 void * mem_fast_base,   size_t mem_fast_size,

@@ -120,11 +135,11 @@ int tiCblasInit(void * mem_vfast_base,  size_t mem_vfast_size,
 

 /**

  *  @ingroup ti_cblas_api

- *  @brief   Function tiCblasDelete() deletes the instance of CBLAS created by

- *           tiCblasNew(). 

+ *  @brief   Function tiCblasDelete() deletes global structures and frees memories

+ *           of CBLAS created by tiCblasNew(). 

  *

- *  @retval     TICBLAS_SUCCESS  @copydoc TICBLAS_SUCCESS

- *  @retval     TICBLAS_ERROR    @copydoc TICBLAS_ERROR

+ *  @retval     TICBLAS_SUCCESS         @copydoc TICBLAS_SUCCESS

+ *  @retval     TICBLAS_ERROR_DELETE    @copydoc TICBLAS_ERROR_DELETE

  */

 int tiCblasDelete();