SWDEV-283515 - Fix crashing in kernel launch on MGPUS

Fix wrong mixing of current device and stream device in
ihipModuleLaunchKernel() and hipLaunchCooperativeKernel().

Fix missing hipSetDevice() in hipMemcpyWithStream* tests.

Change-Id: I09333bb40d239bb42c832df5ea16d17eeaeff5e7


[ROCm/hip commit: a95ff95bf7]

projects/hip/tests/src/runtimeApi/memory/hipMemcpyWithStream.cpp

@@ -394,9 +394,11 @@ void HipMemcpyWithStreamtests::TestkindDefaultForDtoD(void) {
   int *A_h[numDevices], *B_h[numDevices], *C_h[numDevices];
 
   // Initialize and create the host and device elements for first device
+  HIPCHECK(hipSetDevice(0));
   HipTest::initArrays(&A_d[0], &B_d[0], &C_d[0], &A_h[0], &B_h[0], &C_h[0], N, false);
 
   for (int i=1; i < numDevices; ++i) {
+    HIPCHECK(hipSetDevice(i));
     HIPCHECK(hipMalloc(&A_d[i], Nbytes));
     HIPCHECK(hipMalloc(&B_d[i], Nbytes));
     HIPCHECK(hipMalloc(&C_d[i], Nbytes));
@@ -406,27 +408,25 @@ void HipMemcpyWithStreamtests::TestkindDefaultForDtoD(void) {
 
   hipStream_t stream[numDevices];
   for (int i=0; i < numDevices; ++i) {
+    HIPCHECK(hipSetDevice(i));
     HIPCHECK(hipStreamCreate(&stream[i]));
   }
 
-  HIPCHECK(hipSetDevice(0));
   HIPCHECK(hipMemcpyWithStream(A_d[0], A_h[0], Nbytes, hipMemcpyHostToDevice, stream[0]));
   HIPCHECK(hipMemcpyWithStream(B_d[0], B_h[0], Nbytes, hipMemcpyHostToDevice, stream[0]));
 
   for (int i=1; i < numDevices; ++i) {
-    HIPCHECK(hipSetDevice(i));
     HIPCHECK(hipMemcpyWithStream(A_d[i], A_d[0], Nbytes, hipMemcpyDefault, stream[i]));
     HIPCHECK(hipMemcpyWithStream(B_d[i], B_d[0], Nbytes, hipMemcpyDefault, stream[i]));
   }
 
   for (int i=0; i < numDevices; ++i) {
-    HIPCHECK(hipSetDevice(i));
     hipLaunchKernelGGL(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), 0, stream[i],
                       static_cast<const int*>(A_d[i]), static_cast<const int*>(B_d[i]), C_d[i], N);
   }
 
   for (int i=0; i < numDevices; ++i) {
-    HIPCHECK(hipSetDevice(i));
+    HIPCHECK(hipSetDevice(i)); // hipMemcpy will be on this device
     HIPCHECK(hipStreamSynchronize(stream[i]));
     HIPCHECK(hipMemcpy(C_h[i], C_d[i], Nbytes, hipMemcpyDeviceToHost));
     HipTest::checkVectorADD(A_h[0], B_h[0], C_h[i], N);