Jetson orin nx 8Gb
deepstream 6.3
jetpack 5.1.3
tensorRt 8.5
Hello sir,
I am using NVIDIA jetson boards to apply computer vision algorithms on robots and have a serious problem that I have been dealing with this issue for 6 months.
In my previous architecture, i sink coming frames to my C++ code by appsink gstreamer plugin and after that i could process on frames, but this solution causes delay about 8ms. now i could write algorithms in cuda language and wanna run them on frames, my serious problem is that convert frames from CPU after appsink to GPU or cuda memory causes delay again. my frames are in cuda memory (SURFACE_ARRAY (memType = 4)) before sink to code and i could run cuda kernels directly on frames stored in cuda memory and remove sink to C++ code. i used gstdsexample code to do this job from deapstream 6.3 and change transform_ip function like below but i got illegal memory access.
// CUDA kernel launcher
extern "C" void launch_drawRedRectRGBA(void* devPtr,
int width, int height,
int pitch, int thickness);
// Helper macro for CUDA debugging
#define CUDA_CHECK(call) do { \
cudaError_t _status = call; \
if (_status != cudaSuccess) { \
g_printerr("CUDA Error %d: %s at %s:%d\n", _status, cudaGetErrorString(_status), __FILE__, __LINE__); \
} \
} while(0)
// ---------------------------------------------------------------------------
// Simplified gst_dsexample_transform_ip: runs a CUDA kernel on each frame.
// ---------------------------------------------------------------------------
static GstFlowReturn
gst_dsexample_transform_ip(GstBaseTransform *btrans, GstBuffer *inbuf)
{
GstDsExample *dsexample = GST_DSEXAMPLE(btrans);
GstMapInfo in_map_info;
NvBufSurface *surface = NULL;
NvDsBatchMeta *batch_meta = NULL;
NvDsMetaList *l_frame = NULL;
NvDsFrameMeta *frame_meta = NULL;
GstFlowReturn flow_ret = GST_FLOW_ERROR;
CUDA_CHECK(cudaSetDevice(dsexample->gpu_id));
if (!gst_buffer_map(inbuf, &in_map_info, GST_MAP_READWRITE)) {
g_printerr("Error: Failed to map input buffer\n");
return GST_FLOW_ERROR;
}
surface = (NvBufSurface *)in_map_info.data;
if (!surface) {
g_printerr("Error: NvBufSurface not found\n");
goto done;
}
if (CHECK_NVDS_MEMORY_AND_GPUID(dsexample, surface))
goto done;
batch_meta = gst_buffer_get_nvds_batch_meta(inbuf);
if (!batch_meta) {
g_printerr("Error: NvDsBatchMeta not found\n");
goto done;
}
// Process each frame in batch
for (l_frame = batch_meta->frame_meta_list; l_frame != NULL; l_frame = l_frame->next)
{
frame_meta = (NvDsFrameMeta *)(l_frame->data);
guint batch_id = frame_meta->batch_id;
// ------------------------------------------------------------------
// Ensure surface is mapped to CUDA
// ------------------------------------------------------------------
if (surface->memType == NVBUF_MEM_SURFACE_ARRAY) {
if (NvBufSurfaceMap(surface, batch_id, 0, NVBUF_MAP_READ_WRITE) != 0) {
g_printerr("Failed to map surface for CUDA\n");
continue;
}
if (NvBufSurfaceSyncForDevice(surface, batch_id, 0) != 0) {
g_printerr("NvBufSurfaceSyncForDevice failed\n");
NvBufSurfaceUnMap(surface, batch_id, 0);
continue;
}
}
// ------------------------------------------------------------------
// Get CUDA-accessible pointer and parameters
// ------------------------------------------------------------------
int width = surface->surfaceList[batch_id].planeParams.width[0];
int height = surface->surfaceList[batch_id].planeParams.height[0];
int pitch = surface->surfaceList[batch_id].planeParams.pitch[0];
void *devPtr = surface->surfaceList[batch_id].dataPtr;
if (!devPtr) {
g_printerr("Null devPtr for frame %d\n", batch_id);
continue;
}
// ------------------------------------------------------------------
// Launch CUDA kernel
// ------------------------------------------------------------------
launch_drawRedRectRGBA(devPtr, width, height, pitch, 5);
// Check for CUDA errors after kernel
CUDA_CHECK(cudaPeekAtLastError());
CUDA_CHECK(cudaDeviceSynchronize());
// ------------------------------------------------------------------
// Sync back if CPU might access it later
// ------------------------------------------------------------------
if (surface->memType == NVBUF_MEM_SURFACE_ARRAY) {
if (NvBufSurfaceSyncForCpu(surface, batch_id, 0) != 0)
g_printerr("NvBufSurfaceSyncForCpu failed\n");
NvBufSurfaceUnMap(surface, batch_id, 0);
}
}
flow_ret = GST_FLOW_OK;
done:
gst_buffer_unmap(inbuf, &in_map_info);
return flow_ret;
}
and my cuda kernel is so simple:
#include <cuda_runtime.h>
__global__ void drawRedRectRGBA(uchar4* img, int width, int height, int pitch_bytes, int thickness) {
int x = blockIdx.x * blockDim.x + threadIdx.x;
int y = blockIdx.y * blockDim.y + threadIdx.y;
if (x >= width || y >= height) return;
uchar4* row = (uchar4*)((char*)img + y * pitch_bytes); // pitch in bytes
if (x < thickness || x >= width - thickness || y < thickness || y >= height - thickness) {
row[x] = make_uchar4(255, 0, 0, 255); // red
}
}
extern "C" void launch_drawRedRectRGBA(void* devPtr,
int width, int height, int pitch,
int thickness)
{
dim3 block(16, 16);
dim3 grid((width + block.x - 1) / block.x,
(height + block.y - 1) / block.y);
drawRedRectRGBA<<<grid, block>>>((uchar4*)devPtr, width, height, pitch, thickness);
cudaDeviceSynchronize();
}
when i make so file of this plugin and run it with below gstreamer pipeline, i illegal memory access, at this stage frames show correctly but without effect of cuda kernel.
a@a:~/Desktop$ gst-launch-1.0 v4l2src device=/dev/video0 ! “video/x-raw, width=1920, height=1080, framerate=60/1” ! nvvideoconvert compute-hw=2 copy-hw=2 nvbuf-memory-type=4 ! “video/x-raw(memory:NVMM), format=RGBA” ! mux.sink_0 nvstreammux name=mux batch-size=1 width=1920 height=1080 nvbuf-memory-type=4 compute-hw=2 ! dsexample ! nvvideoconvert compute-hw=2 copy-hw=2 nvbuf-memory-type=4 ! fpsdisplaysink sync=false
Setting pipeline to PAUSED …
Pipeline is live and does not need PREROLL …
Setting pipeline to PLAYING …
New clock: GstSystemClock
CUDA Error 700: an illegal memory access was encountered at gstdsexample.cpp:785
CUDA Error 700: an illegal memory access was encountered at gstdsexample.cpp:786
CUDA Error 700: an illegal memory access was encountered at gstdsexample.cpp:785
CUDA Error 700: an illegal memory access was encountered at gstdsexample.cpp:786
CUDA Error 700: an illegal memory access was encountered at gstdsexample.cpp:786
CUDA Error 700: an illegal memory access was encountered at gstdsexample.cpp:785
CUDA Error 700: an illegal memory access was encountered at gstdsexample.cpp:786
^Chandling interrupt.
Interrupt: Stopping pipeline …
Execution ended after 0:00:03.536666880
Setting pipeline to NULL …
^C
In general, I want to run the cuda kernel before the downstream on the frames stored in the Cuda memory. wanna help to write this code or tell me how i must change this codes for applying effect of frames, even simple sample code that what i must do …
TNX alot