From 708031c0274a730dfd99820fd49351785a60e2d7 Mon Sep 17 00:00:00 2001 From: aamir Date: Sat, 2 Jun 2018 19:51:50 -0700 Subject: mma working for type32_32 --- cuda-kernels/Makefile | 4 +- cuda-kernels/tensor_core.cu | 57 +++++----- cuda-kernels/tensorcore_type32_32.cu | 202 +++++++++++++++++++++++++++++++++++ 3 files changed, 230 insertions(+), 33 deletions(-) create mode 100644 cuda-kernels/tensorcore_type32_32.cu (limited to 'cuda-kernels') diff --git a/cuda-kernels/Makefile b/cuda-kernels/Makefile index 673460f..73a4f0c 100755 --- a/cuda-kernels/Makefile +++ b/cuda-kernels/Makefile @@ -1,5 +1,5 @@ -all: tensor_core.cu - nvcc --gpu-architecture=compute_70 --gpu-code=compute_70 -lcudart -g -o tensor_core tensor_core.cu +all: tensorcore_type32_32.cu + nvcc --gpu-architecture=compute_70 --gpu-code=compute_70 -lcudart -g -o tensor_core tensorcore_type32_32.cu # nvcc -arch=sm_70 -lcudart -g -o tensor_core tensor_core.cu .PHONY: diff --git a/cuda-kernels/tensor_core.cu b/cuda-kernels/tensor_core.cu index d5c1b40..b7090c4 100644 --- a/cuda-kernels/tensor_core.cu +++ b/cuda-kernels/tensor_core.cu @@ -132,11 +132,7 @@ int main(int argc, char* argv[]) { cudaErrCheck(cudaEventCreate(&startWMMA)); cudaErrCheck(cudaEventCreate(&stopWMMA)); - - - // Use tensor cores - cudaErrCheck(cudaMalloc((void**)&a_fp32, MATRIX_M * MATRIX_K * sizeof(float))); cudaErrCheck(cudaMalloc((void**)&b_fp32, MATRIX_K * MATRIX_N * sizeof(float))); @@ -151,8 +147,6 @@ int main(int argc, char* argv[]) { a_host_wmma = (float*)malloc(MATRIX_M * MATRIX_K * sizeof(float)); b_host_wmma = (float*)malloc(MATRIX_K * MATRIX_N * sizeof(float)); - - // printf("a_fp32\n"); for(int m=0;m +#include + +// Define some error checking macros. +#define cudaErrCheck(stat) { cudaErrCheck_((stat), __FILE__, __LINE__); } +void cudaErrCheck_(cudaError_t stat, const char *file, int line) { + if (stat != cudaSuccess) { + fprintf(stderr, "CUDA Error: %s %s %d\n", cudaGetErrorString(stat), file, line); + } +} + +#define curandErrCheck(stat) { curandErrCheck_((stat), __FILE__, __LINE__); } +void curandErrCheck_(curandStatus_t stat, const char *file, int line) { + if (stat != CURAND_STATUS_SUCCESS) { + fprintf(stderr, "cuRand Error: %d %s %d\n", stat, file, line); + } +} + +#include +using namespace nvcuda; + +// Must be multiples of 16 for wmma code to work +#define MATRIX_M (16) +#define MATRIX_N (16) +#define MATRIX_K (16) + + +// The only dimensions currently supported by WMMA +const int WMMA_M = 16; +const int WMMA_N = 16; +const int WMMA_K = 16; + +__global__ void wmma_example(half *a, half *b, float *c,float *d_fp16, int M, int N, int K) { + //unsigned int start_time=0,end_time=0; + //start_time=clock(); + + // Declare the fragments + wmma::fragment a_frag; + wmma::fragment b_frag; + wmma::fragment c_frag; + + // Bounds checking + wmma::load_matrix_sync(a_frag, a, K); + wmma::load_matrix_sync(b_frag, b, K); + wmma::load_matrix_sync(c_frag, c, N,wmma::mem_col_major); + wmma::mma_sync(c_frag, a_frag, b_frag, c_frag); + + wmma::store_matrix_sync(d_fp16, c_frag, N, wmma::mem_col_major); + //printf("clock=%d",end_time-start_time); +} + +__global__ void convertFp32ToFp16 (half *out, float *in, int n) { + int idx = blockDim.x * blockIdx.x + threadIdx.x; + if (idx < n) { + out[idx] = in[idx]; + } +} +__global__ void convertFp16ToFp32 (float *out, half *in, int n) { + int idx = blockDim.x * blockIdx.x + threadIdx.x; + if (idx < n) { + out[idx] = in[idx]; + } +} + +int main(int argc, char* argv[]) { + float *a_fp32; + float *b_fp32; + float *c_fp32; + float *d_fp32; + + half *a_fp16; + half *b_fp16; + // half *c_fp16; + // half *d_fp16; + + float *a_host_wmma; + float *b_host_wmma; + float *c_host_wmma; + float *d_host_wmma; + float *d_cal_host_wmma; + + cudaEvent_t startWMMA; + cudaEvent_t stopWMMA; + + + cudaErrCheck(cudaEventCreate(&startWMMA)); + cudaErrCheck(cudaEventCreate(&stopWMMA)); + + // Use tensor cores + cudaErrCheck(cudaMalloc((void**)&a_fp32, MATRIX_M * MATRIX_K * sizeof(float))); + cudaErrCheck(cudaMalloc((void**)&b_fp32, MATRIX_K * MATRIX_N * sizeof(float))); + cudaErrCheck(cudaMalloc((void**)&c_fp32, MATRIX_K * MATRIX_N * sizeof(float))); + cudaErrCheck(cudaMalloc((void**)&d_fp32, MATRIX_K * MATRIX_N * sizeof(float))); + cudaErrCheck(cudaMalloc((void**)&a_fp16, MATRIX_M * MATRIX_K * sizeof(half))); + cudaErrCheck(cudaMalloc((void**)&b_fp16, MATRIX_K * MATRIX_N * sizeof(half))); + //cudaErrCheck(cudaMalloc((void**)&c_fp16, MATRIX_K * MATRIX_N * sizeof(half))); + //cudaErrCheck(cudaMalloc((void**)&d_fp16, MATRIX_K * MATRIX_N * sizeof(half))); + + + a_host_wmma = (float*)malloc(MATRIX_M * MATRIX_K * sizeof(float)); + b_host_wmma = (float*)malloc(MATRIX_K * MATRIX_N * sizeof(float)); + c_host_wmma = (float*)malloc(MATRIX_M * MATRIX_N * sizeof(float)); + d_host_wmma = (float*)malloc(MATRIX_M * MATRIX_N * sizeof(float)); + d_cal_host_wmma = (float*)malloc(MATRIX_M * MATRIX_N * sizeof(float)); + + //printf("a_fp32\n"); + for(int m=0;m>> (a_fp16, a_fp32, MATRIX_M * MATRIX_K); + convertFp32ToFp16 <<< (MATRIX_K * MATRIX_N + 255) / 256, 256 >>> (b_fp16, b_fp32, MATRIX_K * MATRIX_N); + //convertFp32ToFp16 <<< (MATRIX_M * MATRIX_N + 255) / 256, 256 >>> (c_fp16, c_fp32, MATRIX_K * MATRIX_N); + + printf("\nM = %d, N = %d, K = %d. \n", MATRIX_M, MATRIX_N, MATRIX_K); + + printf("Running with wmma...\n"); + cudaErrCheck(cudaEventRecord(startWMMA)); + wmma_example <<< 1, 32>>> (a_fp16, b_fp16, c_fp32, d_fp32 , MATRIX_M, MATRIX_N, MATRIX_K); + cudaErrCheck(cudaEventRecord(stopWMMA)); + cudaErrCheck(cudaEventSynchronize(stopWMMA)); + + // Error checking + printf("\nChecking results...\n"); + cudaErrCheck(cudaMemcpy(d_host_wmma, d_fp32, MATRIX_M * MATRIX_N * sizeof(float), cudaMemcpyDeviceToHost)); + + printf("Results verified: cublas and WMMA agree.\n\n"); + float wmmaTime; + cudaErrCheck(cudaEventElapsedTime(&wmmaTime, startWMMA, stopWMMA)); + printf("wmma took %fms\n", wmmaTime); + + cudaErrCheck(cudaEventDestroy(startWMMA)); + cudaErrCheck(cudaEventDestroy(stopWMMA)); + + printf("D_CALCULATED\n"); + for(int m=0;m