From ff74aadd87f429dbe245da3816b67260322f2bfe Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Mon, 26 Feb 2018 11:59:03 +0530 Subject: [PATCH] Added surface object support [ROCm/hip commit: 5c3b91a0a1a43aeda1df7e508f9e4a5dc7503dc1] --- projects/hip/CMakeLists.txt | 1 + .../hip/include/hip/hcc_detail/hip_runtime.h | 1 + .../include/hip/hcc_detail/hip_runtime_api.h | 5 ++ .../hip/hcc_detail/hip_surface_types.h | 58 ++++++++++++ .../hip/hcc_detail/surface_functions.h | 59 ++++++++++++ projects/hip/src/hip_surface.cpp | 90 +++++++++++++++++++ 6 files changed, 214 insertions(+) create mode 100644 projects/hip/include/hip/hcc_detail/hip_surface_types.h create mode 100644 projects/hip/include/hip/hcc_detail/surface_functions.h create mode 100644 projects/hip/src/hip_surface.cpp diff --git a/projects/hip/CMakeLists.txt b/projects/hip/CMakeLists.txt index 1573ddee5c..cc4b8f1d14 100644 --- a/projects/hip/CMakeLists.txt +++ b/projects/hip/CMakeLists.txt @@ -180,6 +180,7 @@ if(HIP_PLATFORM STREQUAL "hcc") src/hip_db.cpp src/grid_launch.cpp src/hip_texture.cpp + src/hip_surface.cpp src/env.cpp src/program_state.cpp) diff --git a/projects/hip/include/hip/hcc_detail/hip_runtime.h b/projects/hip/include/hip/hcc_detail/hip_runtime.h index 924e774af0..21fa609f53 100644 --- a/projects/hip/include/hip/hcc_detail/hip_runtime.h +++ b/projects/hip/include/hip/hcc_detail/hip_runtime.h @@ -93,6 +93,7 @@ extern int HIP_TRACE_API; #include #include #include +#include // TODO-HCC remove old definitions ; ~1602 hcc supports __HCC_ACCELERATOR__ define. #if defined (__KALMAR_ACCELERATOR__) && !defined (__HCC_ACCELERATOR__) diff --git a/projects/hip/include/hip/hcc_detail/hip_runtime_api.h b/projects/hip/include/hip/hcc_detail/hip_runtime_api.h index 7f159572d7..ccfd56f65b 100644 --- a/projects/hip/include/hip/hcc_detail/hip_runtime_api.h +++ b/projects/hip/include/hip/hcc_detail/hip_runtime_api.h @@ -39,6 +39,7 @@ THE SOFTWARE. #include #include #include +#include #if defined (__HCC__) && (__hcc_workweek__ < 16155) #error("This version of HIP requires a newer version of HCC."); @@ -2407,6 +2408,10 @@ hipError_t hipTexRefSetAddress( size_t* offset, textureReference* tex, hipDevice hipError_t hipTexRefSetAddress2D( textureReference* tex, const HIP_ARRAY_DESCRIPTOR* desc, hipDeviceptr_t devPtr, size_t pitch ); +hipError_t hipCreateSurfaceObject(hipSurfaceObject_t* pSurfObject, const hipResourceDesc* pResDesc); + +hipError_t hipDestroySurfaceObject(hipSurfaceObject_t surfaceObject); + // doxygen end Texture /** * @} diff --git a/projects/hip/include/hip/hcc_detail/hip_surface_types.h b/projects/hip/include/hip/hcc_detail/hip_surface_types.h new file mode 100644 index 0000000000..4abe50f606 --- /dev/null +++ b/projects/hip/include/hip/hcc_detail/hip_surface_types.h @@ -0,0 +1,58 @@ +/* +Copyright (c) 2015- present Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +/** + * @file hcc_detail/hip_surface_types.h + * @brief Defines surface types for HIP runtime. + */ + +#ifndef HIP_INCLUDE_HIP_HCC_DETAIL_HIP_SURFACE_TYPES_H +#define HIP_INCLUDE_HIP_HCC_DETAIL_HIP_SURFACE_TYPES_H + +#include + +/** + * An opaque value that represents a hip surface object + */ +typedef unsigned long long hipSurfaceObject_t; + +/** + * hip surface reference + */ +struct surfaceReference +{ + hipSurfaceObject_t surfaceObject; +}; + +/** + * hip surface boundary modes + */ +enum hipSurfaceBoundaryMode +{ + hipSurfaceBoundaryModeZero = 0, + hipSurfaceBoundaryModeTrap = 1, + hipSurfaceBoundaryModeClamp = 2 +}; + +#endif /* !HIP_INCLUDE_HIP_HCC_DETAIL_HIP_SURFACE_TYPES_H */ + + diff --git a/projects/hip/include/hip/hcc_detail/surface_functions.h b/projects/hip/include/hip/hcc_detail/surface_functions.h new file mode 100644 index 0000000000..ed3af3781d --- /dev/null +++ b/projects/hip/include/hip/hcc_detail/surface_functions.h @@ -0,0 +1,59 @@ +/* +Copyright (c) 2018 - present Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#ifndef HIP_INCLUDE_HIP_HCC_DETAIL_SURFACE_FUNCTIONS_H +#define HIP_INCLUDE_HIP_HCC_DETAIL_SURFACE_FUNCTIONS_H + +#include +#include +#include + +#define __SURFACE_FUNCTIONS_DECL__ static __inline__ __device__ +template +__SURFACE_FUNCTIONS_DECL__ void surf2Dread(T* data, hipSurfaceObject_t surfObj, int x, int y, int boundaryMode = hipSurfaceBoundaryModeZero) +{ + hipArray* temp = (hipArray*) surfObj; + size_t width = temp->width; + size_t height = temp->height; + T* temp1 = (T*) temp->data; + if((x > width) || (x < 0) || (y > height) ||(y < 0)) { + if(boundaryMode == hipSurfaceBoundaryModeZero) { + *data = 0; + } + } else { + *data = *(temp1+ + y*width + x); + } +} + +template +__SURFACE_FUNCTIONS_DECL__ void surf2Dwrite(T data, hipSurfaceObject_t surfObj, int x, int y, int boundaryMode = hipSurfaceBoundaryModeZero) +{ + hipArray* temp = (hipArray*) surfObj; + size_t width = temp->width; + size_t height = temp->height; + T* temp1 = (T*) temp->data; + if(!((x > width) || (x < 0) || (y > height) ||(y < 0))){ + *(temp1 +y*width + x) = data; + } +} + +#endif diff --git a/projects/hip/src/hip_surface.cpp b/projects/hip/src/hip_surface.cpp new file mode 100644 index 0000000000..2b4d1f444c --- /dev/null +++ b/projects/hip/src/hip_surface.cpp @@ -0,0 +1,90 @@ +/* +Copyright (c) 2018 - present Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#include + +#include + +#include "hip/hip_runtime.h" +#include "hip_hcc_internal.h" +#include "trace_helper.h" + +#include "hip_surface.h" + +static std::map surfaceHash; + +void saveSurfaceInfo(const hipSurface* pSurface, + const hipResourceDesc* pResDesc) +{ + if (pResDesc != nullptr) { + memcpy((void*)&(pSurface->resDesc), (void*)pResDesc, sizeof(hipResourceDesc)); + } +} + +// Surface Object APIs +hipError_t hipCreateSurfaceObject(hipSurfaceObject_t* pSurfObject, + const hipResourceDesc* pResDesc) +{ + HIP_INIT_API(pSurfObject, pResDesc); + hipError_t hip_status = hipSuccess; + + auto ctx = ihipGetTlsDefaultCtx(); + if (ctx) { + hipSurface* pSurface = (hipSurface*) malloc(sizeof(hipSurface)); + if (pSurface != nullptr) { + memset(pSurface, 0, sizeof(hipSurface)); + saveSurfaceInfo(pSurface, pResDesc); + } + + switch (pResDesc->resType) { + case hipResourceTypeArray: + pSurface->array = pResDesc->res.array.array; + break; + default: + break; + } + unsigned int* surfObj; + hipMalloc((void **) &surfObj, sizeof(hipArray)); + hipMemcpy(surfObj, (void *)pResDesc->res.array.array, sizeof(hipArray), hipMemcpyHostToDevice); + *pSurfObject = (hipSurfaceObject_t) surfObj; + surfaceHash[*pSurfObject] = pSurface; + } + + return ihipLogStatus(hip_status); +} + +hipError_t hipDestroySurfaceObject(hipSurfaceObject_t surfaceObject) +{ + HIP_INIT_API(surfaceObject); + + hipError_t hip_status = hipSuccess; + + auto ctx = ihipGetTlsDefaultCtx(); + if (ctx) { + hipSurface* pSurface = surfaceHash[surfaceObject]; + if (pSurface != nullptr) { + free(pSurface); + surfaceHash.erase(surfaceObject); + } + } + return ihipLogStatus(hip_status); +}