diff --git a/rocclr/hip_code_object.cpp b/rocclr/hip_code_object.cpp index 95df5173bd..671ccd6d14 100755 --- a/rocclr/hip_code_object.cpp +++ b/rocclr/hip_code_object.cpp @@ -365,7 +365,7 @@ static bool isCodeObjectCompatibleWithDevice(std::string co_triple_target_id, // This will be moved to COMGR eventually hipError_t CodeObject::ExtractCodeObjectFromFile(amd::Os::FileDesc fdesc, size_t fsize, - const std::vector& device_names, + const void ** image, const std::vector& device_names, std::vector>& code_objs) { hipError_t hip_error = hipSuccess; @@ -375,18 +375,14 @@ hipError_t CodeObject::ExtractCodeObjectFromFile(amd::Os::FileDesc fdesc, size_t } // Map the file to memory, with offset 0. - const void* image = nullptr; - if (!amd::Os::MemoryMapFileDesc(fdesc, fsize, 0, &image)) { + //file will be unmapped in ModuleUnload + //const void* image = nullptr; + if (!amd::Os::MemoryMapFileDesc(fdesc, fsize, 0, image)) { return hipErrorInvalidValue; } // retrieve code_objs{binary_image, binary_size} for devices - hip_error = extractCodeObjectFromFatBinary(image, device_names, code_objs); - - // Unmap the file memory after extracting code object. - if (!amd::Os::MemoryUnmapFile(image, fsize)) { - return hipErrorInvalidValue; - } + hip_error = extractCodeObjectFromFatBinary(*image, device_names, code_objs); return hip_error; } diff --git a/rocclr/hip_code_object.hpp b/rocclr/hip_code_object.hpp index bc67b1a414..8bfa79c4df 100755 --- a/rocclr/hip_code_object.hpp +++ b/rocclr/hip_code_object.hpp @@ -52,7 +52,7 @@ class CodeObject { // Given an file desc and file size, extracts to code object for corresponding devices, // return code_objs{binary_ptr, binary_size}, which could be used to determine foffset static hipError_t ExtractCodeObjectFromFile(amd::Os::FileDesc fdesc, size_t fsize, - const std::vector& device_names, + const void ** image, const std::vector& device_names, std::vector>& code_objs); // Given an ptr to memory, extracts to code object for corresponding devices, diff --git a/rocclr/hip_fatbin.cpp b/rocclr/hip_fatbin.cpp index 8304d361fb..f3902ec879 100755 --- a/rocclr/hip_fatbin.cpp +++ b/rocclr/hip_fatbin.cpp @@ -30,13 +30,12 @@ FatBinaryInfo::~FatBinaryInfo() { } if (fdesc_ > 0) { - if (!amd::Os::CloseFileHandle(fdesc_)) { - guarantee(false, "Cannot close file"); - } - if (fsize_ && !amd::Os::MemoryUnmapFile(image_, fsize_)) { guarantee(false, "Cannot unmap file"); } + if (!amd::Os::CloseFileHandle(fdesc_)) { + guarantee(false, "Cannot close file"); + } } fname_ = std::string(); @@ -68,14 +67,9 @@ hipError_t FatBinaryInfo::ExtractFatBinary(const std::vector& devi } // Extract the code object from file - hip_error = CodeObject::ExtractCodeObjectFromFile(fdesc_, fsize_, + hip_error = CodeObject::ExtractCodeObjectFromFile(fdesc_, fsize_, &image_, device_names, code_objs); - // Map the file memory, Later: only map offset, throws error in ElfMagic now. - if (!amd::Os::MemoryMapFileDesc(fdesc_, fsize_, 0, &image_)) { - return hipErrorInvalidValue; - } - } else if (image_ != nullptr) { // We are directly given image pointer directly, try to extract file desc & file Size hip_error = CodeObject::ExtractCodeObjectFromMemory(image_, diff --git a/samples/1_Utils/hipDispatchLatency/hipDispatchEnqueueRateMT.cpp b/samples/1_Utils/hipDispatchLatency/hipDispatchEnqueueRateMT.cpp index 731e9c2770..6584f70b84 100644 --- a/samples/1_Utils/hipDispatchLatency/hipDispatchEnqueueRateMT.cpp +++ b/samples/1_Utils/hipDispatchLatency/hipDispatchEnqueueRateMT.cpp @@ -54,17 +54,6 @@ THE SOFTWARE. __global__ void EmptyKernel() {} -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; -} // Helper to print various timing metrics void print_timing(std::string test, std::array &results, int batch = 1) { @@ -94,7 +83,7 @@ void print_timing(std::string test, std::array &results, } // Measure time taken to enqueue a kernel on the GPU using hipModuleLaunchKernel -void hipModuleLaunchKernel_enqueue_rate(const std::vector buffer, std::atomic_int* shared, int max_threads) +void hipModuleLaunchKernel_enqueue_rate(const std::vector& buffer, std::atomic_int* shared, int max_threads) { //resources necessary for this thread hipStream_t stream; @@ -118,12 +107,13 @@ void hipModuleLaunchKernel_enqueue_rate(const std::vector buffer, std::ato auto stop = std::chrono::high_resolution_clock::now(); results[i] = std::chrono::duration(stop - start).count(); } + HIPCHECK(hipModuleUnload(module)); print_timing("Thread ID : " + std::to_string(tid) + " , " + "hipModuleLaunchKernel enqueue rate", results); HIPCHECK(hipStreamDestroy(stream)); } // Measure time taken to enqueue a kernel on the GPU using hipLaunchKernelGGL -void hipLaunchKernelGGL_enqueue_rate(const std::vector buffer, std::atomic_int* shared, int max_threads) +void hipLaunchKernelGGL_enqueue_rate(const std::vector& buffer, std::atomic_int* shared, int max_threads) { //resources necessary for this thread hipStream_t stream; @@ -146,11 +136,20 @@ void hipLaunchKernelGGL_enqueue_rate(const std::vector buffer, std::atomic // Simple thread pool struct thread_pool { - thread_pool(int total_threads) : max_threads(total_threads) {} - void start(std::function, std::atomic_int*, int)> f) { - auto buffer = load_file(); + thread_pool(int total_threads) : max_threads(total_threads) { + std::ifstream file(FILENAME, std::ios::binary | std::ios::ate); + std::streamsize fsize = file.tellg(); + file.seekg(0, std::ios::beg); + + buffer.resize(fsize); + if (!file.read(buffer.data(), fsize)) { + failed("could not open code object '%s'\n", FILENAME); + } + file.close(); + } + void start(std::function&, std::atomic_int*, int)> f) { for (int i = 0; i < max_threads; ++i) { - threads.push_back(std::async(std::launch::async, f, buffer, &shared, max_threads)); + threads.push_back(std::async(std::launch::async, f, std::ref(buffer), &shared, max_threads)); } } void finish() { @@ -165,6 +164,7 @@ struct thread_pool { } private: std::atomic_int shared {0}; + std::vector buffer; std::vector> threads; int max_threads = 1; }; @@ -176,7 +176,6 @@ int main(int argc, char* argv[]) std::cerr << "Run test as 'hipDispatchEnqueueRateMT <0-hipModuleLaunchKernel /1-hipLaunchKernelGGL>'\n"; return -1; } - int max_threads = atoi(argv[1]); int run_module_test = atoi(argv[2]); if(max_threads < 1 || run_module_test < 0 || run_module_test > 1) { @@ -195,4 +194,3 @@ int main(int argc, char* argv[]) } return 0; } - diff --git a/tests/src/runtimeApi/module/empty_kernel.cpp b/tests/src/runtimeApi/module/empty_kernel.cpp new file mode 100644 index 0000000000..67ae593bf3 --- /dev/null +++ b/tests/src/runtimeApi/module/empty_kernel.cpp @@ -0,0 +1,24 @@ +/* +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 +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 "hip/hip_runtime.h" + +extern "C" __global__ void EmptyKernel() { +} + diff --git a/tests/src/runtimeApi/module/hipModuleLoadMultiThreaded.cpp b/tests/src/runtimeApi/module/hipModuleLoadMultiThreaded.cpp new file mode 100644 index 0000000000..18424314e4 --- /dev/null +++ b/tests/src/runtimeApi/module/hipModuleLoadMultiThreaded.cpp @@ -0,0 +1,115 @@ +/* +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 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: empty_kernel.code %hc --genco %S/empty_kernel.cpp -o empty_kernel.code + * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS -std=c++11 + * TEST: %t + * HIT_END + */ + +#include +#include "hip/hip_runtime.h" +#ifdef __HIP_PLATFORM_AMD__ +#include "hip/hip_ext.h" +#endif +#include +#include +#include +#include +#include +#include +#include +#include +#include +#define THREADS 8 +#define MAX_NUM_THREADS 512 + +#include "test_common.h" + +#define NUM_GROUPS 1 +#define GROUP_SIZE 1 +#define WARMUP_RUN_COUNT 10 +#define TIMING_RUN_COUNT 100 +#define TOTAL_RUN_COUNT WARMUP_RUN_COUNT + TIMING_RUN_COUNT +#define FILENAME "empty_kernel.code" +#define kernel_name "EmptyKernel" + +void hipModuleLaunchKernel_enqueue(std::atomic_int* shared, int max_threads) +{ + //resources necessary for this thread + hipStream_t stream; + HIPCHECK(hipStreamCreate(&stream)); + hipModule_t module; + hipFunction_t function; + + HIPCHECK(hipModuleLoad(&module, FILENAME)); + HIPCHECK(hipModuleGetFunction(&function, module, kernel_name)); + + void* kernel_params = nullptr; + std::array results; + + //synchronize all threads, before running + int tid = shared->fetch_add(1, std::memory_order_release); + while (max_threads != shared->load(std::memory_order_acquire)) {} + + for (auto i = 0; i < TOTAL_RUN_COUNT; ++i) { + HIPCHECK(hipModuleLaunchKernel(function, 1, 1, 1, 1, 1, 1, 0, stream, &kernel_params, nullptr)); + } + HIPCHECK(hipModuleUnload(module)); + HIPCHECK(hipStreamDestroy(stream)); +} + +// thread pool +struct thread_pool { + thread_pool(int total_threads) : max_threads(total_threads) { + } + void start(std::function f) { + for (int i = 0; i < max_threads; ++i) { + threads.push_back(std::async(std::launch::async, f, &shared, max_threads)); + } + } + void finish() { + for (auto&&thread : threads) { + thread.get(); + } + threads.clear(); + shared = 0; + } + ~thread_pool() { + finish(); + } +private: + std::atomic_int shared {0}; + std::vector buffer; + std::vector> threads; + int max_threads = 1; +}; + + +int main(int argc, char* argv[]) +{ + int max_threads = min(THREADS * std::thread::hardware_concurrency(), MAX_NUM_THREADS); + thread_pool task(max_threads); + + task.start(hipModuleLaunchKernel_enqueue); + task.finish(); + + passed(); +}