From a36c3c6d8b4bf2fde6239148ea7a3f5f44fc8088 Mon Sep 17 00:00:00 2001
From: Phaneendr-kumar Lanka <planka@hcBLR-Fiji3.amd.com>
Date: Mon, 30 Oct 2017 14:34:24 +0530
Subject: [PATCH 1/5] [newTests]Adding tests for device APIs

[ROCm/hip commit: 511de63bcb0bc9976069864bcfab16e5d0bb5965]
---
 .../device/hipDeviceComputeCapability.cpp     | 46 ++++++++++++
 .../runtimeApi/device/hipDeviceGetName.cpp    | 45 +++++++++++
 .../device/hipDeviceSynchronize.cpp           | 74 +++++++++++++++++++
 .../runtimeApi/device/hipDeviceTotalMem.cpp   | 45 +++++++++++
 .../src/runtimeApi/device/hipGetDevice.cpp    | 44 +++++++++++
 .../device/hipRuntimeGetVersion.cpp           | 39 ++++++++++
 .../runtimeApi/device/hipSetCachceConfig.cpp  | 36 +++++++++
 .../src/runtimeApi/device/hipSetDevice.cpp    | 36 +++++++++
 .../runtimeApi/device/hipSetDeviceFlags.cpp   | 52 +++++++++++++
 9 files changed, 417 insertions(+)
 create mode 100644 projects/hip/tests/src/runtimeApi/device/hipDeviceComputeCapability.cpp
 create mode 100644 projects/hip/tests/src/runtimeApi/device/hipDeviceGetName.cpp
 create mode 100644 projects/hip/tests/src/runtimeApi/device/hipDeviceSynchronize.cpp
 create mode 100644 projects/hip/tests/src/runtimeApi/device/hipDeviceTotalMem.cpp
 create mode 100644 projects/hip/tests/src/runtimeApi/device/hipGetDevice.cpp
 create mode 100644 projects/hip/tests/src/runtimeApi/device/hipRuntimeGetVersion.cpp
 create mode 100644 projects/hip/tests/src/runtimeApi/device/hipSetCachceConfig.cpp
 create mode 100644 projects/hip/tests/src/runtimeApi/device/hipSetDevice.cpp
 create mode 100644 projects/hip/tests/src/runtimeApi/device/hipSetDeviceFlags.cpp

diff --git a/projects/hip/tests/src/runtimeApi/device/hipDeviceComputeCapability.cpp b/projects/hip/tests/src/runtimeApi/device/hipDeviceComputeCapability.cpp
new file mode 100644
index 0000000000..5c9af3d762
--- /dev/null
+++ b/projects/hip/tests/src/runtimeApi/device/hipDeviceComputeCapability.cpp
@@ -0,0 +1,46 @@
+/*
+Copyright (c) 2015-2017 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 WARRANNTY OF ANY KIND, EXPRESS OR
+IMPLIED, INNCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANNY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+/*
+ * Conformance test for checking functionality of
+ * hipError_t hipDeviceComputeCapability(int* major, int* minor, hipDevice_t device);
+ */
+
+/* HIT_START
+ * BUILD: %t %s ../../test_common.cpp
+ * RUN: %t
+ * HIT_END
+ */
+
+#include "test_common.h"
+
+int main()
+{
+    int numDevices = 0;
+    int major,minor;
+    hipDevice_t device;
+    HIPCHECK(hipGetDeviceCount(&numDevices));
+    for(int i=0;i<numDevices;i++){
+        HIPCHECK(hipDeviceGet(&device,i));
+        HIPCHECK(hipDeviceComputeCapability(&major, &minor, device));
+        HIPASSERT(major >= 0);
+        HIPASSERT(minor >= 0);
+    }
+    passed();
+}
diff --git a/projects/hip/tests/src/runtimeApi/device/hipDeviceGetName.cpp b/projects/hip/tests/src/runtimeApi/device/hipDeviceGetName.cpp
new file mode 100644
index 0000000000..94de10ff4c
--- /dev/null
+++ b/projects/hip/tests/src/runtimeApi/device/hipDeviceGetName.cpp
@@ -0,0 +1,45 @@
+/*
+Copyright (c) 2015-2017 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 WARRANNTY OF ANY KIND, EXPRESS OR
+IMPLIED, INNCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANNY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+/*
+ * Conformance test for checking functionality of
+ * hipError_t hipDeviceGetName(char* name, int len, hipDevice_t device);
+ */
+
+/* HIT_START
+ * BUILD: %t %s ../../test_common.cpp
+ * RUN: %t
+ * HIT_END
+ */
+
+#include "test_common.h"
+
+int main()
+{
+    int numDevices = 0;
+    char name[256];
+    hipDevice_t device;
+    HIPCHECK(hipGetDeviceCount(&numDevices));
+    for(int i=0;i<numDevices;i++){
+        HIPCHECK(hipDeviceGet(&device,i));
+        HIPCHECK(hipDeviceGetName(name,256,device));
+        HIPASSERT(name != "");
+    }
+    passed();
+}
diff --git a/projects/hip/tests/src/runtimeApi/device/hipDeviceSynchronize.cpp b/projects/hip/tests/src/runtimeApi/device/hipDeviceSynchronize.cpp
new file mode 100644
index 0000000000..dac56bf709
--- /dev/null
+++ b/projects/hip/tests/src/runtimeApi/device/hipDeviceSynchronize.cpp
@@ -0,0 +1,74 @@
+/*
+Copyright (c) 2015-2017 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 WARRANNTY OF ANY KIND, EXPRESS OR
+IMPLIED, INNCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANNY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+/*
+ * Test for checking the functionality of
+ * hipError_t hipDeviceSynchronize();
+*/
+
+/* HIT_START
+ * BUILD: %t %s ../../test_common.cpp
+ * RUN: %t
+ * HIT_END
+ */
+
+#include"test_common.h"
+
+#define _SIZE sizeof(int)*1024*1024
+#define NUM_STREAMS 2
+
+__global__ void Iter(hipLaunchParm lp, int *Ad, int num){
+    int tx = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x;
+    // Kernel loop designed to execute very slowly... ... ...   so we can test timing-related behavior below
+    if(tx == 0){
+        for(int i = 0; i<num;i++){
+            Ad[tx] += 1;
+        }
+    }
+}
+
+int main(){
+    int *A[NUM_STREAMS];
+    int *Ad[NUM_STREAMS];
+    hipStream_t stream[NUM_STREAMS];
+    for(int i=0;i<NUM_STREAMS;i++){
+        HIPCHECK(hipHostMalloc((void**)&A[i], _SIZE, hipHostMallocDefault));
+        A[i][0] = 1;
+        HIPCHECK(hipMalloc((void**)&Ad[i], _SIZE));
+        HIPCHECK(hipStreamCreate(&stream[i]));
+    }
+    for(int i=0;i<NUM_STREAMS;i++){
+        HIPCHECK(hipMemcpyAsync(Ad[i], A[i], _SIZE, hipMemcpyHostToDevice, stream[i]));
+    }
+    for(int i=0;i<NUM_STREAMS;i++){
+        hipLaunchKernel(HIP_KERNEL_NAME(Iter), dim3(1), dim3(1), 0, stream[i], Ad[i], 1<<30);
+    }
+    for(int i=0;i<NUM_STREAMS;i++){
+        HIPCHECK(hipMemcpyAsync(A[i], Ad[i], _SIZE, hipMemcpyDeviceToHost, stream[i]));
+    }
+
+
+   // This first check but relies on the kernel running for so long that the D2H async memcopy has not started yet.
+   // This will be true in an optimal asynchronous implementation.
+   // Conservative implementations which synchronize the hipMemcpyAsync will fail, ie if HIP_LAUNCH_BLOCKING=true
+    HIPASSERT(1<<30 != A[NUM_STREAMS-1][0]-1);
+    HIPCHECK(hipDeviceSynchronize());
+    HIPASSERT(1<<30 == A[NUM_STREAMS-1][0]-1);
+    passed();
+}
diff --git a/projects/hip/tests/src/runtimeApi/device/hipDeviceTotalMem.cpp b/projects/hip/tests/src/runtimeApi/device/hipDeviceTotalMem.cpp
new file mode 100644
index 0000000000..e05011e107
--- /dev/null
+++ b/projects/hip/tests/src/runtimeApi/device/hipDeviceTotalMem.cpp
@@ -0,0 +1,45 @@
+/*
+Copyright (c) 2015-2017 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 WARRANNTY OF ANY KIND, EXPRESS OR
+IMPLIED, INNCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANNY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+/*
+ * Conformance test for checking functionality of
+ * hipError_t hipDeviceGetName(char* name, int len, hipDevice_t device);
+ */
+
+/* HIT_START
+ * BUILD: %t %s ../../test_common.cpp
+ * RUN: %t
+ * HIT_END
+ */
+
+#include "test_common.h"
+
+int main()
+{
+    int numDevices = 0;
+    size_t totMem;
+    hipDevice_t device;
+    HIPCHECK(hipGetDeviceCount(&numDevices));
+    for(int i=0;i<numDevices;i++){
+        HIPCHECK(hipDeviceGet(&device,i));
+        HIPCHECK(hipDeviceTotalMem(&totMem,device));
+        HIPASSERT(totMem != 0);
+    }
+    passed();
+}
diff --git a/projects/hip/tests/src/runtimeApi/device/hipGetDevice.cpp b/projects/hip/tests/src/runtimeApi/device/hipGetDevice.cpp
new file mode 100644
index 0000000000..a268ec0acd
--- /dev/null
+++ b/projects/hip/tests/src/runtimeApi/device/hipGetDevice.cpp
@@ -0,0 +1,44 @@
+/*
+Copyright (c) 2015-2017 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 WARRANNTY OF ANY KIND, EXPRESS OR
+IMPLIED, INNCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANNY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+/*
+ * Conformance test for checking functionality of
+ * hipError_t hipGetDevice(int *device);
+ */
+
+/* HIT_START
+ * BUILD: %t %s ../../test_common.cpp
+ * RUN: %t
+ * HIT_END
+ */
+
+#include "test_common.h"
+
+int main()
+{
+    int numDevices = 0;
+    int device;
+    HIPCHECK(hipGetDeviceCount(&numDevices));
+    for(int i=0;i<numDevices;i++){
+        HIPCHECK(hipSetDevice(i));
+        HIPCHECK(hipGetDevice(&device));
+        HIPASSERT(device == i);
+    }
+    passed();
+}
diff --git a/projects/hip/tests/src/runtimeApi/device/hipRuntimeGetVersion.cpp b/projects/hip/tests/src/runtimeApi/device/hipRuntimeGetVersion.cpp
new file mode 100644
index 0000000000..61dceee3ac
--- /dev/null
+++ b/projects/hip/tests/src/runtimeApi/device/hipRuntimeGetVersion.cpp
@@ -0,0 +1,39 @@
+/*
+Copyright (c) 2015-2017 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 WARRANNTY OF ANY KIND, EXPRESS OR
+IMPLIED, INNCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANNY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+/*
+ * Conformance test for checking functionality of
+ * hipError_t hipRuntimeGetVersion(int* runtimeVersion);
+ * On HIP/HCC path this function returns HIP runtime patch version(a 5 digit code) however on HIP/NVCC path this function return CUDA runtime version.
+ */
+
+/* HIT_START
+ * BUILD: %t %s ../../test_common.cpp
+ * RUN: %t
+ * HIT_END
+ */
+
+#include "test_common.h"
+
+int main()
+{
+    int runtimeVersion;
+    HIPCHECK(hipRuntimeGetVersion(&runtimeVersion));
+    passed();
+}
diff --git a/projects/hip/tests/src/runtimeApi/device/hipSetCachceConfig.cpp b/projects/hip/tests/src/runtimeApi/device/hipSetCachceConfig.cpp
new file mode 100644
index 0000000000..e3c3efad3d
--- /dev/null
+++ b/projects/hip/tests/src/runtimeApi/device/hipSetCachceConfig.cpp
@@ -0,0 +1,36 @@
+/*
+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.
+*/
+
+/* HIT_START
+ * BUILD: %t %s ../../test_common.cpp
+ * RUN: %t
+ * HIT_END
+ */
+
+#include<hip/hip_runtime_api.h>
+#include<iostream>
+#include"test_common.h"
+
+int main(){
+  hipFuncCache_t cacheConfig;
+  void *func;
+  hipFuncSetCacheConfig(func, cacheConfig);
+  passed();
+}
+
diff --git a/projects/hip/tests/src/runtimeApi/device/hipSetDevice.cpp b/projects/hip/tests/src/runtimeApi/device/hipSetDevice.cpp
new file mode 100644
index 0000000000..b1b7cac12d
--- /dev/null
+++ b/projects/hip/tests/src/runtimeApi/device/hipSetDevice.cpp
@@ -0,0 +1,36 @@
+/*
+Copyright (c) 2015-2017 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 WARRANNTY OF ANY KIND, EXPRESS OR
+IMPLIED, INNCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANNY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+/* HIT_START
+ * BUILD: %t %s ../../test_common.cpp
+ * RUN: %t EXCLUDE_HIP_PLATFORM
+ * HIT_END
+ */
+
+#include "test_common.h"
+
+int main(){
+    int numDevices = 0;
+    HIPCHECK(hipGetDeviceCount(&numDevices));
+    for(int i=0;i<numDevices;i++){
+        HIPCHECK(hipSetDevice(i));
+    }
+    HIPASSERT(hipErrorInvalidDevice == hipSetDevice(numDevices));
+    passed();
+}
diff --git a/projects/hip/tests/src/runtimeApi/device/hipSetDeviceFlags.cpp b/projects/hip/tests/src/runtimeApi/device/hipSetDeviceFlags.cpp
new file mode 100644
index 0000000000..eacd4d2b63
--- /dev/null
+++ b/projects/hip/tests/src/runtimeApi/device/hipSetDeviceFlags.cpp
@@ -0,0 +1,52 @@
+/*
+Copyright (c) 2015-2017 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.
+*/
+
+/* HIT_START
+ * BUILD: %t %s ../../test_common.cpp
+ * RUN: %t
+ * HIT_END
+ */
+
+#include "test_common.h"
+
+int main()
+{
+    unsigned flag = 0;
+    HIPCHECK(hipDeviceReset());
+
+    int deviceCount = 0;
+    HIPCHECK(hipGetDeviceCount(&deviceCount));
+
+    for(int j=0;j<deviceCount;j++){
+
+        HIPCHECK(hipSetDevice(j));
+
+        for(int i=0;i<4;i++){
+            flag = 1 << i;
+            printf ("Flag=%x\n", flag);
+            HIPCHECK(hipSetDeviceFlags(flag));
+            //HIPCHECK_API(hipSetDeviceFlags(flag), hipErrorInvalidValue);
+        }
+
+        flag = 0;
+
+    }
+
+    passed();
+}

From 6c60773e82745613648242477d73be03949fdf47 Mon Sep 17 00:00:00 2001
From: Ben Sander <ben.sander@amd.com>
Date: Tue, 31 Oct 2017 22:26:26 +0000
Subject: [PATCH 2/5] Add HIP_DB=warn + message if sync on dangerous event.

[ROCm/hip commit: 70c25bdf8e295b013941e62790c4416a362ec932]
---
 projects/hip/src/hip_event.cpp      | 8 ++++++++
 projects/hip/src/hip_hcc_internal.h | 4 +++-
 2 files changed, 11 insertions(+), 1 deletion(-)

diff --git a/projects/hip/src/hip_event.cpp b/projects/hip/src/hip_event.cpp
index d1ee37a45e..93ef63f2e3 100644
--- a/projects/hip/src/hip_event.cpp
+++ b/projects/hip/src/hip_event.cpp
@@ -191,6 +191,10 @@ hipError_t hipEventSynchronize(hipEvent_t event)
 {
     HIP_INIT_SPECIAL_API(TRACE_SYNC, event);
 
+    if (!(event->_flags & hipEventReleaseToSystem)) {
+        tprintf(DB_WARN, "hipEventSynchronize on event without system-scope fence ; consider creating with hipEventReleaseToSystem\n");
+    }
+
     if (event) {
         if (event->_state == hipEventStatusUnitialized) {
             return ihipLogStatus(hipErrorInvalidResourceHandle);
@@ -268,6 +272,10 @@ hipError_t hipEventQuery(hipEvent_t event)
 {
     HIP_INIT_SPECIAL_API(TRACE_QUERY, event);
 
+    if (!(event->_flags & hipEventReleaseToSystem)) {
+        tprintf(DB_WARN, "hipEventQuery on event without system-scope fence ; consider creating with hipEventReleaseToSystem\n");
+    }
+
     if ((event->_state == hipEventStatusRecording) && !event->locked_isReady()) {
         return ihipLogStatus(hipErrorNotReady);
     } else {
diff --git a/projects/hip/src/hip_hcc_internal.h b/projects/hip/src/hip_hcc_internal.h
index 197cd35bfa..27c323887c 100644
--- a/projects/hip/src/hip_hcc_internal.h
+++ b/projects/hip/src/hip_hcc_internal.h
@@ -209,7 +209,8 @@ extern const char *API_COLOR_END;
 #define DB_SYNC   1 /* 0x02 - trace synchronization pieces */
 #define DB_MEM    2 /* 0x04 - trace memory allocation / deallocation */
 #define DB_COPY   3 /* 0x08 - trace memory copy and peer commands. . */
-#define DB_MAX_FLAG 4
+#define DB_WARN   4 /* 0x10 - warn about sub-optimal or shady behavior */
+#define DB_MAX_FLAG 5
 // When adding a new debug flag, also add to the char name table below.
 //
 //
@@ -226,6 +227,7 @@ static const DbName dbName [] =
     {KYEL, "sync"},
     {KCYN, "mem"},
     {KMAG, "copy"},
+    {KRED, "warn"},
 };
 
 

From 0026b0a3d9aa75e2b018c8a2be9dc8a1de91bc03 Mon Sep 17 00:00:00 2001
From: Phaneendr-kumar Lanka <planka@hcBLR-Fiji3.amd.com>
Date: Fri, 3 Nov 2017 10:39:47 +0530
Subject: [PATCH 3/5] [newTests]Modified the test as per review comments

[ROCm/hip commit: a764a0980dc0b45adbd3e9968b40a172e2a356f8]
---
 .../hip/tests/src/runtimeApi/device/hipDeviceGetName.cpp    | 6 ++++--
 1 file changed, 4 insertions(+), 2 deletions(-)

diff --git a/projects/hip/tests/src/runtimeApi/device/hipDeviceGetName.cpp b/projects/hip/tests/src/runtimeApi/device/hipDeviceGetName.cpp
index 94de10ff4c..0a3dce284e 100644
--- a/projects/hip/tests/src/runtimeApi/device/hipDeviceGetName.cpp
+++ b/projects/hip/tests/src/runtimeApi/device/hipDeviceGetName.cpp
@@ -30,15 +30,17 @@ THE SOFTWARE.
 
 #include "test_common.h"
 
+const int len 256;
+
 int main()
 {
     int numDevices = 0;
-    char name[256];
+    char name[len];
     hipDevice_t device;
     HIPCHECK(hipGetDeviceCount(&numDevices));
     for(int i=0;i<numDevices;i++){
         HIPCHECK(hipDeviceGet(&device,i));
-        HIPCHECK(hipDeviceGetName(name,256,device));
+        HIPCHECK(hipDeviceGetName(name,len,device));
         HIPASSERT(name != "");
     }
     passed();

From 9d7c904529a312bc5a8f15513d631254aa0c2cff Mon Sep 17 00:00:00 2001
From: Phaneendr-kumar Lanka <planka@hcBLR-Fiji3.amd.com>
Date: Fri, 3 Nov 2017 14:52:19 +0530
Subject: [PATCH 4/5] [newTests]modified the test

[ROCm/hip commit: 0599d13b0db508e612e1c12151bbbe676afa99f2]
---
 projects/hip/tests/src/runtimeApi/device/hipDeviceGetName.cpp | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/projects/hip/tests/src/runtimeApi/device/hipDeviceGetName.cpp b/projects/hip/tests/src/runtimeApi/device/hipDeviceGetName.cpp
index 0a3dce284e..ba260ac2e5 100644
--- a/projects/hip/tests/src/runtimeApi/device/hipDeviceGetName.cpp
+++ b/projects/hip/tests/src/runtimeApi/device/hipDeviceGetName.cpp
@@ -30,7 +30,7 @@ THE SOFTWARE.
 
 #include "test_common.h"
 
-const int len 256;
+#define len 256
 
 int main()
 {

From 6a4c50cf6c7bc5cc2d162968e6087b41879debf4 Mon Sep 17 00:00:00 2001
From: Ben Sander <ben.sander@amd.com>
Date: Sun, 5 Nov 2017 10:32:24 -0600
Subject: [PATCH 5/5] Set event state AFTER it is recorded.

[ROCm/hip commit: 4c3b65a5cd29451aed08d1f82e2564ddd10ee795]
---
 projects/hip/src/hip_event.cpp  | 5 ++---
 projects/hip/src/hip_stream.cpp | 2 +-
 2 files changed, 3 insertions(+), 4 deletions(-)

diff --git a/projects/hip/src/hip_event.cpp b/projects/hip/src/hip_event.cpp
index d1ee37a45e..91c3f9cb55 100644
--- a/projects/hip/src/hip_event.cpp
+++ b/projects/hip/src/hip_event.cpp
@@ -45,10 +45,10 @@ ihipEvent_t::ihipEvent_t(unsigned flags)
 void ihipEvent_t::attachToCompletionFuture(const hc::completion_future *cf, 
                                            hipStream_t stream, ihipEventType_t eventType)
 {
-    _state  = hipEventStatusRecording;
     _marker = *cf;
     _type   = eventType;
     _stream = stream;
+    _state  = hipEventStatusRecording;
 }
 
 
@@ -157,13 +157,12 @@ hipError_t hipEventRecord(hipEvent_t event, hipStream_t stream)
             event->_state = hipEventStatusComplete;
             return ihipLogStatus(hipSuccess);
         } else {
-            event->_state  = hipEventStatusRecording;
             // Clear timestamps
             event->_timestamp = 0;
 
             // Record the event in the stream:
             stream->locked_recordEvent(event);
-
+            event->_state  = hipEventStatusRecording;
             return ihipLogStatus(hipSuccess);
         }
     } else {
diff --git a/projects/hip/src/hip_stream.cpp b/projects/hip/src/hip_stream.cpp
index f62be31b9e..0cb5e732ab 100644
--- a/projects/hip/src/hip_stream.cpp
+++ b/projects/hip/src/hip_stream.cpp
@@ -103,7 +103,7 @@ hipError_t hipStreamWaitEvent(hipStream_t stream, hipEvent_t event, unsigned int
             event->locked_waitComplete((event->_flags & hipEventBlockingSync) ? hc::hcWaitModeBlocked : hc::hcWaitModeActive);
         } else {
             stream = ihipSyncAndResolveStream(stream);
-            // This will user create_blocking_marker to wait on the specified queue.
+            // This will use create_blocking_marker to wait on the specified queue.
             stream->locked_streamWaitEvent(event);
         }