From 1354a447e1306fce446459d5ad25cab303c865d7 Mon Sep 17 00:00:00 2001 From: Paul Date: Thu, 6 Feb 2020 11:56:17 -0600 Subject: [PATCH 01/17] Use deque instead of vector for code readers so that the iterators and references will be stable --- hipamd/src/program_state.inl | 5 +++-- 1 file changed, 3 insertions(+), 2 deletions(-) diff --git a/hipamd/src/program_state.inl b/hipamd/src/program_state.inl index 272addd053..8df0f1cc0f 100644 --- a/hipamd/src/program_state.inl +++ b/hipamd/src/program_state.inl @@ -26,6 +26,7 @@ #include #include #include +#include #include #include #include @@ -202,7 +203,7 @@ public: std::function>; std::pair< std::mutex, - std::vector>> code_readers; + std::deque>> code_readers; program_state_impl() { // Create placeholder for each agent for the per-agent members. @@ -418,7 +419,7 @@ public: decltype(code_readers.second)::iterator it; { std::lock_guard lck{code_readers.first}; - it = code_readers.second.emplace(code_readers.second.end(), + it = code_readers.second.emplace_back(code_readers.second.end(), move(file), move(tmp)); } From bb26e99c735ca49b88ff23215c80faf89c828650 Mon Sep 17 00:00:00 2001 From: Paul Date: Thu, 6 Feb 2020 12:04:50 -0600 Subject: [PATCH 02/17] Fix compile error --- hipamd/src/program_state.inl | 3 +-- 1 file changed, 1 insertion(+), 2 deletions(-) diff --git a/hipamd/src/program_state.inl b/hipamd/src/program_state.inl index 8df0f1cc0f..0442874900 100644 --- a/hipamd/src/program_state.inl +++ b/hipamd/src/program_state.inl @@ -419,8 +419,7 @@ public: decltype(code_readers.second)::iterator it; { std::lock_guard lck{code_readers.first}; - it = code_readers.second.emplace_back(code_readers.second.end(), - move(file), move(tmp)); + it = code_readers.second.emplace_back(move(file), move(tmp)); } auto check_hsa_error = [](hsa_status_t s) { From 84d6eb985d4983db208e2f3137f19fbdecc6b843 Mon Sep 17 00:00:00 2001 From: Paul Date: Thu, 6 Feb 2020 12:15:29 -0600 Subject: [PATCH 03/17] Assign the iterator --- hipamd/src/program_state.inl | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/hipamd/src/program_state.inl b/hipamd/src/program_state.inl index 0442874900..c8eb2f297c 100644 --- a/hipamd/src/program_state.inl +++ b/hipamd/src/program_state.inl @@ -419,7 +419,8 @@ public: decltype(code_readers.second)::iterator it; { std::lock_guard lck{code_readers.first}; - it = code_readers.second.emplace_back(move(file), move(tmp)); + code_readers.second.emplace_back(move(file), move(tmp)); + it = std::prev(code_readers.second.end()); } auto check_hsa_error = [](hsa_status_t s) { From bf69c7ae32fd86f95b471c38adc91df82bb873e4 Mon Sep 17 00:00:00 2001 From: Paul Date: Thu, 6 Feb 2020 16:21:40 -0600 Subject: [PATCH 04/17] Add multithreaded test --- .../module/hipModuleLoadDataMultThreaded.cpp | 144 ++++++++++++++++++ 1 file changed, 144 insertions(+) create mode 100644 hipamd/tests/src/runtimeApi/module/hipModuleLoadDataMultThreaded.cpp diff --git a/hipamd/tests/src/runtimeApi/module/hipModuleLoadDataMultThreaded.cpp b/hipamd/tests/src/runtimeApi/module/hipModuleLoadDataMultThreaded.cpp new file mode 100644 index 0000000000..f989a14a06 --- /dev/null +++ b/hipamd/tests/src/runtimeApi/module/hipModuleLoadDataMultThreaded.cpp @@ -0,0 +1,144 @@ +/* +Copyright (c) 2015-2016 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 "hip/hip_runtime.h" +#include "hip/hip_runtime_api.h" +#include +#include +#include +#include +#include + +#include "test_common.h" + +#define LEN 64 +#define SIZE LEN << 2 +#define THREADS 64 + +#define FILENAME "vcpy_kernel.code" +#define kernel_name "hello_world" + +using ModuleFunction = std::pair; + +ModuleFunction load() { + hipModule_t Module; + hipFunction_t Function; + 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)) { + HIPCHECK(hipModuleLoadData(&Module, &buffer[0])); + HIPCHECK(hipModuleGetFunction(&Function, Module, kernel_name)); + } + else { + failed("could not open code object '%s'\n", FILENAME); + } + return {Module, Function}; +} + +void run(ModuleFunction mf) { + hipModule_t Module = mf.first; + hipFunction_t Function = mf.second; + 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((void**)&Ad, SIZE)); + HIPCHECK(hipMalloc((void**)&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 = (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)); + + 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]); + } +} + +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) { + std::vector mf(n); + { + std::vector threads; + for (uint32_t i = 0; i < n; i++) { + threads.emplace_back(std::thread{[=, &mf] { + mf[i] = load(); + }}); + } + } + for(auto&& x:mf) + run(x); +} + +int main() { + + HIPCHECK(hipInit(0)); + run_multi_threads(THREADS); + + passed(); +} From fb8c7cb3eca35514dc45349e1d05cd5b79ad7302 Mon Sep 17 00:00:00 2001 From: Paul Date: Thu, 6 Feb 2020 16:23:29 -0600 Subject: [PATCH 05/17] Make threads a multiple of hardware concurrency --- .../src/runtimeApi/module/hipModuleLoadDataMultThreaded.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/hipamd/tests/src/runtimeApi/module/hipModuleLoadDataMultThreaded.cpp b/hipamd/tests/src/runtimeApi/module/hipModuleLoadDataMultThreaded.cpp index f989a14a06..31e930086b 100644 --- a/hipamd/tests/src/runtimeApi/module/hipModuleLoadDataMultThreaded.cpp +++ b/hipamd/tests/src/runtimeApi/module/hipModuleLoadDataMultThreaded.cpp @@ -35,7 +35,7 @@ THE SOFTWARE. #define LEN 64 #define SIZE LEN << 2 -#define THREADS 64 +#define THREADS 4 #define FILENAME "vcpy_kernel.code" #define kernel_name "hello_world" @@ -138,7 +138,7 @@ void run_multi_threads(uint32_t n) { int main() { HIPCHECK(hipInit(0)); - run_multi_threads(THREADS); + run_multi_threads(THREADS * std::thread::hardware_concurrency()); passed(); } From 2405ab236f03e6d22e9ae59f9cd99aadb9702b74 Mon Sep 17 00:00:00 2001 From: Paul Date: Fri, 7 Feb 2020 10:13:28 -0600 Subject: [PATCH 06/17] Output on failure --- hipamd/Jenkinsfile | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/hipamd/Jenkinsfile b/hipamd/Jenkinsfile index b8bd24cd74..734e875e03 100644 --- a/hipamd/Jenkinsfile +++ b/hipamd/Jenkinsfile @@ -177,7 +177,7 @@ def docker_build_inside_image( def build_image, String inside_args, String platf cd ${build_dir_rel} make install -j\$(nproc) make build_tests -i -j\$(nproc) - ctest -E "(hipMultiThreadDevice-pyramid|hipMemoryAllocateCoherentDriver)" + ctest --output-on-failure -E "(hipMultiThreadDevice-pyramid|hipMemoryAllocateCoherentDriver)" """ // If unit tests output a junit or xunit file in the future, jenkins can parse that file // to display test results on the dashboard From 9494d0e3c9619a122b6de3bf19e9126a312c1297 Mon Sep 17 00:00:00 2001 From: Paul Date: Mon, 10 Feb 2020 13:37:45 -0600 Subject: [PATCH 07/17] Add setDevice to try and initialize the context on cuda --- .../src/runtimeApi/module/hipModuleLoadDataMultThreaded.cpp | 1 + 1 file changed, 1 insertion(+) diff --git a/hipamd/tests/src/runtimeApi/module/hipModuleLoadDataMultThreaded.cpp b/hipamd/tests/src/runtimeApi/module/hipModuleLoadDataMultThreaded.cpp index 31e930086b..ff2d5b1ef3 100644 --- a/hipamd/tests/src/runtimeApi/module/hipModuleLoadDataMultThreaded.cpp +++ b/hipamd/tests/src/runtimeApi/module/hipModuleLoadDataMultThreaded.cpp @@ -127,6 +127,7 @@ void run_multi_threads(uint32_t n) { std::vector threads; for (uint32_t i = 0; i < n; i++) { threads.emplace_back(std::thread{[=, &mf] { + hipSetDevice(0); mf[i] = load(); }}); } From bb145e77aedf960346ab97d7920a4672f7feb468 Mon Sep 17 00:00:00 2001 From: Paul Date: Mon, 10 Feb 2020 15:52:34 -0600 Subject: [PATCH 08/17] Create context for cuda --- .../module/hipModuleLoadDataMultThreaded.cpp | 39 +++++++++++++------ 1 file changed, 27 insertions(+), 12 deletions(-) diff --git a/hipamd/tests/src/runtimeApi/module/hipModuleLoadDataMultThreaded.cpp b/hipamd/tests/src/runtimeApi/module/hipModuleLoadDataMultThreaded.cpp index ff2d5b1ef3..92a1cb77c7 100644 --- a/hipamd/tests/src/runtimeApi/module/hipModuleLoadDataMultThreaded.cpp +++ b/hipamd/tests/src/runtimeApi/module/hipModuleLoadDataMultThreaded.cpp @@ -35,28 +35,31 @@ THE SOFTWARE. #define LEN 64 #define SIZE LEN << 2 -#define THREADS 4 +#define THREADS 8 #define FILENAME "vcpy_kernel.code" #define kernel_name "hello_world" using ModuleFunction = std::pair; -ModuleFunction load() { - hipModule_t Module; - hipFunction_t Function; +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)) { - HIPCHECK(hipModuleLoadData(&Module, &buffer[0])); - HIPCHECK(hipModuleGetFunction(&Function, Module, kernel_name)); - } - else { + if (!file.read(buffer.data(), fsize)) { failed("could not open code object '%s'\n", FILENAME); } + return buffer; +} + +ModuleFunction load(const std::vector& buffer) { + hipModule_t Module; + hipFunction_t Function; + HIPCHECK(hipModuleLoadData(&Module, &buffer[0])); + HIPCHECK(hipModuleGetFunction(&Function, Module, kernel_name)); return {Module, Function}; } @@ -121,25 +124,37 @@ struct joinable_thread : std::thread } }; +hipCtx_t create_context() { + hipDevice_t device; + HIPCHECK(hipDeviceGet(&device, 0)); + + hipCtx_t ctx; + HIPCHECK(hipCtxCreate(&ctx, 0, device)); + return ctx; +} + void run_multi_threads(uint32_t n) { std::vector mf(n); { + auto buffer = load_file(); std::vector threads; for (uint32_t i = 0; i < n; i++) { - threads.emplace_back(std::thread{[=, &mf] { - hipSetDevice(0); - mf[i] = load(); + threads.emplace_back(std::thread{[&, i, buffer] { + mf[i] = load(buffer); }}); } } for(auto&& x:mf) run(x); + } int main() { HIPCHECK(hipInit(0)); + auto ctx = create_context(); run_multi_threads(THREADS * std::thread::hardware_concurrency()); + hipCtxDestroy(ctx); passed(); } From 7a6e88480bc6114df5f3bac5b0da43c21940e404 Mon Sep 17 00:00:00 2001 From: Paul Date: Mon, 10 Feb 2020 16:01:53 -0600 Subject: [PATCH 09/17] Set context on each thread --- .../src/runtimeApi/module/hipModuleLoadDataMultThreaded.cpp | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/hipamd/tests/src/runtimeApi/module/hipModuleLoadDataMultThreaded.cpp b/hipamd/tests/src/runtimeApi/module/hipModuleLoadDataMultThreaded.cpp index 92a1cb77c7..cdb4c81c94 100644 --- a/hipamd/tests/src/runtimeApi/module/hipModuleLoadDataMultThreaded.cpp +++ b/hipamd/tests/src/runtimeApi/module/hipModuleLoadDataMultThreaded.cpp @@ -134,27 +134,27 @@ hipCtx_t create_context() { } void run_multi_threads(uint32_t n) { + auto ctx = create_context(); std::vector mf(n); { auto buffer = load_file(); std::vector threads; for (uint32_t i = 0; i < n; i++) { threads.emplace_back(std::thread{[&, i, buffer] { + HIPCHECK(hipCtxSetCurrent(ctx)); mf[i] = load(buffer); }}); } } for(auto&& x:mf) run(x); - + hipCtxDestroy(ctx); } int main() { HIPCHECK(hipInit(0)); - auto ctx = create_context(); run_multi_threads(THREADS * std::thread::hardware_concurrency()); - hipCtxDestroy(ctx); passed(); } From 5d24a2beef3dd5b4e072206f548fc63becd1b302 Mon Sep 17 00:00:00 2001 From: Paul Date: Mon, 10 Feb 2020 16:37:34 -0600 Subject: [PATCH 10/17] Reduce threads on cuda --- .../src/runtimeApi/module/hipModuleLoadDataMultThreaded.cpp | 4 ++++ 1 file changed, 4 insertions(+) diff --git a/hipamd/tests/src/runtimeApi/module/hipModuleLoadDataMultThreaded.cpp b/hipamd/tests/src/runtimeApi/module/hipModuleLoadDataMultThreaded.cpp index cdb4c81c94..3b22176927 100644 --- a/hipamd/tests/src/runtimeApi/module/hipModuleLoadDataMultThreaded.cpp +++ b/hipamd/tests/src/runtimeApi/module/hipModuleLoadDataMultThreaded.cpp @@ -35,7 +35,11 @@ THE SOFTWARE. #define LEN 64 #define SIZE LEN << 2 +#ifdef __CUDACC__ +#define THREADS 1 +#else #define THREADS 8 +#endif #define FILENAME "vcpy_kernel.code" #define kernel_name "hello_world" From accdd882dc6ea1dc75e50218c9f3bab3165cd6a6 Mon Sep 17 00:00:00 2001 From: Paul Date: Mon, 10 Feb 2020 17:23:58 -0600 Subject: [PATCH 11/17] Skip test on cuda --- .../module/hipModuleLoadDataMultThreaded.cpp | 18 +----------------- 1 file changed, 1 insertion(+), 17 deletions(-) diff --git a/hipamd/tests/src/runtimeApi/module/hipModuleLoadDataMultThreaded.cpp b/hipamd/tests/src/runtimeApi/module/hipModuleLoadDataMultThreaded.cpp index 3b22176927..03a2b82b8b 100644 --- a/hipamd/tests/src/runtimeApi/module/hipModuleLoadDataMultThreaded.cpp +++ b/hipamd/tests/src/runtimeApi/module/hipModuleLoadDataMultThreaded.cpp @@ -18,7 +18,7 @@ THE SOFTWARE. */ /* HIT_START - * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS -std=c++11 + * BUILD: %t %s ../../test_common.cpp EXCLUDE_HIP_PLATFORM nvcc * TEST: %t * HIT_END */ @@ -35,11 +35,7 @@ THE SOFTWARE. #define LEN 64 #define SIZE LEN << 2 -#ifdef __CUDACC__ -#define THREADS 1 -#else #define THREADS 8 -#endif #define FILENAME "vcpy_kernel.code" #define kernel_name "hello_world" @@ -128,31 +124,19 @@ struct joinable_thread : std::thread } }; -hipCtx_t create_context() { - hipDevice_t device; - HIPCHECK(hipDeviceGet(&device, 0)); - - hipCtx_t ctx; - HIPCHECK(hipCtxCreate(&ctx, 0, device)); - return ctx; -} - void run_multi_threads(uint32_t n) { - auto ctx = create_context(); std::vector mf(n); { auto buffer = load_file(); std::vector threads; for (uint32_t i = 0; i < n; i++) { threads.emplace_back(std::thread{[&, i, buffer] { - HIPCHECK(hipCtxSetCurrent(ctx)); mf[i] = load(buffer); }}); } } for(auto&& x:mf) run(x); - hipCtxDestroy(ctx); } int main() { From 86cdb30195ed26f4f0c514d15f947dd0397bc366 Mon Sep 17 00:00:00 2001 From: Paul Date: Tue, 11 Feb 2020 11:26:24 -0600 Subject: [PATCH 12/17] Try to initialize the primary context on cuda --- .../runtimeApi/module/hipModuleLoadDataMultThreaded.cpp | 8 +++++++- 1 file changed, 7 insertions(+), 1 deletion(-) diff --git a/hipamd/tests/src/runtimeApi/module/hipModuleLoadDataMultThreaded.cpp b/hipamd/tests/src/runtimeApi/module/hipModuleLoadDataMultThreaded.cpp index 03a2b82b8b..6612392064 100644 --- a/hipamd/tests/src/runtimeApi/module/hipModuleLoadDataMultThreaded.cpp +++ b/hipamd/tests/src/runtimeApi/module/hipModuleLoadDataMultThreaded.cpp @@ -18,7 +18,7 @@ THE SOFTWARE. */ /* HIT_START - * BUILD: %t %s ../../test_common.cpp EXCLUDE_HIP_PLATFORM nvcc + * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS -std=c++11 * TEST: %t * HIT_END */ @@ -127,10 +127,16 @@ struct joinable_thread : std::thread void run_multi_threads(uint32_t n) { std::vector mf(n); { + hipDevice_t device; + HIPCHECK(hipDeviceGet(&device, 0)); + auto buffer = load_file(); std::vector threads; for (uint32_t i = 0; i < n; i++) { threads.emplace_back(std::thread{[&, i, buffer] { + hipCtx_t ctx; + HIPCHECK(hipDevicePrimaryCtxRetain(&ctx, device)); + mf[i] = load(buffer); }}); } From e8a7cc69579405b129b5fe2a70db83d7c12a0d29 Mon Sep 17 00:00:00 2001 From: Paul Date: Tue, 11 Feb 2020 11:46:29 -0600 Subject: [PATCH 13/17] Push ctx to the stack as current --- .../src/runtimeApi/module/hipModuleLoadDataMultThreaded.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/hipamd/tests/src/runtimeApi/module/hipModuleLoadDataMultThreaded.cpp b/hipamd/tests/src/runtimeApi/module/hipModuleLoadDataMultThreaded.cpp index 6612392064..09fedc3660 100644 --- a/hipamd/tests/src/runtimeApi/module/hipModuleLoadDataMultThreaded.cpp +++ b/hipamd/tests/src/runtimeApi/module/hipModuleLoadDataMultThreaded.cpp @@ -136,7 +136,7 @@ void run_multi_threads(uint32_t n) { threads.emplace_back(std::thread{[&, i, buffer] { hipCtx_t ctx; HIPCHECK(hipDevicePrimaryCtxRetain(&ctx, device)); - + HIPCHECK(hipCtxPushCurrent(ctx)); mf[i] = load(buffer); }}); } From cd279cd5748c2667466c147c172ca744fc2efa92 Mon Sep 17 00:00:00 2001 From: Paul Date: Tue, 11 Feb 2020 12:34:10 -0600 Subject: [PATCH 14/17] Revert "Push ctx to the stack as current" This reverts commit e8a7cc69579405b129b5fe2a70db83d7c12a0d29. --- .../src/runtimeApi/module/hipModuleLoadDataMultThreaded.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/hipamd/tests/src/runtimeApi/module/hipModuleLoadDataMultThreaded.cpp b/hipamd/tests/src/runtimeApi/module/hipModuleLoadDataMultThreaded.cpp index 09fedc3660..6612392064 100644 --- a/hipamd/tests/src/runtimeApi/module/hipModuleLoadDataMultThreaded.cpp +++ b/hipamd/tests/src/runtimeApi/module/hipModuleLoadDataMultThreaded.cpp @@ -136,7 +136,7 @@ void run_multi_threads(uint32_t n) { threads.emplace_back(std::thread{[&, i, buffer] { hipCtx_t ctx; HIPCHECK(hipDevicePrimaryCtxRetain(&ctx, device)); - HIPCHECK(hipCtxPushCurrent(ctx)); + mf[i] = load(buffer); }}); } From dd6676dadf886ce6a7a5d87d647dc0403706b20e Mon Sep 17 00:00:00 2001 From: Paul Date: Tue, 11 Feb 2020 12:34:11 -0600 Subject: [PATCH 15/17] Revert "Try to initialize the primary context on cuda" This reverts commit 86cdb30195ed26f4f0c514d15f947dd0397bc366. --- .../runtimeApi/module/hipModuleLoadDataMultThreaded.cpp | 8 +------- 1 file changed, 1 insertion(+), 7 deletions(-) diff --git a/hipamd/tests/src/runtimeApi/module/hipModuleLoadDataMultThreaded.cpp b/hipamd/tests/src/runtimeApi/module/hipModuleLoadDataMultThreaded.cpp index 6612392064..03a2b82b8b 100644 --- a/hipamd/tests/src/runtimeApi/module/hipModuleLoadDataMultThreaded.cpp +++ b/hipamd/tests/src/runtimeApi/module/hipModuleLoadDataMultThreaded.cpp @@ -18,7 +18,7 @@ THE SOFTWARE. */ /* HIT_START - * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS -std=c++11 + * BUILD: %t %s ../../test_common.cpp EXCLUDE_HIP_PLATFORM nvcc * TEST: %t * HIT_END */ @@ -127,16 +127,10 @@ struct joinable_thread : std::thread void run_multi_threads(uint32_t n) { std::vector mf(n); { - hipDevice_t device; - HIPCHECK(hipDeviceGet(&device, 0)); - auto buffer = load_file(); std::vector threads; for (uint32_t i = 0; i < n; i++) { threads.emplace_back(std::thread{[&, i, buffer] { - hipCtx_t ctx; - HIPCHECK(hipDevicePrimaryCtxRetain(&ctx, device)); - mf[i] = load(buffer); }}); } From 627d9a1f46686d0745e96a8bdb399d14401c5c1b Mon Sep 17 00:00:00 2001 From: Satyanvesh Dittakavi Date: Thu, 13 Feb 2020 16:34:05 +0530 Subject: [PATCH 16/17] updated test for nvidia path --- .../module/hipModuleLoadDataMultThreaded.cpp | 55 ++++++++++--------- 1 file changed, 29 insertions(+), 26 deletions(-) diff --git a/hipamd/tests/src/runtimeApi/module/hipModuleLoadDataMultThreaded.cpp b/hipamd/tests/src/runtimeApi/module/hipModuleLoadDataMultThreaded.cpp index 03a2b82b8b..6115125399 100644 --- a/hipamd/tests/src/runtimeApi/module/hipModuleLoadDataMultThreaded.cpp +++ b/hipamd/tests/src/runtimeApi/module/hipModuleLoadDataMultThreaded.cpp @@ -1,5 +1,5 @@ /* -Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved. +Copyright (c) 2015-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 @@ -18,7 +18,7 @@ THE SOFTWARE. */ /* HIT_START - * BUILD: %t %s ../../test_common.cpp EXCLUDE_HIP_PLATFORM nvcc + * BUILD: %t %s ../../test_common.cpp * TEST: %t * HIT_END */ @@ -35,13 +35,11 @@ THE SOFTWARE. #define LEN 64 #define SIZE LEN << 2 -#define THREADS 8 +#define THREADS 2 #define FILENAME "vcpy_kernel.code" #define kernel_name "hello_world" -using ModuleFunction = std::pair; - std::vector load_file() { std::ifstream file(FILENAME, std::ios::binary | std::ios::ate); @@ -55,18 +53,18 @@ std::vector load_file() return buffer; } -ModuleFunction load(const std::vector& 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)); - return {Module, Function}; -} - -void run(ModuleFunction mf) { - hipModule_t Module = mf.first; - hipFunction_t Function = mf.second; - float *A, *B, *Ad, *Bd; + + float *A, *B, *Ad, *Bd; A = new float[LEN]; B = new float[LEN]; @@ -105,6 +103,13 @@ void run(ModuleFunction mf) { for (uint32_t i = 0; i < LEN; i++) { assert(A[i] == B[i]); } + + hipFree(Ad); + hipFree(Bd); + delete A; + delete B; + hipCtxDestroy(context); + } struct joinable_thread : std::thread @@ -124,25 +129,23 @@ struct joinable_thread : std::thread } }; -void run_multi_threads(uint32_t n) { - std::vector mf(n); - { - auto buffer = load_file(); - std::vector threads; - for (uint32_t i = 0; i < n; i++) { - threads.emplace_back(std::thread{[&, i, buffer] { - mf[i] = load(buffer); - }}); - } +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); + }}); } - for(auto&& x:mf) - run(x); + } int main() { HIPCHECK(hipInit(0)); - run_multi_threads(THREADS * std::thread::hardware_concurrency()); + auto buffer = load_file(); + run_multi_threads(THREADS * std::thread::hardware_concurrency(), buffer); passed(); } From 2034ed35614bb4dc05179051dc43f59d1a8a37fc Mon Sep 17 00:00:00 2001 From: Satyanvesh Dittakavi Date: Thu, 13 Feb 2020 19:48:26 +0530 Subject: [PATCH 17/17] Add c++11 option for nvcc --- .../src/runtimeApi/module/hipModuleLoadDataMultThreaded.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/hipamd/tests/src/runtimeApi/module/hipModuleLoadDataMultThreaded.cpp b/hipamd/tests/src/runtimeApi/module/hipModuleLoadDataMultThreaded.cpp index 6115125399..8591a748df 100644 --- a/hipamd/tests/src/runtimeApi/module/hipModuleLoadDataMultThreaded.cpp +++ b/hipamd/tests/src/runtimeApi/module/hipModuleLoadDataMultThreaded.cpp @@ -18,7 +18,7 @@ THE SOFTWARE. */ /* HIT_START - * BUILD: %t %s ../../test_common.cpp + * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS -std=c++11 * TEST: %t * HIT_END */