From 94ea419da747c8f897685b129c10e8b738bdc118 Mon Sep 17 00:00:00 2001
From: haoyuan2 <Hao.Yuan@amd.com>
Date: Fri, 7 Jan 2022 15:15:59 -0800
Subject: [PATCH] SWDEV-316128 - HIP surface API support

Change-Id: I7b5ef4769efb07188915a68caeb4d35360c7aa95


[ROCm/clr commit: 40a6ec42a5370dc52a6f34e861200f9c4cbc7e5d]
---
 projects/clr/hipamd/src/hip_surface.cpp | 64 ++++++++++++++++++++++++-
 1 file changed, 62 insertions(+), 2 deletions(-)

diff --git a/projects/clr/hipamd/src/hip_surface.cpp b/projects/clr/hipamd/src/hip_surface.cpp
index 47c988f1c2..831df321a4 100644
--- a/projects/clr/hipamd/src/hip_surface.cpp
+++ b/projects/clr/hipamd/src/hip_surface.cpp
@@ -23,15 +23,75 @@
 #include "hip_internal.hpp"
 #include <hip/surface_types.h>
 
+hipError_t ihipFree(void* ptr);
+
+struct __hip_surface {
+  uint32_t imageSRD[HIP_IMAGE_OBJECT_SIZE_DWORD];
+  amd::Image* image;
+  hipResourceDesc resDesc;
+
+  __hip_surface(amd::Image* image_, const hipResourceDesc& resDesc_)
+      : image(image_), resDesc(resDesc_) {
+    amd::Context& context = *hip::getCurrentDevice()->asContext();
+    amd::Device& device = *context.devices()[0];
+
+    device::Memory* imageMem = image->getDeviceMemory(device);
+    std::memcpy(imageSRD, imageMem->cpuSrd(), sizeof(imageSRD));
+  }
+};
+
+hipError_t ihipCreateSurfaceObject(hipSurfaceObject_t* pSurfObject,
+                                   const hipResourceDesc* pResDesc) {
+  amd::Device* device = hip::getCurrentDevice()->devices()[0];
+  const device::Info& info = device->info();
+
+  // Validate input params
+  if (pSurfObject == nullptr || pResDesc == nullptr) {
+    return hipErrorInvalidValue;
+  }
+
+  // the type of resource must be a HIP array
+  // hipResourceDesc::res::array::array must be set to a valid HIP array handle.
+  if ((pResDesc->resType != hipResourceTypeArray) || (pResDesc->res.array.array == nullptr)) {
+    return hipErrorInvalidValue;
+  }
+
+  amd::Image* image = nullptr;
+  cl_mem memObj = reinterpret_cast<cl_mem>(pResDesc->res.array.array->data);
+  if (!is_valid(memObj)) {
+    return hipErrorInvalidValue;
+  }
+  image = as_amd(memObj)->asImage();
+
+  void* surfObjectBuffer = nullptr;
+  ihipMalloc(&surfObjectBuffer, sizeof(__hip_surface), CL_MEM_SVM_FINE_GRAIN_BUFFER);
+  if (surfObjectBuffer == nullptr) {
+    return hipErrorOutOfMemory;
+  }
+  *pSurfObject = new (surfObjectBuffer) __hip_surface{image, *pResDesc};
+
+  return hipSuccess;
+}
+
 hipError_t hipCreateSurfaceObject(hipSurfaceObject_t* pSurfObject,
                                   const hipResourceDesc* pResDesc) {
   HIP_INIT_API(hipCreateSurfaceObject, pSurfObject, pResDesc);
 
-  HIP_RETURN(hipErrorNotSupported);
+  HIP_RETURN(ihipCreateSurfaceObject(pSurfObject, pResDesc));
+}
+
+hipError_t ihipDestroySurfaceObject(hipSurfaceObject_t surfaceObject) {
+  HIP_INIT_API(hipDestroySurfaceObject, surfaceObject);
+
+  if (surfaceObject == nullptr) {
+    return hipSuccess;
+  }
+
+  return ihipFree(surfaceObject);
 }
 
 hipError_t hipDestroySurfaceObject(hipSurfaceObject_t surfaceObject) {
   HIP_INIT_API(hipDestroySurfaceObject, surfaceObject);
 
-  HIP_RETURN(hipErrorNotSupported);
+  HIP_RETURN(ihipDestroySurfaceObject(surfaceObject));
 }