This pretty much prevents Orin from being used on 4K and 8K image pipelines - the overhead of cuGraphicsEGLRegisterImage is too big.
One possible work around is to create a buffer 1 byte in size and call cuGraphicsEGLRegisterImage()/cuGraphicsUnregisterResource() on it, which somehow also flushes the buffer on 4K buffer, but I am not sure whether I can rely on this behavior.
I am changed my code one more time to demonstrate how calling cuGraphicsUnregisterResource causes cache flush on the first buffer. Use
width=3840 height=2160 numFrames=10 unregisterSecondBuffer=1 ./testSharedNvBufSurface
/*
/usr/local/cuda/bin/nvcc -x cu \
-I/usr/local/cuda -I/usr/src/jetson_multimedia_api/include \
-I/usr/src/jetson_multimedia_api/samples/common/algorithm/cuda/ \
-I/usr/local/cuda-11.4/targets/aarch64-linux/include/ \
-L/usr/local/cuda/lib64 -L/usr/lib/aarch64-linux-gnu/tegra/ \
-lcuda -lnvbufsurface -lnvrm_surface -lnvrm_mem -lcudart -lnppif -lnppc -lnppisu -lnppidei \
testSharedNvBufSurface.cpp -o testSharedNvBufSurface
width=3840 height=2160 numFrames=10 ./testSharedNvBufSurface # This throws assertion
or
width=3840 height=2160 numFrames=10 eglRegisterForEveryFrame=1 ./testSharedNvBufSurface # This works, but slow
or
This demonstrate useful bug, when flushing second buffer of size 1 erroneously flushes cache on the first large buffer:
width=3840 height=2160 numFrames=10 unregisterSecondBuffer=1 ./testSharedNvBufSurface
# This throws assertions:
width=3840 height=2160 numFrames=10 cuGraphicsMapResourcesEveryFrame=1 useStream=1 ./testSharedNvBufSurface
width=3840 height=2160 numFrames=10 extraSynchronize=1 ./testSharedNvBufSurface
*/
#include <cassert>
#include "NvBufSurface.h"
#include "NvCudaProc.h"
#include "nvosd.h"
#include "cudaEGL.h"
#include <cuda.h>
#include <cuda_runtime.h>
#include <npp.h>
#include <cudaProfiler.h>
//Use thrust only to compare buffers on GPU using thrust::equal
#include <thrust/host_vector.h>
#include <thrust/device_vector.h>
int main(int argc, char ** argv)
{
bool eglRegisterForEveryFrame = getenv("eglRegisterForEveryFrame") != nullptr;
bool cuGraphicsMapResourcesEveryFrame = getenv("cuGraphicsMapResourcesEveryFrame") != nullptr;
bool useStream = getenv("useStream") != nullptr;
bool extraSynchronize = getenv("extraSynchronize") != nullptr;
bool unregisterSecondBuffer = getenv("unregisterSecondBuffer") != nullptr;
int width = strtol(getenv("width"), nullptr, 10);
int height = strtol(getenv("height"), nullptr, 10);
int numFrames = strtol(getenv("numFrames"), nullptr, 10);
printf("eglRegisterForEveryFrame %d cuGraphicsMapResourcesEveryFrame %d extraSynchronize %d width %d height %d numFrames %d\n",
eglRegisterForEveryFrame, cuGraphicsMapResourcesEveryFrame, extraSynchronize, width, height, numFrames);
int bufferSize { width * height * 4 };//For simplicity let's create a 1-D 1 byte array of the same size as 4K RGB image
assert(CUDA_SUCCESS == cuInit(0));
CUdevice cudaDevice;
assert(CUDA_SUCCESS == cuDeviceGet(&cudaDevice, 0));
CUcontext context;
assert(CUDA_SUCCESS == cuCtxCreate(&context, 0, cudaDevice));
CUstream hStream {0};
if(useStream) {
printf("Create stream\n");
assert( cudaSuccess == cudaStreamCreate(&hStream));
}
NvBufSurface * nvbufSurf[2] {};
char * sharedCpuPtr[2] {};
char * sharedGpuPtr[2] {};
CUgraphicsResource cuGraphicsResource[2] {};
EGLImageKHR eglImage[2] {};
CUeglFrame eglFrame[2] {};
int numBuffers { unregisterSecondBuffer? 2 : 1 };
for(int idx = 0; idx < numBuffers; idx++) {
//Create NvBufSurface shared betweren CPU and GPU
NvBufSurfaceAllocateParams inputParams = {{0}};
inputParams.params.width = idx? 1 : bufferSize;//Second buffer of size 1
inputParams.params.height = 1;
inputParams.params.memType = NVBUF_MEM_SURFACE_ARRAY;
inputParams.params.layout = NVBUF_LAYOUT_PITCH;
inputParams.params.colorFormat = NVBUF_COLOR_FORMAT_GRAY8;
inputParams.memtag = NvBufSurfaceTag_CAMERA;
assert(0 == NvBufSurfaceAllocate(&nvbufSurf[idx], 1, &inputParams));
nvbufSurf[idx]->numFilled = 1;
int dmaBuffer = nvbufSurf[idx]->surfaceList[0].bufferDesc;
printf("NvBufSurfaceAllocate ret nvbufSurf %p dmaBuffer %d\n", nvbufSurf[idx], dmaBuffer);
//Map to CPU:
assert(1 == nvbufSurf[idx]->surfaceList->planeParams.num_planes);
assert(0 == NvBufSurfaceMap(nvbufSurf[idx], 0, 0, NVBUF_MAP_READ_WRITE));
sharedCpuPtr[idx] = (char *)nvbufSurf[idx]->surfaceList->mappedAddr.addr[0];
//Map to GPU:
assert(0 == NvBufSurfaceMapEglImage(nvbufSurf[idx], 0));
eglImage[idx] = nvbufSurf[idx]->surfaceList->mappedAddr.eglImage;
printf("NvBufSurfaceMapEglImage ret eglImage %p\n", eglImage[idx]);
assert(CUDA_SUCCESS == cuGraphicsEGLRegisterImage(&cuGraphicsResource[idx], eglImage[idx], CU_GRAPHICS_MAP_RESOURCE_FLAGS_NONE));
assert(CUDA_SUCCESS == cuGraphicsResourceGetMappedEglFrame(&eglFrame[idx], cuGraphicsResource[idx], 0, 0));
sharedGpuPtr[idx] = (char *)eglFrame[idx].frame.pPitch[0];
printf("Shared NvBufSurface mapped to CPU: %p GPU: %p\n", sharedCpuPtr[idx], sharedGpuPtr[idx]);
}
//Create 2 CPU and 2 GPU buffers:
thrust::host_vector<uint8_t> cpuVectors[2] {
thrust::host_vector<uint8_t>(bufferSize),
thrust::host_vector<uint8_t>(bufferSize) };
thrust::device_vector<uint8_t> gpuVectors[2] {
thrust::device_vector<uint8_t>(bufferSize),
thrust::device_vector<uint8_t>(bufferSize) };
printf("Allocated CPU: %p %p GPU: %p %p\n", cpuVectors[0].data(), cpuVectors[1].data(),
gpuVectors[0].data().get(), (void*)gpuVectors[1].data().get());
thrust::fill(cpuVectors[0].begin(), cpuVectors[0].end(), 0);
thrust::fill(cpuVectors[1].begin(), cpuVectors[1].end(), 0xff);
thrust::fill(gpuVectors[0].begin(), gpuVectors[0].end(), 0);
thrust::fill(gpuVectors[1].begin(), gpuVectors[1].end(), 0xff);
for(int frameIdx = 0; frameIdx < numFrames; frameIdx++) {//Simulate frame handling loop
int buffIdx = {frameIdx & 1};
printf("Copy on CPU\n");
memcpy(sharedCpuPtr[0], cpuVectors[buffIdx].data(), bufferSize);//Copy to shared on CPU
if(extraSynchronize) {
printf("cuCtxSynchronize\n");
assert(0 == cuCtxSynchronize());
}
assert(0 == NvBufSurfaceSyncForDevice(nvbufSurf[0], 0, 0));
if(extraSynchronize) {
printf("cuCtxSynchronize\n");
assert(0 == cuCtxSynchronize());
}
if(eglRegisterForEveryFrame && !cuGraphicsResource[0]) {
printf("Register buffer #0\n");
assert(CUDA_SUCCESS == cuGraphicsEGLRegisterImage(&cuGraphicsResource[0], eglImage[0], CU_GRAPHICS_MAP_RESOURCE_FLAGS_NONE));
assert(CUDA_SUCCESS == cuGraphicsResourceGetMappedEglFrame(&eglFrame[0], cuGraphicsResource[0], 0, 0));
}
if(unregisterSecondBuffer) {
printf("Register buffer #1\n");
assert(CUDA_SUCCESS == cuGraphicsEGLRegisterImage(&cuGraphicsResource[1], eglImage[1], CU_GRAPHICS_MAP_RESOURCE_FLAGS_NONE));
assert(CUDA_SUCCESS == cuGraphicsResourceGetMappedEglFrame(&eglFrame[1], cuGraphicsResource[1], 0, 0));
}
if(extraSynchronize) {
printf("cuCtxSynchronize\n");
assert(0 == cuCtxSynchronize());
}
if(cuGraphicsMapResourcesEveryFrame) {
printf("map\n"); assert( cuGraphicsMapResources(1, &cuGraphicsResource[0], hStream) );
}
if(extraSynchronize) {
printf("cuCtxSynchronize\n");
assert(0 == cuCtxSynchronize());
}
printf("Compare on GPU\n");
assert(true == thrust::equal(thrust::cuda::par.on(hStream),
gpuVectors[buffIdx].begin(), gpuVectors[buffIdx].end(), sharedGpuPtr[0]));
if(extraSynchronize) {
printf("cuCtxSynchronize\n");
assert(0 == cuCtxSynchronize());
}
printf("Copy on GPU\n");
buffIdx = !buffIdx;
if(useStream) {
assert(0 == cudaMemcpyAsync(sharedGpuPtr[0], gpuVectors[buffIdx].data().get(), bufferSize, cudaMemcpyDeviceToDevice, hStream));
}
else {
assert(0 == cudaMemcpy(sharedGpuPtr[0], gpuVectors[buffIdx].data().get(), bufferSize, cudaMemcpyDeviceToDevice));
}
printf("NvBufSurfaceSyncForCpu\n");
assert(0 == NvBufSurfaceSyncForCpu(nvbufSurf[0], 0, 0));
printf("cuCtxSynchronize\n");
assert(0 == cuCtxSynchronize());
if(cuGraphicsMapResourcesEveryFrame) {
printf("unmap\n"); assert( cuGraphicsUnmapResources(1, &cuGraphicsResource[0], hStream) );
}
printf("NvBufSurfaceSyncForCpu\n");
assert(0 == NvBufSurfaceSyncForCpu(nvbufSurf[0], 0, 0));
if(eglRegisterForEveryFrame) {
printf("Unregister buffer #0\n");
cuGraphicsUnregisterResource(cuGraphicsResource[0]); cuGraphicsResource[0] = nullptr;
}
if(unregisterSecondBuffer) {
printf("Unregister buffer #1\n");
cuGraphicsUnregisterResource(cuGraphicsResource[1]); cuGraphicsResource[1] = nullptr;
}
if(extraSynchronize) {
printf("cuCtxSynchronize\n");
assert(0 == cuCtxSynchronize());
}
printf("NvBufSurfaceSyncForCpu\n");
assert(0 == NvBufSurfaceSyncForCpu(nvbufSurf[0], 0, 0));
printf("Compare on CPU\n");
assert( !memcmp( cpuVectors[buffIdx].data(), sharedCpuPtr[0], bufferSize));
}
return true;
}