#include #include "utils.cuh" #include "kernel.cuh" float get_sec() { struct timeval time; gettimeofday(&time, NULL); return (1e6 * time.tv_sec + time.tv_usec); } float cpu_elapsed_time(float &beg, float &end) { return 1.0e-6 * (end - beg); } void cudaCheck(cudaError_t error, const char *file, int line) { if (error != cudaSuccess) { printf("[CUDA ERROR] at file %s(line %d):\n%s\n", file, line, cudaGetErrorString(error)); exit(EXIT_FAILURE); } return; }; void CudaDeviceInfo() { int deviceId; cudaGetDevice(&deviceId); cudaDeviceProp props; cudaGetDeviceProperties(&props, deviceId); /* * There should be no need to modify the output string below. */ printf("Device ID: %d\n\ *Number of SMs: %d\n\ Compute Capability Major: %d\n\ Compute Capability Minor: %d\n\ memoryBusWidth: %d\n\ *maxThreadsPerBlock: %d\n\ maxThreadsPerMultiProcessor: %d\n\ *totalGlobalMem: %zuM\n\ sharedMemPerBlock: %zuKB\n\ *sharedMemPerMultiprocessor: %zuKB\n\ totalConstMem: %zuKB\n\ *multiProcessorCount: %d\n\ *Warp Size: %d\n", deviceId, props.multiProcessorCount, props.major, props.minor, props.memoryBusWidth, props.maxThreadsPerBlock, props.maxThreadsPerMultiProcessor, props.totalGlobalMem / 1024 / 1024, props.sharedMemPerBlock / 1024, props.sharedMemPerMultiprocessor / 1024, props.totalConstMem / 1024, props.multiProcessorCount, props.warpSize); }; void randomize_matrix(float *mat, int N) { // NOTICE: 使用gettimeofdays替代srand((unsigned)time(NULL));time精度过低,产生相同随机数 struct timeval time; gettimeofday(&time, NULL); srand(time.tv_usec); for (int i = 0; i < N; i++) { float tmp = (float) (rand() % 5) + 0.01 * (rand() % 5); tmp = (rand() % 2 == 0) ? tmp : tmp * (-1.); mat[i] = tmp; } } void copy_matrix(float *src, float *dest, int N) { int i; for (i = 0; src + i && dest + i && i < N; i++) *(dest + i) = *(src + i); if (i != N) printf("copy failed at %d while there are %d elements in total.\n", i, N); } void print_matrix(const float *A, int M, int N) { int i; printf("["); for (i = 0; i < M * N; i++) { if ((i + 1) % N == 0) printf("%5.2f ", A[i]); else printf("%5.2f, ", A[i]); if ((i + 1) % N == 0) { if (i + 1 < M * N) printf(";\n"); } } printf("]\n"); } bool verify_matrix(float *mat1, float *mat2, int N) { double diff = 0.0; int i; for (i = 0; mat1 + i && mat2 + i && i < N; i++) { diff = fabs((double) mat1[i] - (double) mat2[i]); if (diff > 1e-2) { printf("error. %5.2f,%5.2f,%d\n", mat1[i], mat2[i], i); return false; } } return true; } #define CEIL_DIV(M, N) ((M) + (N)-1) / (N) void test_cublas(cublasHandle_t handle, int M, int N, int K, float alpha, float *A, float *B, float beta, float *C) { //cublas列主序计算:https://www.cnblogs.com/cuancuancuanhao/p/7763256.html cublasSgemm(handle, CUBLAS_OP_N, CUBLAS_OP_N, N, M, K, &alpha, B, N, A, K, &beta, C, N); } void test_mysgemm_v1(int M, int N, int K, float alpha, float *A, float *B, float beta, float *C) { dim3 blockDim(32, 32); dim3 gridDim(CEIL_DIV(M, 32), CEIL_DIV(N, 32)); mysgemm_v1<<>>(M, N, K, alpha, A, B, beta, C); } void test_mysgemm_v2(int M, int N, int K, float alpha, float *A, float *B, float beta, float *C) { dim3 blockDim(1024); dim3 gridDim(CEIL_DIV(M, 32), CEIL_DIV(N, 32)); mysgemm_v2<32><<>>(M, N, K, alpha, A, B, beta, C); } void test_mysgemm_v3(int M, int N, int K, float alpha, float *A, float *B, float beta, float *C) { dim3 blockDim(512); dim3 gridDim(CEIL_DIV(M, 64), CEIL_DIV(N, 64)); mysgemm_v3<64, 64, 8, 8><<>>(M, N, K, alpha, A, B, beta, C); } void test_mysgemm_v4(int M, int N, int K, float alpha, float *A, float *B, float beta, float *C) { dim3 blockDim(256); dim3 gridDim(CEIL_DIV(M, 128), CEIL_DIV(N, 128)); mysgemm_v4<128, 128, 8, 8, 8><<>>(M, N, K, alpha, A, B, beta, C); } void test_mysgemm_v5(int M, int N, int K, float alpha, float *A, float *B, float beta, float *C) { dim3 blockDim(256); dim3 gridDim(CEIL_DIV(M, 128), CEIL_DIV(N, 128)); mysgemm_v5<128, 128, 8, 8, 8><<>>(M, N, K, alpha, A, B, beta, C); } //void test_mysgemm_v6(int M, int N, int K, float alpha, float *A, float *B, float beta, float *C) { // dim3 blockDim(4); // dim3 gridDim(CEIL_DIV(M, 8), CEIL_DIV(N, 8)); // mysgemm_v6<8, 8, 4, 4, 4><<>>(M, N, K, alpha, A, B, beta, C); //} void test_mysgemm_v6(int M, int N, int K, float alpha, float *A, float *B, float beta, float *C) { dim3 blockDim(256); dim3 gridDim(CEIL_DIV(M, 128), CEIL_DIV(N, 128)); mysgemm_v6<128, 128, 8, 8, 8><<>>(M, N, K, alpha, A, B, beta, C); } void test_mysgemm_v7(int M, int N, int K, float alpha, float *A, float *B, float beta, float *C) { dim3 blockDim(256); dim3 gridDim(CEIL_DIV(M, 128), CEIL_DIV(N, 128)); mysgemm_v7<128, 128, 8, 8, 8><<>>(M, N, K, alpha, A, B, beta, C); } void test_kernel(int kernel_num, int M, int N, int K, float alpha, float *A, float *B, float beta, float *C, cublasHandle_t handle) { switch (kernel_num) { case 0: test_cublas(handle, M, N, K, alpha, A, B, beta, C); break; case 1: test_mysgemm_v1(M, N, K, alpha, A, B, beta, C); break; case 2: test_mysgemm_v2(M, N, K, alpha, A, B, beta, C); break; case 3: test_mysgemm_v3(M, N, K, alpha, A, B, beta, C); break; case 4: test_mysgemm_v4(M, N, K, alpha, A, B, beta, C); break; case 5: test_mysgemm_v5(M, N, K, alpha, A, B, beta, C); break; case 6: test_mysgemm_v6(M, N, K, alpha, A, B, beta, C); break; case 7: test_mysgemm_v7(M, N, K, alpha, A, B, beta, C); break; default: break; } }