summary | shortlog | log | commit | commitdiff | tree
raw | patch | inline | side by side (parent: 206ede2)
raw | patch | inline | side by side (parent: 206ede2)
author | Jianzhong Xu <xuj@ti.com> | |
Wed, 25 May 2016 15:35:09 +0000 (15:35 +0000) | ||
committer | Jianzhong Xu <xuj@ti.com> | |
Wed, 25 May 2016 15:35:09 +0000 (15:35 +0000) |
155 files changed:
index 547d98328ae13fad4ada8e194bedd8736449e132..8f896eacaf667c66b619766ffaf867448b17fd1b 100644 (file)
* 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 77d8a473c11a44c4c3b0a3216faa4e8b7c0079f4..907391be0925e1fdddc9b8eb5f13008f2950c818 100644 (file)
-
/******************************************************************************
* Copyright (c) 2013-2015, Texas Instruments Incorporated - http://www.ti.com/
* All rights reserved.
* 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>
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);
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. */
#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);
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 233c69bdb60ccc5ebe51c6bef7243e1473fdf142..cfc47c785548ce49942e2130a5f21d62df05e55d 100644 (file)
#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 2a5fa6b44b4a4cb9617c85ff36ea92c797334b05..88373f15cb1bfdbfa0c36663056c9cf133bb6f5b 100644 (file)
#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 bb8fd34c56a0993ca6a5d32c52475def166db262..5a24501edbc7f41116c074031346bb8b44335596 100644 (file)
#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 2ba90052037e43816c53652d5ec20ca3d5a0659b..95473b10f88a4f729b170b9256d05fd44ba35898 100644 (file)
#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 4a04d07ad56c26c1a5490fd15b4dfa241f8070d0..cc9dd911e6ec4bc97fd492eb16bcaf52c8c6c43d 100644 (file)
#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 9880b6dfc89e374024e4e45d3b31b83cbf47b05b..38bc67c818e4b4f291bcf5057f6470cd4293506a 100644 (file)
#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 c4d35c856a1ba8d9562b82a1bc2e12581782bb27..18e9c5c8a87ec5123981b58f20c92b99fd4642f7 100644 (file)
#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 3c89112142405e95bbe0731daeff73a48e86f41c..5f4cfc59584806692c9703d5a21fbbd3e328aaf8 100644 (file)
#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 3e951aa8f6e579bb73a257e3aee7350f7c21d3e6..1bcc580f9d0891c00a5aefb8a882c2bd515e6edf 100644 (file)
#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 cfdb1f1e1a438b3a65b6601210ce2308285493aa..05c77a8e194ba6a1f0479df19ffefdd0f2dcad9e 100644 (file)
#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 0e8eb77d498a6fea8770dae8730d053e5d4e1083..c02e2db035a1bdf0a05995b2d5d4afccbff8e89a 100644 (file)
#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 7421539a7302fc193bb392ebef13a6ab629a4966..7e72a7840d0c26a5fe75ec3ad6bcb682874f09ac 100644 (file)
#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 fb566fde870116d753abd4e2d6fa1704fc19e8aa..73075dea3398856ff6babec3b075e473b47565cc 100644 (file)
#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 f65b3072b2ab0f4c99594e028010a591f6352f46..3240ed207a6d694c3d916d4ac0277f2872cd1be9 100644 (file)
#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 51ed40667360f4285da4c8795f23de1868b21dd4..f119c3b155eeb94ee097ff075e2b64ee75ebaacb 100644 (file)
#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 5e45c12482c3d4f42cf31098b29fd47d54549602..5b3dab338f894ed95e40a69160f6e02be34b520b 100644 (file)
#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 03e14dd173acf0a5947b29f705f06c90d09cd882..bdbe0cf0575902fec01ae5d93d168b64569458f9 100644 (file)
#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 315d80d9ff23439a75ad4652d4b6d645425d16ed..c68c478b3d98c39e3b4088e48d6a38cab9ba1bcd 100644 (file)
#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 7f0f765e28b88a3b07951f5ee12072139bc4fece..ca0ce5cfb77fc5748639fdbb9d14ff55ee2716c1 100644 (file)
#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 8eb77ee8a817b579e526d73c112b5823a46d634f..492e2794636244f6b7ba843d2ba67c569b3cfe80 100644 (file)
#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 7933cba2d9fc0c7953be17272f5f92b37727ea70..0d0d659399a3273f0088f1d79a8dcbf54590f1e6 100644 (file)
#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 161b6d5f8b8e7c831fc4ecc49e99fc60942c1643..f18b209ba9a1834f882757c8b88542e142083d7c 100644 (file)
#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 1dac9fedd534ce7d44d7130fc68d153488f938ea..c29f25c30d9679b5c354804a979fb0be4851c36d 100644 (file)
#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 41558f8192a7e9d8a27382be64c8ea396b0e2395..5117b67156babb7b76a05f26fa7f6d80f155a7f0 100644 (file)
#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 3a3611f1e7f531e6e7f11d0d859ea5b51d72e4ca..0ca1f1d90b98881eab81b0b41baf8e667533c480 100644 (file)
#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 8a0a4a4224058b659b209c8be272f12137d2028f..d57db3af94f69d161a3d0a7166792326b3e45eef 100644 (file)
#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 237b7febe40b01b01ecb3a8afa78c9ba87f9669c..493897d0032eb7b9f246143a1c868cd01cf1e660 100644 (file)
#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 b03c3166381c64e2851634ffe9d4fa3a5e38f48b..d12a66b0c4e7a2e768bda437ac57f544afe5901b 100644 (file)
#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 816fed24bfa6d2426eabc789356aedafc16846ca..c5ef3a6a3e6ec626f983ee99cec315227a0d2395 100644 (file)
#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 44d25ed595ae7aea9ebba6491e4c57e9709f6377..c5548778717351eb3e68f44a3873ffce79b03bd4 100644 (file)
#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 17713afbbdd62cd2edaa0a410876438a2cbc24cf..6740788641694769a3ed213bf5841e20b2bd47d9 100644 (file)
#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 e2461a537ffc44343583fb0d0737dff32e7717f8..890bbd2ac3ba5481de872645c3e1e4ed58129257 100644 (file)
#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 7279d08ee39073af1f7c9dbd29cbb43171d285d4..1a104b890774ca61b2d73fbfd6c703c996aaeeea 100644 (file)
#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 2ecf16804f5a6d778ca1a9592d216636b7c942cf..1969368040300e9784d406d2510b9155f826979a 100644 (file)
#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 3425c225829197bd64c48ec22c05694b4d1629eb..466478a4cbe1e5f5e152be0d61f8d494a6bbf3a4 100644 (file)
#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 64a6824be27981387aec5cf4fa34656495730dcb..62fea4504b013d4e9a3e30a42a8310378d405c2f 100644 (file)
#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 c145e2da13925a3fe3d34183c8df8174bc6d5ca5..52ab3c78237e0557427b68b1e04e38eed69663bb 100644 (file)
#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 aa1f769a67e9c6f4a7a829b0f303c262b62cf274..17175024b8b15fa8327dc7734382c940fbc28775 100644 (file)
#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 1ce1c80d8f3db4f135bf9dd7d643bc49ad620be3..15f41c9e7adc1e2623ec0389ba101d84be89a6ec 100644 (file)
#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 92ecb14c72cec6f4aafe2597f0cfd6f5edd73b82..58f6abc481c49ff6eeb1559c48cd8ee9b4b1c75c 100644 (file)
#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 2c3a5e2ad743bf11a6bbcfee2900ad824673273e..ba2fefbae9cc83e1a9a985226e1c59eb2dfec20c 100644 (file)
#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 1c911e7245419c126c7177ae96b22483df096502..c30dede51a475220cbdc4258ea68a8d724c3679a 100644 (file)
#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 a3c14c681f8ce132a6d0c4edd4890358a9e6e44a..3b924357668ccc3ccd097b9f6f9d0a05e08732d6 100644 (file)
#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 dd86c24823efd90285b11eb31337c00f356080b0..dfc2b8243428c2fcf3bb7a744b5d9ae0b7928429 100644 (file)
#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 5b3732c187c57927bb951a96b8670f6afc3df8c5..a09995b5292e0c5dcfffa327fa31e6e709ca1350 100644 (file)
#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 0ba101a19437e65e64fa309688cfab33b4b1e32a..30fd5d098d21e741d5441ea16e91fa07833ba515 100644 (file)
#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 efbcc94daccaaaab4b190248e904f5c861c4fdae..921ece43a09b80cf2f77666f4e8755bac096277f 100644 (file)
#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 b8267660872abad4e2f9e799e48183b5e87b7e3c..7532d4898ab781b405d6d23a9caba2864ee2774a 100644 (file)
#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 93d81f0288417f8667eede5160578e686af68ded..dd75f67ddc5c574b4f3d6da547369fe11d9727ef 100644 (file)
#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 8f027918b92f909d41ece6b1671411d4550cb6a4..602414ce35150c0985dc94856a9d0f9c95170e69 100644 (file)
#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 13b403b29b4e710d6081aedadd1c45a87462c351..0c24fee2f612a781f5e47a7cd71ad94cc7ed937f 100644 (file)
#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 c3479696e95e8d5a19ffbda1d0554d250515909f..2f75cf97b11e2ad6ed876773a6a3641bbd9f306f 100644 (file)
#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 812bb3bd7b903889e2de4439deacfde511307d44..38620ae61630556c0967f232e4d9537069a715a1 100644 (file)
#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 d17893dd6a9d3cfaf42c827513869946d7e176c7..9cd70b90cefc36afa926d195d58b638c1896dbcc 100644 (file)
#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 d5c28bc24c720acfaf68bba9cc0be0c52bcb4cf1..4136f9ef29db6544abaa814a1a3256f8461217fc 100644 (file)
#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 80eb1a67ed2955a3c5472a2e2ce5a542a674ffe5..5a5d1c548acae06f8848dd803bdf36b23ee12703 100644 (file)
#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 6b05e31db0cfc8d560f3c3563c31792a3fc33529..6df66381966658402040fe0e3197ceecdc0edd2e 100644 (file)
#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 86a5a48f4bca85f6dc7f60acd45e8cdd958af212..384c68f3c4a8f0220a5ad7dd15cdc04c4dc67dc8 100644 (file)
#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 4caffb7de298583d88eefabe161471684235edaa..35464750527cbf903add05ed455dea3f1d721d41 100644 (file)
#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 341ff82ff53a9e39fb0fef86642fd5c603166927..923a66a3a08bc024349e6652870d411efaf78b9a 100644 (file)
#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 d95c34f09a9124e55768e4a797874627e6d51652..102b8424773c1c66e691573c8f9438a7a8ab7fd5 100644 (file)
#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 c375c4ecf95729de4adca6e47839acd4a8681cba..dde82ef3099af04885d2b453d30fd1b6e2fe8124 100644 (file)
#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 105c178e34796cf8eb303b2c0dec1a392078fb40..0a5da496d5a78ffbbb1da5d2d7e6b9ce42b4cf72 100644 (file)
#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 32311ab056a1f767871ed3cec340bddfb5379323..0f2abb2e7f723e627b595d699a805b3e9c0bb816 100644 (file)
#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 cac7e276f6a73a63abaa363d2392ceab355a9693..28537a3114726d487cceb2f98a01621dfc968506 100644 (file)
#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 d281310446b5ae87da523ff7b950879cfdccee25..14d2000cfd9d1d1e67bc4a3377ac5390d6df776d 100644 (file)
#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 a3bbc06c16a643f46c72233e6e26e57e8b9f5c6b..e53f21f09316c5d1a5d434c5252050be6eeaf8de 100644 (file)
#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 4001c3f22e5b0ced5c9a0faaa6d4fb74a31a07f7..4d773352b6fcf11ee68f3f7648edfe925b93c43b 100644 (file)
#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 5b79f68d7de150abf64c254dfda8f6e99a3bcb14..a1bb3d9403f41ecb21bea44bb68c83e476674332 100644 (file)
#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 81d7e539848459bedf09767664035d4296349c6a..265d6f9ae159342e515608689584d7e7dc12cf99 100644 (file)
#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 7cbcfc3424fd833e3a0f46bef12b76eedbb0a91e..860f4d731941fda001040f6b33b55ed44d47b502 100644 (file)
#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 d393ea2dee1592a4ae1d1a7dcd2a9dba407d6f52..9aab3172ef47a8499d9f2880169012d3e86c2382 100644 (file)
#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 35d6aa6c433dad62251565649f141a6100b1e0f9..b15a9e62d3482363256fcb89907122635251c83c 100644 (file)
#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 f23531903a6a78d7a71e504a4bed8a00134ce271..7fc6d66a220499f2b257b995957e5a33488ad95a 100644 (file)
#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 60b0b88da69eedd7c74e56d7c4841588c316f159..dd46358955332aea7ce0cf101fde69843a094cb5 100644 (file)
#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 2caee520b15c109476508d843704657d17c276ac..cb613055d1ac4db99ca2b0d1d826001e2b1ab68d 100644 (file)
#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 67c18eafd9f856abb22474c2490ef5bbe95fdf4c..3845b05d60937bf48983e73bbe7c693e1a30b752 100644 (file)
#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 a9e271381d91cf4b313ad573d9672107880d246d..733a4e2613f28b8d1f2434d717c3ba50358e52cb 100644 (file)
#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 2c3a8f3215480952678ca68cf5c71beeadc0849d..c3d12fb76a8ce503f24fc59d1bc1f83c13ac17a0 100644 (file)
#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 ab7d8d0eccebab4f57bdef72e11cee33e4419c98..636464d6bd742efdb4e4009a65ed428e2f67aa32 100644 (file)
#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 fd5adb9dea7f0e67464991fada1cdafe62eb7266..ed4822aaeae3a697f6ff717df5549f33b0112bb2 100644 (file)
#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 32d800ae4a9308551d2b7c0d4b5ae3316bfd4a9d..c7f4895d12ea82e6fa4dbd85ef9ec531cd4a08a5 100644 (file)
#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 9599e84db0b23786fddfe3d5ca160697aad4e945..ecd1885403a54d9c31ee68d8b8e7c997f2b072bf 100644 (file)
#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 48a3de602e1fa7501f58a186a2cb7779a56a69b3..ad3adaf9b5bd1b65e5800d3d8e74094818f6aa32 100644 (file)
#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 6780fa907d163c4a7c1032cb7a5632ab156184e8..b6b3680c125e7d50962e8294173544e296ac255d 100644 (file)
#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 ba378a301b3b1015b3b5cbd4bc6f0c36ae93e4bf..92c0a3ada910bf199970b4d2835501e27a942c26 100644 (file)
#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 a7c6bbcbb34ea476515cf25809415394c2a261bf..0a02021535a29f7fd32b62839e51b9f37c7e5cc5 100644 (file)
#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 f399accf8565f87503c65fc5edfad18a7b036d90..38a5d5a98e0ccff4d4d5b076e34694f1c1ca7d73 100644 (file)
#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 845a95436a5af53a86d90e0090d547807252ed17..c823cd20de071d65b01aac9ec0422310e28fc12f 100644 (file)
#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 8a03f42e27dd026caa84c8636b62b3b0685fbf46..68f02d697f22f12d2fb900e7d1fcc682d120c68a 100644 (file)
#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 24f9925a4fc1d2d912644bc98a269147a8b83773..6ca7edd734fcfb1993268ff345a292eb83f44099 100644 (file)
#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 df05602b088e4be22387053c4b045502bc068fc8..7a848153c47a00d94b1c5dce39e03ab415dbd1e0 100644 (file)
#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 562ccec61399d31c7cb66817956295039323a5d9..41c8ef59165711ca2a57b8661805679ab7ed445c 100644 (file)
#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 b2e934244e586b4f52eaa1f244004067451fde73..bae03717282283e83004cd9e1c12ccea47234007 100644 (file)
#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 7b82e54573298f4083c600ffd0b07c5ce36e1efc..2156d5fa658fffb1e7e1d49f24c0dd1bcaad0d6b 100644 (file)
#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 d7590dd121ce2841f7fe3c7d3330513729a15b77..19e5d3896ad7bf66b785d7f0ce583f173449a984 100644 (file)
#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 bc86c8ec103dc7e4bfef3550fd8330b12e66ad03..d62b2ec35800ac816146f9de6a0b5997ad16b468 100644 (file)
#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 bfe09acba07085ca1d1019b0b840b962e76667f0..2a8d6cf8fbc5f3682e10a65b36f7e45fec79eb18 100644 (file)
#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 e6c30c8fda1d3255af0f92fee18dad2a6dabb68f..ac1b3cacda0809e83c3000b12a0ccdc595c21aeb 100644 (file)
#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 fa56ed316d9dc1cfd461ab14a62c98add17d0379..895d1a54a701d26c73fcc5c3a7e58550e9acbb8e 100644 (file)
#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 40d85039f1b95e1ad6d11753cfe049406b9384b7..66ff1011b368eac3f8aaa949b3b80e7256c2af40 100644 (file)
#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 f9c0f31001877212bb62ea3b132cd518711a11de..c407e9729ec596b2b80645347d92dedd571a7657 100644 (file)
#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 607394c273ddd7610f46ebfee76aa49e66f8661c..50fd50d9d45b84ac8694e3dbb0631dfa1de56fff 100644 (file)
#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 96e6cab5bd9f474bf3004b92a4763c388761a819..c6d92a925d1e15d219a776fe4da98881c9f39f65 100644 (file)
#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 b45b877c4c5bfbcc9b77258ca2487c9ef7c27984..950497cf1df9a356c6f9b9fd7a14fe104f2ad1aa 100644 (file)
#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 0681ed8f644185c6cec6ce2797a6f45b4f48e22d..0b7aa2d6a0de10331142c5ee2603c633fc2ccfbb 100644 (file)
#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 3839c8e7956fd284a8086d1aafdcc246fef60725..e818b225eaf8e4ac696002ce87fdaa4d733a02ce 100644 (file)
#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 40ddbc3375d8e86daf218cd3c00ce41a1c16c626..e2b1c0d8ce4b260a70f246860e5176a26b426b5c 100644 (file)
#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 18e329e6287e7d58a1953daa138d4e2b5742703e..a44e53d7ac1770bf50e3a366814a95f7ba88e84d 100644 (file)
#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 26b9a3c8e76547cc2c05a295f6225e59a1245b16..42689ee0e7a2cfd35daad4c541d1594a9adef18d 100644 (file)
#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 23efb904e8bb5b76c347471d593b01b6d4087c38..b24f23938fb453f9f3c29f0850e45b3e5fff464f 100644 (file)
#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 fdbd907ea4e1ac980aea101ff3ed2bb0d8212b98..f17a820b66de4e63c5ba49439a5aaf1af1f72698 100644 (file)
#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 da878b56ffbe71b8951a3b92b2ffde00ceb450e2..a070a5927da12e39810c60925346d5fe772af79e 100644 (file)
#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 4a97d0ba4993c2190e46117a065eb3eb977406d6..aa1459c3078f92edced36fc405534718de48dfb6 100644 (file)
#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 7d7f3e12606d8c51f448f476c2e377d516992759..2f22cb17fd82947fbabbea7b418a4930eb3cf65f 100644 (file)
#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 7294884c36b518c5bba2e1c523505f96430f1dda..db77056a3896a1f070adf928e4785ef00c800765 100644 (file)
#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 3eb84718dd76f48a7489c5e7c6e0d7b9f0903eb2..09ebd59cd3ccd1ac8d68f0ad6c580354b6bcb566 100644 (file)
#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 ac457bd7a69992a2d361731f55a0ea120fa55597..d4b3fccb215f0b5865cd2ed90865057372dd7740 100644 (file)
#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 b4767f5b4858a98d737f2b8ac346262030aad962..02591ee9ec39419f7d605ded1511db8702c87ca4 100644 (file)
#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 58eb984d4af06d5d0498bc7cd5a09dc221dff788..8153ca5f3ce1c5f625853c5074139f18ce51b4e8 100644 (file)
#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 23d4c079cf77b08e356cb59d3f65633f4e76d6c7..7e28de2f7cabeeabfed9be62266f4df943c634f4 100644 (file)
#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 67779e49f77aca7dba32d0f9142c4cf8e8194665..31a34dec91a0bee164c39f848a6d3c6acb50aa51 100644 (file)
#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 7e94b80d08ea4a36ccce275679f50cf53ae15df9..2492139cea0abd0350b206624ac19044d4e739cd 100644 (file)
#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 c6a4cbe1a4ec9edb8bd98f7b0563f13feb08a821..a63a2f9b31b4e7cb9bd04e30abcf46debd5e9cfd 100644 (file)
#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 3fdd79ee8ef72548fff6bbf00d78849217ff0a1d..79f25990a0ec990765e978aa335ce6dd20db9bf2 100644 (file)
#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 0cdf1aba2bbfa961cebe34a3bc4ef7730d608caa..aefd5d6dd3a9b27636f9df2eff4c8f9b944ea834 100644 (file)
#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 e15a8ea81025a0d56c9b4f639238fb2fcddeff5c..55aca4e7d105436ca4fa63264e98f175e889a2cf 100644 (file)
#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 5eb89b4731cdacf20a2791a7b8ad96b40d3ea571..f3f86641bedbb992a83dd5c795ed73b2eb7f0b0a 100644 (file)
#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 bc9cdf9013a2aa36df5382ba41eab1cae8698c6c..d28b86652c525c61d92a29937d482cf1a6c23c48 100644 (file)
#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 1b25b51baf179dfe22d8fb7259e1024a38f92be6..d02acb3650b6889d48c88fa549457ca123dc6af8 100644 (file)
#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 a1aeaae7b6640000a8427109b9da3946300b55d8..95e2e76101fb57d65c080ee33fa03659e9e13e29 100644 (file)
#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 f1e437c882ee4ddc34cce605e8685979c6a98b1e..29e2707cf2afda8dd44d49bc2dabecd8d0d6a025 100644 (file)
#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 0169b07d5a27847529cd3ccef1d8e67a8ad56db7..b9912f7da92b4ce86a9834d13ee6c2f3dd6cba92 100644 (file)
#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 70a29e9e339d900b38d79aad7f40ae273c072d55..f96688e86413d33c494d1191ff68fddf763c88dd 100644 (file)
#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 9175e0ebb8b8a09f02358c07d2077c68fe01b19b..42395010ef590a73e12cf3b34a1bdab35650c75f 100644 (file)
#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 53b704443cb751df55942af3fce43e50b1a1e3ce..1e54fdf572b67ffcaaa16328dc55a53f6a0e7338 100644 (file)
#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 f705eb6ac58e2e84753d16f9bc01d45d8ea95877..e75b9bdc6141afeb71dfe0b0270eee5fdb885090 100644 (file)
#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 c026942373d1e562a45c6b2084be77ce1e034fbd..063f93949993be301a9bd85a6d23c7948278d394 100644 (file)
#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 754eca7b03969c359f69936b9a0ee99a6e246a0a..f9644c4881970b91c7dc50b3e16a697b2d1f2073 100644 (file)
#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 6fd2dd0b29806406dc270141a02a39b91b1a72bc..5361390f230471d59c8f5c64013cb77bbf6799cd 100644 (file)
#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 dcf6100e6b3aa7df19ba96bd6438237b9a94e7fc..c68987d193729fde5cb20ef58dc1deb2494f289b 100644 (file)
#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 c010d894eb1d74182d677139327a1c39a4cb5e1b..51a4a53f103934fee05dc93278d04ede704b06f4 100644 (file)
#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 1cbe9e05c04f0af3a5fcae7f68fc29532a7cc3c5..244a5f9171d0f8673aa6fde7608c0c1e2acae3c8 100644 (file)
#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 4cca72875b944680f67ad2f7324c25d0ff50a4de..d3583e80a363909ad6ccb319fb63b244eee8054d 100644 (file)
#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 044c64e7b5e566bfb5d93b7b0ccd80cb39de0e5b..f7b533c3049cbcf3d9317d6845e6753ba36667fd 100644 (file)
#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 65963d314a9c41065980c76bc60d14e7df366dd7..4ac5ddf2a8d118c52ac9a19e721d882c7f02024a 100644 (file)
* 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>
#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;
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) {
/* 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);
#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 8fa16d0e0071be262aa33f76ce8f46f18bdba3f0..b4368eb6642e0954f9c2f8f7bdbd44c2869f62f7 100644 (file)
* 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 */
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 a506eceaf31bc1771474957c54fc41296f0624be..54d6e276bb3b2dd50f1cf3fcac6400a1fbf90948 100644 (file)
#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();
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
#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 */
/*==============================================================================
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()) {
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 77a49a40eb76383ce835904ba342dcc036d972d6..7c5fe270241c566311baa92dd30f6bd3e7e69528 100755 (executable)
#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
#else
cl_kernel __K;
#endif
- __K = ${namespace}_get_kernel($trampdef, \"ocl_$trampname\");
+ __K = ${namespace}_get_kernel(\"ocl_$trampname\");
#ifdef __cplusplus
try
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;
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;
index 430321744d4b6f45fec03dd0f15654c10440298e..275ff4f0138df76e192c809a06c178bd26e28422 100644 (file)
//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];
-
-
-
void bli_mem_acquire_m( siz_t req_size,
packbuf_t buf_type,
mem_t* mem )
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()
{
index b015f1d313ca76c2894f1cea17679edf42ae0a60..1ea3e169a80448f487b4805381bbe1ec8547b26f 100644 (file)
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,
index 3e45b5fa36146b6b113e3f533f29582ab50868d1..d7795d51b6b0c7286f3e74241f945d228c993402 100644 (file)
#include "../ticblas.h"\r
#include "blis.h"\r
\r
-#define getNextMultiple(x, y) ( ( ((x)+(y)-1)/(y) )* (y) )\r
-\r
-#define BLAS_MEM_SIZE_VFAST ( getNextMultiple(BLIS_MK_POOL_SIZE_L1, BLIS_CACHE_LINE_SIZE) \\r
- + getNextMultiple(BLIS_KN_POOL_SIZE_L1, BLIS_CACHE_LINE_SIZE) \\r
- + getNextMultiple(BLIS_MN_POOL_SIZE_L1, BLIS_CACHE_LINE_SIZE) )\r
-#define BLAS_MEM_SIZE_FAST ( getNextMultiple(BLIS_MK_POOL_SIZE_L2, BLIS_CACHE_LINE_SIZE) \\r
- + getNextMultiple(BLIS_KN_POOL_SIZE_L2, BLIS_CACHE_LINE_SIZE) \\r
- + getNextMultiple(BLIS_MN_POOL_SIZE_L2, BLIS_CACHE_LINE_SIZE) )\r
-#define BLAS_MEM_SIZE_MEDIUM ( getNextMultiple(BLIS_MK_POOL_SIZE_L3, BLIS_CACHE_LINE_SIZE) \\r
- + getNextMultiple(BLIS_KN_POOL_SIZE_L3, BLIS_CACHE_LINE_SIZE) \\r
- + getNextMultiple(BLIS_MN_POOL_SIZE_L3, BLIS_CACHE_LINE_SIZE) )\r
-#define BLAS_MEM_SIZE_SLOW (4804)\r
+/*=============================================================================\r
+ * This file contains the CBLAS API Extension for TI-DSP. This extension is used\r
+ * in the OpenCL DSP layer for ARM+DSP CBLAS library. It may also be used directly\r
+ * by the user for DSP-only applications. \r
+ *===========================================================================*/\r
\r
/* Define memory descriptors for memory management */\r
lib_memdscr_t blas_mem_vfast;\r
&blas_mem_slow\r
};\r
\r
-// note these pointers must be filled if used functions\r
-char *pool_mk_mem_L1;\r
-char *pool_kn_mem_L1;\r
-char *pool_mn_mem_L1;\r
-\r
-char *pool_mk_mem_L2;\r
-char *pool_kn_mem_L2;\r
-char *pool_mn_mem_L2;\r
-\r
-char *pool_mk_mem_L3;\r
-char *pool_kn_mem_L3;\r
-char *pool_mn_mem_L3;\r
-\r
-extern void bli_mem_init();\r
-\r
/*==============================================================================\r
- * This function returns the address of the memory descriptor array\r
+ * This internal function returns the address of the memory descriptor array.\r
*============================================================================*/\r
void * blasGetMemHandle()\r
{\r
} /* blasGetMemHandle */\r
\r
/*==============================================================================\r
- * It returns the size requirement of each of the 4 memory types defined in \r
- * the library framework. \r
+ * External API. Refer to ticblas.h for detailed documentation.\r
*============================================================================*/\r
void tiCblasGetSizes(size_t *smem_size_vfast, size_t *smem_size_fast, \r
size_t *smem_size_medium, size_t *smem_size_slow)\r
-{\r
- *smem_size_vfast = BLAS_MEM_SIZE_VFAST; /* very fast scratch memory */\r
- *smem_size_fast = BLAS_MEM_SIZE_FAST; /* fast scratch memory */\r
- *smem_size_medium = BLAS_MEM_SIZE_MEDIUM; /* medium speed scratch memory */ \r
- *smem_size_slow = BLAS_MEM_SIZE_SLOW; /* slow scratch memory */\r
-/*\r
- printf("BLIS_MK_POOL_SIZE_L1 is %d.\n", BLIS_MK_POOL_SIZE_L1);\r
- printf("BLIS_KN_POOL_SIZE_L1 is %d.\n", BLIS_KN_POOL_SIZE_L1);\r
- printf("BLIS_MN_POOL_SIZE_L1 is %d.\n", BLIS_MN_POOL_SIZE_L1);\r
- printf("BLIS_MK_POOL_SIZE_L2 is %d.\n", BLIS_MK_POOL_SIZE_L2);\r
- printf("BLIS_KN_POOL_SIZE_L2 is %d.\n", BLIS_KN_POOL_SIZE_L2);\r
- printf("BLIS_MN_POOL_SIZE_L2 is %d.\n", BLIS_MN_POOL_SIZE_L2);\r
- printf("BLIS_MK_POOL_SIZE_L3 is %d.\n", BLIS_MK_POOL_SIZE_L3);\r
- printf("BLIS_KN_POOL_SIZE_L3 is %d.\n", BLIS_KN_POOL_SIZE_L3);\r
- printf("BLIS_MN_POOL_SIZE_L3 is %d.\n", BLIS_MN_POOL_SIZE_L3);\r
-*/\r
+{ \r
+ /* get memory requirement information from BLIS */\r
+ bli_get_mem_sizes(smem_size_vfast, smem_size_fast, smem_size_medium, smem_size_slow);\r
} /* tiCblasGetSizes */\r
\r
/*==============================================================================\r
- * It performs necessary initialization through library framework API in order\r
- * to do memory allocations. \r
+ * External API. Refer to ticblas.h for detailed documentation.\r
*============================================================================*/\r
int tiCblasInit(void * mem_vfast_base, size_t mem_vfast_size,\r
void * mem_fast_base, size_t mem_fast_size,\r
void * mem_medium_base, size_t mem_medium_size,\r
void * mem_slow_base, size_t mem_slow_size)\r
{\r
+ size_t mem_vfast_size_req, mem_fast_size_req, mem_medium_size_req, mem_slow_size_req;\r
lib_memdscr_t **blas_mem_handle = blasGetMemHandle();\r
\r
+ /* Get the memory size requirements by BLIS */\r
+ bli_get_mem_sizes(&mem_vfast_size_req, &mem_fast_size_req, \r
+ &mem_medium_size_req, &mem_slow_size_req);\r
+ \r
/* Verify supplied memories meet requirements */ \r
- if( ((mem_vfast_base == NULL) || (mem_vfast_size < BLAS_MEM_SIZE_VFAST) )\r
- ||((mem_fast_base == NULL) || (mem_fast_size < BLAS_MEM_SIZE_FAST) )\r
- ||((mem_medium_base == NULL) || (mem_medium_size < BLAS_MEM_SIZE_MEDIUM) )\r
- ||((mem_slow_base == NULL) || (mem_slow_size < BLAS_MEM_SIZE_SLOW) )\r
+ if( ( (mem_vfast_base == NULL) || (mem_vfast_size < mem_vfast_size_req) )\r
+ ||( (mem_fast_base == NULL) || (mem_fast_size < mem_fast_size_req) )\r
+ ||( (mem_medium_base == NULL) || (mem_medium_size < mem_medium_size_req) )\r
+ ||( (mem_slow_base == NULL) || (mem_slow_size < mem_slow_size_req) )\r
) {\r
return(TICBLAS_ERROR_NOMEM);\r
}\r
else {\r
+ /* Initialize all 4 types of scratch heaps */\r
lib_smem_vinit(blas_mem_handle, mem_vfast_base, mem_vfast_size);\r
lib_smem_finit(blas_mem_handle, mem_fast_base, mem_fast_size);\r
lib_smem_minit(blas_mem_handle, mem_medium_base, mem_medium_size);\r
lib_smem_sinit(blas_mem_handle, mem_slow_base, mem_slow_size); \r
- \r
- pool_mk_mem_L1 = lib_smem_valloc(blas_mem_handle, BLIS_MK_POOL_SIZE_L1, BLIS_CACHE_LINE_SIZE);\r
- pool_kn_mem_L1 = lib_smem_valloc(blas_mem_handle, BLIS_KN_POOL_SIZE_L1, BLIS_CACHE_LINE_SIZE);\r
- pool_mn_mem_L1 = lib_smem_valloc(blas_mem_handle, BLIS_MN_POOL_SIZE_L1, BLIS_CACHE_LINE_SIZE);\r
-\r
- pool_mk_mem_L2 = lib_smem_falloc(blas_mem_handle, BLIS_MK_POOL_SIZE_L2, BLIS_CACHE_LINE_SIZE);\r
- pool_kn_mem_L2 = lib_smem_falloc(blas_mem_handle, BLIS_KN_POOL_SIZE_L2, BLIS_CACHE_LINE_SIZE);\r
- pool_mn_mem_L2 = lib_smem_falloc(blas_mem_handle, BLIS_MN_POOL_SIZE_L2, BLIS_CACHE_LINE_SIZE);\r
\r
- pool_mk_mem_L3 = lib_smem_malloc(blas_mem_handle, BLIS_MK_POOL_SIZE_L3, BLIS_CACHE_LINE_SIZE);\r
- pool_kn_mem_L3 = lib_smem_malloc(blas_mem_handle, BLIS_KN_POOL_SIZE_L3, BLIS_CACHE_LINE_SIZE);\r
- pool_mn_mem_L3 = lib_smem_malloc(blas_mem_handle, BLIS_MN_POOL_SIZE_L3, BLIS_CACHE_LINE_SIZE);\r
-/*\r
- 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);\r
- 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);\r
- 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);\r
- 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);\r
- 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);\r
- 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);\r
- 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);\r
- 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);\r
- 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);\r
-*/ \r
- if( (pool_mk_mem_L1 == NULL)\r
- ||(pool_kn_mem_L1 == NULL) \r
- ||(pool_mn_mem_L1 == NULL) \r
- ||(pool_mk_mem_L2 == NULL) \r
- ||(pool_kn_mem_L2 == NULL) \r
- ||(pool_mn_mem_L2 == NULL) \r
- ||(pool_mk_mem_L3 == NULL) \r
- ||(pool_kn_mem_L3 == NULL) \r
- ||(pool_mn_mem_L3 == NULL) ) {\r
- return(TICBLAS_ERROR_MEMINIT); \r
- } \r
+ /* Make a BLIS call to allocate scratch memory from very fast heap, \r
+ * fast heap and medium speed heap for BLIS. Slow memoris are allocated during \r
+ * BLIS compuation. Refer to blis/frame/base/bli_mem.c for detailed information. \r
+ */\r
+ if(bli_scratch_mem_alloc() == BLI_MEM_ALLOC_ERROR) {\r
+ return(TICBLAS_ERROR_MEMALLOC);\r
+ }\r
else {\r
- bli_mem_init();\r
return(TICBLAS_SUCCESS);\r
}\r
}\r
} /* tiCblasInit */\r
\r
+/*==============================================================================\r
+ * External API. Refer to ticblas.h for detailed documentation.\r
+ *============================================================================*/\r
int tiCblasNew()\r
{\r
if(bli_init() == BLIS_SUCCESS) {\r
}\r
}\r
\r
+/*==============================================================================\r
+ * External API. Refer to ticblas.h for detailed documentation.\r
+ *============================================================================*/\r
int tiCblasDelete()\r
{\r
if(bli_finalize() == BLIS_SUCCESS) {\r
index 18b7a1289f8c08fec3ec8e8471ae76b445287a47..813da186ada12acc6d3b9f1525fe183c1e5f63ee 100644 (file)
#include <stddef.h> \r
\r
/** @defgroup ti_cblas_api CBLAS API Extension for TI-DSP \r
+ * @brief This extension contains the initialization and finalization APIs \r
+ * for proper usage of CBLAS optimized for TI's C66x DSP.\r
+ * - for ARM+DSP applications with CBLAS API on ARM (host), users\r
+ * do not need to use this extension. \r
+ * - for DSP-only applications, uers will need to use this extension\r
+ * to set up before making standard CBLAS calls in their programs\r
+ * running on the DSP, and to tear down before the program exits. \r
* @{\r
*/\r
/** @} */\r
/*@{*/\r
#define TICBLAS_SUCCESS (0) /**< Success. No error. */\r
#define TICBLAS_ERROR_NOMEM (1) /**< Failure. Provided memory is not enough. */\r
-#define TICBLAS_ERROR_MEMINIT (2) /**< Failure. Memory init error. */\r
-#define TICBLAS_ERROR_MEMCFG (3) /**< Failure. L1D/L2 config error. */\r
-#define TICBLAS_ERROR_NEW (4) /**< Failure. tiCblasnew error. */\r
-#define TICBLAS_ERROR_DELETE (5) /**< Failure. tiCblasDelete error. */\r
-#define TICBLAS_ERROR_MEMRECFG (6) /**< Failure. L1D/L2 reconfig error. */\r
+#define TICBLAS_ERROR_MEMALLOC (2) /**< Failure. Memory allocation error. */\r
+#define TICBLAS_ERROR_NEW (3) /**< Failure. CBLAS creation error. */\r
+#define TICBLAS_ERROR_DELETE (4) /**< Failure. CBLAS deletion error. */\r
/*@}*/\r
/** @} */\r
\r
* based on speed:\r
* - very fast memory, e.g. L1D;\r
* - fast memory, e.g. L2;\r
- * - medium memory, e.g. L3/MSMC;\r
+ * - medium speed memory, e.g. L3/MSMC;\r
* - slow memory, e.g. DDR.\r
*\r
* @param[out] smem_size_vfast size of very fast shared memory \r
\r
/**\r
* @ingroup ti_cblas_api\r
- * @brief Function tiCblasNew() creates an instance for CBLAS.\r
+ * @brief Function tiCblasNew() creates and initializes global structures \r
+ * for CBLAS.\r
*\r
* @remarks tiCblasNew() MUST be called before tiCblasInit().\r
*\r
- * @retval TICBLAS_SUCCESS @copydoc TICBLAS_SUCCESS\r
- * @retval TICBLAS_ERROR @copydoc TICBLAS_ERROR\r
+ * @retval TICBLAS_SUCCESS @copydoc TICBLAS_SUCCESS\r
+ * @retval TICBLAS_ERROR_NEW @copydoc TICBLAS_ERROR_NEW\r
*/\r
int tiCblasNew();\r
\r
/**\r
* @ingroup ti_cblas_api\r
- * @brief Function tiCblasInit() performs heap initialization for CBLAS \r
- * to do dynamic memory allocations.\r
- * \r
+ * @brief Function tiCblasInit() performs heap initialization for TI CBLAS \r
+ * to do dynamic memory allocations. \r
+ *\r
* @remarks Users must allocate memories according to the requirements \r
- * given by tiCblasGetSizes() and pass the base and size to this\r
+ * given by tiCblasGetSizes() and pass the bases and sizes to this\r
* function. \r
*\r
- * @remarks tiCblasInit() must NOT be called before tiCblasNew(). \r
+ * @remarks tiCblasNew() MUST be called before tiCblasInit() can be called.\r
+ *\r
+ * @remarks The provided memory can be shared with other modules and thus its\r
+ * content do not need to be preserved from call to call. If the \r
+ * memory bases do not change from call to call, this function only \r
+ * needs to be called just once at boot time. However, if it cannot\r
+ * be guaranteed that memory bases stay the same, this function \r
+ * needs to be called every time a level 3 CBLAS function is called\r
+ * on DSP. \r
*\r
* @param[in] mem_vfast_base base of very fast shared memory \r
* @param[in] mem_vfast_size size of very fast shared memory \r
* @param[in] mem_slow_base base of slow shared memory \r
* @param[in] mem_slow_size size of slow shared memory \r
*\r
- * @retval TICBLAS_SUCCESS @copydoc TICBLAS_SUCCESS\r
- * @retval TICBLAS_ERROR @copydoc TICBLAS_ERROR\r
+ * @retval TICBLAS_SUCCESS @copydoc TICBLAS_SUCCESS\r
+ * @retval TICBLAS_ERROR_NOMEM @copydoc TICBLAS_ERROR_NOMEM\r
+ * @retval TICBLAS_ERROR_MEMALLOC @copydoc TICBLAS_ERROR_MEMALLOC\r
*/\r
int tiCblasInit(void * mem_vfast_base, size_t mem_vfast_size,\r
void * mem_fast_base, size_t mem_fast_size,\r
\r
/**\r
* @ingroup ti_cblas_api\r
- * @brief Function tiCblasDelete() deletes the instance of CBLAS created by\r
- * tiCblasNew(). \r
+ * @brief Function tiCblasDelete() deletes global structures and frees memories\r
+ * of CBLAS created by tiCblasNew(). \r
*\r
- * @retval TICBLAS_SUCCESS @copydoc TICBLAS_SUCCESS\r
- * @retval TICBLAS_ERROR @copydoc TICBLAS_ERROR\r
+ * @retval TICBLAS_SUCCESS @copydoc TICBLAS_SUCCESS\r
+ * @retval TICBLAS_ERROR_DELETE @copydoc TICBLAS_ERROR_DELETE\r
*/\r
int tiCblasDelete();\r
\r