From 7dba6dc28e4310ffba6a365a9c878f204d246a9e Mon Sep 17 00:00:00 2001
From: TomSang <Tao.Sang@amd.com>
Date: Tue, 14 Sep 2021 04:07:30 -0400
Subject: [PATCH] SWDEV-299960 - Fix .hipTestAtomicAdd test on NV (#2338)

1. Add --gpu-architecture=compute_60 to fully support AtomicAdd()
2. Add atomicAddNoRet() for NV.

Change-Id: I5add28ebabcfa21fb3ef99d09004c5d13423a290
---
 tests/src/deviceLib/hipTestAtomicAdd.cpp | 9 ++++++++-
 1 file changed, 8 insertions(+), 1 deletion(-)

diff --git a/tests/src/deviceLib/hipTestAtomicAdd.cpp b/tests/src/deviceLib/hipTestAtomicAdd.cpp
index f371423b73..9d03249e10 100644
--- a/tests/src/deviceLib/hipTestAtomicAdd.cpp
+++ b/tests/src/deviceLib/hipTestAtomicAdd.cpp
@@ -37,7 +37,7 @@ Testcase Scenarios :
 
 
 /* HIT_START
- * BUILD: %t %s ../test_common.cpp EXCLUDE_HIP_PLATFORM nvidia
+ * BUILD: %t %s ../test_common.cpp NVCC_OPTIONS --gpu-architecture=compute_60
  * TEST_NAMED: %t hipTestAtomicnoret-manywaves --atomicnoret --tests 1
  * TEST_NAMED: %t hipTestAtomicnoret-simple --atomicnoret --tests 2
  * TEST_NAMED: %t hipTestAtomic-manywaves --tests 1
@@ -59,6 +59,13 @@ Testcase Scenarios :
 #define LONG_INITIAL_VALUE 10000
 #define UNSIGNED_INITIAL_VALUE 20
 
+#ifdef __HIP_PLATFORM_NVIDIA__
+// atomicAddNoRet is unavailable in cuda
+template <typename T>
+__device__ void atomicAddNoRet(T* x, int y) {
+  atomicAdd(x, static_cast<T>(y));
+}
+#endif
 
 
 /*