From d69edbbb7f51a009f27a4bdfe48a3477d26f7a9c Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Wed, 24 Apr 2019 04:20:03 -0700 Subject: [PATCH] Add hipMallocManaged default functional support (#1036) * Add hipMallocManaged default functional support * Fix build error * Add dtest [ROCm/clr commit: 94769fc8dd42b2fd3cca4257bcd9410f67db908a] --- .../include/hip/hcc_detail/hip_runtime_api.h | 14 +++++ .../clr/hipamd/include/hip/hip_runtime_api.h | 6 +++ .../include/hip/nvcc_detail/hip_runtime_api.h | 7 +++ projects/clr/hipamd/src/hip_memory.cpp | 23 ++++++-- .../runtimeApi/memory/hipMallocManaged.cpp | 52 +++++++++++++++++++ 5 files changed, 99 insertions(+), 3 deletions(-) create mode 100644 projects/clr/hipamd/tests/src/runtimeApi/memory/hipMallocManaged.cpp diff --git a/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime_api.h b/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime_api.h index ccfe7aefc3..863e6c0e4a 100644 --- a/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime_api.h +++ b/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime_api.h @@ -177,6 +177,9 @@ enum hipLimit_t { 0x80000000 ///< Allocate non-coherent memory. Overrides HIP_COHERENT_HOST_ALLOC for specific ///< allocation. +#define hipMemAttachGlobal 0x0 +#define hipMemAttachHost 0x1 + #define hipDeviceMallocDefault 0x0 #define hipDeviceMallocFinegrained 0x1 ///< Memory is allocated in fine grained region of device. @@ -1103,6 +1106,17 @@ hipError_t hipMallocHost(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. + * + * @param[out] ptr Pointer to the allocated managed memory + * @param[in] size Requested memory size + * @param[in] flags must be either hipMemAttachGlobal/hipMemAttachHost + * + * @return #hipSuccess, #hipErrorMemoryAllocation + */ +hipError_t hipMallocManaged(void** devPtr, size_t size, unsigned int flags __dparm(0)); + /** * @brief Allocate device accessible page locked host memory [Deprecated] * diff --git a/projects/clr/hipamd/include/hip/hip_runtime_api.h b/projects/clr/hipamd/include/hip/hip_runtime_api.h index 34363689e4..e7ecede8c1 100644 --- a/projects/clr/hipamd/include/hip/hip_runtime_api.h +++ b/projects/clr/hipamd/include/hip/hip_runtime_api.h @@ -337,6 +337,12 @@ static inline hipError_t hipHostMalloc(T** ptr, size_t size, unsigned int flags = hipHostMallocDefault) { return hipHostMalloc((void**)ptr, size, flags); } + +template +static inline hipError_t hipMallocManaged(T** devPtr, size_t size, + unsigned int flags = hipMemAttachGlobal) { + return hipMallocManaged((void**)devPtr, size, flags); +} #endif #endif diff --git a/projects/clr/hipamd/include/hip/nvcc_detail/hip_runtime_api.h b/projects/clr/hipamd/include/hip/nvcc_detail/hip_runtime_api.h index 4d84cfbc51..73ab91f25c 100644 --- a/projects/clr/hipamd/include/hip/nvcc_detail/hip_runtime_api.h +++ b/projects/clr/hipamd/include/hip/nvcc_detail/hip_runtime_api.h @@ -111,6 +111,9 @@ typedef enum hipChannelFormatKind { #define hipHostMallocCoherent 0x0 #define hipHostMallocNonCoherent 0x0 +#define hipMemAttachGlobal cudaMemAttachGlobal +#define hipMemAttachHost cudaMemAttachHost + #define hipHostRegisterDefault cudaHostRegisterDefault #define hipHostRegisterPortable cudaHostRegisterPortable #define hipHostRegisterMapped cudaHostRegisterMapped @@ -420,6 +423,10 @@ inline static hipError_t hipHostMalloc(void** ptr, size_t size, unsigned int fla return hipCUDAErrorTohipError(cudaHostAlloc(ptr, size, flags)); } +inline static hipError_t hipMallocManaged(void** ptr, size_t size, unsigned int flags) { + return hipCUDAErrorTohipError(cudaMallocManaged(ptr, size, flags)); +} + inline static hipError_t hipMallocArray(hipArray** array, const struct hipChannelFormatDesc* desc, size_t width, size_t height, unsigned int flags __dparm(hipArrayDefault)) { diff --git a/projects/clr/hipamd/src/hip_memory.cpp b/projects/clr/hipamd/src/hip_memory.cpp index b6a1164a4a..b8632970d4 100644 --- a/projects/clr/hipamd/src/hip_memory.cpp +++ b/projects/clr/hipamd/src/hip_memory.cpp @@ -301,9 +301,7 @@ hipError_t hipExtMallocWithFlags(void** ptr, size_t sizeBytes, unsigned int flag return ihipLogStatus(hip_status); } -hipError_t hipHostMalloc(void** ptr, size_t sizeBytes, unsigned int flags) { - HIP_INIT_SPECIAL_API(hipHostMalloc, (TRACE_MEM), ptr, sizeBytes, flags); - HIP_SET_DEVICE(); +hipError_t ihipHostMalloc(void** ptr, size_t sizeBytes, unsigned int flags) { hipError_t hip_status = hipSuccess; if (HIP_SYNC_HOST_ALLOC) { @@ -368,6 +366,25 @@ hipError_t hipHostMalloc(void** ptr, size_t sizeBytes, unsigned int flags) { if (HIP_SYNC_HOST_ALLOC) { hipDeviceSynchronize(); } + return hip_status; +} + +hipError_t hipHostMalloc(void** ptr, size_t sizeBytes, unsigned int flags) { + HIP_INIT_SPECIAL_API(hipHostMalloc, (TRACE_MEM), ptr, sizeBytes, flags); + HIP_SET_DEVICE(); + hipError_t hip_status = hipSuccess; + hip_status = ihipHostMalloc(ptr, sizeBytes, flags); + return ihipLogStatus(hip_status); +} + +hipError_t hipMallocManaged(void** devPtr, size_t size, unsigned int flags) { + HIP_INIT_SPECIAL_API(hipMallocManaged, (TRACE_MEM), devPtr, size, flags); + HIP_SET_DEVICE(); + hipError_t hip_status = hipSuccess; + if(flags != hipMemAttachGlobal) + hip_status = hipErrorInvalidValue; + else + hip_status = ihipHostMalloc(devPtr, size, hipHostMallocDefault); return ihipLogStatus(hip_status); } diff --git a/projects/clr/hipamd/tests/src/runtimeApi/memory/hipMallocManaged.cpp b/projects/clr/hipamd/tests/src/runtimeApi/memory/hipMallocManaged.cpp new file mode 100644 index 0000000000..bf7662f6b7 --- /dev/null +++ b/projects/clr/hipamd/tests/src/runtimeApi/memory/hipMallocManaged.cpp @@ -0,0 +1,52 @@ +#include +#include +#include "test_common.h" + +/* HIT_START + * BUILD: %t %s ../../test_common.cpp + * RUN: %t -N 256M + * HIT_END + */ + +__global__ +void add(int n, float *x, float *y) +{ + int index = blockIdx.x * blockDim.x + threadIdx.x; + int stride = blockDim.x * gridDim.x; + for (int i = index; i < n; i += stride) + y[i] = x[i] + y[i]; +} + +int main(int argc, char *argv[]) +{ + HipTest::parseStandardArguments(argc, argv, true); + int numElements = N; + bool testResult = true; + float *A, *B; + + hipMallocManaged(&A, numElements*sizeof(float)); + hipMallocManaged(&B, numElements*sizeof(float)); + + for (int i = 0; i < numElements; i++) { + A[i] = 1.0f; + B[i] = 2.0f; + } + + int blockSize = 256; + int numBlocks = (numElements + blockSize - 1) / blockSize; + dim3 dimGrid(numBlocks, 1, 1); + dim3 dimBlock(blockSize, 1, 1); + hipLaunchKernelGGL(add, dimGrid, dimBlock, 0, 0, numElements, A, B); + + hipDeviceSynchronize(); + + float maxError = 0.0f; + for (int i = 0; i < numElements; i++) + maxError = fmax(maxError, fabs(B[i]-3.0f)); + + hipFree(A); + hipFree(B); + if(maxError == 0.0f) + passed(); + failed("Output Mismatch\n"); +}