SWDEV-299127 - Merge 'develop' into 'amd-staging'

Change-Id: Id044a4d29629d4cad140bc3e65b301c9f361538a [ROCm/hip commit: c7892daa31]
2021-09-14 19:11:37 -04:00
@@ -271,3 +271,5 @@ HIP version can be queried from HIP API call,
 hipRuntimeGetVersion(&runtimeVersion);

 The version returned will always be greater than the versions in previous ROCm releases.
+
+Note: The version definition of HIP runtime is different from CUDA. On AMD platform, the function returns HIP runtime version, while on NVIDIA platform, it returns CUDA runtime version. And there is no mapping/correlation between HIP version and CUDA version.
@@ -734,6 +734,10 @@ __launch_bounds__ supports two parameters:
 The threads-per-block is the product of (hipBlockDim_x * hipBlockDim_y * hipBlockDim_z).
 - MIN_WARPS_PER_EU - directs the compiler to minimize resource usage so that the requested number of warps can be simultaneously active on a multi-processor. Since active warps compete for the same fixed pool of resources, the compiler must reduce resources required by each warp(primarily registers). MIN_WARPS_PER_EU is optional and defaults to 1 if not specified. Specifying a MIN_WARPS_PER_EU greater than the default 1 effectively constrains the compiler's resource usage.

+When launch kernel with HIP APIs, for example, hipModuleLaunchKernel(), HIP will do validation to make sure input kernel dimension size is not larger than specified launch_bounds.
+In case exceeded, HIP would return launch failure, if AMD_LOG_LEVEL is set with proper value (for details, please refer to docs/markdown/hip_logging.md), detail information will be shown in the error log message, including
+launch parameters of kernel dim size, launch bounds, and the name of the faulting kernel. It's helpful to figure out which is the faulting kernel, besides, the kernel dim size and launch bounds values will also assist in debugging such failures.
+
 ### Compiler Impact
 The compiler uses these parameters as follows:
 - The compiler uses the hints only to manage register usage, and does not automatically reduce shared memory or other resources.
@@ -913,8 +913,10 @@ hipError_t hipDriverGetVersion(int* driverVersion);
 *
 * @returns #hipSuccess, #hipErrorInavlidValue
 *
- * @warning On HIP/HCC path this function returns HIP runtime patch version however on
- * HIP/NVCC path this function return CUDA runtime version.
+ * @warning The version definition of HIP runtime is different from CUDA.
+ * On AMD platform, the function returns HIP runtime version,
+ * while on NVIDIA platform, it returns CUDA runtime version.
+ * And there is no mapping/correlation between HIP version and CUDA version.
 *
 * @see hipDriverGetVersion
 */
@@ -1,9 +1,27 @@
+if(HIP_PLATFORM MATCHES "amd")
 set(TEST_SRC
    hipStreamCreate.cc
    hipStreamGetFlags.cc
    hipStreamGetPriority.cc
    hipMultiStream.cc
+    hipStreamACb_MultiThread.cc
+    hipStreamAddCallback.cc
+    hipStreamCreateWithFlags.cc
+    hipStreamCreateWithPriority.cc
+    hipStreamWithCUMask.cc
 )
+else()
+set(TEST_SRC
+    hipStreamCreate.cc
+    hipStreamGetFlags.cc
+    hipStreamGetPriority.cc
+    hipMultiStream.cc
+    hipStreamACb_MultiThread.cc
+    hipStreamAddCallback.cc
+    hipStreamCreateWithFlags.cc
+    hipStreamCreateWithPriority.cc
+)
+endif()

 # Create shared lib of all tests
 add_library(StreamTest SHARED EXCLUDE_FROM_ALL ${TEST_SRC})
@@ -0,0 +1,168 @@
+/*
+Copyright (c) 2021 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.
+*/
+
+/**
+Testcase Scenario :
+Validate behaviour of HIP when multiple hipStreaAddCallback() are called over
+multiple Threads.
+*/
+
+#include <hip_test_common.hh>
+#include <atomic>
+
+static constexpr size_t N = 4096;
+static constexpr int numThreads = 1000;
+static std::atomic<int> Cb_count{0}, Data_mismatch{0};
+static hipStream_t mystream;
+static float *A1_h, *C1_h;
+
+#if HT_AMD
+#define HIPRT_CB
+#endif
+
+static __global__ void device_function(float* C_d, float* A_d, size_t Num) {
+  size_t gputhread = (blockIdx.x * blockDim.x + threadIdx.x);
+  size_t stride = blockDim.x * gridDim.x;
+
+  for (size_t i = gputhread; i < Num; i += stride) {
+    C_d[i] = A_d[i] * A_d[i];
+  }
+
+  // Delay thread 1 only in the GPU
+  if (gputhread == 1) {
+    uint64_t wait_t = 3200000000, start = clock64(), cur;
+    do {
+      cur = clock64() - start;
+    } while (cur < wait_t);
+  }
+}
+
+
+static void HIPRT_CB Thread1_Callback(hipStream_t stream, hipError_t status,
+                                      void* userData) {
+  HIPASSERT(stream == mystream);
+  HIPASSERT(userData == nullptr);
+  HIPCHECK(status);
+
+  for (size_t i = 0; i < N; i++) {
+    // Validate the data and update Data_mismatch
+    if (C1_h[i] != A1_h[i] * A1_h[i]) {
+      Data_mismatch++;
+    }
+  }
+
+  // Increment the Cb_count to indicate that the callback is processed.
+  ++Cb_count;
+}
+
+static void HIPRT_CB Thread2_Callback(hipStream_t stream, hipError_t status,
+                                      void* userData) {
+  HIPASSERT(stream == mystream);
+  HIPASSERT(userData == nullptr);
+  HIPCHECK(status);
+
+  for (size_t i = 0; i < N; i++) {
+    // Validate the data and update Data_mismatch
+    if (C1_h[i] != A1_h[i] * A1_h[i]) {
+      Data_mismatch++;
+    }
+  }
+
+  // Increment the Cb_count to indicate that the callback is processed.
+  ++Cb_count;
+}
+
+void Thread1_func() {
+  HIPCHECK(hipStreamAddCallback(mystream, Thread1_Callback, nullptr, 0));
+}
+
+void Thread2_func() {
+  HIPCHECK(hipStreamAddCallback(mystream, Thread2_Callback, nullptr, 0));
+}
+
+/**
+ Test multiple hipStreamAddCallback() called over
+ multiple Threads.
+ */
+TEST_CASE("Unit_hipStreamAddCallback_MultipleThreads") {
+  float *A_d, *C_d;
+  size_t Nbytes = (N) * sizeof(float);
+  constexpr float Phi = 1.618f;
+
+  A1_h = reinterpret_cast<float*>(malloc(Nbytes));
+  REQUIRE(A1_h != nullptr);
+  C1_h = reinterpret_cast<float*>(malloc(Nbytes));
+  REQUIRE(C1_h != nullptr);
+
+  // Fill with Phi + i
+  for (size_t i = 0; i < N; i++) {
+    A1_h[i] = Phi + i;
+  }
+
+  HIP_CHECK(hipMalloc(&A_d, Nbytes));
+  HIP_CHECK(hipMalloc(&C_d, Nbytes));
+
+  HIP_CHECK(
+  hipStreamCreateWithFlags(&mystream, hipStreamNonBlocking));
+
+  HIP_CHECK(
+  hipMemcpyAsync(A_d, A1_h, Nbytes, hipMemcpyHostToDevice,
+                 mystream));
+
+  constexpr unsigned threadsPerBlock = 256;
+  constexpr unsigned blocks = (N + 255)/threadsPerBlock;
+
+  hipLaunchKernelGGL((device_function), dim3(blocks),
+                     dim3(threadsPerBlock), 0,
+                     mystream, C_d, A_d, N);
+
+  HIP_CHECK(
+  hipMemcpyAsync(C1_h, C_d, Nbytes,
+                 hipMemcpyDeviceToHost, mystream));
+
+  std::thread *T = new std::thread[numThreads];
+  for (int i = 0; i < numThreads; i++) {
+    // Use different callback for every even thread
+    // The callbacks will be added to same stream from different threads
+    if ((i%2) == 0)
+      T[i] = std::thread(Thread1_func);
+    else
+      T[i] = std::thread(Thread2_func);
+  }
+
+  // Wait until all the threads finish their execution
+  for (int i = 0; i < numThreads; i++) {
+    T[i].join();
+  }
+
+  HIP_CHECK(hipStreamSynchronize(mystream));
+  HIP_CHECK(hipStreamDestroy(mystream));
+
+  HIP_CHECK(hipFree(A_d));
+  HIP_CHECK(hipFree(C_d));
+
+  free(A1_h);
+  free(C1_h);
+
+  // Cb_count should match total number of callbacks added from both threads
+  // Data_mismatch will be updated if there is problem in data validation
+  REQUIRE(Cb_count.load() == numThreads);
+  REQUIRE(Data_mismatch.load() == 0);
+  delete[] T;
+}
@@ -0,0 +1,231 @@
+/*
+Copyright (c) 2021 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.
+*/
+
+/**
+Testcase Scenarios :
+ 1) Validates parameter list of hipStreamAddCallback.
+ 2) Validates hipStreamAddCallback functionality with default stream.
+ 3) Validates hipStreamAddCallback functionality with defined stream.
+*/
+
+#include <hip_test_common.hh>
+#include <hip_test_kernels.hh>
+#include <unistd.h>
+
+#define UNUSED(expr) do { (void)(expr); } while (0)
+
+#ifdef __HIP_PLATFORM_AMD__
+#define HIPRT_CB
+#endif
+
+namespace hipStreaAddCallbackTest {
+size_t NSize = 4 * 1024 * 1024;
+float *A_h, *C_h;
+bool gcbDone = false;
+bool gPassed = true;
+void *ptr0xff = reinterpret_cast<void *>(0xffffffff);
+void *gusrptr;
+hipStream_t gstream;
+
+void HIPRT_CB Callback(hipStream_t stream, hipError_t status,
+                       void* userData) {
+  UNUSED(stream);
+  HIP_CHECK(status);
+  REQUIRE(userData == NULL);
+  gPassed = true;
+  for (size_t i = 0; i < NSize; i++) {
+    if (C_h[i] != A_h[i] * A_h[i]) {
+      UNSCOPED_INFO("Data mismatch :" << i);
+      gPassed = false;
+      break;
+    }
+  }
+  gcbDone = true;
+}
+/**
+ * Validates functionality of hipStreamAddCallback with default/created stream.
+ */
+bool testStreamCallbackFunctionality(bool isDefault) {
+  float *A_d, *C_d;
+  size_t Nbytes = NSize * sizeof(float);
+  gcbDone = false;
+  A_h = reinterpret_cast<float*>(malloc(Nbytes));
+  REQUIRE(A_h != nullptr);
+  C_h = reinterpret_cast<float*>(malloc(Nbytes));
+  REQUIRE(C_h != nullptr);
+
+  // Fill with Phi + i
+  for (size_t i = 0; i < NSize; i++) {
+    A_h[i] = 1.618f + i;
+  }
+
+  HIP_CHECK(hipMalloc(&A_d, Nbytes));
+  HIP_CHECK(hipMalloc(&C_d, Nbytes));
+  if (isDefault) {
+    HIP_CHECK(hipMemcpyAsync(A_d, A_h, Nbytes, hipMemcpyHostToDevice,
+                            0));
+
+    const unsigned blocks = 512;
+    const unsigned threadsPerBlock = 256;
+    hipLaunchKernelGGL((HipTest::vector_square), dim3(blocks),
+                        dim3(threadsPerBlock), 0, 0, A_d, C_d, NSize);
+
+    HIP_CHECK(hipMemcpyAsync(C_h, C_d, Nbytes, hipMemcpyDeviceToHost,
+                            0));
+    HIP_CHECK(hipStreamAddCallback(0, Callback, nullptr, 0));
+    while (!gcbDone) usleep(100000);  // Sleep for 100 ms
+  } else {
+    hipStream_t mystream;
+    HIP_CHECK(hipStreamCreateWithFlags(&mystream, hipStreamNonBlocking));
+
+    HIP_CHECK(hipMemcpyAsync(A_d, A_h, Nbytes, hipMemcpyHostToDevice,
+                            mystream));
+
+    const unsigned blocks = 512;
+    const unsigned threadsPerBlock = 256;
+    hipLaunchKernelGGL((HipTest::vector_square), dim3(blocks),
+                      dim3(threadsPerBlock), 0, mystream, A_d, C_d, NSize);
+
+    HIP_CHECK(hipMemcpyAsync(C_h, C_d, Nbytes, hipMemcpyDeviceToHost,
+                            mystream));
+    HIP_CHECK(hipStreamAddCallback(mystream, Callback, nullptr, 0));
+    while (!gcbDone) usleep(100000);  // Sleep for 100 ms
+    HIP_CHECK(hipStreamDestroy(mystream));
+  }
+  HIP_CHECK(hipFree(reinterpret_cast<void*>(C_d)));
+  HIP_CHECK(hipFree(reinterpret_cast<void*>(A_d)));
+  free(C_h);
+  free(A_h);
+  return gPassed;
+}
+/**
+ * Scenario1: Validates if callback = nullptr returns error code for created stream.
+ * Scenario2: Validates if callback = nullptr returns error code for default stream.
+ * Scenario3: Validates if flag != 0 returns error code for created stream.
+ * Scenario4: Validates if flag != 0 returns error code for default stream.
+ * Scenario5: Validates if userData pointer is passed properly to callback.
+ * Scenario6: Validates if stream value is passed properly to callback.
+ */
+void Callback_ChkUsrdataPtr(hipStream_t stream, hipError_t status,
+                            void* userData) {
+  REQUIRE(stream == gstream);
+  HIP_CHECK(status);
+  gPassed = true;
+  if (gusrptr != userData) {
+    gPassed = false;
+  }
+  gcbDone = true;
+}
+
+void Callback_ChkStreamValue(hipStream_t stream, hipError_t status,
+                             void* userData) {
+  REQUIRE(userData == nullptr);
+  HIP_CHECK(status);
+  gPassed = true;
+  if (stream != gstream) {
+    gPassed = false;
+  }
+  gcbDone = true;
+}
+}  // namespace hipStreaAddCallbackTest
+
+
+using hipStreaAddCallbackTest::gcbDone;
+using hipStreaAddCallbackTest::gPassed;
+using hipStreaAddCallbackTest::ptr0xff;
+using hipStreaAddCallbackTest::gusrptr;
+using hipStreaAddCallbackTest::gstream;
+using hipStreaAddCallbackTest::testStreamCallbackFunctionality;
+using hipStreaAddCallbackTest::Callback;
+using hipStreaAddCallbackTest::Callback_ChkUsrdataPtr;
+using hipStreaAddCallbackTest::Callback_ChkStreamValue;
+
+
+/*
+ * Validates parameter list of hipStreamAddCallback.
+ */
+TEST_CASE("Unit_hipStreamAddCallback_ParamTst") {
+  hipStream_t mystream;
+  HIP_CHECK(hipStreamCreate(&mystream));
+  // Scenario1
+  SECTION("callback is nullptr") {
+    REQUIRE_FALSE(hipSuccess == hipStreamAddCallback(mystream, nullptr,
+                                         nullptr, 0));
+  }
+  // Scenario2
+  SECTION("stream is default") {
+    REQUIRE_FALSE(hipSuccess == hipStreamAddCallback(0, nullptr,
+                                         nullptr, 0));
+  }
+  // Scenario3
+  SECTION("flag is nonzero for non-default stream") {
+    REQUIRE_FALSE(hipSuccess == hipStreamAddCallback(mystream, Callback,
+                                         nullptr, 10));
+  }
+  // Scenario4
+  SECTION("flag is nonzero for default stream") {
+    REQUIRE_FALSE(hipSuccess == hipStreamAddCallback(0, Callback,
+                                         nullptr, 10));
+  }
+  // Scenario5
+  SECTION("userData pointer value validation") {
+    gstream = mystream;
+    gusrptr = ptr0xff;
+    gPassed = true;
+    gcbDone = false;
+    HIP_CHECK(hipStreamAddCallback(mystream, Callback_ChkUsrdataPtr,
+                                  gusrptr, 0));
+    while (!gcbDone) {
+      usleep(100000);  // Sleep for 100 ms
+    }
+    REQUIRE_FALSE(!gPassed);
+  }
+  // Scenario6
+  SECTION("stream value validation") {
+    gstream = mystream;
+    gPassed = true;
+    gcbDone = false;
+    HIP_CHECK(hipStreamAddCallback(mystream, Callback_ChkStreamValue,
+                                  nullptr, 0));
+    while (!gcbDone) {
+      usleep(100000);  // Sleep for 100 ms
+    }
+    REQUIRE_FALSE(!gPassed);
+  }
+  HIP_CHECK(hipStreamDestroy(mystream));
+}
+
+/*
+ * Validates hipStreamAddCallback functionality with default stream.
+ */
+TEST_CASE("Unit_hipStreamAddCallback_WithDefaultStream") {
+  bool TestPassed = true;
+  TestPassed = testStreamCallbackFunctionality(true);
+  REQUIRE(TestPassed);
+}
+
+/*
+ * Validates hipStreamAddCallback functionality with defined stream.
+ */
+TEST_CASE("Unit_hipStreamAddCallback_WithCreatedStream") {
+  bool TestPassed = true;
+  TestPassed = testStreamCallbackFunctionality(false);
+  REQUIRE(TestPassed);
+}
+
@@ -0,0 +1,39 @@
+/*
+Copyright (c) 2021 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.
+*/
+
+/**
+Testcase Scenarios :
+1) Validates functionality of hipStreamCreateWithFlags when stream = nullptr.
+2) Validates functionality of hipStreamCreateWithFlags when flag = 0xffffffff.
+*/
+
+#include <hip_test_common.hh>
+
+
+TEST_CASE("Unit_hipStreamCreateWithFlags_ArgValidation") {
+  // stream = nullptr test
+  SECTION("stream is nullptr") {
+    REQUIRE(hipStreamCreateWithFlags(nullptr, hipStreamDefault) != hipSuccess);
+  }
+  // flag value invalid test
+  SECTION("flag value invalid") {
+    hipStream_t stream;
+    REQUIRE(hipStreamCreateWithFlags(&stream, 0xffffffff) != hipSuccess);
+  }
+}
@@ -0,0 +1,605 @@
+/*
+Copyright (c) 2021 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.
+*/
+
+/**
+Testcase Scenarios :
+
+1)Create streams with default flag for all available priority levels and
+queue tasks in each of these streams, perform device synchronize and validate behavior.
+
+2)Create streams with non-blocking flag for all available priority levels and
+queue tasks in each of these streams, perform stream synchronize and validate behavior.
+
+3)Create streams with default flag for all available priority levels and
+queue tasks in each of these streams, perform stream synchronize and validate behavior.
+
+4)Create streams with non-blocking flag for all available priority levels and
+queue tasks in each of these streams, perform device synchronize and validate behavior.
+
+5)Create a stream for each priority level with default flag, Launch memcpy and kernel
+tasks on these streams from multiple threads. Validate all the results.
+
+6)Create a stream for each priority level with non-blocking flag, Launch memcpy and
+kernel tasks on these streams from multiple threads. Validate all the results.
+
+7) Validate negative scenarios for hipStreamCreateWithPriority api.
+
+8) Validate stream priorities with event after classifying them as low, medium, high.
+*/
+
+#include <hip_test_common.hh>
+#include <hip_test_kernels.hh>
+#include <atomic>
+#include <vector>
+
+#define MEMCPYSIZE 64*1024*1024
+#define MEMCPYSIZE2 1024*1024
+#define NUMITERS   2
+#define GRIDSIZE   1024
+#define BLOCKSIZE  256
+#define TOTALTHREADS 16
+
+namespace hipStreamCreateWithPriorityTest {
+
+std::atomic<int> g_thTestPassed(1);
+// helper rountine to initialize memory
+template <typename T>
+void mem_init(T* buf, size_t n) {
+  for (size_t i = 0; i < n; i++) {
+    buf[i] = i;
+  }
+}
+
+// kernel to copy n elements from src to dst
+template <typename T>
+__global__ void memcpy_kernel(T* dst, T* src, size_t n) {
+  int num = gridDim.x * blockDim.x;
+  int id = blockDim.x * blockIdx.x + threadIdx.x;
+
+  for (size_t i = id; i < n; i += num) {
+    dst[i] = src[i];
+  }
+}
+
+/**
+ * Scenario: Create a stream for all available priority levels
+ * and queue tasks in each of these streams and default stream.
+ * Validate the calculated results.
+ */
+void funcTestsForAllPriorityLevelsWrtNullStrm(unsigned int flags,
+                                        bool deviceSynchronize) {
+  int priority;
+  int priority_low{};
+  int priority_high{};
+  size_t size = MEMCPYSIZE2*sizeof(int);
+  // Test is to get the Stream Priority Range
+  HIP_CHECK(hipDeviceGetStreamPriorityRange(&priority_low, &priority_high));
+
+  // Check if priorities are indeed supported
+  if (priority_low == priority_high) {
+    WARN("Stream priority range not supported. Skipping test.");
+    return;
+  }
+
+  int numOfPriorities = priority_low - priority_high;
+  INFO("numOfPriorities = " << numOfPriorities);
+  const int arr_size = numOfPriorities + 1;
+  // 0 idx is for default stream
+  hipStream_t *stream = reinterpret_cast<hipStream_t*>(
+                       malloc(arr_size*sizeof(hipStream_t)));
+  REQUIRE(stream != nullptr);
+  stream[0] = 0;
+  int count = 1;
+  // Create a stream for each of the priority levels
+  for (priority = priority_high; priority < priority_low; priority++) {
+    HIP_CHECK(hipStreamCreateWithPriority(&stream[count++],
+            flags, priority));
+  }
+  // Allocate memory
+  int **A_d = reinterpret_cast<int**>(malloc(arr_size*sizeof(int*)));
+  int **C_d = reinterpret_cast<int**>(malloc(arr_size*sizeof(int*)));
+  int **A_h = reinterpret_cast<int**>(malloc(arr_size*sizeof(int*)));
+  int **C_h = reinterpret_cast<int**>(malloc(arr_size*sizeof(int*)));
+
+  REQUIRE(A_d != nullptr);
+  REQUIRE(C_d != nullptr);
+  REQUIRE(A_h != nullptr);
+  REQUIRE(C_h != nullptr);
+
+  for (int idx = 0; idx < arr_size; idx++) {
+    A_h[idx] = reinterpret_cast<int*>(malloc(size));
+    REQUIRE(A_h[idx] != nullptr);
+    C_h[idx] = reinterpret_cast<int*>(malloc(size));
+    REQUIRE(C_h[idx] != nullptr);
+    HIP_CHECK(hipMalloc(&A_d[idx], size));
+    HIP_CHECK(hipMalloc(&C_d[idx], size));
+  }
+
+  // Initialize host memory
+  constexpr int initVal = 2;
+  for (int idx = 0; idx < arr_size; idx++) {
+    for (int idy = 0; idy < MEMCPYSIZE2; idy++) {
+      A_h[idx][idy] = initVal;
+    }
+  }
+
+  // Launch task on each stream
+  for (int idx = 0; idx < arr_size; idx++) {
+    HIP_CHECK(hipMemcpyAsync(A_d[idx], A_h[idx], size,
+            hipMemcpyHostToDevice, stream[idx]));
+    hipLaunchKernelGGL((HipTest::vector_square), dim3(GRIDSIZE),
+                        dim3(BLOCKSIZE), 0, stream[idx], A_d[idx],
+                        C_d[idx], MEMCPYSIZE2);
+    HIP_CHECK(hipMemcpyAsync(C_h[idx], C_d[idx], size,
+            hipMemcpyDeviceToHost, stream[idx]));
+  }
+
+  if (deviceSynchronize) {
+    HIP_CHECK(hipDeviceSynchronize());
+  }
+
+  // Validate the output of each queue
+  for (int idx = 0; idx < arr_size; idx++) {
+    if (!deviceSynchronize) {
+      HIP_CHECK(hipStreamSynchronize(stream[idx]));
+    }
+    for (int idy = 0; idy < MEMCPYSIZE2; idy++) {
+      if (C_h[idx][idy] != A_h[idx][idy] * A_h[idx][idy]) {
+        INFO("Data mismatch at idx:" << idx << " idy:" << idy);
+        REQUIRE(false);
+      }
+    }
+  }
+
+  // Deallocate memory
+  for (int idx = 0; idx < arr_size; idx++) {
+    HIP_CHECK(hipFree(reinterpret_cast<void*>(C_d[idx])));
+    HIP_CHECK(hipFree(reinterpret_cast<void*>(A_d[idx])));
+    free(C_h[idx]);
+    free(A_h[idx]);
+  }
+
+  // Destroy the stream for each of the priority levels
+  count = 1;
+  for (priority = priority_high; priority < priority_low; priority++) {
+    HIP_CHECK(hipStreamDestroy(stream[count++]));
+  }
+  free(stream);
+  free(A_d);
+  free(C_d);
+  free(A_h);
+  free(C_h);
+}
+
+/**
+ * Scenario: Queue tasks in each of these streams and default stream.
+ * Validate the calculated results.
+ */
+void queueTasksInStreams(std::vector<hipStream_t> *stream,
+                         const int arrsize) {
+  size_t size = MEMCPYSIZE2*sizeof(int);
+  // Allocate memory
+  int **A_d = reinterpret_cast<int**>(malloc(arrsize*sizeof(int *)));
+  int **C_d = reinterpret_cast<int**>(malloc(arrsize*sizeof(int *)));
+  int **A_h = reinterpret_cast<int**>(malloc(arrsize*sizeof(int *)));
+  int **C_h = reinterpret_cast<int**>(malloc(arrsize*sizeof(int *)));
+
+  HIPASSERT(A_d != nullptr);
+  HIPASSERT(C_d != nullptr);
+  HIPASSERT(A_h != nullptr);
+  HIPASSERT(C_h != nullptr);
+
+  for (int idx = 0; idx < arrsize; idx++) {
+    A_h[idx] = reinterpret_cast<int*>(malloc(size));
+    HIPASSERT(A_h[idx] != nullptr);
+    C_h[idx] = reinterpret_cast<int*>(malloc(size));
+    HIPASSERT(C_h[idx] != nullptr);
+    HIPCHECK(hipMalloc(&A_d[idx], size));
+    HIPCHECK(hipMalloc(&C_d[idx], size));
+  }
+  // Initialize host memory
+  constexpr int initVal = 2;
+  for (int idx = 0; idx < arrsize; idx++) {
+    for (int idy = 0; idy < MEMCPYSIZE2; idy++) {
+      A_h[idx][idy] = initVal;
+    }
+  }
+  // Launch task on each stream
+  for (int idx = 0; idx < arrsize; idx++) {
+    HIPCHECK(hipMemcpyAsync(A_d[idx], A_h[idx], size,
+             hipMemcpyHostToDevice, (*stream)[idx]));
+    hipLaunchKernelGGL((HipTest::vector_square), dim3(GRIDSIZE),
+                        dim3(BLOCKSIZE), 0, (*stream)[idx], A_d[idx],
+                        C_d[idx], MEMCPYSIZE2);
+    HIPCHECK(hipMemcpyAsync(C_h[idx], C_d[idx], size,
+             hipMemcpyDeviceToHost, (*stream)[idx]));
+  }
+
+  bool isPassed = true;
+  // Validate the output of each queue
+  for (int idx = 0; idx < arrsize; idx++) {
+    HIPCHECK(hipStreamSynchronize((*stream)[idx]));
+    for (int idy = 0; idy < MEMCPYSIZE2; idy++) {
+      if (C_h[idx][idy] != A_h[idx][idy] * A_h[idx][idy]) {
+        UNSCOPED_INFO("Data mismatch at idx:" << idx << " idy:" << idy);
+        isPassed = false;
+        break;
+      }
+    }
+    if (false == isPassed) break;
+  }
+  // Deallocate memory
+  for (int idx = 0; idx < arrsize; idx++) {
+    HIPCHECK(hipFree(reinterpret_cast<void*>(C_d[idx])));
+    HIPCHECK(hipFree(reinterpret_cast<void*>(A_d[idx])));
+    free(C_h[idx]);
+    free(A_h[idx]);
+  }
+  free(A_d);
+  free(C_d);
+  free(A_h);
+  free(C_h);
+  g_thTestPassed &= static_cast<int>(isPassed);
+}
+
+/**
+ * Scenario:
+ * Common streams used across multiple threads:Create a stream for each
+ * priority level (flag = hipStreamDefault/hipStreamNonBlocking)
+ * and 1 default stream.
+ * Launch memcpy and kernel tasks on these streams from multiple threads
+ * (use 16 threads). Validate all the results.
+*/
+bool runFuncTestsForAllPriorityLevelsMultThread(unsigned int flags) {
+  bool TestPassed = true;
+  std::thread T[TOTALTHREADS];
+  int priority;
+  int priority_low;
+  int priority_high;
+  // Test is to get the Stream Priority Range
+  HIP_CHECK(hipDeviceGetStreamPriorityRange(&priority_low, &priority_high));
+
+  // Check if priorities are indeed supported
+  if (priority_low == priority_high) {
+    WARN("Stream priority range not supported. Skipping test.");
+    return true;
+  }
+
+  int numOfPriorities = priority_low - priority_high;
+  INFO("numOfPriorities : " << numOfPriorities);
+
+  // 0 idx is for default stream
+  std::vector<hipStream_t> stream(numOfPriorities + 1);
+  stream[0] = 0;
+
+  // Create a stream for each of the priority levels
+  int count = 1;
+  for (priority = priority_high; priority < priority_low; priority++) {
+    HIP_CHECK(hipStreamCreateWithPriority(&stream[count++], flags,
+                                         priority));
+  }
+
+  for (int i = 0; i < TOTALTHREADS; i++) {
+    T[i] = std::thread(queueTasksInStreams,
+                       &stream, numOfPriorities + 1);
+  }
+
+  for (int i=0; i < TOTALTHREADS; i++) {
+    T[i].join();
+  }
+  if (g_thTestPassed) {
+    TestPassed = true;
+  } else {
+    TestPassed = false;
+  }
+
+  // Destroy the stream for each of the priority levels
+  count = 1;
+  for (priority = priority_high; priority < priority_low; priority++) {
+    HIP_CHECK(hipStreamDestroy(stream[count++]));
+  }
+  return TestPassed;
+}
+
+
+template <typename T>
+bool validateStreamPrioritiesWithEvents() {
+  size_t size = NUMITERS*MEMCPYSIZE;
+
+  // get the range of priorities available
+  #define OP(x) \
+     int priority_##x; \
+     bool enable_priority_##x = false;
+  OP(low)
+  OP(normal)
+  OP(high)
+  #undef OP
+  HIP_CHECK(hipDeviceGetStreamPriorityRange(&priority_low, &priority_high));
+
+  INFO("HIP stream priority range - low: " << priority_low << ",high: "
+                                         << priority_high << ",normal: "
+                                         << (priority_low + priority_high)/2);
+  // Check if priorities are indeed supported
+  if (priority_low == priority_high) {
+    WARN("Stream priority range not supported. Skipping test.");
+    return true;
+  }
+
+  // Enable/disable priorities based on number of available priority levels
+  enable_priority_low = true;
+  enable_priority_high = true;
+  if ((priority_low - priority_high) > 1) {
+    enable_priority_normal = true;
+  }
+  if (enable_priority_normal) {
+    priority_normal = ((priority_low + priority_high) / 2);
+  }
+  // create streams with highest and lowest available priorities
+  #define OP(x)\
+    hipStream_t stream_##x;\
+    if (enable_priority_##x) {\
+      HIP_CHECK(hipStreamCreateWithPriority(&stream_##x, \
+              hipStreamDefault, priority_##x));\
+    }
+  OP(low)
+  OP(normal)
+  OP(high)
+  #undef OP
+
+  // allocate and initialise host source and destination buffers
+  #define OP(x) \
+    T* src_h_##x; \
+    T* dst_h_##x; \
+    if (enable_priority_##x) { \
+      src_h_##x = reinterpret_cast<T*>(malloc(size)); \
+      REQUIRE(src_h_##x != nullptr); \
+      mem_init<T>(src_h_##x, (size / sizeof(T))); \
+      dst_h_##x = reinterpret_cast<T*>(malloc(size)); \
+      REQUIRE(dst_h_##x != nullptr); \
+      memset(dst_h_##x, 0, size); \
+    }
+  OP(low)
+  OP(normal)
+  OP(high)
+  #undef OP
+
+  // allocate and initialize device source and destination buffers
+  #define OP(x) \
+    T* src_d_##x; \
+    T* dst_d_##x; \
+    if (enable_priority_##x) { \
+      HIP_CHECK(hipMalloc(&src_d_##x, size)); \
+      HIP_CHECK( \
+        hipMemcpy(src_d_##x, src_h_##x, size, hipMemcpyHostToDevice)); \
+      HIP_CHECK(hipMalloc(&dst_d_##x, size)); \
+    }
+  OP(low)
+  OP(normal)
+  OP(high)
+  #undef OP
+
+  // create events for measuring time spent in kernel execution
+  #define OP(x) \
+    hipEvent_t event_start_##x; \
+    hipEvent_t event_end_##x; \
+    if (enable_priority_##x) { \
+      HIP_CHECK(hipEventCreate(&event_start_##x)); \
+      HIP_CHECK(hipEventCreate(&event_end_##x)); \
+    }
+  OP(low)
+  OP(normal)
+  OP(high)
+  #undef OP
+
+  // record start events for each of the priority streams
+  #define OP(x) \
+    if (enable_priority_##x) { \
+      HIP_CHECK(\
+      hipEventRecord(event_start_##x, stream_##x)); \
+    }
+  OP(low)
+  OP(normal)
+  OP(high)
+  #undef OP
+
+  // launch kernels repeatedly on each of the prioritiy streams
+  for (int i = 0; i < static_cast<int>(size); i += MEMCPYSIZE) {
+    int j = i / sizeof(T);
+    #define OP(x) \
+      if (enable_priority_##x) { \
+        hipLaunchKernelGGL((memcpy_kernel<T>), dim3(GRIDSIZE), \
+        dim3(BLOCKSIZE), 0, stream_##x, dst_d_##x + j, src_d_##x + j, \
+        (MEMCPYSIZE / sizeof(T))); \
+      }
+    OP(low)
+    OP(normal)
+    OP(high)
+    #undef OP
+  }
+
+  // record end events for each of the priority streams
+  #define OP(x) \
+    if (enable_priority_##x) { \
+      HIP_CHECK(hipEventRecord(event_end_##x, stream_##x)); \
+    }
+  OP(low)
+  OP(normal)
+  OP(high)
+  #undef OP
+
+  // synchronize events for each of the priority streams
+  #define OP(x) \
+    if (enable_priority_##x) { \
+      HIP_CHECK(hipEventSynchronize(event_end_##x)); \
+    }
+  OP(low)
+  OP(normal)
+  OP(high)
+  #undef OP
+
+  // compute time spent for memcpy in each stream
+  #define OP(x) \
+    float time_spent_##x; \
+    if (enable_priority_##x) { \
+      HIP_CHECK(hipEventElapsedTime(&time_spent_##x, \
+              event_start_##x, event_end_##x)); \
+      INFO("time spent for memcpy in " << #x << \
+      " priority stream: " << time_spent_##x << " ms"); \
+    }
+  OP(low)
+  OP(normal)
+  OP(high)
+  #undef OP
+
+  // sanity check
+  #define OP(x) \
+    if (enable_priority_##x) { \
+      HIP_CHECK(hipMemcpy(dst_h_##x, dst_d_##x, size, \
+                         hipMemcpyDeviceToHost)); \
+      if (memcmp(dst_h_##x, src_h_##x, size) != 0) { \
+        REQUIRE(false); \
+      } \
+    }
+  OP(low)
+  OP(normal)
+  OP(high)
+  #undef OP
+
+  // validate that stream priorities are working as expected
+  #define OP(x, y) \
+      if (enable_priority_##x && enable_priority_##y) { \
+          if ((1.05f * time_spent_##x) < time_spent_##y) { \
+            INFO("time_spent_##x : " << time_spent_##x << \
+            "time_spent_##y : " << time_spent_##y); \
+            REQUIRE(false); \
+          } \
+      }
+  OP(low, normal)
+  OP(normal, high)
+  OP(low, high)
+  #undef OP
+
+  return true;
+}
+
+}  // namespace hipStreamCreateWithPriorityTest
+
+
+
+/**
+Tests following scenarios.
+ 1)Create streams with default flag for all available priority levels and
+ queue tasks in each of these streams, perform device synchronize and validate behavior.
+ 2)Create streams with non-blocking flag for all available priority levels and
+ queue tasks in each of these streams, perform stream synchronize and validate behavior.
+ 3)Create streams with default flag for all available priority levels and
+ queue tasks in each of these streams, perform stream synchronize and validate behavior.
+ 4)Create streams with non-blocking flag for all available priority levels and
+ queue tasks in each of these streams, perform device synchronize and validate behavior.
+*/
+TEST_CASE("Unit_hipStreamCreateWithPriority_FunctionalForAllPriorities") {
+  SECTION("Default flag and device synchronize") {
+    hipStreamCreateWithPriorityTest::
+    funcTestsForAllPriorityLevelsWrtNullStrm(hipStreamDefault, true);
+  }
+
+  SECTION("Stream non-blocking flag and stream synchronize") {
+    hipStreamCreateWithPriorityTest::
+    funcTestsForAllPriorityLevelsWrtNullStrm(hipStreamNonBlocking, false);
+  }
+
+  SECTION("Default flag and stream synchronize") {
+    hipStreamCreateWithPriorityTest::
+    funcTestsForAllPriorityLevelsWrtNullStrm(hipStreamDefault, false);
+  }
+
+  SECTION("Stream non-blocking flag and device synchronize") {
+    hipStreamCreateWithPriorityTest::
+    funcTestsForAllPriorityLevelsWrtNullStrm(hipStreamNonBlocking, true);
+  }
+}
+
+/**
+ * Create a stream for each priority level with default flag, Launch memcpy and kernel
+ * tasks on these streams from multiple threads. Validate all the results.
+ */
+TEST_CASE("Unit_hipStreamCreateWithPriority_MulthreadDefaultflag") {
+  bool TestPassed = true;
+  TestPassed = hipStreamCreateWithPriorityTest::
+  runFuncTestsForAllPriorityLevelsMultThread(hipStreamDefault);
+  REQUIRE(TestPassed);
+}
+
+/**
+ * Create a stream for each priority level with non-blocking flag, Launch memcpy and
+ * kernel tasks on these streams from multiple threads. Validate all the results.
+ */
+TEST_CASE("Unit_hipStreamCreateWithPriority_MulthreadNonblockingflag") {
+  bool TestPassed = true;
+  TestPassed = hipStreamCreateWithPriorityTest::
+  runFuncTestsForAllPriorityLevelsMultThread(hipStreamNonBlocking);
+  REQUIRE(TestPassed);
+}
+
+
+/**
+ * Scenario1: Validates functionality of hipStreamCreateWithPriority when
+              stream = nullptr.
+ * Scenario2: Validates functionality of hipStreamCreateWithPriority when
+              flag = 0xffffffff.
+*/
+TEST_CASE("Unit_hipStreamCreateWithPriority_NegTst") {
+  hipStream_t stream;
+  int priority_low;
+  int priority_high;
+  hipError_t ret;
+
+  // Test is to get the Stream Priority Range
+  HIP_CHECK(
+          hipDeviceGetStreamPriorityRange(&priority_low, &priority_high));
+  // Check if priorities are indeed supported
+  if (priority_low == priority_high) {
+    WARN("Stream priority range not supported. Skipping test.");
+    return;  // exit the test since priorities are not supported
+  }
+
+  SECTION("stream = nullptr test") {
+    ret = hipStreamCreateWithPriority(nullptr, hipStreamDefault,
+                                    priority_low);
+
+    REQUIRE(ret != hipSuccess);
+  }
+
+  SECTION("flag value invalid test") {
+    ret = hipStreamCreateWithPriority(&stream, 0xffffffff,
+                                 priority_low);
+
+    REQUIRE(ret != hipSuccess);
+  }
+}
+
+/**
+ * Validate stream priorities with event after classifying them as low, medium and high.
+ */
+TEST_CASE("Unit_hipStreamCreateWithPriority_ValidateWithEvents") {
+  bool TestPassed = true;
+  TestPassed = hipStreamCreateWithPriorityTest::
+  validateStreamPrioritiesWithEvents<int>();
+  REQUIRE(TestPassed);
+}
@@ -0,0 +1,350 @@
+/*
+Copyright (c) 2021 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.
+*/
+
+/**
+Testcase Scenarios :
+
+1)Validates functionality of hipStreamAddCallback with created stream.
+
+2)Validates functionality of stream with cu mask.
+
+3)Create a stream with all CU masks disabled (0x00000000).
+Verify that default CU mask is set for the stream.
+
+4)Size is greater than physical CU number. In this case the extra elements
+are ignored and hipExtStreamCreateWithCUMask must return hipSuccess.
+
+5)Negative Testing of hipExtStreamCreateWithCUMask.
+*/
+
+
+#include <hip_test_common.hh>
+#include <hip_test_kernels.hh>
+#include <unistd.h>
+#include <iostream>
+#include <vector>
+
+#define NUM_CU_PARTITIONS 4
+#define CONSTANT 1.618f
+#define SIZE_INBYTES_OF_MB (1024*1024)
+#define GRIDSIZE 512
+#define BLOCKSIZE 256
+#define ZERO_MASK 0x00000000
+
+
+namespace hipExtStreamCreateWithCUMaskTest {
+
+float *A_h, *C_h;
+bool cbDone = false;
+bool isPassed = true;
+size_t N = 4 * SIZE_INBYTES_OF_MB;
+
+
+// Make a default CU mask bit-array where all CUs are active
+// this default mask is expected to be returned when there is no
+// custom or global CU mask defined.
+void createDefaultCUMask(std::vector<uint32_t> *pdefaultCUMask,
+                         int numOfCUs) {
+  uint32_t temp = 0;
+  uint32_t bit_index = 0;
+  for (int i = 0; i < numOfCUs; i++) {
+    temp |= 1UL << bit_index;
+    if (bit_index >= 32) {
+      (*pdefaultCUMask).push_back(temp);
+      temp = 0;
+      bit_index = 0;
+      temp |= 1UL << bit_index;
+    }
+    bit_index += 1;
+  }
+  if (bit_index != 0) {
+    (*pdefaultCUMask).push_back(temp);
+  }
+}
+// Create masks of disabled CU masks.
+void createDisabledCUMask(std::vector<uint32_t> *pdisabledCUMask,
+                         int numOfCUs) {
+  uint32_t temp = ZERO_MASK;
+  uint32_t bit_index = 0;
+  for (int i = 0; i < numOfCUs; i++) {
+    if (bit_index >= 32) {
+      (*pdisabledCUMask).push_back(temp);
+      temp = ZERO_MASK;
+      bit_index = 0;
+    }
+    bit_index += 1;
+  }
+  if (bit_index != 0) {
+    (*pdisabledCUMask).push_back(temp);
+  }
+}
+
+void Callback(hipStream_t stream, hipError_t status,
+              void* userData) {
+  isPassed = true;
+  stream = 0;
+  HIP_CHECK(status);
+  REQUIRE(userData == nullptr);
+  for (size_t i = 0; i < N; i++) {
+    if (C_h[i] != A_h[i] * A_h[i]) {
+      UNSCOPED_INFO("Data mismatch at index: " << i);
+      isPassed = false;
+      break;
+    }
+  }
+  cbDone = true;
+}
+}  // namespace hipExtStreamCreateWithCUMaskTest
+
+using hipExtStreamCreateWithCUMaskTest::A_h;
+using hipExtStreamCreateWithCUMaskTest::C_h;
+using hipExtStreamCreateWithCUMaskTest::cbDone;
+using hipExtStreamCreateWithCUMaskTest::isPassed;
+using hipExtStreamCreateWithCUMaskTest::N;
+using hipExtStreamCreateWithCUMaskTest::createDefaultCUMask;
+using hipExtStreamCreateWithCUMaskTest::createDisabledCUMask;
+using hipExtStreamCreateWithCUMaskTest::Callback;
+
+
+/**
+ * Scenario: Validates functionality of hipStreamAddCallback with created stream.
+ */
+TEST_CASE("Unit_hipExtStreamCreateWithCUMask_ValidateCallbackFunc") {
+  float *A_d, *C_d;
+  size_t Nbytes = N * sizeof(float);
+  cbDone = false;
+  A_h = reinterpret_cast<float*>(malloc(Nbytes));
+  REQUIRE(A_h != nullptr);
+  C_h = reinterpret_cast<float*>(malloc(Nbytes));
+  REQUIRE(C_h != nullptr);
+
+  // Fill with Phi + i
+  for (size_t i = 0; i < N; i++) {
+    A_h[i] = CONSTANT + i;
+  }
+
+  HIP_CHECK(hipMalloc(&A_d, Nbytes));
+  HIP_CHECK(hipMalloc(&C_d, Nbytes));
+
+  hipStream_t mystream;
+  std::vector<uint32_t> defaultCUMask;
+  HIP_CHECK(hipSetDevice(0));
+  hipDeviceProp_t props;
+  HIP_CHECK(hipGetDeviceProperties(&props, 0));
+  createDefaultCUMask(&defaultCUMask, props.multiProcessorCount);
+
+  hipExtStreamCreateWithCUMask(&mystream, defaultCUMask.size(),
+                               defaultCUMask.data());
+  HIP_CHECK(hipMemcpyAsync(A_d, A_h, Nbytes, hipMemcpyHostToDevice,
+          mystream));
+  const unsigned blocks = GRIDSIZE;
+  const unsigned threadsPerBlock = BLOCKSIZE;
+  hipLaunchKernelGGL((HipTest::vector_square), dim3(blocks),
+                     dim3(threadsPerBlock), 0, mystream, A_d, C_d, N);
+  HIP_CHECK(hipMemcpyAsync(C_h, C_d, Nbytes, hipMemcpyDeviceToHost,
+          mystream));
+  HIP_CHECK(hipStreamAddCallback(mystream, Callback, nullptr, 0));
+  while (!cbDone) usleep(100000);  // Sleep for 100 ms
+  HIP_CHECK(hipStreamDestroy(mystream));
+  HIP_CHECK(hipFree(reinterpret_cast<void*>(C_d)));
+  HIP_CHECK(hipFree(reinterpret_cast<void*>(A_d)));
+  free(C_h);
+  free(A_h);
+  REQUIRE(isPassed == true);
+}
+
+/**
+ * Scenario: Validates functionality of stream with cu mask.
+ */
+TEST_CASE("Unit_hipExtStreamCreateWithCUMask_Functionality") {
+  const int KNumPartition = NUM_CU_PARTITIONS;
+  float *dA[KNumPartition], *dC[KNumPartition];
+  float *hA, *hC;
+  size_t N = 25*SIZE_INBYTES_OF_MB;
+  size_t Nbytes = N * sizeof(float);
+  std::vector<hipStream_t> streams(KNumPartition);
+  std::vector<std::vector<uint32_t>> cuMasks(KNumPartition);
+  std::stringstream ss[KNumPartition];
+
+  int nGpu = 0;
+  HIP_CHECK(hipGetDeviceCount(&nGpu));
+  if (nGpu < 1) {
+    WARN("Didn't find any GPU! skipping the test!");
+    return;
+  }
+
+  static int device = 0;
+  HIP_CHECK(hipSetDevice(device));
+  hipDeviceProp_t props;
+  HIP_CHECK(hipGetDeviceProperties(&props, device));
+  INFO("info: running on bus " << "0x" << props.pciBusID << " "
+  << props.name << " with " << props.multiProcessorCount << " CUs");
+
+  hA =  new float[Nbytes];
+  REQUIRE(hA != nullptr);
+  hC =  new float[Nbytes];
+  REQUIRE(hC != nullptr);
+  for (size_t i = 0; i < N; i++) {
+    hA[i] = CONSTANT + i;
+  }
+
+  for (int np = 0; np < KNumPartition; np++) {
+    HIP_CHECK(hipMalloc(&dA[np], Nbytes));
+    HIP_CHECK(hipMalloc(&dC[np], Nbytes));
+    // make unique CU masks in the multiple of dwords for each stream
+    uint32_t temp = 0;
+    uint32_t bit_index = np;
+    for (int i = np; i < props.multiProcessorCount; i = i + 4) {
+      temp |= 1UL << bit_index;
+      if (bit_index >= 32) {
+        cuMasks[np].push_back(temp);
+        temp = 0;
+        bit_index = np;
+        temp |= 1UL << bit_index;
+      }
+      bit_index += 4;
+    }
+    if (bit_index != 0) {
+      cuMasks[np].push_back(temp);
+    }
+
+    HIP_CHECK(hipExtStreamCreateWithCUMask(&streams[np], cuMasks[np].size(),
+            cuMasks[np].data()));
+
+    HIP_CHECK(hipMemcpy(dA[np], hA, Nbytes, hipMemcpyHostToDevice));
+
+    ss[np] << std::hex;
+    for (int i = cuMasks[np].size() - 1; i >= 0; i--) {
+      ss[np] << cuMasks[np][i];
+    }
+  }
+
+  const unsigned blocks = GRIDSIZE;
+  const unsigned threadsPerBlock = BLOCKSIZE;
+
+  auto single_start = std::chrono::steady_clock::now();
+  INFO("info: launch 'vector_square' kernel on one stream " <<
+  streams[0] << " with CU mask: 0x" << ss[0].str().c_str());
+
+  hipLaunchKernelGGL(HipTest::vector_square, dim3(blocks),
+                     dim3(threadsPerBlock), 0, streams[0], dA[0], dC[0], N);
+  hipDeviceSynchronize();
+
+  auto single_end = std::chrono::steady_clock::now();
+  std::chrono::duration<double> single_kernel_time = single_end - single_start;
+
+  HIP_CHECK(hipMemcpy(hC, dC[0], Nbytes, hipMemcpyDeviceToHost));
+
+  for (size_t i = 0; i < N; i++) {
+    REQUIRE(hC[i] == (hA[i] * hA[i]));
+  }
+
+  INFO("info: launch 'vector_square' kernel on "
+  << KNumPartition << " streams:");
+  auto all_start = std::chrono::steady_clock::now();
+  for (int np = 0; np < KNumPartition; np++) {
+    INFO("info: launch 'vector_square' kernel on the stream "
+    << streams[np] << " with CU mask: 0x" << ss[np].str().c_str());
+    hipLaunchKernelGGL(HipTest::vector_square, dim3(blocks),
+    dim3(threadsPerBlock), 0, streams[np], dA[np], dC[np], N);
+  }
+  hipDeviceSynchronize();
+
+  auto all_end = std::chrono::steady_clock::now();
+  std::chrono::duration<double> all_kernel_time = all_end - all_start;
+
+  for (int np = 0; np < KNumPartition; np++) {
+    HIP_CHECK(hipMemcpy(hC, dC[np], Nbytes, hipMemcpyDeviceToHost));
+    for (size_t i = 0; i < N; i++) {
+      REQUIRE(hC[i] == (hA[i] * hA[i]));
+    }
+  }
+
+  INFO("info: kernel launched on one stream took: " <<
+  single_kernel_time.count() << " seconds");
+  INFO("info: kernels launched on " << KNumPartition <<
+  " streams took: " << all_kernel_time.count() << " seconds");
+  INFO("info: launching kernels on " << KNumPartition <<
+  " streams asynchronously is " <<
+  single_kernel_time.count() / (all_kernel_time.count() / KNumPartition)
+  << " times faster per stream than launching on one stream alone");
+
+  delete [] hA;
+  delete [] hC;
+  for (int np = 0; np < KNumPartition; np++) {
+    hipFree(dC[np]);
+    hipFree(dA[np]);
+    HIP_CHECK(hipStreamDestroy(streams[np]));
+  }
+}
+
+/**
+ * Scenario: Create a stream with all CU masks disabled (0x00000000).
+ * Verify that default CU mask is set for the stream.
+ */
+TEST_CASE("Unit_hipExtStreamCreateWithCUMask_AllCUsMasked") {
+  HIP_CHECK(hipSetDevice(0));
+  hipDeviceProp_t props;
+  HIP_CHECK(hipGetDeviceProperties(&props, 0));
+  // make a CU mask with all CUs disabled.
+  std::vector<uint32_t> allCUDisabled;
+  createDisabledCUMask(&allCUDisabled, props.multiProcessorCount);
+  hipStream_t stream;
+  HIP_CHECK(hipExtStreamCreateWithCUMask(&stream, allCUDisabled.size(),
+          allCUDisabled.data()));
+  // Verify whether default CU mask is set for the stream.
+  uint32_t size = (props.multiProcessorCount / 32) + 1;
+  std::vector<uint32_t> cuMask(size);
+  std::vector<uint32_t> defaultCUMask;
+  createDefaultCUMask(&defaultCUMask, props.multiProcessorCount);
+  HIP_CHECK(hipExtStreamGetCUMask(stream, cuMask.size(), &cuMask[0]));
+  for (int i = 0; i < static_cast<int>(defaultCUMask.size()); i++) {
+    REQUIRE(defaultCUMask[i] == cuMask[i]);
+  }
+  HIP_CHECK(hipStreamDestroy(stream));
+}
+
+/**
+ * Scenario: Negative Testing of hipExtStreamCreateWithCUMask.
+ */
+TEST_CASE("Unit_hipExtStreamCreateWithCUMask_NegTst") {
+  std::vector<uint32_t> defaultCUMask;
+  REQUIRE(hipSuccess == hipSetDevice(0));
+  hipDeviceProp_t props;
+  REQUIRE(hipSuccess == hipGetDeviceProperties(&props, 0));
+  createDefaultCUMask(&defaultCUMask, props.multiProcessorCount);
+  hipStream_t stream;
+  // Negative Scenario 1: stream = nullptr
+  SECTION("stream is nullptr") {
+    REQUIRE_FALSE(hipExtStreamCreateWithCUMask(nullptr,
+                  defaultCUMask.size(),
+                  defaultCUMask.data()) == hipSuccess);
+  }
+  // Negative Scenario 2: cuMaskSize = 0
+  SECTION("cuMaskSize is 0") {
+    REQUIRE_FALSE(hipExtStreamCreateWithCUMask(&stream, 0,
+                  defaultCUMask.data()) == hipSuccess);
+  }
+  // Negative Scenario 3: cuMask = nullptr
+  SECTION("cuMask is nullptr") {
+    REQUIRE_FALSE(hipExtStreamCreateWithCUMask(&stream,
+                  defaultCUMask.size(),
+                  nullptr) == hipSuccess);
+  }
+}
@@ -18,7 +18,7 @@
 */

 /* HIT_START
- * BUILD: %t %s ../../src/test_common.cpp EXCLUDE_HIP_PLATFORM nvidia
+ * BUILD: %t %s ../../src/test_common.cpp
 * TEST: %t
 * HIT_END
 */
@@ -18,7 +18,7 @@
 */

 /* HIT_START
- * BUILD: %t %s ../../src/test_common.cpp EXCLUDE_HIP_PLATFORM nvidia
+ * BUILD: %t %s ../../src/test_common.cpp
 * TEST: %t
 * HIT_END
 */
@@ -489,6 +489,7 @@ void hipPerfMandelBrot::double_mandel_unroll(uint *out, uint width, float xPos,
 void hipPerfMandelBrot::run(unsigned int testCase,unsigned int deviceId) {

  unsigned int numStreams = getNumStreams();
+  coordIdx = testCase % numCoords;

  funPtr p[] = {&hipPerfMandelBrot::float_mad, &hipPerfMandelBrot::float_mandel_unroll,
               &hipPerfMandelBrot::double_mad, &hipPerfMandelBrot::double_mandel_unroll};
@@ -555,9 +556,6 @@ void hipPerfMandelBrot::run(unsigned int testCase,unsigned int deviceId) {
  double totalTime = 0.0;

  for (unsigned int k = 0; k < numLoops; k++) {
-
-  coordIdx = testCase % numCoords;
-
  if ((testCase == 0 || testCase == 1 || testCase == 2 ||
                  testCase == 5 || testCase == 6 || testCase == 7 ||
                  testCase == 10 || testCase == 11 || testCase == 12)) {
@@ -653,7 +651,7 @@ void hipPerfMandelBrot::run(unsigned int testCase,unsigned int deviceId) {

  // Free host and device memory
  for (uint i = 0; i < numKernels; i++) {
-    HIPCHECK(hipFree(hPtr[i]));
+    HIPCHECK(hipHostFree(hPtr[i]));
    HIPCHECK(hipFree(dPtr[i]));
  }

@@ -18,7 +18,7 @@
 */

 /* HIT_START
- * BUILD: %t %s ../../src/test_common.cpp ../../src/timer.cpp EXCLUDE_HIP_PLATFORM nvidia
+ * BUILD: %t %s ../../src/test_common.cpp ../../src/timer.cpp
 * TEST: %t
 * HIT_END
 */
@@ -18,7 +18,7 @@ THE SOFTWARE.
 */

 /* HIT_START
- * BUILD: %t %s ../../src/test_common.cpp ../../src/timer.cpp EXCLUDE_HIP_PLATFORM nvidia
+ * BUILD: %t %s ../../src/test_common.cpp ../../src/timer.cpp
 * TEST: %t
 * HIT_END
 */
@@ -18,7 +18,7 @@ THE SOFTWARE.
 */

 /* HIT_START
- * BUILD: %t %s ../../src/test_common.cpp ../../src/timer.cpp EXCLUDE_HIP_PLATFORM nvidia
+ * BUILD: %t %s ../../src/test_common.cpp ../../src/timer.cpp
 * TEST: %t
 * HIT_END
 */
@@ -18,7 +18,7 @@ THE SOFTWARE.
 */

 /* HIT_START
- * BUILD: %t %s ../../src/test_common.cpp EXCLUDE_HIP_PLATFORM nvidia
+ * BUILD: %t %s ../../src/test_common.cpp
 * TEST: %t
 * HIT_END
 */
@@ -18,7 +18,7 @@ THE SOFTWARE.
 */

 /* HIT_START
- * BUILD: %t %s ../../src/test_common.cpp EXCLUDE_HIP_PLATFORM nvidia
+ * BUILD: %t %s ../../src/test_common.cpp
 * TEST: %t
 * HIT_END
 */
@@ -18,7 +18,7 @@
 */

 /* HIT_START
- * BUILD: %t %s ../../src/test_common.cpp EXCLUDE_HIP_PLATFORM nvidia
+ * BUILD: %t %s ../../src/test_common.cpp
 * TEST: %t
 * HIT_END
 */
@@ -136,6 +136,7 @@ class hipPerfMemFill {
    }

    HIPCHECK(hipSetDevice(deviceId));
+    memset(&props_, 0, sizeof(props_));
    HIPCHECK(hipGetDeviceProperties(&props_, deviceId));
    std::cout << "Info: running on device: id: " << deviceId << ", bus: 0x"
        << props_.pciBusID << " " << props_.name << " with "
@@ -397,8 +398,9 @@ class hipPerfMemFill {
    return true;
  }

-  /* This fuction should be via device attribute query*/
+  /* This function should be via device attribute query*/
  bool supportDeviceMallocFinegrained() {
+#ifdef __HIP_PLATFORM_AMD__
    T *A = nullptr;
    hipExtMallocWithFlags((void **)&A, sizeof(T), hipDeviceMallocFinegrained);
    if (!A) {
@@ -406,6 +408,9 @@ class hipPerfMemFill {
    }
    HIPCHECK(hipFree(A));
    return true;
+#else
+    return false;
+#endif
  }

  unsigned int setNumBlocks(size_t size) {
@@ -419,6 +424,7 @@ class hipPerfMemFill {
 #endif
  }

+#ifdef __HIP_PLATFORM_AMD__
  bool testExtDeviceMemoryHostFill(size_t size, unsigned int flags) {
    double GBytes = (double) size / (1024.0 * 1024.0 * 1024.0);

@@ -481,6 +487,7 @@ class hipPerfMemFill {

    return true;
  }
+#endif

  bool run() {
    if (supportLargeBar()) {
@@ -499,11 +506,13 @@ class hipPerfMemFill {
      return false;
    }

+#ifdef __HIP_PLATFORM_AMD__
    if (supportDeviceMallocFinegrained()) {
      if (!testExtDeviceMemory()) {
        return false;
      }
    }
+#endif
    return true;
  }

@@ -22,7 +22,7 @@ THE SOFTWARE.
 #include <time.h>

 /* HIT_START
- * BUILD: %t %s ../../src/test_common.cpp EXCLUDE_HIP_PLATFORM nvidia
+ * BUILD: %t %s ../../src/test_common.cpp
 * TEST: %t
 * HIT_END
 */
@@ -37,10 +37,15 @@ void valSet(int* A, int val, size_t size) {
    }
 }

-void setup(size_t *size, const int num, int **pA) {
+void setup(size_t *size, int &num, int **pA, const size_t totalGlobalMem) {
+
    std::cout << "size: ";
    for (int i = 0; i < num; i++) {
        size[i] = 1 << (i + 6);
+        if((NUM_ITER + 1) * size[i] > totalGlobalMem) {
+          num = i;
+          break;
+        }
        std::cout << size[i] << " ";
    }
    std::cout << std::endl;
@@ -77,11 +82,16 @@ int main() {
    size_t size[NUM_SIZE] = { 0 };
    int *Ad[NUM_ITER] = { nullptr };
    int *A;
+    hipDeviceProp_t props;
+    memset(&props, 0, sizeof(props));
+    HIPCHECK(hipGetDeviceProperties(&props, 0));
+    std::cout << "totalGlobalMem: " << props.totalGlobalMem << std::endl;

-    setup(size, NUM_SIZE, &A);
+    int num = NUM_SIZE;
+    setup(size, num, &A, props.totalGlobalMem);
    testInit(size[0], A);

-    for (int i = 0; i < NUM_SIZE; i++) {
+    for (int i = 0; i < num; i++) {
        std::cout << size[i] << std::endl;
        start = clock();
        for (int j = 0; j < NUM_ITER; j++) {
@@ -18,7 +18,7 @@
 */

 /* HIT_START
- * BUILD: %t %s ../../src/test_common.cpp EXCLUDE_HIP_PLATFORM nvidia
+ * BUILD: %t %s ../../src/test_common.cpp
 * TEST: %t
 * HIT_END
 */
@@ -18,7 +18,7 @@
 */

 /* HIT_START
- * BUILD: %t %s ../../src/test_common.cpp EXCLUDE_HIP_PLATFORM nvidia
+ * BUILD: %t %s ../../src/test_common.cpp
 * TEST: %t
 * HIT_END
 */
@@ -140,13 +140,13 @@ void hipPerfMemset::run1D(unsigned int test, T memsetval, enum MemsetType type,
  HIPCHECK(hipStreamCreate(&stream));

  // Warm-up
-  HIPCHECK(hipMemset((hipDeviceptr_t)A_d, memsetval, bufSize_));
+  HIPCHECK(hipMemset((void *)A_d, memsetval, bufSize_));

  auto start = chrono::steady_clock::now();

  for (uint i = 0; i < NUM_ITER; i++) {
    if (type == hipMemsetTypeDefault && !async) {
-      HIPCHECK(hipMemset((hipDeviceptr_t)A_d, memsetval, bufSize_));
+      HIPCHECK(hipMemset((void *)A_d, memsetval, bufSize_));
    }
    else if (type == hipMemsetTypeDefault && async) {
      HIPCHECK(hipMemsetAsync(A_d, memsetval, bufSize_, stream));
@@ -18,7 +18,7 @@
 */

 /* HIT_START
- * BUILD: %t %s ../../src/test_common.cpp EXCLUDE_HIP_PLATFORM nvidia
+ * BUILD: %t %s ../../src/test_common.cpp
 * TEST: %t
 * HIT_END
 */
@@ -42,6 +42,18 @@ vector<unsigned int> sizes = {1,  2,   4,   8,   16,   32,
 #define NUM_BUFS 6
 #define MAX_BUFS (1 << (NUM_BUFS - 1))

+#ifdef __HIP_PLATFORM_NVIDIA__
+inline __host__ __device__ void operator+=(float2 &a, float2 b)
+{
+  a.x += b.x; a.y += b.y;
+}
+
+inline __host__ __device__ void operator+=(float4 &a, float4 b)
+{
+  a.x += b.x; a.y += b.y; a.z += b.z; a.w += b.w;
+}
+#endif
+
 template <typename T>
 __global__ void sampleRate(T * outBuffer, unsigned int inBufSize, unsigned int writeIt,
                           T **inBuffer, int numBufs) {
@@ -49,7 +61,8 @@ __global__ void sampleRate(T * outBuffer, unsigned int inBufSize, unsigned int w
  uint gid = (blockIdx.x * blockDim.x + threadIdx.x);
  uint inputIdx = gid % inBufSize;

-  T tmp = (T)0.0f;
+  T tmp;
+  memset(&tmp, 0, sizeof(T));
  for(int i = 0; i < numBufs; i++) {
    tmp += *(*(inBuffer+i)+inputIdx);
  }
@@ -264,11 +277,11 @@ void hipPerfSampleRate::run(unsigned int test) {

   // Free host and device memory
   for (uint i = 0; i < numBufs_; i++) {
-    HIPCHECK(hipFree(hInPtr[i]));
+    HIPCHECK(hipHostFree(hInPtr[i]));
    HIPCHECK(hipFree(dInPtr[i]));
    }

-   HIPCHECK(hipFree(hOutPtr));
+   HIPCHECK(hipHostFree(hOutPtr));
   HIPCHECK(hipFree(dPtr));
 }

@@ -18,7 +18,7 @@
 */

 /* HIT_START
- * BUILD: %t %s ../../src/test_common.cpp EXCLUDE_HIP_PLATFORM nvidia
+ * BUILD: %t %s ../../src/test_common.cpp
 * TEST: %t
 * HIT_END
 */
@@ -18,7 +18,7 @@
 */

 /* HIT_START
- * BUILD: %t %s ../../src/test_common.cpp EXCLUDE_HIP_PLATFORM nvidia
+ * BUILD: %t %s ../../src/test_common.cpp
 * TEST: %t
 * HIT_END
 */
@@ -226,7 +226,7 @@ void hipPerfDeviceConcurrency::run(unsigned int testCase, int numGpus) {
  HIPCHECK(hipStreamDestroy(streams[i]));

  // Free host and device memory
-  HIPCHECK(hipFree(hPtr[i]));
+  HIPCHECK(hipHostFree(hPtr[i]));
  HIPCHECK(hipFree(dPtr[i]));
  }

@@ -18,7 +18,7 @@
 */

 /* HIT_START
- * BUILD: %t %s ../../src/test_common.cpp EXCLUDE_HIP_PLATFORM nvidia
+ * BUILD: %t %s ../../src/test_common.cpp
 * TEST: %t
 * HIT_END
 */
@@ -28,6 +28,24 @@
 #include "test_common.h"
 #include <hip/hip_vector_types.h>

+#ifdef __HIP_PLATFORM_NVIDIA__
+inline __device__ float4 operator*(float s, float4 a)
+{
+  return make_float4(a.x * s, a.y * s, a.z * s, a.w * s);
+}
+inline __device__ float4 operator*(float4 a, float4 b)
+{
+  return make_float4(a.x * b.x, a.y * b.y, a.z * b.z, a.w * b.w);
+}
+inline __device__ float4 operator+(float4 a, float4 b)
+{
+  return make_float4(a.x + b.x, a.y + b.y, a.z + b.z, a.w + b.w);
+}
+inline __device__ float4 operator-(float4 a, float4 b)
+{
+  return make_float4(a.x - b.x, a.y - b.y, a.z - b.z, a.w - b.w);
+}
+#endif

 typedef struct {
  double x;
@@ -39,157 +57,130 @@ static coordRec coords[] = {
    {0.0, 0.0, 0.00001},  // All black
 };

-
 static unsigned int numCoords = sizeof(coords) / sizeof(coordRec);

 __global__ void mandelbrot(uint *out, uint width, float xPos, float yPos,
         float xStep, float yStep, uint maxIter) {
-
-
  int tid = (blockIdx.x * blockDim.x + threadIdx.x);
  int i = tid % (width/4);
  int j = tid / (width/4);
  int4 veci = make_int4(4*i, 4*i+1, 4*i+2, 4*i+3);
  int4 vecj = make_int4(j, j, j, j);
  float4 x0;
-  x0.data[0] = (float)(xPos + xStep*veci.data[0]);
-  x0.data[1] = (float)(xPos + xStep*veci.data[1]);
-  x0.data[2] = (float)(xPos + xStep*veci.data[2]);
-  x0.data[3] = (float)(xPos + xStep*veci.data[3]);
+  x0.x = (float)(xPos + xStep*veci.x);
+  x0.y = (float)(xPos + xStep*veci.y);
+  x0.z = (float)(xPos + xStep*veci.z);
+  x0.w = (float)(xPos + xStep*veci.w);
  float4 y0;
-  y0.data[0] = (float)(yPos + yStep*vecj.data[0]);
-  y0.data[1] = (float)(yPos + yStep*vecj.data[1]);
-  y0.data[2] = (float)(yPos + yStep*vecj.data[2]);
-  y0.data[3] = (float)(yPos + yStep*vecj.data[3]);
-
+  y0.x = (float)(yPos + yStep*vecj.x);
+  y0.y = (float)(yPos + yStep*vecj.y);
+  y0.z = (float)(yPos + yStep*vecj.z);
+  y0.w = (float)(yPos + yStep*vecj.w);
  float4 x = x0;
  float4 y = y0;
-
  uint iter = 0;
  float4 tmp;
  int4 stay;
  int4 ccount = make_int4(0, 0, 0, 0);
-
  float4 savx = x;
  float4 savy = y;
-
-  stay.data[0] = (x.data[0]*x.data[0]+y.data[0]*y.data[0]) <= (float)(4.0f);
-  stay.data[1] = (x.data[1]*x.data[1]+y.data[1]*y.data[1]) <= (float)(4.0f);
-  stay.data[2] = (x.data[2]*x.data[2]+y.data[2]*y.data[2]) <= (float)(4.0f);
-  stay.data[3] = (x.data[3]*x.data[3]+y.data[3]*y.data[3]) <= (float)(4.0f);
-
-  for (iter = 0; (stay.data[0] | stay.data[1] | stay.data[2] | stay.data[3]) && (iter < maxIter);
+  stay.x = (x.x*x.x+y.x*y.x) <= (float)(4.0f);
+  stay.y = (x.y*x.y+y.y*y.y) <= (float)(4.0f);
+  stay.z = (x.z*x.z+y.z*y.z) <= (float)(4.0f);
+  stay.w = (x.w*x.w+y.w*y.w) <= (float)(4.0f);
+  for (iter = 0; (stay.x | stay.y | stay.z | stay.w) && (iter < maxIter);
  iter+=16) {
-
-
    x = savx;
    y = savy;
-
    // Two iterations
    tmp = x*x + x0 - y*y;
    y = 2.0f * x * y + y0;
    x = tmp*tmp + x0 - y*y;
    y = 2.0f * tmp * y + y0;
-
    // Two iterations
    tmp = x*x + x0 - y*y;
    y = 2.0f * x * y + y0;
    x = tmp*tmp + x0 - y*y;
    y = 2.0f * tmp * y + y0;
-
    // Two iterations
    tmp = x*x + x0 - y*y;
    y = 2.0f * x * y + y0;
    x = tmp*tmp + x0 - y*y;
    y = 2.0f * tmp * y + y0;
-
    // Two iterations
    tmp = x*x + x0 - y*y;
    y = 2.0f * x * y + y0;
    x = tmp*tmp + x0 - y*y;
    y = 2.0f * tmp * y + y0;
-
    // Two iterations
    tmp = x*x + x0 - y*y;
    y = 2.0f * x * y + y0;
    x = tmp*tmp + x0 - y*y;
    y = 2.0f * tmp * y + y0;
-
    // Two iterations
    tmp = x*x + x0 - y*y;
    y = 2.0f * x * y + y0;
    x = tmp*tmp + x0 - y*y;
    y = 2.0f * tmp * y + y0;
-
    // Two iterations
    tmp = x*x + x0 - y*y;
    y = 2.0f * x * y + y0;
    x = tmp*tmp + x0 - y*y;
    y = 2.0f * tmp * y + y0;
-
-    stay.data[0] = (x.data[0]*x.data[0]+y.data[0]*y.data[0]) <= (float)(4.0f);
-    stay.data[1] = (x.data[1]*x.data[1]+y.data[1]*y.data[1]) <= (float)(4.0f);
-    stay.data[2] = (x.data[2]*x.data[2]+y.data[2]*y.data[2]) <= (float)(4.0f);
-    stay.data[3] = (x.data[3]*x.data[3]+y.data[3]*y.data[3]) <= (float)(4.0f);
-
-    savx.data[0] = (bool)(stay.data[0] ? x.data[0] : savx.data[0]);
-    savx.data[1] = (bool)(stay.data[1] ? x.data[1] : savx.data[1]);
-    savx.data[2] = (bool)(stay.data[2] ? x.data[2] : savx.data[2]);
-    savx.data[3] = (bool)(stay.data[3] ? x.data[3] : savx.data[3]);
-
-    savy.data[0] = (bool)(stay.data[0] ? y.data[0] : savy.data[0]);
-    savy.data[1] = (bool)(stay.data[1] ? y.data[1] : savy.data[1]);
-    savy.data[2] = (bool)(stay.data[2] ? y.data[2] : savy.data[2]);
-    savy.data[3] = (bool)(stay.data[3] ? y.data[3] : savy.data[3]);
-
-    ccount.data[0] -= stay.data[0]*16;
-    ccount.data[1] -= stay.data[1]*16;
-    ccount.data[2] -= stay.data[2]*16;
-    ccount.data[3] -= stay.data[3]*16;
-    }
-
-
+    stay.x = (x.x*x.x+y.x*y.x) <= (float)(4.0f);
+    stay.y = (x.y*x.y+y.y*y.y) <= (float)(4.0f);
+    stay.z = (x.z*x.z+y.z*y.z) <= (float)(4.0f);
+    stay.w = (x.w*x.w+y.w*y.w) <= (float)(4.0f);
+    savx.x = (bool)(stay.x ? x.x : savx.x);
+    savx.y = (bool)(stay.y ? x.y : savx.y);
+    savx.z = (bool)(stay.z ? x.z : savx.z);
+    savx.w = (bool)(stay.w ? x.w : savx.w);
+    savy.x = (bool)(stay.x ? y.x : savy.x);
+    savy.y = (bool)(stay.y ? y.y : savy.y);
+    savy.z = (bool)(stay.z ? y.z : savy.z);
+    savy.w = (bool)(stay.w ? y.w : savy.w);
+    ccount.x -= stay.x*16;
+    ccount.y -= stay.y*16;
+    ccount.z -= stay.z*16;
+    ccount.w -= stay.w*16;
+  }
  // Handle remainder
-  if (!(stay.data[0] & stay.data[1] & stay.data[2] & stay.data[3]))
+  if (!(stay.x & stay.y & stay.z & stay.w))
  {
    iter = 16;
    do
    {
      x = savx;
      y = savy;
-      stay.x = ((x.data[0]*x.data[0]+y.data[0]*y.data[0]) <= 4.0f) && (ccount.data[0] <  maxIter);
-      stay.y = ((x.data[1]*x.data[1]+y.data[1]*y.data[1]) <= 4.0f) && (ccount.data[1] <  maxIter);
-      stay.z = ((x.data[2]*x.data[2]+y.data[2]*y.data[2]) <= 4.0f) && (ccount.data[2] <  maxIter);
-      stay.w = ((x.data[3]*x.data[3]+y.data[3]*y.data[3]) <= 4.0f) && (ccount.data[3] <  maxIter);
+      stay.x = ((x.x*x.x+y.x*y.x) <= 4.0f) && (ccount.x <  maxIter);
+      stay.y = ((x.y*x.y+y.y*y.y) <= 4.0f) && (ccount.y <  maxIter);
+      stay.z = ((x.z*x.z+y.z*y.z) <= 4.0f) && (ccount.z <  maxIter);
+      stay.w = ((x.w*x.w+y.w*y.w) <= 4.0f) && (ccount.w <  maxIter);
      tmp = x;
      x = x*x + x0 - y*y;
      y = 2.0f*tmp*y + y0;
-      ccount.data[0] += stay.data[0];
-      ccount.data[1] += stay.data[1];
-      ccount.data[2] += stay.data[2];
-      ccount.data[3] += stay.data[3];
+      ccount.x += stay.x;
+      ccount.y += stay.y;
+      ccount.z += stay.z;
+      ccount.w += stay.w;
      iter--;
-      savx.data[0] = (stay.data[0] ? x.data[0] : savx.data[0]);
-      savx.data[1] = (stay.data[1] ? x.data[1] : savx.data[1]);
-      savx.data[2] = (stay.data[2] ? x.data[2] : savx.data[2]);
-      savx.data[3] = (stay.data[3] ? x.data[3] : savx.data[3]);
-      savy.data[0] = (stay.data[0] ? y.data[0] : savy.data[0]);
-      savy.data[1] = (stay.data[1] ? y.data[1] : savy.data[1]);
-      savy.data[2] = (stay.data[2] ? y.data[2] : savy.data[2]);
-      savy.data[3] = (stay.data[3] ? y.data[3] : savy.data[3]);
-    } while ((stay.data[0] | stay.data[1] | stay.data[2] | stay.data[3]) && iter);
+      savx.x = (stay.x ? x.x : savx.x);
+      savx.y = (stay.y ? x.y : savx.y);
+      savx.z = (stay.z ? x.z : savx.z);
+      savx.w = (stay.w ? x.w : savx.w);
+      savy.x = (stay.x ? y.x : savy.x);
+      savy.y = (stay.y ? y.y : savy.y);
+      savy.z = (stay.z ? y.z : savy.z);
+      savy.w = (stay.w ? y.w : savy.w);
+    } while ((stay.x | stay.y | stay.z | stay.w) && iter);
  }
-
-
  uint4 *vecOut = (uint4 *)out;
-
-  vecOut[tid].data[0] = (uint)(ccount.data[0]);
-  vecOut[tid].data[1] = (uint)(ccount.data[1]);
-  vecOut[tid].data[2] = (uint)(ccount.data[2]);
-  vecOut[tid].data[3] = (uint)(ccount.data[3]);
+  vecOut[tid].x = (uint)(ccount.x);
+  vecOut[tid].y = (uint)(ccount.y);
+  vecOut[tid].z = (uint)(ccount.z);
+  vecOut[tid].w = (uint)(ccount.w);
 }

-
 class hipPerfStreamConcurrency {
  public:
  hipPerfStreamConcurrency();
@@ -307,7 +298,7 @@ void hipPerfStreamConcurrency::run(unsigned int testCase,unsigned int deviceId)

  // Copy memory asynchronously and concurrently from host to device
  for (uint i = 0; i < numKernels; i++) {
-    HIPCHECK(hipMemcpyHtoDAsync(dPtr[i], hPtr[i], bufSize, streams[i % numStreams]));
+    HIPCHECK(hipMemcpyHtoDAsync(reinterpret_cast<hipDeviceptr_t>(dPtr[i]), hPtr[i], bufSize, streams[i % numStreams]));
  }


@@ -341,7 +332,7 @@ void hipPerfStreamConcurrency::run(unsigned int testCase,unsigned int deviceId)

  // Copy data back from device to the host
  for(uint i = 0; i < numKernels; i++) {
-    HIPCHECK(hipMemcpyDtoHAsync(hPtr[i] ,dPtr[i], bufSize, streams[i % numStreams]));
+    HIPCHECK(hipMemcpyDtoHAsync(hPtr[i], reinterpret_cast<hipDeviceptr_t>(dPtr[i]), bufSize, streams[i % numStreams]));
  }


@@ -361,7 +352,7 @@ void hipPerfStreamConcurrency::run(unsigned int testCase,unsigned int deviceId)

  // Free host and device memory
  for (uint i = 0; i < numKernels; i++) {
-    HIPCHECK(hipFree(hPtr[i]));
+    HIPCHECK(hipHostFree(hPtr[i]));
    HIPCHECK(hipFree(dPtr[i]));
  }

@@ -18,7 +18,7 @@
 */

 /* HIT_START
- * BUILD: %t %s ../../src/test_common.cpp EXCLUDE_HIP_PLATFORM nvidia
+ * BUILD: %t %s ../../src/test_common.cpp
 * TEST: %t
 * HIT_END
 */
@@ -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


 /*
@@ -1,55 +0,0 @@
-/*
-Copyright (c) 2015 - 2021 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 EXCLUDE_HIP_PLATFORM all
- * TEST: %t
- * HIT_END
- */
-
-#include "hip/hip_runtime.h"
-#include "test_common.h"
-#include<stdio.h>
-
-#define ITER 1<<20
-#define SIZE 1024*1024*sizeof(int)
-
-__global__ void Iter(int *Ad){
-    int tx = threadIdx.x + blockIdx.x * blockDim.x;
-    if(tx == 0){
-        for(int i=0;i<ITER;i++){
-            Ad[tx] += 1;
-        }
-    }
-}
-
-int main(){
-    int A=0, *Ad;
-    hipMalloc((void**)&Ad, SIZE);
-    hipMemcpy(Ad, &A, SIZE, hipMemcpyHostToDevice);
-    dim3 dimGrid, dimBlock;
-    dimGrid.x = 1, dimGrid.y =1, dimGrid.z = 1;
-    dimBlock.x = 1, dimBlock.y = 1, dimGrid.z = 1;
-    hipLaunchKernelGGL(HIP_KERNEL_NAME(Iter), dimGrid, dimBlock, 0, 0, Ad);
-    hipMemcpy(&A, Ad, SIZE, hipMemcpyDeviceToHost);
-    passed();
-}
@@ -1,52 +0,0 @@
-/*
-Copyright (c) 2015 - 2021 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.
-*/
-// Test the HCC-specific API extensions for HIP:
-
-/* HIT_START
- * BUILD: %t %s EXCLUDE_HIP_PLATFORM nvidia
- * TEST: %t EXCLUDE_HIP_PLATFORM all
- * HIT_END
- */
-
-#include <stdio.h>
-#include <iostream>
-#include "hip/hip_runtime.h"
-#include "hip/hip_ext.h"
-#include "test_common.h"
-
-#define CHECK(error)                                                                               \
-    if (error != hipSuccess) {                                                                     \
-        fprintf(stderr, "error: '%s'(%d) at %s:%d\n", hipGetErrorString(error), error, __FILE__,   \
-                __LINE__);                                                                         \
-        exit(EXIT_FAILURE);                                                                        \
-    }
-
-
-int main(int argc, char* argv[]) {
-    int deviceId;
-    CHECK(hipGetDevice(&deviceId));
-    hipDeviceProp_t props;
-    CHECK(hipGetDeviceProperties(&props, deviceId));
-    printf("info: running on device #%d %s\n", deviceId, props.name);
-
-    passed();
-};
@@ -20,7 +20,7 @@ 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 EXCLUDE_HIP_PLATFORM nvidia
+ * BUILD: %t %s ../test_common.cpp NVCC_OPTIONS -lnvrtc
 * TEST: %t
 * HIT_END
 */
@@ -81,7 +81,12 @@ int main()
    hipDeviceProp_t props;
    int device = 0;
    hipGetDeviceProperties(&props, device);
+#ifdef __HIP_PLATFORM_AMD__
    std::string sarg = std::string("--gpu-architecture=") + props.gcnArchName;
+#else
+    std::string sarg = std::string("--gpu-architecture=compute_")
+      + std::to_string(props.major) + std::to_string(props.minor);
+#endif
    const char* options[] = {
            sarg.c_str()
    };
@@ -108,11 +113,16 @@ int main()
    hiprtcGetCode(prog, code.data());

    hipModule_t module;
+#ifdef __HIP_PLATFORM_NVIDIA__
+    HIPCHECK(hipInit(0));
+    hipCtx_t ctx;
+    HIPCHECK(hipCtxCreate(&ctx, 0, device));
+#endif
    hipModuleLoadData(&module, code.data());

    hipDeviceptr_t dResult;
    int hResult = 0;
-    hipMalloc(&dResult, sizeof(hResult));
+    hipMalloc((void **)&dResult, sizeof(hResult));
    hipMemcpyHtoD(dResult, &hResult, sizeof(hResult));

    for (decltype(variable_name_vec.size()) i = 0; i != variable_name_vec.size(); ++i) {
@@ -149,10 +159,12 @@ int main()
        if (expected_result[i] != hResult) { failed("Validation failed."); }
    }

-    hipFree(dResult);
+    hipFree((void *)dResult);
    hipModuleUnload(module);
-
    hiprtcDestroyProgram(&prog);

+#ifdef __HIP_PLATFORM_NVIDIA__
+    HIPCHECK(hipCtxDestroy(ctx));
+#endif
    passed();
 }
@@ -20,11 +20,10 @@ 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 EXCLUDE_HIP_PLATFORM nvidia
+ * BUILD: %t %s ../test_common.cpp NVCC_OPTIONS -lnvrtc
 * TEST: %t
 * HIT_END
 */
-
 #include <test_common.h>

 #include <hip/hiprtc.h>
@@ -77,7 +76,12 @@ int main()
    hipDeviceProp_t props;
    int device = 0;
    hipGetDeviceProperties(&props, device);
+#ifdef __HIP_PLATFORM_AMD__
    std::string sarg = std::string("--gpu-architecture=") + props.gcnArchName;
+#else
+    std::string sarg = std::string("--gpu-architecture=compute_")
+      + std::to_string(props.major) + std::to_string(props.minor);
+#endif
    const char* options[] = {
        sarg.c_str()
    };
@@ -106,6 +110,12 @@ int main()

    hipModule_t module;
    hipFunction_t kernel;
+
+#ifdef __HIP_PLATFORM_NVIDIA__
+    HIPCHECK(hipInit(0));
+    hipCtx_t ctx;
+    HIPCHECK(hipCtxCreate(&ctx, 0, device));
+#endif
    hipModuleLoadData(&module, code.data());
    hipModuleGetFunction(&kernel, module, "saxpy");

@@ -123,9 +133,9 @@ int main()
    }

    hipDeviceptr_t dX, dY, dOut;
-    hipMalloc(&dX, bufferSize);
-    hipMalloc(&dY, bufferSize);
-    hipMalloc(&dOut, bufferSize);
+    hipMalloc((void **)&dX, bufferSize);
+    hipMalloc((void **)&dY, bufferSize);
+    hipMalloc((void **)&dOut, bufferSize);
    hipMemcpyHtoD(dX, hX.get(), bufferSize);
    hipMemcpyHtoD(dY, hY.get(), bufferSize);

@@ -150,11 +160,14 @@ int main()
        if (fabs(a * hX[i] + hY[i] - hOut[i]) > fabs(hOut[i])* 1e-6) { failed("Validation failed."); }
    }

-    hipFree(dX);
-    hipFree(dY);
-    hipFree(dOut);
+    hipFree((void *)dX);
+    hipFree((void *)dY);
+    hipFree((void *)dOut);

    hipModuleUnload(module);

+#ifdef __HIP_PLATFORM_NVIDIA__
+    HIPCHECK(hipCtxDestroy(ctx));
+#endif
    passed();
 }
@@ -1,161 +0,0 @@
-/*
-Copyright (c) 2015 - 2021 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 EXCLUDE_HIP_PLATFORM nvidia
- * TEST: %t
- * HIT_END
- */
-
-#include <test_common.h>
-
-#include <hip/hiprtc.h>
-#include <hip/hip_runtime.h>
-
-#include <cassert>
-#include <cstddef>
-#include <memory>
-#include <iostream>
-#include <iterator>
-#include <vector>
-
-static constexpr auto NUM_THREADS{128};
-static constexpr auto NUM_BLOCKS{32};
-
-static constexpr auto saxpy{
-R"(
-#include "test_header.h"
-#include "test_header1.h"
-extern "C"
-__global__
-void saxpy(real a, realptr x, realptr y, realptr out, size_t n)
-{
-    size_t tid = blockIdx.x * blockDim.x + threadIdx.x;
-    if (tid < n) {
-       out[tid] = a * x[tid] + y[tid] ;
-    }
-}
-)"};
-
-int main()
-{
-    using namespace std;
-
-    hiprtcProgram prog;
-    int num_headers = 2;
-    std::vector<const char*> header_names;
-    std::vector<const char*> header_sources;
-    header_names.push_back("test_header.h");
-    header_names.push_back("test_header1.h");
-    header_sources.push_back("#ifndef HIPRTC_TEST_HEADER_H\n#define HIPRTC_TEST_HEADER_H\ntypedef float real;\n#endif //HIPRTC_TEST_HEADER_H\n");
-    header_sources.push_back("#ifndef HIPRTC_TEST_HEADER1_H\n#define HIPRTC_TEST_HEADER1_H\ntypedef float* realptr;\n#endif //HIPRTC_TEST_HEADER1_H\n");
-    hiprtcCreateProgram(&prog,      // prog
-                        saxpy,      // buffer
-                        "saxpy.cu", // name
-                        num_headers,          // numHeaders
-                        &header_sources[0],    // headers
-                        &header_names[0]);   // includeNames
-
-    hipDeviceProp_t props;
-    int device = 0;
-    hipGetDeviceProperties(&props, device);
-    std::string sarg = std::string("--gpu-architecture=") + props.gcnArchName;
-    const char* options[] = {
-        "-hip-pch",
-        sarg.c_str()
-    };
-
-    hiprtcResult compileResult{hiprtcCompileProgram(prog, 2, options)};
-
-    size_t logSize;
-    hiprtcGetProgramLogSize(prog, &logSize);
-
-    if (logSize) {
-        string log(logSize, '\0');
-        hiprtcGetProgramLog(prog, &log[0]);
-
-        cout << log << '\n';
-    }
-
-    if (compileResult != HIPRTC_SUCCESS) { failed("Compilation failed."); }
-
-    size_t codeSize;
-    hiprtcGetCodeSize(prog, &codeSize);
-
-    vector<char> code(codeSize);
-    hiprtcGetCode(prog, code.data());
-
-    hiprtcDestroyProgram(&prog);
-
-    hipModule_t module;
-    hipFunction_t kernel;
-    hipModuleLoadData(&module, code.data());
-    hipModuleGetFunction(&kernel, module, "saxpy");
-
-    size_t n = NUM_THREADS * NUM_BLOCKS;
-    size_t bufferSize = n * sizeof(float);
-
-    float a = 5.1f;
-    unique_ptr<float[]> hX{new float[n]};
-    unique_ptr<float[]> hY{new float[n]};
-    unique_ptr<float[]> hOut{new float[n]};
-
-    for (size_t i = 0; i < n; ++i) {
-        hX[i] = static_cast<float>(i);
-        hY[i] = static_cast<float>(i * 2);
-    }
-
-    hipDeviceptr_t dX, dY, dOut;
-    hipMalloc(&dX, bufferSize);
-    hipMalloc(&dY, bufferSize);
-    hipMalloc(&dOut, bufferSize);
-    hipMemcpyHtoD(dX, hX.get(), bufferSize);
-    hipMemcpyHtoD(dY, hY.get(), bufferSize);
-
-    struct {
-        float a_;
-        hipDeviceptr_t b_;
-        hipDeviceptr_t c_;
-        hipDeviceptr_t d_;
-        size_t e_;
-    } args{a, dX, dY, dOut, n};
-
-    auto size = sizeof(args);
-    void* config[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &args,
-                      HIP_LAUNCH_PARAM_BUFFER_SIZE, &size,
-                      HIP_LAUNCH_PARAM_END};
-
-    hipModuleLaunchKernel(kernel, NUM_BLOCKS, 1, 1, NUM_THREADS, 1, 1,
-                          0, nullptr, nullptr, config);
-    hipMemcpyDtoH(hOut.get(), dOut, bufferSize);
-
-    for (size_t i = 0; i < n; ++i) {
-        if (fabs(a * hX[i] + hY[i] - hOut[i]) > fabs(hOut[i])* 1e-6) { failed("Validation failed."); }
-    }
-
-    hipFree(dX);
-    hipFree(dY);
-    hipFree(dOut);
-
-    hipModuleUnload(module);
-
-    passed();
-}
@@ -21,8 +21,8 @@ THE SOFTWARE.
 */

 /* HIT_START
- * BUILD: %t %s ../../test_common.cpp EXCLUDE_HIP_PLATFORM nvidia
- * TEST: %t EXCLUDE_HIP_PLATFORM amd
+ * BUILD: %t %s ../../test_common.cpp
+ * TEST: %t
 * HIT_END
 */

@@ -30,9 +30,12 @@ THE SOFTWARE.
 #include <thread>
 #include <vector>

-#define THREADS 8
+#define THREADS 2  // threads per core
 #define MAX_NUM_THREADS 512
-#define ITER 100
+#define ITER 5  // total loop number
+
+// 5 loops and 2 threads per core are enough for function verification.
+// You may adjust them for your test purpose.

 extern "C" __global__ void WaitKernel(int *Ad, int clockrate) {
  uint64_t wait_t = 500,
@@ -73,7 +76,9 @@ int main(int argc, char* argv[]) {
  hipEventCreate(&start);
  std::thread t[NUM_THREADS];

+  printf("NUM_THREADS=%d\n", NUM_THREADS);
  for (int i = 0; i < ITER; i++) {
+    printf("loop %d/%d\n", i, ITER);
    for (int j = 0; j < NUM_THREADS; j++) {
       t[j] = std::thread(t1, start, stream1, clkRate, A[j], Ad[j]);
    }
@@ -75,11 +75,17 @@ bool MultiChunkMultiDevice(int NumDevices) {
  const unsigned threadsPerBlock = 256;
  const unsigned blocks = (NUM_ELMS + 255)/256;
  for (int Klaunch = 0; Klaunch < NumDevices; ++Klaunch) {
+
+    // If without setting device, Hmm value will be read as 0 in kernel on
+    // GPU where Hmm isn't allocated by hipMallocManaged(). This looks like
+    // a bug of cuda. The following line is to fix the bug on cuda only.
+    HIPCHECK(hipSetDevice(Klaunch));
+
    vector_sum<float> <<<blocks, threadsPerBlock, 0, stream[Klaunch]>>>
                      (&Hmm[Klaunch * NUM_ELMS], Ad[Klaunch], NUM_ELMS);
  }
-  HIPCHECK(hipDeviceSynchronize());
  for (int m = 0; m < NumDevices; ++m) {
+    HIPCHECK(hipStreamSynchronize(stream[m]));
    HIPCHECK(hipMemcpy(Ah, Ad[m], NUM_ELMS * sizeof(float),
                       hipMemcpyDeviceToHost));
    for (int n = 0; n < NUM_ELMS; ++n) {
@@ -4,7 +4,7 @@

 //Enable test when compiler support is available in mainline
 /* HIT_START
- * BUILD: %t %s ../../test_common.cpp EXCLUDE_HIP_PLATFORM amd
+ * BUILD: %t %s ../../test_common.cpp
 * HIT_END
 */
 #define N 1048576
@@ -125,8 +125,20 @@ bool MemcpyAtoH<T>::hipMemcpyAtoH_PeerDeviceContext() {
      printf("Skipped the test as there is no peer access\n");
    } else {
      HIPCHECK(hipSetDevice(0));
+
+      unsigned int flags = 0;
+      HIPCHECK(hipGetDeviceFlags(&flags));
+
      AllocateMemory();
      HIPCHECK(hipSetDevice(1));
+
+      // hipMemcpyAtoH will invoke cuda driver api cuMemcpyAtoH() which need
+      // the primary context for device 1. The primary context can be
+      // initialized at the first call of a runtime api through hipSetDeviceFlags().
+      // Because of no runtime api called before cuMemcpyAtoH(), we have to
+      // explicitly call hipSetDeviceFlags().
+      HIPCHECK(hipSetDeviceFlags(flags));  // Only cuda driver api need this
+
      HIPCHECK(hipMemcpyAtoH(B_h, A_d, 0, BYTE_COUNT*sizeof(T)));
      TestPassed = ValidateResult(B_h, hData[0]);
      DeAllocateMemory();
@@ -133,8 +133,20 @@ bool MemcpyHtoA<T>::hipMemcpyHtoA_PeerDeviceContext() {
      printf("Skipped the test as there is no peer access\n");
    } else {
      HIPCHECK(hipSetDevice(0));
+
+      unsigned int flags = 0;
+      HIPCHECK(hipGetDeviceFlags(&flags));
+
      AllocateMemory();
      HIPCHECK(hipSetDevice(1));
+
+      // hipMemcpyAtoH will invoke cuda driver api hipMemcpyHtoA() which need
+      // the primary context for device 1. The primary context can be
+      // initialized at the first call of a runtime api through hipSetDeviceFlags().
+      // Because of no runtime api called before hipMemcpyHtoA(), we have to
+      // explicitly call hipSetDeviceFlags().
+      HIPCHECK(hipSetDeviceFlags(flags));  // Only cuda driver api need this
+
      HIPCHECK(hipMemcpyHtoA(A_d, 0, B_h, BYTECOUNT*sizeof(T)));
      HIPCHECK(hipMemcpy2DFromArray(A_h, sizeof(T)*NUM_W, A_d,
      0, 0, sizeof(T)*NUM_W, 1, hipMemcpyDeviceToHost));
@@ -24,7 +24,7 @@ THE SOFTWARE.

 /* HIT_START
 * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS --std=c++11
- * TEST: %t EXCLUDE_HIP_PLATFORM all
+ * TEST: %t
 * HIT_END

 */
@@ -21,8 +21,8 @@ THE SOFTWARE.
 */

 /* HIT_START
- * BUILD_CMD: managed_kernel.code %hc --genco %S/managed_kernel.cpp -o managed_kernel.code EXCLUDE_HIP_PLATFORM amd
- * BUILD: %t %s ../../test_common.cpp EXCLUDE_HIP_PLATFORM amd
+ * BUILD_CMD: managed_kernel.code %hc --genco %S/managed_kernel.cpp -o managed_kernel.code
+ * BUILD: %t %s ../../test_common.cpp
 * TEST: %t
 * HIT_END
 */
@@ -22,7 +22,7 @@ THE SOFTWARE.
 // This test case is disabled currently.

 /* HIT_START
- * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS --std=c++11 EXCLUDE_HIP_PLATFORM all
+ * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS --std=c++11
 * TEST: %t
 * HIT_END
 */
@@ -11,7 +11,7 @@
 #include "test_common.h"

 /* HIT_START
- * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS -std=c++11 EXCLUDE_HIP_PLATFORM 
+ * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS -std=c++11
 * TEST: %t
 * HIT_END
 */
@@ -18,7 +18,7 @@ THE SOFTWARE.
 */

 /* HIT_START
- * BUILD: %t %s ../../test_common.cpp EXCLUDE_HIP_PLATFORM all
+ * BUILD: %t %s ../../test_common.cpp
 * TEST: %t
 * HIT_END
 */
@@ -55,7 +55,7 @@ __global__ void memcpy_kernel(T* dst, T* src, size_t n)
 }

 template <typename T>
-void runTest()
+int runTest()
 {
    size_t size = NUMITERS*MEMCPYSIZE;

@@ -71,13 +71,16 @@ void runTest()
    printf("HIP stream priority range - low: %d to high: %d\n", priority_low, priority_high);

    // Check if priorities are indeed supported
-    if ((priority_low - priority_high) == 0) { passed(); }
-    
+    if (priority_low == 0 && priority_high == 0) {
+      printf("The device doesn't support stream priorities\n");
+      passed();
+    }
+
    // Enable/disable priorities based on number of available priority levels
    enable_priority_low = true;
    enable_priority_high = true;
    if ((priority_low - priority_high) > 1) enable_priority_normal = true;
-    if (enable_priority_normal) priority_normal = ((priority_low - priority_high) / 2);
+    if (enable_priority_normal) priority_normal = ((priority_low + priority_high) / 2);

    // create streams with highest and lowest available priorities
    #define OP(x) \
@@ -216,5 +219,5 @@ void runTest()
 int main(int argc, char **argv)
 {
    HipTest::parseStandardArguments(argc, argv, false);
-    runTest<int>();
+    return runTest<int>();
 }