From 97325ff30338ca51caacac8c6f8eac4509fda021 Mon Sep 17 00:00:00 2001 From: Satyanvesh Dittakavi Date: Fri, 11 Feb 2022 13:45:53 +0000 Subject: [PATCH] SWDEV-314485 - Remove the older implementation of device-side malloc/free Change-Id: I9ba3c6030a6dd7a5e064073d5c21223041282f61 [ROCm/clr commit: e5496b31591729eb1a43a7470aa2339d5106ec89] --- .../include/hip/amd_detail/amd_hip_runtime.h | 21 ---- .../include/hip/amd_detail/hip_memory.h | 114 ------------------ 2 files changed, 135 deletions(-) delete mode 100644 projects/clr/hipamd/include/hip/amd_detail/hip_memory.h diff --git a/projects/clr/hipamd/include/hip/amd_detail/amd_hip_runtime.h b/projects/clr/hipamd/include/hip/amd_detail/amd_hip_runtime.h index 41cdf24ea0..6bb4658c69 100644 --- a/projects/clr/hipamd/include/hip/amd_detail/amd_hip_runtime.h +++ b/projects/clr/hipamd/include/hip/amd_detail/amd_hip_runtime.h @@ -93,11 +93,6 @@ using ::int64_t; } #endif // !defined(__HIPCC_RTC__) -// __hip_malloc is not working. Disable it by default. -#ifndef __HIP_ENABLE_DEVICE_MALLOC__ -#define __HIP_ENABLE_DEVICE_MALLOC__ 0 -#endif - #if __HIP_CLANG_ONLY__ #if !defined(__align__) @@ -192,20 +187,6 @@ __device__ int __hip_move_dpp_N(int src); #endif //__HIP_ARCH_GFX803__ == 1 -#ifndef __OPENMP_AMDGCN__ -#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); -static inline __device__ void* malloc(size_t size) { return __hip_malloc(size); } -static inline __device__ void* free(void* ptr) { return __hip_free(ptr); } -#else -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 // !__OPENMP_AMDGCN__ - // End doxygen API: /** * @} @@ -459,6 +440,4 @@ hc_get_workitem_absolute_id(int dim) #endif // !__CLANG_HIP_RUNTIME_WRAPPER_INCLUDED__ #endif // __HIP_CLANG_ONLY__ -#include - #endif // HIP_AMD_DETAIL_RUNTIME_H diff --git a/projects/clr/hipamd/include/hip/amd_detail/hip_memory.h b/projects/clr/hipamd/include/hip/amd_detail/hip_memory.h deleted file mode 100644 index 585dd2a14a..0000000000 --- a/projects/clr/hipamd/include/hip/amd_detail/hip_memory.h +++ /dev/null @@ -1,114 +0,0 @@ -/* -Copyright (c) 2015 - 2021 Advanced Micro Devices, Inc. All rights reserved. - -Permission is hereby granted, free of charge, to any person obtaining a copy -of this software and associated documentation files (the "Software"), to deal -in the Software without restriction, including without limitation the rights -to use, copy, modify, merge, publish, distribute, sublicense, and/or sell -copies of the Software, and to permit persons to whom the Software is -furnished to do so, subject to the following conditions: - -The above copyright notice and this permission notice shall be included in -all copies or substantial portions of the Software. - -THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR -IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, -FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE -AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER -LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, -OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN -THE SOFTWARE. -*/ - -#ifndef HIP_INCLUDE_HIP_AMD_DETAIL_HIP_MEMORY_H -#define HIP_INCLUDE_HIP_AMD_DETAIL_HIP_MEMORY_H - -// Implementation of malloc and free device functions. -// HIP heap is implemented as a global array with fixed size. Users may define -// __HIP_SIZE_OF_PAGE and __HIP_NUM_PAGES to have a larger heap. - -#if __HIP__ && __HIP_ENABLE_DEVICE_MALLOC__ - -// Size of page in bytes. -#ifndef __HIP_SIZE_OF_PAGE -#define __HIP_SIZE_OF_PAGE 64 -#endif - -// Total number of pages -#ifndef __HIP_NUM_PAGES -#define __HIP_NUM_PAGES (16 * 64 * 64) -#endif - -#define __HIP_SIZE_OF_HEAP (__HIP_NUM_PAGES * __HIP_SIZE_OF_PAGE) - -#if __HIP_DEVICE_COMPILE__ -__attribute__((weak)) __device__ char __hip_device_heap[__HIP_SIZE_OF_HEAP]; -__attribute__((weak)) __device__ - uint32_t __hip_device_page_flag[__HIP_NUM_PAGES]; -#else -extern __device__ char __hip_device_heap[]; -extern __device__ uint32_t __hip_device_page_flag[]; -#endif - -extern "C" inline __device__ void* __hip_malloc(size_t size) { - char* heap = (char*)__hip_device_heap; - if (size > __HIP_SIZE_OF_HEAP) { - return (void*)nullptr; - } - uint32_t totalThreads = - hipBlockDim_x * hipGridDim_x * hipBlockDim_y - * hipGridDim_y * hipBlockDim_z * hipGridDim_z; - uint32_t currentWorkItem = hipThreadIdx_x + hipBlockDim_x * hipBlockIdx_x - + (hipThreadIdx_y + hipBlockDim_y * hipBlockIdx_y) * hipBlockDim_x - + (hipThreadIdx_z + hipBlockDim_z * hipBlockIdx_z) * hipBlockDim_x - * hipBlockDim_y; - - uint32_t numHeapsPerWorkItem = __HIP_NUM_PAGES / totalThreads; - uint32_t heapSizePerWorkItem = __HIP_SIZE_OF_HEAP / totalThreads; - - uint32_t stride = size / __HIP_SIZE_OF_PAGE; - uint32_t start = numHeapsPerWorkItem * currentWorkItem; - - uint32_t k = 0; - - while (__hip_device_page_flag[k] > 0) { - k++; - } - - for (uint32_t i = 0; i < stride - 1; i++) { - __hip_device_page_flag[i + start + k] = 1; - } - - __hip_device_page_flag[start + stride - 1 + k] = 2; - - void* ptr = (void*)(heap - + heapSizePerWorkItem * currentWorkItem + k * __HIP_SIZE_OF_PAGE); - - return ptr; -} - -extern "C" inline __device__ void* __hip_free(void* ptr) { - if (ptr == nullptr) { - return nullptr; - } - - uint32_t offsetByte = (uint64_t)ptr - (uint64_t)__hip_device_heap; - uint32_t offsetPage = offsetByte / __HIP_SIZE_OF_PAGE; - - while (__hip_device_page_flag[offsetPage] != 0) { - if (__hip_device_page_flag[offsetPage] == 2) { - __hip_device_page_flag[offsetPage] = 0; - offsetPage++; - break; - } else { - __hip_device_page_flag[offsetPage] = 0; - offsetPage++; - } - } - - return nullptr; -} - -#endif - -#endif // HIP_INCLUDE_HIP_AMD_DETAIL_HIP_MEMORY_H