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

Change-Id: I1d67c00ff313f7f966358aec3c90455ff6d8ecbd
2022-06-23 12:49:42 -04:00
@@ -38,16 +38,36 @@ foreach $arg (@ARGV) {
 }

 my $SCRIPT_DIR=dirname(__FILE__);
+
 if ($HIPCC_USE_PERL_SCRIPT) {
    #Invoke hipcc.pl
    my $HIPCC_PERL=catfile($SCRIPT_DIR, '/hipcc.pl');
-    exec($^X, $HIPCC_PERL, @ARGV);
+    system($^X, $HIPCC_PERL, @ARGV);
 } else {
-    #Invoke hipcc.bin
-    my $HIPCC_BIN=catfile($SCRIPT_DIR, '/hipcc.bin');
+    $BIN_NAME="/hipcc.bin";
+    if ($isWindows) {
+        $BIN_NAME="/hipcc.bin.exe";
+    }
+    my $HIPCC_BIN=catfile($SCRIPT_DIR, $BIN_NAME);
    if ( -e $HIPCC_BIN ) {
-        exec($HIPCC_BIN, @ARGV);
+        #Invoke hipcc.bin
+        system($HIPCC_BIN, @ARGV);
    } else {
        print "hipcc.bin not present; Install HIPCC binaries before proceeding";
+        exit(-1);
    }
 }
+
+# Because of this wrapper we need to check
+# the output of the system command for perl and bin
+# else the failures are ignored and build fails silently
+if ($? == -1) {
+    exit($?);
+}
+elsif ($? & 127) {
+    exit($?);
+}
+else {
+     $CMD_EXIT_CODE = $? >> 8;
+}
+exit($CMD_EXIT_CODE);
@@ -1,2 +1,2 @@
-@set HIPCC="%~dp0/hipcc"
+@set HIPCC="%~dp0hipcc"
@perl %HIPCC% %*
@@ -22,7 +22,6 @@
 # Need perl > 5.10 to use logic-defined or
 use 5.006; use v5.10.1;

-use strict;
 use warnings;

 use File::Basename;
@@ -30,18 +29,38 @@ use File::Spec::Functions 'catfile';

 #TODO: By default select perl script until change incorporated in HIP build script
 my $HIPCONFIG_USE_PERL_SCRIPT = 1;
-
+my $isWindows =  ($^O eq 'MSWin32' or $^O eq 'msys');
 my $SCRIPT_DIR=dirname(__FILE__);
+
 if ($HIPCONFIG_USE_PERL_SCRIPT) {
    #Invoke hipconfig.pl
    my $HIPCONFIG_PERL=catfile($SCRIPT_DIR, '/hipconfig.pl');
-    exec($^X, $HIPCONFIG_PERL, @ARGV);
+    system($^X, $HIPCONFIG_PERL, @ARGV);
 } else {
-    #Invoke hipconfig.bin
-    my $HIPCONFIG_BIN=catfile($SCRIPT_DIR, '/hipconfig.bin');
+    $BIN_NAME="/hipconfig.bin";
+    if ($isWindows) {
+        $BIN_NAME="/hipconfig.bin.exe";
+    }
+    my $HIPCONFIG_BIN=catfile($SCRIPT_DIR, $BIN_NAME);
    if ( -e $HIPCONFIG_BIN ) {
-        exec($HIPCONFIG_BIN, @ARGV);
+        #Invoke hipconfig.bin
+        system($HIPCONFIG_BIN, @ARGV);
    } else {
        print "hipconfig.bin not present; Install HIPCC binaries before proceeding";
+        exit(-1);
    }
 }
+
+# Because of this wrapper we need to check
+# the output of the system command for perl and bin
+# else the failures are ignored and build fails silently
+if ($? == -1) {
+    exit($?);
+}
+elsif ($? & 127) {
+    exit($?);
+}
+else {
+     $CMD_EXIT_CODE = $? >> 8;
+}
+exit($CMD_EXIT_CODE);
@@ -1,2 +1,2 @@
-@set HIPCONFIG="%~dp0/hipconfig"
+@set HIPCONFIG="%~dp0hipconfig"
@perl %HIPCONFIG% %*
@@ -13,7 +13,7 @@ if(HIP_PLATFORM STREQUAL "amd")
        message(FATAL_ERROR "Unexpected HIP_COMPILER:${HIP_COMPILER} is set for HIP_PLATFOR:amd")
    endif()
 elseif(HIP_PLATFORM STREQUAL "nvidia")
-    if(NOT DEFINED HIP_COMPILER OR NOT HIP_COMPILER STREQUAL "nvcc")
+    if(DEFINED HIP_COMPILER AND NOT HIP_COMPILER STREQUAL "nvcc")
        message(FATAL_ERROR "Unexpected HIP_COMPILER: ${HIP_COMPILER} is set for HIP_PLATFORM:nvidia")
    endif()
 else()
@@ -63,6 +63,7 @@ else()
                    OUTPUT_STRIP_TRAILING_WHITESPACE)
 endif()

+
 string(REPLACE "." ";" VERSION_LIST ${HIP_VERSION})
 list(GET VERSION_LIST 0 HIP_VERSION_MAJOR)
 list(GET VERSION_LIST 1 HIP_VERSION_MINOR)
@@ -115,9 +116,11 @@ if (RTC_TESTING)
 endif()
 add_definitions(-DKERNELS_PATH="${CMAKE_CURRENT_SOURCE_DIR}/kernels/")

-file(COPY ./hipTestMain/config DESTINATION ${CMAKE_CURRENT_BINARY_DIR}/hipTestMain)
-file(COPY ./external/Catch2/cmake/Catch2/CatchAddTests.cmake DESTINATION ${CMAKE_CURRENT_BINARY_DIR}/script)
-set(ADD_SCRIPT_PATH ${CMAKE_CURRENT_BINARY_DIR}/script/CatchAddTests.cmake)
+set(CATCH_BUILD_DIR catch_tests)
+file(COPY ./hipTestMain/config DESTINATION ${CMAKE_CURRENT_BINARY_DIR}/${CATCH_BUILD_DIR}/hipTestMain)
+file(COPY ./external/Catch2/cmake/Catch2/CatchAddTests.cmake
+     DESTINATION ${CMAKE_CURRENT_BINARY_DIR}/${CATCH_BUILD_DIR}/script)
+set(ADD_SCRIPT_PATH ${CMAKE_CURRENT_BINARY_DIR}/${CATCH_BUILD_DIR}/script/CatchAddTests.cmake)

 if (WIN32)
  configure_file(catchProp_in_rc.in ${CMAKE_CURRENT_BINARY_DIR}/catchProp.rc @ONLY)
@@ -150,9 +153,16 @@ message(STATUS "CMAKE HIP ARCHITECTURES: ${CMAKE_HIP_ARCHITECTURES}")
 # This is done due to limitation of rocm_agent_enumerator
 # While building test parallelly, rocm_agent_enumerator can fail and give out an empty target
 # That results in hipcc building the test for gfx803 (the default target)
-if(NOT DEFINED OFFLOAD_ARCH_STR AND EXISTS "${ROCM_PATH}/bin/rocm_agent_enumerator"
+# preference to pass arch -
+# OFFLOAD_ARCH_STR
+# ENV{HCC_AMDGPU_TARGET}
+# rocm_agent_enumerator
+if(NOT DEFINED OFFLOAD_ARCH_STR
+   AND NOT DEFINED ENV{HCC_AMDGPU_TARGET}
+   AND EXISTS "${ROCM_PATH}/bin/rocm_agent_enumerator"
   AND HIP_PLATFORM STREQUAL "amd" AND UNIX)
-    execute_process(COMMAND ${ROCM_PATH}/bin/rocm_agent_enumerator OUTPUT_VARIABLE HIP_GPU_ARCH
+    execute_process(COMMAND ${ROCM_PATH}/bin/rocm_agent_enumerator
+         OUTPUT_VARIABLE HIP_GPU_ARCH
         RESULT_VARIABLE ROCM_AGENT_ENUM_RESULT)
    # Trim out gfx000
    string(REPLACE "gfx000\n" "" HIP_GPU_ARCH ${HIP_GPU_ARCH})
@@ -165,7 +175,6 @@ if(NOT DEFINED OFFLOAD_ARCH_STR AND EXISTS "${ROCM_PATH}/bin/rocm_agent_enumerat
            foreach(_hip_gpu_arch ${HIP_GPU_ARCH_LIST})
                set(OFFLOAD_ARCH_STR " ${OFFLOAD_ARCH_STR} --offload-arch=${_hip_gpu_arch} ")
            endforeach()
-            message(STATUS "Using offload arch string: ${OFFLOAD_ARCH_STR}")
        endif()
    else()
        message(STATUS "ROCm Agent Enumurator found no valid architectures")
@@ -174,22 +183,56 @@ endif()

 if(DEFINED OFFLOAD_ARCH_STR)
    set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} ${OFFLOAD_ARCH_STR} ")
+elseif(DEFINED ENV{HCC_AMDGPU_TARGET})
+    # hipcc pl script appends it to the options
+    set(OFFLOAD_ARCH_STR " --offload-arch=$ENV{HCC_AMDGPU_TARGET}")
+    set(HIP_GPU_ARCH_LIST $ENV{HCC_AMDGPU_TARGET})
 endif()
+message(STATUS "Using offload arch string: ${OFFLOAD_ARCH_STR}")
+
+# prints the catch info to a file
+string(TIMESTAMP _timestamp UTC)
+set(_catchInfo "# Auto-generated by cmake on ${_timestamp} UTC\n")
+set(_catchInfo ${_catchInfo} "HIP_VERSION=${HIP_VERSION}\n")
+set(_catchInfo ${_catchInfo} "HIP_PLATFORM=${HIP_PLATFORM}\n")
+set(_catchInfo ${_catchInfo} "ARCHS=${HIP_GPU_ARCH_LIST}\n")
+file(WRITE ${CMAKE_CURRENT_BINARY_DIR}/${CATCH_BUILD_DIR}/catchInfo.txt ${_catchInfo})

 # Disable CXX extensions (gnu++11 etc)
 set(CMAKE_CXX_EXTENSIONS OFF)

 add_custom_target(build_tests)

+
 # Tests folder
-add_subdirectory(unit)
-add_subdirectory(ABM)
-add_subdirectory(kernels)
-add_subdirectory(hipTestMain)
-add_subdirectory(stress)
-add_subdirectory(TypeQualifiers)
+add_subdirectory(unit ${CATCH_BUILD_DIR}/unit)
+add_subdirectory(ABM ${CATCH_BUILD_DIR}/ABM)
+add_subdirectory(kernels ${CATCH_BUILD_DIR}/kernels)
+add_subdirectory(hipTestMain ${CATCH_BUILD_DIR}/hipTestMain)
+add_subdirectory(stress ${CATCH_BUILD_DIR}/stress)
+add_subdirectory(TypeQualifiers ${CATCH_BUILD_DIR}/TypeQualifiers)
 if(UNIX)
-    add_subdirectory(multiproc)
+   add_subdirectory(multiproc ${CATCH_BUILD_DIR}/multiproc)
 endif()

 cmake_policy(POP)
+
+# packaging the tests
+# make package_test to generate packages for test
+set(BUILD_DIR ${CMAKE_CURRENT_BINARY_DIR}/packages/)
+configure_file(packaging/hip-tests.txt ${BUILD_DIR}/CMakeLists.txt @ONLY)
+if(UNIX)
+add_custom_target(package_test COMMAND ${CMAKE_COMMAND} .
+    COMMAND rm -rf *.deb *.rpm *.tar.gz
+    COMMAND make package
+    COMMAND cp *.deb ${PROJECT_BINARY_DIR}
+    COMMAND cp *.rpm ${PROJECT_BINARY_DIR}
+    COMMAND cp *.tar.gz ${PROJECT_BINARY_DIR}
+    WORKING_DIRECTORY ${BUILD_DIR})
+else()
+file(TO_NATIVE_PATH ${PROJECT_BINARY_DIR} CATCH_BINARY_DIR)
+add_custom_target(package_test COMMAND ${CMAKE_COMMAND} .
+    COMMAND cpack
+    COMMAND copy *.zip ${CATCH_BINARY_DIR}
+    WORKING_DIRECTORY ${BUILD_DIR})
+endif()
@@ -59,6 +59,8 @@
 	   "Unit_hipMemPoolApi_BasicReuse",
 	   "Unit_hipMemPoolApi_Opportunistic",
 	   "Unit_hipMemPoolApi_Default",
-	   "Unit_hipDeviceGetUuid"
+	   "Unit_hipDeviceGetUuid",
+	   "Unit_hipGraphExecEventWaitNodeSetEvent_Negative",
+	   "Unit_hipGraphExecEventWaitNodeSetEvent_SetAndVerifyMemory"
 	]
 }
@@ -0,0 +1,73 @@
+
+# Copyright (c) 2016 - 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.
+cmake_minimum_required(VERSION 3.16.8)
+project(tests)
+install(DIRECTORY @PROJECT_BINARY_DIR@/@CATCH_BUILD_DIR@
+        DESTINATION .
+        USE_SOURCE_PERMISSIONS)
+install(FILES @PROJECT_BINARY_DIR@/CTestTestfile.cmake
+        DESTINATION .)
+
+#############################
+# Packaging steps
+#############################
+set(CPACK_SET_DESTDIR TRUE)
+set(CPACK_INSTALL_PREFIX @ROCM_PATH@/test/hip/)
+set(PKG_NAME catch-@HIP_PLATFORM@)
+set(CPACK_PACKAGE_NAME ${PKG_NAME})
+set(CPACK_PACKAGE_DESCRIPTION_SUMMARY "HIP: Heterogenous-computing Interface for Portability [CATCH TESTS]")
+set(CPACK_PACKAGE_DESCRIPTION "HIP:
+          Heterogenous-computing Interface for Portability [CATCH TESTS]")
+set(CPACK_PACKAGE_VENDOR "Advanced Micro Devices, Inc.")
+set(CPACK_PACKAGE_CONTACT "HIP Support <hip.support@amd.com>")
+set(CPACK_PACKAGE_VERSION @HIP_VERSION_MAJOR@.@HIP_VERSION_MINOR@.@HIP_VERSION_PATCH_GITHASH@)
+
+
+if(NOT WIN32)
+set(CPACK_GENERATOR "TGZ;DEB;RPM" CACHE STRING "Linux package types for catch tests")
+set(CPACK_BINARY_DEB "ON")
+set(CPACK_DEBIAN_FILE_NAME "DEB-DEFAULT")
+set(CPACK_DEBIAN_PACKAGE_PROVIDES "catch")
+set(CPACK_BINARY_RPM "ON")
+set(CPACK_RPM_FILE_NAME "RPM-DEFAULT")
+set(CPACK_RPM_EXCLUDE_FROM_AUTO_FILELIST_ADDITION "/opt")
+if (CPACK_PACKAGE_VERSION MATCHES "local" )
+  #If building locally default value will cause build failure
+  #DEBUG SYMBOL pacaking require SOURCE_DIR to be small
+  set(CPACK_RPM_BUILD_SOURCE_DIRS_PREFIX ${CPACK_INSTALL_PREFIX})
+endif()
+
+set(CPACK_SOURCE_GENERATOR "TGZ")
+# Install license file
+set ( CPACK_RESOURCE_FILE_LICENSE "@CMAKE_CURRENT_LIST_DIR@/../../LICENSE.txt" )
+install(FILES ${CPACK_RESOURCE_FILE_LICENSE} DESTINATION .)
+set(CPACK_RPM_PACKAGE_LICENSE "MIT")
+
+else()
+# windows packaging
+set(CPACK_INSTALL_PREFIX "")
+set(CPACK_GENERATOR "ZIP" CACHE STRING "Windows package types for catch tests")
+set(CPACK_TGZ_FILE_NAME "ZIP-DEFAULT")
+set(CPACK_TEST_ZIP "ON")
+set(CPACK_ZIP_TEST_PACKAGE_NAME "catch")
+endif()
+
+include(CPack)
@@ -37,5 +37,9 @@ TEST_CASE("Unit_hipDeviceGetCacheConfig_FuncTst") {
                (cacheConfig != hipFuncCachePreferEqual)));
  // This code exists to test the dummy implementation of
  // hipDeviceSetCacheConfig.
+#if HT_NVIDIA
  HIP_CHECK(hipDeviceSetCacheConfig(cacheConfig));
+#else
+  REQUIRE_FALSE(hipSuccess == hipDeviceSetCacheConfig(cacheConfig));
+#endif
 }
@@ -62,6 +62,7 @@ set(TEST_SRC
  hipStreamIsCapturing.cc
  hipStreamGetCaptureInfo.cc
  hipStreamEndCapture.cc
+  hipGraphExecEventWaitNodeSetEvent.cc
 )

 hip_add_exe_to_target(NAME GraphsTest
@@ -0,0 +1,313 @@
+/*
+Copyright (c) 2022 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, INCLUDING 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 ANY 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) Create a graph1 with event record nodes as follows:
+    MemcpyH2DNode --> kernel1(x*x) --> event_record_node(event1)
+    Instantiate graph1.
+    Create a graph2 with event record nodes as follows:
+    event_wait_node(event1) --> MemcpyD2HNode.Instantiate graph2.
+    Change the event in event_record_node in graph1 to event2 using
+    hipGraphExecEventRecordNodeSetEvent.
+    Change the event in event_wait_node in graph2 to event2 using
+    hipGraphExecEventWaitNodeSetEvent.Execute graph1 and then graph2.
+    Verify the result matches with x*x.
+ 2) Scenario to verify that hipGraphExecEventWaitNodeSetEvent does not
+    impact the graph and changes only the executable graph.
+    Create an event wait node with event1 and add it to graph. Instantiate
+    the graph to create an executable graph. Change the event in the
+    executable graph to event2. Verify that the event wait node still
+    contains event1.
+ 3) Negative Scenarios
+    - Input executable graph is nullptr.
+    - Input node is nullptr.
+    - Input set event is nullptr.
+    - Input executable graph is uninitialized.
+    - Input node is uninitialized.
+    - Input set event is uninitialized.
+    - Graph does not contain event wait node.
+    - Pass memset node as input node.
+    - Pass event record node as input node.
+*/
+
+#include <hip_test_common.hh>
+#include <hip_test_checkers.hh>
+#include <hip_test_kernels.hh>
+
+#define GRID_DIM 64
+#define BLK_DIM 256
+#define LEN (GRID_DIM * BLK_DIM)
+#define DELAY_IN_MS 2000
+
+/**
+ * Kernel Functions to perform square and introduce delay in device.
+ */
+static __global__ void sqr_ker_func(int* a, int* b, int clockrate) {
+  int tx = hipBlockIdx_x*hipBlockDim_x + hipThreadIdx_x;
+  if (tx < LEN) b[tx] = a[tx]*a[tx];
+  uint64_t wait_t = DELAY_IN_MS,
+  start = clock64()/clockrate, cur;
+  do { cur = clock64()/clockrate - start;}while (cur < wait_t);
+}
+
+/**
+ * Scenario 1: Test to validate setting different events in executable graph.
+ */
+TEST_CASE("Unit_hipGraphExecEventWaitNodeSetEvent_SetAndVerifyMemory") {
+  size_t memsize = LEN*sizeof(int);
+  hipGraph_t graph1, graph2;
+  HIP_CHECK(hipGraphCreate(&graph1, 0));
+  HIP_CHECK(hipGraphCreate(&graph2, 0));
+  // Create events
+  hipEvent_t event1, event2;
+  HIP_CHECK(hipEventCreate(&event1));
+  HIP_CHECK(hipEventCreate(&event2));
+  // Create nodes with event_start and event1_end
+  hipGraphNode_t event_rec;
+  HIP_CHECK(hipGraphAddEventRecordNode(&event_rec, graph1, nullptr, 0,
+                                                    event1));
+  int *inp_h, *inp_d, *out_h, *out_d;
+  // Allocate host buffers
+  inp_h = reinterpret_cast<int*>(malloc(memsize));
+  REQUIRE(inp_h != nullptr);
+  out_h = reinterpret_cast<int*>(malloc(memsize));
+  REQUIRE(out_h != nullptr);
+  // Allocate device buffers
+  HIP_CHECK(hipMalloc(&inp_d, memsize));
+  HIP_CHECK(hipMalloc(&out_d, memsize));
+  // Initialize host buffer
+  for (uint32_t i = 0; i < LEN; i++) {
+    inp_h[i] = i;
+    out_h[i] = 0;
+  }
+  // graph1 creation ...........
+  // Create memcpy and kernel nodes for graph1
+  // MemcpyH2D -> kernel1 -> event_rec
+  hipGraphNode_t memcpyH2D, kernelnode1;
+  hipKernelNodeParams kernelNodeParams1{};
+  HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D, graph1, nullptr, 0, inp_d,
+                inp_h, memsize, hipMemcpyHostToDevice));
+  // Get device clock rate
+  int clkRate = 0;
+  HIPCHECK(hipDeviceGetAttribute(&clkRate, hipDeviceAttributeClockRate, 0));
+  // kernel1
+  void* kernelArgs[] = {&inp_d, &out_d, reinterpret_cast<void *>(&clkRate)};
+  kernelNodeParams1.func = reinterpret_cast<void *>(sqr_ker_func);
+  kernelNodeParams1.gridDim = dim3(GRID_DIM);
+  kernelNodeParams1.blockDim = dim3(BLK_DIM);
+  kernelNodeParams1.sharedMemBytes = 0;
+  kernelNodeParams1.kernelParams = reinterpret_cast<void**>(kernelArgs);
+  kernelNodeParams1.extra = nullptr;
+  HIP_CHECK(hipGraphAddKernelNode(&kernelnode1, graph1, nullptr, 0,
+                                  &kernelNodeParams1));
+  // Create dependencies for graph1
+  HIP_CHECK(hipGraphAddDependencies(graph1, &memcpyH2D,
+                                    &kernelnode1, 1));
+  HIP_CHECK(hipGraphAddDependencies(graph1, &kernelnode1,
+                                    &event_rec, 1));
+  // graph2 creation ...........
+  // waitnode(event1) -> MemcpyD2H
+  hipGraphNode_t event_wait_node, memcpyD2H;
+  HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2H, graph2, nullptr, 0,
+                out_h, out_d, memsize, hipMemcpyDeviceToHost));
+  HIP_CHECK(hipGraphAddEventWaitNode(&event_wait_node, graph2, nullptr, 0,
+                                                            event1));
+  HIP_CHECK(hipGraphAddDependencies(graph2, &event_wait_node,
+                                    &memcpyD2H, 1));
+  // Instantiate graph1 and graph2
+  hipStream_t streamForGraph1, streamForGraph2;
+  hipGraphExec_t graphExec1, graphExec2;
+  HIP_CHECK(hipStreamCreate(&streamForGraph1));
+  HIP_CHECK(hipStreamCreate(&streamForGraph2));
+  HIP_CHECK(hipGraphInstantiate(&graphExec1, graph1, nullptr, nullptr, 0));
+  HIP_CHECK(hipGraphInstantiate(&graphExec2, graph2, nullptr, nullptr, 0));
+  // Launch graph1 and graph2
+  HIP_CHECK(hipGraphLaunch(graphExec1, streamForGraph1));
+  HIP_CHECK(hipGraphLaunch(graphExec2, streamForGraph2));
+  // Wait for graph2 to complete
+  HIP_CHECK(hipStreamSynchronize(streamForGraph2));
+  // Validate output
+  bool btestPassed = true;
+  for (uint32_t i = 0; i < LEN; i++) {
+    if (out_h[i] != (inp_h[i]*inp_h[i])) {
+      btestPassed = false;
+      break;
+    }
+  }
+  REQUIRE(btestPassed == true);
+  // hipGraphExecEventWaitNodeSetEvent() TEST
+  // Change the event at event_wait_node node to event2 and
+  // the event at event_rec node to event2.
+  HIP_CHECK(hipGraphExecEventRecordNodeSetEvent(graphExec1,
+                                            event_rec, event2));
+  HIP_CHECK(hipGraphExecEventWaitNodeSetEvent(graphExec2,
+                                            event_wait_node, event2));
+  // Launch graph1 and graph2
+  HIP_CHECK(hipGraphLaunch(graphExec1, streamForGraph1));
+  HIP_CHECK(hipGraphLaunch(graphExec2, streamForGraph2));
+  // Wait for graph2 to complete
+  HIP_CHECK(hipStreamSynchronize(streamForGraph2));
+  // Validate output
+  btestPassed = true;
+  for (uint32_t i = 0; i < LEN; i++) {
+    if (out_h[i] != (inp_h[i]*inp_h[i])) {
+      btestPassed = false;
+      break;
+    }
+  }
+  REQUIRE(btestPassed == true);
+  // Free resources
+  HIP_CHECK(hipGraphExecDestroy(graphExec1));
+  HIP_CHECK(hipGraphExecDestroy(graphExec2));
+  HIP_CHECK(hipStreamDestroy(streamForGraph1));
+  HIP_CHECK(hipStreamDestroy(streamForGraph2));
+  HIP_CHECK(hipFree(inp_d));
+  HIP_CHECK(hipFree(out_d));
+  free(inp_h);
+  free(out_h);
+  HIP_CHECK(hipEventDestroy(event1));
+  HIP_CHECK(hipEventDestroy(event2));
+  HIP_CHECK(hipGraphDestroy(graph1));
+  HIP_CHECK(hipGraphDestroy(graph2));
+}
+
+/**
+ * Scenario 2: Test to validate setting a different event in an executable
+ * graph does not impact the original graph and nodes.
+ */
+TEST_CASE("Unit_hipGraphExecEventWaitNodeSetEvent_VerifyEventNotChanged") {
+  hipGraph_t graph;
+  HIP_CHECK(hipGraphCreate(&graph, 0));
+  hipEvent_t event1, event2, event_out;
+  HIP_CHECK(hipEventCreate(&event1));
+  HIP_CHECK(hipEventCreate(&event2));
+  hipGraphNode_t eventwait;
+  HIP_CHECK(hipGraphAddEventWaitNode(&eventwait, graph, nullptr, 0,
+                                                          event1));
+  hipGraphExec_t graphExec;
+  HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0));
+  HIP_CHECK(hipGraphExecEventWaitNodeSetEvent(graphExec,
+                                              eventwait, event2));
+  HIP_CHECK(hipGraphEventWaitNodeGetEvent(eventwait, &event_out));
+  // validate set event and get event are same
+  REQUIRE(event1 == event_out);
+  // Instantiate and launch the graph
+  HIP_CHECK(hipGraphExecDestroy(graphExec));
+  HIP_CHECK(hipGraphDestroy(graph));
+  HIP_CHECK(hipEventDestroy(event2));
+  HIP_CHECK(hipEventDestroy(event1));
+}
+
+/**
+ * Scenario 3: Negative and Parameter Tests.
+ */
+TEST_CASE("Unit_hipGraphExecEventWaitNodeSetEvent_Negative") {
+  hipGraph_t graph;
+  HIP_CHECK(hipGraphCreate(&graph, 0));
+  hipEvent_t event1, event2;
+  HIP_CHECK(hipEventCreate(&event1));
+  HIP_CHECK(hipEventCreate(&event2));
+  hipGraphNode_t eventrec, eventwait;
+  HIP_CHECK(hipGraphAddEventRecordNode(&eventrec, graph, nullptr, 0,
+                                                        event1));
+  HIP_CHECK(hipGraphAddEventWaitNode(&eventwait, graph, nullptr, 0,
+                                                        event1));
+  // Create memset
+  constexpr size_t Nbytes = 1024;
+  char *A_d;
+  hipGraphNode_t memset_A;
+  hipMemsetParams memsetParams{};
+  HIP_CHECK(hipMalloc(&A_d, Nbytes));
+  memset(&memsetParams, 0, sizeof(memsetParams));
+  memsetParams.dst = reinterpret_cast<void*>(A_d);
+  memsetParams.value = 0;
+  memsetParams.pitch = 0;
+  memsetParams.elementSize = sizeof(char);
+  memsetParams.width = Nbytes;
+  memsetParams.height = 1;
+
+  hipGraphExec_t graphExec;
+  HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0));
+
+  SECTION("hGraphExec = nullptr") {
+    REQUIRE(hipErrorInvalidValue ==
+        hipGraphExecEventWaitNodeSetEvent(nullptr, eventwait, event2));
+  }
+
+  SECTION("hNode = nullptr") {
+    REQUIRE(hipErrorInvalidValue ==
+        hipGraphExecEventWaitNodeSetEvent(graphExec, nullptr, event2));
+  }
+
+  SECTION("event = nullptr") {
+    REQUIRE(hipErrorInvalidValue ==
+        hipGraphExecEventWaitNodeSetEvent(graphExec, eventwait, nullptr));
+  }
+
+  SECTION("hGraphExec is uninitialized") {
+    hipGraphExec_t graphExec1{};
+    REQUIRE(hipErrorInvalidValue ==
+        hipGraphExecEventWaitNodeSetEvent(graphExec1, eventwait, event2));
+  }
+
+  SECTION("hNode is uninitialized") {
+    hipGraphNode_t dummy{};
+    REQUIRE(hipErrorInvalidValue ==
+        hipGraphExecEventWaitNodeSetEvent(graphExec, dummy, event2));
+  }
+
+  SECTION("event is uninitialized") {
+    hipEvent_t event_dummy{};
+    REQUIRE(hipErrorInvalidValue ==
+        hipGraphExecEventWaitNodeSetEvent(graphExec, eventwait,
+                                          event_dummy));
+  }
+
+  SECTION("event wait node does not exist") {
+    hipGraph_t graph1;
+    HIP_CHECK(hipGraphCreate(&graph1, 0));
+    HIP_CHECK(hipGraphAddMemsetNode(&memset_A, graph1, nullptr, 0,
+                                    &memsetParams));
+    hipGraphExec_t graphExec1;
+    HIP_CHECK(hipGraphInstantiate(&graphExec1, graph1, nullptr, nullptr, 0));
+    REQUIRE(hipErrorInvalidValue ==
+        hipGraphExecEventWaitNodeSetEvent(graphExec1, eventwait, event2));
+    HIP_CHECK(hipGraphExecDestroy(graphExec1));
+    HIP_CHECK(hipGraphDestroy(graph1));
+  }
+
+  SECTION("pass memset node as hNode") {
+    HIP_CHECK(hipGraphAddMemsetNode(&memset_A, graph, nullptr, 0,
+                                    &memsetParams));
+    REQUIRE(hipErrorInvalidValue ==
+        hipGraphExecEventWaitNodeSetEvent(graphExec, memset_A, event2));
+  }
+
+  SECTION("pass event record node as hNode") {
+    REQUIRE(hipErrorInvalidValue ==
+        hipGraphExecEventWaitNodeSetEvent(graphExec, eventrec, event2));
+  }
+
+  HIP_CHECK(hipFree(A_d));
+  HIP_CHECK(hipGraphExecDestroy(graphExec));
+  HIP_CHECK(hipGraphDestroy(graph));
+  HIP_CHECK(hipEventDestroy(event1));
+  HIP_CHECK(hipEventDestroy(event2));
+}
@@ -29,11 +29,11 @@ __device__ float myDeviceGlobalArray[16];


 extern "C" __global__ void hello_world(const float* a, float* b) {
-    int tx = hipThreadIdx_x;
+    int tx = threadIdx.x;
    b[tx] = a[tx];
 }

 extern "C" __global__ void test_globals(const float* a, float* b) {
-    int tx = hipThreadIdx_x;
+    int tx = threadIdx.x;
    b[tx] = a[tx] + myDeviceGlobal + myDeviceGlobalArray[tx % ARRAY_SIZE];
 }
@@ -134,7 +134,7 @@ bool testMultiTargArchCodeObj() {
  std::string CodeObjL1 = "#include \"hip/hip_runtime.h\"\n";
  std::string CodeObjL2 =
       "extern \"C\" __global__ void hello_world(float* a, float* b) {\n";
-  std::string CodeObjL3 = "  int tx = hipThreadIdx_x;\n";
+  std::string CodeObjL3 = "  int tx = threadIdx.x;\n";
  std::string CodeObjL4 = "  b[tx] = a[tx];\n";
  std::string CodeObjL5 = "}";
  // Creating the full code object string
@@ -34,7 +34,7 @@ __device__ int getSquareOfGlobalFloat() {
 }

 extern "C" __global__ void testWeightedCopy(int* a, int* b) {
-  int tx = hipThreadIdx_x;
+  int tx = threadIdx.x;
  b[tx] = deviceGlobalInt1*a[tx] + deviceGlobalInt2 +
  static_cast<int>(deviceGlobalShort) + static_cast<int>(deviceGlobalChar)
  + getSquareOfGlobalFloat();
@@ -31,8 +31,8 @@ __device__ float deviceGlobalFloat;
 extern "C" __global__ void tex2dKernelFloat(float* outputData,
                                       int width, int height) {
 #if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
-  int x = hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x;
-  int y = hipBlockIdx_y * hipBlockDim_y + hipThreadIdx_y;
+  int x = blockIdx.x * blockDim.x + threadIdx.x;
+  int y = blockIdx.y * blockDim.y + threadIdx.y;
  if ((x < width) && (y < width)) {
    outputData[y * width + x] = tex2D(ftex, x, y);
  }
@@ -42,8 +42,8 @@ extern "C" __global__ void tex2dKernelFloat(float* outputData,
 extern "C" __global__ void tex2dKernelInt(int* outputData,
                                          int width, int height) {
 #if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
-  int x = hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x;
-  int y = hipBlockIdx_y * hipBlockDim_y + hipThreadIdx_y;
+  int x = blockIdx.x * blockDim.x + threadIdx.x;
+  int y = blockIdx.y * blockDim.y + threadIdx.y;
  if ((x < width) && (y < width)) {
    outputData[y * width + x] = tex2D(itex, x, y);
  }
@@ -53,8 +53,8 @@ extern "C" __global__ void tex2dKernelInt(int* outputData,
 extern "C" __global__ void tex2dKernelInt16(short* outputData,
                                          int width, int height) {
 #if !defined(__HIP_NO_IMAGE_SUPPORT) || !__HIP_NO_IMAGE_SUPPORT
-  int x = hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x;
-  int y = hipBlockIdx_y * hipBlockDim_y + hipThreadIdx_y;
+  int x = blockIdx.x * blockDim.x + threadIdx.x;
+  int y = blockIdx.y * blockDim.y + threadIdx.y;
  if ((x < width) && (y < width)) {
    outputData[y * width + x] = tex2D(stex, x, y);
  }
@@ -23,6 +23,6 @@ THE SOFTWARE.
 #include "hip/hip_runtime.h"

 extern "C" __global__ void hello_world(float* a, float* b) {
-    int tx = hipThreadIdx_x;
+    int tx = threadIdx.x;
    b[tx] = a[tx];
 }