From 028696bf15efeca3acb3db8c42a96f7b9e0f55ff Mon Sep 17 00:00:00 2001
From: iovodov <b@ovdv.ru>
Date: Thu, 03 May 2018 13:33:46 +0000
Subject: [PATCH] Output improvements for detector results: When printing detector results, output was done in random order, obfuscating results for interpreting. Now: 1. Text output includes coordinates of rects in (left,right,top,bottom in pixels) along with label and score 2. Text output is sorted by rect lefts to simplify finding appropriate rects on image 3. If several class probs are > thresh for some detection, the most probable is written first and coordinates for others are not repeated 4. Rects are imprinted in image in order by their best class prob, so most probable rects are always on top and not overlayed by less probable ones 5. Most probable label for rect is always written first Also: 6. Message about low GPU memory include required amount
---
src/gemm.c | 582 ++++++++++++++++++++++++++--------------------------------
1 files changed, 262 insertions(+), 320 deletions(-)
diff --git a/src/gemm.c b/src/gemm.c
index 9797b85..317fd65 100644
--- a/src/gemm.c
+++ b/src/gemm.c
@@ -1,5 +1,66 @@
-#include "mini_blas.h"
+#include "gemm.h"
#include "utils.h"
+#include "cuda.h"
+#include <stdlib.h>
+#include <stdio.h>
+#include <math.h>
+
+void gemm_bin(int M, int N, int K, float ALPHA,
+ char *A, int lda,
+ float *B, int ldb,
+ float *C, int ldc)
+{
+ int i,j,k;
+ for(i = 0; i < M; ++i){
+ for(k = 0; k < K; ++k){
+ char A_PART = A[i*lda+k];
+ if(A_PART){
+ for(j = 0; j < N; ++j){
+ C[i*ldc+j] += B[k*ldb+j];
+ }
+ } else {
+ for(j = 0; j < N; ++j){
+ C[i*ldc+j] -= B[k*ldb+j];
+ }
+ }
+ }
+ }
+}
+
+float *random_matrix(int rows, int cols)
+{
+ int i;
+ float *m = calloc(rows*cols, sizeof(float));
+ for(i = 0; i < rows*cols; ++i){
+ m[i] = (float)rand()/RAND_MAX;
+ }
+ return m;
+}
+
+void time_random_matrix(int TA, int TB, int m, int k, int n)
+{
+ float *a;
+ if(!TA) a = random_matrix(m,k);
+ else a = random_matrix(k,m);
+ int lda = (!TA)?k:m;
+ float *b;
+ if(!TB) b = random_matrix(k,n);
+ else b = random_matrix(n,k);
+ int ldb = (!TB)?n:k;
+
+ float *c = random_matrix(m,n);
+ int i;
+ clock_t start = clock(), end;
+ for(i = 0; i<10; ++i){
+ gemm_cpu(TA,TB,m,n,k,1,a,lda,b,ldb,1,c,n);
+ }
+ end = clock();
+ printf("Matrix Multiplication %dx%d * %dx%d, TA=%d, TB=%d: %lf ms\n",m,k,k,n, TA, TB, (float)(end-start)/CLOCKS_PER_SEC);
+ free(a);
+ free(b);
+ free(c);
+}
+
void gemm(int TA, int TB, int M, int N, int K, float ALPHA,
float *A, int lda,
@@ -10,22 +71,150 @@
gemm_cpu( TA, TB, M, N, K, ALPHA,A,lda, B, ldb,BETA,C,ldc);
}
-void gemm_nn(int M, int N, int K, float ALPHA,
- float *A, int lda,
- float *B, int ldb,
- float *C, int ldc)
+#if (defined(__AVX__) && defined(__x86_64__)) || defined(_WIN64)
+
+#define OSXSAVEFlag (1UL<<27)
+#define AVXFlag ((1UL<<28)|OSXSAVEFlag)
+#define FMAFlag ((1UL<<12)|AVXFlag|OSXSAVEFlag)
+#define CLMULFlag ((1UL<< 1)|AVXFlag|OSXSAVEFlag)
+#define VAESFlag ((1UL<<25)|AVXFlag|OSXSAVEFlag)
+
+#include <stdint.h>
+
+#ifdef _WIN64
+#include <intrin.h>
+#include <ammintrin.h>
+#include <immintrin.h>
+#include <smmintrin.h>
+
+#else // Linux GCC/Clang
+#include <x86intrin.h>
+#include <ammintrin.h>
+#include <immintrin.h>
+#include <smmintrin.h>
+#include <cpuid.h>
+
+void asm_cpuid(uint32_t* abcd, uint32_t eax)
{
- int i,j,k;
- for(i = 0; i < M; ++i){
- for(k = 0; k < K; ++k){
- register float A_PART = ALPHA*A[i*lda+k];
- for(j = 0; j < N; ++j){
- C[i*ldc+j] += A_PART*B[k*ldb+j];
- }
- }
- }
+ uint32_t ebx = 0, edx = 0, ecx = 0;
+
+ // EBX is saved to EDI and later restored
+ __asm__("movl %%ebx, %%edi;"
+ "cpuid;"
+ "xchgl %%ebx, %%edi;"
+ : "=D"(ebx),
+ "+a"(eax), "+c"(ecx), "=d"(edx));
+
+ abcd[0] = eax;
+ abcd[1] = ebx;
+ abcd[2] = ecx;
+ abcd[3] = edx;
}
+#endif
+
+int simd_detect_x86(unsigned int idFeature)
+{
+ uint32_t regs[4]; // EAX, EBX, ECX, EDX;
+#ifdef _WIN32
+ __cpuid(regs, 0);
+ if (regs[0] > 1U) __cpuid(regs, 1);
+#else
+ __get_cpuid(0, ®s[0], ®s[1], ®s[2], ®s[3]);
+ if(regs[0] > 1U) __get_cpuid(1, ®s[0], ®s[1], ®s[2], ®s[3]);
+#endif
+
+ if ((regs[2] & idFeature) != idFeature)
+ return 0;
+ return 1;
+}
+
+int is_fma_avx() {
+ static int result = -1;
+ if (result == -1) {
+ result = simd_detect_x86(AVXFlag);
+ if (result == 1) printf(" Used AVX \n");
+ else printf(" Not used AVX \n");
+ }
+ return result;
+}
+
+// https://software.intel.com/sites/landingpage/IntrinsicsGuide
+void gemm_nn(int M, int N, int K, float ALPHA,
+ float *A, int lda,
+ float *B, int ldb,
+ float *C, int ldc)
+{
+ int i, j, k;
+ if (is_fma_avx() == 1) { // AVX
+ for (i = 0; i < M; ++i) {
+ for (k = 0; k < K; ++k) {
+ float A_PART = ALPHA*A[i*lda + k];
+ __m256 a256, b256, c256, result256; // AVX
+ a256 = _mm256_set1_ps(A_PART);
+ for (j = 0; j < N - 8; j += 8) {
+ b256 = _mm256_loadu_ps(&B[k*ldb + j]);
+ c256 = _mm256_loadu_ps(&C[i*ldc + j]);
+ // FMA - Intel Haswell (2013), AMD Piledriver (2012)
+ //result256 = _mm256_fmadd_ps(a256, b256, c256);
+ result256 = _mm256_mul_ps(a256, b256);
+ result256 = _mm256_add_ps(result256, c256);
+ _mm256_storeu_ps(&C[i*ldc + j], result256);
+ }
+
+ int prev_end = (N % 8 == 0) ? (N - 8) : (N / 8) * 8;
+ for (j = prev_end; j < N; ++j)
+ C[i*ldc + j] += A_PART*B[k*ldb + j];
+ }
+ }
+ }
+ else {
+ for (i = 0; i < M; ++i) {
+ for (k = 0; k < K; ++k) {
+ register float A_PART = ALPHA*A[i*lda + k];
+ for (j = 0; j < N; ++j) {
+ C[i*ldc + j] += A_PART*B[k*ldb + j];
+ }
+ /* // SSE
+ __m128 a128, b128, c128, result128; // SSE
+ a128 = _mm_set1_ps(A_PART);
+ for (j = 0; j < N - 4; j += 4) {
+ b128 = _mm_loadu_ps(&B[k*ldb + j]);
+ c128 = _mm_loadu_ps(&C[i*ldc + j]);
+ //result128 = _mm_fmadd_ps(a128, b128, c128);
+ result128 = _mm_mul_ps(a128, b128);
+ result128 = _mm_add_ps(result128, c128);
+ _mm_storeu_ps(&C[i*ldc + j], result128);
+ }
+
+ int prev_end = (N % 4 == 0) ? (N - 4) : (N / 4) * 4;
+ for (j = prev_end; j < N; ++j){
+ C[i*ldc + j] += A_PART*B[k*ldb + j];
+ }
+ */
+ }
+ }
+ }
+}
+#else
+
+void gemm_nn(int M, int N, int K, float ALPHA,
+ float *A, int lda,
+ float *B, int ldb,
+ float *C, int ldc)
+{
+ int i, j, k;
+ for (i = 0; i < M; ++i) {
+ for (k = 0; k < K; ++k) {
+ register float A_PART = ALPHA*A[i*lda + k];
+ for (j = 0; j < N; ++j) {
+ C[i*ldc + j] += A_PART*B[k*ldb + j];
+ }
+ }
+ }
+}
+#endif // __x86_64
+
void gemm_nt(int M, int N, int K, float ALPHA,
float *A, int lda,
float *B, int ldb,
@@ -90,188 +279,36 @@
C[i*ldc + j] *= BETA;
}
}
- if(!TA && !TB)
- gemm_nn(M, N, K, ALPHA,A,lda, B, ldb,C,ldc);
- else if(TA && !TB)
- gemm_tn(M, N, K, ALPHA,A,lda, B, ldb,C,ldc);
- else if(!TA && TB)
- gemm_nt(M, N, K, ALPHA,A,lda, B, ldb,C,ldc);
- else
- gemm_tt(M, N, K, ALPHA,A,lda, B, ldb,C,ldc);
+
+ int t;
+ #pragma omp parallel for
+ for (t = 0; t < M; ++t) {
+ if (!TA && !TB)
+ gemm_nn(1, N, K, ALPHA, A + t*lda, lda, B, ldb, C + t*ldc, ldc);
+ else if (TA && !TB)
+ gemm_tn(1, N, K, ALPHA, A + t, lda, B, ldb, C + t*ldc, ldc);
+ else if (!TA && TB)
+ gemm_nt(1, N, K, ALPHA, A + t*lda, lda, B, ldb, C + t*ldc, ldc);
+ else
+ gemm_tt(1, N, K, ALPHA, A + t, lda, B, ldb, C + t*ldc, ldc);
+ }
}
#ifdef GPU
-#include "opencl.h"
#include <math.h>
-#ifdef CLBLAS
-#include <clBLAS.h>
-#endif
-
-#define STR_HELPER(x) #x
-#define STR(x) STR_HELPER(x)
-
-#ifdef __APPLE__
-#define BLOCK 1
-#else
-#define BLOCK 16
-#endif
-
-cl_kernel get_gemm_kernel()
-{
- static int init = 0;
- static cl_kernel gemm_kernel;
- if(!init){
- gemm_kernel = get_kernel("src/gemm.cl", "gemm", "-D BLOCK=" STR(BLOCK) );
- init = 1;
- }
- return gemm_kernel;
-}
-
-cl_kernel get_gemm_nt_kernel()
-{
- static int init = 0;
- static cl_kernel gemm_kernel;
- if(!init){
- gemm_kernel = get_kernel("src/gemm.cl", "gemm_nt", "-D BLOCK=" STR(BLOCK) );
- init = 1;
- }
- return gemm_kernel;
-}
-
-cl_kernel get_gemm_tn_kernel()
-{
- static int init = 0;
- static cl_kernel gemm_kernel;
- if(!init){
- gemm_kernel = get_kernel("src/gemm.cl", "gemm_tn", "-D BLOCK=" STR(BLOCK) );
- init = 1;
- }
- return gemm_kernel;
-}
-
-cl_kernel get_gemm_nn_kernel()
-{
- static int init = 0;
- static cl_kernel gemm_kernel;
- if(!init){
- gemm_kernel = get_kernel("src/gemm.cl", "gemm_nn", "-D BLOCK=" STR(BLOCK) );
- init = 1;
- }
- return gemm_kernel;
-}
-
-#define TILE 64
-#define TILE_K 16
-#define THREADS 64
-
-cl_kernel get_gemm_nn_fast_kernel()
-{
- static int init = 0;
- static cl_kernel gemm_kernel;
- if(!init){
- gemm_kernel = get_kernel("src/gemm_fast.cl", "gemm_nn_fast", "-D TILE=" STR(TILE)
- " -cl-nv-verbose "
- " -D TILE_K=" STR(TILE_K)
- " -D THREADS=" STR(THREADS));
- init = 1;
- }
- return gemm_kernel;
-}
-
void gemm_ongpu(int TA, int TB, int M, int N, int K, float ALPHA,
- cl_mem A_gpu, int lda,
- cl_mem B_gpu, int ldb,
+ float *A_gpu, int lda,
+ float *B_gpu, int ldb,
float BETA,
- cl_mem C_gpu, int ldc)
+ float *C_gpu, int ldc)
{
- gemm_ongpu_offset(TA, TB, M, N, K, ALPHA, A_gpu, 0, lda, B_gpu, 0, ldb, BETA, C_gpu, 0, ldc);
-}
-
-void gemm_ongpu_fast(int TA, int TB, int M, int N, int K, float ALPHA,
- cl_mem A_gpu, int lda,
- cl_mem B_gpu, int ldb,
- float BETA,
- cl_mem C_gpu, int ldc)
-{
- int a_off = 0;
- int b_off = 0;
- int c_off = 0;
- //printf("gpu: %d %d %d %d %d\n",TA, TB, M, N, K);
- cl_kernel gemm_kernel = get_gemm_nn_fast_kernel();
- cl_command_queue queue = cl.queue;
-
- cl_uint i = 0;
- cl.error = clSetKernelArg(gemm_kernel, i++, sizeof(TA), (void*) &TA);
- cl.error = clSetKernelArg(gemm_kernel, i++, sizeof(TB), (void*) &TB);
- cl.error = clSetKernelArg(gemm_kernel, i++, sizeof(M), (void*) &M);
- cl.error = clSetKernelArg(gemm_kernel, i++, sizeof(N), (void*) &N);
- cl.error = clSetKernelArg(gemm_kernel, i++, sizeof(K), (void*) &K);
- cl.error = clSetKernelArg(gemm_kernel, i++, sizeof(ALPHA), (void*) &ALPHA);
- cl.error = clSetKernelArg(gemm_kernel, i++, sizeof(A_gpu), (void*) &A_gpu);
- cl.error = clSetKernelArg(gemm_kernel, i++, sizeof(a_off), (void*) &a_off);
- cl.error = clSetKernelArg(gemm_kernel, i++, sizeof(lda), (void*) &lda);
- cl.error = clSetKernelArg(gemm_kernel, i++, sizeof(B_gpu), (void*) &B_gpu);
- cl.error = clSetKernelArg(gemm_kernel, i++, sizeof(b_off), (void*) &b_off);
- cl.error = clSetKernelArg(gemm_kernel, i++, sizeof(ldb), (void*) &ldb);
- cl.error = clSetKernelArg(gemm_kernel, i++, sizeof(BETA), (void*) &BETA);
- cl.error = clSetKernelArg(gemm_kernel, i++, sizeof(C_gpu), (void*) &C_gpu);
- cl.error = clSetKernelArg(gemm_kernel, i++, sizeof(c_off), (void*) &c_off);
- cl.error = clSetKernelArg(gemm_kernel, i++, sizeof(ldc), (void*) &ldc);
- check_error(cl);
-
- const size_t global_size[] = {THREADS*((N-1)/TILE + 1), (M-1)/TILE + 1};
- const size_t local_size[] = {THREADS, 1};
-
- cl.error = clEnqueueNDRangeKernel(queue, gemm_kernel, 2, 0, global_size, local_size, 0, 0, 0);
- check_error(cl);
-}
-
-void gemm_ongpu_offset(int TA, int TB, int M, int N, int K, float ALPHA,
- cl_mem A_gpu, int a_off, int lda,
- cl_mem B_gpu, int b_off, int ldb,
- 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_kernel gemm_kernel = get_gemm_kernel();
- if(!TA && !TB) gemm_kernel = get_gemm_nn_kernel();
- if(!TA && TB) gemm_kernel = get_gemm_nt_kernel();
- if(TA && !TB) gemm_kernel = get_gemm_tn_kernel();
- cl_command_queue queue = cl.queue;
-
- cl_uint i = 0;
- cl.error = clSetKernelArg(gemm_kernel, i++, sizeof(TA), (void*) &TA);
- cl.error = clSetKernelArg(gemm_kernel, i++, sizeof(TB), (void*) &TB);
- cl.error = clSetKernelArg(gemm_kernel, i++, sizeof(M), (void*) &M);
- cl.error = clSetKernelArg(gemm_kernel, i++, sizeof(N), (void*) &N);
- cl.error = clSetKernelArg(gemm_kernel, i++, sizeof(K), (void*) &K);
- cl.error = clSetKernelArg(gemm_kernel, i++, sizeof(ALPHA), (void*) &ALPHA);
- cl.error = clSetKernelArg(gemm_kernel, i++, sizeof(A_gpu), (void*) &A_gpu);
- cl.error = clSetKernelArg(gemm_kernel, i++, sizeof(a_off), (void*) &a_off);
- cl.error = clSetKernelArg(gemm_kernel, i++, sizeof(lda), (void*) &lda);
- cl.error = clSetKernelArg(gemm_kernel, i++, sizeof(B_gpu), (void*) &B_gpu);
- cl.error = clSetKernelArg(gemm_kernel, i++, sizeof(b_off), (void*) &b_off);
- cl.error = clSetKernelArg(gemm_kernel, i++, sizeof(ldb), (void*) &ldb);
- cl.error = clSetKernelArg(gemm_kernel, i++, sizeof(BETA), (void*) &BETA);
- cl.error = clSetKernelArg(gemm_kernel, i++, sizeof(C_gpu), (void*) &C_gpu);
- cl.error = clSetKernelArg(gemm_kernel, i++, sizeof(c_off), (void*) &c_off);
- cl.error = clSetKernelArg(gemm_kernel, i++, sizeof(ldc), (void*) &ldc);
- check_error(cl);
-
- const size_t global_size[] = {ceil((float)N/BLOCK)*BLOCK, ceil((float)M/BLOCK)*BLOCK};
- const size_t local_size[] = {BLOCK, BLOCK};
-
- cl.error = clEnqueueNDRangeKernel(queue, gemm_kernel, 2, 0, global_size, local_size, 0, 0, 0);
- check_error(cl);
-#endif
+ cublasHandle_t handle = blas_handle();
+ cudaError_t stream_status = cublasSetStream(handle, get_cuda_stream());
+ cudaError_t status = cublasSgemm(handle, (TB ? CUBLAS_OP_T : CUBLAS_OP_N),
+ (TA ? CUBLAS_OP_T : CUBLAS_OP_N), N, M, K, &ALPHA, B_gpu, ldb, A_gpu, lda, &BETA, C_gpu, ldc);
+ check_error(status);
}
void gemm_gpu(int TA, int TB, int M, int N, int K, float ALPHA,
@@ -280,37 +317,16 @@
float BETA,
float *C, int ldc)
{
- cl_context context = cl.context;
- cl_command_queue queue = cl.queue;
+ float *A_gpu = cuda_make_array(A, (TA ? lda*K:lda*M));
+ float *B_gpu = cuda_make_array(B, (TB ? ldb*N : ldb*K));
+ float *C_gpu = cuda_make_array(C, ldc*M);
- size_t size = sizeof(float)*(TA ? lda*K:lda*M);
- cl_mem A_gpu = clCreateBuffer(context,
- CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR,
- size, A, &cl.error);
- check_error(cl);
+ gemm_ongpu(TA, TB, M, N, K, ALPHA, A_gpu, lda, B_gpu, ldb, BETA, C_gpu, ldc);
- size = sizeof(float)*(TB ? ldb*N:ldb*K);
- cl_mem B_gpu = clCreateBuffer(context,
- CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR,
- size, B, &cl.error);
- check_error(cl);
-
- size = sizeof(float)*(ldc*M);
- cl_mem C_gpu = clCreateBuffer(context,
- CL_MEM_READ_WRITE|CL_MEM_COPY_HOST_PTR,
- size, C, &cl.error);
- check_error(cl);
-
- // TODO
- //gemm_ongpu(TA, TB, M, N, K, ALPHA, A_gpu, lda, B_gpu, ldb, BETA, C_gpu, ldc);
- gemm_ongpu_fast(TA, TB, M, N, K, ALPHA, A_gpu, lda, B_gpu, ldb, BETA, C_gpu, ldc);
-
- clEnqueueReadBuffer(queue, C_gpu, CL_TRUE, 0, size, C, 0, 0, 0);
- check_error(cl);
-
- clReleaseMemObject(A_gpu);
- clReleaseMemObject(B_gpu);
- clReleaseMemObject(C_gpu);
+ cuda_pull_array(C_gpu, C, ldc*M);
+ cuda_free(A_gpu);
+ cuda_free(B_gpu);
+ cuda_free(C_gpu);
}
#include <stdio.h>
@@ -353,60 +369,29 @@
float *c = random_matrix(m,n);
- cl_mem a_cl = cl_make_array(a, m*k);
- cl_mem b_cl = cl_make_array(b, k*n);
- cl_mem c_cl = cl_make_array(c, m*n);
+ float *a_cl = cuda_make_array(a, m*k);
+ float *b_cl = cuda_make_array(b, k*n);
+ float *c_cl = cuda_make_array(c, m*n);
int i;
clock_t start = clock(), end;
for(i = 0; i<iter; ++i){
gemm_ongpu(TA,TB,m,n,k,1,a_cl,lda,b_cl,ldb,1,c_cl,n);
+ cudaThreadSynchronize();
}
double flop = ((double)m)*n*(2.*k + 2.)*iter;
double gflop = flop/pow(10., 9);
end = clock();
double seconds = sec(end-start);
printf("Matrix Multiplication %dx%d * %dx%d, TA=%d, TB=%d: %lf s, %lf GFLOPS\n",m,k,k,n, TA, TB, seconds, gflop/seconds);
- clReleaseMemObject(a_cl);
- clReleaseMemObject(b_cl);
- clReleaseMemObject(c_cl);
+ cuda_free(a_cl);
+ cuda_free(b_cl);
+ cuda_free(c_cl);
free(a);
free(b);
free(c);
}
-void time_ongpu_fast(int TA, int TB, int m, int k, int n)
-{
- int iter = 10;
- float *a = random_matrix(m,k);
- float *b = random_matrix(k,n);
-
- int lda = (!TA)?k:m;
- int ldb = (!TB)?n:k;
-
- float *c = random_matrix(m,n);
-
- cl_mem a_cl = cl_make_array(a, m*k);
- cl_mem b_cl = cl_make_array(b, k*n);
- cl_mem c_cl = cl_make_array(c, m*n);
-
- int i;
- clock_t start = clock(), end;
- for(i = 0; i<iter; ++i){
- gemm_ongpu_fast(TA,TB,m,n,k,1,a_cl,lda,b_cl,ldb,1,c_cl,n);
- }
- double flop = ((double)m)*n*(2.*k + 2.)*iter;
- double gflop = flop/pow(10., 9);
- end = clock();
- double seconds = sec(end-start);
- printf("Fast Multiplication %dx%d * %dx%d, TA=%d, TB=%d: %lf s, %lf GFLOPS\n",m,k,k,n, TA, TB, seconds, gflop/seconds);
- clReleaseMemObject(a_cl);
- clReleaseMemObject(b_cl);
- clReleaseMemObject(c_cl);
- free(a);
- free(b);
- free(c);
-}
void test_gpu_accuracy(int TA, int TB, int m, int k, int n)
{
@@ -429,6 +414,7 @@
gemm_gpu(TA,TB,m,n,k,1,a,lda,b,ldb,1,c_gpu,n);
//printf("GPU\n");
//pm(m, n, c_gpu);
+
gemm_cpu(TA,TB,m,n,k,1,a,lda,b,ldb,1,c,n);
//printf("\n\nCPU\n");
//pm(m, n, c);
@@ -444,7 +430,7 @@
free(c_gpu);
}
-void test_gpu_blas()
+int test_gpu_blas()
{
/*
test_gpu_accuracy(0,0,10,576,75);
@@ -458,73 +444,29 @@
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,10,10,10);
+
+ time_ongpu(0,0,64,2916,363);
+ time_ongpu(0,0,64,2916,363);
+ time_ongpu(0,0,64,2916,363);
+ time_ongpu(0,0,192,729,1600);
+ time_ongpu(0,0,384,196,1728);
+ time_ongpu(0,0,256,196,3456);
+ time_ongpu(0,0,256,196,2304);
+ time_ongpu(0,0,128,4096,12544);
+ time_ongpu(0,0,128,4096,4096);
*/
+ time_ongpu(0,0,64,75,12544);
+ time_ongpu(0,0,64,75,12544);
+ time_ongpu(0,0,64,75,12544);
+ time_ongpu(0,0,64,576,12544);
+ time_ongpu(0,0,256,2304,784);
+ time_ongpu(1,1,2304,256,784);
+ time_ongpu(0,0,512,4608,196);
+ time_ongpu(1,1,4608,512,196);
- test_gpu_accuracy(0,0,128,128,128);
-
- time_ongpu(0,0,64,2916,363);
- time_ongpu_fast(0,0,64,2916,363);
- time_ongpu(0,0,64,2916,363);
- time_ongpu_fast(0,0,64,2916,363);
- time_ongpu(0,0,64,2916,363);
- time_ongpu_fast(0,0,64,2916,363);
- time_ongpu(0,0,192,729,1600);
- time_ongpu_fast(0,0,192,729,1600);
- time_ongpu(0,0,384,196,1728);
- time_ongpu_fast(0,0,384,196,1728);
- time_ongpu(0,0,256,196,3456);
- time_ongpu_fast(0,0,256,196,3456);
- time_ongpu(0,0,256,196,2304);
- time_ongpu_fast(0,0,256,196,2304);
- time_ongpu(0,0,128,4096,12544);
- time_ongpu_fast(0,0,128,4096,12544);
- time_ongpu(0,0,128,4096,4096);
- time_ongpu_fast(0,0,128,4096,4096);
-// time_ongpu(1,0,2304,196,256);
-// time_ongpu_fast(1,0,2304,196,256);
-// time_ongpu(0,1,256,2304,196);
-// time_ongpu_fast(0,1,256,2304,196);
-
- time_ongpu(0,0,2048,2048,2048);
- time_ongpu_fast(0,0,2048,2048,2048);
- time_ongpu(0,0,2048,2048,2048);
- time_ongpu_fast(0,0,2048,2048,2048);
- time_ongpu(0,0,2048,2048,2048);
- time_ongpu_fast(0,0,2048,2048,2048);
-
- /*
- 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,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);
- time_random_matrix(0,0,1000,1000,100);
-
- time_gpu_random_matrix(0,1,1000,1000,100);
- time_random_matrix(0,1,1000,1000,100);
-
- time_gpu_random_matrix(1,0,1000,1000,100);
- time_random_matrix(1,0,1000,1000,100);
-
- time_gpu_random_matrix(1,1,1000,1000,100);
- time_random_matrix(1,1,1000,1000,100);
- */
-
+ return 0;
}
#endif
--
Gitblit v1.10.0