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 | 7 +++----
1 files changed, 3 insertions(+), 4 deletions(-)
diff --git a/src/gemm.c b/src/gemm.c
index edffcaf..d1782b1 100644
--- a/src/gemm.c
+++ b/src/gemm.c
@@ -178,14 +178,12 @@
cl_mem C_gpu, int c_off, int ldc)
{
#ifdef CLBLAS
- 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, 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();
@@ -214,7 +212,7 @@
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
}
@@ -225,7 +223,6 @@
float BETA,
float *C, int ldc)
{
- cl_setup();
cl_context context = cl.context;
cl_command_queue queue = cl.queue;
@@ -368,6 +365,7 @@
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);
@@ -377,6 +375,7 @@
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);
--
Gitblit v1.10.0