Hi everyone,
Environment:
Driver version: 595.71.05
CUDA version: 13.2
Red Hat 8
NVIDIA A30
I also reproduced the issue on Jetson Orin with the latest JetPack.
While trying to integrate my pipeline into a CUDA graph, I encountered an issue when updating the parameters of certain nodes, in particular a copy node initially created by capturing a cudaMemcpy2DAsync() call.
In my case, cudaGraphExecMemcpyNodeSetParams() returns cudaErrorInvalidValue when used on a node originating from cudaMemcpy2DAsync(). My understanding is that cudaMemcpy2DAsync() should produce a memcpy node, just like a regular cudaMemcpy, and therefore should be supported.
I’m including below a small sample that reproduces the issue.
It merely create two graphs, each with a memcpy node created either using CudaMemcpyAsync() or cudaMemcpy2DASync(), and then try to update node’s parameters.
#include <cuda_runtime.h>
#include <cstdlib>
#include <cstring>
#include <iostream>
#include <vector>
#define CHECK_CUDA(call) \
do { \
cudaError_t err__ = (call); \
if (err__ != cudaSuccess) { \
std::cerr << "CUDA error at " << __FILE__ << ":" << __LINE__ \
<< " -> " << cudaGetErrorString(err__) \
<< " (" << static_cast<int>(err__) << ")" << std::endl; \
return false; \
} \
} while (0)
static bool testGraphMemcpyUpdates()
{
constexpr size_t N = 1024;
constexpr size_t NEW_N = 512;
constexpr size_t WIDTH = 128; // elements
constexpr size_t HEIGHT = 8; // rows
constexpr size_t NEW_WIDTH = 32;
float* d_src = nullptr;
float* d_dst = nullptr;
cudaStream_t stream = nullptr;
cudaGraph_t graph1 = nullptr;
cudaGraph_t graph2 = nullptr;
cudaGraphExec_t exec1 = nullptr;
cudaGraphExec_t exec2 = nullptr;
bool ok = false;
CHECK_CUDA(cudaMalloc(&d_src, N * sizeof(float)));
CHECK_CUDA(cudaMalloc(&d_dst, N * sizeof(float)));
CHECK_CUDA(cudaMemset(d_src, 0, N * sizeof(float)));
CHECK_CUDA(cudaMemset(d_dst, 0, N * sizeof(float)));
CHECK_CUDA(cudaStreamCreate(&stream));
// ------------------------------------------------------------
// Graph 1: capture a single cudaMemcpyAsync
// ------------------------------------------------------------
CHECK_CUDA(cudaStreamBeginCapture(stream, cudaStreamCaptureModeGlobal));
CHECK_CUDA(cudaMemcpyAsync(
d_dst,
d_src,
N * sizeof(float),
cudaMemcpyDeviceToDevice,
stream));
CHECK_CUDA(cudaStreamEndCapture(stream, &graph1));
// ------------------------------------------------------------
// Graph 2: capture a single cudaMemcpy2DAsync
// ------------------------------------------------------------
CHECK_CUDA(cudaStreamBeginCapture(stream, cudaStreamCaptureModeGlobal));
CHECK_CUDA(cudaMemcpy2DAsync(
d_dst,
WIDTH * sizeof(float), // dpitch
d_src,
WIDTH * sizeof(float), // spitch
WIDTH * sizeof(float), // width in bytes
HEIGHT, // height
cudaMemcpyDeviceToDevice,
stream));
CHECK_CUDA(cudaStreamEndCapture(stream, &graph2));
// Instantiate both graphs
CHECK_CUDA(cudaGraphInstantiate(&exec1, graph1, nullptr, nullptr, 0));
CHECK_CUDA(cudaGraphInstantiate(&exec2, graph2, nullptr, nullptr, 0));
// ------------------------------------------------------------
// Get nodes
// ------------------------------------------------------------
size_t numNodes = 1;
cudaGraphNode_t node1 = nullptr;
cudaGraphNode_t node2 = nullptr;
CHECK_CUDA(cudaGraphGetNodes(graph1, &node1, &numNodes));
CHECK_CUDA(cudaGraphGetNodes(graph2, &node2, &numNodes));
// ------------------------------------------------------------
// Try update for cudaMemcpyAsync-captured node
// Change the copy size from N to NEW_N
// ------------------------------------------------------------
{
cudaMemcpy3DParms p{};
CHECK_CUDA(cudaGraphMemcpyNodeGetParams(node1, &p));
// NOTE: Note even necessary
/*p.srcPtr = make_cudaPitchedPtr(d_src, N * sizeof(float), N, 1);
p.dstPtr = make_cudaPitchedPtr(d_dst, N * sizeof(float), N, 1);
p.extent = make_cudaExtent(NEW_N * sizeof(float), 1, 1);
p.kind = cudaMemcpyDeviceToDevice;*/
cudaError_t err = cudaGraphExecMemcpyNodeSetParams(exec1, node1, &p);
if (err != cudaSuccess) {
std::cerr << "Update failed for cudaMemcpyAsync graph: "
<< cudaGetErrorString(err)
<< " (" << static_cast<int>(err) << ")" << std::endl;
return false;
} else {
std::cout << "Update succeeded for cudaMemcpyAsync graph." << std::endl;
}
}
// ------------------------------------------------------------
// Try update for cudaMemcpy2DAsync-captured node
// Change width from WIDTH to NEW_WIDTH
// ------------------------------------------------------------
{
cudaMemcpy3DParms p{};
CHECK_CUDA(cudaGraphMemcpyNodeGetParams(node2, &p));
// NOTE: Note even necessary
/*p.srcPtr = make_cudaPitchedPtr(d_src, WIDTH * sizeof(float), WIDTH, HEIGHT);
p.dstPtr = make_cudaPitchedPtr(d_dst, WIDTH * sizeof(float), WIDTH, HEIGHT);
p.extent = make_cudaExtent(NEW_WIDTH * sizeof(float), HEIGHT, 1);
p.kind = cudaMemcpyDeviceToDevice;*/
cudaError_t err = cudaGraphExecMemcpyNodeSetParams(exec2, node2, &p);
if (err != cudaSuccess) {
std::cerr << "Update failed for cudaMemcpy2DAsync graph: "
<< cudaGetErrorString(err)
<< " (" << static_cast<int>(err) << ")" << std::endl;
return false;
} else {
std::cout << "Update succeeded for cudaMemcpy2DAsync graph." << std::endl;
}
}
ok = true;
// Cleanup
if (exec1) cudaGraphExecDestroy(exec1);
if (exec2) cudaGraphExecDestroy(exec2);
if (graph1) cudaGraphDestroy(graph1);
if (graph2) cudaGraphDestroy(graph2);
if (stream) cudaStreamDestroy(stream);
if (d_src) cudaFree(d_src);
if (d_dst) cudaFree(d_dst);
return ok;
}
int main()
{
bool ok = testGraphMemcpyUpdates();
std::cout << "Result = " << (ok ? "true" : "false") << std::endl;
return ok ? EXIT_SUCCESS : EXIT_FAILURE;
}
Any help would be much appreciated, thanks!