From cb1f33c6ae840e8dc0f43518daf76e6ed01034f0 Mon Sep 17 00:00:00 2001
From: Joseph Redmon <pjreddie@gmail.com>
Date: Mon, 08 Dec 2014 19:48:57 +0000
Subject: [PATCH] Fixed race condition in server
---
src/gemm.c | 76 ++++++++++++++++++++++---------------
1 files changed, 45 insertions(+), 31 deletions(-)
diff --git a/src/gemm.c b/src/gemm.c
index cc882d5..d1782b1 100644
--- a/src/gemm.c
+++ b/src/gemm.c
@@ -104,7 +104,10 @@
#include "opencl.h"
#include <math.h>
-//#include <clBLAS.h>
+
+#ifdef CLBLAS
+#include <clBLAS.h>
+#endif
#define STR_HELPER(x) #x
#define STR(x) STR_HELPER(x)
@@ -165,13 +168,6 @@
float BETA,
cl_mem C_gpu, int ldc)
{
-/*
- cl_setup();
- cl_command_queue queue = cl.queue;
- cl_event event;
- cl.error = clblasSgemm(clblasRowMajor, TA?clblasTrans:clblasNoTrans, TB?clblasTrans:clblasNoTrans,M, N, K,ALPHA, A_gpu, 0, lda,B_gpu, 0, ldb,BETA, C_gpu, 0, ldc,1, &queue, 0, NULL, &event);
- */
-
gemm_ongpu_offset(TA, TB, M, N, K, ALPHA, A_gpu, 0, lda, B_gpu, 0, ldb, BETA, C_gpu, 0, ldc);
}
@@ -181,8 +177,13 @@
float BETA,
cl_mem C_gpu, int c_off, int ldc)
{
+#ifdef CLBLAS
+ cl_command_queue queue = cl.queue;
+ cl_event event;
+ cl.error = clblasSgemm(clblasRowMajor, TA?clblasTrans:clblasNoTrans, TB?clblasTrans:clblasNoTrans,M, N, K,ALPHA, A_gpu, a_off, lda,B_gpu, b_off, ldb,BETA, C_gpu, c_off, ldc,1, &queue, 0, NULL, &event);
+ check_error(cl);
+#else
//printf("gpu: %d %d %d %d %d\n",TA, TB, M, N, K);
- cl_setup();
cl_kernel gemm_kernel = get_gemm_kernel();
if(!TA && !TB) gemm_kernel = get_gemm_nn_kernel();
if(!TA && TB) gemm_kernel = get_gemm_nt_kernel();
@@ -211,8 +212,9 @@
const size_t global_size[] = {ceil((float)N/BLOCK)*BLOCK, ceil((float)M/BLOCK)*BLOCK};
const size_t local_size[] = {BLOCK, BLOCK};
- clEnqueueNDRangeKernel(queue, gemm_kernel, 2, 0, global_size, local_size, 0, 0, 0);
+ cl.error = clEnqueueNDRangeKernel(queue, gemm_kernel, 2, 0, global_size, local_size, 0, 0, 0);
check_error(cl);
+ #endif
}
void gemm_gpu(int TA, int TB, int M, int N, int K, float ALPHA,
@@ -221,7 +223,6 @@
float BETA,
float *C, int ldc)
{
- cl_setup();
cl_context context = cl.context;
cl_command_queue queue = cl.queue;
@@ -284,7 +285,7 @@
void time_ongpu(int TA, int TB, int m, int k, int n)
{
- int iter = 128;
+ int iter = 10;
float *a = random_matrix(m,k);
float *b = random_matrix(k,n);
@@ -302,7 +303,7 @@
for(i = 0; i<iter; ++i){
gemm_ongpu(TA,TB,m,n,k,1,a_cl,lda,b_cl,ldb,1,c_cl,n);
}
- double flop = m*n*(2.*k+3.)*iter;
+ double flop = m*n*k*iter;
double gflop = flop/pow(10., 9);
end = clock();
double seconds = sec(end-start);
@@ -352,32 +353,45 @@
void test_gpu_blas()
{
/*
- test_gpu_accuracy(0,0,10,576,75);
+ test_gpu_accuracy(0,0,10,576,75);
- test_gpu_accuracy(0,0,17,10,10);
- test_gpu_accuracy(1,0,17,10,10);
- test_gpu_accuracy(0,1,17,10,10);
- test_gpu_accuracy(1,1,17,10,10);
+ test_gpu_accuracy(0,0,17,10,10);
+ test_gpu_accuracy(1,0,17,10,10);
+ test_gpu_accuracy(0,1,17,10,10);
+ test_gpu_accuracy(1,1,17,10,10);
- test_gpu_accuracy(0,0,1000,10,100);
- test_gpu_accuracy(1,0,1000,10,100);
- test_gpu_accuracy(0,1,1000,10,100);
- test_gpu_accuracy(1,1,1000,10,100);
- */
+ test_gpu_accuracy(0,0,1000,10,100);
+ test_gpu_accuracy(1,0,1000,10,100);
+ test_gpu_accuracy(0,1,1000,10,100);
+ test_gpu_accuracy(1,1,1000,10,100);
+ */
+ time_ongpu(0,0,512,256,1152);
+ time_ongpu(0,0,128,1200,4096);
+ time_ongpu(0,0,128,1200,4096);
+ time_ongpu(0,0,128,1200,4096);
+
+ time_ongpu(0,1,128,1200,4096);
+ time_ongpu(1,0,1200,4096,128);
+ time_ongpu(1,0,4096,1200,128);
+ time_ongpu(1,0,1200,128,4096);
+
+ test_gpu_accuracy(0,0,512,256,1152);
test_gpu_accuracy(0,0,131,4093,1199);
test_gpu_accuracy(0,1,131,4093,1199);
test_gpu_accuracy(1,0,131,4093,1199);
test_gpu_accuracy(1,1,131,4093,1199);
+ /*
- time_ongpu(0,0,1024,1024,1024);
- time_ongpu(0,1,1024,1024,1024);
- time_ongpu(1,0,1024,1024,1024);
- time_ongpu(1,1,1024,1024,1024);
+ time_ongpu(0,0,1024,1024,1024);
+ time_ongpu(0,1,1024,1024,1024);
+ time_ongpu(1,0,1024,1024,1024);
+ time_ongpu(1,1,1024,1024,1024);
- time_ongpu(0,0,128,4096,1200);
- time_ongpu(0,1,128,4096,1200);
- time_ongpu(1,0,128,4096,1200);
- time_ongpu(1,1,128,4096,1200);
+ time_ongpu(0,0,128,4096,1200);
+ time_ongpu(0,1,128,4096,1200);
+ time_ongpu(1,0,128,4096,1200);
+ time_ongpu(1,1,128,4096,1200);
+ */
/*
time_gpu_random_matrix(0,0,1000,1000,100);
--
Gitblit v1.10.0