diff --git a/catch/perftests/memory/hipPerfBufferCopyRectSpeed.cc b/catch/perftests/memory/hipPerfBufferCopyRectSpeed.cc index 5fc601b141..61a9bd7b4d 100644 --- a/catch/perftests/memory/hipPerfBufferCopyRectSpeed.cc +++ b/catch/perftests/memory/hipPerfBufferCopyRectSpeed.cc @@ -18,30 +18,31 @@ THE SOFTWARE. */ /** -* @addtogroup hipMemcpy2DAsync hipMemcpy2DAsync -* @{ -* @ingroup perfMemoryTest -* `hipMemcpy2DAsync(void* dst, size_t dpitch, const void* src, size_t spitch, -* size_t width, size_t height, hipMemcpyKind kind, hipStream_t stream = 0)` - -* Copies data between host and device. -*/ + * @addtogroup hipMemcpy2DAsync hipMemcpy2DAsync + * @{ + * @ingroup perfMemoryTest + * `hipMemcpy2DAsync(void* dst, size_t dpitch, const void* src, size_t spitch, + * size_t width, size_t height, hipMemcpyKind kind, hipStream_t stream = 0)` - + * Copies data between host and device. + */ #include +// #define ENABLE_DEBUG 1 #define NUM_SIZES 8 // 4KB, 8KB, 64KB, 256KB, 1 MB, 4MB, 16 MB, 16MB+10 -static const unsigned int Sizes[NUM_SIZES] = - {4096, 8192, 65536, 262144, 1048576, 4194304, 16777216, 16777216+10}; +static const unsigned int Sizes[NUM_SIZES] = {4096, 8192, 65536, 262144, + 1048576, 4194304, 16777216, 16777216 + 10}; static const unsigned int Iterations[2] = {1, 1000}; #define BUF_TYPES 4 // 16 ways to combine 4 different buffer types -#define NUM_SUBTESTS (BUF_TYPES*BUF_TYPES) +#define NUM_SUBTESTS (BUF_TYPES * BUF_TYPES) -static void setData(void *ptr, unsigned int size, char value) { - char *ptr2 = reinterpret_cast(ptr); - for (unsigned int i = 0; i < size ; i++) { +static void setData(void* ptr, unsigned int size, char value) { + char* ptr2 = reinterpret_cast(ptr); + for (unsigned int i = 0; i < size; i++) { ptr2[i] = value; } } @@ -52,17 +53,17 @@ static bool hipPerfBufferCopyRectSpeed_test(int p_tests) { bool hostMalloc[2] = {false}; bool hostRegister[2] = {false}; bool unpinnedMalloc[2] = {false}; - void *memptr[2] = {NULL}; - void *alignedmemptr[2] = {NULL}; - void *srcBuffer = NULL; - void *dstBuffer = NULL; + void* memptr[2] = {NULL}; + void* alignedmemptr[2] = {NULL}; + void* srcBuffer = NULL; + void* dstBuffer = NULL; - int numTests = (p_tests == -1) ? (NUM_SIZES*NUM_SUBTESTS*2 - 1) : p_tests; + int numTests = (p_tests == -1) ? (NUM_SIZES * NUM_SUBTESTS * 2 - 1) : p_tests; int test = (p_tests == -1) ? 0 : p_tests; - for ( ; test <= numTests ; test++ ) { + for (; test <= numTests; test++) { unsigned int srcTest = (test / NUM_SIZES) % BUF_TYPES; - unsigned int dstTest = (test / (NUM_SIZES*BUF_TYPES)) % BUF_TYPES; + unsigned int dstTest = (test / (NUM_SIZES * BUF_TYPES)) % BUF_TYPES; bufSize_ = Sizes[test % NUM_SIZES]; hostMalloc[0] = hostMalloc[1] = false; hostRegister[0] = hostRegister[1] = false; @@ -92,8 +93,7 @@ static bool hipPerfBufferCopyRectSpeed_test(int p_tests) { numIter = Iterations[test / (NUM_SIZES * NUM_SUBTESTS)]; if (hostMalloc[0]) { - HIP_CHECK(hipHostMalloc(reinterpret_cast(&srcBuffer), - bufSize_, 0)); + HIP_CHECK(hipHostMalloc(reinterpret_cast(&srcBuffer), bufSize_, 0)); setData(srcBuffer, bufSize_, 0xd0); } else if (hostRegister[0]) { memptr[0] = malloc(bufSize_ + 4096); @@ -112,8 +112,7 @@ static bool hipPerfBufferCopyRectSpeed_test(int p_tests) { } if (hostMalloc[1]) { - HIP_CHECK(hipHostMalloc(reinterpret_cast(&dstBuffer), - bufSize_, 0)); + HIP_CHECK(hipHostMalloc(reinterpret_cast(&dstBuffer), bufSize_, 0)); } else if (hostRegister[1]) { memptr[1] = malloc(bufSize_ + 4096); alignedmemptr[1] = reinterpret_cast(memptr[0]); @@ -128,15 +127,14 @@ static bool hipPerfBufferCopyRectSpeed_test(int p_tests) { } // warm up - HIP_CHECK(hipMemcpy2D(dstBuffer, width, srcBuffer, - width, width, width, hipMemcpyDefault)); + HIP_CHECK(hipMemcpy2D(dstBuffer, width, srcBuffer, width, width, width, hipMemcpyDefault)); // measure performance based on host time auto all_start = std::chrono::steady_clock::now(); for (unsigned int i = 0; i < numIter; i++) { - HIP_CHECK(hipMemcpy2DAsync(dstBuffer, width, srcBuffer, - width, width, width, hipMemcpyDefault, NULL)); + HIP_CHECK(hipMemcpy2DAsync(dstBuffer, width, srcBuffer, width, width, width, hipMemcpyDefault, + NULL)); } HIP_CHECK(hipDeviceSynchronize()); @@ -144,11 +142,11 @@ static bool hipPerfBufferCopyRectSpeed_test(int p_tests) { std::chrono::duration elapsed_secs = all_end - all_start; // read speed in GB/s - double perf = (static_cast(bufSize_ * numIter) * - static_cast(1e-09)) / elapsed_secs.count(); + double perf = (static_cast(bufSize_ * numIter) * static_cast(1e-09)) / + elapsed_secs.count(); - const char *strSrc = NULL; - const char *strDst = NULL; + const char* strSrc = NULL; + const char* strDst = NULL; if (hostMalloc[0]) strSrc = "hHM"; else if (hostRegister[0]) @@ -170,15 +168,14 @@ static bool hipPerfBufferCopyRectSpeed_test(int p_tests) { // Double results when src and dst are both on device if ((!hostMalloc[0] && !hostRegister[0] && !unpinnedMalloc[0]) && (!hostMalloc[1] && !hostRegister[1] && !unpinnedMalloc[1])) - perf *= 2.0; + perf *= 2.0; // Double results when src and dst are both in sysmem if ((hostMalloc[0] || hostRegister[0] || unpinnedMalloc[0]) && (hostMalloc[1] || hostRegister[1] || unpinnedMalloc[1])) - perf *= 2.0; + perf *= 2.0; - INFO("hipPerfBufferCopyRectSpeed[" << test << "]\t( " << bufSize_ << - ")\ts:" << strSrc << " d:" << strDst << "\ti:" << numIter << - "\t(GB/s) perf\t" << (float)perf); + CONSOLE_PRINT("hipPerfBufferCopyRectSpeed[%d]\t( %u )\ts:%s d:%s\ti:%u\t(GB/s) perf\t%.2f\n", + test, bufSize_, strSrc, strDst, numIter, (float)perf); // Free src if (hostMalloc[0]) { @@ -208,40 +205,42 @@ static bool hipPerfBufferCopyRectSpeed_test(int p_tests) { } /** -* Test Description -* ------------------------ -*  - Verify hipPerfBufferCopy status. -* Test source -* ------------------------ -*  - perftests/memory/hipPerfBufferCopyRectSpeed.cc -* Test requirements -* ------------------------ -*  - HIP_VERSION >= 5.6 -*/ + * Test Description + * ------------------------ + *  - Verify hipPerfBufferCopy status. + * Test source + * ------------------------ + *  - perftests/memory/hipPerfBufferCopyRectSpeed.cc + * Test requirements + * ------------------------ + *  - HIP_VERSION >= 5.6 + */ TEST_CASE("Perf_hipPerfBufferCopyRectSpeed_test") { int numDevices = 0; HIP_CHECK(hipGetDeviceCount(&numDevices)); if (numDevices <= 0) { - SUCCEED("Skipped testcase hipPerfBufferCopyRectSpeed" - "as there is no device to test."); + SUCCEED( + "Skipped testcase hipPerfBufferCopyRectSpeed" + "as there is no device to test."); } else { int deviceId = 0; HIP_CHECK(hipSetDevice(deviceId)); hipDeviceProp_t props; HIP_CHECK(hipGetDeviceProperties(&props, deviceId)); - INFO("hipPerfBufferCopyRectSpeed - info: Set device to " << deviceId - << " : " << props.name << "Legend: unp - unpinned(malloc)," - " hM - hipMalloc(device)\n hHR - hipHostRegister(pinned)," - " hHM - hipHostMalloc(prePinned)\n"); + CONSOLE_PRINT( + "hipPerfBufferCopyRectSpeed - info: Set device to %d : %s Legend: unp - unpinned(malloc), " + "hM - hipMalloc(device)\n hHR - hipHostRegister(pinned), hHM - " + "hipHostMalloc(prePinned)\n", + deviceId, props.name); REQUIRE(true == hipPerfBufferCopyRectSpeed_test(1)); } } /** -* End doxygen group perfMemoryTest. -* @} -*/ + * End doxygen group perfMemoryTest. + * @} + */ diff --git a/catch/perftests/memory/hipPerfDevMemReadSpeed.cc b/catch/perftests/memory/hipPerfDevMemReadSpeed.cc index ae5f631866..9b874b6601 100644 --- a/catch/perftests/memory/hipPerfDevMemReadSpeed.cc +++ b/catch/perftests/memory/hipPerfDevMemReadSpeed.cc @@ -18,13 +18,14 @@ THE SOFTWARE. */ /** -* @addtogroup hipMemcpyKernel hipMemcpyKernel -* @{ -* @ingroup perfMemoryTest -* `hipMemcpy(void* dst, const void* src, size_t count, hipMemcpyKind kind)` - -* Copies data between host and device. -*/ + * @addtogroup hipMemcpyKernel hipMemcpyKernel + * @{ + * @ingroup perfMemoryTest + * `hipMemcpy(void* dst, const void* src, size_t count, hipMemcpyKind kind)` - + * Copies data between host and device. + */ +// #define ENABLE_DEBUG 1 #include #define ARRAY_SIZE 16 @@ -33,7 +34,7 @@ typedef struct d_uint16 { uint data[ARRAY_SIZE]; } d_uint16; -__global__ static void read_kernel(d_uint16 *src, ulong N, uint *dst) { +__global__ static void read_kernel(d_uint16* src, ulong N, uint* dst) { size_t idx = (blockIdx.x * blockDim.x + threadIdx.x); size_t stride = blockDim.x * gridDim.x; @@ -59,8 +60,8 @@ static bool hipPerfDevMemReadSpeed_test() { hipDeviceProp_t props; HIP_CHECK(hipGetDeviceProperties(&props, deviceId)); - INFO("info: running on bus " << "0x" << props.pciBusID << " " << - props.name << " with " << props.multiProcessorCount << " CUs \n"); + CONSOLE_PRINT("info: running on bus 0x%x %s with %d CUs\n", props.pciBusID, props.name, + props.multiProcessorCount); const unsigned threadsPerBlock = 64; const unsigned blocks = props.multiProcessorCount * 4; @@ -70,7 +71,7 @@ static bool hipPerfDevMemReadSpeed_test() { hSrc = new d_uint16[nBytes]; REQUIRE(hSrc != nullptr); - hDst = new uint; + hDst = new uint; REQUIRE(hDst != nullptr); hDst[0] = 0; @@ -88,15 +89,15 @@ static bool hipPerfDevMemReadSpeed_test() { HIP_CHECK(hipMemcpy(dSrc, hSrc, nBytes, hipMemcpyHostToDevice)); HIP_CHECK(hipMemcpy(dDst, hDst, sizeof(uint), hipMemcpyHostToDevice)); - hipLaunchKernelGGL(read_kernel, dim3(blocks), dim3(threadsPerBlock), - 0, stream, dSrc, N, dDst); + hipLaunchKernelGGL(read_kernel, dim3(blocks), dim3(threadsPerBlock), 0, stream, dSrc, N, dDst); HIP_CHECK(hipGetLastError()); HIP_CHECK(hipMemcpy(hDst, dDst, sizeof(uint), hipMemcpyDeviceToHost)); HIP_CHECK(hipDeviceSynchronize()); if (hDst[0] != (nBytes / sizeof(uint))) { - INFO("hipPerfDevMemReadSpeed - Data validation failed for warm up run!" << - " expected " << nBytes / sizeof(uint) << " got " << hDst[0]); + DEBUG_PRINT( + "hipPerfDevMemReadSpeed - Data validation failed for warm up run! expected %u got %u\n", + nBytes / sizeof(uint), hDst[0]); return false; } @@ -104,8 +105,7 @@ static bool hipPerfDevMemReadSpeed_test() { auto all_start = std::chrono::steady_clock::now(); for (int i = 0; i < nIter; i++) { - hipLaunchKernelGGL(read_kernel, dim3(blocks), dim3(threadsPerBlock), - 0, stream, dSrc, N, dDst); + hipLaunchKernelGGL(read_kernel, dim3(blocks), dim3(threadsPerBlock), 0, stream, dSrc, N, dDst); HIP_CHECK(hipGetLastError()); } HIP_CHECK(hipDeviceSynchronize()); @@ -114,14 +114,14 @@ static bool hipPerfDevMemReadSpeed_test() { std::chrono::duration all_kernel_time = all_end - all_start; // read speed in GB/s - double perf = (static_cast(nBytes * nIter * (1e-09))) / - all_kernel_time.count(); + double perf = (static_cast(nBytes * nIter * (1e-09))) / all_kernel_time.count(); - INFO("hipPerfDevMemReadSpeed - info: average read speed of " << - perf << " GB/s " << "achieved for memory size of " << - nBytes / (1024 * 1024) << " MB"); + CONSOLE_PRINT( + "hipPerfDevMemReadSpeed - average read speed of %.2f GB/s achieved for memory size of %u " + "MB\n", + perf, nBytes / (1024 * 1024)); - delete [] hSrc; + delete[] hSrc; delete hDst; HIP_CHECK(hipFree(dSrc)); HIP_CHECK(hipFree(dDst)); @@ -130,30 +130,31 @@ static bool hipPerfDevMemReadSpeed_test() { } /** -* Test Description -* ------------------------ -*  - Verify hipPerfDevMemReadSpeed status. -* Test source -* ------------------------ -*  - perftests/memory/hipPerfDevMemReadSpeed.cc -* Test requirements -* ------------------------ -*  - HIP_VERSION >= 5.6 -*/ + * Test Description + * ------------------------ + *  - Verify hipPerfDevMemReadSpeed status. + * Test source + * ------------------------ + *  - perftests/memory/hipPerfDevMemReadSpeed.cc + * Test requirements + * ------------------------ + *  - HIP_VERSION >= 5.6 + */ TEST_CASE("Perf_hipPerfDevMemReadSpeed_test") { int numDevices = 0; HIP_CHECK(hipGetDeviceCount(&numDevices)); if (numDevices <= 0) { - SUCCEED("Skipped testcase hipPerfDevMemReadSpeed as" - "there is no device to test."); + SUCCEED( + "Skipped testcase hipPerfDevMemReadSpeed as" + "there is no device to test."); } else { REQUIRE(true == hipPerfDevMemReadSpeed_test()); } } /** -* End doxygen group perfMemoryTest. -* @} -*/ + * End doxygen group perfMemoryTest. + * @} + */ diff --git a/catch/perftests/memory/hipPerfDevMemWriteSpeed.cc b/catch/perftests/memory/hipPerfDevMemWriteSpeed.cc index 77eac4b8bb..e96b206c84 100644 --- a/catch/perftests/memory/hipPerfDevMemWriteSpeed.cc +++ b/catch/perftests/memory/hipPerfDevMemWriteSpeed.cc @@ -18,12 +18,12 @@ THE SOFTWARE. */ /** -* @addtogroup hipMemcpyKernel hipMemcpyKernel -* @{ -* @ingroup perfMemoryTest -* `hipMemcpy(void* dst, const void* src, size_t count, hipMemcpyKind kind)` - -* Copies data between host and device. -*/ + * @addtogroup hipMemcpyKernel hipMemcpyKernel + * @{ + * @ingroup perfMemoryTest + * `hipMemcpy(void* dst, const void* src, size_t count, hipMemcpyKind kind)` - + * Copies data between host and device. + */ #include @@ -33,12 +33,12 @@ typedef struct d_uint16 { uint data[ARRAY_SIZE]; } d_uint16; -__global__ void write_kernel(d_uint16 *dst, ulong N, d_uint16 pval) { - size_t idx = (blockIdx.x * blockDim.x + threadIdx.x); - size_t stride = blockDim.x * gridDim.x; - for (size_t i = idx; i < N; i += stride) { - dst[i] = pval; - } +__global__ void write_kernel(d_uint16* dst, ulong N, d_uint16 pval) { + size_t idx = (blockIdx.x * blockDim.x + threadIdx.x); + size_t stride = blockDim.x * gridDim.x; + for (size_t i = idx; i < N; i += stride) { + dst[i] = pval; + } } static bool hipPerfDevMemWriteSpeed_test() { @@ -55,8 +55,8 @@ static bool hipPerfDevMemWriteSpeed_test() { hipDeviceProp_t props; HIP_CHECK(hipGetDeviceProperties(&props, deviceId)); - INFO("info: running on bus " << "0x" << props.pciBusID << " " << - props.name << " with " << props.multiProcessorCount << " CUs \n"); + CONSOLE_PRINT("info: running on bus 0x%x %s with %d CUs\n", props.pciBusID, props.name, + props.multiProcessorCount); const unsigned threadsPerBlock = 64; const unsigned blocks = props.multiProcessorCount * 4; @@ -65,7 +65,7 @@ static bool hipPerfDevMemWriteSpeed_test() { pval.data[i] = inputData; } - hDst = new d_uint16[nBytes]; + hDst = new d_uint16[nBytes]; REQUIRE(hDst != nullptr); for (size_t i = 0; i < N; i++) { @@ -78,18 +78,18 @@ static bool hipPerfDevMemWriteSpeed_test() { HIP_CHECK(hipStreamCreate(&stream)); HIP_CHECK(hipMalloc(&dDst, nBytes)); - hipLaunchKernelGGL(write_kernel, dim3(blocks), dim3(threadsPerBlock), - 0, stream, dDst, N, pval); + hipLaunchKernelGGL(write_kernel, dim3(blocks), dim3(threadsPerBlock), 0, stream, dDst, N, pval); HIP_CHECK(hipGetLastError()); - HIP_CHECK(hipMemcpy(hDst, dDst, nBytes , hipMemcpyDeviceToHost)); + HIP_CHECK(hipMemcpy(hDst, dDst, nBytes, hipMemcpyDeviceToHost)); HIP_CHECK(hipDeviceSynchronize()); for (uint i = 0; i < N; i++) { for (uint j = 0; j < ARRAY_SIZE; j++) { if (hDst[i].data[j] != inputData) { - INFO("hipPerfDevMemWriteSpeed - Data validation failed for warm up run!" - << "at index i: " << i << " element j: " << j << - "expected " << inputData << " but got " << hDst[i].data[j]); + DEBUG_PRINT( + "hipPerfDevMemWriteSpeed - Data validation failed for warm up run! at index i: %u " + "element j: %u expected 0x%x but got 0x%x\n", + i, j, inputData, hDst[i].data[j]); return false; } } @@ -99,8 +99,7 @@ static bool hipPerfDevMemWriteSpeed_test() { auto all_start = std::chrono::steady_clock::now(); for (int i = 0; i < nIter; i++) { - hipLaunchKernelGGL(write_kernel, dim3(blocks), dim3(threadsPerBlock), - 0, stream, dDst, N, pval); + hipLaunchKernelGGL(write_kernel, dim3(blocks), dim3(threadsPerBlock), 0, stream, dDst, N, pval); HIP_CHECK(hipGetLastError()); } HIP_CHECK(hipDeviceSynchronize()); @@ -109,44 +108,45 @@ static bool hipPerfDevMemWriteSpeed_test() { std::chrono::duration all_kernel_time = all_end - all_start; // read speed in GB/s - double perf = (static_cast(nBytes * nIter * (1e-09))) / - all_kernel_time.count(); + double perf = (static_cast(nBytes * nIter * (1e-09))) / all_kernel_time.count(); - INFO("hipPerfDevMemReadSpeed - info: average write speed of " << - perf << " GB/s " << "achieved for memory size of " << - nBytes / (1024 * 1024) << " MB"); + CONSOLE_PRINT( + "hipPerfDevMemWriteSpeed - average write speed of %.2f GB/s achieved for memory size of %u " + "MB\n", + perf, nBytes / (1024 * 1024)); - delete [] hDst; + delete[] hDst; HIP_CHECK(hipFree(dDst)); HIP_CHECK(hipStreamDestroy(stream)); return true; } /** -* Test Description -* ------------------------ -*  - Verify hipPerfDevMemWriteSpeed status. -* Test source -* ------------------------ -*  - perftests/memory/hipPerfDevMemWriteSpeed.cc -* Test requirements -* ------------------------ -*  - HIP_VERSION >= 5.6 -*/ + * Test Description + * ------------------------ + *  - Verify hipPerfDevMemWriteSpeed status. + * Test source + * ------------------------ + *  - perftests/memory/hipPerfDevMemWriteSpeed.cc + * Test requirements + * ------------------------ + *  - HIP_VERSION >= 5.6 + */ TEST_CASE("Perf_hipPerfDevMemWriteSpeed_test") { int numDevices = 0; HIP_CHECK(hipGetDeviceCount(&numDevices)); if (numDevices <= 0) { - SUCCEED("Skipped testcase hipPerfDevMemWriteSpeed as" - "there is no device to test."); + SUCCEED( + "Skipped testcase hipPerfDevMemWriteSpeed as" + "there is no device to test."); } else { REQUIRE(true == hipPerfDevMemWriteSpeed_test()); } } /** -* End doxygen group perfMemoryTest. -* @} -*/ + * End doxygen group perfMemoryTest. + * @} + */ diff --git a/catch/perftests/memory/hipPerfHostNumaAlloc.cc b/catch/perftests/memory/hipPerfHostNumaAlloc.cc index b1f0566814..39a437b681 100644 --- a/catch/perftests/memory/hipPerfHostNumaAlloc.cc +++ b/catch/perftests/memory/hipPerfHostNumaAlloc.cc @@ -18,28 +18,27 @@ THE SOFTWARE. */ /** -* @addtogroup hipMemcpyKernel hipMemcpyKernel -* @{ -* @ingroup perfMemoryTest -* `hipMemcpy(void* dst, const void* src, size_t count, hipMemcpyKind kind)` - -* Copies data between host and device. -*/ + * @addtogroup hipMemcpyKernel hipMemcpyKernel + * @{ + * @ingroup perfMemoryTest + * `hipMemcpy(void* dst, const void* src, size_t count, hipMemcpyKind kind)` - + * Copies data between host and device. + */ #include #include - +// #define ENABLE_DEBUG 1 // To run it correctly, we must not export HIP_VISIBLE_DEVICES. // And we must explicitly link libnuma because of numa api move_pages(). #define NUM_PAGES 4 -char *h = nullptr; -char *d_h = nullptr; -char *m = nullptr; -char *d_m = nullptr; +char* h = nullptr; +char* d_h = nullptr; +char* m = nullptr; +char* d_m = nullptr; int page_size = 1024; -const int mode[] = { MPOL_DEFAULT, MPOL_BIND, MPOL_PREFERRED, MPOL_INTERLEAVE }; -const char* modeStr[] = { "MPOL_DEFAULT", "MPOL_BIND", - "MPOL_PREFERRED", "MPOL_INTERLEAVE" }; +const int mode[] = {MPOL_DEFAULT, MPOL_BIND, MPOL_PREFERRED, MPOL_INTERLEAVE}; +const char* modeStr[] = {"MPOL_DEFAULT", "MPOL_BIND", "MPOL_PREFERRED", "MPOL_INTERLEAVE"}; std::string exeCommand(const char* cmd) { std::array buff; @@ -55,23 +54,22 @@ std::string exeCommand(const char* cmd) { } int getCpuAgentCount() { - const char* cmd = - "cat /proc/cpuinfo | grep \"physical id\" | sort | uniq | wc -l"; + const char* cmd = "cat /proc/cpuinfo | grep \"physical id\" | sort | uniq | wc -l"; int cpuAgentCount = std::atoi(exeCommand(cmd).c_str()); return cpuAgentCount; } bool test(int cpuId, int gpuId, int numaMode, unsigned int hostMallocflags) { - void *pages[NUM_PAGES]; + void* pages[NUM_PAGES]; int status[NUM_PAGES]; int ret_code; - INFO("set cpu " << cpuId << ", gpu " << gpuId << ", numaMode " - << numaMode << ", hostMallocflags " << hostMallocflags << "\n"); + CONSOLE_PRINT("set cpu %d, gpu %d, numaMode %d, hostMallocflags %u\n", cpuId, gpuId, numaMode, + hostMallocflags); if (cpuId >= 0) { - unsigned long nodeMask = 1 << cpuId; //NOLINT - unsigned long maxNode = sizeof(nodeMask) * 8; //NOLINT + unsigned long nodeMask = 1 << cpuId; // NOLINT + unsigned long maxNode = sizeof(nodeMask) * 8; // NOLINT if (set_mempolicy(numaMode, numaMode == MPOL_DEFAULT ? NULL : &nodeMask, numaMode == MPOL_DEFAULT ? 0 : maxNode) == -1) { WARN("set_mempolicy() failed with err " << errno << "\n"); @@ -83,7 +81,7 @@ bool test(int cpuId, int gpuId, int numaMode, unsigned int hostMallocflags) { HIP_CHECK(hipSetDevice(gpuId)); } - posix_memalign(reinterpret_cast(&m), page_size, page_size*NUM_PAGES); + posix_memalign(reinterpret_cast(&m), page_size, page_size * NUM_PAGES); HIP_CHECK(hipHostRegister(m, page_size * NUM_PAGES, hipHostRegisterMapped)); HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast(&d_m), m, 0)); @@ -94,15 +92,13 @@ bool test(int cpuId, int gpuId, int numaMode, unsigned int hostMallocflags) { } ret_code = move_pages(0, NUM_PAGES, pages, NULL, status, 0); - INFO("Memory (malloc) ret " << ret_code << " at " << m << - " (dev " << d_m << "%p) is at node: "); + CONSOLE_PRINT("Memory (malloc) ret %d at %p (dev %p) is at node: ", ret_code, m, d_m); for (int i = 0; i < NUM_PAGES; i++) { - INFO(status[i]); // Don't verify as it's out of our control + CONSOLE_PRINT("%d ", status[i]); // Don't verify as it's out of our control } - INFO("\n"); + CONSOLE_PRINT("\n"); - HIP_CHECK(hipHostMalloc(reinterpret_cast(&h), - page_size*NUM_PAGES, hostMallocflags)); + HIP_CHECK(hipHostMalloc(reinterpret_cast(&h), page_size * NUM_PAGES, hostMallocflags)); pages[0] = h; for (int i = 1; i < NUM_PAGES; i++) { pages[i] = reinterpret_cast(pages[0]) + page_size; @@ -111,16 +107,14 @@ bool test(int cpuId, int gpuId, int numaMode, unsigned int hostMallocflags) { d_h = nullptr; if (hostMallocflags & hipHostMallocMapped) { HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast(&d_h), h, 0)); - INFO("Memory (hipHostMalloc) ret " << ret_code << " at " << h - << " (dev " << d_h << ") is at node: "); + CONSOLE_PRINT("Memory (hipHostMalloc) ret %d at %p (dev %p) is at node: ", ret_code, h, d_h); } else { - INFO("Memory (hipHostMalloc) ret " << ret_code << " at " - << h << " is at node: "); + CONSOLE_PRINT("Memory (hipHostMalloc) ret %d at %p is at node: ", ret_code, h); } for (int i = 0; i < NUM_PAGES; i++) { - INFO(status[i]); // Always print it even if it's wrong. Verify later + CONSOLE_PRINT("%d ", status[i]); // Always print it even if it's wrong. Verify later } - INFO("\n"); + CONSOLE_PRINT("\n"); HIP_CHECK(hipHostFree(reinterpret_cast(h))); HIP_CHECK(hipHostUnregister(m)); @@ -129,8 +123,7 @@ bool test(int cpuId, int gpuId, int numaMode, unsigned int hostMallocflags) { if (cpuId >= 0 && (numaMode == MPOL_BIND || numaMode == MPOL_PREFERRED)) { for (int i = 0; i < NUM_PAGES; i++) { if (status[i] != cpuId) { // Now verify - WARN("Failed at " << i << " status[i] = " << status[i] - << " cpuId " << cpuId << "\n"); + WARN("Failed at " << i << " status[i] = " << status[i] << " cpuId " << cpuId << "\n"); return false; } } @@ -138,12 +131,12 @@ bool test(int cpuId, int gpuId, int numaMode, unsigned int hostMallocflags) { return true; } -bool runTest(const int &cpuCount, const int &gpuCount, - unsigned int hostMallocflags, const std::string &str) { - INFO("Test- " << str.c_str() << "\n"); +bool runTest(const int& cpuCount, const int& gpuCount, unsigned int hostMallocflags, + const std::string& str) { + CONSOLE_PRINT("Test- %s\n", str.c_str()); for (int m = 0; m < sizeof(mode) / sizeof(mode[0]); m++) { - INFO("Testing " << modeStr[m] << "\n"); + CONSOLE_PRINT("Testing %s\n", modeStr[m]); for (int i = 0; i < cpuCount; i++) { for (int j = 0; j < gpuCount; j++) { @@ -157,39 +150,40 @@ bool runTest(const int &cpuCount, const int &gpuCount, } /** -* Test Description -* ------------------------ -*  - Verify hipPerfHostNumaAlloc status. -* Test source -* ------------------------ -*  - perftests/memory/hipPerfHostNumaAlloc.cc -* Test requirements -* ------------------------ -*  - HIP_VERSION >= 5.6 -*/ + * Test Description + * ------------------------ + *  - Verify hipPerfHostNumaAlloc status. + * Test source + * ------------------------ + *  - perftests/memory/hipPerfHostNumaAlloc.cc + * Test requirements + * ------------------------ + *  - HIP_VERSION >= 5.6 + */ TEST_CASE("Perf_hipPerfHostNumaAlloc_test") { int gpuCount = 0; HIP_CHECK(hipGetDeviceCount(&gpuCount)); int cpuCount = getCpuAgentCount(); - INFO("Cpu count " << cpuCount << ", Gpu count " << gpuCount << "\n"); + CONSOLE_PRINT("Cpu count %d, Gpu count %d\n", cpuCount, gpuCount); if (cpuCount < 0 || gpuCount < 0) { - SUCCEED("Skipped testcase hipPerfHostNumaAlloc as " - "there is no device to test.\n"); + SUCCEED( + "Skipped testcase hipPerfHostNumaAlloc as " + "there is no device to test.\n"); return; } - REQUIRE(true == runTest(cpuCount, gpuCount, - hipHostMallocDefault | hipHostMallocNumaUser, - "Testing hipHostMallocDefault | hipHostMallocNumaUser......")); + REQUIRE(true == + runTest(cpuCount, gpuCount, hipHostMallocDefault | hipHostMallocNumaUser, + "Testing hipHostMallocDefault | hipHostMallocNumaUser......")); - REQUIRE(true == runTest(cpuCount, gpuCount, - hipHostMallocMapped | hipHostMallocNumaUser, - "Testing hipHostMallocMapped | hipHostMallocNumaUser.......")); + REQUIRE(true == + runTest(cpuCount, gpuCount, hipHostMallocMapped | hipHostMallocNumaUser, + "Testing hipHostMallocMapped | hipHostMallocNumaUser.......")); } /** -* End doxygen group perfMemoryTest. -* @} -*/ + * End doxygen group perfMemoryTest. + * @} + */ diff --git a/catch/perftests/memory/hipPerfMemFill.cc b/catch/perftests/memory/hipPerfMemFill.cc index 5892c10f86..32bfa622c0 100644 --- a/catch/perftests/memory/hipPerfMemFill.cc +++ b/catch/perftests/memory/hipPerfMemFill.cc @@ -18,20 +18,19 @@ */ /** -* @addtogroup hipMemcpyKernel hipMemcpyKernel -* @{ -* @ingroup perfMemoryTest -* `hipMemcpy(void* dst, const void* src, size_t count, hipMemcpyKind kind)` - -* Copies data between host and device. -*/ + * @addtogroup hipMemcpyKernel hipMemcpyKernel + * @{ + * @ingroup perfMemoryTest + * `hipMemcpy(void* dst, const void* src, size_t count, hipMemcpyKind kind)` - + * Copies data between host and device. + */ #include #define SIMPLY_ASSIGN 0 #define USE_HIPTEST_SETNUMBLOCKS 0 -template -__global__ void vec_fill(T *x, T coef, int N) { +template __global__ void vec_fill(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) { @@ -51,8 +50,7 @@ __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) { +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) { @@ -68,20 +66,17 @@ __global__ void vec_verify(T *x, T coef, int N) { } } -template -__global__ void daxpy(T *__restrict__ x, T *__restrict__ y, - const T coef, int Niter, int N) { +template +__global__ void daxpy(T* __restrict__ x, T* __restrict__ y, const T coef, int Niter, int N) { const int istart = threadIdx.x + blockIdx.x * blockDim.x; const int ishift = blockDim.x * gridDim.x; for (int iter = 0; iter < Niter; ++iter) { T iv = coef * iter; - for (int i = istart; i < N; i += ishift) - y[i] = iv * x[i] + y[i]; + for (int i = istart; i < N; i += ishift) y[i] = iv * x[i] + y[i]; } } -template -class hipPerfMemFill { +template class hipPerfMemFill { private: static constexpr int NUM_START = 27; static constexpr int NUM_SIZE = 4; @@ -96,26 +91,20 @@ class hipPerfMemFill { public: hipPerfMemFill() { for (int i = 0; i < NUM_SIZE; i++) { - // 128M, 256M, 512M, 1024M + // 128M, 256M, 512M, 1024M totalSizes_[i] = 1ull << (i + NUM_START); } } - ~hipPerfMemFill() { } + ~hipPerfMemFill() {} - bool supportLargeBar() { - return props_.isLargeBar != 0; - } + bool supportLargeBar() { return props_.isLargeBar != 0; } - bool supportManagedMemory() { - return props_.managedMemory != 0; - } + bool supportManagedMemory() { return props_.managedMemory != 0; } - const T getCoefficient(double val) { - return static_cast(val); - } + const T getCoefficient(double val) { return static_cast(val); } - void setHostBuffer(T *A, T val, size_t size) { + void setHostBuffer(T* A, T val, size_t size) { size_t len = size / sizeof(T); for (int i = 0; i < len; i++) { A[i] = val; @@ -138,33 +127,29 @@ class hipPerfMemFill { HIP_CHECK(hipGetDeviceProperties(&props_, deviceId)); blocksPerCU_ = props_.multiProcessorCount * 4; - std::cout << "Info: running on device: id: " << deviceId << ", bus: 0x" - << props_.pciBusID << " " << props_.name << " with " - << props_.multiProcessorCount << " CUs, large bar: " - << supportLargeBar() << ", managed memory: " << supportManagedMemory() - << ", DeviceMallocFinegrained: " << supportDeviceMallocFinegrained() - << std::endl; + std::cout << "Info: running on device: id: " << deviceId << ", bus: 0x" << props_.pciBusID + << " " << props_.name << " with " << props_.multiProcessorCount + << " CUs, large bar: " << supportLargeBar() + << ", managed memory: " << supportManagedMemory() + << ", DeviceMallocFinegrained: " << supportDeviceMallocFinegrained() << std::endl; return true; } void log_host(const char* title, double GBytes, double sec) { - std::cout << title << " [" << std::setw(7) << GBytes << " GB]: cost " - << std::setw(10) << sec << " s in bandwidth " << std::setw(10) - << GBytes / sec << " [GB/s]" << std::endl; + std::cout << title << " [" << std::setw(7) << GBytes << " GB]: cost " << std::setw(10) << sec + << " s in bandwidth " << std::setw(10) << GBytes / sec << " [GB/s]" << std::endl; } - void log_kernel(const char* title, double GBytes, double sec, - double sec_hv, double sec_kv) { - std::cout << title << " [" << std::setw(7) << GBytes << " GB]: cost " - << std::setw(10) << sec << " s in bandwidth " << std::setw(10) - << GBytes / sec << " [GB/s]" << ", hostVerify cost " - << std::setw(10) << sec_hv << " s in bandwidth " << std::setw(10) - << GBytes / sec_hv << " [GB/s]" << ", kernelVerify cost " - << std::setw(10) << sec_kv << " s in bandwidth " << std::setw(10) - << GBytes / sec_kv << " [GB/s]" << std::endl; + void log_kernel(const char* title, double GBytes, double sec, double sec_hv, double sec_kv) { + std::cout << title << " [" << std::setw(7) << GBytes << " GB]: cost " << std::setw(10) << sec + << " s in bandwidth " << std::setw(10) << GBytes / sec << " [GB/s]" + << ", hostVerify cost " << std::setw(10) << sec_hv << " s in bandwidth " + << std::setw(10) << GBytes / sec_hv << " [GB/s]" << ", kernelVerify cost " + << std::setw(10) << sec_kv << " s in bandwidth " << std::setw(10) << GBytes / sec_kv + << " [GB/s]" << std::endl; } - void hostFill(size_t size, T *data, T coef, double *sec) { + void hostFill(size_t size, T* data, T coef, double* sec) { size_t num = size / sizeof(T); // Size of elements auto start = std::chrono::steady_clock::now(); for (int i = 0; i < num; ++i) { @@ -179,29 +164,29 @@ class hipPerfMemFill { *sec = diff.count(); } - void kernelFill(size_t size, T *data, T coef, double *sec) { + void kernelFill(size_t size, T* data, T coef, double* sec) { size_t num = size / sizeof(T); // Size of elements unsigned blocks = setNumBlocks(num); // kernel will be loaded first time - hipLaunchKernelGGL(HIP_KERNEL_NAME(vec_fill), dim3(blocks), - dim3(threadsPerBlock_), 0, 0, data, 0, num); + hipLaunchKernelGGL(HIP_KERNEL_NAME(vec_fill), dim3(blocks), dim3(threadsPerBlock_), 0, 0, + data, 0, num); HIP_CHECK(hipDeviceSynchronize()); auto start = std::chrono::steady_clock::now(); for (int iter = 0; iter < NUM_ITER; ++iter) { - hipLaunchKernelGGL(HIP_KERNEL_NAME(vec_fill), dim3(blocks), - dim3(threadsPerBlock_), 0, 0, data, coef, num); + hipLaunchKernelGGL(HIP_KERNEL_NAME(vec_fill), dim3(blocks), dim3(threadsPerBlock_), 0, 0, + data, coef, num); } HIP_CHECK(hipDeviceSynchronize()); auto end = std::chrono::steady_clock::now(); std::chrono::duration diff = end - start; // in second - *sec = diff.count() / NUM_ITER; // in second + *sec = diff.count() / NUM_ITER; // in second } - void hostVerify(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 = std::chrono::steady_clock::now(); for (int i = 0; i < num; ++i) { @@ -224,27 +209,27 @@ class hipPerfMemFill { *sec = diff.count(); } - void kernelVerify(size_t size, T *data, T coef, double *sec) { + void kernelVerify(size_t size, T* data, T coef, double* sec) { size_t num = size / sizeof(T); // Size of elements unsigned blocks = setNumBlocks(num); // kernel will be loaded first time - hipLaunchKernelGGL(HIP_KERNEL_NAME(vec_verify), dim3(blocks), - dim3(threadsPerBlock_), 0, 0, data, coef, num); + hipLaunchKernelGGL(HIP_KERNEL_NAME(vec_verify), dim3(blocks), dim3(threadsPerBlock_), 0, 0, + data, coef, num); HIP_CHECK(hipDeviceSynchronize()); // Now all data verified. The following is to test bandwidth. auto start = std::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); + hipLaunchKernelGGL(HIP_KERNEL_NAME(vec_verify), dim3(blocks), dim3(threadsPerBlock_), 0, 0, + data, coef, num); } HIP_CHECK(hipDeviceSynchronize()); auto end = std::chrono::steady_clock::now(); std::chrono::duration diff = end - start; // in second - *sec = diff.count() / NUM_ITER; // in second + *sec = diff.count() / NUM_ITER; // in second } bool testLargeBarDeviceMemoryHostFill(size_t size) { @@ -254,7 +239,7 @@ class hipPerfMemFill { double GBytes = static_cast(size) / NUM_1GB; - T *A; + T* A; HIP_CHECK(hipMalloc(&A, size)); double sec = 0; hostFill(size, A, coef_, &sec); // Cpu can access device mem in LB @@ -285,7 +270,7 @@ class hipPerfMemFill { } double GBytes = static_cast(size) / NUM_1GB; - T *A; + T* A; HIP_CHECK(hipMallocManaged(&A, size)); double sec = 0; hostFill(size, A, coef_, &sec); // Cpu can access HMM mem @@ -301,7 +286,7 @@ class hipPerfMemFill { } double GBytes = static_cast(size) / NUM_1GB; - T *A; + T* A; HIP_CHECK(hipMallocManaged(&A, size)); double sec = 0, sec_hv = 0, sec_kv = 0; @@ -340,7 +325,7 @@ class hipPerfMemFill { bool testHostMemoryHostFill(size_t size, unsigned int flags) { double GBytes = static_cast(size) / NUM_1GB; - T *A; + T* A; HIP_CHECK(hipHostMalloc(&A, size, flags)); double sec = 0; hostFill(size, A, coef_, &sec); @@ -353,8 +338,8 @@ class hipPerfMemFill { bool testHostMemoryKernelFill(size_t size, unsigned int flags) { double GBytes = static_cast(size) / NUM_1GB; - T *A; - HIP_CHECK(hipHostMalloc(reinterpret_cast(&A), size, flags)); + T* A; + HIP_CHECK(hipHostMalloc(reinterpret_cast(&A), size, flags)); double sec = 0, sec_hv = 0, sec_kv = 0; kernelFill(size, A, coef_, &sec); hostVerify(size, A, coef_, &sec_hv); @@ -400,10 +385,11 @@ class hipPerfMemFill { /* This function should be via device attribute query*/ bool supportDeviceMallocFinegrained() { #ifdef __HIP_PLATFORM_AMD__ - T *A = nullptr; + T* A = nullptr; hipError_t err; - err = hipExtMallocWithFlags(reinterpret_cast(&A), sizeof(T), - hipDeviceMallocFinegrained); + + err = + hipExtMallocWithFlags(reinterpret_cast(&A), sizeof(T), hipDeviceMallocFinegrained); if (err || !A) { return false; } @@ -415,7 +401,7 @@ class hipPerfMemFill { } unsigned int setNumBlocks(size_t size) { - size_t num = size/sizeof(T); + size_t num = size / sizeof(T); #if USE_HIPTEST_SETNUMBLOCKS return HipTest::setNumBlocks(blocksPerCU_, threadsPerBlock_, num); @@ -428,12 +414,11 @@ class hipPerfMemFill { bool testExtDeviceMemoryHostFill(size_t size, unsigned int flags) { double GBytes = static_cast(size) / NUM_1GB; - T *A = nullptr; - HIP_CHECK(hipExtMallocWithFlags(reinterpret_cast(&A), - size, flags)); + T* A = nullptr; + HIP_CHECK(hipExtMallocWithFlags(reinterpret_cast(&A), size, flags)); if (!A) { - std::cout << "failed hipExtMallocWithFlags() with size =" << - size << " flags="<< std::hex << flags << std::endl; + std::cout << "failed hipExtMallocWithFlags() with size =" << size << " flags=" << std::hex + << flags << std::endl; return false; } @@ -448,12 +433,11 @@ class hipPerfMemFill { bool testExtDeviceMemoryKernelFill(size_t size, unsigned int flags) { double GBytes = static_cast(size) / NUM_1GB; - T *A = nullptr; - HIP_CHECK(hipExtMallocWithFlags(reinterpret_cast(&A), - size, flags)); + T* A = nullptr; + HIP_CHECK(hipExtMallocWithFlags(reinterpret_cast(&A), size, flags)); if (!A) { - std::cout << "failed hipExtMallocWithFlags() with size =" << - size << " flags=" << std::hex << flags << std::endl; + std::cout << "failed hipExtMallocWithFlags() with size =" << size << " flags=" << std::hex + << flags << std::endl; return false; } @@ -470,20 +454,16 @@ class hipPerfMemFill { } bool testExtDeviceMemory() { - std::cout << "Test fine grained device memory host filling" - << std::endl; + std::cout << "Test fine grained device memory host filling" << std::endl; for (int i = 0; i < NUM_SIZE; i++) { - if (!testExtDeviceMemoryHostFill(totalSizes_[i], - hipDeviceMallocFinegrained)) { + if (!testExtDeviceMemoryHostFill(totalSizes_[i], hipDeviceMallocFinegrained)) { return false; } } - std::cout << "Test fine grained device memory kernel filling" - << std::endl; + std::cout << "Test fine grained device memory kernel filling" << std::endl; for (int i = 0; i < NUM_SIZE; i++) { - if (!testExtDeviceMemoryKernelFill(totalSizes_[i], - hipDeviceMallocFinegrained)) { + if (!testExtDeviceMemoryKernelFill(totalSizes_[i], hipDeviceMallocFinegrained)) { return false; } } @@ -521,16 +501,16 @@ class hipPerfMemFill { }; /** -* Test Description -* ------------------------ -*  - Verify hipPerfMemFill status. -* Test source -* ------------------------ -*  - perftests/memory/hipPerfMemFill.cc -* Test requirements -* ------------------------ -*  - HIP_VERSION >= 5.6 -*/ + * Test Description + * ------------------------ + *  - Verify hipPerfMemFill status. + * Test source + * ------------------------ + *  - perftests/memory/hipPerfMemFill.cc + * Test requirements + * ------------------------ + *  - HIP_VERSION >= 5.6 + */ TEST_CASE("Perf_hipPerfMemFill_test") { std::cout << "Test int" << std::endl; @@ -545,6 +525,6 @@ TEST_CASE("Perf_hipPerfMemFill_test") { } /** -* End doxygen group perfMemoryTest. -* @} -*/ + * End doxygen group perfMemoryTest. + * @} + */ diff --git a/catch/perftests/memory/hipPerfMemMallocCpyFree.cc b/catch/perftests/memory/hipPerfMemMallocCpyFree.cc index 3960a16bd8..e48bc53e56 100644 --- a/catch/perftests/memory/hipPerfMemMallocCpyFree.cc +++ b/catch/perftests/memory/hipPerfMemMallocCpyFree.cc @@ -18,13 +18,13 @@ THE SOFTWARE. */ /** -* @addtogroup hipMemcpy hipMemcpy -* @{ -* @ingroup perfMemoryTest -* `hipMemcpy(void* dst, const void* src, size_t count, hipMemcpyKind kind)` - -* Copies data between host and device. -*/ - + * @addtogroup hipMemcpy hipMemcpy + * @{ + * @ingroup perfMemoryTest + * `hipMemcpy(void* dst, const void* src, size_t count, hipMemcpyKind kind)` - + * Copies data between host and device. + */ +// #define ENABLE_DEBUG 1 #include #include @@ -38,7 +38,7 @@ void valSet(int* A, int val, size_t size) { } } -void setup(size_t *size, int *num, int **pA, const size_t totalGlobalMem) { +void setup(size_t* size, int* num, int** pA, const size_t totalGlobalMem) { for (int i = 0; i < *num; i++) { size[i] = 1 << (i + 6); if ((NUM_ITER + 1) * size[i] > totalGlobalMem) { @@ -50,39 +50,39 @@ void setup(size_t *size, int *num, int **pA, const size_t totalGlobalMem) { valSet(*pA, 1, size[*num - 1]); } -void testInit(size_t size, int *A) { - int *Ad; +void testInit(size_t size, int* A) { + int* Ad; clock_t start = clock(); - HIP_CHECK(hipMalloc(&Ad, size)); // hip::init() will be called + HIP_CHECK(hipMalloc(&Ad, size)); // hip::init() will be called clock_t end = clock(); double uS = (end - start) * 1000000. / CLOCKS_PER_SEC; - INFO("Initial: hipMalloc(" << size << ") cost " << uS << "us" << "\n"); + CONSOLE_PRINT("Initial: hipMalloc(%zu) cost %.2fus\n", size, uS); start = clock(); HIP_CHECK(hipMemcpy(Ad, A, size, hipMemcpyHostToDevice)); HIP_CHECK(hipDeviceSynchronize()); end = clock(); uS = (end - start) * 1000000. / CLOCKS_PER_SEC; - INFO("hipMemcpy(" << size << ") cost " << uS << "us" << "\n"); + CONSOLE_PRINT("hipMemcpy(%zu) cost %.2fus\n", size, uS); start = clock(); HIP_CHECK(hipFree(Ad)); end = clock(); uS = (end - start) * 1000000. / CLOCKS_PER_SEC; - INFO("hipFree(" << size << ") cost " << uS << "us" << "\n"); + CONSOLE_PRINT("hipFree(%zu) cost %.2fus\n", size, uS); } static bool hipPerfMemMallocCpyFree_test() { double uS; clock_t start, end; - size_t size[NUM_SIZE] = { 0 }; - int *Ad[NUM_ITER] = { nullptr }; - int *A; + size_t size[NUM_SIZE] = {0}; + int* Ad[NUM_ITER] = {nullptr}; + int* A; hipDeviceProp_t props; memset(&props, 0, sizeof(props)); HIP_CHECK(hipGetDeviceProperties(&props, 0)); - INFO("totalGlobalMem: " << props.totalGlobalMem << "\n"); + CONSOLE_PRINT("totalGlobalMem: %zu\n", props.totalGlobalMem); int num = NUM_SIZE; setup(size, &num, &A, props.totalGlobalMem); @@ -91,59 +91,60 @@ static bool hipPerfMemMallocCpyFree_test() { for (int i = 0; i < num; i++) { start = clock(); for (int j = 0; j < NUM_ITER; j++) { - HIP_CHECK(hipMalloc(&Ad[j], size[i])); + HIP_CHECK(hipMalloc(&Ad[j], size[i])); } end = clock(); uS = (end - start) * 1000000. / (NUM_ITER * CLOCKS_PER_SEC); - INFO("hipMalloc(" << size[i] << ") cost " << uS << "us" << "\n"); + CONSOLE_PRINT("hipMalloc(%zu) cost %.2fus\n", size[i], uS); start = clock(); for (int j = 0; j < NUM_ITER; j++) { - HIP_CHECK(hipMemcpy(Ad[j], A, size[i], hipMemcpyHostToDevice)); + HIP_CHECK(hipMemcpy(Ad[j], A, size[i], hipMemcpyHostToDevice)); } HIP_CHECK(hipDeviceSynchronize()); end = clock(); uS = (end - start) * 1000000. / (NUM_ITER * CLOCKS_PER_SEC); - INFO("hipMemcpy(" << size[i] << ") cost " << uS << "us" << "\n"); + CONSOLE_PRINT("hipMemcpy(%zu) cost %.2fus\n", size[i], uS); start = clock(); for (int j = 0; j < NUM_ITER; j++) { - HIP_CHECK(hipFree(Ad[j])); - Ad[j] = nullptr; + HIP_CHECK(hipFree(Ad[j])); + Ad[j] = nullptr; } end = clock(); double uS = (end - start) * 1000000. / (NUM_ITER * CLOCKS_PER_SEC); - INFO("hipFree(" << size[i] << ") cost " << uS << "us" << "\n"); + CONSOLE_PRINT("hipFree(%zu) cost %.2fus\n", size[i], uS); } free(A); return true; } /** -* Test Description -* ------------------------ -*  - Verify hipPerfMemMallocCpyFree status. -* Test source -* ------------------------ -*  - perftests/memory/hipPerfMemMallocCpyFree.cc -* Test requirements -* ------------------------ -*  - HIP_VERSION >= 5.6 -*/ + * Test Description + * ------------------------ + *  - Verify hipPerfMemMallocCpyFree status. + * Test source + * ------------------------ + *  - perftests/memory/hipPerfMemMallocCpyFree.cc + * Test requirements + * ------------------------ + *  - HIP_VERSION >= 5.6 + */ TEST_CASE("Perf_hipPerfMemMallocCpyFree_test") { int numDevices = 0; HIP_CHECK(hipGetDeviceCount(&numDevices)); if (numDevices <= 0) { - SUCCEED("Skipped testcase hipPerfDevMemReadSpeed as" - "there is no device to test."); + SUCCEED( + "Skipped testcase hipPerfDevMemReadSpeed as" + "there is no device to test."); } else { REQUIRE(true == hipPerfMemMallocCpyFree_test()); } } /** -* End doxygen group perfMemoryTest. -* @} -*/ + * End doxygen group perfMemoryTest. + * @} + */ diff --git a/catch/perftests/memory/hipPerfMemcpy.cc b/catch/perftests/memory/hipPerfMemcpy.cc index badc3d43ca..d5ad0786aa 100644 --- a/catch/perftests/memory/hipPerfMemcpy.cc +++ b/catch/perftests/memory/hipPerfMemcpy.cc @@ -18,15 +18,15 @@ */ /** -* @addtogroup hipMemcpy hipMemcpy -* @{ -* @ingroup perfMemoryTest -* `hipMemcpy(void* dst, const void* src, size_t count, hipMemcpyKind kind)` - -* Copies data between host and device. -*/ + * @addtogroup hipMemcpy hipMemcpy + * @{ + * @ingroup perfMemoryTest + * `hipMemcpy(void* dst, const void* src, size_t count, hipMemcpyKind kind)` - + * Copies data between host and device. + */ #include - +// #define ENABLE_DEBUG 1 #define NUM_SIZE 14 #define NUM_ITER 1000 // max BW number for DevicetoDeviceNoCU @@ -35,7 +35,8 @@ class hipPerfMemcpy { private: size_t totalSizes_[NUM_SIZE]; - void setHostBuffer(int *A, int val, size_t size); + void setHostBuffer(int* A, int val, size_t size); + public: hipPerfMemcpy(); ~hipPerfMemcpy() {} @@ -53,7 +54,7 @@ hipPerfMemcpy::hipPerfMemcpy() { } } -void hipPerfMemcpy::setHostBuffer(int *A, int val, size_t size) { +void hipPerfMemcpy::setHostBuffer(int* A, int val, size_t size) { size_t len = size / sizeof(int); for (int i = 0; i < len; i++) { A[i] = val; @@ -61,36 +62,31 @@ void hipPerfMemcpy::setHostBuffer(int *A, int val, size_t size) { } void hipPerfMemcpy::TestResult(unsigned int numTests, - std::chrono::duration diff, hipMemcpyKind type) -{ + std::chrono::duration diff, hipMemcpyKind type) { // BW in GB/s - double perf = (static_cast(totalSizes_[numTests] * NUM_ITER) * - static_cast(1e-03)) / diff.count(); + double perf = + (static_cast(totalSizes_[numTests] * NUM_ITER) * static_cast(1e-03)) / + diff.count(); - const char *typestr = NULL; + const char* typestr = NULL; - if(type == hipMemcpyHostToDevice){ - typestr = "Host to Device"; - } - else if(type == hipMemcpyDeviceToHost){ - typestr = "Device to Host"; - } - else if(type == hipMemcpyDeviceToDevice){ - typestr = "Device to Device"; - perf *= 2.0; - } - else if(type == hipMemcpyDeviceToDeviceNoCU){ - typestr = "Device to Device No CU"; - perf *= 2.0; + if (type == hipMemcpyHostToDevice) { + typestr = "Host to Device"; + } else if (type == hipMemcpyDeviceToHost) { + typestr = "Device to Host"; + } else if (type == hipMemcpyDeviceToDevice) { + typestr = "Device to Device"; + perf *= 2.0; + } else if (type == hipMemcpyDeviceToDeviceNoCU) { + typestr = "Device to Device No CU"; + perf *= 2.0; } - UNSCOPED_INFO("hipPerfMemcpy[" << numTests << "] " << typestr << " copy BW " - << perf << " GB/s for memory size of " << - totalSizes_[numTests] << " Bytes."); - - if(totalSizes_[numTests] == 4194304 && type == hipMemcpyDeviceToDeviceNoCU) - REQUIRE(perf < NOCU_MAX_BW); + CONSOLE_PRINT("hipPerfMemcpy[%d] %s copy BW %.2f GB/s for memory size of %lu Bytes.\n", numTests, + typestr, perf, totalSizes_[numTests]); + if (totalSizes_[numTests] == 4194304 && type == hipMemcpyDeviceToDeviceNoCU) + REQUIRE(perf < NOCU_MAX_BW); } bool hipPerfMemcpy::run_h2d(unsigned int numTests) { @@ -115,7 +111,7 @@ bool hipPerfMemcpy::run_h2d(unsigned int numTests) { TestResult(numTests, diff, hipMemcpyHostToDevice); HIP_CHECK(hipHostUnregister(A)); - delete [] A; + delete[] A; HIP_CHECK(hipFree(Ad)); return true; @@ -143,7 +139,7 @@ bool hipPerfMemcpy::run_d2h(unsigned int numTests) { TestResult(numTests, diff, hipMemcpyDeviceToHost); HIP_CHECK(hipHostUnregister(A)); - delete [] A; + delete[] A; HIP_CHECK(hipFree(Ad)); return true; @@ -186,8 +182,8 @@ bool hipPerfMemcpy::run_d2d_nocu(unsigned int numTests) { auto all_start = std::chrono::steady_clock::now(); for (int j = 0; j < NUM_ITER; j++) { - HIP_CHECK(hipMemcpyAsync(Ad1, Ad2, totalSizes_[numTests], hipMemcpyDeviceToDeviceNoCU, - nullptr)); + HIP_CHECK( + hipMemcpyAsync(Ad1, Ad2, totalSizes_[numTests], hipMemcpyDeviceToDeviceNoCU, nullptr)); } HIP_CHECK(hipDeviceSynchronize()); @@ -204,16 +200,16 @@ bool hipPerfMemcpy::run_d2d_nocu(unsigned int numTests) { } /** -* Test Description -* ------------------------ -*  - Verify hipPerfMemcpy status. -* Test source -* ------------------------ -*  - perftests/memory/hipPerfMemcpy.cc -* Test requirements -* ------------------------ -*  - HIP_VERSION >= 5.6 -*/ + * Test Description + * ------------------------ + *  - Verify hipPerfMemcpy status. + * Test source + * ------------------------ + *  - perftests/memory/hipPerfMemcpy.cc + * Test requirements + * ------------------------ + *  - HIP_VERSION >= 5.6 + */ TEST_CASE("Perf_hipPerfMemcpy_test") { int numDevices = 0; @@ -227,35 +223,34 @@ TEST_CASE("Perf_hipPerfMemcpy_test") { hipDeviceProp_t props; HIP_CHECK(hipGetDeviceProperties(&props, deviceId)); - UNSCOPED_INFO("info: running on bus " << "0x" << props.pciBusID << " " << - props.name << " with " << props.multiProcessorCount << " CUs " - << " and device id: " << deviceId); + CONSOLE_PRINT("info: running on bus 0x%x %s with %d CUs and device id: %d\n", props.pciBusID, + props.name, props.multiProcessorCount, deviceId); hipPerfMemcpy hipPerfMemcpy; - SECTION("Perf test Host Memory to Device Memory"){ + SECTION("Perf test Host Memory to Device Memory") { for (auto testCase = 0; testCase < NUM_SIZE; testCase++) { - REQUIRE(true == hipPerfMemcpy.run_h2d(testCase)); + REQUIRE(true == hipPerfMemcpy.run_h2d(testCase)); } } - SECTION("Perf test Device Memory to Host Memory"){ + SECTION("Perf test Device Memory to Host Memory") { for (auto testCase = 0; testCase < NUM_SIZE; testCase++) { - REQUIRE(true == hipPerfMemcpy.run_d2h(testCase)); + REQUIRE(true == hipPerfMemcpy.run_d2h(testCase)); } } - SECTION("Perf test Device Memory to Device Memory"){ + SECTION("Perf test Device Memory to Device Memory") { for (auto testCase = 0; testCase < NUM_SIZE; testCase++) { - REQUIRE(true == hipPerfMemcpy.run_d2d(testCase)); + REQUIRE(true == hipPerfMemcpy.run_d2d(testCase)); } } - SECTION("Perf test Device Memory to Device Memory No CU"){ + SECTION("Perf test Device Memory to Device Memory No CU") { for (auto testCase = 0; testCase < NUM_SIZE; testCase++) { - REQUIRE(true == hipPerfMemcpy.run_d2d_nocu(testCase)); + REQUIRE(true == hipPerfMemcpy.run_d2d_nocu(testCase)); } } } } /** -* End doxygen group perfMemoryTest. -* @} -*/ + * End doxygen group perfMemoryTest. + * @} + */ diff --git a/catch/perftests/memory/hipPerfMemset.cc b/catch/perftests/memory/hipPerfMemset.cc index 5b84e91942..aaf434245b 100644 --- a/catch/perftests/memory/hipPerfMemset.cc +++ b/catch/perftests/memory/hipPerfMemset.cc @@ -18,30 +18,29 @@ */ /** -* @addtogroup hipMemsetKernel hipMemsetKernel -* @{ -* @ingroup perfMemoryTest -* `hipMemset(void* devPtr, int value, size_t count)` - -* Initializes or sets device memory to a value. -*/ - + * @addtogroup hipMemsetKernel hipMemsetKernel + * @{ + * @ingroup perfMemoryTest + * `hipMemset(void* devPtr, int value, size_t count)` - + * Initializes or sets device memory to a value. + */ +// #define ENABLE_DEBUG 1 #include static unsigned int sizeList[] = { - 256, 512, 1024, 2048, 4096, 8192, + 256, 512, 1024, 2048, 4096, 8192, }; -static unsigned int eleNumList[] = { - 0x100, 0x400, 0x1000, 0x4000, 0x10000, 0x20000, 0x40000, 0x80000, 0x100000, - 0x200000, 0x400000, 0x800000, 0x1000000 -}; +static unsigned int eleNumList[] = {0x100, 0x400, 0x1000, 0x4000, 0x10000, + 0x20000, 0x40000, 0x80000, 0x100000, 0x200000, + 0x400000, 0x800000, 0x1000000}; typedef struct _dataType { char memsetval = 0x42; char memsetD8val = 0xDE; int16_t memsetD16val = 0xDEAD; int memsetD32val = 0xDEADBEEF; -}dataType; +} dataType; #define NUM_ITER 1000 @@ -56,7 +55,7 @@ enum MemsetType { class hipPerfMemset { private: - uint64_t bufSize_; + uint64_t bufSize_; unsigned int num_elements_; unsigned int testNumEle_; unsigned int _numSubTests = 0; @@ -78,25 +77,19 @@ class hipPerfMemset { bool open(int deviceID); - template + template void run1D(unsigned int test, T memsetval, enum MemsetType type, bool async); - template + template void run2D(unsigned int test, T memsetval, enum MemsetType type, bool async); - template + template void run3D(unsigned int test, T memsetval, enum MemsetType type, bool async); - uint getNumTests() { - return _numSubTests; - } + uint getNumTests() { return _numSubTests; } - uint getNumTests2D() { - return _numSubTests2D; - } - uint getNumTests3D() { - return _numSubTests3D; - } + uint getNumTests2D() { return _numSubTests2D; } + uint getNumTests3D() { return _numSubTests3D; } }; bool hipPerfMemset::open(int deviceId) { @@ -109,15 +102,13 @@ bool hipPerfMemset::open(int deviceId) { HIP_CHECK(hipSetDevice(deviceId)); hipDeviceProp_t props; HIP_CHECK(hipGetDeviceProperties(&props, deviceId)); - INFO("info: running on bus " << "0x" << props.pciBusID << " " << props.name - << " with " << props.multiProcessorCount << " CUs and device id: " - << deviceId << "\n"); + CONSOLE_PRINT("info: running on bus 0x%x %s with %d CUs and device id: %d\n", props.pciBusID, + props.name, props.multiProcessorCount, deviceId); return true; } -template -void hipPerfMemset::run1D(unsigned int test, T memsetval, - enum MemsetType type, bool async) { +template +void hipPerfMemset::run1D(unsigned int test, T memsetval, enum MemsetType type, bool async) { T *A_h, *A_d; testNumEle_ = eleNumList[test % num_elements_]; @@ -126,17 +117,17 @@ void hipPerfMemset::run1D(unsigned int test, T memsetval, HIP_CHECK(hipMalloc(&A_d, bufSize_)); - A_h = reinterpret_cast (malloc(bufSize_)); + A_h = reinterpret_cast(malloc(bufSize_)); hipStream_t stream; HIP_CHECK(hipStreamCreateWithFlags(&stream, hipStreamNonBlocking)); // Warm-up if (async) { - HIP_CHECK(hipMemsetAsync((void *)A_d, memsetval, bufSize_, stream)); + HIP_CHECK(hipMemsetAsync((void*)A_d, memsetval, bufSize_, stream)); HIP_CHECK(hipStreamSynchronize(stream)); } else { - HIP_CHECK(hipMemset((void *)A_d, memsetval, bufSize_)); + HIP_CHECK(hipMemset((void*)A_d, memsetval, bufSize_)); HIP_CHECK(hipDeviceSynchronize()); } @@ -144,7 +135,7 @@ void hipPerfMemset::run1D(unsigned int test, T memsetval, for (uint i = 0; i < NUM_ITER; i++) { if (type == hipMemsetTypeDefault && !async) { - HIP_CHECK(hipMemset(reinterpret_cast(A_d), memsetval, bufSize_)); + HIP_CHECK(hipMemset(reinterpret_cast(A_d), memsetval, bufSize_)); } else if (type == hipMemsetTypeDefault && async) { HIP_CHECK(hipMemsetAsync(A_d, memsetval, bufSize_, stream)); } else if (type == hipMemsetTypeD8 && !async) { @@ -152,13 +143,13 @@ void hipPerfMemset::run1D(unsigned int test, T memsetval, } else if (type == hipMemsetTypeD8 && async) { HIP_CHECK(hipMemsetD8Async((hipDeviceptr_t)A_d, memsetval, bufSize_, stream)); } else if (type == hipMemsetTypeD16 && !async) { - HIP_CHECK(hipMemsetD16((hipDeviceptr_t)A_d, memsetval, bufSize_/sizeof(T))); + HIP_CHECK(hipMemsetD16((hipDeviceptr_t)A_d, memsetval, bufSize_ / sizeof(T))); } else if (type == hipMemsetTypeD16 && async) { - HIP_CHECK(hipMemsetD16Async((hipDeviceptr_t)A_d, memsetval, bufSize_/sizeof(T), stream)); + HIP_CHECK(hipMemsetD16Async((hipDeviceptr_t)A_d, memsetval, bufSize_ / sizeof(T), stream)); } else if (type == hipMemsetTypeD32 && !async) { - HIP_CHECK(hipMemsetD32((hipDeviceptr_t)A_d, memsetval, bufSize_/sizeof(T))); + HIP_CHECK(hipMemsetD32((hipDeviceptr_t)A_d, memsetval, bufSize_ / sizeof(T))); } else if (type == hipMemsetTypeD32 && async) { - HIP_CHECK(hipMemsetD32Async((hipDeviceptr_t)A_d, memsetval, bufSize_/sizeof(T), stream)); + HIP_CHECK(hipMemsetD32Async((hipDeviceptr_t)A_d, memsetval, bufSize_ / sizeof(T), stream)); } } if (async) { @@ -169,13 +160,12 @@ void hipPerfMemset::run1D(unsigned int test, T memsetval, auto end = std::chrono::steady_clock::now(); - HIP_CHECK(hipMemcpy(A_h, A_d, bufSize_, hipMemcpyDeviceToHost) ); + HIP_CHECK(hipMemcpy(A_h, A_d, bufSize_, hipMemcpyDeviceToHost)); for (int i = 0; i < bufSize_ / sizeof(T); i++) { if (A_h[i] != memsetval) { - INFO("mismatch at index " << i << " computed: " << - static_cast (A_h[i]) << ", memsetval: " << - static_cast (memsetval) << "\n"); + DEBUG_PRINT("mismatch at index %d computed: %d, memsetval: %d\n", i, static_cast(A_h[i]), + static_cast(memsetval)); REQUIRE(false); } } @@ -188,30 +178,27 @@ void hipPerfMemset::run1D(unsigned int test, T memsetval, auto sec = diff.count(); auto perf = static_cast((bufSize_ * NUM_ITER * (1e-09)) / sec); - std::cout << "[" << std::setw(2) - << test << "] " << std::setw(5) << bufSize_/1024 - << " Kb " << std::setw(4) << " typeSize " << sizeof(T) << " : " - << std::setw(7) << perf << " GB/s \n"; + std::cout << "[" << std::setw(2) << test << "] " << std::setw(5) << bufSize_ / 1024 << " Kb " + << std::setw(4) << " typeSize " << sizeof(T) << " : " << std::setw(7) << perf + << " GB/s \n"; } -template -void hipPerfMemset::run2D(unsigned int test, T memsetval, - enum MemsetType type, bool async) { +template +void hipPerfMemset::run2D(unsigned int test, T memsetval, enum MemsetType type, bool async) { bufSize_ = sizeList[test % num_sizes_]; size_t numH = bufSize_; size_t numW = bufSize_; size_t pitch_A; size_t width = numW * sizeof(char); size_t sizeElements = width * numH; - size_t elements = numW* numH; + size_t elements = numW * numH; - T * A_h, * A_d; + T *A_h, *A_d; - HIP_CHECK(hipMallocPitch(reinterpret_cast(&A_d), - &pitch_A, width, numH)); + HIP_CHECK(hipMallocPitch(reinterpret_cast(&A_d), &pitch_A, width, numH)); A_h = reinterpret_cast(malloc(sizeElements)); - for (size_t i=0; i < elements; i++) { + for (size_t i = 0; i < elements; i++) { A_h[i] = 1; } @@ -244,14 +231,12 @@ void hipPerfMemset::run2D(unsigned int test, T memsetval, auto end = std::chrono::steady_clock::now(); - HIP_CHECK(hipMemcpy2D(A_h, width, A_d, pitch_A, numW, numH, - hipMemcpyDeviceToHost)); + HIP_CHECK(hipMemcpy2D(A_h, width, A_d, pitch_A, numW, numH, hipMemcpyDeviceToHost)); - for (int i=0; i < elements; i++) { + for (int i = 0; i < elements; i++) { if (A_h[i] != memsetval) { - INFO("mismatch at index " << i << " computed: " << - static_cast (A_h[i]) << ", memsetval: " << - static_cast (memsetval) << "\n"); + DEBUG_PRINT("mismatch at index %d computed: %d, memsetval: %d\n", i, static_cast(A_h[i]), + static_cast(memsetval)); REQUIRE(false); } } @@ -259,20 +244,19 @@ void hipPerfMemset::run2D(unsigned int test, T memsetval, std::chrono::duration diff = end - start; auto sec = diff.count(); - auto perf = static_cast((sizeElements* NUM_ITER * (1e-09)) / sec); + auto perf = static_cast((sizeElements * NUM_ITER * (1e-09)) / sec); - std::cout << "hipPerf2DMemset" << (async ? "Async" : " ") << "[" << test << "] " - << " " << "(GB/s) for " << std::setw(5) << bufSize_ - << " x " << std::setw(5) << bufSize_ << " bytes : " << std::setw(7) << perf << "\n"; + std::cout << "hipPerf2DMemset" << (async ? "Async" : " ") << "[" << test << "] " << " " + << "(GB/s) for " << std::setw(5) << bufSize_ << " x " << std::setw(5) << bufSize_ + << " bytes : " << std::setw(7) << perf << "\n"; HIP_CHECK(hipStreamDestroy(stream)); HIP_CHECK(hipFree(A_d)); free(A_h); } -template -void hipPerfMemset::run3D(unsigned int test, T memsetval, - enum MemsetType type, bool async) { +template +void hipPerfMemset::run3D(unsigned int test, T memsetval, enum MemsetType type, bool async) { bufSize_ = sizeList[test % num_sizes_]; size_t numH = bufSize_; @@ -280,12 +264,12 @@ void hipPerfMemset::run3D(unsigned int test, T memsetval, size_t depth = 10; size_t width = numW * sizeof(char); size_t sizeElements = width * numH * depth; - size_t elements = numW* numH* depth; + size_t elements = numW * numH * depth; hipStream_t stream; HIP_CHECK(hipStreamCreateWithFlags(&stream, hipStreamNonBlocking)); - T *A_h; + T* A_h; hipExtent extent = make_hipExtent(width, numH, depth); hipPitchedPtr devPitchedPtr; @@ -325,12 +309,12 @@ void hipPerfMemset::run3D(unsigned int test, T memsetval, auto end = std::chrono::steady_clock::now(); - hipMemcpy3DParms myparms ; + hipMemcpy3DParms myparms; myparms.srcArray = nullptr; myparms.dstArray = nullptr; myparms.srcPos = make_hipPos(0, 0, 0); myparms.dstPos = make_hipPos(0, 0, 0); - myparms.dstPtr = make_hipPitchedPtr(A_h, width , numW, numH); + myparms.dstPtr = make_hipPitchedPtr(A_h, width, numW, numH); myparms.srcPtr = devPitchedPtr; myparms.extent = extent; @@ -338,11 +322,10 @@ void hipPerfMemset::run3D(unsigned int test, T memsetval, HIP_CHECK(hipMemcpy3D(&myparms)); - for (int i=0; i < elements; i++) { + for (int i = 0; i < elements; i++) { if (A_h[i] != memsetval) { - INFO("mismatch at index " << i << " computed: " << - static_cast (A_h[i]) << ", memsetval: " << - static_cast (memsetval) << "\n"); + DEBUG_PRINT("mismatch at index %d computed: %d, memsetval: %d\n", i, static_cast(A_h[i]), + static_cast(memsetval)); REQUIRE(false); } } @@ -352,24 +335,23 @@ void hipPerfMemset::run3D(unsigned int test, T memsetval, auto sec = diff.count(); auto perf = static_cast((sizeElements * NUM_ITER * (1e-09)) / sec); - std::cout << "hipPerf3DMemset" << (async ? "Async" : " ") << "[" << test << "] " << " " - << "(GB/s) for " << std::setw(5) << bufSize_ << " x " << std::setw(5) - << bufSize_ << " x " << depth << " bytes : " << std::setw(7) << perf << "\n"; + CONSOLE_PRINT("hipPerf3DMemset%s[%d] (GB/s) for %5lu x %5lu x %lu bytes : %7.2f\n", + (async ? "Async" : " "), test, bufSize_, bufSize_, depth, perf); HIP_CHECK(hipFree(devPitchedPtr.ptr)); free(A_h); } /** -* Test Description -* ------------------------ -*  - Verify hipPerfMemset status. -* Test source -* ------------------------ -*  - perftests/memory/hipPerfMemset.cc -* Test requirements -* ------------------------ -*  - HIP_VERSION >= 5.6 -*/ + * Test Description + * ------------------------ + *  - Verify hipPerfMemset status. + * Test source + * ------------------------ + *  - perftests/memory/hipPerfMemset.cc + * Test requirements + * ------------------------ + *  - HIP_VERSION >= 5.6 + */ TEST_CASE("Perf_hipPerfMemset_test") { hipPerfMemset hipPerfMemset; @@ -385,44 +367,44 @@ TEST_CASE("Perf_hipPerfMemset_test") { bool async = false; - for (uint i = 0; i < 2 ; i++) { - std::cout << "--------------------- 1D buffer -------------------\n"; + for (uint i = 0; i < 2; i++) { + CONSOLE_PRINT("--------------------- 1D buffer -------------------\n"); for (auto testCase = 0; testCase < numTests; testCase++) { if (testCase < sizeof(eleNumList) / sizeof(uint32_t)) { - std::cout << "hipMemsetD8" << (async ? "Async " : " "); + CONSOLE_PRINT("hipMemsetD8%s", (async ? "Async " : " ")); hipPerfMemset.run1D(testCase, pattern.memsetval, hipMemsetTypeD8, async); } else if (testCase < 2 * sizeof(eleNumList) / sizeof(uint32_t)) { - std::cout << "hipMemsetD16" << (async ? "Async" : " "); + CONSOLE_PRINT("hipMemsetD16%s", (async ? "Async" : " ")); hipPerfMemset.run1D(testCase, pattern.memsetD16val, hipMemsetTypeD16, async); } else if (testCase < 3 * sizeof(eleNumList) / sizeof(uint32_t)) { - std::cout << "hipMemsetD32" << (async ? "Async" : " "); + CONSOLE_PRINT("hipMemsetD32%s", (async ? "Async" : " ")); hipPerfMemset.run1D(testCase, pattern.memsetD32val, hipMemsetTypeD32, async); } else { - std::cout << "hipMemset" << (async ? "Async " : " "); + CONSOLE_PRINT("hipMemset%s", (async ? "Async " : " ")); hipPerfMemset.run1D(testCase, pattern.memsetval, hipMemsetTypeDefault, async); } } async = true; } - INFO("\n"); - std::cout << "------------------ 2D buffer arrays ---------------\n"; + CONSOLE_PRINT("\n"); + CONSOLE_PRINT("\n------------------ 2D buffer arrays ---------------\n"); async = false; for (uint i = 0; i < 2; i++) { - INFO("\n"); + CONSOLE_PRINT("\n"); for (uint test = 0; test < numTests2D; test++) { hipPerfMemset.run2D(test, pattern.memsetval, hipMemsetTypeDefault, async); } async = true; } - INFO("\n"); - std::cout << "------------------ 3D buffer arrays ---------------\n"; + CONSOLE_PRINT("\n"); + CONSOLE_PRINT("\n------------------ 3D buffer arrays ---------------\n"); async = false; for (uint i = 0; i < 2; i++) { - INFO("\n"); + CONSOLE_PRINT("\n"); for (uint test = 0; test < numTests3D; test++) { hipPerfMemset.run3D(test, pattern.memsetval, hipMemsetTypeDefault, async); } @@ -431,6 +413,6 @@ TEST_CASE("Perf_hipPerfMemset_test") { } /** -* End doxygen group perfMemoryTest. -* @} -*/ + * End doxygen group perfMemoryTest. + * @} + */ diff --git a/catch/perftests/memory/hipPerfSampleRate.cc b/catch/perftests/memory/hipPerfSampleRate.cc index 9a083fa3ec..a591f5a091 100644 --- a/catch/perftests/memory/hipPerfSampleRate.cc +++ b/catch/perftests/memory/hipPerfSampleRate.cc @@ -19,66 +19,69 @@ /** -* @addtogroup hipMemcpyKernel hipMemcpyKernel -* @{ -* @ingroup perfMemoryTest -* `hipMemcpy(void* dst, const void* src, size_t count, hipMemcpyKind kind)` - -* Copies data between host and device. -*/ + * @addtogroup hipMemcpyKernel hipMemcpyKernel + * @{ + * @ingroup perfMemoryTest + * `hipMemcpy(void* dst, const void* src, size_t count, hipMemcpyKind kind)` - + * Copies data between host and device. + */ #include - +// #define ENABLE_DEBUG 1 #define NUM_TYPES 3 std::vector types = {"float", "float2", "float4"}; std::vector typeSizes = {4, 8, 16}; #define NUM_SIZES 12 -std::vector sizes = {1, 2, 4, 8, 16, 32, - 64, 128, 256, 512, 1024, 2048}; +std::vector sizes = {1, 2, 4, 8, 16, 32, 64, 128, 256, 512, 1024, 2048}; #define NUM_BUFS 6 #define MAX_BUFS (1 << (NUM_BUFS - 1)) #ifdef __HIP_PLATFORM_NVIDIA__ -__host__ __device__ void operator+=(float2 &a, float2 b) { //NOLINT - a.x += b.x; a.y += b.y; +__host__ __device__ void operator+=(float2& a, float2 b) { // NOLINT + a.x += b.x; + a.y += b.y; } -__host__ __device__ void operator+=(float4 &a, float4 b) { //NOLINT - a.x += b.x; a.y += b.y; a.z += b.z; a.w += b.w; +__host__ __device__ void operator+=(float4& a, float4 b) { // NOLINT + a.x += b.x; + a.y += b.y; + a.z += b.z; + a.w += b.w; } #endif template -__global__ void sampleRate(T * outBuffer, unsigned int inBufSize, - unsigned int writeIt, T **inBuffer, int numBufs) { +__global__ void sampleRate(T* outBuffer, unsigned int inBufSize, unsigned int writeIt, T** inBuffer, + int numBufs) { uint gid = (blockIdx.x * blockDim.x + threadIdx.x); uint inputIdx = gid % inBufSize; T tmp; memset(&tmp, 0, sizeof(T)); for (int i = 0; i < numBufs; i++) { - tmp += *(*(inBuffer+i)+inputIdx); + tmp += *(*(inBuffer + i) + inputIdx); } - if (writeIt*(unsigned int)tmp.x) { + if (writeIt * (unsigned int)tmp.x) { outBuffer[gid] = tmp; } } template -__global__ void sampleRateFloat(T * outBuffer, unsigned int inBufSize, - unsigned int writeIt, T ** inBuffer, int numBufs) { +__global__ void sampleRateFloat(T* outBuffer, unsigned int inBufSize, unsigned int writeIt, + T** inBuffer, int numBufs) { uint gid = (blockIdx.x * blockDim.x + threadIdx.x); uint inputIdx = gid % inBufSize; T tmp = (T)0.0f; for (int i = 0; i < numBufs; i++) { - tmp += *((*inBuffer+i)+inputIdx); + tmp += *((*inBuffer + i) + inputIdx); } - if (writeIt*(unsigned int)tmp) { + if (writeIt * (unsigned int)tmp) { outBuffer[gid] = tmp; } } @@ -93,26 +96,23 @@ class hipPerfSampleRate { void close(void); // array of funtion pointers - typedef void (hipPerfSampleRate::*funPtr)(void * outBuffer, unsigned int - inBufSize, unsigned int writeIt, void **inBuffer, int numBufs, - int grids, int blocks); + typedef void (hipPerfSampleRate::*funPtr)(void* outBuffer, unsigned int inBufSize, + unsigned int writeIt, void** inBuffer, int numBufs, + int grids, int blocks); // Wrappers - void float_kernel(void * outBuffer, unsigned int inBufSize, - unsigned int writeIt, void **inBuffer, int numBufs, - int grids, int blocks); + void float_kernel(void* outBuffer, unsigned int inBufSize, unsigned int writeIt, void** inBuffer, + int numBufs, int grids, int blocks); - void float2_kernel(void * outBuffer, unsigned int inBufSize, - unsigned int writeIt, void **inBuffer, int numBufs, - int grids, int blocks); + void float2_kernel(void* outBuffer, unsigned int inBufSize, unsigned int writeIt, void** inBuffer, + int numBufs, int grids, int blocks); - void float4_kernel(void * outBuffer, unsigned int inBufSize, - unsigned int writeIt, void **inBuffer, int numBufs, - int grids, int blocks); + void float4_kernel(void* outBuffer, unsigned int inBufSize, unsigned int writeIt, void** inBuffer, + int numBufs, int grids, int blocks); private: - void setData(void *ptr, unsigned int value); - void checkData(uint *ptr); + void setData(void* ptr, unsigned int value); + void checkData(uint* ptr); unsigned int width_; unsigned int bufSize_; @@ -139,41 +139,36 @@ bool hipPerfSampleRate::open(void) { hipDeviceProp_t props; HIP_CHECK(hipSetDevice(deviceId)); HIP_CHECK(hipGetDeviceProperties(&props, deviceId)); - INFO("info: running on bus " << "0x" << props.pciBusID << " " << - props.name << " with " << props.multiProcessorCount << - " CUs" << " and device id: " << deviceId << "\n"); + CONSOLE_PRINT("info: running on bus 0x%x %s with %d CUs and device id: %d\n", props.pciBusID, + props.name, props.multiProcessorCount, deviceId); numCUs = props.multiProcessorCount; return true; } // Wrappers for the kernel launches -void hipPerfSampleRate::float_kernel(void * outBuffer, unsigned int inBufSize, - unsigned int writeIt, void **inBuffer, int numBufs, - int grids, int blocks) { - hipLaunchKernelGGL(sampleRateFloat, dim3(grids, grids, grids), - dim3(blocks), 0, 0, reinterpret_cast(outBuffer), - inBufSize, writeIt, reinterpret_cast(inBuffer), numBufs); +void hipPerfSampleRate::float_kernel(void* outBuffer, unsigned int inBufSize, unsigned int writeIt, + void** inBuffer, int numBufs, int grids, int blocks) { + hipLaunchKernelGGL(sampleRateFloat, dim3(grids, grids, grids), dim3(blocks), 0, 0, + reinterpret_cast(outBuffer), inBufSize, writeIt, + reinterpret_cast(inBuffer), numBufs); } -void hipPerfSampleRate::float2_kernel(void * outBuffer, unsigned int inBufSize, - unsigned int writeIt, void **inBuffer, int grids, - int blocks, int numBufs) { - hipLaunchKernelGGL(sampleRate, dim3(grids, grids, grids), - dim3(blocks), 0, 0, reinterpret_cast(outBuffer), - inBufSize, writeIt, reinterpret_cast(inBuffer), numBufs); +void hipPerfSampleRate::float2_kernel(void* outBuffer, unsigned int inBufSize, unsigned int writeIt, + void** inBuffer, int grids, int blocks, int numBufs) { + hipLaunchKernelGGL(sampleRate, dim3(grids, grids, grids), dim3(blocks), 0, 0, + reinterpret_cast(outBuffer), inBufSize, writeIt, + reinterpret_cast(inBuffer), numBufs); } -void hipPerfSampleRate::float4_kernel(void * outBuffer, unsigned int inBufSize, - unsigned int writeIt, void **inBuffer, int grids, - int blocks, int numBufs) { - hipLaunchKernelGGL(sampleRate, dim3(grids, grids, grids), - dim3(blocks), 0, 0, reinterpret_cast(outBuffer), - inBufSize, writeIt, reinterpret_cast(inBuffer), numBufs); +void hipPerfSampleRate::float4_kernel(void* outBuffer, unsigned int inBufSize, unsigned int writeIt, + void** inBuffer, int grids, int blocks, int numBufs) { + hipLaunchKernelGGL(sampleRate, dim3(grids, grids, grids), dim3(blocks), 0, 0, + reinterpret_cast(outBuffer), inBufSize, writeIt, + reinterpret_cast(inBuffer), numBufs); } void hipPerfSampleRate::run(unsigned int test) { - funPtr p[] = {&hipPerfSampleRate::float_kernel, - &hipPerfSampleRate::float2_kernel, + funPtr p[] = {&hipPerfSampleRate::float_kernel, &hipPerfSampleRate::float2_kernel, &hipPerfSampleRate::float4_kernel}; // We compute a square domain @@ -182,35 +177,30 @@ void hipPerfSampleRate::run(unsigned int test) { bufSize_ = width_ * width_ * typeSizes[typeIdx_]; numBufs_ = (1 << (test / (NUM_SIZES * NUM_TYPES))); - void ** dPtr; - void * hOutPtr; - void * dOutPtr; - void ** hInPtr = new void *[numBufs_]; - void ** dInPtr = new void *[numBufs_]; + void** dPtr; + void* hOutPtr; + void* dOutPtr; + void** hInPtr = new void*[numBufs_]; + void** dInPtr = new void*[numBufs_]; - outBufSize_ = - sizes[NUM_SIZES - 1] * sizes[NUM_SIZES - 1] * typeSizes[NUM_TYPES - 1]; + outBufSize_ = sizes[NUM_SIZES - 1] * sizes[NUM_SIZES - 1] * typeSizes[NUM_TYPES - 1]; // Allocate memory on the host and device - HIP_CHECK(hipHostMalloc(reinterpret_cast(&hOutPtr), outBufSize_, - hipHostMallocDefault)); - setData(reinterpret_cast(hOutPtr), 0xdeadbeef); - HIP_CHECK(hipMalloc(reinterpret_cast(&dOutPtr), outBufSize_)); + HIP_CHECK(hipHostMalloc(reinterpret_cast(&hOutPtr), outBufSize_, hipHostMallocDefault)); + setData(reinterpret_cast(hOutPtr), 0xdeadbeef); + HIP_CHECK(hipMalloc(reinterpret_cast(&dOutPtr), outBufSize_)); // Allocate 2D array in Device - HIP_CHECK(hipMalloc(reinterpret_cast(&dPtr), - numBufs_* sizeof(void *))); + HIP_CHECK(hipMalloc(reinterpret_cast(&dPtr), numBufs_ * sizeof(void*))); for (uint i = 0; i < numBufs_; i++) { - HIP_CHECK(hipHostMalloc(reinterpret_cast(&hInPtr[i]), bufSize_, - hipHostMallocDefault)); - HIP_CHECK(hipMalloc(reinterpret_cast(&dInPtr[i]), bufSize_)); + HIP_CHECK(hipHostMalloc(reinterpret_cast(&hInPtr[i]), bufSize_, hipHostMallocDefault)); + HIP_CHECK(hipMalloc(reinterpret_cast(&dInPtr[i]), bufSize_)); setData(hInPtr[i], 0x3f800000); } // Populate array of pointers with array addresses - HIP_CHECK(hipMemcpy(dPtr, dInPtr, numBufs_* sizeof(void *), - hipMemcpyHostToDevice)); + HIP_CHECK(hipMemcpy(dPtr, dInPtr, numBufs_ * sizeof(void*), hipMemcpyHostToDevice)); // Copy memory from host to device for (uint i = 0; i < numBufs_; i++) { @@ -241,20 +231,19 @@ void hipPerfSampleRate::run(unsigned int test) { // Time the kernel execution auto all_start = std::chrono::steady_clock::now(); for (uint i = 0; i < maxIter; i++) { - (this->*p[idx]) (reinterpret_cast(dOutPtr), sizeDW, writeIt, - dPtr, numBufs_, grids, blocks); + (this->*p[idx])(reinterpret_cast(dOutPtr), sizeDW, writeIt, dPtr, numBufs_, grids, + blocks); } HIP_CHECK(hipDeviceSynchronize()); auto all_end = std::chrono::steady_clock::now(); std::chrono::duration all_kernel_time = all_end - all_start; - double perf = (static_cast(outBufSize_ * numBufs_ * - maxIter * (1e-09))) / all_kernel_time.count(); + double perf = + (static_cast(outBufSize_ * numBufs_ * maxIter * (1e-09))) / all_kernel_time.count(); - INFO("Domain " << sizes[NUM_SIZES - 1] << "x"<< sizes[NUM_SIZES - 1] - << " bufs " << numBufs_ << " " << types[typeIdx_] << " " << width_ - << "x" <(numBufs_)) { - INFO("Data validation failed at "<< i << " Got "<< ptr[i] - << ", expected " << (float)numBufs_ << "\n"); + DEBUG_PRINT("Data validation failed at %u Got %u, expected %f\n", i, ptr[i], (float)numBufs_); REQUIRE(false); } } } /** -* Test Description -* ------------------------ -*  - Verify hipPerfSampleRate status. -* Test source -* ------------------------ -*  - perftests/memory/hipPerfSampleRate.cc -* Test requirements -* ------------------------ -*  - HIP_VERSION >= 5.6 -*/ + * Test Description + * ------------------------ + *  - Verify hipPerfSampleRate status. + * Test source + * ------------------------ + *  - perftests/memory/hipPerfSampleRate.cc + * Test requirements + * ------------------------ + *  - HIP_VERSION >= 5.6 + */ TEST_CASE("Perf_hipPerfSampleRate_test") { hipPerfSampleRate sampleTypes; REQUIRE(true == sampleTypes.open()); - for (unsigned int testCase = 0; testCase < 216 ; testCase+=36) { + for (unsigned int testCase = 0; testCase < 216; testCase += 36) { sampleTypes.run(testCase); } } /** -* End doxygen group perfMemoryTest. -* @} -*/ + * End doxygen group perfMemoryTest. + * @} + */ diff --git a/catch/perftests/memory/hipPerfSharedMemReadSpeed.cc b/catch/perftests/memory/hipPerfSharedMemReadSpeed.cc index dbf10a04ae..67f2e59b60 100644 --- a/catch/perftests/memory/hipPerfSharedMemReadSpeed.cc +++ b/catch/perftests/memory/hipPerfSharedMemReadSpeed.cc @@ -18,19 +18,19 @@ */ /** -* @addtogroup hipMemcpyKernel hipMemcpyKernel -* @{ -* @ingroup perfMemoryTest -* `hipMemcpy(void* dst, const void* src, size_t count, hipMemcpyKind kind)` - -* Copies data between host and device. -*/ + * @addtogroup hipMemcpyKernel hipMemcpyKernel + * @{ + * @ingroup perfMemoryTest + * `hipMemcpy(void* dst, const void* src, size_t count, hipMemcpyKind kind)` - + * Copies data between host and device. + */ #include - +// #define ENABLE_DEBUG 1 #define sharedMemSize1 2048 #define sharedMemSize2 256 -__global__ void sharedMemReadSpeed1(float *outBuf, ulong N) { +__global__ void sharedMemReadSpeed1(float* outBuf, ulong N) { size_t gid = (blockIdx.x * blockDim.x + threadIdx.x); size_t lid = threadIdx.x; __shared__ float local[sharedMemSize1]; @@ -84,7 +84,7 @@ __global__ void sharedMemReadSpeed1(float *outBuf, ulong N) { } } -__global__ void sharedMemReadSpeed2(float *outBuf, ulong N) { +__global__ void sharedMemReadSpeed2(float* outBuf, ulong N) { size_t gid = (blockIdx.x * blockDim.x + threadIdx.x); size_t lid = threadIdx.x; __shared__ float local[sharedMemSize2]; @@ -116,8 +116,8 @@ __global__ void sharedMemReadSpeed2(float *outBuf, ulong N) { } static bool hipPerfSharedMemReadSpeed_test() { - float *dDst; - float *hDst; + float* dDst; + float* hDst; hipStream_t stream; constexpr uint numSizes = 4; constexpr uint Sizes[numSizes] = {262144, 1048576, 4194304, 16777216}; @@ -132,8 +132,8 @@ static bool hipPerfSharedMemReadSpeed_test() { HIP_CHECK(hipSetDevice(device)); hipDeviceProp_t props; HIP_CHECK(hipGetDeviceProperties(&props, device)); - INFO("info: running on bus " << "0x" << props.pciBusID << " " << props.name - << " with " << props.multiProcessorCount << " CUs \n"); + CONSOLE_PRINT("info: running on bus 0x%x %s with %d CUs\n", props.pciBusID, props.name, + props.multiProcessorCount); HIP_CHECK(hipStreamCreate(&stream)); @@ -149,8 +149,8 @@ static bool hipPerfSharedMemReadSpeed_test() { HIP_CHECK(hipMalloc(&dDst, nBytes)); HIP_CHECK(hipMemcpy(dDst, hDst, nBytes, hipMemcpyHostToDevice)); - hipLaunchKernelGGL(sharedMemReadSpeed1, dim3(blocks), - dim3(threadsPerBlock), 0, stream, dDst, N); + hipLaunchKernelGGL(sharedMemReadSpeed1, dim3(blocks), dim3(threadsPerBlock), 0, stream, dDst, + N); HIP_CHECK(hipMemcpy(hDst, dDst, nBytes, hipMemcpyDeviceToHost)); HIP_CHECK(hipDeviceSynchronize()); @@ -160,8 +160,7 @@ static bool hipPerfSharedMemReadSpeed_test() { tmp = 0; } if (hDst[i] != tmp) { - INFO("info: Data validation failed for warm up run! \n"); - INFO("info: expected " << tmp << " got " << hDst[i] << " \n"); + DEBUG_PRINT("Data validation failed for warm up run! expected %d got %f\n", tmp, hDst[i]); return false; } tmp += threadsPerBlock / 2; @@ -169,8 +168,8 @@ static bool hipPerfSharedMemReadSpeed_test() { auto all_start = std::chrono::steady_clock::now(); for (int i = 0; i < nIter; i++) { - hipLaunchKernelGGL(sharedMemReadSpeed1, dim3(blocks), - dim3(threadsPerBlock), 0, stream, dDst, N); + hipLaunchKernelGGL(sharedMemReadSpeed1, dim3(blocks), dim3(threadsPerBlock), 0, stream, dDst, + N); } HIP_CHECK(hipDeviceSynchronize()); @@ -178,15 +177,14 @@ static bool hipPerfSharedMemReadSpeed_test() { std::chrono::duration all_kernel_time = all_end - all_start; // read speed in GB/s - double perf = (static_cast(blocks * threadsPerBlock) - * (numReads1 * sizeof(float) + sharedMemSizeBytes1 / 64) - * nIter * (1e-09)) / all_kernel_time.count(); + double perf = (static_cast(blocks * threadsPerBlock) * + (numReads1 * sizeof(float) + sharedMemSizeBytes1 / 64) * nIter * (1e-09)) / + all_kernel_time.count(); - INFO("info: read speed = " << std::setw(8) << perf << " GB/s for " << - sharedMemSizeBytes1 / 1024 << " KB shared memory with " << - std::setw(8) << blocks * threadsPerBlock << " threads, " - << std::setw(4) << numReads1 << - " reads in sharedMemReadSpeed1 kernel \n"); + CONSOLE_PRINT( + "info: read speed = %.2f GB/s for %d KB shared memory with %d threads, %d reads in " + "sharedMemReadSpeed1 kernel\n", + perf, sharedMemSizeBytes1 / 1024, blocks * threadsPerBlock, numReads1); delete[] hDst; HIP_CHECK(hipFree(dDst)); @@ -204,15 +202,15 @@ static bool hipPerfSharedMemReadSpeed_test() { HIP_CHECK(hipMalloc(&dDst, nBytes)); HIP_CHECK(hipMemcpy(dDst, hDst, nBytes, hipMemcpyHostToDevice)); - hipLaunchKernelGGL(sharedMemReadSpeed2, dim3(blocks), - dim3(threadsPerBlock), 0, stream, dDst, N); + hipLaunchKernelGGL(sharedMemReadSpeed2, dim3(blocks), dim3(threadsPerBlock), 0, stream, dDst, + N); HIP_CHECK(hipMemcpy(hDst, dDst, nBytes, hipMemcpyDeviceToHost)); HIP_CHECK(hipDeviceSynchronize()); auto all_start = std::chrono::steady_clock::now(); for (int i = 0; i < nIter; i++) { - hipLaunchKernelGGL(sharedMemReadSpeed2, dim3(blocks), - dim3(threadsPerBlock), 0, stream, dDst, N); + hipLaunchKernelGGL(sharedMemReadSpeed2, dim3(blocks), dim3(threadsPerBlock), 0, stream, dDst, + N); } HIP_CHECK(hipDeviceSynchronize()); @@ -220,15 +218,14 @@ static bool hipPerfSharedMemReadSpeed_test() { std::chrono::duration all_kernel_time = all_end - all_start; // read speed in GB/s - double perf = (static_cast(blocks * threadsPerBlock) - * (numReads2 * sizeof(float) + sharedMemSizeBytes2 / 64) - * nIter * (1e-09)) / all_kernel_time.count(); + double perf = (static_cast(blocks * threadsPerBlock) * + (numReads2 * sizeof(float) + sharedMemSizeBytes2 / 64) * nIter * (1e-09)) / + all_kernel_time.count(); - INFO("info: read speed = " << std::setw(8) << perf << " GB/s for " - << sharedMemSizeBytes2 / 1024 << " KB shared memory with " - << std::setw(8) << blocks * threadsPerBlock << " threads, " - << std::setw(4) << numReads2 << - " reads in sharedMemReadSpeed2 kernel \n"); + CONSOLE_PRINT( + "info: read speed = %.2f GB/s for %d KB shared memory with %d threads, %d reads in " + "sharedMemReadSpeed2 kernel\n", + perf, sharedMemSizeBytes2 / 1024, blocks * threadsPerBlock, numReads2); delete[] hDst; HIP_CHECK(hipFree(dDst)); @@ -238,30 +235,31 @@ static bool hipPerfSharedMemReadSpeed_test() { } /** -* Test Description -* ------------------------ -*  - Verify hipPerfSharedMemReadSpeed status. -* Test source -* ------------------------ -*  - perftests/memory/hipPerfSharedMemReadSpeed.cc -* Test requirements -* ------------------------ -*  - HIP_VERSION >= 5.6 -*/ + * Test Description + * ------------------------ + *  - Verify hipPerfSharedMemReadSpeed status. + * Test source + * ------------------------ + *  - perftests/memory/hipPerfSharedMemReadSpeed.cc + * Test requirements + * ------------------------ + *  - HIP_VERSION >= 5.6 + */ TEST_CASE("Perf_hipPerfSharedMemReadSpeed_test") { int numDevices = 0; HIP_CHECK(hipGetDeviceCount(&numDevices)); if (numDevices <= 0) { - SUCCEED("Skipped testcase hipPerfSharedMemReadSpeed as" - "there is no device to test.\n"); + SUCCEED( + "Skipped testcase hipPerfSharedMemReadSpeed as" + "there is no device to test.\n"); } else { REQUIRE(true == hipPerfSharedMemReadSpeed_test()); } } /** -* End doxygen group perfMemoryTest. -* @} -*/ + * End doxygen group perfMemoryTest. + * @} + */ diff --git a/perftests/memory/hipPerfBufferCopyRectSpeed.cpp b/perftests/memory/hipPerfBufferCopyRectSpeed.cpp deleted file mode 100644 index 78096844f5..0000000000 --- a/perftests/memory/hipPerfBufferCopyRectSpeed.cpp +++ /dev/null @@ -1,300 +0,0 @@ -/* -Copyright (c) 2015 - 2021 Advanced Micro Devices, Inc. All rights reserved. -Permission is hereby granted, free of charge, to any person obtaining a copy -of this software and associated documentation files (the "Software"), to deal -in the Software without restriction, including without limitation the rights -to use, copy, modify, merge, publish, distribute, sublicense, and/or sell -copies of the Software, and to permit persons to whom the Software is -furnished to do so, subject to the following conditions: -The above copyright notice and this permission notice shall be included in -all copies or substantial portions of the Software. -THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR -IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, -FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE -AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER -LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, -OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN -THE SOFTWARE. -*/ - -/* HIT_START - * BUILD: %t %s ../../src/test_common.cpp ../../src/timer.cpp - * TEST: %t - * HIT_END - */ - -#include -#include -#include -#include - -#include "timer.h" -#include "test_common.h" - -// Quiet pesky warnings -#ifdef WIN_OS -#define SNPRINTF sprintf_s -#else -#define SNPRINTF snprintf -#endif - -#define NUM_SIZES 8 -//4KB, 8KB, 64KB, 256KB, 1 MB, 4MB, 16 MB, 16MB+10 -static const unsigned int Sizes[NUM_SIZES] = {4096, 8192, 65536, 262144, 1048576, 4194304, 16777216, 16777216+10}; - -static const unsigned int Iterations[2] = {1, 1000}; - -#define BUF_TYPES 4 -// 16 ways to combine 4 different buffer types -#define NUM_SUBTESTS (BUF_TYPES*BUF_TYPES) - -#define CHECK_RESULT(test, msg) \ - if ((test)) \ - { \ - printf("\n%s\n", msg); \ - abort(); \ - } - -void setData(void *ptr, unsigned int size, char value) -{ - char *ptr2 = (char *)ptr; - for (unsigned int i = 0; i < size ; i++) - { - ptr2[i] = value; - } -} - -void checkData(void *ptr, unsigned int size, char value) -{ - char *ptr2 = (char *)ptr; - for (unsigned int i = 0; i < size; i++) - { - if (ptr2[i] != value) - { - printf("Data validation failed at %d! Got 0x%08x\n", i, ptr2[i]); - printf("Expected 0x%08x\n", value); - CHECK_RESULT(true, "Data validation failed!"); - break; - } - } -} - - -int main(int argc, char* argv[]) { - HipTest::parseStandardArguments(argc, argv, true); - - hipError_t err = hipSuccess; - hipDeviceProp_t props = {0}; - hipGetDeviceProperties(&props, p_gpuDevice); - CHECK_RESULT(err != hipSuccess, "hipGetDeviceProperties failed" ); - printf("Set device to %d : %s\n", p_gpuDevice, props.name); - printf("Legend: unp - unpinned(malloc), hM - hipMalloc(device)\n"); - printf(" hHR - hipHostRegister(pinned), hHM - hipHostMalloc(prePinned)\n"); - err = hipSetDevice(p_gpuDevice); - CHECK_RESULT(err != hipSuccess, "hipSetDevice failed" ); - - unsigned int bufSize_; - bool hostMalloc[2] = {false}; - bool hostRegister[2] = {false}; - bool unpinnedMalloc[2] = {false}; - unsigned int numIter; - void *memptr[2] = {NULL}; - void *alignedmemptr[2] = {NULL}; - void* srcBuffer = NULL; - void* dstBuffer = NULL; - - int numTests = (p_tests == -1) ? (NUM_SIZES*NUM_SUBTESTS*2 - 1) : p_tests; - int test = (p_tests == -1) ? 0 : p_tests; - - for(;test <= numTests; test++) - { - unsigned int srcTest = (test / NUM_SIZES) % BUF_TYPES; - unsigned int dstTest = (test / (NUM_SIZES*BUF_TYPES)) % BUF_TYPES; - bufSize_ = Sizes[test % NUM_SIZES]; - hostMalloc[0] = hostMalloc[1] = false; - hostRegister[0] = hostRegister[1] = false; - unpinnedMalloc[0] = unpinnedMalloc[1] = false; - srcBuffer = dstBuffer = 0; - memptr[0] = memptr[1] = NULL; - alignedmemptr[0] = alignedmemptr[1] = NULL; - - size_t width = static_cast(sqrt(static_cast(bufSize_))); - - if (srcTest == 3) - { - hostRegister[0] = true; - } - else if (srcTest == 2) - { - hostMalloc[0] = true; - } - else if (srcTest == 1) - { - unpinnedMalloc[0] = true; - } - - if (dstTest == 1) - { - unpinnedMalloc[1] = true; - } - else if (dstTest == 2) - { - hostMalloc[1] = true; - } - else if (dstTest == 3) - { - hostRegister[1] = true; - } - - numIter = Iterations[test / (NUM_SIZES * NUM_SUBTESTS)]; - - if (hostMalloc[0]) - { - err = hipHostMalloc((void**)&srcBuffer, bufSize_, 0); - setData(srcBuffer, bufSize_, 0xd0); - CHECK_RESULT(err != hipSuccess, "hipHostMalloc failed"); - } - else if (hostRegister[0]) - { - memptr[0] = malloc(bufSize_ + 4096); - alignedmemptr[0] = (void*)(((size_t)memptr[0] + 4095) & ~4095); - srcBuffer = alignedmemptr[0]; - setData(srcBuffer, bufSize_, 0xd0); - err = hipHostRegister(srcBuffer, bufSize_, 0); - CHECK_RESULT(err != hipSuccess, "hipHostRegister failed"); - } - else if (unpinnedMalloc[0]) - { - memptr[0] = malloc(bufSize_ + 4096); - alignedmemptr[0] = (void*)(((size_t)memptr[0] + 4095) & ~4095); - srcBuffer = alignedmemptr[0]; - setData(srcBuffer, bufSize_, 0xd0); - } - else - { - err = hipMalloc(&srcBuffer, bufSize_); - CHECK_RESULT(err != hipSuccess, "hipMalloc failed"); - err = hipMemset(srcBuffer, 0xd0, bufSize_); - CHECK_RESULT(err != hipSuccess, "hipMemset failed"); - } - - if (hostMalloc[1]) - { - err = hipHostMalloc((void**)&dstBuffer, bufSize_, 0); - CHECK_RESULT(err != hipSuccess, "hipHostMalloc failed"); - } - else if (hostRegister[1]) - { - memptr[1] = malloc(bufSize_ + 4096); - alignedmemptr[1] = (void*)(((size_t)memptr[1] + 4095) & ~4095); - dstBuffer = alignedmemptr[1]; - err = hipHostRegister(dstBuffer, bufSize_, 0); - CHECK_RESULT(err != hipSuccess, "hipHostRegister failed"); - } - else if (unpinnedMalloc[1]) - { - memptr[1] = malloc(bufSize_ + 4096); - alignedmemptr[1] = (void*)(((size_t)memptr[1] + 4095) & ~4095); - dstBuffer = alignedmemptr[1]; - } - else - { - err = hipMalloc(&dstBuffer, bufSize_); - CHECK_RESULT(err != hipSuccess, "hipMalloc failed"); - } - - CPerfCounter timer; - - //warm up - err = hipMemcpy2D(dstBuffer, width, srcBuffer, width, width, width, hipMemcpyDefault); - CHECK_RESULT(err, "hipMemcpy2D failed"); - - timer.Reset(); - timer.Start(); - for (unsigned int i = 0; i < numIter; i++) - { - err = hipMemcpy2DAsync(dstBuffer, width, srcBuffer, width, width, width, hipMemcpyDefault, NULL); - CHECK_RESULT(err, "hipMemcpyAsync2D failed"); - } - err = hipDeviceSynchronize(); - CHECK_RESULT(err, "hipDeviceSynchronize failed"); - timer.Stop(); - double sec = timer.GetElapsedTime(); - - // Buffer copy bandwidth in GB/s - double perf = ((double)bufSize_*numIter*(double)(1e-09)) / sec; - - const char *strSrc = NULL; - const char *strDst = NULL; - if (hostMalloc[0]) - strSrc = "hHM"; - else if (hostRegister[0]) - strSrc = "hHR"; - else if (unpinnedMalloc[0]) - strSrc = "unp"; - else - strSrc = "hM"; - - if (hostMalloc[1]) - strDst = "hHM"; - else if (hostRegister[1]) - strDst = "hHR"; - else if (unpinnedMalloc[1]) - strDst = "unp"; - else - strDst = "hM"; - // Double results when src and dst are both on device - if ((!hostMalloc[0] && !hostRegister[0] && !unpinnedMalloc[0]) && - (!hostMalloc[1] && !hostRegister[1] && !unpinnedMalloc[1])) - perf *= 2.0; - // Double results when src and dst are both in sysmem - if ((hostMalloc[0] || hostRegister[0] || unpinnedMalloc[0]) && - (hostMalloc[1] || hostRegister[1] || unpinnedMalloc[1])) - perf *= 2.0; - - char buf[256]; - SNPRINTF(buf, sizeof(buf), "HIPPerfBufferCopyRectSpeed[%d]\t(%8d bytes)\ts:%s d:%s\ti:%4d\t(GB/s) perf\t%f", - test, bufSize_, strSrc, strDst, numIter, (float)perf); - printf("%s\n", buf); - - //Free src - if (hostMalloc[0]) - { - hipHostFree(srcBuffer); - } - else if (hostRegister[0]) - { - hipHostUnregister(srcBuffer); - free(memptr[0]); - } - else if (unpinnedMalloc[0]) - { - free(memptr[0]); - } - else - { - hipFree(srcBuffer); - } - - //Free dst - if (hostMalloc[1]) - { - hipHostFree(dstBuffer); - } - else if (hostRegister[1]) - { - hipHostUnregister(dstBuffer); - free(memptr[1]); - } - else if (unpinnedMalloc[1]) - { - free(memptr[1]); - } - else - { - hipFree(dstBuffer); - } - } - - passed(); -} diff --git a/perftests/memory/hipPerfDevMemReadSpeed.cpp b/perftests/memory/hipPerfDevMemReadSpeed.cpp deleted file mode 100644 index 6548da94cd..0000000000 --- a/perftests/memory/hipPerfDevMemReadSpeed.cpp +++ /dev/null @@ -1,136 +0,0 @@ -/* -Copyright (c) 2015 - 2021 Advanced Micro Devices, Inc. All rights reserved. -Permission is hereby granted, free of charge, to any person obtaining a copy -of this software and associated documentation files (the "Software"), to deal -in the Software without restriction, including without limitation the rights -to use, copy, modify, merge, publish, distribute, sublicense, and/or sell -copies of the Software, and to permit persons to whom the Software is -furnished to do so, subject to the following conditions: -The above copyright notice and this permission notice shall be included in -all copies or substantial portions of the Software. -THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR -IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, -FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE -AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER -LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, -OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN -THE SOFTWARE. -*/ - -/* HIT_START - * BUILD: %t %s ../../src/test_common.cpp - * TEST: %t - * HIT_END - */ - -#include -#include -#include "test_common.h" - -using namespace std; - -#define arraySize 16 - -typedef struct d_uint16 { - uint data[arraySize]; -} d_uint16; - -__global__ void read_kernel(d_uint16 *src, ulong N, uint *dst) { - - size_t idx = (blockIdx.x * blockDim.x + threadIdx.x); - size_t stride = blockDim.x * gridDim.x ; - - uint tmp = 0; - for (size_t i = idx; i < N; i += stride) { - for (size_t j = 0; j < arraySize; j++) { - tmp += src[i].data[j]; - } - } - - atomicAdd(dst, tmp); -} - -int main(int argc, char* argv[]) { - d_uint16 *dSrc; - d_uint16 *hSrc; - uint *dDst; - uint *hDst; - hipStream_t stream; - ulong N = 4 * 1024 * 1024; - uint nBytes = N * sizeof(d_uint16); - - int nGpu = 0; - HIPCHECK(hipGetDeviceCount(&nGpu)); - if (nGpu < 1) { - cout << "info: didn't find any GPU! skipping the test!\n"; - passed(); - return 0; - } - - static int device = 0; - HIPCHECK(hipSetDevice(device)); - hipDeviceProp_t props; - HIPCHECK(hipGetDeviceProperties(&props, device)); - cout << "info: running on bus " << "0x" << props.pciBusID << " " << props.name << - " with " << props.multiProcessorCount << " CUs" << endl; - - const unsigned threadsPerBlock = 64; - const unsigned blocks = props.multiProcessorCount * 4; - - uint inputData = 0x1; - int nIter = 1000; - - hSrc = new d_uint16[nBytes]; - HIPCHECK(hSrc == 0 ? hipErrorOutOfMemory : hipSuccess); - hDst = new uint; - hDst[0] = 0; - HIPCHECK(hDst == 0 ? hipErrorOutOfMemory : hipSuccess); - for (size_t i = 0; i < N; i++) { - for (int j = 0; j < arraySize; j++) { - hSrc[i].data[j] = inputData; - } - } - - HIPCHECK(hipMalloc(&dSrc, nBytes)); - HIPCHECK(hipMalloc(&dDst, sizeof(uint))); - - HIPCHECK(hipStreamCreate(&stream)); - - HIPCHECK(hipMemcpy(dSrc, hSrc, nBytes, hipMemcpyHostToDevice)); - HIPCHECK(hipMemcpy(dDst, hDst, sizeof(uint), hipMemcpyHostToDevice)); - - hipLaunchKernelGGL(read_kernel, dim3(blocks), dim3(threadsPerBlock), 0, stream, dSrc, N, dDst); - HIPCHECK(hipMemcpy(hDst, dDst, sizeof(uint), hipMemcpyDeviceToHost)); - hipDeviceSynchronize(); - - if (hDst[0] != (nBytes / sizeof(uint))) { - cout << "info: Data validation failed for warm up run!" << endl; - cout << "info: expected " << nBytes / sizeof(uint) << " got " << hDst[0] << endl; - HIPCHECK(hipErrorUnknown); - } - - // measure performance based on host time - auto all_start = chrono::steady_clock::now(); - - for(int i = 0; i < nIter; i++) { - hipLaunchKernelGGL(read_kernel, dim3(blocks), dim3(threadsPerBlock), 0, stream, dSrc, N, dDst); - } - hipDeviceSynchronize(); - - auto all_end = chrono::steady_clock::now(); - chrono::duration all_kernel_time = all_end - all_start; - - // read speed in GB/s - double perf = ((double)nBytes * nIter * (double)(1e-09)) / all_kernel_time.count(); - - cout << "info: average read speed of " << perf << " GB/s " << "achieved for memory size of " << - nBytes / (1024 * 1024) << " MB" << endl; - - delete [] hSrc; - delete hDst; - hipFree(dSrc); - hipFree(dDst); - HIPCHECK(hipStreamDestroy(stream)); - - passed(); -} diff --git a/perftests/memory/hipPerfDevMemWriteSpeed.cpp b/perftests/memory/hipPerfDevMemWriteSpeed.cpp deleted file mode 100644 index cc48836604..0000000000 --- a/perftests/memory/hipPerfDevMemWriteSpeed.cpp +++ /dev/null @@ -1,126 +0,0 @@ -/* -Copyright (c) 2015 - 2021 Advanced Micro Devices, Inc. All rights reserved. -Permission is hereby granted, free of charge, to any person obtaining a copy -of this software and associated documentation files (the "Software"), to deal -in the Software without restriction, including without limitation the rights -to use, copy, modify, merge, publish, distribute, sublicense, and/or sell -copies of the Software, and to permit persons to whom the Software is -furnished to do so, subject to the following conditions: -The above copyright notice and this permission notice shall be included in -all copies or substantial portions of the Software. -THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR -IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, -FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE -AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER -LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, -OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN -THE SOFTWARE. -*/ - -/* HIT_START - * BUILD: %t %s ../../src/test_common.cpp - * TEST: %t - * HIT_END - */ - -#include -#include -#include "test_common.h" - -using namespace std; - -#define arraySize 16 - -typedef struct d_uint16 { - uint data[arraySize]; -} d_uint16; - -__global__ void write_kernel(d_uint16 *dst, ulong N, d_uint16 pval) { - size_t idx = (blockIdx.x * blockDim.x + threadIdx.x); - size_t stride = blockDim.x * gridDim.x; - for (size_t i = idx; i < N; i += stride) { - dst[i] = pval; - } -}; - -int main(int argc, char* argv[]) { - d_uint16 *dDst; - d_uint16 *hDst; - hipStream_t stream; - ulong N = 4 * 1024 * 1024; - uint nBytes = N * sizeof(d_uint16); - d_uint16 pval; - - for (int i = 0; i < arraySize; i++) { - pval.data[i] = 0xabababab; - } - - int nGpu = 0; - HIPCHECK(hipGetDeviceCount(&nGpu)); - if (nGpu < 1) { - cout << "info: didn't find any GPU! skipping the test!\n"; - passed(); - return 0; - } - - static int device = 0; - HIPCHECK(hipSetDevice(device)); - hipDeviceProp_t props; - HIPCHECK(hipGetDeviceProperties(&props, device)); - cout << "info: running on bus " << "0x" << props.pciBusID << " " << props.name << - " with " << props.multiProcessorCount << " CUs" << endl; - - size_t threadsPerBlock = 64; - size_t blocks = props.multiProcessorCount * 4; - - uint inputData = 0xabababab; - int nIter = 1000; - - hDst = new d_uint16[nBytes]; - HIPCHECK(hDst == 0 ? hipErrorOutOfMemory : hipSuccess); - for (size_t i = 0; i < N; i++) { - for (size_t j = 0; j < arraySize; j++) { - hDst[i].data[j] = 0; - } - } - - HIPCHECK(hipMalloc(&dDst, nBytes)); - - HIPCHECK(hipStreamCreate(&stream)); - - hipLaunchKernelGGL(write_kernel, dim3(blocks), dim3(threadsPerBlock), 0, stream, dDst, N, pval); - HIPCHECK(hipMemcpy(hDst, dDst, nBytes , hipMemcpyDeviceToHost)); - hipDeviceSynchronize(); - - for (uint i = 0; i < N; i++) { - for (uint j = 0; j < arraySize; j++) { - if (hDst[i].data[j] != inputData) { - cout << "info: Data validation failed for warm up run! " << endl; - cout << "at index i: " << i << " element j: " << j << endl; - cout << hex << "expected 0x" << inputData << " but got 0x" << hDst[i].data[j] << endl; - HIPCHECK(hipErrorUnknown); - } - } - } - - auto all_start = chrono::steady_clock::now(); - for(int i = 0; i < nIter; i++) { - hipLaunchKernelGGL(write_kernel, dim3(blocks), dim3(threadsPerBlock), 0, stream, dDst, N, pval); - } - hipDeviceSynchronize(); - auto all_end = chrono::steady_clock::now(); - chrono::duration all_kernel_time = all_end - all_start; - - // read speed in GB/s - double perf = ((double)nBytes * nIter * (double)(1e-09)) / all_kernel_time.count(); - - cout << "info: average write speed of " << perf << " GB/s " << "achieved for memory size of " << - nBytes / (1024 * 1024) << " MB" << endl; - - - delete [] hDst; - hipFree(dDst); - HIPCHECK(hipStreamDestroy(stream)); - - passed(); -} diff --git a/perftests/memory/hipPerfHostNumaAlloc.cpp b/perftests/memory/hipPerfHostNumaAlloc.cpp deleted file mode 100644 index 93cd71ace7..0000000000 --- a/perftests/memory/hipPerfHostNumaAlloc.cpp +++ /dev/null @@ -1,190 +0,0 @@ -/* -Copyright (c) 2015 - 2021 Advanced Micro Devices, Inc. All rights reserved. -Permission is hereby granted, free of charge, to any person obtaining a copy -of this software and associated documentation files (the "Software"), to deal -in the Software without restriction, including without limitation the rights -to use, copy, modify, merge, publish, distribute, sublicense, and/or sell -copies of the Software, and to permit persons to whom the Software is -furnished to do so, subject to the following conditions: -The above copyright notice and this permission notice shall be included in -all copies or substantial portions of the Software. -THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR -IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, -FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE -AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER -LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, -OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN -THE SOFTWARE. -*/ - -/* HIT_START - * BUILD_CMD: hipPerfHostNumaAlloc %hc -I%S/../../src %S/%s %S/../../src/test_common.cpp -lnuma -o %T/%t EXCLUDE_HIP_PLATFORM nvidia - * TEST: %t - * HIT_END - */ - -#include "test_common.h" -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include "hip/hip_runtime.h" - -// To run it correctly, we must not export HIP_VISIBLE_DEVICES. -// And we must explicitly link libnuma because of numa api move_pages(). -#define NUM_PAGES 4 -char *h = nullptr; -char *d_h = nullptr; -char *m = nullptr; -char *d_m = nullptr; -int page_size = 0; -const int mode[] = { MPOL_DEFAULT, MPOL_BIND, MPOL_PREFERRED, MPOL_INTERLEAVE }; -const char* modeStr[] = { "MPOL_DEFAULT", "MPOL_BIND", "MPOL_PREFERRED", "MPOL_INTERLEAVE" }; - -std::string exeCommand(const char* cmd) { - std::array buff; - std::string result; - std::unique_ptr pipe(popen(cmd, "r"), pclose); - if (!pipe) { - return result; - } - while (fgets(buff.data(), buff.size(), pipe.get()) != nullptr) { - result += buff.data(); - } - return result; -} - -int getCpuAgentCount() { - const char* cmd = "cat /proc/cpuinfo | grep \"physical id\" | sort | uniq | wc -l"; - int cpuAgentCount = std::atoi(exeCommand(cmd).c_str()); - return cpuAgentCount; -} - -bool test(int cpuId, int gpuId, int numaMode, unsigned int hostMallocflags) { - void *pages[NUM_PAGES]; - int status[NUM_PAGES]; - int nodes[NUM_PAGES]; - int ret_code; - - printf("set cpu %d, gpu %d, numaMode %d, hostMallocflags 0x%x\n", cpuId, - gpuId, numaMode, hostMallocflags); - - if (cpuId >= 0) { - unsigned long nodeMask = 1 << cpuId; - unsigned long maxNode = sizeof(nodeMask) * 8; - if (set_mempolicy(numaMode, numaMode == MPOL_DEFAULT ? NULL : &nodeMask, - numaMode == MPOL_DEFAULT ? 0 : maxNode) == -1) { - printf("set_mempolicy() failed with err %d\n", errno); - return false; - } - } - - if (gpuId >= 0) { - HIPCHECK(hipSetDevice(gpuId)); - } - - posix_memalign((void**) &m, page_size, page_size * NUM_PAGES); - hipHostRegister(m, page_size * NUM_PAGES, hipHostRegisterMapped); - hipHostGetDevicePointer((void**) &d_m, m, 0); - - status[0] = -1; - pages[0] = m; - for (int i = 1; i < NUM_PAGES; i++) { - pages[i] = (char*) pages[0] + page_size; - } - ret_code = move_pages(0, NUM_PAGES, pages, NULL, status, 0); - printf("Memory (malloc) ret %d at %p (dev %p) is at node: ", ret_code, m, d_m); - for (int i = 0; i < NUM_PAGES; i++) { - printf("%d ", status[i]); // Don't verify as it's out of our control - } - printf("\n"); - - HIPCHECK(hipHostMalloc((void**) &h, page_size*NUM_PAGES, hostMallocflags)); - pages[0] = h; - for (int i = 1; i < NUM_PAGES; i++) { - pages[i] = (char*) pages[0] + page_size; - } - ret_code = move_pages(0, NUM_PAGES, pages, NULL, status, 0); - d_h = nullptr; - if (hostMallocflags & hipHostMallocMapped) { - hipHostGetDevicePointer((void**) &d_h, h, 0); - printf("Memory (hipHostMalloc) ret %d at %p (dev %p) is at node: ", - ret_code, h, d_h); - } else { - printf("Memory (hipHostMalloc) ret %d at %p is at node: ", ret_code, h); - } - for (int i = 0; i < NUM_PAGES; i++) { - printf("%d ", status[i]); // Always print it even if it's wrong. Verify later - } - printf("\n"); - - HIPCHECK(hipHostFree((void* )h)); - hipHostUnregister(m); - free(m); - - if (cpuId >= 0 && (numaMode == MPOL_BIND || numaMode == MPOL_PREFERRED)) { - for (int i = 0; i < NUM_PAGES; i++) { - if (status[i] != cpuId) { // Now verify - printf("Failed at %d", i); - return false; - } - } - } - return true; -} - -bool runTest(const int &cpuCount, const int &gpuCount, - const unsigned int &hostMallocflags, const std::string &str) { - printf("%s\n", str.c_str()); - - for (int m = 0; m < sizeof(mode) / sizeof(mode[0]); m++) { - printf("Testing %s\n", modeStr[m]); - - for (int i = 0; i < cpuCount; i++) { - for (int j = 0; j < gpuCount; j++) { - if (!test(i, j, mode[m], hostMallocflags)) { - return false; - } - } - } - } - return true; -} - -int main(int argc, char *argv[]) { - int gpuCount = 0; - HIPCHECK(hipGetDeviceCount(&gpuCount)); - int cpuCount = getCpuAgentCount(); - page_size = getpagesize(); - printf("Cpu count %d, Gpu count %d, Page size %d\n", cpuCount, gpuCount, - page_size); - - if (cpuCount < 0 || gpuCount < 0) { - failed("Bad device count\n"); - return -1; - } - - if (!runTest(cpuCount, gpuCount, hipHostMallocDefault | hipHostMallocNumaUser, - "Testing hipHostMallocDefault | hipHostMallocNumaUser........................")) { - failed("Failed testing hipHostMallocDefault | hipHostMallocNumaUser\n"); - return -1; - } - - if (!runTest(cpuCount, gpuCount, hipHostMallocMapped | hipHostMallocNumaUser, - "Testing hipHostMallocMapped | hipHostMallocNumaUser.........................")) { - failed("Failed testing hipHostMallocMapped | hipHostMallocNumaUser\n"); - return -1; - } - - passed(); -} diff --git a/perftests/memory/hipPerfMemFill.cpp b/perftests/memory/hipPerfMemFill.cpp deleted file mode 100644 index dd54ec6855..0000000000 --- a/perftests/memory/hipPerfMemFill.cpp +++ /dev/null @@ -1,534 +0,0 @@ -/* - Copyright (c) 2015 - 2021 Advanced Micro Devices, Inc. All rights reserved. - Permission is hereby granted, free of charge, to any person obtaining a copy - of this software and associated documentation files (the "Software"), to deal - in the Software without restriction, including without limitation the rights - to use, copy, modify, merge, publish, distribute, sublicense, and/or sell - copies of the Software, and to permit persons to whom the Software is - furnished to do so, subject to the following conditions: - The above copyright notice and this permission notice shall be included in - all copies or substantial portions of the Software. - THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR - IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, - FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE - AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER - LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, - OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN - THE SOFTWARE. - */ - -/* HIT_START - * BUILD: %t %s ../../src/test_common.cpp - * TEST: %t - * HIT_END - */ - -#include "test_common.h" -#include -#include -#include -#include - -#define SIMPLY_ASSIGN 0 -#define USE_HIPTEST_SETNUMBLOCKS 0 - -using namespace std; - -template -__global__ void vec_fill(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 - x[i] = coef; -#else - x[i] = coef * i; -#endif - } -} - -__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) { - const int istart = threadIdx.x + blockIdx.x * blockDim.x; - const int ishift = blockDim.x * gridDim.x; - for (int iter = 0; iter < Niter; ++iter) { - T iv = coef * iter; - for (int i = istart; i < N; i += ishift) - y[i] = iv * x[i] + y[i]; - } -} - -template -class hipPerfMemFill { - private: - static constexpr int NUM_START = 27; - static constexpr int NUM_SIZE = 5; - static constexpr int NUM_ITER = 10; - size_t totalSizes_[NUM_SIZE]; - hipDeviceProp_t props_; - const T coef_ = getCoefficient(3.14159); - const unsigned int blocksPerCU_; - const unsigned int threadsPerBlock_; - - public: - hipPerfMemFill(unsigned int blocksPerCU, unsigned int threadsPerBlock) : - blocksPerCU_(blocksPerCU), threadsPerBlock_(threadsPerBlock) { - for (int i = 0; i < NUM_SIZE; i++) { - totalSizes_[i] = 1ull << (i + NUM_START); // 128M, 256M, 512M, 1024M, 2048M - } - } - - ~hipPerfMemFill() { - } - - bool supportLargeBar() { - return props_.isLargeBar != 0; - } - - bool supportManagedMemory() { - return props_.managedMemory != 0; - } - - const T getCoefficient(double val) { - return static_cast(val); - } - - void setHostBuffer(T *A, T val, size_t size) { - size_t len = size / sizeof(T); - for (int i = 0; i < len; i++) { - A[i] = val; - } - } - - void open(int deviceId) { - int nGpu = 0; - HIPCHECK(hipGetDeviceCount(&nGpu)); - if (nGpu < 1) { - failed("No GPU!"); - } else if (deviceId >= nGpu) { - failed("Info: wrong GPU Id %d\n", deviceId); - } - - HIPCHECK(hipSetDevice(deviceId)); - memset(&props_, 0, sizeof(props_)); - HIPCHECK(hipGetDeviceProperties(&props_, deviceId)); - std::cout << "Info: running on device: id: " << deviceId << ", bus: 0x" - << props_.pciBusID << " " << props_.name << " with " - << props_.multiProcessorCount << " CUs, large bar: " - << supportLargeBar() << ", managed memory: " << supportManagedMemory() - << ", DeviceMallocFinegrained: " << supportDeviceMallocFinegrained() - << 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(); - for (int i = 0; i < num; ++i) { -#if SIMPLY_ASSIGN - data[i] = coef; -#else - data[i] = coef * i; -#endif - } - auto end = chrono::steady_clock::now(); - chrono::duration diff = end - start; // in second - sec = diff.count(); - } - - void kernelFill(size_t size, T *data, T coef, double &sec) { - size_t num = size / sizeof(T); // Size of elements - unsigned blocks = setNumBlocks(num); - - hipLaunchKernelGGL(HIP_KERNEL_NAME(vec_fill), dim3(blocks), - dim3(threadsPerBlock), 0, 0, data, 0, num); // kernel will be loaded first time - HIPCHECK(hipDeviceSynchronize()); - - auto start = chrono::steady_clock::now(); - - for (int iter = 0; iter < NUM_ITER; ++iter) { - hipLaunchKernelGGL(HIP_KERNEL_NAME(vec_fill), 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 - } - - 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 << "hostVerify failed: i=" << i << ", data[i]=" << data[i] << ", expected=" << coef << endl; - failed("failed\n"); - } -#else - if(data[i] != coef * i) { - cout << "hostVerify failed: i=" << i << ", data[i]=" << data[i] << ", expected=" << coef * i << endl; - failed("failed\n"); - } -#endif - } - auto end = chrono::steady_clock::now(); - chrono::duration diff = end - start; // in second - sec = diff.count(); - } - - 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; - } - - double GBytes = (double) size / (1024.0 * 1024.0 * 1024.0); - - T *A; - HIPCHECK(hipMalloc(&A, size)); - double sec = 0; - hostFill(size, A, coef_, sec); // Cpu can access device mem in LB - HIPCHECK(hipFree(A)); - - log_host("Largebar: host fill", GBytes, sec); - return true; - } - - bool testLargeBar() { - if (!supportLargeBar()) { - return false; - } - - cout << "Test large bar device memory host filling" << endl; - for (int i = 0; i < NUM_SIZE; i++) { - if (!testLargeBarDeviceMemoryHostFill(totalSizes_[i])) { - return false; - } - } - - return true; - } - - bool testManagedMemoryHostFill(size_t size) { - if (!supportManagedMemory()) { - return false; - } - double GBytes = (double) size / (1024.0 * 1024.0 * 1024.0); - - T *A; - HIPCHECK(hipMallocManaged(&A, size)); - double sec = 0; - hostFill(size, A, coef_, sec); // Cpu can access HMM mem - HIPCHECK(hipFree(A)); - - log_host("Managed: host fill", GBytes, sec); - return true; - } - - bool testManagedMemoryKernelFill(size_t size) { - if (!supportManagedMemory()) { - return false; - } - double GBytes = (double) size / (1024.0 * 1024.0 * 1024.0); - - T *A; - HIPCHECK(hipMallocManaged(&A, size)); - - double sec = 0, sec_hv = 0, sec_kv = 0; - kernelFill(size, A, coef_, sec); - hostVerify(size, A, coef_, sec_hv); // Managed memory can be verified by host - kernelVerify(size, A, coef_, sec_kv); - HIPCHECK(hipFree(A)); - - log_kernel("Managed: kernel fill", GBytes, sec, sec_hv, sec_kv); - - return true; - } - - bool testManagedMemory() { - if (!supportManagedMemory()) { - return false; - } - - cout << "Test managed memory host filling" << endl; - for (int i = 0; i < NUM_SIZE; i++) { - if (!testManagedMemoryHostFill(totalSizes_[i])) { - return false; - } - } - - cout << "Test managed memory kernel filling" << endl; - for (int i = 0; i < NUM_SIZE; i++) { - if (!testManagedMemoryKernelFill(totalSizes_[i])) { - return false; - } - } - - return true; - } - - bool testHostMemoryHostFill(size_t size, unsigned int flags) { - double GBytes = (double) size / (1024.0 * 1024.0 * 1024.0); - T *A; - HIPCHECK(hipHostMalloc(&A, size, flags)); - double sec = 0; - hostFill(size, A, coef_, sec); - HIPCHECK(hipHostFree(A)); - - log_host("Host: host fill", GBytes, sec); - return true; - } - - bool testHostMemoryKernelFill(size_t size, unsigned int flags) { - double GBytes = (double) size / (1024.0 * 1024.0 * 1024.0); - - T *A; - HIPCHECK(hipHostMalloc((void** ) &A, size, flags)); - double sec = 0, sec_hv = 0, sec_kv = 0; - kernelFill(size, A, coef_, sec); - hostVerify(size, A, coef_, sec_hv); - kernelVerify(size, A, coef_, sec_kv); - HIPCHECK(hipHostFree(A)); - - log_kernel("Host: kernel fill", GBytes, sec, sec_hv, sec_kv); - return true; - } - - bool testHostMemory() { - cout << "Test coherent host memory host filling" << endl; - for (int i = 0; i < NUM_SIZE; i++) { - if (!testHostMemoryHostFill(totalSizes_[i], hipHostMallocCoherent)) { - return false; - } - } - - cout << "Test non-coherent host memory host filling" << endl; - for (int i = 0; i < NUM_SIZE; i++) { - if (!testHostMemoryHostFill(totalSizes_[i], hipHostMallocNonCoherent)) { - return false; - } - } - - cout << "Test coherent host memory kernel filling" << endl; - for (int i = 0; i < NUM_SIZE; i++) { - if (!testHostMemoryKernelFill(totalSizes_[i], hipHostMallocCoherent)) { - return false; - } - } - - cout << "Test non-coherent host memory kernel filling" << endl; - for (int i = 0; i < NUM_SIZE; i++) { - if (!testHostMemoryKernelFill(totalSizes_[i], hipHostMallocNonCoherent)) { - return false; - } - } - - return true; - } - - /* This function should be via device attribute query*/ - bool supportDeviceMallocFinegrained() { -#ifdef __HIP_PLATFORM_AMD__ - T *A = nullptr; - hipExtMallocWithFlags((void **)&A, sizeof(T), hipDeviceMallocFinegrained); - if (!A) { - return false; - } - HIPCHECK(hipFree(A)); - return true; -#else - return false; -#endif - } - - unsigned int setNumBlocks(size_t size) { - size_t num = size/sizeof(T); - -#if USE_HIPTEST_SETNUMBLOCKS - return HipTest::setNumBlocks(blocksPerCU_, threadsPerBlock_, - num); -#else - return (num + threadsPerBlock_ - 1) / threadsPerBlock_; -#endif - } - -#ifdef __HIP_PLATFORM_AMD__ - 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; - } - - double sec = 0; - hostFill(size, A, coef_, sec); // Cpu can access this mem - HIPCHECK(hipFree(A)); - - log_host("ExtDevice: host fill", GBytes, sec); - return true; - } - - bool testExtDeviceMemoryKernelFill(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; - } - - double sec = 0, sec_hv = 0, sec_kv = 0; - kernelFill(size, A, coef_, sec); - hostVerify(size, A, coef_, sec_hv); // Fine grained device memory can be verified by host - kernelVerify(size, A, coef_, sec_kv); - HIPCHECK(hipFree(A)); - - log_kernel("ExtDevice: kernel fill", GBytes, sec, sec_hv, sec_kv); - - return true; - } - - bool testExtDeviceMemory() { - cout << "Test fine grained device memory host filling" - << endl; - for (int i = 0; i < NUM_SIZE; i++) { - if (!testExtDeviceMemoryHostFill(totalSizes_[i], - hipDeviceMallocFinegrained)) { - return false; - } - } - - cout << "Test fine grained device memory kernel filling" - << endl; - for (int i = 0; i < NUM_SIZE; i++) { - if (!testExtDeviceMemoryKernelFill(totalSizes_[i], - hipDeviceMallocFinegrained)) { - return false; - } - } - - return true; - } -#endif - - bool run() { - if (supportLargeBar()) { - if (!testLargeBar()) { - return false; - } - } - - if (supportManagedMemory()) { - if (!testManagedMemory()) { - return false; - } - } - - if (!testHostMemory()) { - return false; - } - -#ifdef __HIP_PLATFORM_AMD__ - if (supportDeviceMallocFinegrained()) { - if (!testExtDeviceMemory()) { - return false; - } - } -#endif - return true; - } - -}; - -int main(int argc, char *argv[]) { - HipTest::parseStandardArguments(argc, argv, true); // For ::p_gpuDevice, ::blocksPerCU, ::threadsPerBlock - cout << "Test int" << endl; - hipPerfMemFill hipPerfMemFillInt(::blocksPerCU, ::threadsPerBlock); - hipPerfMemFillInt.open(::p_gpuDevice); - HIPASSERT(hipPerfMemFillInt.run()); - - cout << "Test double" << endl; - hipPerfMemFill hipPerfMemFillDouble(::blocksPerCU, ::threadsPerBlock); - hipPerfMemFillDouble.open(::p_gpuDevice); - HIPASSERT(hipPerfMemFillDouble.run()); - - passed(); -} diff --git a/perftests/memory/hipPerfMemMallocCpyFree.cpp b/perftests/memory/hipPerfMemMallocCpyFree.cpp deleted file mode 100644 index 94ceb68cba..0000000000 --- a/perftests/memory/hipPerfMemMallocCpyFree.cpp +++ /dev/null @@ -1,124 +0,0 @@ -/* -Copyright (c) 2015 - 2021 Advanced Micro Devices, Inc. All rights reserved. -Permission is hereby granted, free of charge, to any person obtaining a copy -of this software and associated documentation files (the "Software"), to deal -in the Software without restriction, including without limitation the rights -to use, copy, modify, merge, publish, distribute, sublicense, and/or sell -copies of the Software, and to permit persons to whom the Software is -furnished to do so, subject to the following conditions: -The above copyright notice and this permission notice shall be included in -all copies or substantial portions of the Software. -THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR -IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, -FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE -AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER -LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, -OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN -THE SOFTWARE. -*/ - -#include "test_common.h" -#include -#include - -/* HIT_START - * BUILD: %t %s ../../src/test_common.cpp - * TEST: %t - * HIT_END - */ - -#define NUM_SIZE 19 //size up to 16M -#define NUM_ITER 500 //Total GPU memory up to 16M*500=8G - -void valSet(int* A, int val, size_t size) { - size_t len = size / sizeof(int); - for (int i = 0; i < len; i++) { - A[i] = val; - } -} - -void setup(size_t *size, int &num, int **pA, const size_t totalGlobalMem) { - - std::cout << "size: "; - for (int i = 0; i < num; i++) { - size[i] = 1 << (i + 6); - if((NUM_ITER + 1) * size[i] > totalGlobalMem) { - num = i; - break; - } - std::cout << size[i] << " "; - } - std::cout << std::endl; - *pA = (int*)malloc(size[num - 1]); - valSet(*pA, 1, size[num - 1]); -} - -void testInit(size_t size, int *A) { - int *Ad; - clock_t start = clock(); - hipMalloc(&Ad, size); //hip::init() will be called - clock_t end = clock(); - double uS = (end - start) * 1000000. / CLOCKS_PER_SEC; - std::cout << "Initial" << std::endl; - std::cout << "hipMalloc(" << size << ") cost " << uS << "us" << std::endl; - - start = clock(); - hipMemcpy(Ad, A, size, hipMemcpyHostToDevice); - hipDeviceSynchronize(); - end = clock(); - uS = (end - start) * 1000000. / CLOCKS_PER_SEC; - std::cout << "hipMemcpy(" << size << ") cost " << uS << "us" << std::endl; - - start = clock(); - hipFree(Ad); - end = clock(); - uS = (end - start) * 1000000. / CLOCKS_PER_SEC; - std::cout << "hipFree(" << size << ") cost " << uS << "us" << std::endl; -} - -int main() { - double uS; - clock_t start, end; - size_t size[NUM_SIZE] = { 0 }; - int *Ad[NUM_ITER] = { nullptr }; - int *A; - hipDeviceProp_t props; - memset(&props, 0, sizeof(props)); - HIPCHECK(hipGetDeviceProperties(&props, 0)); - std::cout << "totalGlobalMem: " << props.totalGlobalMem << std::endl; - - int num = NUM_SIZE; - setup(size, num, &A, props.totalGlobalMem); - testInit(size[0], A); - - for (int i = 0; i < num; i++) { - std::cout << size[i] << std::endl; - start = clock(); - for (int j = 0; j < NUM_ITER; j++) { - HIPCHECK(hipMalloc(&Ad[j], size[i])); - } - end = clock(); - uS = (end - start) * 1000000. / (NUM_ITER * CLOCKS_PER_SEC); - std::cout << "hipMalloc(" << size[i] << ") cost " << uS << "us" << std::endl; - - start = clock(); - for (int j = 0; j < NUM_ITER; j++) { - HIPCHECK(hipMemcpy(Ad[j], A, size[i], hipMemcpyHostToDevice)); - } - hipDeviceSynchronize(); - end = clock(); - uS = (end - start) * 1000000. / (NUM_ITER * CLOCKS_PER_SEC); - std::cout << "hipMemcpy(" << size[i] << ") cost " << uS << "us" << std::endl; - - start = clock(); - for (int j = 0; j < NUM_ITER; j++) { - HIPCHECK(hipFree(Ad[j])); - Ad[j] = nullptr; - } - end = clock(); - double uS = (end - start) * 1000000. / (NUM_ITER * CLOCKS_PER_SEC); - std::cout << "hipFree(" << size[i] << ") cost " << uS << "us" << std::endl; - } - free(A); - passed(); -} diff --git a/perftests/memory/hipPerfMemcpy.cpp b/perftests/memory/hipPerfMemcpy.cpp deleted file mode 100644 index 9751117eca..0000000000 --- a/perftests/memory/hipPerfMemcpy.cpp +++ /dev/null @@ -1,114 +0,0 @@ -/* - Copyright (c) 2015 - 2021 Advanced Micro Devices, Inc. All rights reserved. - Permission is hereby granted, free of charge, to any person obtaining a copy - of this software and associated documentation files (the "Software"), to deal - in the Software without restriction, including without limitation the rights - to use, copy, modify, merge, publish, distribute, sublicense, and/or sell - copies of the Software, and to permit persons to whom the Software is - furnished to do so, subject to the following conditions: - The above copyright notice and this permission notice shall be included in - all copies or substantial portions of the Software. - THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR - IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, - FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE - AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER - LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, - OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN - THE SOFTWARE. - */ - -/* HIT_START - * BUILD: %t %s ../../src/test_common.cpp - * TEST: %t - * HIT_END - */ - -#include "test_common.h" -#include -#include - -#define NUM_SIZE 8 -#define NUM_ITER 0x40000 - - -using namespace std; - -class hipPerfMemcpy { - private: - unsigned int numBuffers_; - size_t totalSizes_[NUM_SIZE]; - void setHostBuffer(int *A, int val, size_t size); - public: - hipPerfMemcpy(); - ~hipPerfMemcpy() {}; - void open(int deviceID); - void run(unsigned int testNumber); -}; - -hipPerfMemcpy::hipPerfMemcpy() : numBuffers_(0) { - for (int i = 0; i < NUM_SIZE; i++) { - totalSizes_[i] = 1 << (i + 6); - } -}; - -void hipPerfMemcpy::setHostBuffer(int *A, int val, size_t size) { - size_t len = size / sizeof(int); - for (int i = 0; i < len; i++) { - A[i] = val; - } -} - -void hipPerfMemcpy::open(int deviceId) { - int nGpu = 0; - HIPCHECK(hipGetDeviceCount(&nGpu)); - if (nGpu < 1) { - failed("No GPU!"); - } - - HIPCHECK(hipSetDevice(deviceId)); - hipDeviceProp_t props = {0}; - HIPCHECK(hipGetDeviceProperties(&props, deviceId)); - std::cout << "info: running on bus " << "0x" << props.pciBusID << " " << props.name - << " with " << props.multiProcessorCount << " CUs" << " and device id: " << deviceId << std::endl; -} - -void hipPerfMemcpy::run(unsigned int testNumber) { - int *A, *Ad; - A = new int[totalSizes_[testNumber]]; - setHostBuffer(A, 1, totalSizes_[testNumber]); - hipMalloc(&Ad, totalSizes_[testNumber]); - - auto start = chrono::steady_clock::now(); - - for (int j = 0; j < NUM_ITER; j++) { - hipMemcpy(Ad, A, totalSizes_[testNumber], hipMemcpyHostToDevice); - } - - hipDeviceSynchronize(); - - auto end = chrono::steady_clock::now(); - chrono::duration diff = end - start; - - cout << "hipPerfMemcpy[" << testNumber << "] " << "Host to Device copy took " - << diff.count() / NUM_ITER << " us for memory size of " << totalSizes_[testNumber] - << " Bytes" << endl; - - delete [] A; - HIPCHECK(hipFree(Ad)); - -} - - -int main() { - hipPerfMemcpy hipPerfMemcpy; - - int deviceId = 0; - hipPerfMemcpy.open(deviceId); - - for (auto testCase = 0; testCase < NUM_SIZE; testCase++) { - hipPerfMemcpy.run(testCase); - } - - passed(); - -} diff --git a/perftests/memory/hipPerfMemset.cpp b/perftests/memory/hipPerfMemset.cpp deleted file mode 100644 index 2df0c9727b..0000000000 --- a/perftests/memory/hipPerfMemset.cpp +++ /dev/null @@ -1,437 +0,0 @@ -/* - Copyright (c) 2015 - 2023 Advanced Micro Devices, Inc. All rights reserved. - Permission is hereby granted, free of charge, to any person obtaining a copy - of this software and associated documentation files (the "Software"), to deal - in the Software without restriction, including without limitation the rights - to use, copy, modify, merge, publish, distribute, sublicense, and/or sell - copies of the Software, and to permit persons to whom the Software is - furnished to do so, subject to the following conditions: - The above copyright notice and this permission notice shall be included in - all copies or substantial portions of the Software. - THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR - IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, - FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE - AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER - LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, - OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN - THE SOFTWARE. - */ - -/* HIT_START - * BUILD: %t %s ../../src/test_common.cpp - * TEST: %t - * HIT_END - */ - -#include "test_common.h" -#include -#include - -static unsigned int sizeList[] = { - 256, 512, 1024, 2048, 4096, 8192, -}; - -static unsigned int eleNumList[] = { - 0x100, 0x400, 0x1000, 0x4000, 0x10000, 0x20000, 0x40000, 0x80000, 0x100000, - 0x200000, 0x400000, 0x800000, 0x1000000 -}; - -typedef struct _dataType { -char memsetval = 0x42; -char memsetD8val = 0xDE; -int16_t memsetD16val = 0xDEAD; -int memsetD32val = 0xDEADBEEF; -}dataType; - -#define NUM_ITER 1000 - -enum MemsetType { - hipMemsetTypeDefault, - hipMemsetTypeD8, - hipMemsetTypeD16, - hipMemsetTypeD32, - hipMemsetTypeMax - -}; - -using namespace std; - -class hipPerfMemset { - private: - uint64_t bufSize_; - unsigned int num_elements_; - unsigned int testNumEle_; - unsigned int _numSubTests = 0; - unsigned int _numSubTests2D = 0; - unsigned int _numSubTests3D = 0; - unsigned int num_sizes_ =0; - - public: - hipPerfMemset() { - num_elements_ = sizeof(eleNumList) / sizeof(unsigned int); - _numSubTests = num_elements_ * hipMemsetTypeMax; - - num_sizes_ = sizeof(sizeList) / sizeof(unsigned int); - _numSubTests2D = num_sizes_; - _numSubTests3D = _numSubTests2D; - }; - - ~hipPerfMemset() {}; - - void open(int deviceID); - - template - void run1D(unsigned int test, T memsetval, enum MemsetType type, bool async); - - template - void run2D(unsigned int test, T memsetval, enum MemsetType type, bool async); - - template - void run3D(unsigned int test, T memsetval, enum MemsetType type, bool async); - - uint getNumTests() { - return _numSubTests; - } - - uint getNumTests2D() { - return _numSubTests2D; - } - uint getNumTests3D() { - return _numSubTests3D; - } -}; - - -void hipPerfMemset::open(int deviceId) { - int nGpu = 0; - HIPCHECK(hipGetDeviceCount(&nGpu)); - if (nGpu < 1) { - failed("No GPU!"); - } - - HIPCHECK(hipSetDevice(deviceId)); - hipDeviceProp_t props = {0}; - HIPCHECK(hipGetDeviceProperties(&props, deviceId)); - std::cout << "info: running on bus " << "0x" << props.pciBusID << " " << props.name - << " with " << props.multiProcessorCount << " CUs" << " and device id: " << deviceId - << std::endl; -} - -template -void hipPerfMemset::run1D(unsigned int test, T memsetval, enum MemsetType type, bool async) { - - T * A_h; - T * A_d; - - testNumEle_ = eleNumList[test % num_elements_]; - - bufSize_ = testNumEle_ * sizeof(uint32_t); - - HIPCHECK(hipMalloc(&A_d, bufSize_)); - - A_h = reinterpret_cast (malloc(bufSize_)); - - hipStream_t stream; - HIPCHECK(hipStreamCreateWithFlags(&stream, hipStreamNonBlocking)); - - // Warm-up - if (async) { - HIPCHECK(hipMemsetAsync((void *)A_d, memsetval, bufSize_, stream)); - HIPCHECK(hipStreamSynchronize(stream)); - } else { - HIPCHECK(hipMemset((void *)A_d, memsetval, bufSize_)); - HIPCHECK(hipDeviceSynchronize()); - } - - auto start = chrono::high_resolution_clock::now(); - for (uint i = 0; i < NUM_ITER; i++) { - if (type == hipMemsetTypeDefault && !async) { - HIPCHECK(hipMemset((void *)A_d, memsetval, bufSize_)); - } - else if (type == hipMemsetTypeDefault && async) { - HIPCHECK(hipMemsetAsync(A_d, memsetval, bufSize_, stream)); - } - else if (type == hipMemsetTypeD8 && !async){ - HIPCHECK(hipMemsetD8((hipDeviceptr_t)A_d, memsetval, bufSize_)); - } - else if (type == hipMemsetTypeD8 && async) { - HIPCHECK(hipMemsetD8Async((hipDeviceptr_t)A_d, memsetval, bufSize_, stream)); - } - else if (type == hipMemsetTypeD16 && !async) { - HIPCHECK(hipMemsetD16((hipDeviceptr_t)A_d, memsetval, bufSize_/sizeof(T))); - } - else if (type == hipMemsetTypeD16 && async) { - HIPCHECK(hipMemsetD16Async((hipDeviceptr_t)A_d, memsetval, bufSize_/sizeof(T), stream)); - } - else if (type == hipMemsetTypeD32 && !async) { - HIPCHECK(hipMemsetD32((hipDeviceptr_t)A_d, memsetval, bufSize_/sizeof(T))); - } - else if (type == hipMemsetTypeD32 && async) { - HIPCHECK(hipMemsetD32Async((hipDeviceptr_t)A_d, memsetval, bufSize_/sizeof(T), stream)); - } - } - if (async) { - HIPCHECK(hipStreamSynchronize(stream)); - } else { - HIPCHECK(hipDeviceSynchronize()); - } - - auto end = chrono::high_resolution_clock::now(); - - HIPCHECK(hipMemcpy(A_h, A_d, bufSize_, hipMemcpyDeviceToHost) ); - - for (int i = 0; i < bufSize_ / sizeof(T); i++) { - if (A_h[i] != memsetval) { - cout << "mismatch at index " << i << " computed: " << static_cast (A_h[i]) - << ", memsetval: " << static_cast (memsetval) << endl; - break; - } - } - - HIPCHECK(hipFree(A_d)); - free(A_h); - - auto diff = std::chrono::duration(end - start); - auto sec = diff.count(); - - auto perf = static_cast((bufSize_ * NUM_ITER * (double)(1e-09)) / sec); - - cout << "[" << setw(2) << test << "] " << setw(5) << bufSize_/1024 << " Kb " << setw(4) - << " typeSize " << (int)sizeof(T) << " : " << setw(7) << perf << " GB/s " << endl; -} - -template -void hipPerfMemset::run2D(unsigned int test, T memsetval, enum MemsetType type, bool async) { - - bufSize_ = sizeList[test % num_sizes_]; - - size_t numH = bufSize_; - size_t numW = bufSize_; - size_t pitch_A; - size_t width = numW * sizeof(char); - size_t sizeElements = width * numH; - size_t elements = numW* numH; - - T * A_h; - T * A_d; - - HIPCHECK(hipMallocPitch(reinterpret_cast(&A_d), &pitch_A, width , - numH)); - A_h = reinterpret_cast(malloc(sizeElements)); - - for (size_t i=0; i < elements; i++) { - A_h[i] = 1; - } - - hipStream_t stream; - HIPCHECK(hipStreamCreateWithFlags(&stream, hipStreamNonBlocking)); - - // Warm-up - if (async) { - HIPCHECK(hipMemset2DAsync(A_d, pitch_A, memsetval, numW, numH, stream)); - HIPCHECK(hipStreamSynchronize(stream)); - } else { - HIPCHECK(hipMemset2D(A_d, pitch_A, memsetval, numW, numH)); - HIPCHECK(hipDeviceSynchronize()); - } - - auto start = chrono::steady_clock::now(); - - for (uint i = 0; i < NUM_ITER; i++) { - if (type == hipMemsetTypeDefault && !async) { - HIPCHECK(hipMemset2D(A_d, pitch_A, memsetval, numW, numH)); - } - else if (type == hipMemsetTypeDefault && async) { - HIPCHECK(hipMemset2DAsync(A_d, pitch_A, memsetval, numW, numH, stream)); - } - } - - if (async) { - HIPCHECK(hipStreamSynchronize(stream)); - } else { - HIPCHECK(hipDeviceSynchronize()); - } - - auto end = chrono::steady_clock::now(); - - HIPCHECK(hipMemcpy2D(A_h, width, A_d, pitch_A, numW, numH, - hipMemcpyDeviceToHost)); - - for (int i=0; i < elements; i++) { - if (A_h[i] != memsetval) { - cout << "mismatch at index " << i << " computed: " << static_cast (A_h[i]) - << ", memsetval: " << static_cast (memsetval) << endl; - break; - } - } - - chrono::duration diff = end - start; - - auto sec = diff.count(); - - auto perf = static_cast((sizeElements* NUM_ITER * (double)(1e-09)) / sec); - - cout << " hipPerf2DMemset" << (async ? "Async" : " ") << "[" << test << "] " - << " " << "(GB/s) for " << setw(5) << bufSize_ - << " x " << setw(5) << bufSize_ << " bytes : " << setw(7) << perf << endl; - - HIPCHECK(hipStreamDestroy(stream)); - HIPCHECK(hipFree(A_d)); - free(A_h); -} - -template -void hipPerfMemset::run3D(unsigned int test, T memsetval, enum MemsetType type, bool async) { - - bufSize_ = sizeList[test % num_sizes_]; - - size_t numH = bufSize_; - size_t numW = bufSize_; - size_t depth = 10; - size_t width = numW * sizeof(char); - size_t sizeElements = width * numH * depth; - size_t elements = numW* numH* depth; - - hipStream_t stream; - HIPCHECK(hipStreamCreateWithFlags(&stream, hipStreamNonBlocking)); - - T *A_h; - - hipExtent extent = make_hipExtent(width, numH, depth); - hipPitchedPtr devPitchedPtr; - - HIPCHECK(hipMalloc3D(&devPitchedPtr, extent)); - A_h = (char*)malloc(sizeElements); - HIPASSERT(A_h != NULL); - - for (size_t i=0; i (A_h[i]) - << ", memsetval: " << static_cast (memsetval) << endl; - break; - } - } - - chrono::duration diff = end - start; - - auto sec = diff.count(); - - auto perf = static_cast((sizeElements * NUM_ITER * (double)(1e-09)) / sec); - - cout << " hipPerf3DMemset" << (async ? "Async" : " ") << "[" << test << "] " << " " - << "(GB/s) for " << setw(5) << bufSize_ << " x " << setw(5) - << bufSize_ << " x " << depth << " bytes : " << setw(7) << perf << endl; - HIPCHECK(hipFree(devPitchedPtr.ptr)); - free(A_h); -} - -int main() { - hipPerfMemset hipPerfMemset; - - dataType pattern; - int deviceId = 0; - hipPerfMemset.open(deviceId); - MemsetType type; - - int numTests = hipPerfMemset.getNumTests(); - int numTests2D = hipPerfMemset.getNumTests2D(); - int numTests3D = hipPerfMemset.getNumTests3D(); - - - cout << "--------------------- 1D buffer -------------------" << endl; - bool async= false; - for (uint i = 0; i < 2 ; i++) { - cout << endl; - - for (auto testCase = 0; testCase < numTests; testCase++) { - if (testCase < sizeof(eleNumList) / sizeof(uint32_t)) { - cout << "API: hipMemsetD8" << (async ? "Async " : " "); - hipPerfMemset.run1D(testCase, pattern.memsetval, hipMemsetTypeD8, async); - } - - else if (testCase < 2 * sizeof(eleNumList) / sizeof(uint32_t)) { - cout << "API: hipMemsetD16" << (async ? "Async" : " "); - hipPerfMemset.run1D(testCase,pattern.memsetD16val, hipMemsetTypeD16, async); - } - - else if (testCase < 3 * sizeof(eleNumList) / sizeof(uint32_t)) { - cout << "API: hipMemsetD32" << (async ? "Async" : " "); - hipPerfMemset.run1D(testCase,pattern.memsetD32val, hipMemsetTypeD32, async); - } - - else { - cout << "API: hipMemset" << (async ? "Async " : " "); - hipPerfMemset.run1D(testCase,pattern.memsetval, hipMemsetTypeDefault, async); - } - } - async = true; - } - - cout << endl; - cout << "------------------ 2D buffer arrays ---------------" << endl; - - async = false; - for (uint i = 0; i < 2; i++) { - cout << endl; - for (uint test = 0; test < numTests2D; test++) { - hipPerfMemset.run2D(test, pattern.memsetval, hipMemsetTypeDefault, async); - } - async = true; - } - - cout << endl; - cout << "------------------ 3D buffer arrays ---------------" << endl; - - async = false; - for (uint i = 0; i < 2; i++) { - cout << endl; - for (uint test =0; test < numTests3D; test++) { - hipPerfMemset.run3D(test, pattern.memsetval, hipMemsetTypeDefault, async); - } - async = true; - } - - passed(); -} diff --git a/perftests/memory/hipPerfSampleRate.cpp b/perftests/memory/hipPerfSampleRate.cpp deleted file mode 100644 index 1ecadfe749..0000000000 --- a/perftests/memory/hipPerfSampleRate.cpp +++ /dev/null @@ -1,319 +0,0 @@ -/* - Copyright (c) 2015 - 2021 Advanced Micro Devices, Inc. All rights reserved. - Permission is hereby granted, free of charge, to any person obtaining a copy - of this software and associated documentation files (the "Software"), to deal - in the Software without restriction, including without limitation the rights - to use, copy, modify, merge, publish, distribute, sublicense, and/or sell - copies of the Software, and to permit persons to whom the Software is - furnished to do so, subject to the following conditions: - The above copyright notice and this permission notice shall be included in - all copies or substantial portions of the Software. - THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR - IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, - FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE - AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER - LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, - OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN - THE SOFTWARE. - */ - -/* HIT_START - * BUILD: %t %s ../../src/test_common.cpp - * TEST: %t - * HIT_END - */ - -#include -#include -#include "test_common.h" -#include -#include - -using namespace std; - -#define NUM_TYPES 3 -vector types= {"float", "float2", "float4"}; -vector typeSizes = {4, 8, 16}; - -#define NUM_SIZES 12 -vector sizes = {1, 2, 4, 8, 16, 32, - 64, 128, 256, 512, 1024, 2048}; - -#define NUM_BUFS 6 -#define MAX_BUFS (1 << (NUM_BUFS - 1)) - -#ifdef __HIP_PLATFORM_NVIDIA__ -inline __host__ __device__ void operator+=(float2 &a, float2 b) -{ - a.x += b.x; a.y += b.y; -} - -inline __host__ __device__ void operator+=(float4 &a, float4 b) -{ - a.x += b.x; a.y += b.y; a.z += b.z; a.w += b.w; -} -#endif - -template -__global__ void sampleRate(T * outBuffer, unsigned int inBufSize, unsigned int writeIt, - T **inBuffer, int numBufs) { - - uint gid = (blockIdx.x * blockDim.x + threadIdx.x); - uint inputIdx = gid % inBufSize; - - T tmp; - memset(&tmp, 0, sizeof(T)); - for(int i = 0; i < numBufs; i++) { - tmp += *(*(inBuffer+i)+inputIdx); - } - - if (writeIt*(unsigned int)tmp.x) { - outBuffer[gid] = tmp; - } -}; - -template -__global__ void sampleRateFloat(T * outBuffer, unsigned int inBufSize, unsigned int writeIt, - T ** inBuffer, int numBufs) { - - uint gid = (blockIdx.x * blockDim.x + threadIdx.x); - uint inputIdx = gid % inBufSize; - - T tmp = (T)0.0f; - - for(int i = 0; i < numBufs; i++) { - tmp += *((*inBuffer+i)+inputIdx); - } - - if (writeIt*(unsigned int)tmp) { - outBuffer[gid] = tmp; - } -}; - -class hipPerfSampleRate { - public: - hipPerfSampleRate(); - ~hipPerfSampleRate(); - - void open(void); - void run(unsigned int testCase); - void close(void); - - // array of funtion pointers - typedef void (hipPerfSampleRate::*funPtr)(void * outBuffer, unsigned int - inBufSize, unsigned int writeIt, void **inBuffer, int numBufs, int grids, int blocks, - int threads_per_block); - - // Wrappers - void float_kernel(void * outBuffer, unsigned int - inBufSize, unsigned int writeIt, void **inBuffer, int numBufs, int grids, int blocks, - int threads_per_block); - - void float2_kernel(void * outBuffer, unsigned int - inBufSize, unsigned int writeIt, void **inBuffer, int numBufs, int grids, int blocks, - int threads_per_block); - - void float4_kernel(void * outBuffer, unsigned int - inBufSize, unsigned int writeIt, void **inBuffer, int numBufs, int grids, int blocks, - int threads_per_block); - - private: - void setData(void *ptr, unsigned int value); - void checkData(uint *ptr); - - unsigned int width_; - unsigned int bufSize_; - unsigned long long totalIters = 0; - int numCUs; - - unsigned int outBufSize_; - static const unsigned int MAX_ITERATIONS = 25; - unsigned int numBufs_; - unsigned int typeIdx_; -}; - - -hipPerfSampleRate::hipPerfSampleRate() {} - -hipPerfSampleRate::~hipPerfSampleRate() {} - -void hipPerfSampleRate::open(void) { - - int nGpu = 0; - HIPCHECK(hipGetDeviceCount(&nGpu)); - if (nGpu < 1) { - failed("No GPU!"); - } - - int deviceId = 0; - hipDeviceProp_t props = {0}; - props = {0}; - HIPCHECK(hipSetDevice(deviceId)); - HIPCHECK(hipGetDeviceProperties(&props, deviceId)); - std::cout << "info: running on bus " << "0x" << props.pciBusID << " " << props.name - << " with " << props.multiProcessorCount << " CUs" << " and device id: " << deviceId - << std::endl; - numCUs = props.multiProcessorCount; - } - - -void hipPerfSampleRate::close() { - -} - - -// Wrappers for the kernel launches -void hipPerfSampleRate::float_kernel(void * outBuffer, unsigned int inBufSize, - unsigned int writeIt, void **inBuffer, - int numBufs, int grids, int blocks, int threads_per_block) { - - hipLaunchKernelGGL(sampleRateFloat, dim3(grids, grids, grids), dim3 (blocks), 0, 0, - (float*)outBuffer, inBufSize, writeIt, (float**)inBuffer, numBufs); - -} - -void hipPerfSampleRate::float2_kernel(void * outBuffer, unsigned int inBufSize, - unsigned int writeIt, void **inBuffer, - int grids, int blocks, int threads_per_block, int numBufs) { - - hipLaunchKernelGGL(sampleRate, dim3(grids, grids, grids), dim3(blocks), 0, 0, - (float2 *)outBuffer, inBufSize, writeIt, (float2**)inBuffer, numBufs); -} - -void hipPerfSampleRate::float4_kernel(void * outBuffer, unsigned int inBufSize, - unsigned int writeIt, void **inBuffer, - int grids, int blocks, int threads_per_block, int numBufs) { - - hipLaunchKernelGGL(sampleRate, dim3(grids, grids, grids), dim3(blocks), 0, 0, - (float4 *) outBuffer, inBufSize, writeIt, (float4**)inBuffer, numBufs); -} - -void hipPerfSampleRate::run(unsigned int test) { - - funPtr p[] = {&hipPerfSampleRate::float_kernel, &hipPerfSampleRate::float2_kernel, - &hipPerfSampleRate::float4_kernel}; - - // We compute a square domain - width_ = sizes[test % NUM_SIZES]; - typeIdx_ = (test / NUM_SIZES) % NUM_TYPES; - bufSize_ = width_ * width_ * typeSizes[typeIdx_]; - numBufs_ = (1 << (test / (NUM_SIZES * NUM_TYPES))); - - void * hOutPtr; - void * dOutPtr; - void * hInPtr[numBufs_]; - void ** dPtr; - void * dInPtr[numBufs_]; - - outBufSize_ = - sizes[NUM_SIZES - 1] * sizes[NUM_SIZES - 1] * typeSizes[NUM_TYPES - 1]; - - // Allocate memory on the host and device - HIPCHECK(hipHostMalloc((void **)&hOutPtr, outBufSize_, hipHostMallocDefault)); - setData((void *)hOutPtr, 0xdeadbeef); - HIPCHECK(hipMalloc((uint **)&dOutPtr, outBufSize_)); - - // Allocate 2D array in Device - hipMalloc((void **)&dPtr, numBufs_* sizeof(void *)); - - for (uint i = 0; i < numBufs_; i++) { - HIPCHECK(hipHostMalloc((void **)&hInPtr[i], bufSize_, hipHostMallocDefault)); - HIPCHECK(hipMalloc((uint **)&dInPtr[i], bufSize_)); - setData(hInPtr[i], 0x3f800000); - } - - // Populate array of pointers with array addresses - hipMemcpy(dPtr, dInPtr, numBufs_* sizeof(void *), hipMemcpyHostToDevice); - - // Copy memory from host to device - for (uint i = 0; i < numBufs_; i++) { - HIPCHECK(hipMemcpy(dInPtr[i], hInPtr[i], bufSize_, hipMemcpyHostToDevice)); - } - - HIPCHECK(hipMemcpy(dOutPtr, hOutPtr, outBufSize_, hipMemcpyHostToDevice)); - - // Prepare kernel launch parameters - // outBufSize_/sizeof(uint) - Grid size in 3D - int grids = 64; - int blocks = 64; - int threads_per_block = 1; - - unsigned int maxIter = MAX_ITERATIONS * (MAX_BUFS / numBufs_); - unsigned int sizeDW = width_ * width_; - unsigned int writeIt = 0; - - int idx = 0; - - if (!types[typeIdx_].compare("float")) { - idx = 0; - } - else if(!types[typeIdx_].compare("float2")) { - idx = 1; - } - else if(!types[typeIdx_].compare("float4")) { - idx = 2; - } - - - // Time the kernel execution - auto all_start = std::chrono::steady_clock::now(); - for (uint i = 0; i < maxIter; i++) { - (this->*p[idx]) ((void *)dOutPtr, sizeDW, writeIt, dPtr, numBufs_, grids, blocks, - threads_per_block); - } - - hipDeviceSynchronize(); - auto all_end = std::chrono::steady_clock::now(); - std::chrono::duration all_kernel_time = all_end - all_start; - - double perf = ((double)outBufSize_ * numBufs_ * (double)maxIter * (double)(1e-09)) / - all_kernel_time.count(); - - cout << "Domain " << sizes[NUM_SIZES - 1] << "x"<< sizes[NUM_SIZES - 1] << " bufs " - << numBufs_ << " " << types[typeIdx_] << " " << width_<<"x"< -#include -#include "test_common.h" - -using namespace std; - -#define sharedMemSize1 2048 -#define sharedMemSize2 256 - -__global__ void sharedMemReadSpeed1(float *outBuf, ulong N) { - - size_t gid = (blockIdx.x * blockDim.x + threadIdx.x); - size_t lid = threadIdx.x; - __shared__ float local[sharedMemSize1]; - - float val1 = 0; - float val2 = 0; - float val3 = 0; - float val4 = 0; - - for (int i = 0; i < (sharedMemSize1 / 64); i++) { - local[lid + i * 64] = lid; - } - - __syncthreads(); - - val1 += local[lid]; - val2 += local[lid + 64]; - val3 += local[lid + 128]; - val4 += local[lid + 192]; - val1 += local[lid + 256]; - val2 += local[lid + 320]; - val3 += local[lid + 384]; - val4 += local[lid + 448]; - val1 += local[lid + 512]; - val2 += local[lid + 576]; - val3 += local[lid + 640]; - val4 += local[lid + 704]; - val1 += local[lid + 768]; - val2 += local[lid + 832]; - val3 += local[lid + 896]; - val4 += local[lid + 960]; - val1 += local[lid + 1024]; - val2 += local[lid + 1088]; - val3 += local[lid + 1152]; - val4 += local[lid + 1216]; - val1 += local[lid + 1280]; - val2 += local[lid + 1344]; - val3 += local[lid + 1408]; - val4 += local[lid + 1472]; - val1 += local[lid + 1536]; - val2 += local[lid + 1600]; - val3 += local[lid + 1664]; - val4 += local[lid + 1728]; - val1 += local[lid + 1792]; - val2 += local[lid + 1856]; - val3 += local[lid + 1920]; - val4 += local[lid + 1984]; - - if (gid < N) { - outBuf[gid] = val1 + val2 + val3 + val4; - } -}; - -__global__ void sharedMemReadSpeed2(float *outBuf, ulong N) { - size_t gid = (blockIdx.x * blockDim.x + threadIdx.x); - size_t lid = threadIdx.x; - __shared__ float local[sharedMemSize2]; - - float val0 = 0.0f; - float val1 = 0.0f; - - for (int i = 0; i < (sharedMemSize2 / 64); i++) { - local[lid + i * 64] = lid; - } - - __syncthreads(); - -#pragma nounroll - for (uint i = 0; i < 32; i++) { - val0 += local[8 * i + 0]; - val1 += local[8 * i + 1]; - val0 += local[8 * i + 2]; - val1 += local[8 * i + 3]; - val0 += local[8 * i + 4]; - val1 += local[8 * i + 5]; - val0 += local[8 * i + 6]; - val1 += local[8 * i + 7]; - } - - if (gid < N) { - outBuf[gid] = val0 + val1; - } -}; - -int main(int argc, char *argv[]) { - float *dDst; - float *hDst; - hipStream_t stream; - constexpr uint numSizes = 4; - constexpr uint Sizes[numSizes] = {262144, 1048576, 4194304, 16777216}; - uint numReads1 = 32; - uint numReads2 = 256; - uint sharedMemSizeBytes1 = sharedMemSize1 * sizeof(float); - uint sharedMemSizeBytes2 = sharedMemSize2 * sizeof(float); - int nIter = 1000; - const unsigned threadsPerBlock = 64; - - int nGpu = 0; - HIPCHECK(hipGetDeviceCount(&nGpu)); - if (nGpu < 1) { - cout << "info: didn't find any GPU! skipping the test!\n"; - passed(); - return 0; - } - - static int device = 0; - HIPCHECK(hipSetDevice(device)); - hipDeviceProp_t props; - HIPCHECK(hipGetDeviceProperties(&props, device)); - cout << "info: running on bus " << "0x" << props.pciBusID << " " << props.name - << " with " << props.multiProcessorCount << " CUs" << endl; - - HIPCHECK(hipStreamCreate(&stream)); - - for (int nTest = 0; nTest < numSizes; nTest++) { - uint nBytes = Sizes[nTest % numSizes]; - ulong N = nBytes / sizeof(float); - const unsigned blocks = N / threadsPerBlock; - - hDst = new float[nBytes]; - HIPCHECK(hDst == 0 ? hipErrorOutOfMemory : hipSuccess); - memset(hDst, 0, nBytes); - - HIPCHECK(hipMalloc(&dDst, nBytes)); - HIPCHECK(hipMemcpy(dDst, hDst, nBytes, hipMemcpyHostToDevice)); - - hipLaunchKernelGGL(sharedMemReadSpeed1, dim3(blocks), dim3(threadsPerBlock), - 0, stream, dDst, N); - HIPCHECK(hipMemcpy(hDst, dDst, nBytes, hipMemcpyDeviceToHost)); - hipDeviceSynchronize(); - - int tmp = 0; - for (int i = 0; i < N; i++) { - if (i % threadsPerBlock == 0) { - tmp = 0; - } - if (hDst[i] != tmp) { - cout << "info: Data validation failed for warm up run!" << endl; - cout << "info: expected " << tmp << " got " << hDst[i] << endl; - HIPCHECK (hipErrorUnknown); - } - tmp += threadsPerBlock / 2; - } - - auto all_start = chrono::steady_clock::now(); - for (int i = 0; i < nIter; i++) { - hipLaunchKernelGGL(sharedMemReadSpeed1, dim3(blocks), - dim3(threadsPerBlock), 0, stream, dDst, N); - } - hipDeviceSynchronize(); - - auto all_end = chrono::steady_clock::now(); - chrono::duration all_kernel_time = all_end - all_start; - - // read speed in GB/s - double perf = ((double) blocks * threadsPerBlock - * (numReads1 * sizeof(float) + sharedMemSizeBytes1 / 64) * nIter - * (double) (1e-09)) / all_kernel_time.count(); - - cout << "info: read speed = " << setw(8) << perf << " GB/s for " - << sharedMemSizeBytes1 / 1024 << " KB shared memory" - " with " << setw(8) << blocks * threadsPerBlock << " threads, " - << setw(4) << numReads1 << " reads in sharedMemReadSpeed1 kernel" << endl; - - delete[] hDst; - hipFree(dDst); - } - - - for (int nTest = 0; nTest < numSizes; nTest++) { - uint nBytes = Sizes[nTest % numSizes]; - ulong N = nBytes / sizeof(float); - const unsigned blocks = N / threadsPerBlock; - - hDst = new float[nBytes]; - HIPCHECK(hDst == 0 ? hipErrorOutOfMemory : hipSuccess); - memset(hDst, 0, nBytes); - - HIPCHECK(hipMalloc(&dDst, nBytes)); - HIPCHECK(hipMemcpy(dDst, hDst, nBytes, hipMemcpyHostToDevice)); - - hipLaunchKernelGGL(sharedMemReadSpeed2, dim3(blocks), dim3(threadsPerBlock), - 0, stream, dDst, N); - HIPCHECK(hipMemcpy(hDst, dDst, nBytes, hipMemcpyDeviceToHost)); - hipDeviceSynchronize(); - - auto all_start = chrono::steady_clock::now(); - for (int i = 0; i < nIter; i++) { - hipLaunchKernelGGL(sharedMemReadSpeed2, dim3(blocks), - dim3(threadsPerBlock), 0, stream, dDst, N); - } - hipDeviceSynchronize(); - - auto all_end = chrono::steady_clock::now(); - chrono::duration all_kernel_time = all_end - all_start; - - // read speed in GB/s - double perf = ((double) blocks * threadsPerBlock - * (numReads2 * sizeof(float) + sharedMemSizeBytes2 / 64) * nIter - * (double) (1e-09)) / all_kernel_time.count(); - - cout << "info: read speed = " << setw(8) << perf << " GB/s for " - << sharedMemSizeBytes2 / 1024 << " KB shared memory" - " with " << setw(8) << blocks * threadsPerBlock << " threads, " - << setw(4) << numReads2 << " reads in sharedMemReadSpeed2 kernel" << endl; - - delete[] hDst; - hipFree(dDst); - } - - HIPCHECK(hipStreamDestroy(stream)); - - passed(); -}