Merge pull request #363 from gargrahul/surface_object_api

Added surface object support [ROCm/clr commit: eee7fa6072]
2018-03-09 11:50:30 +05:30
@@ -187,6 +187,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)

@@ -93,6 +93,7 @@ extern int HIP_TRACE_API;
 #include <hip/hcc_detail/math_functions.h>
 #include <hip/hcc_detail/device_functions.h>
 #include <hip/hcc_detail/texture_functions.h>
+#include <hip/hcc_detail/surface_functions.h>

 // TODO-HCC remove old definitions ; ~1602 hcc supports __HCC_ACCELERATOR__ define.
 #if defined (__KALMAR_ACCELERATOR__) && !defined (__HCC_ACCELERATOR__)
@@ -39,6 +39,7 @@ THE SOFTWARE.
 #include <hip/hip_runtime_api.h>
 #include <hip/hcc_detail/driver_types.h>
 #include <hip/hcc_detail/hip_texture_types.h>
+#include <hip/hcc_detail/hip_surface_types.h>

 #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
 /**
 * @}
@@ -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<hip/hcc_detail/driver_types.h>
+
+/**
+ * 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
+{
+    hipBoundaryModeZero   = 0,
+    hipBoundaryModeTrap   = 1,
+    hipBoundaryModeClamp  = 2
+};
+
+#endif /* !HIP_INCLUDE_HIP_HCC_DETAIL_HIP_SURFACE_TYPES_H */
+
+
@@ -0,0 +1,61 @@
+/*
+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 <hc.hpp>
+#include <hc_short_vector.hpp>
+#include <hip/hcc_detail/hip_surface_types.h>
+
+#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 = hipBoundaryModeZero)
+{
+    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 = *(dataPtr + y*width + xOffset);
+    }
+}
+
+template <class T>
+__SURFACE_FUNCTIONS_DECL__ void surf2Dwrite(T data, hipSurfaceObject_t surfObj, int x, int y, int boundaryMode = hipBoundaryModeZero)
+{
+    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;
+    }
+}
+
+#endif
@@ -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
@@ -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 <map>
+
+#include <string.h>
+
+#include "hip/hip_runtime.h"
+#include "hip_hcc_internal.h"
+#include "trace_helper.h"
+
+#include "hip_surface.h"
+
+static std::map<hipSurfaceObject_t, hipSurface*> 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);
+}
@@ -0,0 +1,32 @@
+/*
+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.
+*/
+
+#ifndef HIP_INCLUDE_HIP_HCC_DETAIL_HIP_SURFACE_H
+#define HIP_INCLUDE_HIP_HCC_DETAIL_HIP_SURFACE_H
+
+#include <hip/hcc_detail/hip_surface_types.h>
+struct hipSurface {
+    hipArray* array;
+    hipResourceDesc resDesc;
+};
+
+#endif
@@ -0,0 +1,115 @@
+/* HIT_START
+ * BUILD: %t %s ../test_common.cpp
+ * RUN: %t
+ * HIT_END
+ */
+#include <stdlib.h>
+#include <stdio.h>
+#include <string.h>
+
+#include <hip/hip_runtime.h>
+#include "test_common.h"
+
+bool testResult = true;
+
+__global__ void tex2DKernel(hipSurfaceObject_t surfaceObject,
+                            hipSurfaceObject_t outputSurfObj,
+                            int width,
+                            int height)
+{
+    int x = blockIdx.x*blockDim.x + threadIdx.x;
+    int y = blockIdx.y*blockDim.y + threadIdx.y;
+    float data;
+    surf2Dread(&data, surfaceObject, x*4, y, hipBoundaryModeZero);
+    surf2Dwrite(data, outputSurfObj, x*4, y, hipBoundaryModeZero);
+}
+
+void runTest(int argc, char **argv);
+
+int main(int argc, char **argv)
+{
+    runTest(argc, argv);
+
+    if(testResult) {
+        passed();
+    } else {
+        exit(EXIT_FAILURE);
+    }
+
+}
+
+void runTest(int argc, char **argv)
+{
+    unsigned int width = 256;
+    unsigned int height = 256;
+    unsigned int size = width * height * sizeof(float);
+    float* hData = (float*) malloc(size);
+    memset(hData, 0, size);
+    for (int i = 0; i < height; i++) {
+        for (int j = 0; j < width; j++) {
+            hData[i*width+j] = i*width+j;
+        }
+    }
+    printf("hData: ");
+    for (int i = 0; i < 64; i++) {
+        printf("%f  ", hData[i]);
+    }
+    printf("\n");
+
+    hipChannelFormatDesc channelDesc = hipCreateChannelDesc(32, 0, 0, 0, hipChannelFormatKindFloat);
+    hipArray *hipArray, *hipOutArray;
+    hipMallocArray(&hipArray, &channelDesc, width, height);
+
+    hipMemcpyToArray(hipArray, 0, 0, hData, size, hipMemcpyHostToDevice);
+
+    struct hipResourceDesc resDesc;
+    memset(&resDesc, 0, sizeof(resDesc));
+    resDesc.resType = hipResourceTypeArray;
+    resDesc.res.array.array = hipArray;
+    // Create surface object
+    hipSurfaceObject_t surfaceObject = 0;
+    hipCreateSurfaceObject(&surfaceObject, &resDesc);
+
+    hipMallocArray(&hipOutArray, &channelDesc, width, height);
+    struct hipResourceDesc resOutDesc;
+    memset(&resOutDesc, 0, sizeof(resOutDesc));
+    resOutDesc.resType = hipResourceTypeArray;
+    resOutDesc.res.array.array = hipOutArray;
+    hipSurfaceObject_t outSurfaceObject = 0;
+    hipCreateSurfaceObject(&outSurfaceObject, &resOutDesc);
+
+    float* dData = NULL;
+    hipMalloc((void **) &dData, size);
+
+    dim3 dimBlock(16, 16, 1);
+    dim3 dimGrid(width / dimBlock.x, height / dimBlock.y, 1);
+
+    hipLaunchKernelGGL(tex2DKernel, dim3(dimGrid), dim3(dimBlock), 0, 0, surfaceObject,outSurfaceObject, width, height);
+
+    hipDeviceSynchronize();
+
+    float *hOutputData = (float *) malloc(size);
+    memset(hOutputData, 0,  size);
+    hipMemcpyFromArray(hOutputData, hipOutArray, 0, 0, size,hipMemcpyDeviceToHost);
+
+    printf("dData: ");
+    for (int i = 0; i < 64; i++) {
+        printf("%f  ", hOutputData[i]);
+    }
+    printf("\n");
+    for (int i = 0; i < height; i++) {
+        for (int j = 0; j < width; j++) {
+            if (hData[i*width+j] != hOutputData[i*width+j]) {
+                printf("Difference [ %d %d ]:%f ----%f\n",i, j, hData[i*width+j] , hOutputData[i*width+j]);
+                testResult = false;
+                break;
+            }
+        }
+    }
+    hipDestroySurfaceObject(surfaceObject);
+    hipDestroySurfaceObject(outSurfaceObject);
+    hipFree(dData);
+    hipFreeArray(hipArray);
+    hipFreeArray(hipOutArray);
+}
+