From 1ab7f22334234e66fdffec0eb7bfec7a789bdedf Mon Sep 17 00:00:00 2001
From: Satyanvesh Dittakavi <Satyanvesh.Dittakavi@amd.com>
Date: Tue, 19 Apr 2022 12:31:13 +0000
Subject: [PATCH] SWDEV-332606 - SWDEV-333585 - Fix hipMemcpySymbol* and
 hipGetSymbol* APIs

- hipMemcpy[From/To]Symbol APIs to return hipErrorInvalidMemcpyDirection
if the copy kind is incorrect
- hipGetSymbolSize and hipGetSymbolAddress APIs to return hipErrorInvalidValue
if the pointer is null

Change-Id: I32f428045c804b1e4f648e91b507f749f60c09d0


[ROCm/clr commit: f1fb6127748817f4baec546451d7d3460f4bba8c]
---
 projects/clr/hipamd/src/hip_memory.cpp   | 18 ++++++++++++++++++
 projects/clr/hipamd/src/hip_platform.cpp |  6 ++++++
 2 files changed, 24 insertions(+)

diff --git a/projects/clr/hipamd/src/hip_memory.cpp b/projects/clr/hipamd/src/hip_memory.cpp
index 510307b2ba..bfe7a53cc2 100644
--- a/projects/clr/hipamd/src/hip_memory.cpp
+++ b/projects/clr/hipamd/src/hip_memory.cpp
@@ -1087,6 +1087,11 @@ inline hipError_t ihipMemcpySymbol_validate(const void* symbol, size_t sizeBytes
 hipError_t hipMemcpyToSymbol_common(const void* symbol, const void* src, size_t sizeBytes,
                              size_t offset, hipMemcpyKind kind, hipStream_t stream=nullptr) {
   CHECK_STREAM_CAPTURING();
+
+  if (kind != hipMemcpyHostToDevice) {
+    HIP_RETURN(hipErrorInvalidMemcpyDirection);
+  }
+
   size_t sym_size = 0;
   hipDeviceptr_t device_ptr = nullptr;
 
@@ -1115,6 +1120,11 @@ hipError_t hipMemcpyToSymbol_spt(const void* symbol, const void* src, size_t siz
 hipError_t hipMemcpyFromSymbol_common(void* dst, const void* symbol, size_t sizeBytes,
                                size_t offset, hipMemcpyKind kind, hipStream_t stream=nullptr) {
   CHECK_STREAM_CAPTURING();
+
+  if (kind != hipMemcpyDeviceToHost) {
+    HIP_RETURN(hipErrorInvalidMemcpyDirection);
+  }
+
   size_t sym_size = 0;
   hipDeviceptr_t device_ptr = nullptr;
 
@@ -1146,6 +1156,10 @@ hipError_t hipMemcpyToSymbolAsync(const void* symbol, const void* src, size_t si
 
   STREAM_CAPTURE(hipMemcpyToSymbolAsync, stream, symbol, src, sizeBytes, offset, kind);
 
+  if (kind != hipMemcpyHostToDevice) {
+    HIP_RETURN(hipErrorInvalidMemcpyDirection);
+  }
+
   size_t sym_size = 0;
   hipDeviceptr_t device_ptr = nullptr;
 
@@ -1163,6 +1177,10 @@ hipError_t hipMemcpyFromSymbolAsync(void* dst, const void* symbol, size_t sizeBy
 
   STREAM_CAPTURE(hipMemcpyFromSymbolAsync, stream, dst, symbol, sizeBytes, offset, kind);
 
+  if (kind != hipMemcpyDeviceToHost) {
+    HIP_RETURN(hipErrorInvalidMemcpyDirection);
+  }
+
   size_t sym_size = 0;
   hipDeviceptr_t device_ptr = nullptr;
 
diff --git a/projects/clr/hipamd/src/hip_platform.cpp b/projects/clr/hipamd/src/hip_platform.cpp
index 26de428030..9c791f3a10 100644
--- a/projects/clr/hipamd/src/hip_platform.cpp
+++ b/projects/clr/hipamd/src/hip_platform.cpp
@@ -278,6 +278,9 @@ hipError_t hipGetSymbolAddress(void** devPtr, const void* symbol) {
   HIP_INIT_API(hipGetSymbolAddress, devPtr, symbol);
 
   hipError_t hip_error = hipSuccess;
+  if (devPtr == nullptr) {
+    HIP_RETURN(hipErrorInvalidValue);
+  }
   size_t sym_size = 0;
 
   HIP_RETURN_ONFAIL(PlatformState::instance().getStatGlobalVar(symbol, ihipGetDevice(), devPtr, &sym_size));
@@ -288,6 +291,9 @@ hipError_t hipGetSymbolAddress(void** devPtr, const void* symbol) {
 hipError_t hipGetSymbolSize(size_t* sizePtr, const void* symbol) {
   HIP_INIT_API(hipGetSymbolSize, sizePtr, symbol);
 
+  if (sizePtr == nullptr) {
+    HIP_RETURN(hipErrorInvalidValue);
+  }
   hipDeviceptr_t device_ptr = nullptr;
   HIP_RETURN_ONFAIL(PlatformState::instance().getStatGlobalVar(symbol, ihipGetDevice(), &device_ptr, sizePtr));