From 7aad531c84663bf4da9b2f4112844e2ef284fcf5 Mon Sep 17 00:00:00 2001 From: Tao Sang Date: Wed, 26 May 2021 23:14:45 -0400 Subject: [PATCH] SWDEV-276853 - Add kernelVerify Add kernelVerify for data verification and memory reading performance checking in kernel. Change-Id: Id3f9bcad75d643f493daf9d5f47b3a012a427179 [ROCm/clr commit: f7e3347215165213f0f2482a91dab280705180b6] --- .../performance/memory/hipPerfMemFill.cpp | 133 +++++++++++++----- .../hipamd/tests/src/printf/printf_common.h | 15 ++ 2 files changed, 114 insertions(+), 34 deletions(-) diff --git a/projects/clr/hipamd/tests/performance/memory/hipPerfMemFill.cpp b/projects/clr/hipamd/tests/performance/memory/hipPerfMemFill.cpp index 81845b886e..1570c84301 100644 --- a/projects/clr/hipamd/tests/performance/memory/hipPerfMemFill.cpp +++ b/projects/clr/hipamd/tests/performance/memory/hipPerfMemFill.cpp @@ -24,6 +24,7 @@ */ #include "test_common.h" +#include #include #include #include @@ -46,6 +47,31 @@ __global__ void vec_fill(T *x, T coef, int N) { } } +__device__ void print_log(int i, double value, double expected) { + printf("failed at %d: val=%g, expected=%g\n", i, value, expected); +} + +__device__ void print_log(int i, int value, int expected) { + printf("failed at %d: val=%d, expected=%d\n", i, value, expected); +} + +template +__global__ void vec_verify(T *x, T coef, int N) { + const int istart = threadIdx.x + blockIdx.x * blockDim.x; + const int ishift = blockDim.x * gridDim.x; + for (int i = istart; i < N; i += ishift) { +#if SIMPLY_ASSIGN + if(x[i] != coef) { + print_log(i, x[i], coef); + } +#else + if(x[i] != coef * i) { + print_log(i, x[i], coef * i); + } +#endif + } +} + template __global__ void daxpy(T *__restrict__ x, T *__restrict__ y, const T coef, int Niter, int N) { @@ -120,6 +146,19 @@ class hipPerfMemFill { << std::endl; } + void log_host(const char* title, double GBytes, double sec) { + cout << title << " [" << setw(7) << GBytes << " GB]: cost " << setw(10) << sec + << " s in bandwidth " << setw(10) << GBytes / sec << " [GB/s]" << endl; + } + + void log_kernel(const char* title, double GBytes, double sec, double sec_hv, double sec_kv) { + cout << title << " [" << setw(7) << GBytes << " GB]: cost " << setw(10) << sec + << " s in bandwidth " << setw(10) << GBytes / sec << " [GB/s]" << ", hostVerify cost " + << setw(10) << sec_hv << " s in bandwidth " << setw(10) << GBytes / sec_hv << " [GB/s]" + << ", kernelVerify cost "<< setw(10) << sec_kv << " s in bandwidth " << setw(10) + << GBytes / sec_kv << " [GB/s]" << endl; + } + void hostFill(size_t size, T *data, T coef, double &sec) { size_t num = size / sizeof(T); // Size of elements auto start = chrono::steady_clock::now(); @@ -156,18 +195,18 @@ class hipPerfMemFill { sec = diff.count() / NUM_ITER; // in second } - void verify(size_t size, T *data, T coef, double &sec) { + void hostVerify(size_t size, T *data, T coef, double &sec) { size_t num = size / sizeof(T); // Size of elements auto start = chrono::steady_clock::now(); for (int i = 0; i < num; ++i) { #if SIMPLY_ASSIGN if(data[i] != coef) { - cout << "verify failed: i=" << i << ", data[i]=" << data[i] << ", expected=" << coef << endl; + cout << "hostVerify failed: i=" << i << ", data[i]=" << data[i] << ", expected=" << coef << endl; failed("failed\n"); } #else if(data[i] != coef * i) { - cout << "verify failed: i=" << i << ", data[i]=" << data[i] << ", expected=" << coef * i << endl; + cout << "hostVerify failed: i=" << i << ", data[i]=" << data[i] << ", expected=" << coef * i << endl; failed("failed\n"); } #endif @@ -177,7 +216,40 @@ class hipPerfMemFill { sec = diff.count(); } - bool testLargeBarHostFill(size_t size) { + void kernelVerify(size_t size, T *data, T coef, double &sec) { + size_t num = size / sizeof(T); // Size of elements + unsigned blocks = setNumBlocks(num); + + CaptureStream *capture = new CaptureStream(stdout); + capture->Begin(); + + hipLaunchKernelGGL(HIP_KERNEL_NAME(vec_verify), dim3(blocks), + dim3(threadsPerBlock), 0, 0, data, coef, num); // kernel will be loaded first time + HIPCHECK(hipDeviceSynchronize()); + + capture->End(); + capture->Truncate(1000); // Don't want too long log if existing + std::string device_output = capture->getData(); + delete capture; + if (device_output.length() > 0) { + failed("kernelVerify failed:\n%s\n", device_output.c_str()); + } + + // Now all data verified. The following is to test bandwidth. + auto start = chrono::steady_clock::now(); + + for (int iter = 0; iter < NUM_ITER; ++iter) { + hipLaunchKernelGGL(HIP_KERNEL_NAME(vec_verify), dim3(blocks), + dim3(threadsPerBlock), 0, 0, data, coef, num); + } + HIPCHECK(hipDeviceSynchronize()); + + auto end = chrono::steady_clock::now(); + chrono::duration diff = end - start; // in second + sec = diff.count() / NUM_ITER; // in second + } + + bool testLargeBarDeviceMemoryHostFill(size_t size) { if (!supportLargeBar()) { return false; } @@ -190,8 +262,7 @@ class hipPerfMemFill { hostFill(size, A, coef_, sec); // Cpu can access device mem in LB HIPCHECK(hipFree(A)); - cout << "Largebar: host fill [" << setw(7) << GBytes << " GB]: cost " << setw(10) << sec - << " s in bandwidth " << setw(10) << GBytes / sec << " [GB/s]" << endl; + log_host("Largebar: host fill", GBytes, sec); return true; } @@ -202,7 +273,7 @@ class hipPerfMemFill { cout << "Test large bar device memory host filling" << endl; for (int i = 0; i < NUM_SIZE; i++) { - if (!testLargeBarHostFill(totalSizes_[i])) { + if (!testLargeBarDeviceMemoryHostFill(totalSizes_[i])) { return false; } } @@ -222,9 +293,7 @@ class hipPerfMemFill { hostFill(size, A, coef_, sec); // Cpu can access HMM mem HIPCHECK(hipFree(A)); - cout << "Managed: host fill [" << setw(7) << GBytes << " GB]: cost " << setw(10) << sec - << " s in bandwidth " << setw(10) << GBytes / sec << " [GB/s]" << endl; - + log_host("Managed: host fill", GBytes, sec); return true; } @@ -236,15 +305,14 @@ class hipPerfMemFill { T *A; HIPCHECK(hipMallocManaged(&A, size)); - double sec = 0, sec_v = 0; + + double sec = 0, sec_hv = 0, sec_kv = 0; kernelFill(size, A, coef_, sec); - verify(size, A, coef_, sec_v); // Managed memory can be verified by host + hostVerify(size, A, coef_, sec_hv); // Managed memory can be verified by host + kernelVerify(size, A, coef_, sec_kv); HIPCHECK(hipFree(A)); - cout << "Managed: kernel fill [" << setw(7) << GBytes << " GB]: cost " << setw(10) << sec - << " s in bandwidth " << setw(10) << GBytes / sec << " [GB/s]" << ", verify: cost " - << setw(10) << sec_v << " s in bandwidth " << setw(10) << GBytes / sec_v << " [GB/s]" - << endl; + log_kernel("Managed: kernel fill", GBytes, sec, sec_hv, sec_kv); return true; } @@ -279,9 +347,7 @@ class hipPerfMemFill { hostFill(size, A, coef_, sec); HIPCHECK(hipHostFree(A)); - cout << "Host: host fill [" << setw(7) << GBytes << " GB]: cost " << setw(10) << sec - << " s in bandwidth " << setw(10) << GBytes / sec << " [GB/s]" << endl; - + log_host("Host: host fill", GBytes, sec); return true; } @@ -290,16 +356,13 @@ class hipPerfMemFill { T *A; HIPCHECK(hipHostMalloc((void** ) &A, size, flags)); - double sec = 0, sec_v = 0; + double sec = 0, sec_hv = 0, sec_kv = 0; kernelFill(size, A, coef_, sec); - verify(size, A, coef_, sec_v); + hostVerify(size, A, coef_, sec_hv); + kernelVerify(size, A, coef_, sec_kv); HIPCHECK(hipHostFree(A)); - cout << "Host: kernel fill [" << setw(7) << GBytes << " GB]: cost " << setw(10) << sec - << " s in bandwidth " << setw(10) << GBytes / sec << " [GB/s]" << ", verify cost " - << setw(10) << sec_v << " s in bandwidth " << setw(10) << GBytes / sec_v << " [GB/s]" - << endl; - + log_kernel("Host: kernel fill", GBytes, sec, sec_hv, sec_kv); return true; } @@ -356,12 +419,15 @@ class hipPerfMemFill { return (num + threadsPerBlock_ - 1) / threadsPerBlock_; #endif } + bool testExtDeviceMemoryHostFill(size_t size, unsigned int flags) { double GBytes = (double) size / (1024.0 * 1024.0 * 1024.0); T *A = nullptr; HIPCHECK(hipExtMallocWithFlags((void **)&A, size, flags)); if (!A) { + cout << "failed hipExtMallocWithFlags() with size =" << size << " flags=" + << std::hex << flags << endl; return false; } @@ -369,8 +435,7 @@ class hipPerfMemFill { hostFill(size, A, coef_, sec); // Cpu can access this mem HIPCHECK(hipFree(A)); - cout << "ExtDevice: host fill [" << setw(7) << GBytes << " GB]: cost " << setw(10) << sec - << " s in bandwidth " << setw(10) << GBytes / sec << " [GB/s]" << endl; + log_host("ExtDevice: host fill", GBytes, sec); return true; } @@ -380,18 +445,18 @@ class hipPerfMemFill { T *A = nullptr; HIPCHECK(hipExtMallocWithFlags((void **)&A, size, flags)); if (!A) { + cout << "failed hipExtMallocWithFlags() with size =" << size << " flags=" + << std::hex << flags << endl; return false; } - double sec = 0, sec_v = 0; + double sec = 0, sec_hv = 0, sec_kv = 0; kernelFill(size, A, coef_, sec); - verify(size, A, coef_, sec_v); // Fine grained device memory can be verified by host + hostVerify(size, A, coef_, sec_hv); // Fine grained device memory can be verified by host + kernelVerify(size, A, coef_, sec_kv); HIPCHECK(hipFree(A)); - cout << "ExtDevice: kernel fill [" << setw(7) << GBytes << " GB]: cost " << setw(10) << sec - << " s in bandwidth " << setw(10) << GBytes / sec << " [GB/s]" << ", verify cost " - << setw(10) << sec_v << " s in bandwidth " << setw(10) << GBytes / sec_v << " [GB/s]" - << endl; + log_kernel("ExtDevice: kernel fill", GBytes, sec, sec_hv, sec_kv); return true; } diff --git a/projects/clr/hipamd/tests/src/printf/printf_common.h b/projects/clr/hipamd/tests/src/printf/printf_common.h index 4f75b13d62..d8cce83f0d 100644 --- a/projects/clr/hipamd/tests/src/printf/printf_common.h +++ b/projects/clr/hipamd/tests/src/printf/printf_common.h @@ -134,6 +134,21 @@ struct CaptureStream { assert(false); } } + + // Truncate the file up to size if we don't want too long log + void Truncate(size_t size) { + struct stat sb = { 0 }; + if (::stat(tempname, &sb) == -1) { + failed("failed lstat(%s) with error: %s \n", tempname, ::strerror(errno)); + return; + } + if (sb.st_size > size) { + if (::truncate(tempname, static_cast(size)) == -1) { + failed("failed truncate(%s) with error: %s \n", tempname, ::strerror(errno)); + return; + } + } + } }; #endif