Open linehill opened 4 months ago
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
The reproducer:
Level0 backend produces incorrect output:
OpenCL backend produces the expected output (order may vary):