From 413c9ac9cb5cb48edaa93aae9b5fe38c7262a0ab Mon Sep 17 00:00:00 2001
From: Rahul Garg <rahul.garg@amd.com>
Date: Sun, 4 Mar 2018 19:05:37 +0530
Subject: [PATCH] Fixed byte offset issue Added HIP/NVCC support

[ROCm/hip commit: d697739d57b84d94638ed1564807f33195a004ef]
---
 .../hip/hcc_detail/hip_surface_types.h        |  6 ++--
 .../hip/hcc_detail/surface_functions.h        | 32 ++++++++++---------
 .../include/hip/nvcc_detail/hip_runtime_api.h | 17 ++++++++++
 3 files changed, 37 insertions(+), 18 deletions(-)

diff --git a/projects/hip/include/hip/hcc_detail/hip_surface_types.h b/projects/hip/include/hip/hcc_detail/hip_surface_types.h
index 4abe50f606..c1bea64f97 100644
--- a/projects/hip/include/hip/hcc_detail/hip_surface_types.h
+++ b/projects/hip/include/hip/hcc_detail/hip_surface_types.h
@@ -48,9 +48,9 @@ struct surfaceReference
  */
 enum hipSurfaceBoundaryMode
 {
-    hipSurfaceBoundaryModeZero   = 0,
-    hipSurfaceBoundaryModeTrap   = 1,
-    hipSurfaceBoundaryModeClamp  = 2
+    hipBoundaryModeZero   = 0,
+    hipBoundaryModeTrap   = 1,
+    hipBoundaryModeClamp  = 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
index ed3af3781d..2d4c6e4f6d 100644
--- a/projects/hip/include/hip/hcc_detail/surface_functions.h
+++ b/projects/hip/include/hip/hcc_detail/surface_functions.h
@@ -29,30 +29,32 @@ THE SOFTWARE.
 
 #define __SURFACE_FUNCTIONS_DECL__ static __inline__ __device__
 template <class T>
-__SURFACE_FUNCTIONS_DECL__ void surf2Dread(T* data, hipSurfaceObject_t surfObj, int x, int y, int boundaryMode = hipSurfaceBoundaryModeZero)
+__SURFACE_FUNCTIONS_DECL__ void surf2Dread(T* data, hipSurfaceObject_t surfObj, int x, int y, int boundaryMode = hipBoundaryModeZero)
 {
-    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) {
+    hipArray* arrayPtr = (hipArray*) surfObj;
+    size_t width = arrayPtr->width;
+    size_t height = arrayPtr->height;
+    int32_t xOffset = x / sizeof(T);
+    T* dataPtr = (T*) arrayPtr->data;
+    if((xOffset > width) || (xOffset < 0) || (y > height) ||(y < 0)) {
+        if(boundaryMode == hipBoundaryModeZero) {
             *data = 0;
         }
     } else {
-        *data = *(temp1+ + y*width + x);
+        *data = *(dataPtr + y*width + xOffset);
     }
 }
 
 template <class T>
-__SURFACE_FUNCTIONS_DECL__ void surf2Dwrite(T data, hipSurfaceObject_t surfObj, int x, int y, int boundaryMode = hipSurfaceBoundaryModeZero)
+__SURFACE_FUNCTIONS_DECL__ void surf2Dwrite(T data, hipSurfaceObject_t surfObj, int x, int y, int boundaryMode = hipBoundaryModeZero)
 {
-    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;
+    hipArray* arrayPtr = (hipArray*) surfObj;
+    size_t width = arrayPtr->width;
+    size_t height = arrayPtr->height;
+    int32_t xOffset = x / sizeof(T);
+    T* dataPtr = (T*) arrayPtr->data;
+    if(!((xOffset > width) || (xOffset < 0) || (y > height) ||(y < 0))){
+        *(dataPtr +y*width + xOffset) = data;
     }
 }
 
diff --git a/projects/hip/include/hip/nvcc_detail/hip_runtime_api.h b/projects/hip/include/hip/nvcc_detail/hip_runtime_api.h
index 902e3620fa..e1050e21e8 100644
--- a/projects/hip/include/hip/nvcc_detail/hip_runtime_api.h
+++ b/projects/hip/include/hip/nvcc_detail/hip_runtime_api.h
@@ -75,6 +75,11 @@ typedef enum hipChannelFormatKind {
     hipChannelFormatKindNone = 3
 }hipChannelFormatKind;
 
+#define hipSurfaceBoundaryMode cudaSurfaceBoundaryMode
+#define hipBoundaryModeZero cudaBoundaryModeZero
+#define hipBoundaryModeTrap cudaBoundaryModeTrap
+#define hipBoundaryModeClamp cudaBoundaryModeClamp
+
 //hipResourceType
 #define hipResourceType cudaResourceType
 #define hipResourceTypeArray cudaResourceTypeArray
@@ -149,6 +154,7 @@ typedef struct cudaArray* hipArray_const_t;
 #define hipArrayDefault cudaArrayDefault
 
 typedef cudaTextureObject_t hipTextureObject_t;
+typedef cudaSurfaceObject_t hipSurfaceObject_t;
 #define hipTextureType2D cudaTextureType2D;
 #define hipDeviceMapHost cudaDeviceMapHost
 
@@ -1143,6 +1149,17 @@ inline static hipError_t hipDestroyTextureObject(hipTextureObject_t textureObjec
 {
     return hipCUDAErrorTohipError(cudaDestroyTextureObject(textureObject));
 }
+
+inline static hipError_t hipCreateSurfaceObject(hipSurfaceObject_t* pSurfObject, const hipResourceDesc* pResDesc)
+{
+    return hipCUDAErrorTohipError(cudaCreateSurfaceObject(pSurfObject, pResDesc));
+}
+
+inline static hipError_t hipDestroySurfaceObject(hipSurfaceObject_t surfaceObject)
+{
+    return hipCUDAErrorTohipError(cudaDestroySurfaceObject(surfaceObject));
+}
+
 #endif //__CUDACC__
 
 #endif //HIP_INCLUDE_HIP_NVCC_DETAIL_HIP_RUNTIME_API_H