Re-sync with upstream.

[ROCm/hip commit: fe959f7bd7]

projects/hip/tests/src/runtimeApi/stream/hipNullStream.cpp

@@ -37,7 +37,7 @@ int p_db = 0;
 using namespace std;
 
 template <typename T>
-__global__ void vectorADDRepeat(hipLaunchParm lp, const T* A_d, const T* B_d, T* C_d, size_t NELEM,
+__global__ void vectorADDRepeat(const T* A_d, const T* B_d, T* C_d, size_t NELEM,
                                 int repeat) {
     size_t offset = (blockIdx.x * blockDim.x + threadIdx.x);
     size_t stride = blockDim.x * gridDim.x;
@@ -117,7 +117,7 @@ void Streamer<T>::enqueAsync() {
     printf("testing: %s  numElements=%zu size=%6.2fMB\n", __func__, _numElements,
            _numElements * sizeof(T) / 1024.0 / 1024.0);
     unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, _numElements);
-    hipLaunchKernel(vectorADDRepeat, dim3(blocks), dim3(threadsPerBlock), 0, _stream,
+    hipLaunchKernelGGL(vectorADDRepeat, dim3(blocks), dim3(threadsPerBlock), 0, _stream,
                     static_cast<const T*>(_A_d), static_cast<const T*>(_B_d), _C_d, _numElements,
                     p_repeat);
 }
@@ -210,7 +210,7 @@ int main(int argc, char* argv[]) {
 
             // Dispatch to NULL stream, should wait for prior async activity to complete before
             // beginning:
-            hipLaunchKernel(vectorADDRepeat, dim3(blocks), dim3(threadsPerBlock), 0,
+            hipLaunchKernelGGL(vectorADDRepeat, dim3(blocks), dim3(threadsPerBlock), 0,
                             0 /*nullstream*/, static_cast<const int*>(lastStreamer->_C_d),
                             static_cast<const int*>(lastStreamer->_C_d), nullStreamer->_C_d,
                             numElements, 1 /*repeat*/);
@@ -246,7 +246,7 @@ int main(int argc, char* argv[]) {
 
             // Dispatch to NULL stream, should wait for prior async activity to complete before
             // beginning:
-            hipLaunchKernel(vectorADDRepeat, dim3(blocks), dim3(threadsPerBlock), 0,
+            hipLaunchKernelGGL(vectorADDRepeat, dim3(blocks), dim3(threadsPerBlock), 0,
                             0 /*nullstream*/, static_cast<const int*>(lastStreamer->_C_d),
                             static_cast<const int*>(lastStreamer->_C_d), nullStreamer->_C_d,
                             numElements, 1 /*repeat*/);