Shared Code
main.cu
1234567891011121314151617181920212223242526272829303132333435363738394041424344454647484950515253545556575859606162636465666768697071727374757677787980818283848586878889909192939495969798
#include <stdio.h> #define A(i,j) a[(i)*n+(j)] #define B(i,j) b[(i)*n+(j)] float compare_matrices(int m, int n, float *a, float *b) { int i, j; float max_diff = 0.0, diff; int printed = 0; for (i = 0; i < m; i++) { for (j = 0; j < n; j++) { diff = abs(A(i, j) - B(i, j)); max_diff = (diff > max_diff ? diff : max_diff); if (0 == printed) { if (max_diff > 0.5F || max_diff < -0.5F) { printf("\n error: i %d j %d diff %f got %f expect %f ", i, j, max_diff, A(i, j), B(i, j)); printed = 1; } } } } return max_diff; } void random_matrix(int m,int n, float *a){ for(int i=0;i<m;i++) for(int j=0;j<n;j++) A(i,j)=2.0*(float)drand48()-1.0; } void cpu_sgemm(float *A_ptr,float *B_ptr,float *C_ptr,const int M,const int N,const int K){ for(int m=0;m<M;m++){ for(int n=0;n<N;n++){ float tmp = 0.f; for(int k=0;k<K; k++){ tmp+=A_ptr[K*m+k]*B_ptr[k*N+n]; } C_ptr[m*N+n]=tmp; } } } int main() { int m=2048; int n=2048; int k=2048; const size_t mem_size_A = m*k*sizeof(float); const size_t mem_size_B = k*n*sizeof(float); const size_t mem_size_C = m*n*sizeof(float); float* maxtrix_A_host = (float *)malloc(mem_size_A); float* maxtrix_B_host = (float *)malloc(mem_size_B); float* maxtrix_C_host_gpu_calc = (float *)malloc(mem_size_C); float* maxtrix_C_host_cpu_calc = (float *)malloc(mem_size_C); random_matrix(m,k,maxtrix_A_host); random_matrix(k,n,maxtrix_B_host); memset(maxtrix_C_host_gpu_calc,0,mem_size_C); memset(maxtrix_C_host_cpu_calc,0,mem_size_C); float *maxtrix_A_device,*maxtrix_B_device,*maxtrix_C_device; cudaMalloc((void **)&maxtrix_A_device,mem_size_A); cudaMalloc((void **)&maxtrix_B_device,mem_size_B); cudaMalloc((void **)&maxtrix_C_device,mem_size_C); cudaMemcpy(maxtrix_A_device,maxtrix_A_host,mem_size_A,cudaMemcpyHostToDevice); cudaMemcpy(maxtrix_B_device,maxtrix_B_host,mem_size_B,cudaMemcpyHostToDevice); cpu_sgemm(maxtrix_A_host,maxtrix_B_host,maxtrix_C_host_gpu_calc,m,n,k ); constexpr int BLOCK = 16; dim3 block(BLOCK,BLOCK); dim3 grid((m+BLOCK-1)/BLOCK,(n+BLOCK-1)/BLOCK) cucudaMemcpy(maxtrix_C_host_gpu_calc,maxtrix_C_device,mem_size_C,cucudaMemcpyDeviceToHost); float diff = compare_matrices( m, n, matrix_C_host_gpu_calc, matrix_C_host_cpu_calc); if (diff > 0.5F || diff < -0.5F) { printf("diff too big !\n"); exit(-1); } else { printf("right\n"); } printf("sgemm\n"); return 0; }
Comments (0)
to post comments and vote
No comments yet. Be the first to share your thoughts!