关于cudaGraphInstantiateFlagUseNodePriority的行为疑问与验证

我对cudaGraphInstantiateFlagUseNodePriority的理解是它用于对CUDA图中的核函数调用进行优先级排序。例如在一个包含三个独立核函数first、second、third的CUDA图中，每个核函数会等待1秒并打印自身名称：

• 通过cudaGraphKernelNodeSetAttribute（属性为cudaLaunchAttributePriority）分别将三个核节点的优先级设为0、2、1，启用cudaGraphInstantiateFlagUseNodePriority执行图时，应按优先级从高到低的顺序执行：second→third→first。

但实际测试中发现，设置优先级后调用cudaGraphKernelNodeGetAttribute获取该属性，返回值始终为0，这与预期不符。请问我的理解是否正确？

以下是验证用测试代码：

#include <cuda_runtime.h>
#include <vector>
#include <cstdio>
#include <chrono>

#define CUDA_CHECK(error)                                                                            \
    {                                                                                              \
        cudaError_t localError = error;                                                             \
        if (localError != cudaSuccess) {                                                             \
           printf("error: '%s'(%d) from %s at %s:%d\n",  cudaGetErrorString(localError),   \
                   localError, #error, __FUNCTION__, __LINE__);                              \
                exit(0);\
        }                                                                                          \
    }

__global__ void first(uint32_t interval, const uint32_t ticks_per_ms) {
  size_t globaltid = blockIdx.x * blockDim.x + threadIdx.x;
  if (globaltid == 0) {
    printf("\nfirst..");
    while (interval--) {
      uint64_t start = clock64();
      while (clock64() - start < ticks_per_ms) {
      }
    }
    printf("first\n");
  }
}
__global__ void second(uint32_t interval, const uint32_t ticks_per_ms) {
  size_t globaltid = blockIdx.x * blockDim.x + threadIdx.x;
  if (globaltid == 0) {
    printf("\nsecond..");
    while (interval--) {
      uint64_t start = clock64();
      while (clock64() - start < ticks_per_ms) {
      }
    }
    printf("second\n");
  }
}
__global__ void third(uint32_t interval, const uint32_t ticks_per_ms) {
  size_t globaltid = blockIdx.x * blockDim.x + threadIdx.x;
  if (globaltid == 0) {
    printf("\nthird..");
    while (interval--) {
      uint64_t start = clock64();
      while (clock64() - start < ticks_per_ms) {
      }
    }
    printf("third\n");
  }
}

void cudaGraphsManual() {
  cudaStream_t streamForGraph;
  cudaGraph_t graph;
  cudaGraphNode_t kernelNode;
  CUDA_CHECK(cudaStreamCreate(&streamForGraph));
  cudaKernelNodeParams kernelNodeParams = {0};
  CUDA_CHECK(cudaGraphCreate(&graph, 0));

  int ticks_per_ms = 0;
  CUDA_CHECK(cudaDeviceGetAttribute(&ticks_per_ms, cudaDevAttrClockRate, 0));
  uint32_t interval = std::chrono::milliseconds(1000).count();

  void *kernelArgs[2] = {&interval,
                         &ticks_per_ms};

  kernelNodeParams.func = (void *)first;
  kernelNodeParams.gridDim = dim3(1, 1, 1);
  kernelNodeParams.blockDim = dim3(1, 1, 1);
  kernelNodeParams.sharedMemBytes = 0;
  kernelNodeParams.kernelParams = kernelArgs;
  kernelNodeParams.extra = NULL;

  CUDA_CHECK(cudaGraphAddKernelNode(&kernelNode, graph, NULL, 0, &kernelNodeParams));
  union cudaKernelNodeAttrValue p1; p1.priority = 0;
  CUDA_CHECK(cudaGraphKernelNodeSetAttribute(kernelNode, cudaLaunchAttributePriority, &p1));
  
  union cudaKernelNodeAttrValue p4;
  CUDA_CHECK(cudaGraphKernelNodeGetAttribute(kernelNode, cudaLaunchAttributePriority, &p4));
  printf("\nPriority : %d\n", p4.priority);
  
  kernelNodeParams.func = (void *)second;
  CUDA_CHECK(cudaGraphAddKernelNode(&kernelNode, graph,  NULL, 0, &kernelNodeParams));
  union cudaKernelNodeAttrValue p2; p2.priority = 2;
  CUDA_CHECK(cudaGraphKernelNodeSetAttribute(kernelNode, cudaLaunchAttributePriority, &p2));

  CUDA_CHECK(cudaGraphKernelNodeGetAttribute(kernelNode, cudaLaunchAttributePriority, &p4));
  printf("\nPriority : %d\n", p4.priority);

  kernelNodeParams.func = (void *)third;
  CUDA_CHECK(cudaGraphAddKernelNode(&kernelNode, graph,  NULL, 0, &kernelNodeParams));
  union cudaKernelNodeAttrValue p3; p3.priority = 1;
  CUDA_CHECK(cudaGraphKernelNodeSetAttribute(kernelNode, cudaLaunchAttributePriority, &p3));
  
  CUDA_CHECK(cudaGraphKernelNodeGetAttribute(kernelNode, cudaLaunchAttributePriority, &p4));
  printf("\nPriority : %d\n", p4.priority);
  
  cudaGraphExec_t graphExec;
  CUDA_CHECK(cudaGraphInstantiateWithFlags(&graphExec, graph, cudaGraphInstantiateFlagUseNodePriority));
  CUDA_CHECK(cudaGraphLaunch(graphExec, streamForGraph));
  CUDA_CHECK(cudaStreamSynchronize(streamForGraph));
  CUDA_CHECK(cudaGraphExecDestroy(graphExec));
  CUDA_CHECK(cudaGraphDestroy(graph));
  CUDA_CHECK(cudaStreamDestroy(streamForGraph));
}

int main(int argc, char **argv) {
  cudaGraphsManual();
  return EXIT_SUCCESS;
}

内容的提问来源于stack exchange，提问作者user9684153