for almost 2 weeks i´m trying to use the cuda opengl interop interface. I want to read from a GL_TEXTURE_2D (the oldColor), do a [newColor = vec4(1.f) - oldColor] operation in cuda and write to a GL_TEXTURE_2D (the newColor). Just as a simple example that can be extended to some further operations.
But all i achieved so far is a lot frustration, sometimes access violation errors but never a working example. All i see when i use the texture that has to be written by cuda, is the old gpu memory content.
I first tried to write a fixed color (rgba - 0 0.6 0.8 1) to a texture, then i wanted to read from a texture and store it instead of the fixed color.
What i did is this:
I created a texture during the init phase (writeTex is a GLuint):
blue(cudaSurfaceObject_t s, dim3 texDim)
{
unsigned int x = blockIdx.x * blockDim.x + threadIdx.x;
unsigned int y = blockIdx.y * blockDim.y + threadIdx.y;
if(x >= texDim.x || y >= texDim.y)
{
return;
}
float4 data = make_float4(.0f, .6f, .8f, 1.0f);
surf2Dwrite(data, s, x * sizeof(float4), y);
}
I don´t get errorValues when checking the error-Variable (e).
So my question is, am i doing it fundamentally wrong or why do i see the old memory content all the time instead of the (0 .6 .8 1) color when i render the texture?
But i wouldn´t be surprised if it is not possible.
I didn´t find a working example regarding this topic so far, but if it isn´t possible at all, the whole cuda approach is a dead end. My plan is to render solid geometry to a texture_3D but if it fails already with texture_2D…
What Compute Capability does your card provide? You are using Surface Objects API, which needs a card with CC 3.5 I guess…
If this is the point you’d have to use Surface Reference API. Otherwise I have no clue.
Thank you for your help so far! Then it seems to be a conceptual misunderstanding, right?
My current state is, that the texture contains the old garbage but in the second rendering iteration i get an access violation when i create a new cudaSurfaceObject_t.
After the first rendering iteration i get to the cudaGraphicsGLRegisterImage(…) (line 3) and it reports error 30 (unknown error). If i step through the code until cudaCreateSurfaceObject(…) (line 26) I get an access violation at 0xccc…
I´m pretty clueless but i suppose it won´t get easier with OpenCL. And i start running out of time because it´s a university project.
I mean the whole if statement but I can limit the problem to line 10. The surf2Dwrite call seems to produce the errors. Cuda-Memcheck reports with active line 10 a varying number of errors (61 in the last run).
A single cudaGraphicsSubResourceGetMappedArray hits error 30 message and a large list of invalid global read of size 4 in the kernel. All of them try to access address 0x40100008 which is out of bounds.
The interesting part is that the access is independent of the blocks. The threads x component seems to follow a pattern (29, 28, 27, 26 and so on) but the y component of the threads is fix to 29. Would a dump in a textfile be helpful?
Dunno will it help (may be a bit of off-topic), but when I want to write stuff to a texture I use pixel buffer object ‘trick’ - create it and register in CUDA, map it as kernel input/output, do the kernel work on it, and finally copy the outcome from PBO to bound GL texture.
You can write to a OpenGL texture directly without the need of a pixel buffer object if you do it like described above. The texture is generated as a GL_RGBA texture containing GL_FLOAT values. That is why I tried to write a float4 value into the texture. It turns out that the OpenGL texture seems to assume 4 byte for each color so writing 4 floats (or 16 bytes) into the texture is just wrong. Although i assumed the texture to be semantically “4 floats each color channel” i had to write a uchar4 (4 byte) to each pixel (see code below). Unfortunately i didn´t find any documentation about it. But i suppose i should have known sigh
__global__ void blue(cudaSurfaceObject_t target, dim3 dimension)
{
unsigned int x = blockIdx.x * blockDim.x + threadIdx.x;
unsigned int y = blockIdx.y * blockDim.y + threadIdx.y;
if(x < dimension.x && y < dimension.y)
{
//float4 data = make_float4(0.f, .6f, .8f, 1.f);
//surf2Dwrite(data, target, x * sizeof(float4), y);
uchar4 data = make_uchar4(0x00, 0x99, 0xcc, 0xff);
surf2Dwrite(data, target, x * sizeof(uchar4), y);
}
}
And do You recon is it faster then PBO approach? The fact that one need to convert each channel to unsigned byte isn’t that good, though. What about precision? Will it be as good as with PBO trick? Are there any better ways, with current CUDA version?
i didn´t implement the PBO approach so i don´t know if it is faster, but as far as i know it requires a glSubImage2D after the Cuda Kernel to copy the content of the the PBO into the texture. And i still have to write the data into the PBO…
The GL_FLOAT specified by the glTexImage2D is refering to the data type of the data pointer (in this case the nullptr), so it has nothing to do with the internal representation of the texture data. You don´t have to convert the whole channel into ubyte i guess… i forgot about the meaning of ths function parameter (gl has too many parameters, i suppose)
When I´m done with this semester I could do some research about the performance of this solution but i´d expect a better performance with the current solution. My current kernel time for a 1920x1080 texture is about 1.8 ms (read from texture, do a color negation (1 - old) and write to a texture, so 2 texture accesses are required)
Hi,
An other solution to this problem is to put
glTexImage2D(GL_TEXTURE_2D, 0, GL_RGBA32F, width, height, 0, GL_RGBA, GL_FLOAT, nullptr);
instead of
glTexImage2D(GL_TEXTURE_2D, 0, GL_RGBA, width, height, 0, GL_RGBA, GL_FLOAT, nullptr);
and you will be able to use float4 instead of uchar4.
as discussed here, when writing to an openGL resource from CUDA, the cudaGraphicsRegisterFlagsWriteDiscard can be used to indicate a 1-way type of delivery.
This is (at least) applicable to writing an OpenGL resource that has been exposed to CUDA as a linear buffer.
the cudaGraphicsRegisterFlagsWriteSurfaceLoadStore is used when writing to a surface, (or reading from a surface) and such write and read activity is always to a cudaArray, not a linear buffer.
The two cases are not necessarily the same. But I believe you can probably combine those two flags when doing writes (only) to a surface, although I don’t know what affect it would have, if any.
