From 7272f1608c2a45a3ef1a12282a588d476fde62ad Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Mon, 26 Feb 2018 11:59:03 +0530 Subject: [PATCH 1/4] Added surface object support [ROCm/clr commit: 92283d24d048ebd829c509722cb90979f5950a50] --- projects/clr/hipamd/CMakeLists.txt | 1 + .../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/clr/hipamd/src/hip_surface.cpp | 90 +++++++++++++++++++ 6 files changed, 214 insertions(+) create mode 100644 projects/clr/hipamd/include/hip/hcc_detail/hip_surface_types.h create mode 100644 projects/clr/hipamd/include/hip/hcc_detail/surface_functions.h create mode 100644 projects/clr/hipamd/src/hip_surface.cpp diff --git a/projects/clr/hipamd/CMakeLists.txt b/projects/clr/hipamd/CMakeLists.txt index 1573ddee5c..cc4b8f1d14 100644 --- a/projects/clr/hipamd/CMakeLists.txt +++ b/projects/clr/hipamd/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/clr/hipamd/include/hip/hcc_detail/hip_runtime.h b/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime.h index 924e774af0..21fa609f53 100644 --- a/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime.h +++ b/projects/clr/hipamd/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/clr/hipamd/include/hip/hcc_detail/hip_runtime_api.h b/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime_api.h index 7f159572d7..ccfd56f65b 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 @@ -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/clr/hipamd/include/hip/hcc_detail/hip_surface_types.h b/projects/clr/hipamd/include/hip/hcc_detail/hip_surface_types.h new file mode 100644 index 0000000000..4abe50f606 --- /dev/null +++ b/projects/clr/hipamd/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/clr/hipamd/include/hip/hcc_detail/surface_functions.h b/projects/clr/hipamd/include/hip/hcc_detail/surface_functions.h new file mode 100644 index 0000000000..ed3af3781d --- /dev/null +++ b/projects/clr/hipamd/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/clr/hipamd/src/hip_surface.cpp b/projects/clr/hipamd/src/hip_surface.cpp new file mode 100644 index 0000000000..2b4d1f444c --- /dev/null +++ b/projects/clr/hipamd/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); +} From 685260a08a97f0939948c7ebb4c254f62ae9c950 Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Thu, 1 Mar 2018 12:22:56 +0530 Subject: [PATCH 2/4] Added hip_surface header file [ROCm/clr commit: 7c50ae3b04a9c47015f495045e1cc2b4b624942a] --- projects/clr/hipamd/src/hip_surface.h | 32 +++++++++++++++++++++++++++ 1 file changed, 32 insertions(+) create mode 100644 projects/clr/hipamd/src/hip_surface.h diff --git a/projects/clr/hipamd/src/hip_surface.h b/projects/clr/hipamd/src/hip_surface.h new file mode 100644 index 0000000000..8b30c95f2b --- /dev/null +++ b/projects/clr/hipamd/src/hip_surface.h @@ -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 +struct hipSurface { + hipArray* array; + hipResourceDesc resDesc; +}; + +#endif From 10ee7a6a8063c048b1d6b1d53e92b8aa88801edd Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Sun, 4 Mar 2018 19:05:37 +0530 Subject: [PATCH 3/4] Fixed byte offset issue Added HIP/NVCC support [ROCm/clr commit: d2426e1b9a08d208855e0f479b1765353f53a86b] --- .../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/clr/hipamd/include/hip/hcc_detail/hip_surface_types.h b/projects/clr/hipamd/include/hip/hcc_detail/hip_surface_types.h index 4abe50f606..c1bea64f97 100644 --- a/projects/clr/hipamd/include/hip/hcc_detail/hip_surface_types.h +++ b/projects/clr/hipamd/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/clr/hipamd/include/hip/hcc_detail/surface_functions.h b/projects/clr/hipamd/include/hip/hcc_detail/surface_functions.h index ed3af3781d..2d4c6e4f6d 100644 --- a/projects/clr/hipamd/include/hip/hcc_detail/surface_functions.h +++ b/projects/clr/hipamd/include/hip/hcc_detail/surface_functions.h @@ -29,30 +29,32 @@ THE SOFTWARE. #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) +__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 -__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/clr/hipamd/include/hip/nvcc_detail/hip_runtime_api.h b/projects/clr/hipamd/include/hip/nvcc_detail/hip_runtime_api.h index 902e3620fa..e1050e21e8 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 @@ -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 From 5f632673270cf2fb726b56524971c8add3acae94 Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Sun, 4 Mar 2018 22:49:23 +0530 Subject: [PATCH 4/4] Add surface object API 2D test [ROCm/clr commit: 919eb74a93edcce22c27fae124d7ca82233ff047] --- .../tests/src/surface/hipSurfaceObj2D.cpp | 115 ++++++++++++++++++ 1 file changed, 115 insertions(+) create mode 100644 projects/clr/hipamd/tests/src/surface/hipSurfaceObj2D.cpp diff --git a/projects/clr/hipamd/tests/src/surface/hipSurfaceObj2D.cpp b/projects/clr/hipamd/tests/src/surface/hipSurfaceObj2D.cpp new file mode 100644 index 0000000000..cd87bc6701 --- /dev/null +++ b/projects/clr/hipamd/tests/src/surface/hipSurfaceObj2D.cpp @@ -0,0 +1,115 @@ +/* HIT_START + * BUILD: %t %s ../test_common.cpp + * RUN: %t + * HIT_END + */ +#include +#include +#include + +#include +#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); +} +