Discussion :

[CeluiQuiCode]

Bonjour / Bonsoir,

je suis actuellement en train de travailler sur un logiciel d'imagerie holographique. Ce logiciel existe depuis Septembre 2014 et je ne suis pas le premier dev a passer dessus.
Il est codé en C++ et utilise principalement Qt, OpenGL et Cuda.

Ma premiere tache sur ce logiciel consiste a une refonte des Widgets OpenGL, la seule classe existante est un fourre tout de fonctionnalité, différents types de rendu, fonctions de zoom, d'events, etc ...
Et ma plus grosse frayeur fut la vision des fonctions glBegin() et glEnd()

Le premier changement fut de mettre en place une classe mere afin de derivée vers differentes classes fille pour eviter le fourre tout.
Le deuxieme fut de passer d'une classe derivée de QGLWidget (deprecated) à une classe derivée de QOpenGLWidget.
Le troisieme fut de passer du pipeline fixe au pipeline programmable d'OpenGL.

Pour ceux qui ne connaissent pas les widgets opengl de qt, ils ont 3 fonctions a implementer pour son bon fonctionnement : initializeGL(), paintGL() et resizeGL(). Je ne parlerai pas de resizeGL().

Voila le code des deux fonctions de la "vielle" classe :

initializeGL() :

Code :

Sélectionner tout - Visualiser dans une fenêtre à part

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
 
void GLWidget::initializeGL()
{
	initializeOpenGLFunctions();
	glClearColor(0.0f, 0.0f, 0.0f, 1.0f);
	glEnable(GL_TEXTURE_2D);
 
	glGenBuffers(1, &buffer_);
	glBindBuffer(GL_TEXTURE_BUFFER, buffer_);
	unsigned int size = frame_desc_.frame_size();
	if (frame_desc_.depth == 4 || frame_desc_.depth == 8)
		size /= 2;
 
	glBufferData(GL_TEXTURE_BUFFER, size, nullptr, GL_DYNAMIC_DRAW);
	glBindBuffer(GL_TEXTURE_BUFFER, 0);
	cudaGraphicsGLRegisterBuffer(
		&cuda_buffer_,
		buffer_,
		cudaGraphicsMapFlags::cudaGraphicsMapFlagsNone);
}

paintGL() :

Code :

Sélectionner tout - Visualiser dans une fenêtre à part

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
 
void GLWidget::paintGL()
{
	glEnable(GL_TEXTURE_2D);
	glClear(GL_COLOR_BUFFER_BIT);
 
	const void* frame = queue_.get_last_images(1);
 
	cudaGraphicsMapResources(1, &cuda_buffer_, cuda_stream_);
	size_t	buffer_size;
	void*	buffer_ptr;
	cudaGraphicsResourceGetMappedPointer(&buffer_ptr, &buffer_size, cuda_buffer_);
 
	cudaMemcpy(buffer_ptr, frame, buffer_size, cudaMemcpyKind::cudaMemcpyDeviceToDevice);
 
	cudaGraphicsUnmapResources(1, &cuda_buffer_, cuda_stream_);
 
	glBindBuffer(GL_PIXEL_UNPACK_BUFFER, buffer_);
 
	if (frame_desc_.endianness == camera::BIG_ENDIAN)
		glPixelStorei(GL_UNPACK_SWAP_BYTES, GL_TRUE);
	else
		glPixelStorei(GL_UNPACK_SWAP_BYTES, GL_FALSE);
 
	auto depth = GL_UNSIGNED_SHORT;
	if (frame_desc_.depth == 1)
		depth = GL_UNSIGNED_BYTE;
 
	auto kind = GL_RED;
	if (frame_desc_.depth == 8)
		kind = GL_RG;
 
	glTexImage2D(GL_TEXTURE_2D, 0, kind, frame_desc_.width, frame_desc_.height, 0, kind, depth, nullptr);
	glGenerateMipmap(GL_TEXTURE_2D);
	glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_WRAP_S, GL_CLAMP_TO_BORDER);
	glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_WRAP_T, GL_CLAMP_TO_BORDER);
	glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MAG_FILTER, GL_NEAREST);
	glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MIN_FILTER, GL_LINEAR_MIPMAP_LINEAR);
	glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_SWIZZLE_G, GL_RED);
	glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_SWIZZLE_B, GL_RED);
 
	glBindBuffer(GL_PIXEL_UNPACK_BUFFER, 0);
 
	glBegin(GL_QUADS);
	glColor4f(1.0f, 1.0f, 1.0f, 1.0f);
	glTexCoord2d(0.0 + px_, 0.0 + py_); glVertex2d(-1.0, +1.0);
	glTexCoord2d(1.0 + px_, 0.0 + py_); glVertex2d(+1.0, +1.0);
	glTexCoord2d(1.0 + px_, 1.0 + py_); glVertex2d(+1.0, -1.0);
	glTexCoord2d(0.0 + px_, 1.0 + py_); glVertex2d(-1.0, -1.0);
	glEnd();
 
	glDisable(GL_TEXTURE_2D);
}

Avec ce logiciel, c'etait la premiere fois que je touchais du Cuda pour de vrai, et je me suis donc plongé dans la doc pour savoir keskecé que ces fonctions :
* cudaGraphicsGLRegisterBuffer()
* cudaGraphicsMapResources() & cudaGraphicsUnmapResources()
* cudaGraphicsResourceGetMappedPointer()

Je me suis rendu compte que cudaGraphicsGLRegisterBuffer() et cudaGraphicsResourceGetMappedPointer() sont utilisées uniquement pour des Buffer Objects. Dans notre cas, un Pixel Buffer Object.
J'ai besoin de faire la meme chose mais pour des textures. Et j'ai donc utilisé cudaGraphicsGLRegisterImage() et cudaGraphicsSubResourceGetMappedArray().

Voici mes deux fonctions :

initializeGL():

Code :

Sélectionner tout - Visualiser dans une fenêtre à part

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
 
void SliceWidget::initializeGL()
{
	makeCurrent();
	initializeOpenGLFunctions();
	glClearColor(0.0f, 0.0f, 0.0f, 1.0f);
	glClear(GL_COLOR_BUFFER_BIT);
 
	Vertex = new QOpenGLShader(QOpenGLShader::Vertex);
	Vertex->compileSourceCode(
		"#version 450\n"
		"layout(location = 0) in vec2 xy;\n"
		"layout(location = 1) in vec2 uv;\n"
		"out vec2	texCoord;\n"
		"void main()"
		"{\n"
		"	texCoord = uv;\n"
		"	gl_Position = vec4(xy, 0.0f, 1.0f);\n"
		"}\n"
	);
	if (!Vertex->isCompiled())
			std::cerr << "[Error] Vertex Shader is not compiled\n";
	Fragment = new QOpenGLShader(QOpenGLShader::Fragment);
	Fragment->compileSourceCode(
		"#version 450\n"
		"in vec2	texCoord;\n"
		"out vec4	out_color;\n"
		"uniform sampler2D	tex;\n"
		"void main()"
		"{\n"
		"	out_color = texture(tex, texCoord).rgba;\n"
		"}\n"
	);
	if (!Fragment->isCompiled())
		std::cerr << "[Error] Fragment Shader is not compiled\n";
 
	Program = new QOpenGLShaderProgram();
	Program->addShader(Vertex);
	Program->addShader(Fragment);
	if (!Program->bind())
		std::cerr << "[Error] " << Program->log().toStdString() << '\n';
 
	if (!Vao.create())
		std::cerr << "[Error] Vao create() fail\n";
	Vao.bind();
 
	glGenTextures(1, &Tex);
	glBindTexture(GL_TEXTURE_2D, Tex);
 
	glTexImage2D(GL_TEXTURE_2D, 0,
		GL_RGBA,
		Fd.width, Fd.height, 0,
		GL_RGBA, GL_UNSIGNED_BYTE, nullptr);
 
	glUniform1i(glGetUniformLocation(Program->programId(), "tex"), 0);
	glGenerateMipmap(GL_TEXTURE_2D);
	glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_WRAP_S, GL_CLAMP_TO_BORDER);
	glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_WRAP_T, GL_CLAMP_TO_BORDER);
	glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MAG_FILTER, GL_NEAREST);
	glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MIN_FILTER, GL_LINEAR_MIPMAP_LINEAR);
 
	glBindTexture(GL_TEXTURE_2D, 0);
	cudaGraphicsGLRegisterImage(&cuResource, Tex, GL_TEXTURE_2D,
		cudaGraphicsMapFlags::cudaGraphicsMapFlagsNone);
 
	const float	data[16] = {
		// Top-left
		-1.0f, 1.0,	//xy coord
		0.0f, 0.0f,	// uv coord
		// Top-right
		1.0f, 1.0f,
		1.0f, 0.0f,
		// Bottom-right
		1.0f, -1.0f,
		1.0f, 1.0f,
		// Bottom-left
		-1.0, -1.0,
		0.0f, 1.0f
	};
	glGenBuffers(1, &Vbo);
	glBindBuffer(GL_ARRAY_BUFFER, Vbo);
	glBufferData(GL_ARRAY_BUFFER, 16 * sizeof(GLfloat), data, GL_STATIC_DRAW);
 
	glEnableVertexAttribArray(0);
	glVertexAttribPointer(0, 2, GL_FLOAT, GL_FALSE, 4 * sizeof(float), 0);
 
	glEnableVertexAttribArray(1);
	glVertexAttribPointer(1, 2, GL_FLOAT, GL_FALSE, 4 * sizeof(float),
		reinterpret_cast<void*>(2 * sizeof(float)));
 
	glDisableVertexAttribArray(1);
	glDisableVertexAttribArray(0);
	glBindBuffer(GL_ARRAY_BUFFER, 0);
 
	const GLuint elements[6] = {
		0, 1, 2,
		2, 3, 0
	};
	glGenBuffers(1, &Ebo);
	glBindBuffer(GL_ELEMENT_ARRAY_BUFFER, Ebo);
	glBufferData(GL_ELEMENT_ARRAY_BUFFER, 6 * sizeof(GLuint), elements, GL_STATIC_DRAW);
 
	Vao.release();
	Program->release();
	glViewport(0, 0, Width, Height);
	doneCurrent();
}

paintGL() :

Code :

Sélectionner tout - Visualiser dans une fenêtre à part

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
 
void SliceWidget::paintGL()
{
	void* frame = HQueue.get_last_images(1);
	makeCurrent();
	glClear(GL_COLOR_BUFFER_BIT);
 
	cudaGraphicsMapResources(1, &cuResource, cuStream);
	cudaArray_t cuArr = nullptr;
 
	cudaGraphicsSubResourceGetMappedArray(&cuArr, cuResource, 0, 0);
	cudaMemcpyToArray(cuArr, 0, 0, frame, Fd.frame_size(), cudaMemcpyDeviceToDevice);
	cudaGraphicsUnmapResources(1, &cuResource, cuStream);
 
	glBindTexture(GL_TEXTURE_2D, Tex);
	Program->bind();
	Vao.bind();
 
	glEnableVertexAttribArray(0);
	glEnableVertexAttribArray(1);
 
	glDrawElements(GL_TRIANGLES, 6, GL_UNSIGNED_INT, 0);
 
	glDisableVertexAttribArray(1);
	glDisableVertexAttribArray(0);
	glBindTexture(GL_TEXTURE_2D, 0);
 
	Vao.release();
	Program->release();
	doneCurrent();
}

Forcement, ma version ne fonctionne pas, ma fenetre est completement noir.
J'ai bien essayer d'insérer une fonction glTexSubImage2D() (avec le bind de la texture juste avant) apres la cudaGraphicsUnmapResources(), mais non, un segfault sauvage est apparu
J'ai aussi vu qu'il y avait une methode avec des surface object, j'ai testé vite fait, et rien de vraimenet de probant
Si vous avez besoin de plus de precisions, de voir les shaders, meme si ils sont assez basique (un vertex et un fragment), dites moi, je rajouterai ca.
Edit: Code des fonctions au complet

Y a t il des gens de competents dans la salle ?

Cordialement o/

Edit: Je savais pas vraiment si je devais poster dans les topic OpenGL ou Cuda (ici) vu que les deux sont lié pour mon probleme. Désolé si je me suis planté ^^'

[dragonjoker59]

Salut!

J'ai un doute sur ton utilisation des VBO et vertex attributes.
Si tu utilises un contexte OpenGL 3, il faut avoir un VAO qui va s'occuper du bind de ceux-ci, et il est possible que ton soucis vienne de là (un simple changement de la couleur de fond devrait suffire à confirmer)

[CeluiQuiCode]

Bonjour,

Le contexte OpenGL est en 4.5
J'ai mis le code entier de mes fonctions sur le premier post, ca sera plus simple
J'avais deja fait mes petits test pour voir si mes vbo/vao/shaders fonctionnent correctement. J'ai créé une texture, avec des pixels rgb qui s'alternent, a la main pour confirmer tout ca :

Je me rend de plus en plus compte que j'aurais du faire cette discution dans la partie Cuda du forum

[dragonjoker59]

Ca n'a peut-être pas d'influence, mais comme tu intiialises ta texture avec pour min filter GL_LINEAR_MIPMAP_LINEAR, as-tu essayé en ajoutant un appel à glGenerateMipmaps après l'avoir bindé, lors de l'affichage ?

[CeluiQuiCode]

J'ai deja mis un glGenerateMipmap() dans initializeGL(), juste avant les glTexParameteri(). Cela suffit normalement :o
J'ai testé, au cas où, mais ca ne change rien :/

Je continue mes bidouilles, je vous tiens au courant

Merci pour votre temps !

[dragonjoker59]

Non, ça ne suffit pas, il faut rappeler glGenerateMipmap à chaque fois que le contenu de a texture a changé, pour mettre à jour les mipmaps

[CeluiQuiCode]

Bonsoir,

Merci pour la precision, je laisse le glGenerateMipmap() dans paintGL()

J'ai reussi à sortir un rendu sympa depuis hier :

Avec ceci dans paintGL(), :

Code :

Sélectionner tout - Visualiser dans une fenêtre à part

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
 
	cudaGraphicsMapResources(1, &cuResource, cuStream);
	cudaArray_t cuArr = nullptr;
 
	cudaGraphicsSubResourceGetMappedArray(&cuArr, cuResource, 0, 0);
	cudaResourceDesc cuArrRD;
	{
		cuArrRD.resType = cudaResourceTypeArray;
		cuArrRD.res.array.array = cuArr;
	}
	cudaSurfaceObject_t cuSurface;
	cudaCreateSurfaceObject(&cuSurface, &cuArrRD);
	{
		textureUpdate(cuSurface, frame, Fd.width, Fd.height);
	}
	cudaDestroySurfaceObject(cuSurface);
 
	cudaGraphicsUnmapResources(1, &cuResource, cuStream);
	cudaStreamSynchronize(cuStream);

et voila la partie cuda:

Code :

Sélectionner tout - Visualiser dans une fenêtre à part

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
 
__global__
void kernelTextureUpdate(	unsigned short* frame,
					cudaSurfaceObject_t cuSurface,
					dim3 texDim)
{
	unsigned int x = blockIdx.x * blockDim.x + threadIdx.x;
	unsigned int y = blockIdx.y * blockDim.y + threadIdx.y;
 
	surf2Dwrite(frame[(y * texDim.x + x)], cuSurface, x * 4, y);
}
 
void textureUpdate(	cudaSurfaceObject_t cuSurface,
				void *frame,
				unsigned short width,
				unsigned short height)
{
	dim3 threads(32, 32);
	dim3 blocks(width / threads.x, height / threads.y);
 
	kernelTextureUpdate <<< blocks, threads >>>(reinterpret_cast<unsigned short*>(frame),
		cuSurface, dim3(width, height));
}

Me manque plus qu'à bidouiller mon kernel et comprendre pourquoi j'ai une image juste en rouge et vert.

Edit : Bon, meme pas 5 minutes apres, je me retrouve avec ceci :

Je ne sais pas pour vous, mais je trouve ca stylé !

Voila le kernel modifié :

Code :

Sélectionner tout - Visualiser dans une fenêtre à part

1
2
3
4
5
6
7
8
9
10
11
12
13
14
 
__global__
void kernelTextureUpdate(	unsigned short* frame,
					cudaSurfaceObject_t cuSurface,
					dim3 texDim)
{
	unsigned int x = blockIdx.x * blockDim.x + threadIdx.x;
	unsigned int y = blockIdx.y * blockDim.y + threadIdx.y;
 
	unsigned short pix = frame[(y * texDim.x + x)];
	float pix_f = pix / 65536.f * 256;
 
	surf2Dwrite(static_cast<unsigned short>(pix_f), cuSurface, x * 4, y);
}

Je vais garder ce rendu là, il ressort bien.

Bonne soirée