Add CULiP sbmv

src/cublas.cu

@@ -559,6 +559,26 @@ cublasStatus_t cublasGemmStridedBatchedEx(cublasHandle_t handle,
 #undef CULIP_FUNC_ENUM_NAME
 #undef CULIP_TYPE
 
+// -------------------------------------------------
+// SBMV
+// -------------------------------------------------
+
+#define CULIP_FUNC_NAME cublasSsbmv
+#define CULIP_FUNC_ENUM_NAME CULiP_cublasSsbmv
+#define CULIP_TYPE float
+#include "cublas.sbmv.template.h"
+#undef CULIP_FUNC_NAME
+#undef CULIP_FUNC_ENUM_NAME
+#undef CULIP_TYPE
+
+#define CULIP_FUNC_NAME cublasDsbmv
+#define CULIP_FUNC_ENUM_NAME CULiP_cublasDsbmv
+#define CULIP_TYPE double
+#include "cublas.sbmv.template.h"
+#undef CULIP_FUNC_NAME
+#undef CULIP_FUNC_ENUM_NAME
+#undef CULIP_TYPE
+
 // -------------------------------------------------
 // SYRK
 // -------------------------------------------------

src/cublas.sbmv.template.h

@@ -0,0 +1,39 @@
+cublasStatus_t CULIP_FUNC_NAME(cublasHandle_t handle,
+                           cublasFillMode_t uplo, int m, int n,
+                           const CULIP_TYPE *alpha, const CULIP_TYPE *A, int lda,
+                           const CULIP_TYPE *x, int incx, const CULIP_TYPE *beta, CULIP_TYPE *y,
+                           int incy) {
+	const int profiling_flag = (CULiP_profiling_control_array[CULIP_FUNC_ENUM_NAME] == 0) && CULiP_is_profiling_enabled(CULIP_CUBLAS_DISABLE_ENV_NAME);
+
+	// Get the function pointer
+	cublasStatus_t (*cublas_lib_func)(cublasHandle_t, cublasFillMode_t, int, int, const CULIP_TYPE*, const CULIP_TYPE*, int, const CULIP_TYPE*, int, const CULIP_TYPE*, CULIP_TYPE*, int);
+	*(void**)(&cublas_lib_func) = CULiP_get_function_pointer(CULIP_CUBLAS_LIBRARY_NAME, CULIP_CUBLAS_ENV_NAME, __func__, &CULiP_cublas_lib_handle_cache);
+
+	cudaStream_t cuda_stream;
+	struct CULiP_profile_result profile_result;
+
+	if (profiling_flag) {
+		// Get current cuda stream
+		cublasGetStream(handle, &cuda_stream);
+
+		// Profile result structure
+		snprintf(profile_result.function_name, profile_result.function_name_length - 1, "%s-%s-m%d-n%d", __func__, CULiP_get_cublasFillMode_t_string(uplo), m, n);
+
+		// Record start rimestamp
+		CULiP_launch_function(cuda_stream, &CULiP_record_timestamp, (void*)&profile_result.start_timestamp);
+	}
+
+	// Call the function
+	const cublasStatus_t result = (*cublas_lib_func)(handle, uplo, m, n, alpha, A, lda, x, incx, beta, y, incy);
+	CULIBPROFILER_DEBUG_PRINT(printf("[CULiP Debug][%s] executed\n", __func__));
+
+	if (profiling_flag) {
+		// Record end rimestamp
+		CULiP_launch_function(cuda_stream, &CULiP_record_timestamp, (void*)&profile_result.end_timestamp);
+
+		// Print result
+		CULiP_launch_function(cuda_stream, &CULiP_print_profile_result, (void*)&profile_result);
+	}
+
+	return result;
+}