From db1ce3ba84264f6b41eb038f091a9237dcf6bf8e Mon Sep 17 00:00:00 2001 From: Aditya Atluri Date: Tue, 8 Mar 2016 12:57:22 -0600 Subject: [PATCH] Added hipHostRegister for hip with tests and added copyright [ROCm/hip commit: 102f1733967b342490fcdf9faa46c5b755969837] --- .../hip/include/hcc_detail/hip_runtime_api.h | 24 ++++++++++ projects/hip/src/hip_hcc.cpp | 45 +++++++++++++++++++ projects/hip/tests/src/hipHostRegister.cpp | 26 ++++++++++- projects/hip/tests/src/hipKernel.cpp | 19 ++++++++ projects/hip/tests/src/hipMemcpyAsync.cpp | 19 ++++++++ .../hip/tests/src/hipRandomMemcpyAsync.cpp | 19 ++++++++ .../hip/tests/src/hipSimpleAtomicsTest.cpp | 19 ++++++++ projects/hip/tests/src/hipStream.h | 19 ++++++++ projects/hip/tests/src/hip_ballot.cpp | 19 ++++++++ projects/hip/tests/src/specialFunc.cu | 19 ++++++++ projects/hip/tests/src/test_common.h | 19 ++++++++ 11 files changed, 246 insertions(+), 1 deletion(-) diff --git a/projects/hip/include/hcc_detail/hip_runtime_api.h b/projects/hip/include/hcc_detail/hip_runtime_api.h index 0d83dedd64..9f0b55251a 100644 --- a/projects/hip/include/hcc_detail/hip_runtime_api.h +++ b/projects/hip/include/hcc_detail/hip_runtime_api.h @@ -61,6 +61,11 @@ extern "C" { #define hipHostAllocMapped 0x2 #define hipHostAllocWriteCombined 0x4 +#define hipHostRegisterDefault 0x0 +#define hipHostRegisterPortable 0x1 +#define hipHostRegisterMapped 0x2 +#define hipHostRegisterIoMemory 0x4 + /** * @warning On AMD devices and recent Nvidia devices, these hints and controls are ignored. */ @@ -696,6 +701,25 @@ hipError_t hipHostGetDevicePointer(void** devPtr, void* hstPtr, size_t size) ; */ hipError_t hipHostGetFlags(unsigned int* flagsPtr, void* hostPtr) ; +/** + * @brief Pin host memory + * + * @param[out] hostPtr Pointer to host memory to be pinned + * @param[in] sizeBytes size of the host memory + * @param[in] flags Type of pinning the the host memory + * @return Error code + */ +hipError_t hipHostRegister(void* hostPtr, size_t sizeBytes, unsigned int flags) ; + +/** + * @brief Un-pin host pointer + * + * @param[in] hostPtr Pinned Host Pointer + * @return Error code + */ +hipError_t hipHostUnregister(void* hostPtr) ; + + /** * @brief Free memory allocated by the hcc hip memory allocation API. * This API performs an implicit hipDeviceSynchronize() call. diff --git a/projects/hip/src/hip_hcc.cpp b/projects/hip/src/hip_hcc.cpp index 313d0ebbd0..4a4dd31b0a 100644 --- a/projects/hip/src/hip_hcc.cpp +++ b/projects/hip/src/hip_hcc.cpp @@ -2090,6 +2090,51 @@ hipError_t hipHostGetFlags(unsigned int* flagsPtr, void* hostPtr) return ihipLogStatus(hip_status); } +hipError_t hipHostRegister(void *hostPtr, size_t sizeBytes, unsigned int flags) +{ + std::call_once(hip_initialized, ihipInit); + hipError_t hip_status = hipSuccess; + + auto device = ihipGetTlsDefaultDevice(); + void* srcPtr; + if(hostPtr == NULL){ + return ihipLogStatus(hipErrorInvalidValue); + } + if(device){ + if(flags == hipHostAllocDefault){ + hsa_status_t hsa_status = hsa_amd_memory_lock(hostPtr, sizeBytes, &device->_hsa_agent, 1, &srcPtr); + if(hsa_status == HSA_STATUS_SUCCESS){ + hip_status = hipSuccess; + }else{ + hip_status = hipErrorMemoryAllocation; + } + } + else if (flags | hipHostRegisterMapped){ + hsa_status_t hsa_status = hsa_amd_memory_lock(hostPtr, sizeBytes, &device->_hsa_agent, 1, &srcPtr); + //TODO: Added feature for actual host pointer being tracked + if(hsa_status != HSA_STATUS_SUCCESS){ + hip_status = hipErrorMemoryAllocation; + } + } + } + return ihipLogStatus(hip_status); +} + +hipError_t hipHostUnregister(void *hostPtr){ + std::call_once(hip_initialized, ihipInit); + hipError_t hip_status = hipSuccess; + if(hostPtr == NULL){ + hip_status = hipErrorInvalidValue; + }else{ + hsa_status_t hsa_status = hsa_amd_memory_unlock(hostPtr); + if(hsa_status != HSA_STATUS_SUCCESS){ + hip_status = hipErrorInvalidValue; +// TODO: Add a different return error. This is not true + } + } + return ihipLogStatus(hip_status); +} + //--- hipError_t hipMemcpyToSymbol(const char* symbolName, const void *src, size_t count, size_t offset, hipMemcpyKind kind) { diff --git a/projects/hip/tests/src/hipHostRegister.cpp b/projects/hip/tests/src/hipHostRegister.cpp index 241d37a802..b9e4632369 100644 --- a/projects/hip/tests/src/hipHostRegister.cpp +++ b/projects/hip/tests/src/hipHostRegister.cpp @@ -1,4 +1,24 @@ +/* +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. +*/ + #include"test_common.h" +#include __global__ void Inc(hipLaunchParm lp, float *Ad){ int tx = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x; @@ -8,7 +28,11 @@ Ad[tx] = Ad[tx] + float(1); int main(){ float *A, *Ad; const size_t size = N * sizeof(float); - A = (float*)malloc(size); +#ifdef __HIP_PLATFORM_NVCC__ + A = (float*)malloc(size*2); +#else + A = (float*)memalign(64, size); +#endif HIPCHECK(hipHostRegister(A, size, 0)); for(int i=0;i #include #include diff --git a/projects/hip/tests/src/hipSimpleAtomicsTest.cpp b/projects/hip/tests/src/hipSimpleAtomicsTest.cpp index 975b94d18b..1be32f6679 100644 --- a/projects/hip/tests/src/hipSimpleAtomicsTest.cpp +++ b/projects/hip/tests/src/hipSimpleAtomicsTest.cpp @@ -1,3 +1,22 @@ +/* +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. +*/ + // includes, system #include #include diff --git a/projects/hip/tests/src/hipStream.h b/projects/hip/tests/src/hipStream.h index f9ec3472d0..3cf1284671 100644 --- a/projects/hip/tests/src/hipStream.h +++ b/projects/hip/tests/src/hipStream.h @@ -1,3 +1,22 @@ +/* +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. +*/ + #ifndef HIPSTREAM_H #define HIPSTREAM_H #include diff --git a/projects/hip/tests/src/hip_ballot.cpp b/projects/hip/tests/src/hip_ballot.cpp index 5af4c32d7b..e1adb3095d 100644 --- a/projects/hip/tests/src/hip_ballot.cpp +++ b/projects/hip/tests/src/hip_ballot.cpp @@ -1,3 +1,22 @@ +/* +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. +*/ + #include #include diff --git a/projects/hip/tests/src/specialFunc.cu b/projects/hip/tests/src/specialFunc.cu index c5c1931024..085be062d9 100644 --- a/projects/hip/tests/src/specialFunc.cu +++ b/projects/hip/tests/src/specialFunc.cu @@ -1,3 +1,22 @@ +/* +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. +*/ + //Test to ensure hipify runs correctly. // Hipify may report warnings for some missing/unsupported functions diff --git a/projects/hip/tests/src/test_common.h b/projects/hip/tests/src/test_common.h index e37eec7e86..0a2ba96d22 100644 --- a/projects/hip/tests/src/test_common.h +++ b/projects/hip/tests/src/test_common.h @@ -1,3 +1,22 @@ +/* +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. +*/ + #include #include #include