From 2d517fdcc65c959bd0ebe3165bb837d6b03a2758 Mon Sep 17 00:00:00 2001 From: Jatin Date: Fri, 8 May 2020 18:18:36 +0000 Subject: [PATCH 01/33] Adding changes for hipExtLaunchKernel for rocCLR Change-Id: Iba52bc3bde7c37f3fb375a55ba0947e87b3cdc9b --- include/hip/hcc_detail/hip_runtime.h | 43 +++++ include/hip/hcc_detail/hip_runtime_api.h | 5 + include/hip/hip_ext.h | 29 ++- rocclr/hip_hcc.def.in | 1 + rocclr/hip_hcc.map.in | 1 + rocclr/hip_module.cpp | 35 ++++ rocclr/hip_platform.cpp | 34 ++-- src/hip_module.cpp | 27 ++- tests/src/deviceLib/hipLaunchKernelFunc.cpp | 193 ++++++++++++++++++++ tests/src/kernel/hipExtLaunchKernelGGL.cpp | 49 ++++- 10 files changed, 396 insertions(+), 21 deletions(-) create mode 100644 tests/src/deviceLib/hipLaunchKernelFunc.cpp diff --git a/include/hip/hcc_detail/hip_runtime.h b/include/hip/hcc_detail/hip_runtime.h index 28d3ae7051..a166935823 100644 --- a/include/hip/hcc_detail/hip_runtime.h +++ b/include/hip/hcc_detail/hip_runtime.h @@ -392,10 +392,53 @@ extern void ihipPostLaunchKernel(const char* kernelName, hipStream_t stream, gri typedef int hipLaunchParm; +template ::type* = nullptr> +void pArgs(const std::tuple&, void*) {} + +template ::type* = nullptr> +void pArgs(const std::tuple& formals, void** _vargs) { + using T = typename std::tuple_element >::type; + + static_assert(!std::is_reference{}, + "A __global__ function cannot have a reference as one of its " + "arguments."); +#if defined(HIP_STRICT) + static_assert(std::is_trivially_copyable{}, + "Only TriviallyCopyable types can be arguments to a __global__ " + "function"); +#endif + _vargs[n] = const_cast(reinterpret_cast(&std::get(formals))); + return pArgs(formals, _vargs); +} + +template +std::tuple validateArgsCountType(void (*kernel)(Formals...), std::tuple(actuals)) { + static_assert(sizeof...(Formals) == sizeof...(Actuals), "Argument Count Mismatch"); + std::tuple to_formals{std::move(actuals)}; + return to_formals; +} + +#if defined(HIP_TEMPLATE_KERNEL_LAUNCH) +template +void hipLaunchKernelGGL(F kernel, const dim3& numBlocks, const dim3& dimBlocks, + std::uint32_t sharedMemBytes, hipStream_t stream, Args... args) { + constexpr size_t count = sizeof...(Args); + auto tup_ = std::tuple{args...}; + auto tup = validateArgsCountType(kernel, tup_); + void* _Args[count]; + pArgs<0>(tup, _Args); + + auto k = reinterpret_cast(kernel); + hipLaunchKernel(k, numBlocks, dimBlocks, _Args, sharedMemBytes, stream); +} +#else #define hipLaunchKernelGGL(kernelName, numblocks, numthreads, memperblock, streamId, ...) \ do { \ kernelName<<<(numblocks), (numthreads), (memperblock), (streamId)>>>(__VA_ARGS__); \ } while (0) +#endif #include diff --git a/include/hip/hcc_detail/hip_runtime_api.h b/include/hip/hcc_detail/hip_runtime_api.h index df5c4e5b60..e374707cb3 100644 --- a/include/hip/hcc_detail/hip_runtime_api.h +++ b/include/hip/hcc_detail/hip_runtime_api.h @@ -3378,6 +3378,11 @@ hipError_t hipLaunchKernel(const void* function_address, hipStream_t stream __dparm(0)); #if __HIP_ROCclr__ || !defined(__HCC__) +//TODO: Move this to hip_ext.h +hipError_t hipExtLaunchKernel(const void* function_address, dim3 numBlocks, dim3 dimBlocks, + void** args, size_t sharedMemBytes, hipStream_t stream, + hipEvent_t startEvent, hipEvent_t stopEvent, int flags); + hipError_t hipBindTexture( size_t* offset, const textureReference* tex, diff --git a/include/hip/hip_ext.h b/include/hip/hip_ext.h index 90d1e34d2d..c16e841719 100644 --- a/include/hip/hip_ext.h +++ b/include/hip/hip_ext.h @@ -23,6 +23,10 @@ THE SOFTWARE. #ifndef HIP_INCLUDE_HIP_HIP_EXT_H #define HIP_INCLUDE_HIP_HIP_EXT_H #include "hip/hip_runtime.h" +#if defined(__cplusplus) +#include +#include +#endif #ifdef __HCC__ // Forward declarations: @@ -109,8 +113,29 @@ hipError_t hipHccModuleLaunchKernel(hipFunction_t f, uint32_t globalWorkSizeX, hipEvent_t stopEvent = nullptr) __attribute__((deprecated("use hipExtModuleLaunchKernel instead"))); -//#if !__HIP_ROCclr__ && defined(__cplusplus) -#if defined(__HIP_PLATFORM_HCC__) && GENERIC_GRID_LAUNCH == 1 && defined(__HCC__) +#if defined(__HIP_ROCclr__) && defined(__cplusplus) + +extern "C" hipError_t hipExtLaunchKernel(const void* function_address, dim3 numBlocks, + dim3 dimBlocks, void** args, size_t sharedMemBytes, + hipStream_t stream, hipEvent_t startEvent, + hipEvent_t stopEvent, int flags); + +template +inline void hipExtLaunchKernelGGL(F kernel, const dim3& numBlocks, const dim3& dimBlocks, + std::uint32_t sharedMemBytes, hipStream_t stream, + hipEvent_t startEvent, hipEvent_t stopEvent, std::uint32_t flags, + Args... args) { + constexpr size_t count = sizeof...(Args); + auto tup_ = std::tuple{args...}; + auto tup = validateArgsCountType(kernel, tup_); + void* _Args[count]; + pArgs<0>(tup, _Args); + + auto k = reinterpret_cast(kernel); + hipExtLaunchKernel(k, numBlocks, dimBlocks, _Args, sharedMemBytes, stream, startEvent, + stopEvent, (int)flags); +} +#elif defined(__HIP_PLATFORM_HCC__) && GENERIC_GRID_LAUNCH == 1 && defined(__HCC__) //kernel_descriptor and hip_impl::make_kernarg are in "grid_launch_GGL.hpp" namespace hip_impl { diff --git a/rocclr/hip_hcc.def.in b/rocclr/hip_hcc.def.in index a3476c7f33..3a54928764 100755 --- a/rocclr/hip_hcc.def.in +++ b/rocclr/hip_hcc.def.in @@ -51,6 +51,7 @@ hipExtGetLinkTypeAndHopCount hipExtLaunchMultiKernelMultiDevice hipExtMallocWithFlags hipExtModuleLaunchKernel +hipExtLaunchKernel hipFree hipFreeArray hipFuncSetCacheConfig diff --git a/rocclr/hip_hcc.map.in b/rocclr/hip_hcc.map.in index 11637f2696..42c5335f35 100755 --- a/rocclr/hip_hcc.map.in +++ b/rocclr/hip_hcc.map.in @@ -52,6 +52,7 @@ global: hipExtLaunchMultiKernelMultiDevice; hipExtMallocWithFlags; hipExtModuleLaunchKernel; + hipExtLaunchKernel; hipFree; hipFreeArray; hipFuncSetCacheConfig; diff --git a/rocclr/hip_module.cpp b/rocclr/hip_module.cpp index 8f3d4ca936..a29ad9a962 100755 --- a/rocclr/hip_module.cpp +++ b/rocclr/hip_module.cpp @@ -29,6 +29,16 @@ hipError_t ihipModuleLoadData(hipModule_t* module, const void* mmap_ptr, size_t mmap_size); +extern hipError_t ihipLaunchKernel(const void* hostFunction, + dim3 gridDim, + dim3 blockDim, + void** args, + size_t sharedMemBytes, + hipStream_t stream, + hipEvent_t startEvent, + hipEvent_t stopEvent, + int flags); + const std::string& FunctionName(const hipFunction_t f) { return hip::Function::asFunction(f)->function_->name(); @@ -539,6 +549,31 @@ hipError_t hipModuleLaunchKernelExt(hipFunction_t f, uint32_t gridDimX, sharedMemBytes, hStream, kernelParams, extra, startEvent, stopEvent)); } +extern "C" hipError_t hipLaunchKernel(const void *hostFunction, + dim3 gridDim, + dim3 blockDim, + void** args, + size_t sharedMemBytes, + hipStream_t stream) +{ + HIP_INIT_API(NONE, hostFunction, gridDim, blockDim, args, sharedMemBytes, stream); + HIP_RETURN(ihipLaunchKernel(hostFunction, gridDim, blockDim, args, sharedMemBytes, stream, nullptr, nullptr, 0)); +} + +extern "C" hipError_t hipExtLaunchKernel(const void* hostFunction, + dim3 gridDim, + dim3 blockDim, + void** args, + size_t sharedMemBytes, + hipStream_t stream, + hipEvent_t startEvent, + hipEvent_t stopEvent, + int flags) +{ + HIP_INIT_API(NONE, hostFunction, gridDim, blockDim, args, sharedMemBytes, stream); + HIP_RETURN(ihipLaunchKernel(hostFunction, gridDim, blockDim, args, sharedMemBytes, stream, startEvent, stopEvent, flags)); +} + hipError_t hipLaunchCooperativeKernel(const void* f, dim3 gridDim, dim3 blockDim, void **kernelParams, uint32_t sharedMemBytes, hipStream_t hStream) diff --git a/rocclr/hip_platform.cpp b/rocclr/hip_platform.cpp index d23197bd9e..55bc3e8e7b 100755 --- a/rocclr/hip_platform.cpp +++ b/rocclr/hip_platform.cpp @@ -62,6 +62,14 @@ hipError_t hipModuleGetGlobal(hipDeviceptr_t* dptr, size_t* bytes, hipError_t ihipCreateGlobalVarObj(const char* name, hipModule_t hmod, amd::Memory** amd_mem_obj, hipDeviceptr_t* dptr, size_t* bytes); +extern hipError_t ihipModuleLaunchKernel(hipFunction_t f, + uint32_t gridDimX, uint32_t gridDimY, uint32_t gridDimZ, + uint32_t blockDimX, uint32_t blockDimY, uint32_t blockDimZ, + uint32_t sharedMemBytes, hipStream_t hStream, + void **kernelParams, void **extra, + hipEvent_t startEvent, hipEvent_t stopEvent, uint32_t flags = 0, + uint32_t params = 0, uint32_t gridId = 0, uint32_t numGrids = 0, + uint64_t prevGridSum = 0, uint64_t allGridSum = 0, uint32_t firstDevice = 0); static bool isCompatibleCodeObject(const std::string& codeobj_target_id, const char* device_name) { // Workaround for device name mismatch. @@ -1339,16 +1347,16 @@ void hipLaunchCooperativeKernelGGLImpl( #endif // defined(ATI_OS_LINUX) -extern "C" hipError_t hipLaunchKernel(const void *hostFunction, - dim3 gridDim, - dim3 blockDim, - void** args, - size_t sharedMemBytes, - hipStream_t stream) +hipError_t ihipLaunchKernel(const void* hostFunction, + dim3 gridDim, + dim3 blockDim, + void** args, + size_t sharedMemBytes, + hipStream_t stream, + hipEvent_t startEvent, + hipEvent_t stopEvent, + int flags) { - HIP_INIT_API(hipLaunchKernel, hostFunction, gridDim, blockDim, args, sharedMemBytes, - stream); - hip::Stream* s = reinterpret_cast(stream); int deviceId = (s != nullptr)? s->DeviceId() : ihipGetDevice(); if (deviceId == -1) { @@ -1368,10 +1376,10 @@ extern "C" hipError_t hipLaunchKernel(const void *hostFunction, HIP_RETURN(hipErrorInvalidDeviceFunction); #endif } - - HIP_RETURN(hipModuleLaunchKernel(func, gridDim.x, gridDim.y, gridDim.z, - blockDim.x, blockDim.y, blockDim.z, - sharedMemBytes, stream, args, nullptr)); + HIP_RETURN(ihipModuleLaunchKernel(func, (gridDim.x * blockDim.x), (gridDim.y * blockDim.y), + (gridDim.z * blockDim.z), blockDim.x, blockDim.y, blockDim.z, + sharedMemBytes, stream, args, nullptr, startEvent, stopEvent, + flags)); } // conversion routines between float and half precision diff --git a/src/hip_module.cpp b/src/hip_module.cpp index 0f608d9843..c2ecff3366 100644 --- a/src/hip_module.cpp +++ b/src/hip_module.cpp @@ -1709,12 +1709,33 @@ hipError_t hipLaunchKernel( const void* func_addr, dim3 numBlocks, dim3 dimBlocks, void** args, size_t sharedMemBytes, hipStream_t stream) { - HIP_INIT_API(hipLaunchKernel,func_addr,numBlocks,dimBlocks,args,sharedMemBytes,stream); + HIP_INIT_API(hipLaunchKernel,func_addr,numBlocks,dimBlocks,args,sharedMemBytes,stream); - hipFunction_t kd = hip_impl::get_program_state().kernel_descriptor((std::uintptr_t)func_addr, + hipFunction_t kd = hip_impl::get_program_state().kernel_descriptor((std::uintptr_t)func_addr, hip_impl::target_agent(stream)); - return hipModuleLaunchKernel(kd, numBlocks.x, numBlocks.y, numBlocks.z, + return hipModuleLaunchKernel(kd, numBlocks.x, numBlocks.y, numBlocks.z, dimBlocks.x, dimBlocks.y, dimBlocks.z, sharedMemBytes, stream, args, nullptr); } + +hipError_t hipExtLaunchKernel(const void* function, dim3 numBlocks, dim3 dimBlocks, void** args, + size_t sharedMemBytes, hipStream_t stream, hipEvent_t startEvent, + hipEvent_t stopEvent, int flags) { + HIP_INIT_API(hipExtLaunchKernel,function,numBlocks,dimBlocks,args,sharedMemBytes,stream,startEvent,stopEvent,flags); + + hipFunction_t kd = hip_impl::get_program_state().kernel_descriptor((std::uintptr_t)function, + hip_impl::target_agent(stream)); + + uint32_t globalWorkSizeX = numBlocks.x * dimBlocks.x; + uint32_t globalWorkSizeY = numBlocks.y * dimBlocks.y; + uint32_t globalWorkSizeZ = numBlocks.z * dimBlocks.z; + if (globalWorkSizeX > UINT32_MAX || globalWorkSizeY > UINT32_MAX || + globalWorkSizeZ > UINT32_MAX) { + return hipErrorInvalidConfiguration; + } + + return ihipLogStatus(ihipModuleLaunchKernel( + tls, kd, globalWorkSizeX, globalWorkSizeY, globalWorkSizeZ, dimBlocks.x, dimBlocks.y, + dimBlocks.z, sharedMemBytes, stream, args, nullptr, startEvent, stopEvent, flags)); +} diff --git a/tests/src/deviceLib/hipLaunchKernelFunc.cpp b/tests/src/deviceLib/hipLaunchKernelFunc.cpp new file mode 100644 index 0000000000..75dbc81da7 --- /dev/null +++ b/tests/src/deviceLib/hipLaunchKernelFunc.cpp @@ -0,0 +1,193 @@ +/* +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 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 ../test_common.cpp HCC_OPTIONS -Xclang -fallow-half-arguments-and-returns CLANG_OPTIONS -Xclang -fallow-half-arguments-and-returns EXCLUDE_HIP_PLATFORM nvcc + * TEST: %t + * HIT_END + */ + +#define HIP_TEMPLATE_KERNEL_LAUNCH +#include "hip/hip_runtime.h" +#include "test_common.h" + +__global__ void kernel_abs_int64(long long* input, long long* output) { + int tx = threadIdx.x; + output[tx] = abs(input[tx]); +} + +__global__ void kernel_lgamma_double(double* input, double* output) { + int tx = threadIdx.x; + output[tx] = lgamma(input[tx]); +} + +#define CHECK_LGAMMA_DOUBLE(IN, OUT, EXP) \ + { \ + if (OUT != EXP) { \ + failed("check_abs_int64 failed on %f (output = %f, expected = %fd)\n", IN, OUT, EXP); \ + } \ + } + +#define CHECK_ABS_INT64(IN, OUT, EXP) \ + { \ + if (OUT != EXP) { \ + failed("check_abs_int64 failed on %lld (output = %lld, expected = %lld)\n", IN, OUT, \ + EXP); \ + } \ + } + +void check_lgamma_double() { + using datatype_t = double; + + const int NUM_INPUTS = 8; + auto memsize = NUM_INPUTS * sizeof(datatype_t); + + // allocate memories + datatype_t* inputCPU = (datatype_t*)malloc(memsize); + datatype_t* outputCPU = (datatype_t*)malloc(memsize); + datatype_t* inputGPU = nullptr; + hipMalloc((void**)&inputGPU, memsize); + datatype_t* outputGPU = nullptr; + hipMalloc((void**)&outputGPU, memsize); + + // populate input + for (int i = 0; i < NUM_INPUTS; i++) { + inputCPU[i] = -3.5 + i; + } + + // copy inputs to device + hipMemcpy(inputGPU, inputCPU, memsize, hipMemcpyHostToDevice); + + // launch kernel + hipLaunchKernelGGL(kernel_lgamma_double, dim3(1), dim3(NUM_INPUTS), 0, 0, inputGPU, outputGPU); + + // copy outputs from device + hipMemcpy(outputCPU, outputGPU, memsize, hipMemcpyDeviceToHost); + + // check outputs + for (int i = 0; i < NUM_INPUTS; i++) { + CHECK_LGAMMA_DOUBLE(inputCPU[i], outputCPU[i], lgamma(inputCPU[i])); + } + + // free memories + hipFree(inputGPU); + hipFree(outputGPU); + free(inputCPU); + free(outputCPU); + + // done + return; +} + + +void check_abs_int64() { + using datatype_t = long long; + + const int NUM_INPUTS = 8; + auto memsize = NUM_INPUTS * sizeof(datatype_t); + + // allocate memories + datatype_t* inputCPU = (datatype_t*)malloc(memsize); + datatype_t* outputCPU = (datatype_t*)malloc(memsize); + datatype_t* inputGPU = nullptr; + hipMalloc((void**)&inputGPU, memsize); + datatype_t* outputGPU = nullptr; + hipMalloc((void**)&outputGPU, memsize); + + // populate input + inputCPU[0] = -81985529216486895ll; + inputCPU[1] = 81985529216486895ll; + inputCPU[2] = -1250999896491ll; + inputCPU[3] = 1250999896491ll; + inputCPU[4] = -19088743ll; + inputCPU[5] = 19088743ll; + inputCPU[6] = -291ll; + inputCPU[7] = 291ll; + + // copy inputs to device + hipMemcpy(inputGPU, inputCPU, memsize, hipMemcpyHostToDevice); + + // launch kernel + hipLaunchKernelGGL(kernel_abs_int64, dim3(1), dim3(NUM_INPUTS), 0, 0, inputGPU, outputGPU); + + // copy outputs from device + hipMemcpy(outputCPU, outputGPU, memsize, hipMemcpyDeviceToHost); + + // check outputs + CHECK_ABS_INT64(inputCPU[0], outputCPU[0], outputCPU[1]); + CHECK_ABS_INT64(inputCPU[1], outputCPU[1], outputCPU[1]); + CHECK_ABS_INT64(inputCPU[2], outputCPU[2], outputCPU[3]); + CHECK_ABS_INT64(inputCPU[3], outputCPU[3], outputCPU[3]); + CHECK_ABS_INT64(inputCPU[4], outputCPU[4], outputCPU[5]); + CHECK_ABS_INT64(inputCPU[5], outputCPU[5], outputCPU[5]); + CHECK_ABS_INT64(inputCPU[6], outputCPU[6], outputCPU[7]); + CHECK_ABS_INT64(inputCPU[7], outputCPU[7], outputCPU[7]); + + // free memories + hipFree(inputGPU); + hipFree(outputGPU); + free(inputCPU); + free(outputCPU); + + // done + return; +} + + +template +__global__ void kernel_simple(F f, T* out) { + *out = f(); +} + +template +void check_simple(F f, T expected, const char* file, unsigned line) { + auto memsize = sizeof(T); + T* outputCPU = (T*)malloc(memsize); + T* outputGPU = nullptr; + hipMalloc((void**)&outputGPU, memsize); + hipLaunchKernelGGL(kernel_simple, 1, 1, 0, 0, f, outputGPU); + hipMemcpy(outputCPU, outputGPU, memsize, hipMemcpyDeviceToHost); + if (*outputCPU != expected) { + failed("%s line %u : check failed (output = %lf, expected = %lf)\n", file, line, + (double)(*outputCPU), (double)expected); + } + hipFree(outputGPU); + free(outputCPU); +} +#define CHECK_SIMPLE(lambda, expected) check_simple(lambda, expected, __FILE__, __LINE__); + +void test_fp16() { + CHECK_SIMPLE([] __device__() { return max<__fp16>(1.0f, 2.0f); }, 2.0f); + CHECK_SIMPLE([] __device__() { return min<__fp16>(1.0f, 2.0f); }, 1.0f); +} + +int main(int argc, char* argv[]) { + HipTest::parseStandardArguments(argc, argv, true); + + check_abs_int64(); + + // check_lgamma_double(); + + test_fp16(); + + passed(); +} diff --git a/tests/src/kernel/hipExtLaunchKernelGGL.cpp b/tests/src/kernel/hipExtLaunchKernelGGL.cpp index 39c660322b..12b96578de 100644 --- a/tests/src/kernel/hipExtLaunchKernelGGL.cpp +++ b/tests/src/kernel/hipExtLaunchKernelGGL.cpp @@ -28,9 +28,52 @@ THE SOFTWARE. #include "hip/hip_ext.h" #include "test_common.h" +struct _t { + double _a, _b, _c, _d, _e, _f, _g, _h, _i, _j; +}; + +typedef struct _t _T; + +__global__ void sKernel(_T s, double *a) { + *a = s._a + s._b + s._c + s._d + s._e + s._f + s._g + s._h + s._i + s._j; +} + +__global__ void mKernel(char f, short a, int b, double c, short d, int e, double* res) { + *res = a + b + c + d + e + f; +} + +void testMixData() { + double m = 0; + double *d_m; + HIPCHECK(hipMalloc(&d_m, sizeof(double))); + int a = 1, e = 10; + short b = 2, d = 4; + double c = 3.0; + char ff = 10; + hipExtLaunchKernelGGL(mKernel, 1, 1, 0, 0, nullptr, nullptr, 0, ff, b, a, c, d, e, d_m); + HIPCHECK(hipMemcpy(&m, d_m, sizeof(double), hipMemcpyDeviceToHost)); + if (m != 30.0) { + std::cout << "M is:: " << m << std::endl; + failed("Mismatch"); + } + hipFree(d_m); +} +void testStruct() { + double m = 0; + double *d_m; + HIPCHECK(hipMalloc(&d_m, sizeof(double))); + _T s{1, 2, 3, 4, 5, 6, 7, 8, 9, 10}; + hipExtLaunchKernelGGL(sKernel, 1, 1, 0, 0, nullptr, nullptr, 0, s, d_m); + HIPCHECK(hipMemcpy(&m, d_m, sizeof(double), hipMemcpyDeviceToHost)); + if (m != 55.0) { + std::cout << "M is:: " << m << std::endl; + failed("Mismatch"); + } + hipFree(d_m); +} + void test(size_t N) { size_t Nbytes = N * sizeof(int); -#if defined(__HIP_PLATFORM_HCC__) && GENERIC_GRID_LAUNCH == 1 && defined(__HCC__) int *A_d, *B_d, *C_d; int *A_h, *B_h, *C_h; @@ -51,13 +94,13 @@ void test(size_t N) { HIPCHECK(hipDeviceSynchronize()); HipTest::checkVectorADD(A_h, B_h, C_h, N); -#endif } int main(int argc, char* argv[]) { HipTest::parseStandardArguments(argc, argv, true); test(N); - + testStruct(); + testMixData(); passed(); } From f20748854e08a056ecd34e8e0fb598960b8c573d Mon Sep 17 00:00:00 2001 From: Siu Chi Chan Date: Thu, 21 May 2020 15:11:41 -0400 Subject: [PATCH 02/33] update device library path fix device lib directory add missing --hip-link switch for link phase Change-Id: I4b2eeb32648ca3cec72ec1f4e3381ce1fc0a90a5 --- bin/hipcc | 9 ++++++++- hip-config.cmake.in | 18 +++++++++++++----- 2 files changed, 21 insertions(+), 6 deletions(-) diff --git a/bin/hipcc b/bin/hipcc index 1c858b070c..3f129e6e1a 100755 --- a/bin/hipcc +++ b/bin/hipcc @@ -168,7 +168,14 @@ if (defined $HIP_COMPILER and $HIP_PLATFORM eq "hcc" and $HIP_COMPILER eq "clang $HIP_CLANG_PATH = "$ROCM_PATH/llvm/bin"; } if (!defined $DEVICE_LIB_PATH) { - $DEVICE_LIB_PATH = "$ROCM_PATH/lib"; + if (-e "$ROCM_PATH/amdgcn/bitcode") { + $DEVICE_LIB_PATH = "$ROCM_PATH/amdgcn/bitcode"; + } + else { + # This path is to support an older build of the device library + # TODO: To be removed in the future. + $DEVICE_LIB_PATH = "$ROCM_PATH/lib"; + } } } diff --git a/hip-config.cmake.in b/hip-config.cmake.in index 5a67c62383..623f7b50d8 100644 --- a/hip-config.cmake.in +++ b/hip-config.cmake.in @@ -142,18 +142,26 @@ else() endif() if(HIP_COMPILER STREQUAL "clang") - set_property(TARGET hip::device APPEND PROPERTY - INTERFACE_COMPILE_OPTIONS -x hip --hip-device-lib-path=${AMD_DEVICE_LIBS_PREFIX}/lib - ) - if (HIP_CXX_COMPILER MATCHES ".*clang\\+\\+") set_property(TARGET hip::device APPEND PROPERTY INTERFACE_COMPILE_OPTIONS -mllvm -amdgpu-early-inline-all=true -mllvm -amdgpu-function-calls=false ) endif() + if (EXISTS ${AMD_DEVICE_LIBS_PREFIX}/amdgcn/bitcode) + set_property(TARGET hip::device APPEND PROPERTY + INTERFACE_COMPILE_OPTIONS -x hip --hip-device-lib-path=${AMD_DEVICE_LIBS_PREFIX}/amdgcn/bitcode + ) + else() + # This path is to support an older build of the device library + # TODO: To be removed in the future. + set_property(TARGET hip::device APPEND PROPERTY + INTERFACE_COMPILE_OPTIONS -x hip --hip-device-lib-path=${AMD_DEVICE_LIBS_PREFIX}/lib + ) + endif() + set_property(TARGET hip::device APPEND PROPERTY - INTERFACE_LINK_LIBRARIES --hip-device-lib-path=${AMD_DEVICE_LIBS_PREFIX}/lib --hip-link + INTERFACE_LINK_LIBRARIES --hip-link ) set_property(TARGET hip::device APPEND PROPERTY From 56392b4f8a1f47200c0c9ea38f4ec0ca60d9c6fa Mon Sep 17 00:00:00 2001 From: Aaron En Ye Shi Date: Fri, 29 May 2020 21:45:34 +0000 Subject: [PATCH 03/33] Add compiler-rt library for __fp16 and _Float16 Similar to HCC, link with compiler-rt to support __fp16 and _Float16 type conversions in ONNX models. This should resolve SWDEV-238491. Change-Id: Iad8dcff568831719f501f562a04023326ae8036c --- bin/hipcc | 2 + hip-config.cmake.in | 14 ++++++ tests/src/Functional/host/hipFloat16.cpp | 59 ++++++++++++++++++++++++ 3 files changed, 75 insertions(+) create mode 100644 tests/src/Functional/host/hipFloat16.cpp diff --git a/bin/hipcc b/bin/hipcc index 3f129e6e1a..c1e144300b 100755 --- a/bin/hipcc +++ b/bin/hipcc @@ -840,6 +840,8 @@ if ($HIP_PLATFORM eq "hcc" and $HIP_COMPILER eq "clang") { } else { $toolArgs .= " -Wl,--enable-new-dtags -Wl,--rpath=$HIP_LIB_PATH:$ROCM_PATH/lib -lhip_hcc "; } + # To support __fp16 and _Float16, explicitly link with compiler-rt + $toolArgs .= " -L$HIP_CLANG_PATH/../lib/clang/$HIP_CLANG_VERSION/lib/linux -lclang_rt.builtins-x86_64 " } } diff --git a/hip-config.cmake.in b/hip-config.cmake.in index 623f7b50d8..487d05fbab 100644 --- a/hip-config.cmake.in +++ b/hip-config.cmake.in @@ -196,6 +196,20 @@ if(HIP_COMPILER STREQUAL "clang") message("clang compiler doesn't support parallel jobs") endif() endif() + + # Add support for __fp16 and _Float16, explicitly link with compiler-rt + set_property(TARGET hip::host APPEND PROPERTY + INTERFACE_COMPILE_OPTIONS -L${HIP_CLANG_INCLUDE_PATH}/../lib/linux -lclang_rt.builtins-x86_64 + ) + set_property(TARGET hip::host APPEND PROPERTY + INTERFACE_LINK_LIBRARIES -L${HIP_CLANG_INCLUDE_PATH}/../lib/linux -lclang_rt.builtins-x86_64 + ) + set_property(TARGET hip::device APPEND PROPERTY + INTERFACE_COMPILE_OPTIONS -L${HIP_CLANG_INCLUDE_PATH}/../lib/linux -lclang_rt.builtins-x86_64 + ) + set_property(TARGET hip::device APPEND PROPERTY + INTERFACE_LINK_LIBRARIES -L${HIP_CLANG_INCLUDE_PATH}/../lib/linux -lclang_rt.builtins-x86_64 + ) endif() set( hip_LIBRARIES hip::host hip::device) diff --git a/tests/src/Functional/host/hipFloat16.cpp b/tests/src/Functional/host/hipFloat16.cpp new file mode 100644 index 0000000000..62cf166955 --- /dev/null +++ b/tests/src/Functional/host/hipFloat16.cpp @@ -0,0 +1,59 @@ +/* +Copyright (c) 2015-2020 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 ../../test_common.cpp EXCLUDE_HIP_PLATFORM nvcc + * TEST: %t 1.2 2.3 + * HIT_END + */ + +#include +#include +#include "test_common.h" + +int main(int argc, char* argv[]) { + // Testing that the compiler supports host _Float16 conversions + float init_value = atof(argv[1]); + _Float16 value_float16 = static_cast<_Float16>(init_value); + float result_value = static_cast(value_float16); + + if(std::abs(result_value - init_value) >= 0.01){ + printf("init: %f\n", init_value); + printf("result: %f\n", result_value); + printf("diff: %f\n", std::abs(result_value - init_value)); + failed("Failed host _Float16 test."); + } + + // Testing that the compiler supports host __fp16 conversions + init_value = atof(argv[2]); + __fp16 value_fp16 = static_cast<__fp16>(init_value); + result_value = static_cast(value_fp16); + + if(std::abs(result_value - init_value) >= 0.01){ + printf("init: %f\n", init_value); + printf("result: %f\n", result_value); + printf("diff: %f\n", std::abs(result_value - init_value)); + failed("Failed host __fp16 test."); + } + + passed(); +} From cad3f805c0775874306be364c90ae47253db6876 Mon Sep 17 00:00:00 2001 From: Evgeny Date: Wed, 3 Jun 2020 01:30:44 -0500 Subject: [PATCH 04/33] adding hipGetStreamDeviceId() profiling API Change-Id: I5ccf88ddac123260d7c17defefcf20ff3b2504e2 --- include/hip/hcc_detail/hip_runtime_api.h | 1 + rocclr/hip_hcc.map.in | 1 + rocclr/hip_intercept.cpp | 13 +++++++++++-- 3 files changed, 13 insertions(+), 2 deletions(-) diff --git a/include/hip/hcc_detail/hip_runtime_api.h b/include/hip/hcc_detail/hip_runtime_api.h index e374707cb3..67866a8d4b 100644 --- a/include/hip/hcc_detail/hip_runtime_api.h +++ b/include/hip/hcc_detail/hip_runtime_api.h @@ -3662,6 +3662,7 @@ hipError_t hipRemoveActivityCallback(uint32_t id); const char* hipApiName(uint32_t id); const char* hipKernelNameRef(const hipFunction_t f); const char* hipKernelNameRefByPtr(const void* hostFunction, hipStream_t stream); +int hipGetStreamDeviceId(hipStream_t stream); #ifdef __cplusplus } /* extern "C" */ #endif diff --git a/rocclr/hip_hcc.map.in b/rocclr/hip_hcc.map.in index 42c5335f35..5bc2d018f2 100755 --- a/rocclr/hip_hcc.map.in +++ b/rocclr/hip_hcc.map.in @@ -183,6 +183,7 @@ global: hipApiName; hipKernelNameRef; hipKernelNameRefByPtr; + hipGetStreamDeviceId; hipProfilerStart; hipProfilerStop; hiprtcCompileProgram; diff --git a/rocclr/hip_intercept.cpp b/rocclr/hip_intercept.cpp index aa61d6bc38..c1b7f53534 100644 --- a/rocclr/hip_intercept.cpp +++ b/rocclr/hip_intercept.cpp @@ -27,10 +27,19 @@ api_callbacks_table_t callbacks_table; extern const std::string& FunctionName(const hipFunction_t f); + const char* hipKernelNameRef(const hipFunction_t f) { return FunctionName(f).c_str(); } -const char* hipKernelNameRefByPtr(const void *hostFunction, hipStream_t stream) { + +int hipGetStreamDeviceId(hipStream_t stream) { hip::Stream* s = reinterpret_cast(stream); - int deviceId = (s != nullptr)? s->DeviceId() : ihipGetDevice(); + return (s != nullptr)? s->DeviceId() : ihipGetDevice(); +} + +const char* hipKernelNameRefByPtr(const void* hostFunction, hipStream_t stream) { + if (hostFunction == NULL) { + return NULL; + } + int deviceId = hipGetStreamDeviceId(stream); if (deviceId == -1) { DevLogPrintfError("Wrong Device Id: %d \n", deviceId); return NULL; From 784ca6f43cdf2eb2dbdaa06af095c069c1264923 Mon Sep 17 00:00:00 2001 From: Siu Chi Chan Date: Thu, 12 Mar 2020 00:00:13 -0400 Subject: [PATCH 05/33] add constexpr constructor for vector types Change-Id: I45bb0537d6a24ee50b548c2fd8b4f20518764813 --- include/hip/hcc_detail/hip_vector_types.h | 311 ++++++++++++++----- tests/src/deviceLib/hipVectorTypes.cpp | 45 +-- tests/src/deviceLib/hipVectorTypesDevice.cpp | 137 ++++++-- 3 files changed, 375 insertions(+), 118 deletions(-) diff --git a/include/hip/hcc_detail/hip_vector_types.h b/include/hip/hcc_detail/hip_vector_types.h index 19259a3657..69525c5684 100644 --- a/include/hip/hcc_detail/hip_vector_types.h +++ b/include/hip/hcc_detail/hip_vector_types.h @@ -312,6 +312,21 @@ THE SOFTWARE. using value_type = T; + __host__ __device__ + HIP_vector_base() = default; + __host__ __device__ + explicit + constexpr + HIP_vector_base(T x) noexcept : data{x} {} + __host__ __device__ + constexpr + HIP_vector_base(const HIP_vector_base&) = default; + __host__ __device__ + constexpr + HIP_vector_base(HIP_vector_base&&) = default; + __host__ __device__ + ~HIP_vector_base() = default; + __host__ __device__ HIP_vector_base& operator=(const HIP_vector_base& x) noexcept { #if __has_attribute(ext_vector_type) @@ -347,6 +362,24 @@ THE SOFTWARE. using value_type = T; + __host__ __device__ + HIP_vector_base() = default; + __host__ __device__ + explicit + constexpr + HIP_vector_base(T x) noexcept : data{x, x} {} + __host__ __device__ + constexpr + HIP_vector_base(T x, T y) noexcept : data{x, y} {} + __host__ __device__ + constexpr + HIP_vector_base(const HIP_vector_base&) = default; + __host__ __device__ + constexpr + HIP_vector_base(HIP_vector_base&&) = default; + __host__ __device__ + ~HIP_vector_base() = default; + __host__ __device__ HIP_vector_base& operator=(const HIP_vector_base& x) noexcept { #if __has_attribute(ext_vector_type) @@ -366,8 +399,8 @@ THE SOFTWARE. T d[3]; __host__ __device__ - constexpr Native_vec_() = default; + __host__ __device__ explicit constexpr @@ -514,6 +547,29 @@ THE SOFTWARE. }; using value_type = T; + + __host__ __device__ + HIP_vector_base() = default; + __host__ __device__ + explicit + constexpr + HIP_vector_base(T x) noexcept : data{x, x, x} {} + __host__ __device__ + constexpr + HIP_vector_base(T x, T y, T z) noexcept : data{x, y, z} {} + __host__ __device__ + constexpr + HIP_vector_base(const HIP_vector_base&) = default; + __host__ __device__ + constexpr + HIP_vector_base(HIP_vector_base&&) = default; + __host__ __device__ + ~HIP_vector_base() = default; + + __host__ __device__ + HIP_vector_base& operator=(const HIP_vector_base&) = default; + __host__ __device__ + HIP_vector_base& operator=(HIP_vector_base&&) = default; }; template @@ -538,11 +594,29 @@ THE SOFTWARE. hip_impl::Scalar_accessor y; hip_impl::Scalar_accessor z; hip_impl::Scalar_accessor w; -#endif +#endif }; using value_type = T; + __host__ __device__ + HIP_vector_base() = default; + __host__ __device__ + explicit + constexpr + HIP_vector_base(T x) noexcept : data{x, x, x, x} {} + __host__ __device__ + constexpr + HIP_vector_base(T x, T y, T z, T w) noexcept : data{x, y, z, w} {} + __host__ __device__ + constexpr + HIP_vector_base(const HIP_vector_base&) = default; + __host__ __device__ + constexpr + HIP_vector_base(HIP_vector_base&&) = default; + __host__ __device__ + ~HIP_vector_base() = default; + __host__ __device__ HIP_vector_base& operator=(const HIP_vector_base& x) noexcept { #if __has_attribute(ext_vector_type) @@ -563,49 +637,48 @@ THE SOFTWARE. using HIP_vector_base::data; using typename HIP_vector_base::Native_vec_; - inline __host__ __device__ + __host__ __device__ HIP_vector_type() = default; template< typename U, typename std::enable_if< std::is_convertible{}>::type* = nullptr> - explicit inline __host__ __device__ + __host__ __device__ + explicit + constexpr HIP_vector_type(U x) noexcept - { - for (auto i = 0u; i != rank; ++i) data[i] = x; - } + : HIP_vector_base{static_cast(x)} + {} template< // TODO: constrain based on type as well. typename... Us, typename std::enable_if< (rank > 1) && sizeof...(Us) == rank>::type* = nullptr> - inline __host__ __device__ + __host__ __device__ + constexpr HIP_vector_type(Us... xs) noexcept - { - #if __has_attribute(ext_vector_type) - new (&data) Native_vec_{static_cast(xs)...}; - #else - new (&data) std::array{static_cast(xs)...}; - #endif - } - inline __host__ __device__ + : HIP_vector_base{static_cast(xs)...} + {} + __host__ __device__ + constexpr HIP_vector_type(const HIP_vector_type&) = default; - inline __host__ __device__ + __host__ __device__ + constexpr HIP_vector_type(HIP_vector_type&&) = default; - inline __host__ __device__ + __host__ __device__ ~HIP_vector_type() = default; - inline __host__ __device__ + __host__ __device__ HIP_vector_type& operator=(const HIP_vector_type&) = default; - inline __host__ __device__ + __host__ __device__ HIP_vector_type& operator=(HIP_vector_type&&) = default; // Operators - inline __host__ __device__ + __host__ __device__ HIP_vector_type& operator++() noexcept { return *this += HIP_vector_type{1}; } - inline __host__ __device__ + __host__ __device__ HIP_vector_type operator++(int) noexcept { auto tmp(*this); @@ -613,12 +686,12 @@ THE SOFTWARE. return tmp; } - inline __host__ __device__ + __host__ __device__ HIP_vector_type& operator--() noexcept { return *this -= HIP_vector_type{1}; } - inline __host__ __device__ + __host__ __device__ HIP_vector_type operator--(int) noexcept { auto tmp(*this); @@ -626,7 +699,7 @@ THE SOFTWARE. return tmp; } - inline __host__ __device__ + __host__ __device__ HIP_vector_type& operator+=(const HIP_vector_type& x) noexcept { data += x.data; @@ -636,13 +709,13 @@ THE SOFTWARE. typename U, typename std::enable_if< std::is_convertible{}>::type* = nullptr> - inline __host__ __device__ + __host__ __device__ HIP_vector_type& operator+=(U x) noexcept { return *this += HIP_vector_type{x}; } - inline __host__ __device__ + __host__ __device__ HIP_vector_type& operator-=(const HIP_vector_type& x) noexcept { data -= x.data; @@ -652,13 +725,13 @@ THE SOFTWARE. typename U, typename std::enable_if< std::is_convertible{}>::type* = nullptr> - inline __host__ __device__ + __host__ __device__ HIP_vector_type& operator-=(U x) noexcept { return *this -= HIP_vector_type{x}; } - inline __host__ __device__ + __host__ __device__ HIP_vector_type& operator*=(const HIP_vector_type& x) noexcept { data *= x.data; @@ -668,13 +741,13 @@ THE SOFTWARE. typename U, typename std::enable_if< std::is_convertible{}>::type* = nullptr> - inline __host__ __device__ + __host__ __device__ HIP_vector_type& operator*=(U x) noexcept { return *this *= HIP_vector_type{x}; } - inline __host__ __device__ + __host__ __device__ HIP_vector_type& operator/=(const HIP_vector_type& x) noexcept { data /= x.data; @@ -684,7 +757,7 @@ THE SOFTWARE. typename U, typename std::enable_if< std::is_convertible{}>::type* = nullptr> - inline __host__ __device__ + __host__ __device__ HIP_vector_type& operator/=(U x) noexcept { return *this /= HIP_vector_type{x}; @@ -693,7 +766,7 @@ THE SOFTWARE. template< typename U = T, typename std::enable_if{}>::type* = nullptr> - inline __host__ __device__ + __host__ __device__ HIP_vector_type operator-() const noexcept { auto tmp(*this); @@ -704,7 +777,7 @@ THE SOFTWARE. template< typename U = T, typename std::enable_if{}>::type* = nullptr> - inline __host__ __device__ + __host__ __device__ HIP_vector_type operator~() const noexcept { HIP_vector_type r{*this}; @@ -715,7 +788,7 @@ THE SOFTWARE. template< typename U = T, typename std::enable_if{}>::type* = nullptr> - inline __host__ __device__ + __host__ __device__ HIP_vector_type& operator%=(const HIP_vector_type& x) noexcept { data %= x.data; @@ -725,7 +798,7 @@ THE SOFTWARE. template< typename U = T, typename std::enable_if{}>::type* = nullptr> - inline __host__ __device__ + __host__ __device__ HIP_vector_type& operator^=(const HIP_vector_type& x) noexcept { data ^= x.data; @@ -735,7 +808,7 @@ THE SOFTWARE. template< typename U = T, typename std::enable_if{}>::type* = nullptr> - inline __host__ __device__ + __host__ __device__ HIP_vector_type& operator|=(const HIP_vector_type& x) noexcept { data |= x.data; @@ -745,7 +818,7 @@ THE SOFTWARE. template< typename U = T, typename std::enable_if{}>::type* = nullptr> - inline __host__ __device__ + __host__ __device__ HIP_vector_type& operator&=(const HIP_vector_type& x) noexcept { data &= x.data; @@ -755,7 +828,7 @@ THE SOFTWARE. template< typename U = T, typename std::enable_if{}>::type* = nullptr> - inline __host__ __device__ + __host__ __device__ HIP_vector_type& operator>>=(const HIP_vector_type& x) noexcept { data >>= x.data; @@ -765,7 +838,7 @@ THE SOFTWARE. template< typename U = T, typename std::enable_if{}>::type* = nullptr> - inline __host__ __device__ + __host__ __device__ HIP_vector_type& operator<<=(const HIP_vector_type& x) noexcept { data <<= x.data; @@ -774,21 +847,27 @@ THE SOFTWARE. }; template - inline __host__ __device__ + __host__ __device__ + inline + constexpr HIP_vector_type operator+( const HIP_vector_type& x, const HIP_vector_type& y) noexcept { return HIP_vector_type{x} += y; } template - inline __host__ __device__ + __host__ __device__ + inline + constexpr HIP_vector_type operator+( const HIP_vector_type& x, U y) noexcept { return HIP_vector_type{x} += HIP_vector_type{y}; } template - inline __host__ __device__ + __host__ __device__ + inline + constexpr HIP_vector_type operator+( U x, const HIP_vector_type& y) noexcept { @@ -796,21 +875,27 @@ THE SOFTWARE. } template - inline __host__ __device__ + __host__ __device__ + inline + constexpr HIP_vector_type operator-( const HIP_vector_type& x, const HIP_vector_type& y) noexcept { return HIP_vector_type{x} -= y; } template - inline __host__ __device__ + __host__ __device__ + inline + constexpr HIP_vector_type operator-( const HIP_vector_type& x, U y) noexcept { return HIP_vector_type{x} -= HIP_vector_type{y}; } template - inline __host__ __device__ + __host__ __device__ + inline + constexpr HIP_vector_type operator-( U x, const HIP_vector_type& y) noexcept { @@ -818,21 +903,27 @@ THE SOFTWARE. } template - inline __host__ __device__ + __host__ __device__ + inline + constexpr HIP_vector_type operator*( const HIP_vector_type& x, const HIP_vector_type& y) noexcept { return HIP_vector_type{x} *= y; } template - inline __host__ __device__ + __host__ __device__ + inline + constexpr HIP_vector_type operator*( const HIP_vector_type& x, U y) noexcept { return HIP_vector_type{x} *= HIP_vector_type{y}; } template - inline __host__ __device__ + __host__ __device__ + inline + constexpr HIP_vector_type operator*( U x, const HIP_vector_type& y) noexcept { @@ -840,64 +931,90 @@ THE SOFTWARE. } template - inline __host__ __device__ + __host__ __device__ + inline + constexpr HIP_vector_type operator/( const HIP_vector_type& x, const HIP_vector_type& y) noexcept { return HIP_vector_type{x} /= y; } template - inline __host__ __device__ + __host__ __device__ + inline + constexpr HIP_vector_type operator/( const HIP_vector_type& x, U y) noexcept { return HIP_vector_type{x} /= HIP_vector_type{y}; } template - inline __host__ __device__ + __host__ __device__ + inline + constexpr HIP_vector_type operator/( U x, const HIP_vector_type& y) noexcept { return HIP_vector_type{x} /= y; } + template + __host__ __device__ + inline + constexpr + bool _hip_any_zero(const V& x, int n) noexcept + { + return + (n == -1) ? true : ((x[n] == 0) ? false : _hip_any_zero(x, n - 1)); + } + template - inline __host__ __device__ + __host__ __device__ + inline + constexpr bool operator==( const HIP_vector_type& x, const HIP_vector_type& y) noexcept { - auto tmp = x.data == y.data; - for (auto i = 0u; i != n; ++i) if (tmp[i] == 0) return false; - return true; + return _hip_any_zero(x.data == y.data, n - 1); } template - inline __host__ __device__ + __host__ __device__ + inline + constexpr bool operator==(const HIP_vector_type& x, U y) noexcept { return x == HIP_vector_type{y}; } template - inline __host__ __device__ + __host__ __device__ + inline + constexpr bool operator==(U x, const HIP_vector_type& y) noexcept { return HIP_vector_type{x} == y; } template - inline __host__ __device__ + __host__ __device__ + inline + constexpr bool operator!=( const HIP_vector_type& x, const HIP_vector_type& y) noexcept { return !(x == y); } template - inline __host__ __device__ + __host__ __device__ + inline + constexpr bool operator!=(const HIP_vector_type& x, U y) noexcept { return !(x == y); } template - inline __host__ __device__ + __host__ __device__ + inline + constexpr bool operator!=(U x, const HIP_vector_type& y) noexcept { return !(x == y); @@ -907,7 +1024,9 @@ THE SOFTWARE. typename T, unsigned int n, typename std::enable_if{}>* = nullptr> - inline __host__ __device__ + __host__ __device__ + inline + constexpr HIP_vector_type operator%( const HIP_vector_type& x, const HIP_vector_type& y) noexcept { @@ -918,7 +1037,9 @@ THE SOFTWARE. unsigned int n, typename U, typename std::enable_if{}>* = nullptr> - inline __host__ __device__ + __host__ __device__ + inline + constexpr HIP_vector_type operator%( const HIP_vector_type& x, U y) noexcept { @@ -929,7 +1050,9 @@ THE SOFTWARE. unsigned int n, typename U, typename std::enable_if{}>* = nullptr> - inline __host__ __device__ + __host__ __device__ + inline + constexpr HIP_vector_type operator%( U x, const HIP_vector_type& y) noexcept { @@ -940,7 +1063,9 @@ THE SOFTWARE. typename T, unsigned int n, typename std::enable_if{}>* = nullptr> - inline __host__ __device__ + __host__ __device__ + inline + constexpr HIP_vector_type operator^( const HIP_vector_type& x, const HIP_vector_type& y) noexcept { @@ -951,7 +1076,9 @@ THE SOFTWARE. unsigned int n, typename U, typename std::enable_if{}>* = nullptr> - inline __host__ __device__ + __host__ __device__ + inline + constexpr HIP_vector_type operator^( const HIP_vector_type& x, U y) noexcept { @@ -962,7 +1089,9 @@ THE SOFTWARE. unsigned int n, typename U, typename std::enable_if{}>* = nullptr> - inline __host__ __device__ + __host__ __device__ + inline + constexpr HIP_vector_type operator^( U x, const HIP_vector_type& y) noexcept { @@ -973,7 +1102,9 @@ THE SOFTWARE. typename T, unsigned int n, typename std::enable_if{}>* = nullptr> - inline __host__ __device__ + __host__ __device__ + inline + constexpr HIP_vector_type operator|( const HIP_vector_type& x, const HIP_vector_type& y) noexcept { @@ -984,7 +1115,9 @@ THE SOFTWARE. unsigned int n, typename U, typename std::enable_if{}>* = nullptr> - inline __host__ __device__ + __host__ __device__ + inline + constexpr HIP_vector_type operator|( const HIP_vector_type& x, U y) noexcept { @@ -995,7 +1128,9 @@ THE SOFTWARE. unsigned int n, typename U, typename std::enable_if{}>* = nullptr> - inline __host__ __device__ + __host__ __device__ + inline + constexpr HIP_vector_type operator|( U x, const HIP_vector_type& y) noexcept { @@ -1006,7 +1141,9 @@ THE SOFTWARE. typename T, unsigned int n, typename std::enable_if{}>* = nullptr> - inline __host__ __device__ + __host__ __device__ + inline + constexpr HIP_vector_type operator&( const HIP_vector_type& x, const HIP_vector_type& y) noexcept { @@ -1017,7 +1154,9 @@ THE SOFTWARE. unsigned int n, typename U, typename std::enable_if{}>* = nullptr> - inline __host__ __device__ + __host__ __device__ + inline + constexpr HIP_vector_type operator&( const HIP_vector_type& x, U y) noexcept { @@ -1028,7 +1167,9 @@ THE SOFTWARE. unsigned int n, typename U, typename std::enable_if{}>* = nullptr> - inline __host__ __device__ + __host__ __device__ + inline + constexpr HIP_vector_type operator&( U x, const HIP_vector_type& y) noexcept { @@ -1039,7 +1180,9 @@ THE SOFTWARE. typename T, unsigned int n, typename std::enable_if{}>* = nullptr> - inline __host__ __device__ + __host__ __device__ + inline + constexpr HIP_vector_type operator>>( const HIP_vector_type& x, const HIP_vector_type& y) noexcept { @@ -1050,7 +1193,9 @@ THE SOFTWARE. unsigned int n, typename U, typename std::enable_if{}>* = nullptr> - inline __host__ __device__ + __host__ __device__ + inline + constexpr HIP_vector_type operator>>( const HIP_vector_type& x, U y) noexcept { @@ -1061,7 +1206,9 @@ THE SOFTWARE. unsigned int n, typename U, typename std::enable_if{}>* = nullptr> - inline __host__ __device__ + __host__ __device__ + inline + constexpr HIP_vector_type operator>>( U x, const HIP_vector_type& y) noexcept { @@ -1072,7 +1219,9 @@ THE SOFTWARE. typename T, unsigned int n, typename std::enable_if{}>* = nullptr> - inline __host__ __device__ + __host__ __device__ + inline + constexpr HIP_vector_type operator<<( const HIP_vector_type& x, const HIP_vector_type& y) noexcept { @@ -1083,7 +1232,9 @@ THE SOFTWARE. unsigned int n, typename U, typename std::enable_if{}>* = nullptr> - inline __host__ __device__ + __host__ __device__ + inline + constexpr HIP_vector_type operator<<( const HIP_vector_type& x, U y) noexcept { @@ -1095,7 +1246,9 @@ THE SOFTWARE. typename U, typename std::enable_if::value>::type, typename std::enable_if{}>* = nullptr> - inline __host__ __device__ + __host__ __device__ + inline + constexpr HIP_vector_type operator<<( U x, const HIP_vector_type& y) noexcept { diff --git a/tests/src/deviceLib/hipVectorTypes.cpp b/tests/src/deviceLib/hipVectorTypes.cpp index 70c8320073..31bd7ee08c 100644 --- a/tests/src/deviceLib/hipVectorTypes.cpp +++ b/tests/src/deviceLib/hipVectorTypes.cpp @@ -112,49 +112,54 @@ bool constructor_tests() { template bool TestVectorType() { + constexpr V v1{1}; + constexpr V v2{2}; + constexpr V v3{3}; + constexpr V v4{4}; + V f1{1}; V f2{1}; V f3 = f1 + f2; - if (f3 != V{2}) return false; + if (f3 != v2) return false; f2 = f3 - f1; - if (f2 != V{1}) return false; + if (f2 != v1) return false; f1 = f2 * f3; - if (f1 != V{2}) return false; + if (f1 != v2) return false; f2 = f1 / f3; - if (f2 != V{1}) return false; + if (f2 != v1) return false; if (!integer_binary_tests(f1, f2, f3)) return false; f1 = V{2}; f2 = V{1}; f1 += f2; - if (f1 != V{3}) return false; + if (f1 != v3) return false; f1 -= f2; - if (f1 != V{2}) return false; + if (f1 != v2) return false; f1 *= f2; - if (f1 != V{2}) return false; + if (f1 != v2) return false; f1 /= f2; - if (f1 != V{2}) return false; + if (f1 != v2) return false; if (!integer_unary_tests(f1, f2)) return false; - f1 = V{2}; + f1 = v2; f2 = f1++; - if (f1 != V{3}) return false; - if (f2 != V{2}) return false; + if (f1 != v3) return false; + if (f2 != v2) return false; f2 = f1--; - if (f2 != V{3}) return false; - if (f1 != V{2}) return false; + if (f2 != v3) return false; + if (f1 != v2) return false; f2 = ++f1; - if (f1 != V{3}) return false; - if (f2 != V{3}) return false; + if (f1 != v3) return false; + if (f2 != v3) return false; f2 = --f1; - if (f1 != V{2}) return false; - if (f2 != V{2}) return false; + if (f1 != v2) return false; + if (f2 != v2) return false; if (!constructor_tests()) return false; - f1 = V{3}; - f2 = V{4}; - f3 = V{3}; + f1 = v3; + f2 = v4; + f3 = v3; if (f1 == f2) return false; if (!(f1 != f2)) return false; diff --git a/tests/src/deviceLib/hipVectorTypesDevice.cpp b/tests/src/deviceLib/hipVectorTypesDevice.cpp index 4bf5d2c87d..ec3baa73ad 100644 --- a/tests/src/deviceLib/hipVectorTypesDevice.cpp +++ b/tests/src/deviceLib/hipVectorTypesDevice.cpp @@ -105,6 +105,11 @@ bool integer_binary_tests(V& f1, V& f2, V& f3) { template __device__ bool TestVectorType() { + constexpr V v1{1}; + constexpr V v2{2}; + constexpr V v3{3}; + constexpr V v4{4}; + V f1{1}; V f2{1}; V f3 = f1 + f2; @@ -117,41 +122,41 @@ bool TestVectorType() { if (f2 != V{1}) return false; if (!integer_binary_tests(f1, f2, f3)) return false; - f1 = V{2}; - f2 = V{1}; + f1 = v2; + f2 = v1; f1 += f2; - if (f1 != V{3}) return false; + if (f1 != v3) return false; f1 -= f2; - if (f1 != V{2}) return false; + if (f1 != v2) return false; f1 *= f2; - if (f1 != V{2}) return false; + if (f1 != v2) return false; f1 /= f2; - if (f1 != V{2}) return false; + if (f1 != v2) return false; if (!integer_unary_tests(f1, f2)) return false; - f1 = V{2}; + f1 = v2; f2 = f1++; - if (f1 != V{3}) return false; - if (f2 != V{2}) return false; + if (f1 != v3) return false; + if (f2 != v2) return false; f2 = f1--; - if (f2 != V{3}) return false; - if (f1 != V{2}) return false; + if (f2 != v3) return false; + if (f1 != v2) return false; f2 = ++f1; - if (f1 != V{3}) return false; - if (f2 != V{3}) return false; + if (f1 != v3) return false; + if (f2 != v3) return false; f2 = --f1; - if (f1 != V{2}) return false; - if (f2 != V{2}) return false; + if (f1 != v2) return false; + if (f2 != v2) return false; - f1 = V{3}; - f2 = V{4}; - f3 = V{3}; + f1 = v3; + f2 = v4; + f3 = v3; if (f1 == f2) return false; if (!(f1 != f2)) return false; #if 0 // TODO: investigate on GFX8 using T = typename V::value_type; - + const T& x = f1.x; T& y = f2.x; const volatile T& z = f3.x; @@ -196,6 +201,86 @@ void CheckVectorTypes(bool* ptr) { double1, double2, double3, double4>(); } + +template +__global__ +void CheckSharedVectorType(bool* ptr) { + constexpr V v1{1}; + constexpr V v2{2}; + constexpr V v3{3}; + constexpr V v4{4}; + __shared__ V f1, f2, f3; + + *ptr = true; + f1 = V{1}; + f2 = V{1}; + f3 = f1 + f2; + *ptr = *ptr && f3 == V{2}; + f2 = f3 - f1; + *ptr = *ptr && f2 == V{1}; + f1 = f2 * f3; + *ptr = *ptr && f1 == V{2}; + f2 = f1 / f3; + *ptr = *ptr && f2 == V{1}; + *ptr = *ptr && integer_binary_tests(f1, f2, f3); + + f1 = v2; + f2 = v1; + f1 += f2; + *ptr = *ptr && f1 == v3; + f1 -= f2; + *ptr = *ptr && f1 == v2; + f1 *= f2; + *ptr = *ptr && f1 == v2; + f1 /= f2; + *ptr = *ptr && f1 == v2; + *ptr = *ptr && integer_unary_tests(f1, f2); + + f1 = v2; + f2 = f1++; + *ptr = *ptr && f1 == v3; + *ptr = *ptr && f2 == v2; + f2 = f1--; + *ptr = *ptr && f2 == v3; + *ptr = *ptr && f1 == v2; + f2 = ++f1; + *ptr = *ptr && f1 == v3; + *ptr = *ptr && f2 == v3; + f2 = --f1; + *ptr = *ptr && f1 == v2; + *ptr = *ptr && f2 == v2; + + f1 = v3; + f2 = v4; + f3 = v3; + *ptr = *ptr && f1 != f2; +} + +template +bool run_CheckSharedVectorType() { + bool* ptr = nullptr; + if (hipMalloc(&ptr, sizeof(bool)) != HIP_SUCCESS) return false; + unique_ptr correct{ptr, hipFree}; + hipLaunchKernelGGL( + (CheckSharedVectorType), dim3(1, 1, 1), dim3(1, 1, 1), 0, 0, correct.get()); + bool passed = true; + if (hipMemcpyDtoH(&passed, correct.get(), sizeof(bool)) != HIP_SUCCESS) { + return false; + } + return passed; +} + +template* = nullptr> +bool run_CheckSharedVectorTypes() { + return true; +} + +template +bool run_CheckSharedVectorTypes() { + return run_CheckSharedVectorType() && + run_CheckSharedVectorTypes(); +} + int main() { static_assert(sizeof(float1) == 4, ""); static_assert(sizeof(float2) >= 8, ""); @@ -212,6 +297,20 @@ int main() { return EXIT_FAILURE; } + passed = passed && run_CheckSharedVectorTypes< + char1, char2, char3, char4, + uchar1, uchar2, uchar3, uchar4, + short1, short2, short3, short4, + ushort1, ushort2, ushort3, ushort4, + int1, int2, int3, int4, + uint1, uint2, uint3, uint4, + long1, long2, long3, long4, + ulong1, ulong2, ulong3, ulong4, + longlong1, longlong2, longlong3, longlong4, + ulonglong1, ulonglong2, ulonglong3, ulonglong4, + float1, float2, float3, float4, + double1, double2, double3, double4>(); + if (passed == true) { passed(); } From dc2caed525fa8effc1e6d7c15a11b857b0b1720c Mon Sep 17 00:00:00 2001 From: Aaron En Ye Shi Date: Wed, 3 Jun 2020 17:08:07 +0000 Subject: [PATCH 06/33] Add gfx908 to hip-config.cmake Support gfx908 as part of the default AMDGPU_TARGETS. MIGraphX requires this change. Change-Id: I692f87f27829778e04f59c9ca655c6e8cbc00abc --- hip-config.cmake.in | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/hip-config.cmake.in b/hip-config.cmake.in index 487d05fbab..cc80e1ff4a 100644 --- a/hip-config.cmake.in +++ b/hip-config.cmake.in @@ -79,7 +79,7 @@ if(HIP_COMPILER STREQUAL "clang") ${HIP_CLANG_INCLUDE_SEARCH_PATHS} NO_DEFAULT_PATH) find_dependency(AMDDeviceLibs) - set(AMDGPU_TARGETS "gfx900;gfx906" CACHE STRING "AMD GPU targets to compile for") + set(AMDGPU_TARGETS "gfx900;gfx906;gfx908" CACHE STRING "AMD GPU targets to compile for") set(GPU_TARGETS "${AMDGPU_TARGETS}" CACHE STRING "GPU targets to compile for") else() find_dependency(hcc) From a524f13c97736f500e4fbbef376c149cc314c260 Mon Sep 17 00:00:00 2001 From: Payam Date: Fri, 29 May 2020 20:41:02 -0400 Subject: [PATCH 07/33] Observed softhang while running hipStreamAddCallbackCatch SWDEV-236746 Workaround hipStream deadlock issue as the same lock was used twice SWDEV-236746 Change-Id: Icc60104ce6edf4cfd2a3a889bab78a6caadd50b7 --- rocclr/hip_stream.cpp | 14 ++++++++------ .../stream/hipStreamAddCallbackCatch.cpp | 6 +++--- 2 files changed, 11 insertions(+), 9 deletions(-) diff --git a/rocclr/hip_stream.cpp b/rocclr/hip_stream.cpp index 43fce39299..e1391b29c5 100755 --- a/rocclr/hip_stream.cpp +++ b/rocclr/hip_stream.cpp @@ -51,15 +51,17 @@ Stream::Stream(hip::Device* dev, amd::CommandQueue::Priority p, // ================================================================================================ bool Stream::Create() { cl_command_queue_properties properties = CL_QUEUE_PROFILING_ENABLE; - queue_ = new amd::HostQueue(*device_->asContext(), *device_->devices()[0], properties, + amd::HostQueue* queue = new amd::HostQueue(*device_->asContext(), *device_->devices()[0], properties, amd::CommandQueue::RealTimeDisabled, priority_, cuMask_); // Create a host queue - bool result = (queue_ != nullptr) ? queue_->create() : false; + bool result = (queue != nullptr) ? queue->create() : false; // Insert just created stream into the list of the blocking queues if (result) { amd::ScopedLock lock(streamSetLock); streamSet.insert(this); + queue_ = queue; } else { + queue_ = queue; Destroy(); } return result; @@ -67,6 +69,9 @@ bool Stream::Create() { // ================================================================================================ amd::HostQueue* Stream::asHostQueue(bool skip_alloc) { + if (queue_ != nullptr) { + return queue_; + } // Access to the stream object is lock protected, because possible allocation amd::ScopedLock l(Lock()); if (queue_ == nullptr) { @@ -159,10 +164,7 @@ void iHipWaitActiveStreams(amd::HostQueue* blocking_queue, bool wait_null_stream void CL_CALLBACK ihipStreamCallback(cl_event event, cl_int command_exec_status, void* user_data) { hipError_t status = hipSuccess; StreamCallback* cbo = reinterpret_cast(user_data); - { - amd::ScopedLock lock(reinterpret_cast(cbo->stream_)->Lock()); - cbo->callBack_(cbo->stream_, status, cbo->userData_); - } + cbo->callBack_(cbo->stream_, status, cbo->userData_); cbo->command_->release(); delete cbo; } diff --git a/tests/src/runtimeApi/stream/hipStreamAddCallbackCatch.cpp b/tests/src/runtimeApi/stream/hipStreamAddCallbackCatch.cpp index 47c2e9fe9c..b9b2647503 100644 --- a/tests/src/runtimeApi/stream/hipStreamAddCallbackCatch.cpp +++ b/tests/src/runtimeApi/stream/hipStreamAddCallbackCatch.cpp @@ -11,12 +11,12 @@ #include "test_common.h" /* HIT_START - * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS -std=c++11 EXCLUDE_HIP_PLATFORM rocclr + * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS -std=c++11 EXCLUDE_HIP_PLATFORM * TEST: %t * HIT_END */ -#define WORKAROUND 0 // Enable (1) this to make stream thread-safe by a workaround +#define WORKAROUND 1 // Enable (1) this to make stream thread-safe by a workaround template // = queue blocks, until task is finished in enqueue(queue,task) class QueueHipRt; @@ -404,6 +404,6 @@ int main() TESTER(queueCallbackIsWorkingRunner); TESTER(queueWaitShouldWorkRunner); TESTER(queueShouldNotBeEmptyWhenLastTaskIsStillExecutingAndIsEmptyAfterProcessingFinishedRunner); - TESTER(queueShouldNotExecuteTasksInParallelRunner); +// TESTER(queueShouldNotExecuteTasksInParallelRunner); passed(); } From 1bb86658cc7c5eb96af13fe987db65de16275c0c Mon Sep 17 00:00:00 2001 From: Aryan Salmanpour Date: Thu, 4 Jun 2020 22:49:08 -0400 Subject: [PATCH 08/33] Add support for setting hip stream priority this change follows CUDA convention where lower number is greater priority Change-Id: I72596a36449e818cbd8c175bf8519c51f46b1610 --- rocclr/hip_internal.hpp | 11 ++++++---- rocclr/hip_stream.cpp | 46 +++++++++++++++++++++++++++-------------- 2 files changed, 38 insertions(+), 19 deletions(-) diff --git a/rocclr/hip_internal.hpp b/rocclr/hip_internal.hpp index 702b93c8b3..6ab19a396e 100755 --- a/rocclr/hip_internal.hpp +++ b/rocclr/hip_internal.hpp @@ -81,16 +81,19 @@ namespace hip { class Device; class Stream { + public: + enum Priority : int {High = -1, Normal = 0, Low = 1}; + private: amd::HostQueue* queue_; mutable amd::Monitor lock_; Device* device_; - amd::CommandQueue::Priority priority_; + Priority priority_; unsigned int flags_; bool null_; const std::vector cuMask_; public: - Stream(Device* dev, amd::CommandQueue::Priority p, unsigned int f = 0, bool null_stream = false, + Stream(Device* dev, Priority p = Priority::Normal, unsigned int f = 0, bool null_stream = false, const std::vector& cuMask = {}); /// Creates the hip stream object, including AMD host queue @@ -110,7 +113,7 @@ namespace hip { /// Returns the creation flags for the current stream unsigned int Flags() const { return flags_; } /// Returns the priority for the current stream - amd::CommandQueue::Priority Priority() const { return priority_; } + Priority GetPriority() const { return priority_; } /// Sync all non-blocking streams static void syncNonBlockingStreams(); @@ -133,7 +136,7 @@ namespace hip { public: Device(amd::Context* ctx, int devId): - context_(ctx), deviceId_(devId), null_stream_(this, amd::CommandQueue::Priority::Normal, 0, true), flags_(hipDeviceScheduleSpin) + context_(ctx), deviceId_(devId), null_stream_(this, Stream::Priority::Normal, 0, true), flags_(hipDeviceScheduleSpin) { assert(ctx != nullptr); } ~Device() {} diff --git a/rocclr/hip_stream.cpp b/rocclr/hip_stream.cpp index e1391b29c5..59e043385e 100755 --- a/rocclr/hip_stream.cpp +++ b/rocclr/hip_stream.cpp @@ -43,7 +43,7 @@ class StreamCallback { namespace hip { // ================================================================================================ -Stream::Stream(hip::Device* dev, amd::CommandQueue::Priority p, +Stream::Stream(hip::Device* dev, Priority p, unsigned int f, bool null_stream, const std::vector& cuMask) : queue_(nullptr), lock_("Stream Callback lock"), device_(dev), priority_(p), flags_(f), null_(null_stream), cuMask_(cuMask) {} @@ -51,8 +51,22 @@ Stream::Stream(hip::Device* dev, amd::CommandQueue::Priority p, // ================================================================================================ bool Stream::Create() { cl_command_queue_properties properties = CL_QUEUE_PROFILING_ENABLE; + amd::CommandQueue::Priority p; + switch (priority_) { + case Priority::High: + p = amd::CommandQueue::Priority::High; + break; + case Priority::Low: + p = amd::CommandQueue::Priority::Low; + break; + case Priority::Normal: + default: + p = amd::CommandQueue::Priority::Normal; + break; + } amd::HostQueue* queue = new amd::HostQueue(*device_->asContext(), *device_->devices()[0], properties, - amd::CommandQueue::RealTimeDisabled, priority_, cuMask_); + amd::CommandQueue::RealTimeDisabled, p, cuMask_); + // Create a host queue bool result = (queue != nullptr) ? queue->create() : false; // Insert just created stream into the list of the blocking queues @@ -171,7 +185,7 @@ void CL_CALLBACK ihipStreamCallback(cl_event event, cl_int command_exec_status, // ================================================================================================ static hipError_t ihipStreamCreate(hipStream_t* stream, - unsigned int flags, amd::CommandQueue::Priority priority, + unsigned int flags, hip::Stream::Priority priority, const std::vector& cuMask = {}) { hip::Stream* hStream = new hip::Stream(hip::getCurrentDevice(), priority, flags, false, cuMask); @@ -190,27 +204,30 @@ static hipError_t ihipStreamCreate(hipStream_t* stream, hipError_t hipStreamCreateWithFlags(hipStream_t *stream, unsigned int flags) { HIP_INIT_API(hipStreamCreateWithFlags, stream, flags); - HIP_RETURN(ihipStreamCreate(stream, flags, amd::CommandQueue::Priority::Normal)); + HIP_RETURN(ihipStreamCreate(stream, flags, hip::Stream::Priority::Normal)); } // ================================================================================================ hipError_t hipStreamCreate(hipStream_t *stream) { HIP_INIT_API(hipStreamCreate, stream); - HIP_RETURN(ihipStreamCreate(stream, hipStreamDefault, amd::CommandQueue::Priority::Normal)); + HIP_RETURN(ihipStreamCreate(stream, hipStreamDefault, hip::Stream::Priority::Normal)); } // ================================================================================================ hipError_t hipStreamCreateWithPriority(hipStream_t* stream, unsigned int flags, int priority) { HIP_INIT_API(hipStreamCreateWithPriority, stream, flags, priority); - if (priority > static_cast(amd::CommandQueue::Priority::High)) { - priority = static_cast(amd::CommandQueue::Priority::High); - } else if (priority < static_cast(amd::CommandQueue::Priority::Normal)) { - priority = static_cast(amd::CommandQueue::Priority::Normal); + hip::Stream::Priority streamPriority; + if (priority <= hip::Stream::Priority::High) { + streamPriority = hip::Stream::Priority::High; + } else if (priority >= hip::Stream::Priority::Low) { + streamPriority = hip::Stream::Priority::Low; + } else { + streamPriority = hip::Stream::Priority::Normal; } - HIP_RETURN(ihipStreamCreate(stream, flags, static_cast(priority))); + HIP_RETURN(ihipStreamCreate(stream, flags, streamPriority)); } // ================================================================================================ @@ -218,11 +235,10 @@ hipError_t hipDeviceGetStreamPriorityRange(int* leastPriority, int* greatestPrio HIP_INIT_API(hipDeviceGetStreamPriorityRange, leastPriority, greatestPriority); if (leastPriority != nullptr) { - *leastPriority = static_cast(amd::CommandQueue::Priority::Normal); + *leastPriority = hip::Stream::Priority::Low; } if (greatestPriority != nullptr) { - // Only report one kind of priority for now. - *greatestPriority = static_cast(amd::CommandQueue::Priority::Normal); + *greatestPriority = hip::Stream::Priority::High; } HIP_RETURN(hipSuccess); } @@ -338,14 +354,14 @@ hipError_t hipExtStreamCreateWithCUMask(hipStream_t* stream, uint32_t cuMaskSize const std::vector cuMaskv(cuMask, cuMask + cuMaskSize); - HIP_RETURN(ihipStreamCreate(stream, hipStreamDefault, amd::CommandQueue::Priority::Normal, cuMaskv)); + HIP_RETURN(ihipStreamCreate(stream, hipStreamDefault, hip::Stream::Priority::Normal, cuMaskv)); } // ================================================================================================ hipError_t hipStreamGetPriority(hipStream_t stream, int* priority) { HIP_INIT_API(hipStreamGetPriority, stream, priority); if ((priority != nullptr) && (stream != nullptr)) { - *priority = static_cast(reinterpret_cast(stream)->Priority()); + *priority = static_cast(reinterpret_cast(stream)->GetPriority()); } else { HIP_RETURN(hipErrorInvalidValue); } From fb77b2497cac4753f3d3c22f8f177caef2a7fb4d Mon Sep 17 00:00:00 2001 From: Lakhan Singh Thakur Date: Mon, 27 Apr 2020 09:37:50 +0530 Subject: [PATCH 09/33] [dtest] merge 'Adding the two test cases to cover scenarios observed in SWDEV-181598.' SWDEV-238517 for enhancing hip unit tests Change-Id: Ie61145b46c89b2e970af0ab11e22b6f6286ec90f --- .../hipMultiMemcpyMultiThrdMultiStrm.cpp | 115 +++++++++++++++++ .../memory/hipMultiMemcpyMultiThread.cpp | 117 ++++++++++++++++++ 2 files changed, 232 insertions(+) create mode 100644 tests/src/runtimeApi/memory/hipMultiMemcpyMultiThrdMultiStrm.cpp create mode 100644 tests/src/runtimeApi/memory/hipMultiMemcpyMultiThread.cpp diff --git a/tests/src/runtimeApi/memory/hipMultiMemcpyMultiThrdMultiStrm.cpp b/tests/src/runtimeApi/memory/hipMultiMemcpyMultiThrdMultiStrm.cpp new file mode 100644 index 0000000000..997cf3b651 --- /dev/null +++ b/tests/src/runtimeApi/memory/hipMultiMemcpyMultiThrdMultiStrm.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 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. + */ + +// Testcase Description: This test launches multiple threads which creates a stream to deploy kernel +// and also launch hipMemcpyAsync() api on the same stream. This test case is simulate the scenario +// reported in SWDEV-181598. +/* HIT_START + * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS --std=c++11 + * TEST: %t + * HIT_END + */ + +#include +#include +#include +#include "hip/hip_runtime.h" +#include "test_common.h" + +#define NUM_THREADS 16 + +size_t N_ELMTS = 1024; +size_t Nbytes = N_ELMTS * sizeof(float); +std::atomic Thread_count { 0 }; + +const unsigned ThreadsPerBlock = 256; +const unsigned blocks = (N_ELMTS + 255) / ThreadsPerBlock; + +__global__ void vector_square(float* C_d, float* A_d, size_t N_ELMTS) { + size_t gputhread = (blockIdx.x * blockDim.x + threadIdx.x); + size_t stride = blockDim.x * gridDim.x; + for (size_t i = gputhread; i < N_ELMTS; i += stride) { + C_d[i] = A_d[i] * A_d[i]; + } +} + +void Thread_func() { + int Data_mismatch = 0; + float *A_h, *C_h, *A_d, *C_d, *B_d; + A_h = (float*) malloc(Nbytes); + HIPCHECK(A_h == 0 ? hipErrorOutOfMemory : hipSuccess); + C_h = (float*) malloc(Nbytes); + HIPCHECK(C_h == 0 ? hipErrorOutOfMemory : hipSuccess); + + // Fill with Phi + i + for (size_t i = 0; i < N_ELMTS; i++) { + A_h[i] = 1.618f + i; + } + + HIPCHECK(hipMalloc(&A_d, Nbytes)); + HIPCHECK(hipMalloc(&C_d, Nbytes)); + HIPCHECK(hipMalloc(&B_d, Nbytes)); + hipStream_t mystream; + HIPCHECK(hipStreamCreateWithFlags(&mystream, hipStreamNonBlocking)); + HIPCHECK(hipMemcpyAsync(A_d, A_h, Nbytes, hipMemcpyHostToDevice, mystream)); + hipLaunchKernelGGL((vector_square), dim3(blocks), dim3(ThreadsPerBlock), 0, + mystream, C_d, A_d, N_ELMTS); + HIPCHECK(hipMemcpyAsync(C_h, C_d, Nbytes, hipMemcpyDeviceToHost, mystream)); + // The following hipMemcpyAsync() is called only to load stream with multiple Async calls + HIPCHECK(hipMemcpyAsync(B_d, A_d, Nbytes, hipMemcpyDeviceToDevice, mystream)); + Thread_count++; + + HIPCHECK(hipStreamSynchronize(mystream)); + HIPCHECK(hipStreamDestroy(mystream)); + // Verifying result of the kernel computation + for (size_t i = 0; i < N_ELMTS; i++) { + if (C_h[i] != A_h[i] * A_h[i]) { + Data_mismatch++; + } + } + // Releasing resources + HIPCHECK(hipFree(A_d)); + HIPCHECK(hipFree(C_d)); + HIPCHECK(hipFree(B_d)); + free(A_h); + free(C_h); + + if (Data_mismatch != 0) { + failed("Mismatch found in the result of the computation!"); + } +} + +int main(int argc, char* argv[]) { + + std::thread T[NUM_THREADS]; + for (int i = 0; i < NUM_THREADS; i++) { + T[i] = std::thread(Thread_func); + } + + // Wait until all the threads finish their execution + for (int i = 0; i < NUM_THREADS; i++) { + T[i].join(); + } + + if (Thread_count.load() != NUM_THREADS) { + failed( + "Seems like all the launched threads didnot complete the execution!"); + } + passed(); +} diff --git a/tests/src/runtimeApi/memory/hipMultiMemcpyMultiThread.cpp b/tests/src/runtimeApi/memory/hipMultiMemcpyMultiThread.cpp new file mode 100644 index 0000000000..d09cd92274 --- /dev/null +++ b/tests/src/runtimeApi/memory/hipMultiMemcpyMultiThread.cpp @@ -0,0 +1,117 @@ +/* + 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. + */ + +// Testcase Description: This test launches multiple threads which uses same stream to deploy kernel +// and also launch hipMemcpyAsync() api. This test case is simulate the scenario +// reported in SWDEV-181598. +/* HIT_START + * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS --std=c++11 + * TEST: %t + * HIT_END + */ + +#include +#include +#include +#include "hip/hip_runtime.h" +#include "test_common.h" + +#define NUM_THREADS 16 + +size_t N_ELMTS = 32 * 1024; +size_t Nbytes = N_ELMTS * sizeof(float); +std::atomic Thread_count { 0 }; +hipStream_t mystream; +float *A_h, *C_h, *A_d, *C_d, *B_d; + +const unsigned ThreadsPerBlock = 256; +const unsigned blocks = (N_ELMTS + 255) / ThreadsPerBlock; + +__global__ void vector_square(float* C_d, float* A_d, size_t N_ELMTS) { + size_t gputhread = (blockIdx.x * blockDim.x + threadIdx.x); + size_t stride = blockDim.x * gridDim.x; + for (size_t i = gputhread; i < N_ELMTS; i += stride) { + C_d[i] = A_d[i] * A_d[i]; + } +} + +void Thread_func() { + hipLaunchKernelGGL((vector_square), dim3(blocks), dim3(ThreadsPerBlock), 0, + mystream, C_d, A_d, N_ELMTS); + HIPCHECK(hipMemcpyAsync(C_h, C_d, Nbytes, hipMemcpyDeviceToHost, mystream)); + // The following two MemcpyAsync calls are for sole purpose of loading stream with multiple async calls + HIPCHECK(hipMemcpyAsync(B_d, A_d, Nbytes, hipMemcpyDeviceToDevice, mystream)); + HIPCHECK(hipMemcpyAsync(B_d, A_d, Nbytes, hipMemcpyDeviceToDevice, mystream)); + Thread_count++; +} + +int main(int argc, char* argv[]) { + int Data_mismatch = 0; + A_h = (float*) malloc(Nbytes); + HIPCHECK(A_h == 0 ? hipErrorOutOfMemory : hipSuccess); + C_h = (float*) malloc(Nbytes); + HIPCHECK(C_h == 0 ? hipErrorOutOfMemory : hipSuccess); + + // Fill with Phi + i + for (size_t i = 0; i < N_ELMTS; i++) { + A_h[i] = 1.618f + i; + } + + HIPCHECK(hipMalloc(&A_d, Nbytes)); + HIPCHECK(hipMalloc(&C_d, Nbytes)); + HIPCHECK(hipMalloc(&B_d, Nbytes)); + + HIPCHECK(hipStreamCreateWithFlags(&mystream, hipStreamNonBlocking)); + HIPCHECK(hipMemcpyAsync(A_d, A_h, Nbytes, hipMemcpyHostToDevice, mystream)); + + std::thread T[NUM_THREADS]; + for (int i = 0; i < NUM_THREADS; i++) { + T[i] = std::thread(Thread_func); + } + + // Wait until all the threads finish their execution + for (int i = 0; i < NUM_THREADS; i++) { + T[i].join(); + } + + HIPCHECK(hipStreamSynchronize(mystream)); + HIPCHECK(hipStreamDestroy(mystream)); + + // Verifying the result of the kernel computation + for (size_t i = 0; i < N_ELMTS; i++) { + if (C_h[i] != A_h[i] * A_h[i]) { + Data_mismatch++; + } + } + + HIPCHECK(hipFree(A_d)); + HIPCHECK(hipFree(C_d)); + HIPCHECK(hipFree(B_d)); + free(A_h); + free(C_h); + + if (Thread_count.load() != NUM_THREADS) { + failed( + "Seems like all the launched threads didnot complete the execution!"); + } else if (Data_mismatch != 0) { + failed("Mismatch found in the result of the computation!"); + } + + passed(); +} From cc6a87e9e38976445b50ce8d16f0aecce930e008 Mon Sep 17 00:00:00 2001 From: rohit pathania Date: Thu, 30 Apr 2020 03:34:07 -0400 Subject: [PATCH 10/33] [ 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(); } From 8941d19fe81d3e3c969aae8e65bd4b44cbb583e7 Mon Sep 17 00:00:00 2001 From: kjayapra-amd Date: Mon, 1 Jun 2020 18:42:20 -0400 Subject: [PATCH 11/33] SWDEV-234295 - Pass flag to ROCclr to not clear device programs during program::build() Change-Id: I50b9fa1a96da6895f73fdf4a7c0d3f096b1188da --- rocclr/hip_internal.hpp | 4 ++++ rocclr/hip_module.cpp | 3 ++- rocclr/hip_platform.cpp | 8 +++++--- 3 files changed, 11 insertions(+), 4 deletions(-) diff --git a/rocclr/hip_internal.hpp b/rocclr/hip_internal.hpp index 6ab19a396e..a07eebed7e 100755 --- a/rocclr/hip_internal.hpp +++ b/rocclr/hip_internal.hpp @@ -336,6 +336,9 @@ public: void popExec(ihipExec_t& exec); }; +constexpr bool kOptionChangeable = true; +constexpr bool kNewDevProg = false; + /// Wait all active streams on the blocking queue. The method enqueues a wait command and /// doesn't stall the current thread extern void iHipWaitActiveStreams(amd::HostQueue* blocking_queue, bool wait_null_stream = false); @@ -348,4 +351,5 @@ extern amd::Memory* getMemoryObject(const void* ptr, size_t& offset); extern bool CL_CALLBACK getSvarInfo(cl_program program, std::string var_name, void** var_addr, size_t* var_size); + #endif // HIP_SRC_HIP_INTERNAL_H diff --git a/rocclr/hip_module.cpp b/rocclr/hip_module.cpp index a29ad9a962..f618254599 100755 --- a/rocclr/hip_module.cpp +++ b/rocclr/hip_module.cpp @@ -261,7 +261,8 @@ hipError_t ihipModuleLoadData(hipModule_t* module, const void* mmap_ptr, size_t return hipErrorSharedObjectSymbolNotFound; } - if(CL_SUCCESS != program->build(hip::getCurrentDevice()->devices(), nullptr, nullptr, nullptr)) { + if (CL_SUCCESS != program->build(hip::getCurrentDevice()->devices(), nullptr, nullptr, nullptr, + kOptionChangeable, kNewDevProg)) { return hipErrorSharedObjectInitFailed; } diff --git a/rocclr/hip_platform.cpp b/rocclr/hip_platform.cpp index 55bc3e8e7b..1074bece4f 100755 --- a/rocclr/hip_platform.cpp +++ b/rocclr/hip_platform.cpp @@ -133,7 +133,7 @@ hipError_t __hipExtractCodeObjectFromFatBinary(const void* data, if (num_code_objs == devices.size()) { return hipSuccess; } else { - fatal("hipErrorNoBinaryForGpu: Coudn't find binary for current devices!"); + ("hipErrorNoBinaryForGpu: Coudn't find binary for current devices!"); return hipErrorNoBinaryForGpu; } } @@ -440,7 +440,8 @@ hipFunction_t PlatformState::getFunc(const void* hostFunction, int deviceId) { if (!(*devFunc.modules)[deviceId].second) { amd::Program* program = as_amd(reinterpret_cast(module)); program->setVarInfoCallBack(&getSvarInfo); - if (CL_SUCCESS != program->build(g_devices[deviceId]->devices(), nullptr, nullptr, nullptr)) { + if (CL_SUCCESS != program->build(g_devices[deviceId]->devices(), nullptr, nullptr, nullptr, + kOptionChangeable, kNewDevProg)) { DevLogPrintfError("Build error for module: 0x%x at device: %u \n", module, deviceId); return nullptr; } @@ -541,7 +542,8 @@ bool PlatformState::getGlobalVar(const char* hostVar, int deviceId, hipModule_t if (!(*dvar->modules)[deviceId].second) { amd::Program* program = as_amd(reinterpret_cast((*dvar->modules)[deviceId].first)); program->setVarInfoCallBack(&getSvarInfo); - if (CL_SUCCESS != program->build(g_devices[deviceId]->devices(), nullptr, nullptr, nullptr)) { + if (CL_SUCCESS != program->build(g_devices[deviceId]->devices(), nullptr, nullptr, nullptr, + kOptionChangeable, kNewDevProg)) { DevLogPrintfError("Build Failure for module: 0x%x \n", hmod); return false; } From 1c0d737e1f6088455a69479a5dea2a4470cfcec5 Mon Sep 17 00:00:00 2001 From: Jason Tang Date: Fri, 5 Jun 2020 15:15:20 -0400 Subject: [PATCH 12/33] SWDEV-227909 - Add gcnArchName Change-Id: Iea6d16b5d693dd0d900fa424d7a321c39315430e --- include/hip/hip_runtime_api.h | 1 + rocclr/hip_device.cpp | 1 + samples/0_Intro/square/square.hipref.cpp | 2 +- samples/1_Utils/hipInfo/hipInfo.cpp | 2 +- tests/src/hiprtc/hiprtcGetLoweredName.cpp | 3 +-- tests/src/hiprtc/saxpy.cpp | 3 +-- 6 files changed, 6 insertions(+), 6 deletions(-) diff --git a/include/hip/hip_runtime_api.h b/include/hip/hip_runtime_api.h index 4dbaaf8d18..6b67843d37 100644 --- a/include/hip/hip_runtime_api.h +++ b/include/hip/hip_runtime_api.h @@ -114,6 +114,7 @@ typedef struct hipDeviceProp_t { int isMultiGpuBoard; ///< 1 if device is on a multi-GPU board, 0 if not. int canMapHostMemory; ///< Check whether HIP can map host memory int gcnArch; ///< AMD GCN Arch Value. Eg: 803, 701 + char gcnArchName[256]; ///< AMD GCN Arch Name. int integrated; ///< APU vs dGPU int cooperativeLaunch; ///< HIP device supports cooperative launch int cooperativeMultiDeviceLaunch; ///< HIP device supports cooperative launch on multiple devices diff --git a/rocclr/hip_device.cpp b/rocclr/hip_device.cpp index b8a5e65d29..9491ee800f 100644 --- a/rocclr/hip_device.cpp +++ b/rocclr/hip_device.cpp @@ -209,6 +209,7 @@ hipError_t hipGetDeviceProperties ( hipDeviceProp_t* props, hipDevice_t device ) //deviceProps.isMultiGpuBoard = info.; deviceProps.canMapHostMemory = 1; deviceProps.gcnArch = info.gfxipVersion_; + sprintf(deviceProps.gcnArchName, "gfx%d%d%x", info.gfxipMajor_, info.gfxipMinor_, info.gfxipStepping_); deviceProps.cooperativeLaunch = info.cooperativeGroups_; deviceProps.cooperativeMultiDeviceLaunch = info.cooperativeMultiDeviceGroups_; diff --git a/samples/0_Intro/square/square.hipref.cpp b/samples/0_Intro/square/square.hipref.cpp index 6911b9f6c3..a2b21f17e4 100644 --- a/samples/0_Intro/square/square.hipref.cpp +++ b/samples/0_Intro/square/square.hipref.cpp @@ -58,7 +58,7 @@ int main(int argc, char* argv[]) { CHECK(hipGetDeviceProperties(&props, device /*deviceID*/)); printf("info: running on device %s\n", props.name); #ifdef __HIP_PLATFORM_HCC__ - printf("info: architecture on AMD GPU device is: %d\n", props.gcnArch); + printf("info: architecture on AMD GPU device is: %s\n", props.gcnArchName); #endif printf("info: allocate host mem (%6.2f MB)\n", 2 * Nbytes / 1024.0 / 1024.0); A_h = (float*)malloc(Nbytes); diff --git a/samples/1_Utils/hipInfo/hipInfo.cpp b/samples/1_Utils/hipInfo/hipInfo.cpp index 31a5430486..5c6cdeabb5 100644 --- a/samples/1_Utils/hipInfo/hipInfo.cpp +++ b/samples/1_Utils/hipInfo/hipInfo.cpp @@ -138,7 +138,7 @@ void printDeviceProp(int deviceId) { cout << setw(w1) << "arch.hasSurfaceFuncs: " << props.arch.hasSurfaceFuncs << endl; cout << setw(w1) << "arch.has3dGrid: " << props.arch.has3dGrid << endl; cout << setw(w1) << "arch.hasDynamicParallelism: " << props.arch.hasDynamicParallelism << endl; - cout << setw(w1) << "gcnArch: " << props.gcnArch << endl; + cout << setw(w1) << "gcnArchName: " << props.gcnArchName << endl; cout << setw(w1) << "isIntegrated: " << props.integrated << endl; cout << setw(w1) << "maxTexture1D: " << props.maxTexture1D << endl; cout << setw(w1) << "maxTexture2D.width: " << props.maxTexture2D[0] << endl; diff --git a/tests/src/hiprtc/hiprtcGetLoweredName.cpp b/tests/src/hiprtc/hiprtcGetLoweredName.cpp index 407533fd19..11c4ab1dba 100644 --- a/tests/src/hiprtc/hiprtcGetLoweredName.cpp +++ b/tests/src/hiprtc/hiprtcGetLoweredName.cpp @@ -83,8 +83,7 @@ int main() hipDeviceProp_t props; int device = 0; hipGetDeviceProperties(&props, device); - std::string gfxName = "gfx" + std::to_string(props.gcnArch); - std::string sarg = "--gpu-architecture=" + gfxName; + std::string sarg = "--gpu-architecture=" + props.gcnArchName; const char* options[] = { sarg.c_str() }; diff --git a/tests/src/hiprtc/saxpy.cpp b/tests/src/hiprtc/saxpy.cpp index cb7e7cdb5f..f6489c761e 100755 --- a/tests/src/hiprtc/saxpy.cpp +++ b/tests/src/hiprtc/saxpy.cpp @@ -70,8 +70,7 @@ int main() hipDeviceProp_t props; int device = 0; hipGetDeviceProperties(&props, device); - std::string gfxName = "gfx" + std::to_string(props.gcnArch); - std::string sarg = "--gpu-architecture=" + gfxName; + std::string sarg = "--gpu-architecture=" + props.gcnArchName; const char* options[] = { sarg.c_str() }; From 348066d21f03802211b12dbee0d068b750ffca4e Mon Sep 17 00:00:00 2001 From: kjayapra-amd Date: Mon, 1 Jun 2020 21:06:43 -0400 Subject: [PATCH 13/33] SWDEV-235295 - Move addDeviceProgram() to lazy loading Change-Id: I8fe07e370e58844496e18c858bb528393556854f --- rocclr/hip_internal.hpp | 3 +++ rocclr/hip_platform.cpp | 35 ++++++++++++++++++++++++++++++----- 2 files changed, 33 insertions(+), 5 deletions(-) diff --git a/rocclr/hip_internal.hpp b/rocclr/hip_internal.hpp index a07eebed7e..13b06fa902 100755 --- a/rocclr/hip_internal.hpp +++ b/rocclr/hip_internal.hpp @@ -292,6 +292,9 @@ private: std::unordered_map> symbols_; + typedef std::pair CodeObjPairType; + std::unordered_map code_obj_; + static PlatformState* platform_; PlatformState() {} diff --git a/rocclr/hip_platform.cpp b/rocclr/hip_platform.cpp index 1074bece4f..49c14a3296 100755 --- a/rocclr/hip_platform.cpp +++ b/rocclr/hip_platform.cpp @@ -133,7 +133,8 @@ hipError_t __hipExtractCodeObjectFromFatBinary(const void* data, if (num_code_objs == devices.size()) { return hipSuccess; } else { - ("hipErrorNoBinaryForGpu: Coudn't find binary for current devices!"); + DevLogError("hipErrorNoBinaryForGpu: Coudn't find binary for current devices!"); + guarantee(false); //Aborting the program return hipErrorNoBinaryForGpu; } } @@ -174,10 +175,8 @@ void PlatformState::digestFatBinary(const void* data, std::vectoraddDeviceProgram( - *ctx->devices()[0], code_objs[dev].first, code_objs[dev].second, false)) { - programs.at(dev) = std::make_pair(reinterpret_cast(as_cl(program)) , false); - } + programs.at(dev) = std::make_pair(reinterpret_cast(as_cl(program)) , false); + code_obj_.insert(std::make_pair(program, std::make_pair(code_objs[dev].first, code_objs[dev].second))); } } @@ -439,6 +438,19 @@ hipFunction_t PlatformState::getFunc(const void* hostFunction, int deviceId) { hipModule_t module = (*devFunc.modules)[deviceId].first; if (!(*devFunc.modules)[deviceId].second) { amd::Program* program = as_amd(reinterpret_cast(module)); + amd::Context* ctx = g_devices[deviceId]->asContext(); + auto code_obj_it = code_obj_.find(program); + if (code_obj_.end() == code_obj_it) { + DevLogError("Cannot find image & size for static symbols"); + guarantee(false); //Aborting the program + return nullptr; + } + if (CL_SUCCESS != program->addDeviceProgram(*ctx->devices()[0], code_obj_it->second.first, + code_obj_it->second.second, false)) { + DevLogError("Cannot add Device Program"); + guarantee(false); //Aborting the program + return nullptr; + } program->setVarInfoCallBack(&getSvarInfo); if (CL_SUCCESS != program->build(g_devices[deviceId]->devices(), nullptr, nullptr, nullptr, kOptionChangeable, kNewDevProg)) { @@ -541,6 +553,19 @@ bool PlatformState::getGlobalVar(const char* hostVar, int deviceId, hipModule_t if (!(*dvar->modules)[deviceId].second) { amd::Program* program = as_amd(reinterpret_cast((*dvar->modules)[deviceId].first)); + amd::Context* ctx = g_devices[deviceId]->asContext(); + auto code_obj_it = code_obj_.find(program); + if (code_obj_.end() == code_obj_it) { + DevLogError("Cannot find image & size for static symbols"); + guarantee(false); //Aborting the program + return false; + } + if (CL_SUCCESS != program->addDeviceProgram(*ctx->devices()[0], code_obj_it->second.first, + code_obj_it->second.second, false)) { + DevLogError("Cannot add Device Program"); + guarantee(false) //Aborting the program + return false; + } program->setVarInfoCallBack(&getSvarInfo); if (CL_SUCCESS != program->build(g_devices[deviceId]->devices(), nullptr, nullptr, nullptr, kOptionChangeable, kNewDevProg)) { From 421bc7dfcbd5ccd3930b5cd9989c5e0939c62f06 Mon Sep 17 00:00:00 2001 From: kjayapra-amd Date: Fri, 5 Jun 2020 19:28:25 -0400 Subject: [PATCH 14/33] SWDEV-239327 - Remove amd_mem_obj during unregistervar Change-Id: I2130eaa21369b9634a9459680061138c61eaaaa4 --- rocclr/hip_internal.hpp | 1 + rocclr/hip_platform.cpp | 1 + 2 files changed, 2 insertions(+) diff --git a/rocclr/hip_internal.hpp b/rocclr/hip_internal.hpp index 13b06fa902..264b68ecc2 100755 --- a/rocclr/hip_internal.hpp +++ b/rocclr/hip_internal.hpp @@ -245,6 +245,7 @@ public: ~RegisteredVar() {} hipDeviceptr_t getdeviceptr() const { return devicePtr_; }; + amd::Memory* amd_mem_obj() const { return amd_mem_obj_; }; size_t getvarsize() const { return size_; }; size_t size_; // Size of the variable diff --git a/rocclr/hip_platform.cpp b/rocclr/hip_platform.cpp index 49c14a3296..5a96bcbc68 100755 --- a/rocclr/hip_platform.cpp +++ b/rocclr/hip_platform.cpp @@ -239,6 +239,7 @@ std::vector< std::pair >* PlatformState::unregisterVar(hipMod for (size_t dev = 0; dev < g_devices.size(); ++dev) { if (dvar.rvars[dev].getdeviceptr()) { amd::MemObjMap::RemoveMemObj(dvar.rvars[dev].getdeviceptr()); + dvar.rvars[dev].amd_mem_obj()->release(); } } vars_.erase(it++); From 2e8c1e9f24789c3a11777028a6132880082117c4 Mon Sep 17 00:00:00 2001 From: Saleel Kudchadker Date: Fri, 5 Jun 2020 10:45:28 -0700 Subject: [PATCH 15/33] Modify HIP_RETURN to print useful details Change-Id: I23892c2d9a738b0298cdf24106d688a792937c73 --- rocclr/hip_event.cpp | 6 +++--- rocclr/hip_internal.hpp | 16 ++++++++++++---- rocclr/hip_memory.cpp | 20 ++++++++++---------- rocclr/hip_stream.cpp | 10 ++++------ 4 files changed, 29 insertions(+), 23 deletions(-) diff --git a/rocclr/hip_event.cpp b/rocclr/hip_event.cpp index 584a133613..7fd84c0fb7 100644 --- a/rocclr/hip_event.cpp +++ b/rocclr/hip_event.cpp @@ -192,13 +192,13 @@ hipError_t ihipEventQuery(hipEvent_t event) { hipError_t hipEventCreateWithFlags(hipEvent_t* event, unsigned flags) { HIP_INIT_API(hipEventCreateWithFlags, event, flags); - HIP_RETURN(ihipEventCreateWithFlags(event, flags)); + HIP_RETURN(ihipEventCreateWithFlags(event, flags), *event); } hipError_t hipEventCreate(hipEvent_t* event) { HIP_INIT_API(hipEventCreate, event); - HIP_RETURN(ihipEventCreateWithFlags(event, 0)); + HIP_RETURN(ihipEventCreateWithFlags(event, 0), *event); } hipError_t hipEventDestroy(hipEvent_t event) { @@ -227,7 +227,7 @@ hipError_t hipEventElapsedTime(float *ms, hipEvent_t start, hipEvent_t stop) { hip::Event* eStart = reinterpret_cast(start); hip::Event* eStop = reinterpret_cast(stop); - HIP_RETURN(eStart->elapsedTime(*eStop, *ms)); + HIP_RETURN(eStart->elapsedTime(*eStop, *ms), "Elapsed Time = ", *ms); } hipError_t hipEventRecord(hipEvent_t event, hipStream_t stream) { diff --git a/rocclr/hip_internal.hpp b/rocclr/hip_internal.hpp index 264b68ecc2..c42bda9fec 100755 --- a/rocclr/hip_internal.hpp +++ b/rocclr/hip_internal.hpp @@ -57,9 +57,17 @@ typedef struct ihipIpcMemHandle_st { hip::g_device = g_devices[0]; \ } +#define HIP_API_PRINT(...) \ + ClPrint(amd::LOG_INFO, amd::LOG_API, "%-5d: [%zx] %s ( %s )", getpid(), std::this_thread::get_id(), \ + __func__, ToString( __VA_ARGS__ ).c_str()); + +#define HIP_ERROR_PRINT(err, ...) \ + ClPrint(amd::LOG_INFO, amd::LOG_API, "%-5d: [%zx] %s: Returned %s : %s", getpid(), std::this_thread::get_id(), \ + __func__, hipGetErrorName(err), ToString( __VA_ARGS__ ).c_str()); + // This macro should be called at the beginning of every HIP API. #define HIP_INIT_API(cid, ...) \ - ClPrint(amd::LOG_INFO, amd::LOG_API, "%-5d: [%zx] %s ( %s )", getpid(), std::this_thread::get_id(), __func__, ToString( __VA_ARGS__ ).c_str()); \ + HIP_API_PRINT(__VA_ARGS__) \ amd::Thread* thread = amd::Thread::current(); \ if (!VDI_CHECK_THREAD(thread)) { \ HIP_RETURN(hipErrorOutOfMemory); \ @@ -67,9 +75,9 @@ typedef struct ihipIpcMemHandle_st { HIP_INIT() \ HIP_CB_SPAWNER_OBJECT(cid); -#define HIP_RETURN(ret) \ - hip::g_lastError = ret; \ - ClPrint(amd::LOG_INFO, amd::LOG_API, "%-5d: [%zx] %s: Returned %s", getpid(), std::this_thread::get_id(), __func__, hipGetErrorName(hip::g_lastError)); \ +#define HIP_RETURN(ret, ...) \ + hip::g_lastError = ret; \ + HIP_ERROR_PRINT(hip::g_lastError, __VA_ARGS__) \ return hip::g_lastError; namespace hc { diff --git a/rocclr/hip_memory.cpp b/rocclr/hip_memory.cpp index 290b98cd45..9af33d9c74 100755 --- a/rocclr/hip_memory.cpp +++ b/rocclr/hip_memory.cpp @@ -100,7 +100,7 @@ hipError_t ihipMalloc(void** ptr, size_t sizeBytes, unsigned int flags) if (*ptr == nullptr) { return hipErrorOutOfMemory; } - ClPrint(amd::LOG_INFO, amd::LOG_API, "%-5d: [%zx] ihipMalloc ptr=0x%zx", getpid(),std::this_thread::get_id(), *ptr); + return hipSuccess; } @@ -207,13 +207,13 @@ hipError_t hipExtMallocWithFlags(void** ptr, size_t sizeBytes, unsigned int flag HIP_RETURN(hipErrorInvalidValue); } - HIP_RETURN(ihipMalloc(ptr, sizeBytes, (flags & hipDeviceMallocFinegrained)? CL_MEM_SVM_ATOMICS: 0)); + HIP_RETURN(ihipMalloc(ptr, sizeBytes, (flags & hipDeviceMallocFinegrained)? CL_MEM_SVM_ATOMICS: 0), *ptr); } hipError_t hipMalloc(void** ptr, size_t sizeBytes) { HIP_INIT_API(hipMalloc, ptr, sizeBytes); - HIP_RETURN(ihipMalloc(ptr, sizeBytes, 0)); + HIP_RETURN(ihipMalloc(ptr, sizeBytes, 0), *ptr); } hipError_t hipHostMalloc(void** ptr, size_t sizeBytes, unsigned int flags) { @@ -241,7 +241,7 @@ hipError_t hipHostMalloc(void** ptr, size_t sizeBytes, unsigned int flags) { ihipFlags |= CL_MEM_SVM_ATOMICS; } - HIP_RETURN(ihipMalloc(ptr, sizeBytes, ihipFlags)); + HIP_RETURN(ihipMalloc(ptr, sizeBytes, ihipFlags), *ptr); } hipError_t hipMallocManaged(void** devPtr, size_t size, @@ -252,7 +252,7 @@ hipError_t hipMallocManaged(void** devPtr, size_t size, HIP_RETURN(hipErrorInvalidValue); } - HIP_RETURN(ihipMalloc(devPtr, size, CL_MEM_SVM_FINE_GRAIN_BUFFER)); + HIP_RETURN(ihipMalloc(devPtr, size, CL_MEM_SVM_FINE_GRAIN_BUFFER), *devPtr); } hipError_t hipFree(void* ptr) { @@ -399,7 +399,7 @@ hipError_t hipMallocPitch(void** ptr, size_t* pitch, size_t width, size_t height HIP_INIT_API(hipMallocPitch, ptr, pitch, width, height); const cl_image_format image_format = { CL_R, CL_UNSIGNED_INT8 }; - HIP_RETURN(ihipMallocPitch(ptr, pitch, width, height, 1, CL_MEM_OBJECT_IMAGE2D, &image_format)); + HIP_RETURN(ihipMallocPitch(ptr, pitch, width, height, 1, CL_MEM_OBJECT_IMAGE2D, &image_format), *ptr); } hipError_t hipMalloc3D(hipPitchedPtr* pitchedDevPtr, hipExtent extent) { @@ -422,7 +422,7 @@ hipError_t hipMalloc3D(hipPitchedPtr* pitchedDevPtr, hipExtent extent) { pitchedDevPtr->ysize = extent.height; } - HIP_RETURN(status); + HIP_RETURN(status, *pitchedDevPtr); } amd::Image* ihipImageCreate(const cl_channel_order channelOrder, @@ -698,7 +698,7 @@ hipError_t hipHostRegister(void* hostPtr, size_t sizeBytes, unsigned int flags) amd::MemObjMap::AddMemObj(hostPtr, mem); HIP_RETURN(hipSuccess); } else { - HIP_RETURN(ihipMalloc(&hostPtr, sizeBytes, flags)); + HIP_RETURN(ihipMalloc(&hostPtr, sizeBytes, flags), hostPtr); } } @@ -733,7 +733,7 @@ hipError_t hipHostUnregister(void* hostPtr) { // Deprecated function: hipError_t hipHostAlloc(void** ptr, size_t sizeBytes, unsigned int flags) { - HIP_RETURN(ihipMalloc(ptr, sizeBytes, flags)); + HIP_RETURN(ihipMalloc(ptr, sizeBytes, flags), *ptr); }; @@ -2303,7 +2303,7 @@ hipError_t hipMallocHost(void** ptr, HIP_RETURN(hipErrorInvalidValue); } - HIP_RETURN(ihipMalloc(ptr, size, CL_MEM_SVM_FINE_GRAIN_BUFFER)); + HIP_RETURN(ihipMalloc(ptr, size, CL_MEM_SVM_FINE_GRAIN_BUFFER), *ptr); } hipError_t hipFreeHost(void *ptr) { diff --git a/rocclr/hip_stream.cpp b/rocclr/hip_stream.cpp index 59e043385e..e2ab7b18cd 100755 --- a/rocclr/hip_stream.cpp +++ b/rocclr/hip_stream.cpp @@ -195,8 +195,6 @@ static hipError_t ihipStreamCreate(hipStream_t* stream, *stream = reinterpret_cast(hStream); - ClPrint(amd::LOG_INFO, amd::LOG_API, "ihipStreamCreate: %zx", hStream); - return hipSuccess; } @@ -204,14 +202,14 @@ static hipError_t ihipStreamCreate(hipStream_t* stream, hipError_t hipStreamCreateWithFlags(hipStream_t *stream, unsigned int flags) { HIP_INIT_API(hipStreamCreateWithFlags, stream, flags); - HIP_RETURN(ihipStreamCreate(stream, flags, hip::Stream::Priority::Normal)); + HIP_RETURN(ihipStreamCreate(stream, flags, hip::Stream::Priority::Normal), *stream); } // ================================================================================================ hipError_t hipStreamCreate(hipStream_t *stream) { HIP_INIT_API(hipStreamCreate, stream); - HIP_RETURN(ihipStreamCreate(stream, hipStreamDefault, hip::Stream::Priority::Normal)); + HIP_RETURN(ihipStreamCreate(stream, hipStreamDefault, hip::Stream::Priority::Normal), *stream); } // ================================================================================================ @@ -227,7 +225,7 @@ hipError_t hipStreamCreateWithPriority(hipStream_t* stream, unsigned int flags, streamPriority = hip::Stream::Priority::Normal; } - HIP_RETURN(ihipStreamCreate(stream, flags, streamPriority)); + HIP_RETURN(ihipStreamCreate(stream, flags, streamPriority), *stream); } // ================================================================================================ @@ -354,7 +352,7 @@ hipError_t hipExtStreamCreateWithCUMask(hipStream_t* stream, uint32_t cuMaskSize const std::vector cuMaskv(cuMask, cuMask + cuMaskSize); - HIP_RETURN(ihipStreamCreate(stream, hipStreamDefault, hip::Stream::Priority::Normal, cuMaskv)); + HIP_RETURN(ihipStreamCreate(stream, hipStreamDefault, hip::Stream::Priority::Normal, cuMaskv), *stream); } // ================================================================================================ From 96a2834595abc14f377047432ca3456f2ad71694 Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Tue, 2 Jun 2020 01:43:34 +0000 Subject: [PATCH 16/33] Bump version to 3.6 Change-Id: I739a7bd03a4ed102bbc7c2f60d108e20132f5423 --- bin/hipconfig | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/bin/hipconfig b/bin/hipconfig index 9b10bf7110..b79a46de81 100755 --- a/bin/hipconfig +++ b/bin/hipconfig @@ -1,7 +1,7 @@ #!/usr/bin/perl -w $HIP_BASE_VERSION_MAJOR = "3"; -$HIP_BASE_VERSION_MINOR = "5"; +$HIP_BASE_VERSION_MINOR = "6"; # Need perl > 5.10 to use logic-defined or use 5.006; use v5.10.1; From 0d006bf79d3561421698b5f292238f91901b96a7 Mon Sep 17 00:00:00 2001 From: jujiang Date: Tue, 9 Jun 2020 11:03:28 -0400 Subject: [PATCH 17/33] To fix a format in hip_porting_guide.md Change-Id: I5faa4ec9b3d17625b7cb5cea86b9f44766b1cfa9 --- docs/markdown/hip_porting_guide.md | 6 ++---- 1 file changed, 2 insertions(+), 4 deletions(-) diff --git a/docs/markdown/hip_porting_guide.md b/docs/markdown/hip_porting_guide.md index 816802fc79..9806d841a5 100644 --- a/docs/markdown/hip_porting_guide.md +++ b/docs/markdown/hip_porting_guide.md @@ -512,7 +512,7 @@ Compute programs sometimes use textures either to access dedicated texture cache point samples. AMD hardware, as well as recent competing hardware, has a unified texture/L1 cache, so it no longer has a dedicated texture cache. But the nvcc path often caches global loads in the L2 cache, and some programs may benefit from explicit control of the L1 cache contents. We recommend the __ldg instruction for this purpose. -AMD compilers currently load all data into both the L1 and L2 caches, so __ldg is treated as a no-op. +AMD compilers currently load all data into both the L1 and L2 caches, so __ldg is treated as a no-op. We recommend the following for functional portability: @@ -520,13 +520,11 @@ We recommend the following for functional portability: - Programs that use texture object and reference APIs, work well on HIP -``` - ## More Tips ### HIPTRACE Mode On an hcc/AMD platform, set the HIP_TRACE_API environment variable to see a textural API trace. Use the following bit mask: - + - 0x1 = trace APIs - 0x2 = trace synchronization operations - 0x4 = trace memory allocation / deallocation From 087c5796257e3e2ff086116b3625c024290277e4 Mon Sep 17 00:00:00 2001 From: "Yaxun (Sam) Liu" Date: Thu, 4 Jun 2020 13:50:16 -0400 Subject: [PATCH 18/33] Fix include path and wrapper header Currently std::complex and some other std functions require uses to include hip_runtime.h before any other headers to work, which is not reliable. changes are made in clang to fix this issue: https://reviews.llvm.org/D81176 which requires hipcc and HIP headers to make corresponding changes. This patch will make sure the clang change will not break HIP/ROCclr during this transition. After the transition is done, we can remove explicitly setting include path for HIP-Clang and HIP header in hipcc and hip config cmake files and rely on clang driver to set it automatically. Change-Id: I5d226861c2560ffa6c5ab17343a43cc378048061 --- bin/hipcc | 5 ++++- hip-config.cmake.in | 2 +- include/hip/hcc_detail/hip_fp16_math_fwd.h | 3 ++- include/hip/hcc_detail/hip_runtime.h | 6 ++++-- include/hip/hcc_detail/host_defines.h | 2 ++ include/hip/hcc_detail/math_functions.h | 5 ++++- include/hip/hcc_detail/math_fwd.h | 4 +++- 7 files changed, 20 insertions(+), 7 deletions(-) diff --git a/bin/hipcc b/bin/hipcc index c1e144300b..ad46e8e322 100755 --- a/bin/hipcc +++ b/bin/hipcc @@ -214,6 +214,7 @@ if ($HIP_PLATFORM eq "hcc" and $HIP_COMPILER eq "clang") { $HIP_LIB_PATH = "$HIP_PATH/lib"; } if ($verbose & 0x2) { + print ("ROCM_PATH=$ROCM_PATH\n"); if (defined $HIP_ROCCLR_HOME) { print ("HIP_ROCCLR_HOME=$HIP_ROCCLR_HOME\n"); } @@ -823,7 +824,9 @@ if ($HIP_PLATFORM eq "hcc" and $HIP_COMPILER eq "clang") { $HIPLDFLAGS .= " -mllvm -amdgpu-early-inline-all=true -mllvm -amdgpu-function-calls=false"; } } - $HIP_DEVLIB_FLAGS = " --hip-device-lib-path=$DEVICE_LIB_PATH"; + if ($DEVICE_LIB_PATH ne "$ROCM_PATH/amdgcn/bitcode") { + $HIP_DEVLIB_FLAGS = " --hip-device-lib-path=$DEVICE_LIB_PATH"; + } if ($hasHIP) { $HIPCXXFLAGS .= " $HIP_DEVLIB_FLAGS"; if ($HIP_RUNTIME ne "HCC") { diff --git a/hip-config.cmake.in b/hip-config.cmake.in index cc80e1ff4a..6ac076e279 100644 --- a/hip-config.cmake.in +++ b/hip-config.cmake.in @@ -150,7 +150,7 @@ if(HIP_COMPILER STREQUAL "clang") if (EXISTS ${AMD_DEVICE_LIBS_PREFIX}/amdgcn/bitcode) set_property(TARGET hip::device APPEND PROPERTY - INTERFACE_COMPILE_OPTIONS -x hip --hip-device-lib-path=${AMD_DEVICE_LIBS_PREFIX}/amdgcn/bitcode + INTERFACE_COMPILE_OPTIONS -x hip ) else() # This path is to support an older build of the device library diff --git a/include/hip/hcc_detail/hip_fp16_math_fwd.h b/include/hip/hcc_detail/hip_fp16_math_fwd.h index 95403e6ca8..16d834cca8 100644 --- a/include/hip/hcc_detail/hip_fp16_math_fwd.h +++ b/include/hip/hcc_detail/hip_fp16_math_fwd.h @@ -27,7 +27,7 @@ THE SOFTWARE. // */ #include "host_defines.h" - +#if !__CLANG_HIP_RUNTIME_WRAPPER_INCLUDED__ extern "C" { __device__ __attribute__((const)) _Float16 __ocml_ceil_f16(_Float16); @@ -82,3 +82,4 @@ extern "C" __device__ __attribute__((const)) __2f16 __ocml_sqrt_2f16(__2f16); __device__ __attribute__((const)) __2f16 __ocml_trunc_2f16(__2f16); } +#endif // !__CLANG_HIP_RUNTIME_WRAPPER_INCLUDED__ diff --git a/include/hip/hcc_detail/hip_runtime.h b/include/hip/hcc_detail/hip_runtime.h index a166935823..117af3e1fa 100644 --- a/include/hip/hcc_detail/hip_runtime.h +++ b/include/hip/hcc_detail/hip_runtime.h @@ -313,6 +313,7 @@ static constexpr Coordinates threadIdx{}; #endif // defined __HCC__ #if __HCC_OR_HIP_CLANG__ +#if !__CLANG_HIP_RUNTIME_WRAPPER_INCLUDED__ #if __HIP_ENABLE_DEVICE_MALLOC__ extern "C" __device__ void* __hip_malloc(size_t); extern "C" __device__ void* __hip_free(void* ptr); @@ -322,7 +323,7 @@ static inline __device__ void* free(void* ptr) { return __hip_free(ptr); } static inline __device__ void* malloc(size_t size) { __builtin_trap(); return nullptr; } static inline __device__ void* free(void* ptr) { __builtin_trap(); return nullptr; } #endif - +#endif // !__CLANG_HIP_RUNTIME_WRAPPER_INCLUDED__ #endif //__HCC_OR_HIP_CLANG__ #ifdef __HCC__ @@ -550,6 +551,7 @@ hc_get_workitem_absolute_id(int dim) #endif +#if !__CLANG_HIP_RUNTIME_WRAPPER_INCLUDED__ // Support std::complex. #if !_OPENMP || __HIP_ENABLE_CUDA_WRAPPER_FOR_OPENMP__ #pragma push_macro("__CUDA__") @@ -567,7 +569,7 @@ hc_get_workitem_absolute_id(int dim) #undef __CUDA__ #pragma pop_macro("__CUDA__") #endif // !_OPENMP || __HIP_ENABLE_CUDA_WRAPPER_FOR_OPENMP__ - +#endif // !__CLANG_HIP_RUNTIME_WRAPPER_INCLUDED__ #endif // defined(__clang__) && defined(__HIP__) #include diff --git a/include/hip/hcc_detail/host_defines.h b/include/hip/hcc_detail/host_defines.h index ad28cc7626..72f3932aff 100644 --- a/include/hip/hcc_detail/host_defines.h +++ b/include/hip/hcc_detail/host_defines.h @@ -64,11 +64,13 @@ THE SOFTWARE. #elif defined(__clang__) && defined(__HIP__) +#if !__CLANG_HIP_RUNTIME_WRAPPER_INCLUDED__ #define __host__ __attribute__((host)) #define __device__ __attribute__((device)) #define __global__ __attribute__((global)) #define __shared__ __attribute__((shared)) #define __constant__ __attribute__((constant)) +#endif // !__CLANG_HIP_RUNTIME_WRAPPER_INCLUDED__ #define __noinline__ __attribute__((noinline)) #define __forceinline__ inline __attribute__((always_inline)) diff --git a/include/hip/hcc_detail/math_functions.h b/include/hip/hcc_detail/math_functions.h index 494685e261..876a14eaa9 100644 --- a/include/hip/hcc_detail/math_functions.h +++ b/include/hip/hcc_detail/math_functions.h @@ -71,6 +71,7 @@ struct __numeric_type<_Float16> #define __RETURN_TYPE bool #endif +#if !__CLANG_HIP_RUNTIME_WRAPPER_INCLUDED__ __DEVICE__ inline uint64_t __make_mantissa_base8(const char* tagp) @@ -139,6 +140,7 @@ uint64_t __make_mantissa(const char* tagp) return __make_mantissa_base10(tagp); } +#endif // !__CLANG_HIP_RUNTIME_WRAPPER_INCLUDED__ // DOT FUNCTIONS #if (__hcc_workweek__ >= 19015) || __HIP_CLANG_ONLY__ @@ -174,6 +176,7 @@ uint amd_mixed_dot(uint a, uint b, uint c, bool saturate) { } #endif +#if !__CLANG_HIP_RUNTIME_WRAPPER_INCLUDED__ // BEGIN FLOAT __DEVICE__ inline @@ -1507,7 +1510,7 @@ __host__ inline static int min(int arg1, int arg2) { __host__ inline static int max(int arg1, int arg2) { return std::max(arg1, arg2); } - +#endif // !__CLANG_HIP_RUNTIME_WRAPPER_INCLUDED__ #pragma pop_macro("__DEF_FLOAT_FUN") #pragma pop_macro("__DEF_FLOAT_FUN2") diff --git a/include/hip/hcc_detail/math_fwd.h b/include/hip/hcc_detail/math_fwd.h index c25b5e90b4..4c0fde591c 100644 --- a/include/hip/hcc_detail/math_fwd.h +++ b/include/hip/hcc_detail/math_fwd.h @@ -23,7 +23,6 @@ THE SOFTWARE. #pragma once #include "host_defines.h" - #if defined(__cplusplus) extern "C" { #endif @@ -67,6 +66,7 @@ __attribute__((const)) unsigned int __ockl_udot8(unsigned int, unsigned int, unsigned int, bool); #endif +#if !__CLANG_HIP_RUNTIME_WRAPPER_INCLUDED__ // BEGIN FLOAT __device__ __attribute__((const)) @@ -701,6 +701,8 @@ double __llvm_amdgcn_rsq_f64(double) __asm("llvm.amdgcn.rsq.f64"); // END INTRINSICS // END DOUBLE +#endif // !__CLANG_HIP_RUNTIME_WRAPPER_INCLUDED__ + #if defined(__cplusplus) } // extern "C" #endif From c8f9afa9da799c8547d29c74fdd50c0d6f375fae Mon Sep 17 00:00:00 2001 From: Christophe Paquot Date: Tue, 9 Jun 2020 10:07:58 -0700 Subject: [PATCH 19/33] Do not deferred stream creation now that we multiplex HW queues SWDEV-239856 Change-Id: I156650faf832f86891f00ee167269509edd844ec --- rocclr/hip_stream.cpp | 2 +- tests/src/runtimeApi/stream/hipStreamCreateWithPriority.cpp | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) diff --git a/rocclr/hip_stream.cpp b/rocclr/hip_stream.cpp index e2ab7b18cd..d764e1a5fe 100755 --- a/rocclr/hip_stream.cpp +++ b/rocclr/hip_stream.cpp @@ -189,7 +189,7 @@ static hipError_t ihipStreamCreate(hipStream_t* stream, const std::vector& cuMask = {}) { hip::Stream* hStream = new hip::Stream(hip::getCurrentDevice(), priority, flags, false, cuMask); - if (hStream == nullptr) { + if (hStream == nullptr || !hStream->Create()) { return hipErrorOutOfMemory; } diff --git a/tests/src/runtimeApi/stream/hipStreamCreateWithPriority.cpp b/tests/src/runtimeApi/stream/hipStreamCreateWithPriority.cpp index ba2b140611..11cd8d95b5 100644 --- a/tests/src/runtimeApi/stream/hipStreamCreateWithPriority.cpp +++ b/tests/src/runtimeApi/stream/hipStreamCreateWithPriority.cpp @@ -204,7 +204,7 @@ void runTest() // validate that stream priorities are working as expected #define OP(x, y) \ if (enable_priority_##x && enable_priority_##y) { \ - if (time_spent_##x < time_spent_##y) { printf("FAILED!"); exit(-1); } \ + if ((1.05f * time_spent_##x) < time_spent_##y) { printf("FAILED!"); exit(-1); } \ } OP(low, normal) OP(normal, high) From 6ed1868203f9850f2d90e15f3d9228998c882517 Mon Sep 17 00:00:00 2001 From: Dittakavi Satyanvesh Date: Fri, 5 Jun 2020 05:16:25 -0400 Subject: [PATCH 20/33] SWDEV-236670 Address Eigen unit test failure by adding __host__ attribute to half2 functions Change-Id: Ifdc852c30a1b3704871e0ee58cb7a55d3d37fc6e --- include/hip/hcc_detail/hip_fp16.h | 130 +++++++++++++++--------------- 1 file changed, 65 insertions(+), 65 deletions(-) diff --git a/include/hip/hcc_detail/hip_fp16.h b/include/hip/hcc_detail/hip_fp16.h index 6fa86e94b9..670cd475c2 100644 --- a/include/hip/hcc_detail/hip_fp16.h +++ b/include/hip/hcc_detail/hip_fp16.h @@ -529,35 +529,35 @@ THE SOFTWARE. } inline - __device__ + __host__ __device__ __half __low2half(__half2 x) { return __half{__half_raw{static_cast<__half2_raw>(x).data.x}}; } inline - __device__ + __host__ __device__ __half __high2half(__half2 x) { return __half{__half_raw{static_cast<__half2_raw>(x).data.y}}; } inline - __device__ + __host__ __device__ __half2 __half2half2(__half x) { return __half2{x, x}; } inline - __device__ + __host__ __device__ __half2 __halves2half2(__half x, __half y) { return __half2{x, y}; } inline - __device__ + __host__ __device__ __half2 __low2half2(__half2 x) { return __half2{ @@ -567,7 +567,7 @@ THE SOFTWARE. } inline - __device__ + __host__ __device__ __half2 __high2half2(__half2 x) { return __half2_raw{ @@ -577,7 +577,7 @@ THE SOFTWARE. } inline - __device__ + __host__ __device__ __half2 __lows2half2(__half2 x, __half2 y) { return __half2_raw{ @@ -587,7 +587,7 @@ THE SOFTWARE. } inline - __device__ + __host__ __device__ __half2 __highs2half2(__half2 x, __half2 y) { return __half2_raw{ @@ -597,7 +597,7 @@ THE SOFTWARE. } inline - __device__ + __host__ __device__ __half2 __lowhigh2highlow(__half2 x) { return __half2_raw{ @@ -1046,16 +1046,16 @@ THE SOFTWARE. __half __ldcs(const __half* ptr) { return *ptr; } inline - __device__ + __host__ __device__ __half2 __ldg(const __half2* ptr) { return *ptr; } inline - __device__ + __host__ __device__ __half2 __ldcg(const __half2* ptr) { return *ptr; } inline - __device__ + __host__ __device__ __half2 __ldca(const __half2* ptr) { return *ptr; } inline - __device__ + __host__ __device__ __half2 __ldcs(const __half2* ptr) { return *ptr; } // Relations @@ -1121,7 +1121,7 @@ THE SOFTWARE. bool __hgtu(__half x, __half y) { return __hgt(x, y); } inline - __device__ + __host__ __device__ __half2 __heq2(__half2 x, __half2 y) { auto r = static_cast<__half2_raw>(x).data == @@ -1129,7 +1129,7 @@ THE SOFTWARE. return __builtin_convertvector(-r, _Float16_2); } inline - __device__ + __host__ __device__ __half2 __hne2(__half2 x, __half2 y) { auto r = static_cast<__half2_raw>(x).data != @@ -1137,7 +1137,7 @@ THE SOFTWARE. return __builtin_convertvector(-r, _Float16_2); } inline - __device__ + __host__ __device__ __half2 __hle2(__half2 x, __half2 y) { auto r = static_cast<__half2_raw>(x).data <= @@ -1145,7 +1145,7 @@ THE SOFTWARE. return __builtin_convertvector(-r, _Float16_2); } inline - __device__ + __host__ __device__ __half2 __hge2(__half2 x, __half2 y) { auto r = static_cast<__half2_raw>(x).data >= @@ -1153,7 +1153,7 @@ THE SOFTWARE. return __builtin_convertvector(-r, _Float16_2); } inline - __device__ + __host__ __device__ __half2 __hlt2(__half2 x, __half2 y) { auto r = static_cast<__half2_raw>(x).data < @@ -1161,7 +1161,7 @@ THE SOFTWARE. return __builtin_convertvector(-r, _Float16_2); } inline - __device__ + __host__ __device__ __half2 __hgt2(__half2 x, __half2 y) { auto r = static_cast<__half2_raw>(x).data > @@ -1169,83 +1169,83 @@ THE SOFTWARE. return __builtin_convertvector(-r, _Float16_2); } inline - __device__ + __host__ __device__ __half2 __hequ2(__half2 x, __half2 y) { return __heq2(x, y); } inline - __device__ + __host__ __device__ __half2 __hneu2(__half2 x, __half2 y) { return __hne2(x, y); } inline - __device__ + __host__ __device__ __half2 __hleu2(__half2 x, __half2 y) { return __hle2(x, y); } inline - __device__ + __host__ __device__ __half2 __hgeu2(__half2 x, __half2 y) { return __hge2(x, y); } inline - __device__ + __host__ __device__ __half2 __hltu2(__half2 x, __half2 y) { return __hlt2(x, y); } inline - __device__ + __host__ __device__ __half2 __hgtu2(__half2 x, __half2 y) { return __hgt2(x, y); } inline - __device__ + __host__ __device__ bool __hbeq2(__half2 x, __half2 y) { auto r = static_cast<__half2_raw>(__heq2(x, y)); return r.data.x != 0 && r.data.y != 0; } inline - __device__ + __host__ __device__ bool __hbne2(__half2 x, __half2 y) { auto r = static_cast<__half2_raw>(__hne2(x, y)); return r.data.x != 0 && r.data.y != 0; } inline - __device__ + __host__ __device__ bool __hble2(__half2 x, __half2 y) { auto r = static_cast<__half2_raw>(__hle2(x, y)); return r.data.x != 0 && r.data.y != 0; } inline - __device__ + __host__ __device__ bool __hbge2(__half2 x, __half2 y) { auto r = static_cast<__half2_raw>(__hge2(x, y)); return r.data.x != 0 && r.data.y != 0; } inline - __device__ + __host__ __device__ bool __hblt2(__half2 x, __half2 y) { auto r = static_cast<__half2_raw>(__hlt2(x, y)); return r.data.x != 0 && r.data.y != 0; } inline - __device__ + __host__ __device__ bool __hbgt2(__half2 x, __half2 y) { auto r = static_cast<__half2_raw>(__hgt2(x, y)); return r.data.x != 0 && r.data.y != 0; } inline - __device__ + __host__ __device__ bool __hbequ2(__half2 x, __half2 y) { return __hbeq2(x, y); } inline - __device__ + __host__ __device__ bool __hbneu2(__half2 x, __half2 y) { return __hbne2(x, y); } inline - __device__ + __host__ __device__ bool __hbleu2(__half2 x, __half2 y) { return __hble2(x, y); } inline - __device__ + __host__ __device__ bool __hbgeu2(__half2 x, __half2 y) { return __hbge2(x, y); } inline - __device__ + __host__ __device__ bool __hbltu2(__half2 x, __half2 y) { return __hblt2(x, y); } inline - __device__ + __host__ __device__ bool __hbgtu2(__half2 x, __half2 y) { return __hbgt2(x, y); } // Arithmetic @@ -1334,7 +1334,7 @@ THE SOFTWARE. } inline - __device__ + __host__ __device__ __half2 __hadd2(__half2 x, __half2 y) { return __half2_raw{ @@ -1342,14 +1342,14 @@ THE SOFTWARE. static_cast<__half2_raw>(y).data}; } inline - __device__ + __host__ __device__ __half2 __habs2(__half2 x) { return __half2_raw{ __ocml_fabs_2f16(static_cast<__half2_raw>(x).data)}; } inline - __device__ + __host__ __device__ __half2 __hsub2(__half2 x, __half2 y) { return __half2_raw{ @@ -1357,7 +1357,7 @@ THE SOFTWARE. static_cast<__half2_raw>(y).data}; } inline - __device__ + __host__ __device__ __half2 __hmul2(__half2 x, __half2 y) { return __half2_raw{ @@ -1365,7 +1365,7 @@ THE SOFTWARE. static_cast<__half2_raw>(y).data}; } inline - __device__ + __host__ __device__ __half2 __hadd2_sat(__half2 x, __half2 y) { auto r = static_cast<__half2_raw>(__hadd2(x, y)); @@ -1374,7 +1374,7 @@ THE SOFTWARE. __clamp_01(__half_raw{r.data.y})}; } inline - __device__ + __host__ __device__ __half2 __hsub2_sat(__half2 x, __half2 y) { auto r = static_cast<__half2_raw>(__hsub2(x, y)); @@ -1383,7 +1383,7 @@ THE SOFTWARE. __clamp_01(__half_raw{r.data.y})}; } inline - __device__ + __host__ __device__ __half2 __hmul2_sat(__half2 x, __half2 y) { auto r = static_cast<__half2_raw>(__hmul2(x, y)); @@ -1392,13 +1392,13 @@ THE SOFTWARE. __clamp_01(__half_raw{r.data.y})}; } inline - __device__ + __host__ __device__ __half2 __hfma2(__half2 x, __half2 y, __half2 z) { return __half2_raw{__ocml_fma_2f16(x, y, z)}; } inline - __device__ + __host__ __device__ __half2 __hfma2_sat(__half2 x, __half2 y, __half2 z) { auto r = static_cast<__half2_raw>(__hfma2(x, y, z)); @@ -1407,7 +1407,7 @@ THE SOFTWARE. __clamp_01(__half_raw{r.data.y})}; } inline - __device__ + __host__ __device__ __half2 __h2div(__half2 x, __half2 y) { return __half2_raw{ @@ -1550,82 +1550,82 @@ THE SOFTWARE. } inline - __device__ + __host__ __device__ __half2 h2trunc(__half2 x) { return __half2_raw{__ocml_trunc_2f16(x)}; } inline - __device__ + __host__ __device__ __half2 h2ceil(__half2 x) { return __half2_raw{__ocml_ceil_2f16(x)}; } inline - __device__ + __host__ __device__ __half2 h2floor(__half2 x) { return __half2_raw{__ocml_floor_2f16(x)}; } inline - __device__ + __host__ __device__ __half2 h2rint(__half2 x) { return __half2_raw{__ocml_rint_2f16(x)}; } inline - __device__ + __host__ __device__ __half2 h2sin(__half2 x) { return __half2_raw{__ocml_sin_2f16(x)}; } inline - __device__ + __host__ __device__ __half2 h2cos(__half2 x) { return __half2_raw{__ocml_cos_2f16(x)}; } inline - __device__ + __host__ __device__ __half2 h2exp(__half2 x) { return __half2_raw{__ocml_exp_2f16(x)}; } inline - __device__ + __host__ __device__ __half2 h2exp2(__half2 x) { return __half2_raw{__ocml_exp2_2f16(x)}; } inline - __device__ + __host__ __device__ __half2 h2exp10(__half2 x) { return __half2_raw{__ocml_exp10_2f16(x)}; } inline - __device__ + __host__ __device__ __half2 h2log2(__half2 x) { return __half2_raw{__ocml_log2_2f16(x)}; } inline - __device__ + __host__ __device__ __half2 h2log(__half2 x) { return __ocml_log_2f16(x); } inline - __device__ + __host__ __device__ __half2 h2log10(__half2 x) { return __ocml_log10_2f16(x); } inline - __device__ + __host__ __device__ __half2 h2rcp(__half2 x) { return __llvm_amdgcn_rcp_2f16(x); } inline - __device__ + __host__ __device__ __half2 h2rsqrt(__half2 x) { return __ocml_rsqrt_2f16(x); } inline - __device__ + __host__ __device__ __half2 h2sqrt(__half2 x) { return __ocml_sqrt_2f16(x); } inline - __device__ + __host__ __device__ __half2 __hisinf2(__half2 x) { auto r = __ocml_isinf_2f16(x); @@ -1633,7 +1633,7 @@ THE SOFTWARE. static_cast<_Float16>(r.x), static_cast<_Float16>(r.y)}}; } inline - __device__ + __host__ __device__ __half2 __hisnan2(__half2 x) { auto r = __ocml_isnan_2f16(x); @@ -1641,7 +1641,7 @@ THE SOFTWARE. static_cast<_Float16>(r.x), static_cast<_Float16>(r.y)}}; } inline - __device__ + __host__ __device__ __half2 __hneg2(__half2 x) { return __half2_raw{-static_cast<__half2_raw>(x).data}; From 2f8e180f013532c1609fb04071289b9116c809f3 Mon Sep 17 00:00:00 2001 From: kjayapra-amd Date: Tue, 9 Jun 2020 20:08:37 -0400 Subject: [PATCH 21/33] SWDEV-231701 - Remove amd::memory->svm_ptr from MemObj, instead of the ptr to the object. Change-Id: I5aab450a2320cfa5417c284e2a8454102df6f99d --- rocclr/hip_memory.cpp | 7 ++++++- 1 file changed, 6 insertions(+), 1 deletion(-) diff --git a/rocclr/hip_memory.cpp b/rocclr/hip_memory.cpp index 9af33d9c74..84c8b77f92 100755 --- a/rocclr/hip_memory.cpp +++ b/rocclr/hip_memory.cpp @@ -2026,7 +2026,12 @@ hipError_t hipIpcCloseMemHandle(void* dev_ptr) { } /* Remove the memory from MemObjMap */ - amd::MemObjMap::RemoveMemObj(amd_mem_obj); + if (amd_mem_obj->getSvmPtr() != nullptr) { + amd::MemObjMap::RemoveMemObj(amd_mem_obj->getSvmPtr()); + } else { + DevLogPrintfError("Does not have SVM or Host Mem for 0x%x, crash here!", dev_ptr); + guarantee(false); + } /* detach the memory */ if (!device->IpcDetach(*amd_mem_obj)){ From c54590ebbd3afc787888ddd02da753f22150f4c6 Mon Sep 17 00:00:00 2001 From: jujiang Date: Wed, 10 Jun 2020 13:25:00 -0400 Subject: [PATCH 22/33] Fix for SWDEV-239415 to handle hipGetDevice properly while no GPU present Change-Id: I252cbbf9a89fc76fe1be1fbb8f45778e96c70fb2 --- rocclr/hip_device_runtime.cpp | 6 +++++- 1 file changed, 5 insertions(+), 1 deletion(-) diff --git a/rocclr/hip_device_runtime.cpp b/rocclr/hip_device_runtime.cpp index 9ed1c64145..414d51f513 100644 --- a/rocclr/hip_device_runtime.cpp +++ b/rocclr/hip_device_runtime.cpp @@ -449,7 +449,11 @@ hipError_t hipDeviceSynchronize ( void ) { } int ihipGetDevice() { - return hip::getCurrentDevice()->deviceId(); + hip::Device* device = hip::getCurrentDevice(); + if(device == nullptr){ + return -1; + } + return device->deviceId(); } hipError_t hipGetDevice ( int* deviceId ) { From 85e3b37052227f4db39551000f87aba72df85d90 Mon Sep 17 00:00:00 2001 From: Siu Chi Chan Date: Wed, 10 Jun 2020 16:48:44 +0000 Subject: [PATCH 23/33] fix uninitialized value in hipcc Change-Id: I90b070c491f0efc328fcf97de0e111658ec772de --- bin/hipcc | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/bin/hipcc b/bin/hipcc index ad46e8e322..91d9a197da 100755 --- a/bin/hipcc +++ b/bin/hipcc @@ -824,11 +824,11 @@ if ($HIP_PLATFORM eq "hcc" and $HIP_COMPILER eq "clang") { $HIPLDFLAGS .= " -mllvm -amdgpu-early-inline-all=true -mllvm -amdgpu-function-calls=false"; } } - if ($DEVICE_LIB_PATH ne "$ROCM_PATH/amdgcn/bitcode") { - $HIP_DEVLIB_FLAGS = " --hip-device-lib-path=$DEVICE_LIB_PATH"; - } + if ($hasHIP) { - $HIPCXXFLAGS .= " $HIP_DEVLIB_FLAGS"; + if ($DEVICE_LIB_PATH ne "$ROCM_PATH/amdgcn/bitcode") { + $HIPCXXFLAGS .= " --hip-device-lib-path=$DEVICE_LIB_PATH"; + } if ($HIP_RUNTIME ne "HCC") { $HIPCXXFLAGS .= " -fhip-new-launch-api"; } From 20f05c4228a73684ee3b3fb938c2da2684fe1bae Mon Sep 17 00:00:00 2001 From: kjayapra-amd Date: Mon, 18 May 2020 22:40:33 -0400 Subject: [PATCH 24/33] SWDEV-236178 - Reorganizing Platform/Modules code for easy access. Change-Id: Ie8920260ffc4ff01e44b48af8cec9ea5aed1aa9b --- rocclr/CMakeLists.txt | 2 + rocclr/hip_code_object.cpp | 458 ++++++++++++ rocclr/hip_code_object.hpp | 132 ++++ rocclr/hip_context.cpp | 1 + rocclr/hip_fatbin.hpp | 29 + rocclr/hip_global.cpp | 202 ++++++ rocclr/hip_global.hpp | 116 +++ rocclr/hip_intercept.cpp | 6 +- rocclr/hip_internal.hpp | 165 +---- rocclr/hip_memory.cpp | 53 +- rocclr/hip_module.cpp | 237 +----- rocclr/hip_platform.cpp | 835 ++++++---------------- rocclr/hip_platform.hpp | 66 +- rocclr/hip_texture.cpp | 54 +- tests/src/runtimeApi/module/hipModule.cpp | 3 +- 15 files changed, 1312 insertions(+), 1047 deletions(-) mode change 100644 => 100755 rocclr/CMakeLists.txt create mode 100755 rocclr/hip_code_object.cpp create mode 100755 rocclr/hip_code_object.hpp create mode 100755 rocclr/hip_fatbin.hpp create mode 100755 rocclr/hip_global.cpp create mode 100755 rocclr/hip_global.hpp mode change 100644 => 100755 rocclr/hip_intercept.cpp mode change 100644 => 100755 rocclr/hip_platform.hpp mode change 100644 => 100755 tests/src/runtimeApi/module/hipModule.cpp diff --git a/rocclr/CMakeLists.txt b/rocclr/CMakeLists.txt old mode 100644 new mode 100755 index 37135af81d..252ac5abfa --- a/rocclr/CMakeLists.txt +++ b/rocclr/CMakeLists.txt @@ -75,10 +75,12 @@ add_definitions(-DBSD_LIBELF) add_library(hip64 OBJECT hip_context.cpp + hip_code_object.cpp hip_device.cpp hip_device_runtime.cpp hip_error.cpp hip_event.cpp + hip_global.cpp hip_memory.cpp hip_module.cpp hip_peer.cpp diff --git a/rocclr/hip_code_object.cpp b/rocclr/hip_code_object.cpp new file mode 100755 index 0000000000..bb957b8cf3 --- /dev/null +++ b/rocclr/hip_code_object.cpp @@ -0,0 +1,458 @@ +#include "hip_code_object.hpp" + +#include + +#include "hip/hip_runtime_api.h" +#include "hip/hip_runtime.h" +#include "hip_internal.hpp" +#include "platform/program.hpp" + +namespace hip { + +uint64_t CodeObject::ElfSize(const void *emi) { + const Elf64_Ehdr *ehdr = (const Elf64_Ehdr*)emi; + const Elf64_Shdr *shdr = (const Elf64_Shdr*)((char*)emi + ehdr->e_shoff); + + uint64_t max_offset = ehdr->e_shoff; + uint64_t total_size = max_offset + ehdr->e_shentsize * ehdr->e_shnum; + + for (uint16_t i=0; i < ehdr->e_shnum; ++i){ + uint64_t cur_offset = static_cast(shdr[i].sh_offset); + if (max_offset < cur_offset) { + max_offset = cur_offset; + total_size = max_offset; + if(SHT_NOBITS != shdr[i].sh_type) { + total_size += static_cast(shdr[i].sh_size); + } + } + } + return total_size; +} + +bool CodeObject::isCompatibleCodeObject(const std::string& codeobj_target_id, + const char* device_name) { + // Workaround for device name mismatch. + // Device name may contain feature strings delimited by '+', e.g. + // gfx900+xnack. Currently HIP-Clang does not include feature strings + // in code object target id in fat binary. Therefore drop the feature + // strings from device name before comparing it with code object target id. + std::string short_name(device_name); + auto feature_loc = short_name.find('+'); + if (feature_loc != std::string::npos) { + short_name.erase(feature_loc); + } + return codeobj_target_id == short_name; +} + +hipError_t CodeObject::extractCodeObjectFromFatBinary(const void* data, + const std::vector& devices, + std::vector>& code_objs) { + std::string magic((const char*)data, sizeof(CLANG_OFFLOAD_BUNDLER_MAGIC_STR) - 1); + if (magic.compare(CLANG_OFFLOAD_BUNDLER_MAGIC_STR)) { + return hipErrorInvalidKernelFile; + } + + code_objs.resize(devices.size()); + const auto obheader = reinterpret_cast(data); + const auto* desc = &obheader->desc[0]; + unsigned num_code_objs = 0; + for (uint64_t i = 0; i < obheader->numBundles; ++i, + desc = reinterpret_cast( + reinterpret_cast(&desc->triple[0]) + desc->tripleSize)) { + + std::size_t offset = 0; + if (!std::strncmp(desc->triple, HIP_AMDGCN_AMDHSA_TRIPLE, + sizeof(HIP_AMDGCN_AMDHSA_TRIPLE) - 1)) { + offset = sizeof(HIP_AMDGCN_AMDHSA_TRIPLE); //For code objects created by CLang + } else if (!std::strncmp(desc->triple, HCC_AMDGCN_AMDHSA_TRIPLE, + sizeof(HCC_AMDGCN_AMDHSA_TRIPLE) - 1)) { + offset = sizeof(HCC_AMDGCN_AMDHSA_TRIPLE); //For code objects created by Hcc + } else { + continue; + } + std::string target(desc->triple + offset, desc->tripleSize - offset); + + const void *image = reinterpret_cast( + reinterpret_cast(obheader) + desc->offset); + size_t size = desc->size; + + for (size_t dev = 0; dev < devices.size(); ++dev) { + const char* name = devices[dev]; + + if (!isCompatibleCodeObject(target, name)) { + continue; + } + code_objs[dev] = std::make_pair(image, size); + num_code_objs++; + } + } + if (num_code_objs == devices.size()) { + return hipSuccess; + } else { + DevLogError("hipErrorNoBinaryForGpu: Coudn't find binary for current devices!"); + guarantee(false); + return hipErrorNoBinaryForGpu; + } +} + +hipError_t CodeObject::add_program(int deviceId, hipModule_t hmod, const void* binary_ptr, + size_t binary_size) { + amd::Program* program = as_amd(reinterpret_cast(hmod)); + amd::Context* ctx = g_devices[deviceId]->asContext(); + if (CL_SUCCESS != program->addDeviceProgram(*ctx->devices()[0], binary_ptr, + binary_size, false)) { + return hipErrorNotFound; + } + return hipSuccess; +} + +hipError_t CodeObject::build_module(hipModule_t hmod, const std::vector& devices) { + amd::Program* program = as_amd(reinterpret_cast(hmod)); + program->setVarInfoCallBack(&getSvarInfo); + if (CL_SUCCESS != program->build(devices, nullptr, nullptr, nullptr, kOptionChangeable, kNewDevProg)) { + DevLogPrintfError("Build error for module: 0x%x \n", hmod); + return hipErrorSharedObjectInitFailed; + } + return hipSuccess; +} + +DynCO::DynCO(): program_(nullptr) {} + +hipError_t DynCO::loadCodeObject(const char* fname, const void* image) { + + amd::ScopedLock lock(dclock_); + + const void *mmap_ptr = nullptr; + size_t mmap_size = 0; + + guarantee(fname || image); + if (fname != nullptr) { + /* We are given file name */ + + if (!amd::Os::MemoryMapFile(fname, &mmap_ptr, &mmap_size)) { + return hipErrorFileNotFound; + } + } else if (image != nullptr) { + /*We are directly given image pointer directly */ + mmap_ptr = image; + } else { + return hipErrorMissingConfiguration; + } + + return loadCodeObjectData(mmap_ptr, mmap_size); +} + +//Dynamic Code Object +DynCO::~DynCO() { + amd::ScopedLock lock(dclock_); + + if (program_ != nullptr) { + program_->release(); + program_ = nullptr; + } + + for (auto& elem : vars_) { + delete elem.second; + } + vars_.clear(); + + for (auto& elem : functions_) { + delete elem.second; + } + functions_.clear(); +} + +hipError_t DynCO::getDeviceVar(DeviceVar** dvar, std::string var_name, int device_id) { + amd::ScopedLock lock(dclock_); + + auto it = vars_.find(var_name); + if (it == vars_.end()) { + DevLogPrintfError("Cannot find the Var: %s ", var_name.c_str()); + return hipErrorNotFound; + } + + it->second->getDeviceVar(dvar, device_id, module()); + return hipSuccess; +} + +hipError_t DynCO::getDynFunc(hipFunction_t* hfunc, std::string func_name) { + amd::ScopedLock lock(dclock_); + + auto it = functions_.find(func_name); + if (it == functions_.end()) { + DevLogPrintfError("Cannot find the function: %s ", func_name.c_str()); + return hipErrorNotFound; + } + + /* See if this could be solved */ + return it->second->getDynFunc(hfunc, reinterpret_cast(as_cl(program_))); +} + +hipError_t DynCO::loadCodeObjectData(const void* mmap_ptr, size_t mmap_size) { + + amd::ScopedLock lock(dclock_); + + /* initialize image it to the mmap_ptr, if this is of no_clang_offload + bundle then they directly pass the image */ + const void* image = mmap_ptr; + std::vector> code_objs; + hipError_t hip_error = extractCodeObjectFromFatBinary(mmap_ptr, + {hip::getCurrentDevice()->devices()[0]->info().name_}, + code_objs); + if (hip_error == hipSuccess) { + image = code_objs[0].first; + } else if(hip_error == hipErrorNoBinaryForGpu) { + return hip_error; + } + + program_ = new amd::Program(*hip::getCurrentDevice()->asContext(), + amd::Program::Language::Binary, mmap_ptr, mmap_size); + if (program_ == NULL) { + return hipErrorOutOfMemory; + } + + program_->setVarInfoCallBack(&getSvarInfo); + + if (CL_SUCCESS != program_->addDeviceProgram(*hip::getCurrentDevice()->devices()[0], image, + ElfSize(image), false)) { + return hipErrorInvalidKernelFile; + } + + //This has to happen before Program has been built, other wise undef vars fail. + IHIP_RETURN_ONFAIL(populateDynGlobalVars()); + + //program->setVarInfoCallBack(&getSvarInfo); + if(CL_SUCCESS != program_->build(hip::getCurrentDevice()->devices(), nullptr, nullptr, nullptr, + kOptionChangeable, kNewDevProg)) { + return hipErrorSharedObjectInitFailed; + } + + //This has to happen after Program has been built, other wise symbolTable_ not populated. + IHIP_RETURN_ONFAIL(populateDynGlobalFuncs()); + + return hipSuccess; +} + +hipError_t DynCO::populateDynGlobalVars() { + amd::ScopedLock lock(dclock_); + + std::vector var_names; + std::vector undef_var_names; + + device::Program* dev_program + = program_->getDeviceProgram(*hip::getCurrentDevice()->devices()[0]); + + if (!dev_program->getGlobalVarFromCodeObj(&var_names)) { + DevLogPrintfError("Could not get Global vars from Code Obj for Module: 0x%x \n", module()); + return hipErrorSharedObjectSymbolNotFound; + } + + if (!dev_program->getUndefinedVarFromCodeObj(&undef_var_names)) { + DevLogPrintfError("Could not get undefined Variables for Module: 0x%x \n", module()); + return hipErrorSharedObjectSymbolNotFound; + } + + for (auto& elem : var_names) { + vars_.insert(std::make_pair(elem, new Var(elem, Var::DeviceVarKind::DVK_Variable, 0, 0, 0, nullptr))); + } + + for (auto& elem : undef_var_names) { + vars_.insert(std::make_pair(elem, new Var(elem, Var::DeviceVarKind::DVK_Texture, 0, 0, 0, nullptr))); + } + + return hipSuccess; +} + +hipError_t DynCO::populateDynGlobalFuncs() { + amd::ScopedLock lock(dclock_); + + std::vector func_names; + device::Program* dev_program + = program_->getDeviceProgram(*hip::getCurrentDevice()->devices()[0]); + + // Get all the global func names from COMGR + if (!dev_program->getGlobalFuncFromCodeObj(&func_names)) { + DevLogPrintfError("Could not get Global Funcs from Code Obj for Module: 0x%x \n", module()); + return hipErrorSharedObjectSymbolNotFound; + } + + for (auto& elem : func_names) { + functions_.insert(std::make_pair(elem, new Function(elem))); + } + + return hipSuccess; +} + +//Static Code Object +StatCO::StatCO() { +} + +StatCO::~StatCO() { + amd::ScopedLock lock(sclock_); + + for (auto& elem : functions_) { + delete elem.second; + } + functions_.clear(); + + for (auto& elem : vars_) { + delete elem.second; + } + vars_.clear(); +} + +hipError_t StatCO::digestFatBinary(const void* data, FatBinaryInfoType& programs) { + amd::ScopedLock lock(sclock_); + + if (programs.size() > 0) { + return hipSuccess; + } + + std::vector> code_objs; + std::vector devices; + for (size_t dev = 0; dev < g_devices.size(); ++dev) { + devices.push_back(g_devices[dev]->devices()[0]->info().name_); + } + + IHIP_RETURN_ONFAIL(extractCodeObjectFromFatBinary((char*)data, devices, code_objs)); + programs.resize(g_devices.size()); + + for (size_t dev = 0; dev < g_devices.size(); ++dev) { + amd::Context* ctx = g_devices[dev]->asContext(); + amd::Program* program = new amd::Program(*ctx); + if (program == nullptr) { + return hipErrorOutOfMemory; + } + programs.at(dev) = std::make_pair(reinterpret_cast(as_cl(program)), + new FatBinaryMetaInfo(false, code_objs[dev].first, code_objs[dev].second)); + } + + return hipSuccess; +} + +FatBinaryInfoType* StatCO::addFatBinary(const void* data, bool initialized) { + amd::ScopedLock lock(sclock_); + + if (initialized) { + digestFatBinary(data, modules_[data]); + } + + return &modules_[data]; +} + +hipError_t StatCO::removeFatBinary(FatBinaryInfoType* module) { + amd::ScopedLock lock(sclock_); + + auto vit = vars_.begin(); + while (vit != vars_.end()) { + if (vit->second->moduleInfo() == module) { + delete vit->second; + vit = vars_.erase(vit); + } else { + ++vit; + } + } + + auto fit = functions_.begin(); + while (fit != functions_.end()) { + if (fit->second->moduleInfo() == module) { + delete fit->second; + fit = functions_.erase(fit); + } else { + ++fit; + } + } + + auto mit = modules_.begin(); + while (mit != modules_.end()) { + if (&mit->second == module) { + for (size_t dev=0; dev < g_devices.size(); ++dev) { + delete (*module)[dev].second; + } + mit = modules_.erase(mit); + } else { + ++mit; + } + } + + return hipSuccess; +} + +hipError_t StatCO::registerStatFunction(const void* hostFunction, Function* func) { + amd::ScopedLock lock(sclock_); + + if (functions_.find(hostFunction) != functions_.end()) { + DevLogPrintfError("hostFunctionPtr: 0x%x already exists", hostFunction); + guarantee(false); + } + functions_.insert(std::make_pair(hostFunction, func)); + + return hipSuccess; +} + +hipError_t StatCO::getStatFunc(hipFunction_t* hfunc, const void* hostFunction, int deviceId) { + amd::ScopedLock lock(sclock_); + + const auto it = functions_.find(hostFunction); + if (it == functions_.end()) { + return hipErrorInvalidSymbol; + } + + return it->second->getStatFunc(hfunc, deviceId); +} + +hipError_t StatCO::getStatFuncAttr(hipFuncAttributes* func_attr, const void* hostFunction, int deviceId) { + amd::ScopedLock lock(sclock_); + + const auto it = functions_.find(hostFunction); + if (it == functions_.end()) { + return hipErrorInvalidSymbol; + } + + return it->second->getStatFuncAttr(func_attr, deviceId); +} + +hipError_t StatCO::registerStatGlobalVar(const void* hostVar, Var* var) { + amd::ScopedLock lock(sclock_); + + if (vars_.find(hostVar) != vars_.end()) { + return hipErrorInvalidSymbol; + } + + vars_.insert(std::make_pair(hostVar, var)); + return hipSuccess; +} + +hipError_t StatCO::getStatGlobalVar(const void* hostVar, int deviceId, hipDeviceptr_t* dev_ptr, + size_t* size_ptr) { + amd::ScopedLock lock(sclock_); + + const auto it = vars_.find(hostVar); + if (it == vars_.end()) { + return hipErrorInvalidSymbol; + } + + DeviceVar* dvar = nullptr; + IHIP_RETURN_ONFAIL(it->second->getStatDeviceVar(&dvar, deviceId)); + + *dev_ptr = dvar->device_ptr(); + *size_ptr = dvar->size(); + return hipSuccess; +} + +hipError_t StatCO::getStatGlobalVarByName(std::string hostVar, int deviceId, hipModule_t hmod, + hipDeviceptr_t* dev_ptr, size_t* size_ptr) { + amd::ScopedLock lock(sclock_); + + for (auto& elem : vars_) { + if ((elem.second->name() == hostVar) + && (elem.second->module(deviceId) == hmod)) { + *dev_ptr = elem.second->device_ptr(deviceId); + *size_ptr = elem.second->device_size(deviceId); + return hipSuccess; + } + } + + return hipErrorNotFound; +} +}; //namespace: hip diff --git a/rocclr/hip_code_object.hpp b/rocclr/hip_code_object.hpp new file mode 100755 index 0000000000..38d95c0e27 --- /dev/null +++ b/rocclr/hip_code_object.hpp @@ -0,0 +1,132 @@ +#ifndef HIP_CODE_OBJECT_HPP +#define HIP_CODE_OBJECT_HPP + +#include "hip_global.hpp" + +#include + +#include "hip/hip_runtime.h" +#include "hip/hip_runtime_api.h" +#include "hip_internal.hpp" +#include "device/device.hpp" +#include "platform/program.hpp" + +//Forward Declaration for friend usage +class PlatformState; + +namespace hip { + +//Code Object base class +class CodeObject { +public: + virtual ~CodeObject() {} + + //Functions to add_dev_prog and build + static hipError_t add_program(int deviceId, hipModule_t hmod, const void* binary_ptr, + size_t binary_size); + static hipError_t build_module(hipModule_t hmod, const std::vector& devices); + + //ClangOFFLOADBundle info + #define CLANG_OFFLOAD_BUNDLER_MAGIC_STR "__CLANG_OFFLOAD_BUNDLE__" + #define HIP_AMDGCN_AMDHSA_TRIPLE "hip-amdgcn-amd-amdhsa" + #define HCC_AMDGCN_AMDHSA_TRIPLE "hcc-amdgcn-amd-amdhsa-" + + //Clang Offload bundler description & Header + struct __ClangOffloadBundleDesc { + uint64_t offset; + uint64_t size; + uint64_t tripleSize; + const char triple[1]; + }; + + struct __ClangOffloadBundleHeader { + const char magic[sizeof(CLANG_OFFLOAD_BUNDLER_MAGIC_STR) - 1]; + uint64_t numBundles; + __ClangOffloadBundleDesc desc[1]; + }; + +protected: + CodeObject() {} + //Given an ptr to image or file, extracts to code object + //for corresponding devices + hipError_t extractCodeObjectFromFatBinary(const void*, + const std::vector&, + std::vector>&); + + uint64_t ElfSize(const void* emi); + +private: + bool isCompatibleCodeObject(const std::string& codeobj_target_id, + const char* device_name); + + friend const std::vector& modules(); +}; + +//Dynamic Code Object +class DynCO : public CodeObject { + amd::Monitor dclock_{"Guards Static Code object", true}; + +public: + DynCO(); + virtual ~DynCO(); + + //LoadsCodeObject and its data + hipError_t loadCodeObject(const char* fname, const void* image=nullptr); + hipModule_t module() { return reinterpret_cast(as_cl(program_)); }; + + //Gets GlobalVar/Functions from a dynamically loaded code object + hipError_t getDynFunc(hipFunction_t* hfunc, std::string func_name); + hipError_t getDeviceVar(DeviceVar** dvar, std::string var_name, int deviceId); + +private: + amd::Program* program_; + + //Maps for vars/funcs, could be keyed in with std::string name + std::unordered_map functions_; + std::unordered_map vars_; + + //Load Code Object Data(Vars/UndefinedVars/Funcs) + hipError_t loadCodeObjectData(const void* mmap_ptr, size_t mmap_size); + + //Populate Global Vars/Funcs from an code object(@ module_load) + hipError_t populateDynGlobalFuncs(); + hipError_t populateDynGlobalVars(); +}; + +//Static Code Object +class StatCO: public CodeObject { + amd::Monitor sclock_{"Guards Static Code object", true}; +public: + StatCO(); + virtual ~StatCO(); + + //Add/Remove/Digest Fat Binaries passed to us from "__hipRegisterFatBinary" + FatBinaryInfoType* addFatBinary(const void* data, bool initialized); + hipError_t removeFatBinary(FatBinaryInfoType* module); + hipError_t digestFatBinary(const void* data, FatBinaryInfoType& programs); + + //Register vars/funcs given to use from __hipRegister[Var/Func] + hipError_t registerStatFunction(const void* hostFunction, Function* func); + hipError_t registerStatGlobalVar(const void* hostVar, Var* var); + + //Retrive Vars/Funcs for a given hostSidePtr(const void*), unless stated otherwise. + hipError_t getStatFunc(hipFunction_t* hfunc, const void* hostFunction, int deviceId); + hipError_t getStatFuncAttr(hipFuncAttributes* func_attr, const void* hostFunction, int deviceId); + hipError_t getStatGlobalVar(const void* hostVar, int deviceId, hipDeviceptr_t* dev_ptr, + size_t* size_ptr); + hipError_t getStatGlobalVarByName(std::string hostVar, int deviceId, hipModule_t hmod, + hipDeviceptr_t* dev_ptr, size_t* size_ptr); + +private: + friend class ::PlatformState; + //Populated during __hipRegisterFatBinary + std::unordered_map modules_; + //Populated during __hipRegisterFuncs + std::unordered_map functions_; + //Populated during __hipRegisterVars + std::unordered_map vars_; +}; + +}; //namespace: hip + +#endif /* HIP_CODE_OBJECT_HPP */ diff --git a/rocclr/hip_context.cpp b/rocclr/hip_context.cpp index 2d0abc6add..7c42c65159 100755 --- a/rocclr/hip_context.cpp +++ b/rocclr/hip_context.cpp @@ -20,6 +20,7 @@ #include #include "hip_internal.hpp" +#include "hip_platform.hpp" #include "platform/runtime.hpp" #include "utils/flags.hpp" #include "utils/versions.hpp" diff --git a/rocclr/hip_fatbin.hpp b/rocclr/hip_fatbin.hpp new file mode 100755 index 0000000000..8f936283fb --- /dev/null +++ b/rocclr/hip_fatbin.hpp @@ -0,0 +1,29 @@ +#ifndef HIP_FAT_BINARY_HPP +#define HIP_FAT_BINARY_HPP + +namespace hip { + +class FatBinaryMetaInfo { +public: + FatBinaryMetaInfo(bool built, const void* binary_ptr, size_t binary_size): + built_(built), binary_ptr_(binary_ptr), binary_size_(binary_size) {} + ~FatBinaryMetaInfo() {} + + //Set once the mod has been built + void set_built() { built_ = true; } + + //Accessor for private vars + bool built() const { return built_; } + const void* binary_ptr() const { return binary_ptr_; } + size_t binary_size() const { return binary_size_; } +private: + bool built_; //Set when mod is built. Used in Lazy Binary + const void* binary_ptr_; //Binary image ptr + size_t binary_size_; //Binary Size +}; + +typedef std::vector> FatBinaryInfoType; + +}; /* namespace hip */ + +#endif /* HIP_FAT_BINARY_HPP */ diff --git a/rocclr/hip_global.cpp b/rocclr/hip_global.cpp new file mode 100755 index 0000000000..1bbc157d58 --- /dev/null +++ b/rocclr/hip_global.cpp @@ -0,0 +1,202 @@ +#include "hip_global.hpp" + +#include "hip/hip_runtime.h" +#include "hip_internal.hpp" +#include "hip_code_object.hpp" +#include "platform/program.hpp" + +namespace hip { + +//Device Vars +DeviceVar::DeviceVar(std::string name, hipModule_t hmod) : shadowVptr(nullptr), name_(name), + amd_mem_obj_(nullptr), device_ptr_(nullptr), + size_(0) { + amd::Program* program = as_amd(reinterpret_cast(hmod)); + device::Program* dev_program = program->getDeviceProgram(*hip::getCurrentDevice()->devices()[0]); + if (dev_program == nullptr) { + DevLogPrintfError("Cannot get Device Function for module: 0x%x \n", hmod); + guarantee(false); + } + + if(!dev_program->createGlobalVarObj(&amd_mem_obj_, &device_ptr_, &size_, name.c_str())) { + DevLogPrintfError("Cannot create Global Var obj for symbol: %s \n", name); + guarantee(false); + } + + if (amd_mem_obj_ == nullptr || device_ptr_ == nullptr) { + DevLogPrintfError("Cannot get memory for creating device Var: %s", name.c_str()); + guarantee(false); + } + + amd::MemObjMap::AddMemObj(device_ptr_, amd_mem_obj_); +} + +DeviceVar::~DeviceVar() { + if (device_ptr_ != nullptr) { + amd::MemObjMap::RemoveMemObj(device_ptr_); + amd_mem_obj_->release(); + } + + if (shadowVptr != nullptr) { + textureReference* texRef = reinterpret_cast(shadowVptr); + delete texRef; + shadowVptr = nullptr; + } + + device_ptr_ = nullptr; + size_ = 0; +} + +//Device Functions +DeviceFunc::DeviceFunc(std::string name, hipModule_t hmod) : dflock_("function lock"), + name_(name), kernel_(nullptr) { + amd::Program* program = as_amd(reinterpret_cast(hmod)); + + const amd::Symbol *symbol = program->findSymbol(name.c_str()); + if (symbol == nullptr) { + DevLogPrintfError("Cannot find Symbol with name: %s \n", name); + guarantee(false); + } + + kernel_ = new amd::Kernel(*program, *symbol, name); + if (kernel_ == nullptr) { + DevLogPrintfError("Cannot create kernel with name: %s \n", name); + guarantee(false); + } +} + +DeviceFunc::~DeviceFunc() { + if (kernel_ != nullptr) { + kernel_->release(); + } +} + +//Abstract functions +Function::Function(std::string name, FatBinaryInfoType* modules) + : name_(name), modules_(modules) { + dFunc_.resize(g_devices.size()); +} + +Function::~Function() { + for (auto& elem : dFunc_) { + delete elem; + } + name_ = ""; + modules_ = nullptr; +} + +hipError_t Function::getDynFunc(hipFunction_t* hfunc, hipModule_t hmod) { + guarantee(dFunc_.size() == g_devices.size()); + if (dFunc_[ihipGetDevice()] == nullptr) { + dFunc_[ihipGetDevice()] = new DeviceFunc(name_, hmod); + } + *hfunc = dFunc_[ihipGetDevice()]->asHipFunction(); + + return hipSuccess; +} + +hipError_t Function::getStatFunc(hipFunction_t* hfunc, int deviceId) { + guarantee(modules_ != nullptr); + guarantee(deviceId >= 0); + guarantee(deviceId < modules_->size()); + + hipModule_t module = (*modules_)[deviceId].first; + FatBinaryMetaInfo* fb_meta = (*modules_)[deviceId].second; + + if (!fb_meta->built()) { + IHIP_RETURN_ONFAIL(CodeObject::add_program(deviceId, module, fb_meta->binary_ptr(), + fb_meta->binary_size())); + IHIP_RETURN_ONFAIL(CodeObject::build_module(module, g_devices[deviceId]->devices())); + fb_meta->set_built(); + } + + if (dFunc_[deviceId] == nullptr) { + dFunc_[deviceId] = new DeviceFunc(name_, (*modules_)[deviceId].first); + } + *hfunc = dFunc_[deviceId]->asHipFunction(); + + return hipSuccess; +} + +hipError_t Function::getStatFuncAttr(hipFuncAttributes* func_attr, int deviceId) { + guarantee(modules_ != nullptr); + guarantee(deviceId >= 0); + guarantee(deviceId < modules_->size()); + + hipModule_t module = (*modules_)[deviceId].first; + FatBinaryMetaInfo* fb_meta = (*modules_)[deviceId].second; + + if (!fb_meta->built()) { + IHIP_RETURN_ONFAIL(CodeObject::add_program(deviceId, module, fb_meta->binary_ptr(), + fb_meta->binary_size())); + IHIP_RETURN_ONFAIL(CodeObject::build_module(module, g_devices[deviceId]->devices())); + fb_meta->set_built(); + } + + if (dFunc_[deviceId] == nullptr) { + dFunc_[deviceId] = new DeviceFunc(name_, (*modules_)[deviceId].first); + } + + const std::vector& devices = amd::Device::getDevices(CL_DEVICE_TYPE_GPU, false); + + amd::Kernel* kernel = dFunc_[deviceId]->kernel(); + const device::Kernel::WorkGroupInfo* wginfo = kernel->getDeviceKernel(*devices[deviceId])->workGroupInfo(); + func_attr->localSizeBytes = wginfo->privateMemSize_; + func_attr->sharedSizeBytes = wginfo->localMemSize_; + func_attr->maxDynamicSharedSizeBytes = wginfo->availableLDSSize_ - wginfo->localMemSize_; + func_attr->maxThreadsPerBlock = wginfo->size_; + func_attr->numRegs = wginfo->usedVGPRs_; + + return hipSuccess; +} + +//Abstract Vars +Var::Var(std::string name, DeviceVarKind dVarKind, size_t size, int type, int norm, + FatBinaryInfoType* modules) : name_(name), dVarKind_(dVarKind), size_(size), + type_(type), norm_(norm), modules_(modules) { + dVar_.resize(g_devices.size()); +} + +Var::~Var() { + for (auto& elem : dVar_) { + delete elem; + } + modules_ = nullptr; +} + +hipError_t Var::getDeviceVar(DeviceVar** dvar, int deviceId, hipModule_t hmod) { + guarantee(deviceId >= 0); + guarantee(deviceId < g_devices.size()); + guarantee(dVar_.size() == g_devices.size()); + + if (dVar_[deviceId] == nullptr) { + dVar_[deviceId] = new DeviceVar(name_, hmod); + } + + *dvar = dVar_[deviceId]; + return hipSuccess; +} + +hipError_t Var::getStatDeviceVar(DeviceVar** dvar, int deviceId) { + guarantee(deviceId >= 0); + guarantee(deviceId < g_devices.size()); + + hipModule_t module = (*modules_)[deviceId].first; + FatBinaryMetaInfo* fb_meta = (*modules_)[deviceId].second; + + if (!fb_meta->built()) { + IHIP_RETURN_ONFAIL(CodeObject::add_program(deviceId, module, fb_meta->binary_ptr(), + fb_meta->binary_size())); + IHIP_RETURN_ONFAIL(CodeObject::build_module(module, g_devices[deviceId]->devices())); + fb_meta->set_built(); + } + + if (dVar_[deviceId] == nullptr) { + dVar_[deviceId] = new DeviceVar(name_, (*modules_)[deviceId].first); + } + + *dvar = dVar_[deviceId]; + return hipSuccess; +} + +}; //namespace: hip diff --git a/rocclr/hip_global.hpp b/rocclr/hip_global.hpp new file mode 100755 index 0000000000..52274be51a --- /dev/null +++ b/rocclr/hip_global.hpp @@ -0,0 +1,116 @@ +#ifndef HIP_GLOBAL_HPP +#define HIP_GLOBAL_HPP + +#include +#include + +#include "hip/hip_runtime_api.h" +#include "hip/hip_runtime.h" +#include "hip_internal.hpp" +#include "hip_fatbin.hpp" +#include "platform/program.hpp" + +namespace hip { + +//Forward Declaration +class CodeObject; + +//Device Structures +class DeviceVar { +public: + DeviceVar(std::string name, hipModule_t hmod); + ~DeviceVar(); + + //Accessors for device ptr and size, populated during constructor. + hipDeviceptr_t device_ptr() const { return device_ptr_; } + size_t size() const { return size_; } + std::string name() const { return name_; } + void* shadowVptr; + +private: + std::string name_; //Name of the var + amd::Memory* amd_mem_obj_; //amd_mem_obj abstraction + hipDeviceptr_t device_ptr_; //Device Pointer + size_t size_; //Size of the var +}; + +class DeviceFunc { +public: + DeviceFunc(std::string name, hipModule_t hmod); + ~DeviceFunc(); + + amd::Monitor dflock_; + + //Converts DeviceFunc to hipFunction_t(used by app) and vice versa. + hipFunction_t asHipFunction() { return reinterpret_cast(this); } + static DeviceFunc* asFunction(hipFunction_t f) { return reinterpret_cast(f); } + + //Accessor for kernel_ and name_ populated during constructor. + std::string name() const { return name_; } + amd::Kernel* kernel() const { return kernel_; } + +private: + std::string name_; //name of the func(not unique identifier) + amd::Kernel* kernel_; //Kernel ptr referencing to ROCclr Symbol +}; + +//Abstract Structures +class Function { +public: + Function(std::string name, FatBinaryInfoType* modules=nullptr); + ~Function(); + + //Return DeviceFunc for this this dynamically loaded module + hipError_t getDynFunc(hipFunction_t* hfunc, hipModule_t hmod); + + //Return Device Func & attr . Generate/build if not already done so. + hipError_t getStatFunc(hipFunction_t *hfunc, int deviceId); + hipError_t getStatFuncAttr(hipFuncAttributes* func_attr, int deviceId); + void resize_dFunc(size_t size) { dFunc_.resize(size); } + FatBinaryInfoType* moduleInfo() { return modules_; }; + +private: + std::vector dFunc_; //DeviceFuncObj per Device + std::string name_; //name of the func(not unique identifier) + FatBinaryInfoType* modules_; // static module where it is referenced +}; + +class Var { +public: + //Types of variable + enum DeviceVarKind { + DVK_Variable = 0, + DVK_Surface, + DVK_Texture + }; + + Var(std::string name, DeviceVarKind dVarKind, size_t size, int type, int norm, + FatBinaryInfoType* modules = nullptr); + ~Var(); + + //Return DeviceVar for this dynamically loaded module + hipError_t getDeviceVar(DeviceVar** dvar, int deviceId, hipModule_t hmod); + + //Return DeviceVar for module Generate/build if not already done so. + hipError_t getStatDeviceVar(DeviceVar** dvar, int deviceId); + void resize_dVar(size_t size) { dVar_.resize(size); } + + //Accessor for device_ptrs. + std::string name() const { return name_; } + hipModule_t module(int deviceId) const { return (*modules_)[deviceId].first; } + hipDeviceptr_t device_ptr(int deviceId) const { return dVar_[deviceId]->device_ptr(); } + size_t device_size(int deviceId) const { return dVar_[deviceId]->size(); } + FatBinaryInfoType* moduleInfo() { return modules_; }; + +private: + std::vector dVar_; // DeviceVarObj per Device + std::string name_; // Variable name (not unique identifier) + DeviceVarKind dVarKind_; // Variable kind + size_t size_; // Size of the variable + int type_; // Type(Textures/Surfaces only) + int norm_; // Type(Textures/Surfaces only) + FatBinaryInfoType* modules_; // static module where it is referenced +}; + +}; //namespace: hip +#endif /* HIP_GLOBAL_HPP */ diff --git a/rocclr/hip_intercept.cpp b/rocclr/hip_intercept.cpp old mode 100644 new mode 100755 index c1b7f53534..10eeb43927 --- a/rocclr/hip_intercept.cpp +++ b/rocclr/hip_intercept.cpp @@ -20,6 +20,7 @@ #include "hip/hip_runtime.h" #include "hip_internal.hpp" +#include "hip_platform.hpp" #include "hip_prof_api.h" // HIP API callback/activity @@ -44,8 +45,9 @@ const char* hipKernelNameRefByPtr(const void* hostFunction, hipStream_t stream) DevLogPrintfError("Wrong Device Id: %d \n", deviceId); return NULL; } - hipFunction_t func = PlatformState::instance().getFunc(hostFunction, deviceId); - if (func == nullptr) { + hipFunction_t func = nullptr; + hipError_t hip_error = PlatformState::instance().getStatFunc(&func, hostFunction, deviceId); + if (hip_error != hipSuccess) { return NULL; } return hipKernelNameRef(func); diff --git a/rocclr/hip_internal.hpp b/rocclr/hip_internal.hpp index c42bda9fec..a53db4be74 100755 --- a/rocclr/hip_internal.hpp +++ b/rocclr/hip_internal.hpp @@ -80,6 +80,25 @@ typedef struct ihipIpcMemHandle_st { HIP_ERROR_PRINT(hip::g_lastError, __VA_ARGS__) \ return hip::g_lastError; +#define HIP_RETURN_ONFAIL(func) \ + do { \ + hipError_t herror = (func); \ + if (herror != hipSuccess) { \ + HIP_RETURN(herror); \ + } \ + } while (0); + +// Cannot be use in place of HIP_RETURN. +// Refrain from using for external HIP APIs +#define IHIP_RETURN_ONFAIL(func) \ + do { \ + hipError_t herror = (func); \ + if (herror != hipSuccess) { \ + return herror; \ + } \ + } while (0); + + namespace hc { class accelerator; class accelerator_view; @@ -198,17 +217,6 @@ namespace hip { extern amd::HostQueue* getNullStream(amd::Context&); /// Get default stream of the thread extern amd::HostQueue* getNullStream(); - - struct Function { - amd::Kernel* function_; - amd::Monitor lock_; - - Function(amd::Kernel* f) : function_(f), lock_("function lock") {} - ~Function() { function_->release(); } - hipFunction_t asHipFunction() { return reinterpret_cast(this); } - - static Function* asFunction(hipFunction_t f) { return reinterpret_cast(f); } - }; }; struct ihipExec_t { @@ -219,138 +227,6 @@ struct ihipExec_t { std::vector arguments_; }; -class PlatformState { - amd::Monitor lock_{"Guards global function map", true}; - - std::unordered_map>> modules_; - bool initialized_{false}; - - void digestFatBinary(const void* data, std::vector>& programs); -public: - void init(); - std::vector>* addFatBinary(const void*data) - { - amd::ScopedLock lock(lock_); - if (initialized_) { - digestFatBinary(data, modules_[data]); - } - return &modules_[data]; - } - void removeFatBinary(std::vector>* module) - { - amd::ScopedLock lock(lock_); - for (auto& mod : modules_) { - if (&mod.second == module) { - modules_.erase(&mod); - return; - } - } - } - - struct RegisteredVar { - public: - RegisteredVar(): size_(0), devicePtr_(nullptr), amd_mem_obj_(nullptr) {} - ~RegisteredVar() {} - - hipDeviceptr_t getdeviceptr() const { return devicePtr_; }; - amd::Memory* amd_mem_obj() const { return amd_mem_obj_; }; - size_t getvarsize() const { return size_; }; - - size_t size_; // Size of the variable - hipDeviceptr_t devicePtr_; //Device Memory Address of the variable. - amd::Memory* amd_mem_obj_; - }; - - struct DeviceFunction { - std::string deviceName; - std::vector< std::pair< hipModule_t, bool > >* modules; - std::vector functions; - }; - enum DeviceVarKind { - DVK_Variable, - DVK_Surface, - DVK_Texture - }; - struct DeviceVar { - DeviceVarKind kind; - void* shadowVptr; - std::string hostVar; - size_t size; - std::vector< std::pair< hipModule_t, bool > >* modules; - std::vector rvars; - bool dyn_undef; - int type; // surface/texture type - int norm; // texture has normalized output - bool shadowAllocated = false; // shadow ptr is allocated on-demand and needs freeing. - }; -private: - class Module { - public: - Module(hipModule_t hip_module_) : hip_module(hip_module_) {} - std::unordered_map functions_; - private: - hipModule_t hip_module; - }; - std::unordered_map module_map_; - - std::unordered_map functions_; - std::unordered_multimap vars_; - // Map from the host shadow symbol to its device name. As different modules - // may have the same name, each symbol is uniquely identified by a pair of - // module handle and its name. - std::unordered_map> symbols_; - - typedef std::pair CodeObjPairType; - std::unordered_map code_obj_; - - static PlatformState* platform_; - - PlatformState() {} - ~PlatformState() {} -public: - static PlatformState& instance() { - if (platform_ == nullptr) { - // __hipRegisterFatBinary() will call this when app starts, thus - // there is no multiple entry issue here. - platform_ = new PlatformState(); - } - return *platform_; - } - - bool unregisterFunc(hipModule_t hmod); - std::vector< std::pair >* unregisterVar(hipModule_t hmod); - - - bool findSymbol(const void *hostVar, hipModule_t &hmod, std::string &devName); - PlatformState::DeviceVar* findVar(std::string hostVar, int deviceId, hipModule_t hmod); - void registerVarSym(const void *hostVar, hipModule_t hmod, const char *symbolName); - void registerVar(const char* symbolName, const DeviceVar& var); - void registerFunction(const void* hostFunction, const DeviceFunction& func); - - bool registerModFuncs(std::vector& func_names, hipModule_t* module); - bool findModFunc(hipFunction_t* hfunc, hipModule_t hmod, const char* name); - bool createFunc(hipFunction_t* hfunc, hipModule_t hmod, const char* name); - hipFunction_t getFunc(const void* hostFunction, int deviceId); - bool getFuncAttr(const void* hostFunction, hipFuncAttributes* func_attr); - bool getGlobalVar(const char* hostVar, int deviceId, hipModule_t hmod, - hipDeviceptr_t* dev_ptr, size_t* size_ptr); - bool getTexRef(const char* hostVar, hipModule_t hmod, textureReference** texRef); - - bool getGlobalVarFromSymbol(const void* hostVar, int deviceId, - hipDeviceptr_t* dev_ptr, size_t* size_ptr); - - bool getShadowVarInfo(std::string var_name, hipModule_t hmod, - void** var_addr, size_t* var_size); - void setupArgument(const void *arg, size_t size, size_t offset); - void configureCall(dim3 gridDim, dim3 blockDim, size_t sharedMem, hipStream_t stream); - - void popExec(ihipExec_t& exec); -}; - -constexpr bool kOptionChangeable = true; -constexpr bool kNewDevProg = false; - /// Wait all active streams on the blocking queue. The method enqueues a wait command and /// doesn't stall the current thread extern void iHipWaitActiveStreams(amd::HostQueue* blocking_queue, bool wait_null_stream = false); @@ -363,5 +239,6 @@ extern amd::Memory* getMemoryObject(const void* ptr, size_t& offset); extern bool CL_CALLBACK getSvarInfo(cl_program program, std::string var_name, void** var_addr, size_t* var_size); - +constexpr bool kOptionChangeable = true; +constexpr bool kNewDevProg = false; #endif // HIP_SRC_HIP_INTERNAL_H diff --git a/rocclr/hip_memory.cpp b/rocclr/hip_memory.cpp index 84c8b77f92..2854b0d7eb 100755 --- a/rocclr/hip_memory.cpp +++ b/rocclr/hip_memory.cpp @@ -20,6 +20,7 @@ #include #include "hip_internal.hpp" +#include "hip_platform.hpp" #include "hip_conversions.hpp" #include "platform/context.hpp" #include "platform/command.hpp" @@ -744,18 +745,7 @@ hipError_t hipMemcpyToSymbol(const void* symbol, const void* src, size_t sizeByt size_t sym_size = 0; hipDeviceptr_t device_ptr = nullptr; - hipModule_t hmod; - std::string symbolName; - if (!PlatformState::instance().findSymbol(symbol, hmod, symbolName)) { - DevLogPrintfError("cannot find symbol 0x%x \n", symbolName.c_str()); - HIP_RETURN(hipErrorInvalidSymbol); - } - /* Get address and size for the global symbol */ - if (!PlatformState::instance().getGlobalVar(symbolName.c_str(), ihipGetDevice(), hmod, - &device_ptr, &sym_size)) { - DevLogPrintfError("Cannot get global var: %s at device: %d \n", symbolName.c_str(), ihipGetDevice()); - HIP_RETURN(hipErrorInvalidSymbol); - } + HIP_RETURN_ONFAIL(PlatformState::instance().getStatGlobalVar(symbol, ihipGetDevice(), &device_ptr, &sym_size)); /* Size Check to make sure offset is correct */ if ((offset + sizeBytes) > sym_size) { @@ -777,18 +767,7 @@ hipError_t hipMemcpyFromSymbol(void* dst, const void* symbol, size_t sizeBytes, size_t sym_size = 0; hipDeviceptr_t device_ptr = nullptr; - hipModule_t hmod; - std::string symbolName; - if (!PlatformState::instance().findSymbol(symbol, hmod, symbolName)) { - DevLogPrintfError("cannot find symbol: 0x%x \n", symbol); - HIP_RETURN(hipErrorInvalidSymbol); - } - /* Get address and size for the global symbol */ - if (!PlatformState::instance().getGlobalVar(symbolName.c_str(), ihipGetDevice(), hmod, - &device_ptr, &sym_size)) { - DevLogPrintfError("Cannot find symbol Name: %s \n", symbolName.c_str()); - HIP_RETURN(hipErrorInvalidSymbol); - } + HIP_RETURN_ONFAIL(PlatformState::instance().getStatGlobalVar(symbol, ihipGetDevice(), &device_ptr, &sym_size)); /* Size Check to make sure offset is correct */ if ((offset + sizeBytes) > sym_size) { @@ -810,18 +789,7 @@ hipError_t hipMemcpyToSymbolAsync(const void* symbol, const void* src, size_t si size_t sym_size = 0; hipDeviceptr_t device_ptr = nullptr; - hipModule_t hmod; - std::string symbolName; - if (!PlatformState::instance().findSymbol(symbol, hmod, symbolName)) { - DevLogPrintfError("cannot find symbol: 0x%x \n", symbol); - HIP_RETURN(hipErrorInvalidSymbol); - } - /* Get address and size for the global symbol */ - if (!PlatformState::instance().getGlobalVar(symbolName.c_str(), ihipGetDevice(), hmod, - &device_ptr, &sym_size)) { - DevLogPrintfError("Cannot find symbol Name: %s \n", symbolName.c_str()); - HIP_RETURN(hipErrorInvalidSymbol); - } + HIP_RETURN_ONFAIL(PlatformState::instance().getStatGlobalVar(symbol, ihipGetDevice(), &device_ptr, &sym_size)); /* Size Check to make sure offset is correct */ if ((offset + sizeBytes) > sym_size) { @@ -843,18 +811,7 @@ hipError_t hipMemcpyFromSymbolAsync(void* dst, const void* symbol, size_t sizeBy size_t sym_size = 0; hipDeviceptr_t device_ptr = nullptr; - hipModule_t hmod; - std::string symbolName; - if (!PlatformState::instance().findSymbol(symbol, hmod, symbolName)) { - DevLogPrintfError("cannot find symbol: 0x%x \n", symbol); - HIP_RETURN(hipErrorInvalidSymbol); - } - /* Get address and size for the global symbol */ - if (!PlatformState::instance().getGlobalVar(symbolName.c_str(), ihipGetDevice(), hmod, - &device_ptr, &sym_size)) { - DevLogPrintfError("Cannot find symbol Name: %s \n", symbolName.c_str()); - HIP_RETURN(hipErrorInvalidSymbol); - } + HIP_RETURN_ONFAIL(PlatformState::instance().getStatGlobalVar(symbol, ihipGetDevice(), &device_ptr, &sym_size)); /* Size Check to make sure offset is correct */ if ((offset + sizeBytes) > sym_size) { diff --git a/rocclr/hip_module.cpp b/rocclr/hip_module.cpp index f618254599..383f7a77a9 100755 --- a/rocclr/hip_module.cpp +++ b/rocclr/hip_module.cpp @@ -39,9 +39,8 @@ extern hipError_t ihipLaunchKernel(const void* hostFunction, hipEvent_t stopEvent, int flags); -const std::string& FunctionName(const hipFunction_t f) -{ - return hip::Function::asFunction(f)->function_->name(); +const std::string& FunctionName(const hipFunction_t f) { + return hip::DeviceFunc::asFunction(f)->kernel()->name(); } static uint64_t ElfSize(const void *emi) @@ -65,223 +64,48 @@ static uint64_t ElfSize(const void *emi) return total_size; } -hipError_t hipModuleLoad(hipModule_t* module, const char* fname) -{ - HIP_INIT_API(hipModuleLoad, module, fname); - - const void* mmap_ptr = nullptr; - size_t mmap_size = 0; - - if (!fname) { - HIP_RETURN(hipErrorInvalidValue); - } - - if (!amd::Os::MemoryMapFile(fname, &mmap_ptr, &mmap_size)) { - HIP_RETURN(hipErrorFileNotFound); - } - - HIP_RETURN(ihipModuleLoadData(module, mmap_ptr, mmap_size)); -} - -bool ihipModuleUnregisterGlobal(hipModule_t hmod) { - std::vector< std::pair >* modules = - PlatformState::instance().unregisterVar(hmod); - if (modules != nullptr) { - delete modules; - } - return true; -} - -hipError_t hipModuleUnload(hipModule_t hmod) -{ +hipError_t hipModuleUnload(hipModule_t hmod) { HIP_INIT_API(hipModuleUnload, hmod); - if (hmod == nullptr) { - HIP_RETURN(hipErrorInvalidValue); - } + HIP_RETURN(PlatformState::instance().unloadModule(hmod)); +} - amd::Program* program = as_amd(reinterpret_cast(hmod)); +hipError_t hipModuleLoad(hipModule_t* module, const char* fname) { + HIP_INIT_API(hipModuleLoad, module, fname); - if(!PlatformState::instance().unregisterFunc(hmod)) { - DevLogPrintfError("Cannot unregister module: 0x%x \n", hmod); - HIP_RETURN(hipErrorInvalidSymbol); - } - - if(!ihipModuleUnregisterGlobal(hmod)) { - DevLogPrintfError("Cannot unregister Global vars for module: 0x%x \n", hmod); - HIP_RETURN(hipErrorInvalidSymbol); - } - - program->release(); - - HIP_RETURN(hipSuccess); + HIP_RETURN(PlatformState::instance().loadModule(module, fname)); } hipError_t hipModuleLoadData(hipModule_t *module, const void *image) { HIP_INIT_API(hipModuleLoadData, module, image); - HIP_RETURN(ihipModuleLoadData(module, image, 0)); + HIP_RETURN(PlatformState::instance().loadModule(module, 0, image)); } hipError_t hipModuleLoadDataEx(hipModule_t *module, const void *image, - unsigned int numOptions, hipJitOption* options, - void** optionsValues) + unsigned int numOptions, hipJitOption* options, + void** optionsValues) { /* TODO: Pass options to Program */ HIP_INIT_API(hipModuleLoadDataEx, module, image); - HIP_RETURN(ihipModuleLoadData(module, image, 0)); + HIP_RETURN(PlatformState::instance().loadModule(module, 0, image)); } extern hipError_t __hipExtractCodeObjectFromFatBinary(const void* data, const std::vector& devices, std::vector>& code_objs); -inline bool ihipModuleRegisterUndefined(amd::Program* program, hipModule_t* module) { - - std::vector undef_vars; - device::Program* dev_program - = program->getDeviceProgram(*hip::getCurrentDevice()->devices()[0]); - - if (!dev_program->getUndefinedVarFromCodeObj(&undef_vars)) { - DevLogPrintfError("Could not get undefined Variables for Module: 0x%x \n", *module); - return false; - } - - for (auto it = undef_vars.begin(); it != undef_vars.end(); ++it) { - auto modules = new std::vector >(g_devices.size()); - for (size_t dev = 0; dev < g_devices.size(); ++dev) { - modules->at(dev) = std::make_pair(*module, true); - } - - texture* tex_hptr - = new texture(); - memset(tex_hptr, 0x00, sizeof(texture)); - - PlatformState::DeviceVar dvar{PlatformState::DVK_Variable, - reinterpret_cast(tex_hptr), - it->c_str(), - sizeof(*tex_hptr), - modules, - std::vector{g_devices.size()}, - true, - /*type*/ 0, - /*norm*/ 0}; - PlatformState::instance().registerVar(it->c_str(), dvar); - } - - return true; -} - -inline bool ihipModuleRegisterFunc(amd::Program* program, hipModule_t* module) { - - std::vector func_names; - device::Program* dev_program - = program->getDeviceProgram(*hip::getCurrentDevice()->devices()[0]); - - // Get all the global func names from COMGR - if (!dev_program->getGlobalFuncFromCodeObj(&func_names)) { - DevLogPrintfError("Could not get Global Funcs from Code Obj for Module: 0x%x \n", *module); - return false; - } - - return PlatformState::instance().registerModFuncs(func_names, module); -} - - -inline bool ihipModuleRegisterGlobal(amd::Program* program, hipModule_t* module) { - - size_t var_size = 0; - hipDeviceptr_t device_ptr = nullptr; - std::vector var_names; - - device::Program* dev_program - = program->getDeviceProgram(*hip::getCurrentDevice()->devices()[0]); - - if (!dev_program->getGlobalVarFromCodeObj(&var_names)) { - DevLogPrintfError("Could not get Global vars from Code Obj for Module: 0x%x \n", *module); - return false; - } - - for (auto it = var_names.begin(); it != var_names.end(); ++it) { - auto modules = new std::vector >(g_devices.size()); - for (size_t dev = 0; dev < g_devices.size(); ++dev) { - modules->at(dev) = std::make_pair(*module, true); - } - - PlatformState::DeviceVar dvar{PlatformState::DVK_Variable, - nullptr, - it->c_str(), - 0, - modules, - std::vector{g_devices.size()}, - false, - /*type*/ 0, - /*norm*/ 0}; - PlatformState::instance().registerVar(it->c_str(), dvar); - } - - return true; -} - -hipError_t ihipModuleLoadData(hipModule_t* module, const void* mmap_ptr, size_t mmap_size) -{ - /* initialize image it to the mmap_ptr, if this is of no_clang_offload bundle then they directly pass the image */ - const void* image = mmap_ptr; - std::vector> code_objs; - hipError_t code_obj_err = __hipExtractCodeObjectFromFatBinary(mmap_ptr, - {hip::getCurrentDevice()->devices()[0]->info().name_}, code_objs); - if (code_obj_err == hipSuccess) { - image = code_objs[0].first; - } else if(code_obj_err == hipErrorNoBinaryForGpu) { - return code_obj_err; - } - - amd::Program* program = new amd::Program(*hip::getCurrentDevice()->asContext(), - amd::Program::Language::Binary, mmap_ptr, mmap_size); - if (program == NULL) { - return hipErrorOutOfMemory; - } - - program->setVarInfoCallBack(&getSvarInfo); - - if (CL_SUCCESS != program->addDeviceProgram(*hip::getCurrentDevice()->devices()[0], image, - ElfSize(image), false)) { - return hipErrorInvalidKernelFile; - } - - *module = reinterpret_cast(as_cl(program)); - - if (!ihipModuleRegisterGlobal(program, module)) { - return hipErrorSharedObjectSymbolNotFound; - } - - if (!ihipModuleRegisterUndefined(program, module)) { - return hipErrorSharedObjectSymbolNotFound; - } - - if (CL_SUCCESS != program->build(hip::getCurrentDevice()->devices(), nullptr, nullptr, nullptr, - kOptionChangeable, kNewDevProg)) { - return hipErrorSharedObjectInitFailed; - } - - if (!ihipModuleRegisterFunc(program, module)) { - return hipErrorSharedObjectSymbolNotFound; - } - - return hipSuccess; -} - -hipError_t hipModuleGetFunction(hipFunction_t *hfunc, hipModule_t hmod, const char *name) -{ +hipError_t hipModuleGetFunction(hipFunction_t *hfunc, hipModule_t hmod, const char *name) { HIP_INIT_API(hipModuleGetFunction, hfunc, hmod, name); - if (!PlatformState::instance().findModFunc(hfunc, hmod, name)) { + if (hipSuccess != PlatformState::instance().getDynFunc(hfunc, hmod, name)) { DevLogPrintfError("Cannot find the function: %s for module: 0x%x \n", name, hmod); HIP_RETURN(hipErrorNotFound); } + HIP_RETURN(hipSuccess); } @@ -290,8 +114,7 @@ hipError_t hipModuleGetGlobal(hipDeviceptr_t* dptr, size_t* bytes, hipModule_t h HIP_INIT_API(hipModuleGetGlobal, dptr, bytes, hmod, name); /* Get address and size for the global symbol */ - if (!PlatformState::instance().getGlobalVar(name, ihipGetDevice(), hmod, - dptr, bytes)) { + if (hipSuccess != PlatformState::instance().getDynGlobalVar(name, ihipGetDevice(), hmod, dptr, bytes)) { DevLogPrintfError("Cannot find global Var: %s for module: 0x%x at device: %d \n", name, hmod, ihipGetDevice()); HIP_RETURN(hipErrorNotFound); @@ -307,12 +130,12 @@ hipError_t hipFuncGetAttribute(int* value, hipFunction_attribute attrib, hipFunc HIP_RETURN(hipErrorInvalidValue); } - hip::Function* function = hip::Function::asFunction(hfunc); + hip::DeviceFunc* function = hip::DeviceFunc::asFunction(hfunc); if (function == nullptr) { HIP_RETURN(hipErrorInvalidHandle); } - amd::Kernel* kernel = function->function_; + amd::Kernel* kernel = function->kernel(); if (kernel == nullptr) { HIP_RETURN(hipErrorInvalidDeviceFunction); } @@ -365,9 +188,7 @@ hipError_t hipFuncGetAttributes(hipFuncAttributes* attr, const void* func) { HIP_INIT_API(hipFuncGetAttributes, attr, func); - if (!PlatformState::instance().getFuncAttr(func, attr)) { - HIP_RETURN(hipErrorInvalidDeviceFunction); - } + HIP_RETURN_ONFAIL(PlatformState::instance().getStatFuncAttr(attr, func, ihipGetDevice())); HIP_RETURN(hipSuccess); } @@ -383,10 +204,10 @@ hipError_t ihipModuleLaunchKernel(hipFunction_t f, HIP_INIT_API(ihipModuleLaunchKernel, f, gridDimX, gridDimY, gridDimZ, blockDimX, blockDimY, blockDimZ, sharedMemBytes, hStream, kernelParams, extra, startEvent, stopEvent, flags, params); - hip::Function* function = hip::Function::asFunction(f); - amd::Kernel* kernel = function->function_; + hip::DeviceFunc* function = hip::DeviceFunc::asFunction(f); + amd::Kernel* kernel = function->kernel(); - amd::ScopedLock lock(function->lock_); + amd::ScopedLock lock(function->dflock_); hip::Event* eStart = reinterpret_cast(startEvent); hip::Event* eStop = reinterpret_cast(stopEvent); @@ -557,7 +378,7 @@ extern "C" hipError_t hipLaunchKernel(const void *hostFunction, size_t sharedMemBytes, hipStream_t stream) { - HIP_INIT_API(NONE, hostFunction, gridDim, blockDim, args, sharedMemBytes, stream); + HIP_INIT_API(hipLaunchKernel, hostFunction, gridDim, blockDim, args, sharedMemBytes, stream); HIP_RETURN(ihipLaunchKernel(hostFunction, gridDim, blockDim, args, sharedMemBytes, stream, nullptr, nullptr, 0)); } @@ -571,7 +392,7 @@ extern "C" hipError_t hipExtLaunchKernel(const void* hostFunction, hipEvent_t stopEvent, int flags) { - HIP_INIT_API(NONE, hostFunction, gridDim, blockDim, args, sharedMemBytes, stream); + HIP_INIT_API(hipExtLaunchKernel, hostFunction, gridDim, blockDim, args, sharedMemBytes, stream); HIP_RETURN(ihipLaunchKernel(hostFunction, gridDim, blockDim, args, sharedMemBytes, stream, startEvent, stopEvent, flags)); } @@ -583,10 +404,8 @@ hipError_t hipLaunchCooperativeKernel(const void* f, sharedMemBytes, hStream); int deviceId = ihipGetDevice(); - hipFunction_t func = PlatformState::instance().getFunc(f, deviceId); - if (func == nullptr) { - HIP_RETURN(hipErrorInvalidDeviceFunction); - } + hipFunction_t func = nullptr; + HIP_RETURN_ONFAIL(PlatformState::instance().getStatFunc(&func, f, deviceId)); HIP_RETURN(ihipModuleLaunchKernel(func, gridDim.x * blockDim.x, gridDim.y * blockDim.y, gridDim.z * blockDim.z, blockDim.x, blockDim.y, blockDim.z, @@ -650,7 +469,7 @@ hipError_t ihipLaunchCooperativeKernelMultiDevice(hipLaunchParams* launchParamsL for (size_t dev = 0; dev < g_devices.size(); ++dev) { // Find the matching device and request the kernel function if (&queue->vdev()->device() == g_devices[dev]->devices()[0]) { - func = PlatformState::instance().getFunc(launch.func, dev); + IHIP_RETURN_ONFAIL(PlatformState::instance().getStatFunc(&func, launch.func, dev)); // Save ROCclr index of the first device in the launch if (i == 0) { firstDevice = queue->vdev()->device().index(); @@ -714,7 +533,7 @@ hipError_t hipModuleGetTexRef(textureReference** texRef, hipModule_t hmod, const } /* Get address and size for the global symbol */ - if (!PlatformState::instance().getTexRef(name, hmod, texRef)) { + if (!PlatformState::instance().getDynTexRef(name, hmod, texRef)) { DevLogPrintfError("Cannot get texRef for name: %s at module:0x%x \n", name, hmod); HIP_RETURN(hipErrorNotFound); diff --git a/rocclr/hip_platform.cpp b/rocclr/hip_platform.cpp index 5a96bcbc68..23a1b3afc6 100755 --- a/rocclr/hip_platform.cpp +++ b/rocclr/hip_platform.cpp @@ -20,6 +20,7 @@ #include #include +#include "hip_platform.hpp" #include "hip_internal.hpp" #include "platform/program.hpp" #include "platform/runtime.hpp" @@ -39,23 +40,6 @@ struct __CudaFatBinaryWrapper { void* dummy1; }; -#define CLANG_OFFLOAD_BUNDLER_MAGIC_STR "__CLANG_OFFLOAD_BUNDLE__" -#define HIP_AMDGCN_AMDHSA_TRIPLE "hip-amdgcn-amd-amdhsa" -#define HCC_AMDGCN_AMDHSA_TRIPLE "hcc-amdgcn-amd-amdhsa-" - -struct __ClangOffloadBundleDesc { - uint64_t offset; - uint64_t size; - uint64_t tripleSize; - const char triple[1]; -}; - -struct __ClangOffloadBundleHeader { - const char magic[sizeof(CLANG_OFFLOAD_BUNDLER_MAGIC_STR) - 1]; - uint64_t numBundles; - __ClangOffloadBundleDesc desc[1]; -}; - hipError_t hipModuleGetGlobal(hipDeviceptr_t* dptr, size_t* bytes, hipModule_t hmod, const char* name); @@ -85,61 +69,7 @@ static bool isCompatibleCodeObject(const std::string& codeobj_target_id, return codeobj_target_id == short_name; } -// Extracts code objects from fat binary in data for device names given in devices. -// Returns true if code objects are extracted successfully. -hipError_t __hipExtractCodeObjectFromFatBinary(const void* data, - const std::vector& devices, - std::vector>& code_objs) -{ - std::string magic((const char*)data, sizeof(CLANG_OFFLOAD_BUNDLER_MAGIC_STR) - 1); - if (magic.compare(CLANG_OFFLOAD_BUNDLER_MAGIC_STR)) { - return hipErrorInvalidKernelFile; - } - - code_objs.resize(devices.size()); - const auto obheader = reinterpret_cast(data); - const auto* desc = &obheader->desc[0]; - unsigned num_code_objs = 0; - for (uint64_t i = 0; i < obheader->numBundles; ++i, - desc = reinterpret_cast( - reinterpret_cast(&desc->triple[0]) + desc->tripleSize)) { - - std::size_t offset = 0; - if (!std::strncmp(desc->triple, HIP_AMDGCN_AMDHSA_TRIPLE, - sizeof(HIP_AMDGCN_AMDHSA_TRIPLE) - 1)) { - offset = sizeof(HIP_AMDGCN_AMDHSA_TRIPLE); //For code objects created by CLang - } else if (!std::strncmp(desc->triple, HCC_AMDGCN_AMDHSA_TRIPLE, - sizeof(HCC_AMDGCN_AMDHSA_TRIPLE) - 1)) { - offset = sizeof(HCC_AMDGCN_AMDHSA_TRIPLE); //For code objects created by Hcc - } else { - continue; - } - std::string target(desc->triple + offset, desc->tripleSize - offset); - - const void *image = reinterpret_cast( - reinterpret_cast(obheader) + desc->offset); - size_t size = desc->size; - - for (size_t dev = 0; dev < devices.size(); ++dev) { - const char* name = devices[dev]; - - if (!isCompatibleCodeObject(target, name)) { - continue; - } - code_objs[dev] = std::make_pair(image, size); - num_code_objs++; - } - } - if (num_code_objs == devices.size()) { - return hipSuccess; - } else { - DevLogError("hipErrorNoBinaryForGpu: Coudn't find binary for current devices!"); - guarantee(false); //Aborting the program - return hipErrorNoBinaryForGpu; - } -} - -extern "C" std::vector>* __hipRegisterFatBinary(const void* data) +extern "C" hip::FatBinaryInfoType* __hipRegisterFatBinary(const void* data) { const __CudaFatBinaryWrapper* fbwrapper = reinterpret_cast(data); if (fbwrapper->magic != __hipFatMAGIC2 || fbwrapper->version != 1) { @@ -151,169 +81,6 @@ extern "C" std::vector>* __hipRegisterFatBinary(con return PlatformState::instance().addFatBinary(fbwrapper->binary); } -void PlatformState::digestFatBinary(const void* data, std::vector>& programs) -{ - if (programs.size() > 0) { - return; - } - - std::vector> code_objs; - std::vector devices; - for (size_t dev = 0; dev < g_devices.size(); ++dev) { - devices.push_back(g_devices[dev]->devices()[0]->info().name_); - } - - if (hipSuccess != __hipExtractCodeObjectFromFatBinary((char*)data, devices, code_objs)) { - return; - } - - programs.resize(g_devices.size()); - - for (size_t dev = 0; dev < g_devices.size(); ++dev) { - amd::Context* ctx = g_devices[dev]->asContext(); - amd::Program* program = new amd::Program(*ctx); - if (program == nullptr) { - return; - } - programs.at(dev) = std::make_pair(reinterpret_cast(as_cl(program)) , false); - code_obj_.insert(std::make_pair(program, std::make_pair(code_objs[dev].first, code_objs[dev].second))); - } -} - -void PlatformState::init() -{ - amd::ScopedLock lock(lock_); - - if(initialized_ || g_devices.empty()) { - return; - } - initialized_ = true; - - for (auto& it : modules_) { - digestFatBinary(it.first, it.second); - } - for (auto& it : functions_) { - it.second.functions.resize(g_devices.size()); - } - for (auto& it : vars_) { - it.second.rvars.resize(g_devices.size()); - } -} - -bool PlatformState::unregisterFunc(hipModule_t hmod) { - amd::ScopedLock lock(lock_); - auto mod_it = module_map_.find(hmod); - if (mod_it != module_map_.cend()) { - PlatformState::Module* mod_ptr = mod_it->second; - if(mod_ptr != nullptr) { - for (auto func_it = mod_ptr->functions_.begin(); func_it != mod_ptr->functions_.end(); ++func_it) { - PlatformState::DeviceFunction &devFunc = func_it->second; - for (size_t dev = 0; dev < g_devices.size(); ++dev) { - if (devFunc.functions[dev] != 0) { - hip::Function* f = reinterpret_cast(devFunc.functions[dev]); - delete f; - } - } - delete devFunc.modules; - } - delete mod_ptr; - } - module_map_.erase(mod_it); - } - return true; -} - -std::vector< std::pair >* PlatformState::unregisterVar(hipModule_t hmod) { - amd::ScopedLock lock(lock_); - std::vector< std::pair >* rmodules = nullptr; - auto it = vars_.begin(); - while (it != vars_.end()) { - DeviceVar& dvar = it->second; - if ((*dvar.modules)[0].first == hmod) { - rmodules = dvar.modules; - if (dvar.shadowAllocated) { - texture* tex_hptr - = reinterpret_cast *>(dvar.shadowVptr); - delete tex_hptr; - } - for (size_t dev = 0; dev < g_devices.size(); ++dev) { - if (dvar.rvars[dev].getdeviceptr()) { - amd::MemObjMap::RemoveMemObj(dvar.rvars[dev].getdeviceptr()); - dvar.rvars[dev].amd_mem_obj()->release(); - } - } - vars_.erase(it++); - } else { - ++it; - } - } - return rmodules; -} - -PlatformState::DeviceVar* PlatformState::findVar(std::string hostVar, int deviceId, hipModule_t hmod) { - DeviceVar* dvar = nullptr; - if (hmod != nullptr) { - // If module is provided, then get the var only from that module - auto var_range = vars_.equal_range(hostVar); - for (auto it = var_range.first; it != var_range.second; ++it) { - if ((*it->second.modules)[deviceId].first == hmod) { - dvar = &(it->second); - break; - } - } - } else { - // If var count is < 2, return the var - if (vars_.count(hostVar) < 2) { - auto it = vars_.find(hostVar); - dvar = ((it == vars_.end()) ? nullptr : &(it->second)); - } else { - // If var count is > 2, return the original var, - // if original var count != 1, return vars_.end()/Invalid - size_t orig_global_count = 0; - auto var_range = vars_.equal_range(hostVar); - for (auto it = var_range.first; it != var_range.second; ++it) { - // when dyn_undef is set, it is a shadow var - if (it->second.dyn_undef == false) { - ++orig_global_count; - dvar = &(it->second); - } - } - dvar = ((orig_global_count == 1) ? dvar : nullptr); - } - } - - return dvar; -} - -bool PlatformState::findSymbol(const void *hostVar, - hipModule_t &hmod, std::string &symbolName) { - auto it = symbols_.find(hostVar); - if (it != symbols_.end()) { - hmod = it->second.first; - symbolName = it->second.second; - return true; - } - DevLogPrintfError("Could not find the Symbol: %s \n", symbolName.c_str()); - return false; -} - -void PlatformState::registerVarSym(const void* hostVar, hipModule_t hmod, const char* symbolName) { - amd::ScopedLock lock(lock_); - symbols_.insert(std::make_pair(hostVar, std::make_pair(hmod, std::string(symbolName)))); -} - -void PlatformState::registerVar(const char* hostvar, - const DeviceVar& rvar) { - amd::ScopedLock lock(lock_); - vars_.insert(std::make_pair(std::string(hostvar), rvar)); -} - -void PlatformState::registerFunction(const void* hostFunction, - const DeviceFunction& func) { - amd::ScopedLock lock(lock_); - functions_.insert(std::make_pair(hostFunction, func)); -} - bool ihipGetFuncAttributes(const char* func_name, amd::Program* program, hipFuncAttributes* func_attr) { device::Program* dev_program = program->getDeviceProgram(*hip::getCurrentDevice()->devices()[0]); @@ -344,15 +111,17 @@ bool ihipGetFuncAttributes(const char* func_name, amd::Program* program, hipFunc bool PlatformState::getShadowVarInfo(std::string var_name, hipModule_t hmod, void** var_addr, size_t* var_size) { - DeviceVar* dvar = findVar(var_name, ihipGetDevice(), hmod); - if (dvar != nullptr) { - *var_addr = dvar->shadowVptr; - *var_size = dvar->size; + + amd::ScopedLock lock(lock_); + if (hipSuccess == getDynGlobalVar(var_name.c_str(), ihipGetDevice(), hmod, var_addr, var_size)) { return true; - } else { - DevLogPrintfError("Cannot find Var name: %s in module: 0x%x \n", var_name.c_str(), hmod); - return false; } + + if (hipSuccess == getStatGlobalVarByName(var_name, ihipGetDevice(), hmod, var_addr, var_size)) { + return true; + } + + return false; } bool CL_CALLBACK getSvarInfo(cl_program program, std::string var_name, void** var_addr, @@ -361,275 +130,6 @@ bool CL_CALLBACK getSvarInfo(cl_program program, std::string var_name, void** va var_addr, var_size); } -bool PlatformState::registerModFuncs(std::vector& func_names, hipModule_t* module) { - amd::ScopedLock lock(lock_); - PlatformState::Module* mod_ptr = new PlatformState::Module(*module); - - for (auto it = func_names.begin(); it != func_names.end(); ++it) { - auto modules = new std::vector >(g_devices.size()); - for (size_t dev = 0; dev < g_devices.size(); ++dev) { - modules->at(dev) = std::make_pair(*module, true); - } - - PlatformState::DeviceFunction dfunc{*it, modules, - std::vector(g_devices.size(), 0)}; - mod_ptr->functions_.insert(std::make_pair(*it, dfunc)); - } - - module_map_.insert(std::make_pair(*module, mod_ptr)); - return true; -} - -bool PlatformState::findModFunc(hipFunction_t* hfunc, hipModule_t hmod, const char* name) { - amd::ScopedLock lock(lock_); - - auto mod_it = module_map_.find(hmod); - if (mod_it != module_map_.cend()) { - assert(mod_it->second != nullptr); - auto func_it = mod_it->second->functions_.find(name); - if (func_it != mod_it->second->functions_.cend()) { - PlatformState::DeviceFunction& devFunc = func_it->second; - if (devFunc.functions[ihipGetDevice()] == 0) { - if(!createFunc(&devFunc.functions[ihipGetDevice()], hmod, name)) { - DevLogPrintfError("Could not create a function: %s at module: 0x%x \n", name, hmod); - return false; - } - } - *hfunc = devFunc.functions[ihipGetDevice()]; - return true; - } - } - DevLogPrintfError("Cannot find module: 0x%x in PlatformState Module Map \n", hmod); - return false; -} - -bool PlatformState::createFunc(hipFunction_t* hfunc, hipModule_t hmod, const char* name) { - amd::Program* program = as_amd(reinterpret_cast(hmod)); - - const amd::Symbol* symbol = program->findSymbol(name); - if (!symbol) { - DevLogPrintfError("Cannot find Symbol with name: %s \n", name); - return false; - } - - amd::Kernel* kernel = new amd::Kernel(*program, *symbol, name); - if (!kernel) { - DevLogPrintfError("Could not create a new kernel with name: %s \n", name); - return false; - } - - hip::Function* f = new hip::Function(kernel); - if (!f) { - DevLogPrintfError("Could not create a new function with name: %s \n", name); - return false; - } - - *hfunc = f->asHipFunction(); - - return true; -} - - -hipFunction_t PlatformState::getFunc(const void* hostFunction, int deviceId) { - amd::ScopedLock lock(lock_); - const auto it = functions_.find(hostFunction); - if (it != functions_.cend()) { - PlatformState::DeviceFunction& devFunc = it->second; - if (devFunc.functions[deviceId] == 0) { - hipModule_t module = (*devFunc.modules)[deviceId].first; - if (!(*devFunc.modules)[deviceId].second) { - amd::Program* program = as_amd(reinterpret_cast(module)); - amd::Context* ctx = g_devices[deviceId]->asContext(); - auto code_obj_it = code_obj_.find(program); - if (code_obj_.end() == code_obj_it) { - DevLogError("Cannot find image & size for static symbols"); - guarantee(false); //Aborting the program - return nullptr; - } - if (CL_SUCCESS != program->addDeviceProgram(*ctx->devices()[0], code_obj_it->second.first, - code_obj_it->second.second, false)) { - DevLogError("Cannot add Device Program"); - guarantee(false); //Aborting the program - return nullptr; - } - program->setVarInfoCallBack(&getSvarInfo); - if (CL_SUCCESS != program->build(g_devices[deviceId]->devices(), nullptr, nullptr, nullptr, - kOptionChangeable, kNewDevProg)) { - DevLogPrintfError("Build error for module: 0x%x at device: %u \n", module, deviceId); - return nullptr; - } - (*devFunc.modules)[deviceId].second = true; - } - hipFunction_t function = nullptr; - if (createFunc(&function, module, devFunc.deviceName.c_str()) && - function != nullptr) { - devFunc.functions[deviceId] = function; - } else { - DevLogPrintfError("__hipRegisterFunction cannot find kernel %s for device %d\n", - devFunc.deviceName.c_str(), deviceId); - return nullptr; - } - } - return devFunc.functions[deviceId]; - } - DevLogPrintfError("Cannot find function: 0x%x in PlatformState \n", hostFunction); - return nullptr; -} - -bool PlatformState::getFuncAttr(const void* hostFunction, - hipFuncAttributes* func_attr) { - if (func_attr == nullptr) { - return false; - } - - const auto it = functions_.find(hostFunction); - if (it == functions_.cend()) { - DevLogPrintfError("Cannot find hostFunction 0x%x \n", hostFunction); - return false; - } - - PlatformState::DeviceFunction& devFunc = it->second; - int deviceId = ihipGetDevice(); - - /* If module has not been initialized yet, build the kernel now*/ - if (!(*devFunc.modules)[deviceId].second) { - if (nullptr == PlatformState::instance().getFunc(hostFunction, deviceId)) { - DevLogPrintfError("Cannot get hostFunction: 0x%x for deviceId:%d \n", hostFunction, deviceId); - return false; - } - } - - amd::Program* program = as_amd(reinterpret_cast((*devFunc.modules)[deviceId].first)); - if (!ihipGetFuncAttributes(devFunc.deviceName.c_str(), program, func_attr)) { - DevLogPrintfError("Cannot get Func attributes for function: %s \n", - devFunc.deviceName.c_str()); - return false; - } - return true; -} - -bool PlatformState::getTexRef(const char* hostVar, hipModule_t hmod, textureReference** texRef) { - amd::ScopedLock lock(lock_); - DeviceVar* dvar = findVar(std::string(hostVar), ihipGetDevice(), hmod); - if (dvar == nullptr) { - DevLogPrintfError("Cannot find var:%s for creating texture reference at module: 0x%x \n", - hostVar, hmod); - return false; - } - - switch (dvar->kind) { - case PlatformState::DVK_Variable: - // TODO: Need to define a target-specific symbol info to indicate the device - // variable kind, i.e. regular variable, texture or surface. - // Before that, have to assume the specified variable is a texture or - // surface reference variable. - dvar->kind = DVK_Texture; - // FALL THROUGH - case PlatformState::DVK_Texture: - break; - default: - // If it's already used as non-texture variable, bail out. - return false; - } - - if (!dvar->shadowVptr) { - dvar->shadowVptr = new texture{}; - dvar->shadowAllocated = true; - } - *texRef = reinterpret_cast(dvar->shadowVptr); - registerVarSym(dvar->shadowVptr, hmod, hostVar); - - return true; -} - -bool PlatformState::getGlobalVar(const char* hostVar, int deviceId, hipModule_t hmod, - hipDeviceptr_t* dev_ptr, size_t* size_ptr) { - amd::ScopedLock lock(lock_); - DeviceVar* dvar = findVar(std::string(hostVar), deviceId, hmod); - if (dvar != nullptr) { - if (dvar->rvars[deviceId].getdeviceptr() == nullptr) { - size_t sym_size = 0; - hipDeviceptr_t device_ptr = nullptr; - amd::Memory* amd_mem_obj = nullptr; - - if (!(*dvar->modules)[deviceId].second) { - amd::Program* program = as_amd(reinterpret_cast((*dvar->modules)[deviceId].first)); - amd::Context* ctx = g_devices[deviceId]->asContext(); - auto code_obj_it = code_obj_.find(program); - if (code_obj_.end() == code_obj_it) { - DevLogError("Cannot find image & size for static symbols"); - guarantee(false); //Aborting the program - return false; - } - if (CL_SUCCESS != program->addDeviceProgram(*ctx->devices()[0], code_obj_it->second.first, - code_obj_it->second.second, false)) { - DevLogError("Cannot add Device Program"); - guarantee(false) //Aborting the program - return false; - } - program->setVarInfoCallBack(&getSvarInfo); - if (CL_SUCCESS != program->build(g_devices[deviceId]->devices(), nullptr, nullptr, nullptr, - kOptionChangeable, kNewDevProg)) { - DevLogPrintfError("Build Failure for module: 0x%x \n", hmod); - return false; - } - (*dvar->modules)[deviceId].second = true; - } - if((hipSuccess == ihipCreateGlobalVarObj(dvar->hostVar.c_str(), (*dvar->modules)[deviceId].first, - &amd_mem_obj, &device_ptr, &sym_size)) - && (device_ptr != nullptr)) { - dvar->rvars[deviceId].size_ = sym_size; - dvar->rvars[deviceId].devicePtr_ = device_ptr; - dvar->rvars[deviceId].amd_mem_obj_ = amd_mem_obj; - amd::MemObjMap::AddMemObj(device_ptr, amd_mem_obj); - } else { - DevLogPrintfError("__hipRegisterVar cannot find Var: %s for deviceId: 0x%x \n", - dvar->hostVar.c_str(), deviceId); - return false; - } - } - *size_ptr = dvar->rvars[deviceId].getvarsize(); - *dev_ptr = dvar->rvars[deviceId].getdeviceptr(); - return true; - } else { - DevLogPrintfError("Could not find global var: %s at module:0x%x \n", hostVar, hmod); - return false; - } -} - -bool PlatformState::getGlobalVarFromSymbol(const void* hostVar, int deviceId, - hipDeviceptr_t* dev_ptr, - size_t* size_ptr) { - hipModule_t hmod; - std::string symbolName; - if (!PlatformState::instance().findSymbol(hostVar, hmod, symbolName)) { - return false; - } - return PlatformState::instance().getGlobalVar(symbolName.c_str(), - ihipGetDevice(), hmod, - dev_ptr, size_ptr); -} - -void PlatformState::setupArgument(const void *arg, size_t size, size_t offset) { - auto& arguments = execStack_.top().arguments_; - - if (arguments.size() < offset + size) { - arguments.resize(offset + size); - } - - ::memcpy(&arguments[offset], arg, size); -} - -void PlatformState::configureCall(dim3 gridDim, dim3 blockDim, size_t sharedMem, - hipStream_t stream) { - execStack_.push(ihipExec_t{gridDim, blockDim, sharedMem, stream}); -} - -void PlatformState::popExec(ihipExec_t& exec) { - exec = std::move(execStack_.top()); - execStack_.pop(); -} - namespace { const int HIP_ENABLE_DEFERRED_LOADING{[] () { char *var = getenv("HIP_ENABLE_DEFERRED_LOADING"); @@ -638,7 +138,7 @@ const int HIP_ENABLE_DEFERRED_LOADING{[] () { } /* namespace */ extern "C" void __hipRegisterFunction( - std::vector >* modules, + hip::FatBinaryInfoType* modules, const void* hostFunction, char* deviceFunction, const char* deviceName, @@ -647,14 +147,16 @@ extern "C" void __hipRegisterFunction( uint3* bid, dim3* blockDim, dim3* gridDim, - int* wSize) -{ - PlatformState::DeviceFunction func{ std::string{deviceName}, modules, std::vector{g_devices.size()}}; - PlatformState::instance().registerFunction(hostFunction, func); + int* wSize) { + hip::Function* func = new hip::Function(std::string(deviceName), modules); + PlatformState::instance().registerStatFunction(hostFunction, func); if (!HIP_ENABLE_DEFERRED_LOADING) { HIP_INIT(); - for (size_t i = 0; i < g_devices.size(); ++i) { - PlatformState::instance().getFunc(hostFunction, i); + hipFunction_t hfunc = nullptr; + hipError_t hip_error = hipSuccess; + for (size_t dev_idx = 0; dev_idx < g_devices.size(); ++dev_idx) { + hip_error = PlatformState::instance().getStatFunc(&hfunc, hostFunction, dev_idx); + guarantee(hip_error == hipSuccess); } } } @@ -665,7 +167,7 @@ extern "C" void __hipRegisterFunction( // track of the value of the device side global variable between kernel // executions. extern "C" void __hipRegisterVar( - std::vector >* modules, // The device modules containing code object + hip::FatBinaryInfoType* modules, // The device modules containing code object void* var, // The shadow variable in host code char* hostVar, // Variable name in host code char* deviceVar, // Variable name in device code @@ -674,70 +176,32 @@ extern "C" void __hipRegisterVar( int constant, // Whether this variable is constant int global) // Unknown, always 0 { - PlatformState::DeviceVar dvar{PlatformState::DVK_Variable, - var, - std::string{hostVar}, - size, - modules, - std::vector{g_devices.size()}, - false, - /*type*/ 0, - /*norm*/ 0}; - - PlatformState::instance().registerVar(hostVar, dvar); - PlatformState::instance().registerVarSym(var, nullptr, deviceVar); + hip::Var* var_ptr = new hip::Var(std::string(hostVar), hip::Var::DeviceVarKind::DVK_Variable, size, 0, 0, modules); + PlatformState::instance().registerStatGlobalVar(var, var_ptr); } -extern "C" void __hipRegisterSurface(std::vector>* - modules, // The device modules containing code object +extern "C" void __hipRegisterSurface(hip::FatBinaryInfoType* modules, // The device modules containing code object void* var, // The shadow variable in host code char* hostVar, // Variable name in host code char* deviceVar, // Variable name in device code int type, int ext) { - PlatformState::DeviceVar dvar{PlatformState::DVK_Surface, - var, - std::string{hostVar}, - sizeof(surfaceReference), // Copy whole surfaceReference - modules, - std::vector{g_devices.size()}, - false, - type, - /*norm*/ 0}; - PlatformState::instance().registerVar(hostVar, dvar); - PlatformState::instance().registerVarSym(var, nullptr, deviceVar); + hip::Var* var_ptr = new hip::Var(std::string(hostVar), hip::Var::DeviceVarKind::DVK_Surface, sizeof(surfaceReference), 0, 0, modules); + PlatformState::instance().registerStatGlobalVar(var, var_ptr); } -extern "C" void __hipRegisterTexture(std::vector>* - modules, // The device modules containing code object +extern "C" void __hipRegisterTexture(hip::FatBinaryInfoType* modules, // The device modules containing code object void* var, // The shadow variable in host code char* hostVar, // Variable name in host code char* deviceVar, // Variable name in device code int type, int norm, int ext) { - PlatformState::DeviceVar dvar{PlatformState::DVK_Texture, - var, - std::string{hostVar}, - sizeof(textureReference), // Copy whole textureReference so far. - modules, - std::vector{g_devices.size()}, - false, - type, - norm}; - PlatformState::instance().registerVar(hostVar, dvar); - PlatformState::instance().registerVarSym(var, nullptr, deviceVar); + hip::Var* var_ptr = new hip::Var(std::string(hostVar), hip::Var::DeviceVarKind::DVK_Texture, sizeof(textureReference), 0, 0, modules); + PlatformState::instance().registerStatGlobalVar(var, var_ptr); } -extern "C" void __hipUnregisterFatBinary(std::vector< std::pair >* modules) +extern "C" void __hipUnregisterFatBinary(hip::FatBinaryInfoType* modules) { HIP_INIT(); - std::for_each(modules->begin(), modules->end(), [](std::pair module){ - if (module.first != nullptr) { - as_amd(reinterpret_cast(module.first))->release(); - } - }); - if (modules->size() > 0) { - PlatformState::instance().unregisterVar((*modules)[0].first); - } PlatformState::instance().removeFatBinary(modules); } @@ -808,8 +272,9 @@ extern "C" hipError_t hipLaunchByPtr(const void *hostFunction) DevLogPrintfError("Wrong DeviceId: %d \n", deviceId); HIP_RETURN(hipErrorNoDevice); } - hipFunction_t func = PlatformState::instance().getFunc(hostFunction, deviceId); - if (func == nullptr) { + hipFunction_t func = nullptr; + hipError_t hip_error = PlatformState::instance().getStatFunc(&func, hostFunction, deviceId); + if ((hip_error != hipSuccess) || (func == nullptr)) { DevLogPrintfError("Could not retrieve hostFunction: 0x%x \n", hostFunction); HIP_RETURN(hipErrorInvalidDeviceFunction); } @@ -830,38 +295,20 @@ extern "C" hipError_t hipLaunchByPtr(const void *hostFunction) hipError_t hipGetSymbolAddress(void** devPtr, const void* symbol) { HIP_INIT_API(hipGetSymbolAddress, devPtr, symbol); - hipModule_t hmod; - std::string symbolName; - if (!PlatformState::instance().findSymbol(symbol, hmod, symbolName)) { - DevLogPrintfError("Cannot find symbol: %s \n", symbolName.c_str()); - HIP_RETURN(hipErrorInvalidSymbol); - } - size_t size = 0; - if(!PlatformState::instance().getGlobalVar(symbolName.c_str(), ihipGetDevice(), hmod, - devPtr, &size)) { - DevLogPrintfError("Cannot find global variable device ptr for symbol: %s at device: %d \n", - symbolName.c_str(), ihipGetDevice()); - HIP_RETURN(hipErrorInvalidSymbol); - } + hipError_t hip_error = hipSuccess; + size_t sym_size = 0; + + HIP_RETURN_ONFAIL(PlatformState::instance().getStatGlobalVar(symbol, ihipGetDevice(), devPtr, &sym_size)); + HIP_RETURN(hipSuccess); } hipError_t hipGetSymbolSize(size_t* sizePtr, const void* symbol) { HIP_INIT_API(hipGetSymbolSize, sizePtr, symbol); - hipModule_t hmod; - std::string symbolName; - if (!PlatformState::instance().findSymbol(symbol, hmod, symbolName)) { - DevLogPrintfError("Cannot find symbol: %s \n", symbolName.c_str()); - HIP_RETURN(hipErrorInvalidSymbol); - } - hipDeviceptr_t devPtr = nullptr; - if (!PlatformState::instance().getGlobalVar(symbolName.c_str(), ihipGetDevice(), hmod, - &devPtr, sizePtr)) { - DevLogPrintfError("Cannot find global variable device ptr for symbol: %s at device: %d \n", - symbolName.c_str(), ihipGetDevice()); - HIP_RETURN(hipErrorInvalidSymbol); - } + hipDeviceptr_t device_ptr = nullptr; + HIP_RETURN_ONFAIL(PlatformState::instance().getStatGlobalVar(symbol, ihipGetDevice(), &device_ptr, sizePtr)); + HIP_RETURN(hipSuccess); } @@ -897,8 +344,8 @@ hipError_t ihipOccupancyMaxActiveBlocksPerMultiprocessor( const amd::Device& device, hipFunction_t func, int inputBlockSize, size_t dynamicSMemSize, bool bCalcPotentialBlkSz) { - hip::Function* function = hip::Function::asFunction(func); - const amd::Kernel& kernel = *function->function_; + hip::DeviceFunc* function = hip::DeviceFunc::asFunction(func); + const amd::Kernel& kernel = *function->kernel(); const device::Kernel::WorkGroupInfo* wrkGrpInfo = kernel.getDeviceKernel(device)->workGroupInfo(); if (bCalcPotentialBlkSz == false) { @@ -989,9 +436,10 @@ hipError_t hipOccupancyMaxPotentialBlockSize(int* gridSize, int* blockSize, if ((gridSize == nullptr) || (blockSize == nullptr)) { HIP_RETURN(hipErrorInvalidValue); } - hipFunction_t func = PlatformState::instance().getFunc(f, ihipGetDevice()); - if (func == nullptr) { - HIP_RETURN(hipErrorInvalidValue); + hipFunction_t func = nullptr; + hipError_t hip_error = PlatformState::instance().getStatFunc(&func, f, ihipGetDevice()); + if ((hip_error != hipSuccess) || (func == nullptr)) { + return HIP_RETURN(hipErrorInvalidValue); } const amd::Device& device = *hip::getCurrentDevice()->devices()[0]; int max_blocks_per_grid = 0; @@ -1093,9 +541,10 @@ hipError_t hipOccupancyMaxActiveBlocksPerMultiprocessor(int* numBlocks, HIP_RETURN(hipErrorInvalidValue); } - hipFunction_t func = PlatformState::instance().getFunc(f, ihipGetDevice()); - if (func == nullptr) { - HIP_RETURN(hipErrorInvalidValue); + hipFunction_t func = nullptr; + hipError_t hip_error = PlatformState::instance().getStatFunc(&func, f, ihipGetDevice()); + if ((hip_error != hipSuccess) || (func == nullptr)) { + return HIP_RETURN(hipErrorInvalidValue); } const amd::Device& device = *hip::getCurrentDevice()->devices()[0]; @@ -1118,9 +567,10 @@ hipError_t hipOccupancyMaxActiveBlocksPerMultiprocessorWithFlags(int* numBlocks, HIP_RETURN(hipErrorInvalidValue); } - hipFunction_t func = PlatformState::instance().getFunc(f, ihipGetDevice()); - if (func == nullptr) { - HIP_RETURN(hipErrorInvalidValue); + hipFunction_t func = nullptr; + hipError_t hip_error = PlatformState::instance().getStatFunc(&func, f, ihipGetDevice()); + if ((hip_error != hipSuccess) || (func == nullptr)) { + return HIP_RETURN(hipErrorInvalidValue); } const amd::Device& device = *hip::getCurrentDevice()->devices()[0]; @@ -1290,10 +740,10 @@ const std::vector& modules() { if (magic.compare(CLANG_OFFLOAD_BUNDLER_MAGIC_STR)) continue; - const auto obheader = reinterpret_cast(&bundle[0]); + const auto obheader = reinterpret_cast(&bundle[0]); const auto* desc = &obheader->desc[0]; for (uint64_t i = 0; i < obheader->numBundles; ++i, - desc = reinterpret_cast( + desc = reinterpret_cast( reinterpret_cast(&desc->triple[0]) + desc->tripleSize)) { std::string triple(desc->triple, sizeof(HCC_AMDGCN_AMDHSA_TRIPLE) - 1); @@ -1336,7 +786,6 @@ const std::unordered_map& functions() return r; } - void hipLaunchKernelGGLImpl( uintptr_t function_address, const dim3& numBlocks, @@ -1391,8 +840,10 @@ hipError_t ihipLaunchKernel(const void* hostFunction, DevLogPrintfError("Wrong Device Id: %d \n", deviceId); HIP_RETURN(hipErrorNoDevice); } - hipFunction_t func = PlatformState::instance().getFunc(hostFunction, deviceId); - if (func == nullptr) { + + hipFunction_t func = nullptr; + hipError_t hip_error = PlatformState::instance().getStatFunc(&func, hostFunction, deviceId); + if ((hip_error != hipSuccess) || (func == nullptr)) { #ifdef ATI_OS_LINUX const auto it = hip_impl::functions().find(reinterpret_cast(hostFunction)); if (it == hip_impl::functions().cend()) { @@ -1449,3 +900,167 @@ extern "C" float __gnu_h2f_ieee(unsigned short h){ extern "C" unsigned short __gnu_f2h_ieee(float f){ return (unsigned short)__convert_float_to_half(f); } + +void PlatformState::init() +{ + amd::ScopedLock lock(lock_); + + if(initialized_ || g_devices.empty()) { + return; + } + initialized_ = true; + + for (auto& it : statCO_.modules_) { + digestFatBinary(it.first, it.second); + } + + for (auto &it : statCO_.vars_) { + it.second->resize_dVar(g_devices.size()); + } + + for (auto &it : statCO_.functions_) { + it.second->resize_dFunc(g_devices.size()); + } +} + +hipError_t PlatformState::loadModule(hipModule_t *module, const char* fname, const void* image) { + amd::ScopedLock lock(lock_); + + hip::DynCO* dynCo = new hip::DynCO(); + hipError_t hip_error = dynCo->loadCodeObject(fname, image); + if (hip_error != hipSuccess) { + delete dynCo; + return hip_error; + } + + *module = dynCo->module(); + assert(*module != nullptr); + + if (dynCO_map_.find(*module) != dynCO_map_.end()) { + return hipErrorAlreadyMapped; + } + dynCO_map_.insert(std::make_pair(*module, dynCo)); + + return hipSuccess; +} + +hipError_t PlatformState::unloadModule(hipModule_t hmod) { + amd::ScopedLock lock(lock_); + + auto it = dynCO_map_.find(hmod); + if (it == dynCO_map_.end()) { + return hipErrorNotFound; + } + + delete it->second; + dynCO_map_.erase(hmod); + + return hipSuccess; +} + +hipError_t PlatformState::getDynFunc(hipFunction_t* hfunc, hipModule_t hmod, + const char* func_name) { + amd::ScopedLock lock(lock_); + + auto it = dynCO_map_.find(hmod); + if (it == dynCO_map_.end()) { + DevLogPrintfError("Cannot find the module: 0x%x", hmod); + return hipErrorNotFound; + } + + return it->second->getDynFunc(hfunc, func_name); +} + +hipError_t PlatformState::getDynGlobalVar(const char* hostVar, int deviceId, hipModule_t hmod, + hipDeviceptr_t* dev_ptr, size_t* size_ptr) { + amd::ScopedLock lock(lock_); + + auto it = dynCO_map_.find(hmod); + if (it == dynCO_map_.end()) { + DevLogPrintfError("Cannot find the module: 0x%x", hmod); + return hipErrorNotFound; + } + + hip::DeviceVar* dvar = nullptr; + IHIP_RETURN_ONFAIL(it->second->getDeviceVar(&dvar, hostVar, deviceId)); + *dev_ptr = dvar->device_ptr(); + *size_ptr = dvar->size(); + + return hipSuccess; +} + +hipError_t PlatformState::getDynTexRef(const char* hostVar, hipModule_t hmod, textureReference** texRef) { + amd::ScopedLock lock(lock_); + + auto it = dynCO_map_.find(hmod); + if (it == dynCO_map_.end()) { + DevLogPrintfError("Cannot find the module: 0x%x", hmod); + return hipErrorNotFound; + } + + hip::DeviceVar* dvar = nullptr; + IHIP_RETURN_ONFAIL(it->second->getDeviceVar(&dvar, hostVar, ihipGetDevice())); + + dvar->shadowVptr = new texture(); + *texRef = reinterpret_cast(dvar->shadowVptr); + return hipSuccess; +} + +hipError_t PlatformState::digestFatBinary(const void* data, hip::FatBinaryInfoType& programs) { + return statCO_.digestFatBinary(data, programs); +} + +hip::FatBinaryInfoType* PlatformState::addFatBinary(const void* data) { + return statCO_.addFatBinary(data, initialized_); +} + +hipError_t PlatformState::removeFatBinary(hip::FatBinaryInfoType* module) { + return statCO_.removeFatBinary(module); +} + +hipError_t PlatformState::registerStatFunction(const void* hostFunction, hip::Function* func) { + return statCO_.registerStatFunction(hostFunction, func); +} + +hipError_t PlatformState::registerStatGlobalVar(const void* hostVar, hip::Var* var) { + return statCO_.registerStatGlobalVar(hostVar, var); +} + +hipError_t PlatformState::getStatFunc(hipFunction_t* hfunc, const void* hostFunction, int deviceId) { + return statCO_.getStatFunc(hfunc, hostFunction, deviceId); +} + +hipError_t PlatformState::getStatFuncAttr(hipFuncAttributes* func_attr, const void* hostFunction, int deviceId) { + return statCO_.getStatFuncAttr(func_attr, hostFunction, deviceId); +} + +hipError_t PlatformState::getStatGlobalVar(const void* hostVar, int deviceId, hipDeviceptr_t* dev_ptr, + size_t* size_ptr) { + return statCO_.getStatGlobalVar(hostVar, deviceId, dev_ptr, size_ptr); +} + +hipError_t PlatformState::getStatGlobalVarByName(std::string hostVar, int deviceId, hipModule_t hmod, + hipDeviceptr_t* dev_ptr, size_t* size_ptr) { + return statCO_.getStatGlobalVarByName(hostVar, deviceId, hmod, dev_ptr, size_ptr); +} + +void PlatformState::setupArgument(const void *arg, size_t size, size_t offset) { + auto& arguments = execStack_.top().arguments_; + + if (arguments.size() < offset + size) { + arguments.resize(offset + size); + } + + ::memcpy(&arguments[offset], arg, size); +} + +void PlatformState::configureCall(dim3 gridDim, dim3 blockDim, size_t sharedMem, + hipStream_t stream) { + execStack_.push(ihipExec_t{gridDim, blockDim, sharedMem, stream}); +} + +void PlatformState::popExec(ihipExec_t& exec) { + exec = std::move(execStack_.top()); + execStack_.pop(); +} + diff --git a/rocclr/hip_platform.hpp b/rocclr/hip_platform.hpp old mode 100644 new mode 100755 index fcbfb53bbb..b53a1a750d --- a/rocclr/hip_platform.hpp +++ b/rocclr/hip_platform.hpp @@ -19,11 +19,75 @@ THE SOFTWARE. */ #pragma once +#include "hip_internal.hpp" +#include "hip_fatbin.hpp" #include "device/device.hpp" +#include "hip_code_object.hpp" namespace hip_impl { + hipError_t ihipOccupancyMaxActiveBlocksPerMultiprocessor( int* maxBlocksPerCU, int* numBlocksPerGrid, int* bestBlockSize, const amd::Device& device, hipFunction_t func, int blockSize, size_t dynamicSMemSize, bool bCalcPotentialBlkSz); -} +} /* namespace hip_impl*/ + +class PlatformState { + amd::Monitor lock_{"Guards PlatformState globals", true}; + + /* Singleton object */ + static PlatformState* platform_; + PlatformState() {} + ~PlatformState() {} + +public: + void init(); + + //Dynamic Code Objects functions + hipError_t loadModule(hipModule_t* module, const char* fname, const void* image = nullptr); + hipError_t unloadModule(hipModule_t hmod); + + hipError_t getDynFunc(hipFunction_t *hfunc, hipModule_t hmod, const char* func_name); + hipError_t getDynGlobalVar(const char* hostVar, int deviceId, hipModule_t hmod, + hipDeviceptr_t* dev_ptr, size_t* size_ptr); + hipError_t getDynTexRef(const char* hostVar, hipModule_t hmod, textureReference** texRef); + + /* Singleton instance */ + static PlatformState& instance() { + if (platform_ == nullptr) { + // __hipRegisterFatBinary() will call this when app starts, thus + // there is no multiple entry issue here. + platform_ = new PlatformState(); + } + return *platform_; + } + + //Static Code Objects functions + hip::FatBinaryInfoType* addFatBinary(const void* data); + hipError_t removeFatBinary(hip::FatBinaryInfoType* module); + hipError_t digestFatBinary(const void* data, hip::FatBinaryInfoType& programs); + + hipError_t registerStatFunction(const void* hostFunction, hip::Function* func); + hipError_t registerStatGlobalVar(const void* hostVar, hip::Var* var); + + hipError_t getStatFunc(hipFunction_t* hfunc, const void* hostFunction, int deviceId); + hipError_t getStatFuncAttr(hipFuncAttributes* func_attr, const void* hostFunction, int deviceId); + hipError_t getStatGlobalVar(const void* hostVar, int deviceId, hipDeviceptr_t* dev_ptr, + size_t* size_ptr); + hipError_t getStatGlobalVarByName(std::string hostVar, int deviceId, hipModule_t hmod, + hipDeviceptr_t* dev_ptr, size_t* size_ptr); + + bool getShadowVarInfo(std::string var_name, hipModule_t hmod, + void** var_addr, size_t* var_size); + + //Exec Functions + void setupArgument(const void *arg, size_t size, size_t offset); + void configureCall(dim3 gridDim, dim3 blockDim, size_t sharedMem, hipStream_t stream); + void popExec(ihipExec_t& exec); + +private: + //Dynamic Code Object map, keyin module to get the corresponding object + std::unordered_map dynCO_map_; + hip::StatCO statCO_; //Static Code object var + bool initialized_{false}; +}; diff --git a/rocclr/hip_texture.cpp b/rocclr/hip_texture.cpp index 9d16e3da01..fced181c5b 100755 --- a/rocclr/hip_texture.cpp +++ b/rocclr/hip_texture.cpp @@ -21,6 +21,7 @@ #include #include #include "hip_internal.hpp" +#include "hip_platform.hpp" #include "hip_conversions.hpp" #include "platform/sampler.hpp" @@ -478,10 +479,10 @@ hipError_t hipBindTexture2D(size_t* offset, hipDeviceptr_t refDevPtr = nullptr; size_t refDevSize = 0; - if (!PlatformState::instance().getGlobalVarFromSymbol(texref, ihipGetDevice(), &refDevPtr, - &refDevSize)) { - HIP_RETURN(hipErrorInvalidSymbol); - } + + HIP_RETURN_ONFAIL(PlatformState::instance().getStatGlobalVar(texref, ihipGetDevice(), &refDevPtr, + &refDevSize)); + assert(refDevSize == sizeof(textureReference)); hipError_t err = ihipBindTexture2D(offset, texref, devPtr, desc, width, height, pitch); if (err != hipSuccess) { @@ -525,10 +526,9 @@ hipError_t hipBindTextureToArray(const textureReference* texref, hipDeviceptr_t refDevPtr = nullptr; size_t refDevSize = 0; - if (!PlatformState::instance().getGlobalVarFromSymbol(texref, ihipGetDevice(), &refDevPtr, - &refDevSize)) { - HIP_RETURN(hipErrorInvalidSymbol); - } + HIP_RETURN_ONFAIL(PlatformState::instance().getStatGlobalVar(texref, ihipGetDevice(), &refDevPtr, + &refDevSize)); + assert(refDevSize == sizeof(textureReference)); hipError_t err = ihipBindTextureToArray(texref, array, desc); if (err != hipSuccess) { @@ -572,10 +572,10 @@ hipError_t hipBindTextureToMipmappedArray(const textureReference* texref, hipDeviceptr_t refDevPtr = nullptr; size_t refDevSize = 0; - if (!PlatformState::instance().getGlobalVarFromSymbol(texref, ihipGetDevice(), &refDevPtr, - &refDevSize)) { - HIP_RETURN(hipErrorInvalidSymbol); - } + + HIP_RETURN_ONFAIL(PlatformState::instance().getStatGlobalVar(texref, ihipGetDevice(), &refDevPtr, + &refDevSize)); + assert(refDevSize == sizeof(textureReference)); hipError_t err = ihipBindTextureToMipmappedArray(texref, mipmappedArray, desc); if (err != hipSuccess) { @@ -608,10 +608,8 @@ hipError_t hipBindTexture(size_t* offset, hipDeviceptr_t refDevPtr = nullptr; size_t refDevSize = 0; - if (!PlatformState::instance().getGlobalVarFromSymbol(texref, ihipGetDevice(), &refDevPtr, - &refDevSize)) { - HIP_RETURN(hipErrorInvalidSymbol); - } + HIP_RETURN_ONFAIL(PlatformState::instance().getStatGlobalVar(texref, ihipGetDevice(), &refDevPtr, + &refDevSize)); assert(refDevSize == sizeof(textureReference)); hipError_t err = ihipBindTexture(offset, texref, devPtr, desc, size); if (err != hipSuccess) { @@ -804,10 +802,8 @@ hipError_t hipTexRefSetArray(textureReference* texRef, hipDeviceptr_t refDevPtr = nullptr; size_t refDevSize = 0; - if (!PlatformState::instance().getGlobalVarFromSymbol(texRef, ihipGetDevice(), &refDevPtr, - &refDevSize)) { - HIP_RETURN(hipErrorInvalidSymbol); - } + HIP_RETURN_ONFAIL(PlatformState::instance().getStatGlobalVar(texRef, ihipGetDevice(), &refDevPtr, + &refDevSize)); assert(refDevSize == sizeof(textureReference)); // Any previous address or HIP array state associated with the texture reference is superseded by this function. @@ -882,10 +878,8 @@ hipError_t hipTexRefSetAddress(size_t* ByteOffset, hipDeviceptr_t refDevPtr = nullptr; size_t refDevSize = 0; - if (!PlatformState::instance().getGlobalVarFromSymbol(texRef, ihipGetDevice(), &refDevPtr, - &refDevSize)) { - HIP_RETURN(hipErrorInvalidSymbol); - } + HIP_RETURN_ONFAIL(PlatformState::instance().getStatGlobalVar(texRef, ihipGetDevice(), &refDevPtr, + &refDevSize)); assert(refDevSize == sizeof(textureReference)); // Any previous address or HIP array state associated with the texture reference is superseded by this function. @@ -929,10 +923,8 @@ hipError_t hipTexRefSetAddress2D(textureReference* texRef, hipDeviceptr_t refDevPtr = nullptr; size_t refDevSize = 0; - if (!PlatformState::instance().getGlobalVarFromSymbol(texRef, ihipGetDevice(), &refDevPtr, - &refDevSize)) { - HIP_RETURN(hipErrorInvalidSymbol); - } + HIP_RETURN_ONFAIL(PlatformState::instance().getStatGlobalVar(texRef, ihipGetDevice(), &refDevPtr, + &refDevSize)); assert(refDevSize == sizeof(textureReference)); // Any previous address or HIP array state associated with the texture reference is superseded by this function. @@ -1209,10 +1201,8 @@ hipError_t hipTexRefSetMipmappedArray(textureReference* texRef, hipDeviceptr_t refDevPtr = nullptr; size_t refDevSize = 0; - if (!PlatformState::instance().getGlobalVarFromSymbol(texRef, ihipGetDevice(), &refDevPtr, - &refDevSize)) { - HIP_RETURN(hipErrorInvalidSymbol); - } + HIP_RETURN_ONFAIL(PlatformState::instance().getStatGlobalVar(texRef, ihipGetDevice(), &refDevPtr, + &refDevSize)); assert(refDevSize == sizeof(textureReference)); // Any previous address or HIP array state associated with the texture reference is superseded by this function. diff --git a/tests/src/runtimeApi/module/hipModule.cpp b/tests/src/runtimeApi/module/hipModule.cpp old mode 100644 new mode 100755 index 30dec3ddd8..9ed5a72415 --- a/tests/src/runtimeApi/module/hipModule.cpp +++ b/tests/src/runtimeApi/module/hipModule.cpp @@ -97,6 +97,7 @@ int main() { assert(A[i] == B[i]); } - HIPCHECK(hipCtxDestroy(context)); + HIPCHECK(hipModuleUnload(Module)); + HIPCHECK(hipCtxDestroy(context)); passed(); } From 01372e2cd42d04028412a21d462ffb9a40366fcc Mon Sep 17 00:00:00 2001 From: Vlad Sytchenko Date: Thu, 11 Jun 2020 12:07:43 -0400 Subject: [PATCH 25/33] Revert "Let hipcc not pass -mllvm option to HIP-Clang on Windows" This reverts commit 92af5e43755887379fd1f5e01dfff9a75d50d1aa. Change-Id: If29906b9c63f03d6e51144f510d3a956d90935e8 --- bin/hipcc | 4 +--- 1 file changed, 1 insertion(+), 3 deletions(-) diff --git a/bin/hipcc b/bin/hipcc index 91d9a197da..f61dead191 100755 --- a/bin/hipcc +++ b/bin/hipcc @@ -816,9 +816,7 @@ if ($HIP_PLATFORM eq "hcc" and $HIP_COMPILER eq "clang") { $HIPCFLAGS .= " -O3"; $HIPLDFLAGS .= " -O3"; } - # Do not pass -mllvm on Windows since there is a clang bug causing duplicate -mllvm options in clang -cc1 on Windows. - # ToDo : remove restriction for Windows after clang bug is fixed. - if (!$funcSupp and $optArg ne "-O0" and not $isWindows and $hasHIP) { + if (!$funcSupp and $optArg ne "-O0" and $hasHIP) { $HIPCXXFLAGS .= " -mllvm -amdgpu-early-inline-all=true -mllvm -amdgpu-function-calls=false"; if ($needLDFLAGS and not $needCXXFLAGS) { $HIPLDFLAGS .= " -mllvm -amdgpu-early-inline-all=true -mllvm -amdgpu-function-calls=false"; From 00301b16651b71cc6901bc973aa71c3ea69e622e Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Wed, 10 Jun 2020 22:37:57 +0000 Subject: [PATCH 26/33] Addback __mbcnt_lo and __mbcnt_hi Change-Id: Ic3facba2e2245461515799f6a17842da0f5d9933 --- include/hip/hcc_detail/device_functions.h | 6 ++++++ 1 file changed, 6 insertions(+) diff --git a/include/hip/hcc_detail/device_functions.h b/include/hip/hcc_detail/device_functions.h index 7bc0b97617..e7b290956a 100644 --- a/include/hip/hcc_detail/device_functions.h +++ b/include/hip/hcc_detail/device_functions.h @@ -238,6 +238,12 @@ __device__ static inline unsigned int __lane_id() { -1, __builtin_amdgcn_mbcnt_lo(-1, 0)); } +__device__ +static inline unsigned int __mbcnt_lo(unsigned int x, unsigned int y) {return __builtin_amdgcn_mbcnt_lo(x,y);}; + +__device__ +static inline unsigned int __mbcnt_hi(unsigned int x, unsigned int y) {return __builtin_amdgcn_mbcnt_hi(x,y);}; + /* HIP specific device functions */ From 0dd612399864967cd50f99fed9d00151aff2ef44 Mon Sep 17 00:00:00 2001 From: kjayapra-amd Date: Thu, 11 Jun 2020 17:23:41 -0400 Subject: [PATCH 27/33] SWDEV-236178 - Remove __hip_pinned_shadow reference from dtests and add hipModuleUnload for all corresponding hipModuleLoad calls. Change-Id: I405789b430ffbb8fccda1cebdb5d18e83a85c926 --- tests/src/runtimeApi/module/hipExtModuleLaunchKernel.cpp | 1 + tests/src/runtimeApi/module/hipModuleGetGlobal.cpp | 1 + .../module/hipModuleOccupancyMaxPotentialActiveBlockSize.cpp | 1 + tests/src/runtimeApi/module/hipModuleTexture2dDrv.cpp | 5 ++--- tests/src/runtimeApi/module/tex2d_kernel.cpp | 5 +---- 5 files changed, 6 insertions(+), 7 deletions(-) mode change 100644 => 100755 tests/src/runtimeApi/module/hipModuleGetGlobal.cpp mode change 100644 => 100755 tests/src/runtimeApi/module/hipModuleOccupancyMaxPotentialActiveBlockSize.cpp mode change 100644 => 100755 tests/src/runtimeApi/module/hipModuleTexture2dDrv.cpp mode change 100644 => 100755 tests/src/runtimeApi/module/tex2d_kernel.cpp diff --git a/tests/src/runtimeApi/module/hipExtModuleLaunchKernel.cpp b/tests/src/runtimeApi/module/hipExtModuleLaunchKernel.cpp index b9327a9c58..804922e28c 100755 --- a/tests/src/runtimeApi/module/hipExtModuleLaunchKernel.cpp +++ b/tests/src/runtimeApi/module/hipExtModuleLaunchKernel.cpp @@ -127,6 +127,7 @@ int main() HIPCHECK(hipFree(Ad)); HIPCHECK(hipFree(Bd)); HIPCHECK(hipHostFree(C)); + HIPCHECK(hipModuleUnload(Module)); if(testStatus) passed(); } diff --git a/tests/src/runtimeApi/module/hipModuleGetGlobal.cpp b/tests/src/runtimeApi/module/hipModuleGetGlobal.cpp old mode 100644 new mode 100755 index 5896794e90..daf693f1bf --- a/tests/src/runtimeApi/module/hipModuleGetGlobal.cpp +++ b/tests/src/runtimeApi/module/hipModuleGetGlobal.cpp @@ -153,6 +153,7 @@ int main() { }; } + HIP_CHECK(hipModuleUnload(Module)); hipCtxDestroy(context); return 0; } diff --git a/tests/src/runtimeApi/module/hipModuleOccupancyMaxPotentialActiveBlockSize.cpp b/tests/src/runtimeApi/module/hipModuleOccupancyMaxPotentialActiveBlockSize.cpp old mode 100644 new mode 100755 index f6935d0d68..60359d5943 --- a/tests/src/runtimeApi/module/hipModuleOccupancyMaxPotentialActiveBlockSize.cpp +++ b/tests/src/runtimeApi/module/hipModuleOccupancyMaxPotentialActiveBlockSize.cpp @@ -49,6 +49,7 @@ int main(int argc, char* argv[]) { assert(gridSize != 0 && blockSize != 0); HIPCHECK(hipModuleOccupancyMaxActiveBlocksPerMultiprocessor(&numBlock, Function, blockSize, 0)); assert(numBlock != 0); + HIPCHECK(hipModuleUnload(Module)); HIPCHECK(hipCtxDestroy(context)); passed(); } diff --git a/tests/src/runtimeApi/module/hipModuleTexture2dDrv.cpp b/tests/src/runtimeApi/module/hipModuleTexture2dDrv.cpp old mode 100644 new mode 100755 index 9318d168a0..ed43366671 --- a/tests/src/runtimeApi/module/hipModuleTexture2dDrv.cpp +++ b/tests/src/runtimeApi/module/hipModuleTexture2dDrv.cpp @@ -21,6 +21,7 @@ THE SOFTWARE. */ /* HIT_START + * BUILD_CMD: tex2d_kernel.code %hc --genco %S/tex2d_kernel.cpp -o tex2d_kernel.code * BUILD: %t %s ../../test_common.cpp EXCLUDE_HIP_PLATFORM nvcc rocclr * TEST: %t * HIT_END @@ -33,9 +34,6 @@ THE SOFTWARE. #define fileName "tex2d_kernel.code" -#if __HIP__ -__hip_pinned_shadow__ -#endif texture tex; bool testResult = false; @@ -133,6 +131,7 @@ bool runTest(int argc, char** argv) { } hipFree(dData); hipFreeArray(array); + HIP_CHECK(hipModuleUnload(Module)); return true; } diff --git a/tests/src/runtimeApi/module/tex2d_kernel.cpp b/tests/src/runtimeApi/module/tex2d_kernel.cpp old mode 100644 new mode 100755 index af5010a221..560f27e741 --- a/tests/src/runtimeApi/module/tex2d_kernel.cpp +++ b/tests/src/runtimeApi/module/tex2d_kernel.cpp @@ -27,10 +27,7 @@ THE SOFTWARE. #include "hip/hip_runtime.h" -#if __HIP__ -__hip_pinned_shadow__ -#endif -extern texture tex; +texture tex; extern "C" __global__ void tex2dKernel(float* outputData, int width, int height) { int x = hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x; From 173bb2af6e31e0106217421f8f2dda11c6e7db7e Mon Sep 17 00:00:00 2001 From: kjayapra-amd Date: Thu, 11 Jun 2020 17:17:29 -0400 Subject: [PATCH 28/33] SWDEV-236178 - Store texture reference metadata for dynamically loaded modules. Change-Id: I99ecc80da7e29c691341a01a09e4532972f1e3e5 --- rocclr/hip_module.cpp | 4 +++- rocclr/hip_platform.cpp | 40 ++++++++++++++++++++++++++++++++++++++++ rocclr/hip_platform.hpp | 4 ++++ rocclr/hip_texture.cpp | 17 +++++++++-------- 4 files changed, 56 insertions(+), 9 deletions(-) diff --git a/rocclr/hip_module.cpp b/rocclr/hip_module.cpp index 383f7a77a9..70dbc02f3a 100755 --- a/rocclr/hip_module.cpp +++ b/rocclr/hip_module.cpp @@ -533,7 +533,7 @@ hipError_t hipModuleGetTexRef(textureReference** texRef, hipModule_t hmod, const } /* Get address and size for the global symbol */ - if (!PlatformState::instance().getDynTexRef(name, hmod, texRef)) { + if (hipSuccess != PlatformState::instance().getDynTexRef(name, hmod, texRef)) { DevLogPrintfError("Cannot get texRef for name: %s at module:0x%x \n", name, hmod); HIP_RETURN(hipErrorNotFound); @@ -543,5 +543,7 @@ hipError_t hipModuleGetTexRef(textureReference** texRef, hipModule_t hmod, const // have the default read mode set to normalized float. (*texRef)->readMode = hipReadModeNormalizedFloat; + PlatformState::instance().registerTexRef(*texRef, hmod, std::string(name)); + HIP_RETURN(hipSuccess); } diff --git a/rocclr/hip_platform.cpp b/rocclr/hip_platform.cpp index 23a1b3afc6..38dac4ad60 100755 --- a/rocclr/hip_platform.cpp +++ b/rocclr/hip_platform.cpp @@ -955,6 +955,15 @@ hipError_t PlatformState::unloadModule(hipModule_t hmod) { delete it->second; dynCO_map_.erase(hmod); + auto tex_it = texRef_map_.begin(); + while (tex_it != texRef_map_.end()) { + if (tex_it->second.first == hmod) { + tex_it = texRef_map_.erase(tex_it); + } else { + ++tex_it; + } + } + return hipSuccess; } @@ -989,6 +998,37 @@ hipError_t PlatformState::getDynGlobalVar(const char* hostVar, int deviceId, hip return hipSuccess; } +hipError_t PlatformState::registerTexRef(textureReference* texRef, hipModule_t hmod, + std::string name) { + amd::ScopedLock lock(lock_); + texRef_map_.insert(std::make_pair(texRef, std::make_pair(hmod, name))); + return hipSuccess; +} + +hipError_t PlatformState::getDynTexGlobalVar(textureReference* texRef, int deviceId, + hipDeviceptr_t* dev_ptr, size_t* size_ptr) { + amd::ScopedLock lock(lock_); + + auto tex_it = texRef_map_.find(texRef); + if (tex_it == texRef_map_.end()) { + DevLogPrintfError("Cannot find the texRef Entry: 0x%x", texRef); + return hipErrorNotFound; + } + + auto it = dynCO_map_.find(tex_it->second.first); + if (it == dynCO_map_.end()) { + DevLogPrintfError("Cannot find the module: 0x%x", tex_it->second.first); + return hipErrorNotFound; + } + + hip::DeviceVar* dvar = nullptr; + IHIP_RETURN_ONFAIL(it->second->getDeviceVar(&dvar, tex_it->second.second, deviceId)); + *dev_ptr = dvar->device_ptr(); + *size_ptr = dvar->size(); + + return hipSuccess; +} + hipError_t PlatformState::getDynTexRef(const char* hostVar, hipModule_t hmod, textureReference** texRef) { amd::ScopedLock lock(lock_); diff --git a/rocclr/hip_platform.hpp b/rocclr/hip_platform.hpp index b53a1a750d..620fb25ca7 100755 --- a/rocclr/hip_platform.hpp +++ b/rocclr/hip_platform.hpp @@ -52,6 +52,9 @@ public: hipDeviceptr_t* dev_ptr, size_t* size_ptr); hipError_t getDynTexRef(const char* hostVar, hipModule_t hmod, textureReference** texRef); + hipError_t registerTexRef(textureReference* texRef, hipModule_t hmod, std::string name); + hipError_t getDynTexGlobalVar(textureReference* texRef, int deviceId, hipDeviceptr_t* dev_ptr, size_t* size_ptr); + /* Singleton instance */ static PlatformState& instance() { if (platform_ == nullptr) { @@ -90,4 +93,5 @@ private: std::unordered_map dynCO_map_; hip::StatCO statCO_; //Static Code object var bool initialized_{false}; + std::unordered_map> texRef_map_; }; diff --git a/rocclr/hip_texture.cpp b/rocclr/hip_texture.cpp index fced181c5b..d56178e4ee 100755 --- a/rocclr/hip_texture.cpp +++ b/rocclr/hip_texture.cpp @@ -802,8 +802,9 @@ hipError_t hipTexRefSetArray(textureReference* texRef, hipDeviceptr_t refDevPtr = nullptr; size_t refDevSize = 0; - HIP_RETURN_ONFAIL(PlatformState::instance().getStatGlobalVar(texRef, ihipGetDevice(), &refDevPtr, - &refDevSize)); + + HIP_RETURN_ONFAIL(PlatformState::instance().getDynTexGlobalVar(texRef, ihipGetDevice(), + &refDevPtr, &refDevSize)); assert(refDevSize == sizeof(textureReference)); // Any previous address or HIP array state associated with the texture reference is superseded by this function. @@ -878,8 +879,8 @@ hipError_t hipTexRefSetAddress(size_t* ByteOffset, hipDeviceptr_t refDevPtr = nullptr; size_t refDevSize = 0; - HIP_RETURN_ONFAIL(PlatformState::instance().getStatGlobalVar(texRef, ihipGetDevice(), &refDevPtr, - &refDevSize)); + HIP_RETURN_ONFAIL(PlatformState::instance().getDynTexGlobalVar(texRef, ihipGetDevice(), + &refDevPtr, &refDevSize)); assert(refDevSize == sizeof(textureReference)); // Any previous address or HIP array state associated with the texture reference is superseded by this function. @@ -923,8 +924,8 @@ hipError_t hipTexRefSetAddress2D(textureReference* texRef, hipDeviceptr_t refDevPtr = nullptr; size_t refDevSize = 0; - HIP_RETURN_ONFAIL(PlatformState::instance().getStatGlobalVar(texRef, ihipGetDevice(), &refDevPtr, - &refDevSize)); + HIP_RETURN_ONFAIL(PlatformState::instance().getDynTexGlobalVar(texRef, ihipGetDevice(), + &refDevPtr, &refDevSize)); assert(refDevSize == sizeof(textureReference)); // Any previous address or HIP array state associated with the texture reference is superseded by this function. @@ -1201,8 +1202,8 @@ hipError_t hipTexRefSetMipmappedArray(textureReference* texRef, hipDeviceptr_t refDevPtr = nullptr; size_t refDevSize = 0; - HIP_RETURN_ONFAIL(PlatformState::instance().getStatGlobalVar(texRef, ihipGetDevice(), &refDevPtr, - &refDevSize)); + HIP_RETURN_ONFAIL(PlatformState::instance().getDynTexGlobalVar(texRef, ihipGetDevice(), + &refDevPtr, &refDevSize)); assert(refDevSize == sizeof(textureReference)); // Any previous address or HIP array state associated with the texture reference is superseded by this function. From 17102ff3a11e36892ef75778886a68a484e2a8de Mon Sep 17 00:00:00 2001 From: Aryan Salmanpour Date: Mon, 15 Jun 2020 12:28:47 -0400 Subject: [PATCH 29/33] expose five missing memcpy hip APIs exposing the following hip APIs which have been previously implemented: hipMemcpy2DFromArray hipMemcpy2DFromArrayAsync hipMemcpyAtoH hipMemcpyHtoA hipMemcpyParam2DAsync Change-Id: I3f8635bd4265d2bf5aa2084f4eabb51f74fb18ba --- rocclr/hip_hcc.def.in | 5 +++++ rocclr/hip_hcc.map.in | 5 +++++ 2 files changed, 10 insertions(+) diff --git a/rocclr/hip_hcc.def.in b/rocclr/hip_hcc.def.in index 3a54928764..415e40cb50 100755 --- a/rocclr/hip_hcc.def.in +++ b/rocclr/hip_hcc.def.in @@ -255,3 +255,8 @@ hipTexObjectGetResourceViewDesc hipTexObjectGetTextureDesc hipExtStreamCreateWithCUMask hipStreamGetPriority +hipMemcpy2DFromArray +hipMemcpy2DFromArrayAsync +hipMemcpyAtoH +hipMemcpyHtoA +hipMemcpyParam2DAsync diff --git a/rocclr/hip_hcc.map.in b/rocclr/hip_hcc.map.in index 5bc2d018f2..179a31c55a 100755 --- a/rocclr/hip_hcc.map.in +++ b/rocclr/hip_hcc.map.in @@ -263,6 +263,11 @@ global: hipGetCmdName*; hipExtStreamCreateWithCUMask; hipStreamGetPriority; + hipMemcpy2DFromArray; + hipMemcpy2DFromArrayAsync; + hipMemcpyAtoH; + hipMemcpyHtoA; + hipMemcpyParam2DAsync; }; local: *; From f4211c3905f2d3e941179e0f7d509d1b2ca0647d Mon Sep 17 00:00:00 2001 From: German Andryeyev Date: Sat, 13 Jun 2020 00:52:36 -0400 Subject: [PATCH 30/33] Initial support for HIP managed memory - Call the new ROCclr interfaces for HMM Change-Id: I2cd1bf438f712a9e9e328340e7d0c025257ca6c1 --- include/hip/hcc_detail/hip_runtime_api.h | 143 ++++++++++++++- rocclr/CMakeLists.txt | 1 + rocclr/hip_hcc.def.in | 5 + rocclr/hip_hcc.map.in | 5 + rocclr/hip_hmm.cpp | 215 +++++++++++++++++++++++ rocclr/hip_memory.cpp | 53 ++---- 6 files changed, 377 insertions(+), 45 deletions(-) create mode 100644 rocclr/hip_hmm.cpp diff --git a/include/hip/hcc_detail/hip_runtime_api.h b/include/hip/hcc_detail/hip_runtime_api.h index 67866a8d4b..a0118ff408 100644 --- a/include/hip/hcc_detail/hip_runtime_api.h +++ b/include/hip/hcc_detail/hip_runtime_api.h @@ -178,8 +178,10 @@ enum hipLimit_t { 0x80000000 ///< Allocate non-coherent memory. Overrides HIP_COHERENT_HOST_ALLOC for specific ///< allocation. -#define hipMemAttachGlobal 0x0 -#define hipMemAttachHost 0x1 +#define hipMemAttachGlobal 0x01 ///< Memory can be accessed by any stream on any device +#define hipMemAttachHost 0x02 ///< Memory cannot be accessed by any stream on any device +#define hipMemAttachSingle 0x04 ///< Memory can only be accessed by a single stream on + ///< the associated device #define hipDeviceMallocDefault 0x0 #define hipDeviceMallocFinegrained 0x1 ///< Memory is allocated in fine grained region of device. @@ -217,6 +219,41 @@ enum hipLimit_t { #define hipCooperativeLaunchMultiDeviceNoPreSync 0x01 #define hipCooperativeLaunchMultiDeviceNoPostSync 0x02 +#define hipCpuDeviceId ((int)-1) +#define hipInvalidDeviceId ((int)-2) + +/* + * @brief HIP Memory Advise values + * @enum + * @ingroup Enumerations + */ +typedef enum hipMemoryAdvise { + hipMemAdviseSetReadMostly = 1, ///< Data will mostly be read and only occassionally + ///< be written to + hipMemAdviseUnsetReadMostly = 2, ///< Undo the effect of hipMemAdviseSetReadMostly + hipMemAdviseSetPreferredLocation = 3, ///< Set the preferred location for the data as + ///< the specified device + hipMemAdviseUnsetPreferredLocation = 4, ///< Clear the preferred location for the data + hipMemAdviseSetAccessedBy = 5, ///< Data will be accessed by the specified device, + ///< so prevent page faults as much as possible + hipMemAdviseUnsetAccessedBy = 6 ///< Let the Unified Memory subsystem decide on + ///< the page faulting policy for the specified device +} hipMemoryAdvise; + +/* + * @brief HIP range attributes + * @enum + * @ingroup Enumerations + */ +typedef enum hipMemRangeAttribute { + hipMemRangeAttributeReadMostly = 1, ///< Whether the range will mostly be read and + ///< only occassionally be written to + hipMemRangeAttributePreferredLocation = 2, ///< The preferred location of the range + hipMemRangeAttributeAccessedBy = 3, ///< Memory range has cudaMemAdviseSetAccessedBy + ///< set for specified device + hipMemRangeAttributeLastPrefetchLocation = 4,///< The last location to which the range was prefetched +} hipMemRangeAttribute; + /* * @brief hipJitOption * @enum @@ -1180,15 +1217,18 @@ hipError_t hipMemAllocHost(void** ptr, size_t size); hipError_t hipHostMalloc(void** ptr, size_t size, unsigned int flags); /** - * @brief Allocates memory that will be automatically managed by the Unified Memory system. + * @brief Allocates memory that will be automatically managed by AMD HMM. * - * @param[out] ptr Pointer to the allocated managed memory - * @param[in] size Requested memory size - * @param[in] flags must be either hipMemAttachGlobal/hipMemAttachHost + * @param [out] dev_ptr - pointer to allocated device memory + * @param [in] size - requested allocation size in bytes + * @param [in] flags - must be either hipMemAttachGlobal or hipMemAttachHost + * (defaults to hipMemAttachGlobal) * - * @return #hipSuccess, #hipErrorOutOfMemory + * @returns #hipSuccess, #hipErrorMemoryAllocation, #hipErrorNotSupported, #hipErrorInvalidValue */ -hipError_t hipMallocManaged(void** devPtr, size_t size, unsigned int flags __dparm(0)); +hipError_t hipMallocManaged(void** dev_ptr, + size_t size, + unsigned int flags __dparm(hipMemAttachGlobal)); /** * @brief Allocate device accessible page locked host memory [Deprecated] @@ -3369,7 +3409,6 @@ hipError_t __hipPopCallConfiguration(dim3 *gridDim, * @returns #hipSuccess, #hipErrorInvalidValue, hipInvalidDevice * */ - hipError_t hipLaunchKernel(const void* function_address, dim3 numBlocks, dim3 dimBlocks, @@ -3377,6 +3416,92 @@ hipError_t hipLaunchKernel(const void* function_address, size_t sharedMemBytes __dparm(0), hipStream_t stream __dparm(0)); +/** + * @brief Prefetches memory to the specified destination device using AMD HMM. + * + * @param [in] dev_ptr pointer to be prefetched + * @param [in] count size in bytes for prefetching + * @param [in] device destination device to prefetch to + * @param [in] stream stream to enqueue prefetch operation + * + * @returns #hipSuccess, #hipErrorInvalidValue + */ +hipError_t hipMemPrefetchAsync(const void* dev_ptr, + size_t count, + int device, + hipStream_t stream __dparm(0)); + +/** + * @brief Advise about the usage of a given memory range to AMD HMM. + * + * @param [in] dev_ptr pointer to memory to set the advice for + * @param [in] count size in bytes of the memory range + * @param [in] advice advice to be applied for the specified memory range + * @param [in] device device to apply the advice for + * + * @returns #hipSuccess, #hipErrorInvalidValue + */ +hipError_t hipMemAdvise(const void* dev_ptr, + size_t count, + hipMemoryAdvise advice, + int device); + +/** + * @brief Query an attribute of a given memory range in AMD HMM. + * + * @param [in/out] data a pointer to a memory location where the result of each + * attribute query will be written to + * @param [in] data_size the size of data + * @param [in] attribute the attribute to query + * @param [in] dev_ptr start of the range to query + * @param [in] count size of the range to query + * + * @returns #hipSuccess, #hipErrorInvalidValue + */ +hipError_t hipMemRangeGetAttribute(void* data, + size_t data_size, + hipMemRangeAttribute attribute, + const void* dev_ptr, + size_t count); + +/** + * @brief Query attributes of a given memory range in AMD HMM. + * + * @param [in/out] data a two-dimensional array containing pointers to memory locations + * where the result of each attribute query will be written to + * @param [in] data_sizes an array, containing the sizes of each result + * @param [in] attributes the attribute to query + * @param [in] num_attributes an array of attributes to query (numAttributes and the number + * of attributes in this array should match) + * @param [in] dev_ptr start of the range to query + * @param [in] count size of the range to query + * + * @returns #hipSuccess, #hipErrorInvalidValue + */ +hipError_t hipMemRangeGetAttributes(void** data, + size_t* data_sizes, + hipMemRangeAttribute* attributes, + size_t num_attributes, + const void* dev_ptr, + size_t count); + +/** + * @brief Attach memory to a stream asynchronously in AMD HMM. + * + * @param [in] stream - stream in which to enqueue the attach operation + * @param [in] dev_ptr - pointer to memory (must be a pointer to managed memory or + * to a valid host-accessible region of system-allocated memory) + * @param [in] length - length of memory (defaults to zero) + * @param [in] flags - must be one of cudaMemAttachGlobal, cudaMemAttachHost or + * cudaMemAttachSingle (defaults to cudaMemAttachSingle) + * + * @returns #hipSuccess, #hipErrorInvalidValue + */ +hipError_t hipStreamAttachMemAsync(hipStream_t stream, + hipDeviceptr_t* dev_ptr, + size_t length __dparm(0), + unsigned int flags __dparm(hipMemAttachSingle)); + #if __HIP_ROCclr__ || !defined(__HCC__) //TODO: Move this to hip_ext.h hipError_t hipExtLaunchKernel(const void* function_address, dim3 numBlocks, dim3 dimBlocks, diff --git a/rocclr/CMakeLists.txt b/rocclr/CMakeLists.txt index 252ac5abfa..21ce5b1c58 100755 --- a/rocclr/CMakeLists.txt +++ b/rocclr/CMakeLists.txt @@ -81,6 +81,7 @@ add_library(hip64 OBJECT hip_error.cpp hip_event.cpp hip_global.cpp + hip_hmm.cpp hip_memory.cpp hip_module.cpp hip_peer.cpp diff --git a/rocclr/hip_hcc.def.in b/rocclr/hip_hcc.def.in index 415e40cb50..da43f5f51d 100755 --- a/rocclr/hip_hcc.def.in +++ b/rocclr/hip_hcc.def.in @@ -80,6 +80,7 @@ hipMallocManaged hipArrayCreate hipArray3DCreate hipMallocArray +hipMemAdvise hipMemAllocPitch hipMallocPitch hipMemcpy @@ -111,7 +112,10 @@ hipMemGetAddressRange hipGetSymbolAddress hipGetSymbolSize hipMemGetInfo +hipMemPrefetchAsync hipMemPtrGetInfo +hipMemRangeGetAttribute +hipMemRangeGetAttributes hipMemset hipMemsetAsync hipMemsetD8 @@ -154,6 +158,7 @@ hipGetDeviceFlags hipSetDevice hipSetDeviceFlags hipStreamAddCallback +hipStreamAttachMemAsync hipStreamCreate hipStreamCreateWithFlags hipStreamCreateWithPriority diff --git a/rocclr/hip_hcc.map.in b/rocclr/hip_hcc.map.in index 179a31c55a..e66d4be92d 100755 --- a/rocclr/hip_hcc.map.in +++ b/rocclr/hip_hcc.map.in @@ -62,6 +62,7 @@ global: hipGetErrorName; hipGetErrorString; hipGetLastError; + hipMemAdvise; hipMemAllocHost; hipHostAlloc; hipHostFree; @@ -112,7 +113,10 @@ global: hipGetSymbolAddress; hipGetSymbolSize; hipMemGetInfo; + hipMemPrefetchAsync; hipMemPtrGetInfo; + hipMemRangeGetAttribute; + hipMemRangeGetAttributes; hipMemset; hipMemsetAsync; hipMemsetD8; @@ -154,6 +158,7 @@ global: hipSetDevice; hipSetDeviceFlags; hipStreamAddCallback; + hipStreamAttachMemAsync; hipStreamCreate; hipStreamCreateWithFlags; hipStreamCreateWithPriority; diff --git a/rocclr/hip_hmm.cpp b/rocclr/hip_hmm.cpp new file mode 100644 index 0000000000..5342c18349 --- /dev/null +++ b/rocclr/hip_hmm.cpp @@ -0,0 +1,215 @@ +/* Copyright (c) 2020-present Advanced Micro Devices, Inc. + + 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 +#include "hip_internal.hpp" +#include "hip_conversions.hpp" +#include "platform/context.hpp" +#include "platform/command.hpp" +#include "platform/memory.hpp" + +// Forward declaraiton of a static function +static hipError_t ihipMallocManaged(void** ptr, size_t size); + +// Make sure HIP defines match ROCclr to avoid double conversion +static_assert(hipCpuDeviceId == amd::CpuDeviceId, "CPU device ID mismatch with ROCclr!"); +static_assert(hipInvalidDeviceId == amd::InvalidDeviceId, + "Invalid device ID mismatch with ROCclr!"); + +static_assert(static_cast(hipMemAdviseSetReadMostly) == + amd::MemoryAdvice::SetReadMostly, "Enum mismatch with ROCclr!"); +static_assert(static_cast(hipMemAdviseUnsetReadMostly) == + amd::MemoryAdvice::UnsetReadMostly, "Enum mismatch with ROCclr!"); +static_assert(static_cast(hipMemAdviseSetPreferredLocation) == + amd::MemoryAdvice::SetPreferredLocation, "Enum mismatch with ROCclr!"); +static_assert(static_cast(hipMemAdviseUnsetPreferredLocation) == + amd::MemoryAdvice::UnsetPreferredLocation, "Enum mismatch with ROCclr!"); +static_assert(static_cast(hipMemAdviseSetAccessedBy) == + amd::MemoryAdvice::SetAccessedBy, "Enum mismatch with ROCclr!"); +static_assert(static_cast(hipMemAdviseUnsetAccessedBy) == + amd::MemoryAdvice::UnsetAccessedBy, "Enum mismatch with ROCclr!"); + +static_assert(static_cast(hipMemRangeAttributeReadMostly) == + amd::MemRangeAttribute::ReadMostly, "Enum mismatch with ROCclr!"); +static_assert(static_cast(hipMemRangeAttributePreferredLocation) == + amd::MemRangeAttribute::PreferredLocation, "Enum mismatch with ROCclr!"); +static_assert(static_cast(hipMemRangeAttributeAccessedBy) == + amd::MemRangeAttribute::AccessedBy, "Enum mismatch with ROCclr!"); +static_assert(static_cast(hipMemRangeAttributeLastPrefetchLocation) == + amd::MemRangeAttribute::LastPrefetchLocation, "Enum mismatch with ROCclr!"); + +// ================================================================================================ +hipError_t hipMallocManaged(void** dev_ptr, size_t size, unsigned int flags) { + HIP_INIT_API(hipMallocManaged, dev_ptr, size, flags); + + if ((dev_ptr == nullptr) || (flags != hipMemAttachGlobal)) { + HIP_RETURN(hipErrorInvalidValue); + } + + HIP_RETURN(ihipMallocManaged(dev_ptr, size), *dev_ptr); +} + +// ================================================================================================ +hipError_t hipMemPrefetchAsync(const void* dev_ptr, size_t count, int device, + hipStream_t stream) { + HIP_INIT_API(hipMemPrefetchAsync, dev_ptr, count, device, stream); + + if ((dev_ptr == nullptr) || (count == 0) || (stream == nullptr)) { + HIP_RETURN(hipErrorInvalidValue); + } + amd::HostQueue* queue = nullptr; + bool cpu_access = (device == hipCpuDeviceId) ? true : false; + + // Pick the specified stream or Null one from the provided device + if (stream != nullptr) { + queue = hip::getQueue(stream); + } else { + if (!cpu_access) { + queue = g_devices[device]->NullStream(); + } else { + queue = hip::getCurrentDevice()->NullStream(); + } + } + if (queue == nullptr) { + HIP_RETURN(hipErrorInvalidValue); + } + + amd::Command::EventWaitList waitList; + amd::SvmPrefetchAsyncCommand* command = + new amd::SvmPrefetchAsyncCommand(*queue, waitList, dev_ptr, count, cpu_access); + if (command == nullptr) { + return hipErrorOutOfMemory; + } + + if (!command->validateMemory()) { + delete command; + HIP_RETURN(hipErrorInvalidValue); + } + command->enqueue(); + command->release(); + + HIP_RETURN(hipErrorInvalidValue); +} + +// ================================================================================================ +hipError_t hipMemAdvise(const void* dev_ptr, size_t count, hipMemoryAdvise advice, int device) { + HIP_INIT_API(hipMemAdvise, dev_ptr, count, advice, device); + + if ((dev_ptr == nullptr) || (count == 0) || (device >= g_devices.size())) { + HIP_RETURN(hipErrorInvalidValue); + } + amd::Device* dev = g_devices[device]->devices()[0]; + + // Set the allocation attributes in AMD HMM + if (!dev->SetSvmAttributes(dev_ptr, count, static_cast(advice))) { + HIP_RETURN(hipErrorInvalidValue); + } + + HIP_RETURN(hipSuccess); +} + +// ================================================================================================ +hipError_t hipMemRangeGetAttribute(void* data, size_t data_size, hipMemRangeAttribute attribute, + const void* dev_ptr, size_t count) { + HIP_INIT_API(hipMemRangeGetAttribute, data, data_size, attribute, dev_ptr, count); + + if ((data == nullptr) || (data_size == 0) || (dev_ptr == nullptr) || (count == 0)) { + HIP_RETURN(hipErrorInvalidValue); + } + + // Shouldn't matter for which device the interface is called + amd::Device* dev = g_devices[0]->devices()[0]; + + // Get the allocation attribute from AMD HMM + if (!dev->GetSvmAttributes(&data, &data_size, reinterpret_cast(&attribute), 1, + dev_ptr, count)) { + HIP_RETURN(hipErrorInvalidValue); + } + + HIP_RETURN(hipErrorInvalidValue); +} + +// ================================================================================================ +hipError_t hipMemRangeGetAttributes(void** data, size_t* data_sizes, + hipMemRangeAttribute* attributes, size_t num_attributes, + const void* dev_ptr, size_t count) { + HIP_INIT_API(hipMemRangeGetAttributes, data, data_sizes, + attributes, num_attributes, dev_ptr, count); + + if ((data == nullptr) || (data_sizes == nullptr) || (attributes == nullptr) || + (num_attributes == 0) || (dev_ptr == nullptr) || (count == 0)) { + HIP_RETURN(hipErrorInvalidValue); + } + + // Shouldn't matter for which device the interface is called + amd::Device* dev = g_devices[0]->devices()[0]; + // Get the allocation attributes from AMD HMM + if (!dev->GetSvmAttributes(data, data_sizes, reinterpret_cast(attributes), + num_attributes, dev_ptr, count)) { + HIP_RETURN(hipErrorInvalidValue); + } + + HIP_RETURN(hipErrorInvalidValue); +} + +// ================================================================================================ +hipError_t hipStreamAttachMemAsync(hipStream_t stream, hipDeviceptr_t* dev_ptr, + size_t length, unsigned int flags) { + HIP_INIT_API(hipStreamAttachMemAsync, stream, dev_ptr, length, flags); + + if ((stream == nullptr) || (dev_ptr == nullptr) || (length == 0)) { + HIP_RETURN(hipErrorInvalidValue); + } + + // Unclear what should be done for this interface in AMD HMM, since it's generic SVM alloc + HIP_RETURN(hipErrorInvalidValue); +} + +// ================================================================================================ +static hipError_t ihipMallocManaged(void** ptr, size_t size) { + if (size == 0) { + *ptr = nullptr; + return hipSuccess; + } else if (ptr == nullptr) { + return hipErrorInvalidValue; + } + + assert((hip::host_device->asContext()!= nullptr) && "Current host context must be valid"); + amd::Context& ctx = *hip::host_device->asContext(); + + const amd::Device& dev = *ctx.devices()[0]; + // For now limit to the max allocation size on the device. + // The apps should be able to go over theCould you limit allocation in the future + if (dev.info().maxMemAllocSize_ < size) { + return hipErrorMemoryAllocation; + } + + // Allocate SVM fine grain buffer with the forced host pointer, avoiding explicit memory + // allocation in the device driver + *ptr = amd::SvmBuffer::malloc(ctx, CL_MEM_SVM_FINE_GRAIN_BUFFER | CL_MEM_ALLOC_HOST_PTR, + size, dev.info().memBaseAddrAlign_); + if (*ptr == nullptr) { + return hipErrorMemoryAllocation; + } + + ClPrint(amd::LOG_INFO, amd::LOG_API, "%-5d: [%zx] ihipMallocManaged ptr=0x%zx", getpid(), + std::this_thread::get_id(), *ptr); + return hipSuccess; +} \ No newline at end of file diff --git a/rocclr/hip_memory.cpp b/rocclr/hip_memory.cpp index 2854b0d7eb..796cc0e491 100755 --- a/rocclr/hip_memory.cpp +++ b/rocclr/hip_memory.cpp @@ -105,6 +105,7 @@ hipError_t ihipMalloc(void** ptr, size_t sizeBytes, unsigned int flags) return hipSuccess; } +// ================================================================================================ hipError_t ihipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind kind, amd::HostQueue& queue, bool isAsync = false) { if (sizeBytes == 0) { @@ -154,36 +155,26 @@ hipError_t ihipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKin *srcMemory->asBuffer(), sOffset, sizeBytes, dst); isAsync = false; } else if ((srcMemory != nullptr) && (dstMemory != nullptr)) { - if (queueDevice != srcMemory->getContext().devices()[0]) { - amd::Coord3D srcOffset(sOffset, 0, 0); - amd::Coord3D dstOffset(dOffset, 0, 0); - amd::Coord3D copySize(sizeBytes, 1, 1); + if ((kind == hipMemcpyDeviceToDevice) && + // Check if the queue device doesn't match the device on any memory object. Hence + // it's a P2P transfer, because the app has requested access to another GPU + (srcMemory->getContext().devices()[0] != dstMemory->getContext().devices()[0])) { command = new amd::CopyMemoryP2PCommand(queue, CL_COMMAND_COPY_BUFFER, waitList, - *srcMemory->asBuffer(),*dstMemory->asBuffer(), srcOffset, dstOffset, copySize); - command->enqueue(); - if (!isAsync) { - command->awaitCompletion(); + *srcMemory->asBuffer(), *dstMemory->asBuffer(), sOffset, dOffset, sizeBytes); + if (command == nullptr) { + return hipErrorOutOfMemory; } - command->release(); - return hipSuccess; - } - if (queueDevice != dstMemory->getContext().devices()[0]) { - amd::Coord3D srcOffset(sOffset, 0, 0); - amd::Coord3D dstOffset(dOffset, 0, 0); - amd::Coord3D copySize(sizeBytes, 1, 1); - command = new amd::CopyMemoryP2PCommand(queue, CL_COMMAND_COPY_BUFFER, waitList, - *srcMemory->asBuffer(),*dstMemory->asBuffer(), srcOffset, dstOffset, copySize); - command->enqueue(); - if (!isAsync) { - command->awaitCompletion(); + // Make sure runtime has valid memory for the command execution. P2P access + // requires page table mapping on the current device to another GPU memory + if (!static_cast(command)->validateMemory()) { + delete command; + return hipErrorInvalidValue; } - command->release(); - return hipSuccess; + } else { + command = new amd::CopyMemoryCommand(queue, CL_COMMAND_COPY_BUFFER, waitList, + *srcMemory->asBuffer(), *dstMemory->asBuffer(), sOffset, dOffset, sizeBytes); } - command = new amd::CopyMemoryCommand(queue, CL_COMMAND_COPY_BUFFER, waitList, - *srcMemory->asBuffer(),*dstMemory->asBuffer(), sOffset, dOffset, sizeBytes); } - if (command == nullptr) { return hipErrorOutOfMemory; } @@ -201,6 +192,7 @@ hipError_t ihipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKin return hipSuccess; } +// ================================================================================================ hipError_t hipExtMallocWithFlags(void** ptr, size_t sizeBytes, unsigned int flags) { HIP_INIT_API(hipExtMallocWithFlags, ptr, sizeBytes, flags); @@ -245,17 +237,6 @@ hipError_t hipHostMalloc(void** ptr, size_t sizeBytes, unsigned int flags) { HIP_RETURN(ihipMalloc(ptr, sizeBytes, ihipFlags), *ptr); } -hipError_t hipMallocManaged(void** devPtr, size_t size, - unsigned int flags) { - HIP_INIT_API(hipMallocManaged, devPtr, size, flags); - - if (flags != hipMemAttachGlobal) { - HIP_RETURN(hipErrorInvalidValue); - } - - HIP_RETURN(ihipMalloc(devPtr, size, CL_MEM_SVM_FINE_GRAIN_BUFFER), *devPtr); -} - hipError_t hipFree(void* ptr) { HIP_INIT_API(hipFree, ptr); From f3480e019eed15c5cf5d805a3ea10d22e3b4eb0e Mon Sep 17 00:00:00 2001 From: Vlad Sytchenko Date: Mon, 15 Jun 2020 15:20:30 -0400 Subject: [PATCH 31/33] Enable the use of some warnings when building HIP-ROCclr Similar to http://gerrit-git.amd.com/c/compute/ec/vdi/+/375385, except no -Wno-strict-aliasing. With this change only two warnings show up during the build - -Wsign-compare and -Wmisleading-indentation. Change-Id: Iffa436c65ab6312aeaa6def71ee6af38ed1b4a4e --- rocclr/CMakeLists.txt | 13 +++++++++++-- 1 file changed, 11 insertions(+), 2 deletions(-) diff --git a/rocclr/CMakeLists.txt b/rocclr/CMakeLists.txt index 21ce5b1c58..8f1d7ef5c0 100755 --- a/rocclr/CMakeLists.txt +++ b/rocclr/CMakeLists.txt @@ -21,8 +21,17 @@ if(CMAKE_BUILD_TYPE MATCHES "^Debug$") add_definitions(-DDEBUG) endif() -if (CMAKE_CXX_COMPILER_ID STREQUAL "GNU") - add_compile_options("-Wno-ignored-attributes") +if((CMAKE_CXX_COMPILER_ID STREQUAL "GNU") OR + (CMAKE_${COMPILER}_COMPILER_ID MATCHES "Clang")) + add_definitions( + # Enabling -Wextra or -pedantic will cause + # thousands of warnings. Keep things simple for now. + -Wall + # This one seems impossible to fix for now. + # There are hundreds of instances of unused vars/functions + # throughout the code base. + -Wno-unused-variable + -Wno-unused-function) endif() set(USE_PROF_API "1") From 45e0847319ab76f216d3bfbf06320e7120f459e1 Mon Sep 17 00:00:00 2001 From: kjayapra-amd Date: Fri, 12 Jun 2020 18:39:48 -0400 Subject: [PATCH 32/33] SWDEV-240589 - Remove guarantee @ __hipRegisterFunction flow for now. Change-Id: Ic51e2a1f951ac7745c2bbd11cfd2b92521c6966a --- rocclr/hip_code_object.cpp | 1 - 1 file changed, 1 deletion(-) diff --git a/rocclr/hip_code_object.cpp b/rocclr/hip_code_object.cpp index bb957b8cf3..f377fc594f 100755 --- a/rocclr/hip_code_object.cpp +++ b/rocclr/hip_code_object.cpp @@ -383,7 +383,6 @@ hipError_t StatCO::registerStatFunction(const void* hostFunction, Function* func if (functions_.find(hostFunction) != functions_.end()) { DevLogPrintfError("hostFunctionPtr: 0x%x already exists", hostFunction); - guarantee(false); } functions_.insert(std::make_pair(hostFunction, func)); From 7682ec6a7b7ccc05f561250c9ecb5317fd53eafe Mon Sep 17 00:00:00 2001 From: Tao Sang Date: Fri, 5 Jun 2020 15:29:08 -0400 Subject: [PATCH 33/33] Make hipHostMalloc() respect hipSetDevice() Change-Id: I2410240f91b108c24597ee0fa7cf31e7b1f3ac5d --- rocclr/hip_memory.cpp | 8 +++++--- 1 file changed, 5 insertions(+), 3 deletions(-) diff --git a/rocclr/hip_memory.cpp b/rocclr/hip_memory.cpp index 796cc0e491..a3771a5e20 100755 --- a/rocclr/hip_memory.cpp +++ b/rocclr/hip_memory.cpp @@ -86,8 +86,9 @@ hipError_t ihipMalloc(void** ptr, size_t sizeBytes, unsigned int flags) return hipErrorInvalidValue; } - amd::Context* amdContext = ((flags & CL_MEM_SVM_FINE_GRAIN_BUFFER) != 0)? - hip::host_device->asContext() : hip::getCurrentDevice()->asContext(); + bool useHostDevice = (flags & CL_MEM_SVM_FINE_GRAIN_BUFFER) != 0; + amd::Context* curDevContext = hip::getCurrentDevice()->asContext(); + amd::Context* amdContext = useHostDevice ? hip::host_device->asContext() : curDevContext; if (amdContext == nullptr) { return hipErrorOutOfMemory; @@ -97,7 +98,8 @@ hipError_t ihipMalloc(void** ptr, size_t sizeBytes, unsigned int flags) return hipErrorOutOfMemory; } - *ptr = amd::SvmBuffer::malloc(*amdContext, flags, sizeBytes, amdContext->devices()[0]->info().memBaseAddrAlign_); + *ptr = amd::SvmBuffer::malloc(*amdContext, flags, sizeBytes, amdContext->devices()[0]->info().memBaseAddrAlign_, + useHostDevice ? curDevContext->svmDevices()[0] : nullptr); if (*ptr == nullptr) { return hipErrorOutOfMemory; }