L0: Graph kernel nodes targeting the same kernel overwrite other one's arguments

[linehill]

The reproducer:

#include <hip/hip_runtime.h>
__global__ void k(int Val) { printf("Val=%d\n", Val); }

int main() {

  int A0 = 123, A1 = 321;
  void *K0Args[] = {&A0};
  void *K1Args[] = {&A1};

  hipKernelNodeParams K0Params;
  K0Params.func = (void *)k;
  K0Params.gridDim = dim3(1);
  K0Params.blockDim = dim3(1);
  K0Params.sharedMemBytes = 0;
  K0Params.kernelParams = K0Args;
  K0Params.extra = nullptr;

  hipKernelNodeParams K1Params = K0Params;
  K1Params.kernelParams = K1Args;

  hipGraph_t Graph;
  (void)hipGraphCreate(&Graph, 0);

  hipGraphNode_t KN0, KN1;
  (void)hipGraphAddKernelNode(&KN0, Graph, nullptr, 0, &K0Params);
  (void)hipGraphAddKernelNode(&KN1, Graph, nullptr, 0, &K1Params);

  hipGraphExec_t GraphExec;
  (void)hipGraphInstantiate(&GraphExec, Graph, nullptr, nullptr, 0);
  (void)hipGraphLaunch(GraphExec, 0);
  (void)hipDeviceSynchronize();

  return 0;
}

Level0 backend produces incorrect output:

$ ./repro 


Val=321
Val=321

OpenCL backend produces the expected output (order may vary):

$ ./repro 


Val=123
Val=321