From cc6a87e9e38976445b50ce8d16f0aecce930e008 Mon Sep 17 00:00:00 2001 From: rohit pathania Date: Thu, 30 Apr 2020 03:34:07 -0400 Subject: [PATCH] [ dtest ] hipModuleLaunchKernel multiThreaded n multiGPU scenarios 1.Added hipModuleLaunchKernel multithreaded multi GPU scenario. 2.removed hipCtxCreate API from earlier test as it is deprecated. SWDEV-238517 for enhancing hip unit tests Change-Id: Id102d80887b6ff61a59938dbeb9fa2a26a3275b2 --- .../hipModuleLoadDataMultThreadOnMultGPU.cpp | 145 +++++++++++++++++ .../module/hipModuleLoadDataMultThreaded.cpp | 152 ++++++++---------- 2 files changed, 213 insertions(+), 84 deletions(-) create mode 100644 tests/src/runtimeApi/module/hipModuleLoadDataMultThreadOnMultGPU.cpp diff --git a/tests/src/runtimeApi/module/hipModuleLoadDataMultThreadOnMultGPU.cpp b/tests/src/runtimeApi/module/hipModuleLoadDataMultThreadOnMultGPU.cpp new file mode 100644 index 0000000000..ce78590147 --- /dev/null +++ b/tests/src/runtimeApi/module/hipModuleLoadDataMultThreadOnMultGPU.cpp @@ -0,0 +1,145 @@ +/* +Copyright (c) 2020-Present 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 WARRANNTY OF ANY KIND, EXPRESS OR +IMPLIED, INNCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANNY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +/* HIT_START + * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS -std=c++11 + * TEST: %t + * HIT_END + */ + + +#include +#include +#include +#include +#include + +#include "hip/hip_runtime.h" +#include "hip/hip_runtime_api.h" + + +#include "test_common.h" + +#define LEN 64 +#define SIZE LEN << 2 +#define THREADS 8 +#define MAX_THREADS 512 + +#define FILENAME "vcpy_kernel.code" +#define kernel_name "hello_world" + +std::vector load_file() { + std::ifstream file(FILENAME, std::ios::binary | std::ios::ate); + std::streamsize fsize = file.tellg(); + file.seekg(0, std::ios::beg); + + std::vector buffer(fsize); + if (!file.read(buffer.data(), fsize)) { + failed("could not open code object '%s'\n", FILENAME); + } + return buffer; +} + +void run(const std::vector& buffer, int deviceNo) { + hipSetDevice(deviceNo); + hipModule_t Module; + hipFunction_t Function; + HIPCHECK(hipModuleLoadData(&Module, &buffer[0])); + HIPCHECK(hipModuleGetFunction(&Function, Module, kernel_name)); + + float *A, *B, *Ad, *Bd; + A = new float[LEN]; + B = new float[LEN]; + + for (uint32_t i = 0; i < LEN; i++) { + A[i] = i * 1.0f; + B[i] = 0.0f; + } + + HIPCHECK(hipMalloc(&Ad, SIZE)); + HIPCHECK(hipMalloc(&Bd, SIZE)); + + HIPCHECK(hipMemcpy(Ad, A, SIZE, hipMemcpyHostToDevice)); + HIPCHECK(hipMemcpy(Bd, B, SIZE, hipMemcpyHostToDevice)); + + hipStream_t stream; + HIPCHECK(hipStreamCreate(&stream)); + + struct { + void* _Ad; + void* _Bd; + } args; + args._Ad = static_cast(Ad); + args._Bd = static_cast(Bd); + size_t size = sizeof(args); + + void* config[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &args, HIP_LAUNCH_PARAM_BUFFER_SIZE, &size, + HIP_LAUNCH_PARAM_END}; + HIPCHECK(hipModuleLaunchKernel(Function, 1, 1, 1, LEN, 1, 1, 0, stream, NULL, (void**)&config)); + + HIPCHECK(hipStreamDestroy(stream)); + + HIPCHECK(hipModuleUnload(Module)); + + HIPCHECK(hipMemcpy(B, Bd, SIZE, hipMemcpyDeviceToHost)); + + for (uint32_t i = 0; i < LEN; i++) { + assert(A[i] == B[i]); + } + hipFree(Ad); + hipFree(Bd); + delete[] A; + delete[] B; +} + +struct joinable_thread : std::thread { + template + joinable_thread(Xs&&... xs) : std::thread(std::forward(xs)...) {} // NOLINT + + joinable_thread& operator=(joinable_thread&& other) = default; + joinable_thread(joinable_thread&& other) = default; + + ~joinable_thread() { + if (this->joinable()) + this->join(); + } +}; + +void run_multi_threads(uint32_t n, const std::vector& buffer) { + int numDevices = 0; + HIPCHECK(hipGetDeviceCount(&numDevices)); + + std::vector threads; + + for (int deviceNo=0; deviceNo < numDevices; ++deviceNo) { + for (uint32_t i = 0; i < n; i++) { + threads.emplace_back(std::thread{[&, buffer] { + run(buffer, deviceNo); + }}); + } + } +} + +int main() { + HIPCHECK(hipInit(0)); + auto buffer = load_file(); + run_multi_threads(min(THREADS * std::thread::hardware_concurrency(), MAX_THREADS), buffer); + + passed(); +} diff --git a/tests/src/runtimeApi/module/hipModuleLoadDataMultThreaded.cpp b/tests/src/runtimeApi/module/hipModuleLoadDataMultThreaded.cpp index eef367ab70..6ae1b92ab3 100644 --- a/tests/src/runtimeApi/module/hipModuleLoadDataMultThreaded.cpp +++ b/tests/src/runtimeApi/module/hipModuleLoadDataMultThreaded.cpp @@ -35,118 +35,102 @@ THE SOFTWARE. #define LEN 64 #define SIZE LEN << 2 -#define THREADS 2 -#define MAX_THREADS 16 +#define THREADS 8 +#define MAX_THREADS 512 #define FILENAME "vcpy_kernel.code" #define kernel_name "hello_world" -std::vector load_file() -{ - std::ifstream file(FILENAME, std::ios::binary | std::ios::ate); - std::streamsize fsize = file.tellg(); - file.seekg(0, std::ios::beg); +std::vector load_file() { + std::ifstream file(FILENAME, std::ios::binary | std::ios::ate); + std::streamsize fsize = file.tellg(); + file.seekg(0, std::ios::beg); - std::vector buffer(fsize); - if (!file.read(buffer.data(), fsize)) { - failed("could not open code object '%s'\n", FILENAME); - } - return buffer; + std::vector buffer(fsize); + if (!file.read(buffer.data(), fsize)) { + failed("could not open code object '%s'\n", FILENAME); + } + return buffer; } void run(const std::vector& buffer) { - hipDevice_t device; - HIPCHECK(hipDeviceGet(&device, 0)); - hipCtx_t context; - HIPCHECK(hipCtxCreate(&context, 0, device)); + hipModule_t Module; + hipFunction_t Function; + HIPCHECK(hipModuleLoadData(&Module, &buffer[0])); + HIPCHECK(hipModuleGetFunction(&Function, Module, kernel_name)); - hipModule_t Module; - hipFunction_t Function; - HIPCHECK(hipModuleLoadData(&Module, &buffer[0])); - HIPCHECK(hipModuleGetFunction(&Function, Module, kernel_name)); - - float *A, *B, *Ad, *Bd; - A = new float[LEN]; - B = new float[LEN]; + float *A, *B, *Ad, *Bd; + A = new float[LEN]; + B = new float[LEN]; - for (uint32_t i = 0; i < LEN; i++) { - A[i] = i * 1.0f; - B[i] = 0.0f; - } + for (uint32_t i = 0; i < LEN; i++) { + A[i] = i * 1.0f; + B[i] = 0.0f; + } - HIPCHECK(hipMalloc((void**)&Ad, SIZE)); - HIPCHECK(hipMalloc((void**)&Bd, SIZE)); + HIPCHECK(hipMalloc((void**)&Ad, SIZE)); + HIPCHECK(hipMalloc((void**)&Bd, SIZE)); - HIPCHECK(hipMemcpy(Ad, A, SIZE, hipMemcpyHostToDevice)); - HIPCHECK(hipMemcpy(Bd, B, SIZE, hipMemcpyHostToDevice)); + HIPCHECK(hipMemcpy(Ad, A, SIZE, hipMemcpyHostToDevice)); + HIPCHECK(hipMemcpy(Bd, B, SIZE, hipMemcpyHostToDevice)); - hipStream_t stream; - HIPCHECK(hipStreamCreate(&stream)); + hipStream_t stream; + HIPCHECK(hipStreamCreate(&stream)); - struct { - void* _Ad; - void* _Bd; - } args; - args._Ad = (void*) Ad; - args._Bd = (void*) Bd; - size_t size = sizeof(args); + struct { + void* _Ad; + void* _Bd; + } args; + args._Ad = (void*) Ad; + args._Bd = (void*) Bd; + size_t size = sizeof(args); - void* config[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &args, HIP_LAUNCH_PARAM_BUFFER_SIZE, &size, - HIP_LAUNCH_PARAM_END}; - HIPCHECK(hipModuleLaunchKernel(Function, 1, 1, 1, LEN, 1, 1, 0, stream, NULL, (void**)&config)); + void* config[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &args, HIP_LAUNCH_PARAM_BUFFER_SIZE, &size, + HIP_LAUNCH_PARAM_END}; + HIPCHECK(hipModuleLaunchKernel(Function, 1, 1, 1, LEN, 1, 1, 0, stream, NULL, (void**)&config)); - HIPCHECK(hipStreamDestroy(stream)); + HIPCHECK(hipStreamDestroy(stream)); - HIPCHECK(hipModuleUnload(Module)); + HIPCHECK(hipModuleUnload(Module)); - HIPCHECK(hipMemcpy(B, Bd, SIZE, hipMemcpyDeviceToHost)); + HIPCHECK(hipMemcpy(B, Bd, SIZE, hipMemcpyDeviceToHost)); - for (uint32_t i = 0; i < LEN; i++) { - assert(A[i] == B[i]); - } - - hipFree(Ad); - hipFree(Bd); - delete[] A; - delete[] B; - hipCtxDestroy(context); - + for (uint32_t i = 0; i < LEN; i++) { + assert(A[i] == B[i]); + } + + hipFree(Ad); + hipFree(Bd); + delete[] A; + delete[] B; } -struct joinable_thread : std::thread -{ - template - joinable_thread(Xs&&... xs) : std::thread(std::forward(xs)...) // NOLINT - { - } +struct joinable_thread : std::thread { + template + joinable_thread(Xs&&... xs) : std::thread(std::forward(xs)...) {} // NOLINT - joinable_thread& operator=(joinable_thread&& other) = default; - joinable_thread(joinable_thread&& other) = default; + joinable_thread& operator=(joinable_thread&& other) = default; + joinable_thread(joinable_thread&& other) = default; - ~joinable_thread() - { - if(this->joinable()) - this->join(); - } + ~joinable_thread() { + if (this->joinable()) + this->join(); + } }; void run_multi_threads(uint32_t n, const std::vector& buffer) { - - std::vector threads; - - for (uint32_t i = 0; i < n; i++) { - threads.emplace_back(std::thread{[&, buffer] { - run(buffer); - }}); - } - + std::vector threads; + for (uint32_t i = 0; i < n; i++) { + threads.emplace_back(std::thread{[&, buffer] { + run(buffer); + }}); + } } int main() { + HIPCHECK(hipInit(0)); + auto buffer = load_file(); + run_multi_threads(min(THREADS * std::thread::hardware_concurrency(), MAX_THREADS), buffer); - HIPCHECK(hipInit(0)); - auto buffer = load_file(); - run_multi_threads(min(THREADS * std::thread::hardware_concurrency(), MAX_THREADS), buffer); - - passed(); + passed(); }