From 8e8fb993de41e6adf1801701e2e3ac7c5c12e4f3 Mon Sep 17 00:00:00 2001
From: Jaydeep Patel <jaydeepkumar.patel@amd.com>
Date: Tue, 17 Oct 2023 17:13:54 +0000
Subject: [PATCH] SWDEV-425248 - numerous bug fixes.

Set flag with hipCtxCreate so that get flags works.
Validate hipHostGetDevicePointer for flags!=0.
Validate mem cpy kind and accommodate new type hipMemcpyDeviceToDeviceNoCU.
Match error code for hipGetChannelDesc.

Change-Id: If09a635ac01bc53f1fe2b7df3f3f9c1b0d69a0ab
---
 hipamd/src/hip_context.cpp |  1 +
 hipamd/src/hip_memory.cpp  | 22 ++++++++++++++++++----
 hipamd/src/hip_texture.cpp |  5 ++++-
 3 files changed, 23 insertions(+), 5 deletions(-)

diff --git a/hipamd/src/hip_context.cpp b/hipamd/src/hip_context.cpp
index 391be153f8..6ee0f31d44 100644
--- a/hipamd/src/hip_context.cpp
+++ b/hipamd/src/hip_context.cpp
@@ -159,6 +159,7 @@ hipError_t hipCtxCreate(hipCtx_t *ctx, unsigned int flags,  hipDevice_t device)
 
   // Increment ref count for device primary context
   g_devices[device]->retain();
+  g_devices[device]->setFlags(flags);
   tls.ctxt_stack_.push(g_devices[device]);
 
   HIP_RETURN(hipSuccess);
diff --git a/hipamd/src/hip_memory.cpp b/hipamd/src/hip_memory.cpp
index bd7f005fad..08e8f1aa75 100644
--- a/hipamd/src/hip_memory.cpp
+++ b/hipamd/src/hip_memory.cpp
@@ -358,7 +358,9 @@ hipError_t ihipMemcpy_validate(void* dst, const void* src, size_t sizeBytes,
   if (dst == nullptr || src == nullptr) {
     return hipErrorInvalidValue;
   }
-
+  if (static_cast<uint32_t>(kind) > hipMemcpyDefault && kind != hipMemcpyDeviceToDeviceNoCU) {
+    return hipErrorInvalidMemcpyDirection;
+  }
   size_t sOffset = 0;
   amd::Memory* srcMemory = getMemoryObject(src, sOffset);
   size_t dOffset = 0;
@@ -1478,6 +1480,9 @@ hipError_t hipMemcpyAsync_common(void* dst, const void* src, size_t sizeBytes,
                           hipMemcpyKind kind, hipStream_t stream) {
   STREAM_CAPTURE(hipMemcpyAsync, stream, dst, src, sizeBytes, kind);
 
+  if (static_cast<uint32_t>(kind) > hipMemcpyDefault && kind != hipMemcpyDeviceToDeviceNoCU) {
+    return hipErrorInvalidMemcpyDirection;
+  }
   hip::Stream* hip_stream = hip::getStream(stream);
   if (hip_stream == nullptr) {
     return hipErrorInvalidValue;
@@ -1503,6 +1508,9 @@ hipError_t hipMemcpyHtoDAsync(hipDeviceptr_t dstDevice, void* srcHost, size_t By
   HIP_INIT_API(hipMemcpyHtoDAsync, dstDevice, srcHost, ByteCount, stream);
   hipMemcpyKind kind = hipMemcpyHostToDevice;
   STREAM_CAPTURE(hipMemcpyHtoDAsync, stream, dstDevice, srcHost, ByteCount, kind);
+  if (static_cast<uint32_t>(kind) > hipMemcpyDefault && kind != hipMemcpyDeviceToDeviceNoCU) {
+    return hipErrorInvalidMemcpyDirection;
+  }
   hip::Stream* hip_stream = hip::getStream(stream);
   if (hip_stream == nullptr) {
     HIP_RETURN(hipErrorInvalidValue);
@@ -1516,6 +1524,9 @@ hipError_t hipMemcpyDtoDAsync(hipDeviceptr_t dstDevice, hipDeviceptr_t srcDevice
   HIP_INIT_API(hipMemcpyDtoDAsync, dstDevice, srcDevice, ByteCount, stream);
   hipMemcpyKind kind = hipMemcpyDeviceToDevice;
   STREAM_CAPTURE(hipMemcpyDtoDAsync, stream, dstDevice, srcDevice, ByteCount, kind);
+  if (static_cast<uint32_t>(kind) > hipMemcpyDefault && kind != hipMemcpyDeviceToDeviceNoCU) {
+    return hipErrorInvalidMemcpyDirection;
+  }
   hip::Stream* hip_stream = hip::getStream(stream);
   if (hip_stream == nullptr) {
     HIP_RETURN(hipErrorInvalidValue);
@@ -1529,6 +1540,9 @@ hipError_t hipMemcpyDtoHAsync(void* dstHost, hipDeviceptr_t srcDevice, size_t By
   HIP_INIT_API(hipMemcpyDtoHAsync, dstHost, srcDevice, ByteCount, stream);
   hipMemcpyKind kind = hipMemcpyDeviceToHost;
   STREAM_CAPTURE(hipMemcpyDtoHAsync, stream, dstHost, srcDevice, ByteCount, kind);
+  if (static_cast<uint32_t>(kind) > hipMemcpyDefault && kind != hipMemcpyDeviceToDeviceNoCU) {
+    return hipErrorInvalidMemcpyDirection;
+  }
   hip::Stream* hip_stream = hip::getStream(stream);
   if (hip_stream == nullptr) {
     HIP_RETURN(hipErrorInvalidValue);
@@ -2357,7 +2371,7 @@ hipError_t hipMemcpyParam2D(const hip_Memcpy2D* pCopy) {
 
 hipError_t hipMemcpy2DValidateParams(hipMemcpyKind kind, hipStream_t stream = nullptr) {
 
-  if (kind < hipMemcpyHostToHost || kind > hipMemcpyDefault) {
+  if (static_cast<uint32_t>(kind) > hipMemcpyDefault && kind != hipMemcpyDeviceToDeviceNoCU) {
     return hipErrorInvalidMemcpyDirection;
   }
 
@@ -2788,7 +2802,7 @@ hipError_t ihipMemcpy3D_validate(const hipMemcpy3DParms* p) {
     }
   }
 
-  if (p->kind < hipMemcpyHostToHost || p->kind > hipMemcpyDefault) {
+  if (static_cast<uint32_t>(p->kind) > hipMemcpyDefault && p->kind != hipMemcpyDeviceToDeviceNoCU) {
     return hipErrorInvalidMemcpyDirection;
   }
   //If src and dst ptr are null then kind must be either h2h or def.
@@ -3505,7 +3519,7 @@ hipError_t hipIpcCloseMemHandle(void* dev_ptr) {
 hipError_t hipHostGetDevicePointer(void** devicePointer, void* hostPointer, unsigned flags) {
   HIP_INIT_API(hipHostGetDevicePointer, devicePointer, hostPointer, flags);
 
-  if (devicePointer == nullptr) {
+  if (devicePointer == nullptr || flags != 0) {
     HIP_RETURN(hipErrorInvalidValue);
   }
 
diff --git a/hipamd/src/hip_texture.cpp b/hipamd/src/hip_texture.cpp
index 9eba4bd0f6..3c831b108d 100644
--- a/hipamd/src/hip_texture.cpp
+++ b/hipamd/src/hip_texture.cpp
@@ -744,9 +744,12 @@ hipError_t hipGetChannelDesc(hipChannelFormatDesc* desc,
                              hipArray_const_t array) {
   HIP_INIT_API(hipGetChannelDesc, desc, array);
 
-  if ((desc == nullptr) || (array == nullptr)) {
+  if (desc == nullptr) {
     HIP_RETURN(hipErrorInvalidValue);
   }
+  if (array == nullptr) {
+    HIP_RETURN(hipErrorInvalidHandle);
+  }
   amd::Device* device = hip::getCurrentDevice()->devices()[0];
   const device::Info& info = device->info();
   if (!info.imageSupport_) {