Cannot get node type of conditional nodes with `cudaGraphNodeGetType`

Accelerated Computing CUDA CUDA NVCC Compiler
1

While integrating conditional nodes into a CUDA graph framework for signal processing I discovered that getting the node type of a conditional node with cudaGraphNodeGetType results in a CUDA error cudaErrorUnknown with code 999.

Here is a minimal example:

#include <assert.h>
#include <cuda_runtime.h>
#include <stdexcept>

#define CUDA_CHECK(err)                                                                    \
   do {                                                                                    \
      cudaError_t err_ = (err);                                                            \
      if(err_ != cudaSuccess) {                                                            \
         fprintf(                                                                          \
            stderr,                                                                        \
            "CUDA error at %s:%d code=%d(%s) \"%s\" \n",                                   \
            __FILE__,                                                                      \
            __LINE__,                                                                      \
            static_cast<unsigned int>(err_),                                               \
            cudaGetErrorName(err_),                                                        \
            #err);                                                                         \
         throw std::runtime_error("CUDA error: " + std::string(cudaGetErrorString(err_))); \
      }                                                                                    \
   } while(0)

int main()
{
   cudaGraphNode_t node;
   cudaGraph_t graph;
   CUDA_CHECK(cudaGraphCreate(&graph, 0));

   // Attach a memset node
   int *mem;
   CUDA_CHECK(cudaMallocManaged(&mem, sizeof(int)));
   cudaMemsetParams memsetParams = {};
   memsetParams.dst              = mem;
   memsetParams.pitch            = 0;
   memsetParams.value            = 0;
   memsetParams.elementSize      = sizeof(int);
   memsetParams.width            = 1;
   memsetParams.height           = 1;
   CUDA_CHECK(cudaGraphAddMemsetNode(&node, graph, nullptr, 0, &memsetParams));

   // Check that the node type is `Memset`
   {
      cudaGraphNodeType type;
      CUDA_CHECK(cudaGraphNodeGetType(node, &type));
      assert(type == cudaGraphNodeTypeMemset);
   }

   // Attach a memset node via generic params
   cudaGraphNodeParams genericParamsMemset = {};
   genericParamsMemset.type                = cudaGraphNodeTypeMemset;
   genericParamsMemset.memset.dst          = mem;
   genericParamsMemset.memset.pitch        = 0;
   genericParamsMemset.memset.value        = 0;
   genericParamsMemset.memset.elementSize  = sizeof(int);
   genericParamsMemset.memset.width        = 1;
   genericParamsMemset.memset.height       = 1;
   CUDA_CHECK(cudaGraphAddNode(&node, graph, nullptr, 0, &genericParamsMemset));

   // Check that the node type is still `Memset`
   {
      cudaGraphNodeType type;
      CUDA_CHECK(cudaGraphNodeGetType(node, &type));
      assert(type == cudaGraphNodeTypeMemset);
   }

   // Create conditional handle and attach a conditional node to the graph
   cudaGraphConditionalHandle conditionalHandle;
   CUDA_CHECK(
      cudaGraphConditionalHandleCreate(&conditionalHandle, graph, 1, cudaGraphCondAssignDefault));
   // As of CUDA 12.6 it seems like adding a conditional node is only possible with generic node parameters
   cudaGraphNodeParams genericParamsConditional = {};
   genericParamsConditional.type                = cudaGraphNodeTypeConditional;
   genericParamsConditional.conditional.handle  = conditionalHandle;
   genericParamsConditional.conditional.type    = cudaGraphCondTypeIf;
   genericParamsConditional.conditional.size    = 1;
   CUDA_CHECK(cudaGraphAddNode(&node, graph, nullptr, 0, &genericParamsConditional));

   // Check that the node type is `Conditional`
   {
      cudaGraphNodeType type;
      CUDA_CHECK(cudaGraphNodeGetType(node, &type));
      assert(type == cudaGraphNodeTypeConditional);
   }

   CUDA_CHECK(cudaGraphDestroy(graph));
   CUDA_CHECK(cudaFree(mem));
}

Compiling and executing this on my Ampere GPU with CUDA 12.6:

nvcc conditional_node_type.cu
./a.out

The expected result would be no output, but I get:

CUDA error at conditional_node_type.cu:79 code=999(cudaErrorUnknown) "cudaGraphNodeGetType(node, &type)" 
terminate called after throwing an instance of 'std::runtime_error'
  what():  CUDA error: unknown error
Aborted

Is this strictly not supported or am I doing something wrong here? I’m trying to find out whether a graph has conditional nodes so that I can forbid it being attached as a child to another graph, which is not supported as documented in the driver_types.h:3100 header file (only reference I found about this): “The following restrictions apply to graphs which contain conditional nodes: The graph cannot be used in a child node.”

PS: I also do not see a way to get the graph behind a conditional node using the CUDA runtime API when only the node pointer is known. I did not look into the CUDA driver API yet.