e276548b |
1 | // Created on: 2013-08-27 |
2 | // Created by: Denis BOGOLEPOV |
3 | // Copyright (c) 2013 OPEN CASCADE SAS |
4 | // |
973c2be1 |
5 | // This file is part of Open CASCADE Technology software library. |
e276548b |
6 | // |
d5f74e42 |
7 | // This library is free software; you can redistribute it and/or modify it under |
8 | // the terms of the GNU Lesser General Public License version 2.1 as published |
973c2be1 |
9 | // by the Free Software Foundation, with special exception defined in the file |
10 | // OCCT_LGPL_EXCEPTION.txt. Consult the file LICENSE_LGPL_21.txt included in OCCT |
11 | // distribution for complete text of the license and disclaimer of any warranty. |
e276548b |
12 | // |
973c2be1 |
13 | // Alternatively, this file may be used under the terms of Open CASCADE |
14 | // commercial license or contractual agreement. |
e276548b |
15 | |
16 | #ifdef HAVE_CONFIG_H |
17 | #include <config.h> |
18 | #endif |
19 | |
20 | #ifdef HAVE_OPENCL |
21 | |
618d8e61 |
22 | #include <OpenGl_Cl.hxx> |
23 | |
e276548b |
24 | #if defined(_WIN32) |
25 | |
26 | #include <windows.h> |
27 | #include <wingdi.h> |
28 | |
29 | #pragma comment (lib, "DelayImp.lib") |
30 | #pragma comment (lib, "OpenCL.lib") |
31 | |
32 | #elif defined(__APPLE__) && !defined(MACOSX_USE_GLX) |
33 | #include <OpenGL/CGLCurrent.h> |
34 | #else |
35 | #include <GL/glx.h> |
36 | #endif |
37 | |
38 | #include <OpenGl_Context.hxx> |
39 | #include <OpenGl_Texture.hxx> |
40 | #include <OpenGl_View.hxx> |
41 | #include <OpenGl_Workspace.hxx> |
265d4508 |
42 | #include <Standard_Assert.hxx> |
e276548b |
43 | |
44 | using namespace OpenGl_Raytrace; |
45 | |
46 | //! Use this macro to output ray-tracing debug info |
d5af8626 |
47 | //#define RAY_TRACE_PRINT_INFO |
e276548b |
48 | |
49 | #ifdef RAY_TRACE_PRINT_INFO |
50 | #include <OSD_Timer.hxx> |
51 | #endif |
52 | |
53 | //! OpenCL source of ray-tracing kernels. |
54 | extern const char THE_RAY_TRACE_OPENCL_SOURCE[]; |
55 | |
56 | // ======================================================================= |
57 | // function : MatVecMult |
58 | // purpose : Multiples 4x4 matrix by 4D vector |
59 | // ======================================================================= |
265d4508 |
60 | template<typename T> |
61 | BVH_Vec4f MatVecMult (const T m[16], const BVH_Vec4f& v) |
e276548b |
62 | { |
265d4508 |
63 | return BVH_Vec4f ( |
e276548b |
64 | static_cast<float> (m[ 0] * v.x() + m[ 4] * v.y() + |
65 | m[ 8] * v.z() + m[12] * v.w()), |
66 | static_cast<float> (m[ 1] * v.x() + m[ 5] * v.y() + |
67 | m[ 9] * v.z() + m[13] * v.w()), |
68 | static_cast<float> (m[ 2] * v.x() + m[ 6] * v.y() + |
69 | m[10] * v.z() + m[14] * v.w()), |
70 | static_cast<float> (m[ 3] * v.x() + m[ 7] * v.y() + |
71 | m[11] * v.z() + m[15] * v.w())); |
72 | } |
73 | |
74 | // ======================================================================= |
75 | // function : UpdateRaytraceEnvironmentMap |
76 | // purpose : Updates environment map for ray-tracing |
77 | // ======================================================================= |
78 | Standard_Boolean OpenGl_Workspace::UpdateRaytraceEnvironmentMap() |
79 | { |
80 | if (myView.IsNull()) |
81 | return Standard_False; |
82 | |
83 | if (myViewModificationStatus == myView->ModificationState()) |
84 | return Standard_True; |
85 | |
86 | cl_int anError = CL_SUCCESS; |
87 | |
88 | if (myRaytraceEnvironment != NULL) |
89 | clReleaseMemObject (myRaytraceEnvironment); |
90 | |
68333c8f |
91 | Standard_Integer aSizeX = 1; |
92 | Standard_Integer aSizeY = 1; |
e276548b |
93 | |
94 | if (!myView->TextureEnv().IsNull() && myView->SurfaceDetail() != Visual3d_TOD_NONE) |
95 | { |
96 | aSizeX = (myView->TextureEnv()->SizeX() <= 0) ? 1 : myView->TextureEnv()->SizeX(); |
97 | aSizeY = (myView->TextureEnv()->SizeY() <= 0) ? 1 : myView->TextureEnv()->SizeY(); |
98 | } |
99 | |
265d4508 |
100 | cl_image_format anImageFormat; |
e276548b |
101 | |
265d4508 |
102 | anImageFormat.image_channel_order = CL_RGBA; |
103 | anImageFormat.image_channel_data_type = CL_FLOAT; |
e276548b |
104 | |
265d4508 |
105 | myRaytraceEnvironment = clCreateImage2D (myComputeContext, |
106 | CL_MEM_READ_ONLY, &anImageFormat, aSizeX, aSizeY, 0, NULL, &anError); |
e276548b |
107 | |
108 | cl_float* aPixelData = new cl_float[aSizeX * aSizeY * 4]; |
109 | |
110 | // Note: texture format is not compatible with OpenCL image |
111 | // (it's not possible to create image directly from texture) |
112 | |
113 | if (!myView->TextureEnv().IsNull() && myView->SurfaceDetail() != Visual3d_TOD_NONE) |
114 | { |
115 | myView->TextureEnv()->Bind (GetGlContext()); |
116 | |
117 | glGetTexImage (GL_TEXTURE_2D, |
118 | 0, |
119 | GL_RGBA, |
120 | GL_FLOAT, |
121 | aPixelData); |
122 | |
123 | myView->TextureEnv()->Unbind (GetGlContext()); |
124 | } |
125 | else |
126 | { |
68333c8f |
127 | for (Standard_Integer aPixel = 0; aPixel < aSizeX * aSizeY * 4; ++aPixel) |
e276548b |
128 | aPixelData[aPixel] = 0.f; |
129 | } |
130 | |
131 | size_t anImageOffset[] = { 0, |
132 | 0, |
133 | 0 }; |
134 | |
135 | size_t anImageRegion[] = { aSizeX, |
136 | aSizeY, |
137 | 1 }; |
138 | |
265d4508 |
139 | anError |= clEnqueueWriteImage (myComputeQueue, myRaytraceEnvironment, |
140 | CL_TRUE, anImageOffset, anImageRegion, 0, 0, aPixelData, 0, NULL, NULL); |
141 | |
e276548b |
142 | #ifdef RAY_TRACE_PRINT_INFO |
143 | if (anError != CL_SUCCESS) |
144 | std::cout << "Error! Failed to write environment map image!" << std::endl; |
145 | #endif |
146 | |
147 | delete[] aPixelData; |
148 | |
149 | myViewModificationStatus = myView->ModificationState(); |
150 | |
151 | return (anError == CL_SUCCESS); |
152 | } |
153 | |
154 | // ======================================================================= |
155 | // function : UpdateRaytraceGeometry |
156 | // purpose : Updates 3D scene geometry for ray tracing |
157 | // ======================================================================= |
158 | Standard_Boolean OpenGl_Workspace::UpdateRaytraceGeometry (Standard_Boolean theCheck) |
159 | { |
160 | if (myView.IsNull()) |
161 | return Standard_False; |
162 | |
163 | // Note: In 'check' mode the scene geometry is analyzed for modifications |
164 | // This is light-weight procedure performed for each frame |
165 | |
166 | if (!theCheck) |
167 | { |
265d4508 |
168 | myRaytraceGeometry.Clear(); |
e276548b |
169 | |
170 | myIsRaytraceDataValid = Standard_False; |
171 | } |
172 | else |
173 | { |
174 | if (myLayersModificationStatus != myView->LayerList().ModificationState()) |
175 | { |
176 | return UpdateRaytraceGeometry (Standard_False); |
177 | } |
178 | } |
179 | |
265d4508 |
180 | Standard_ShortReal* aTransform (NULL); |
e276548b |
181 | |
182 | // The set of processed structures (reflected to ray-tracing) |
183 | // This set is used to remove out-of-date records from the |
184 | // hash map of structures |
185 | std::set<const OpenGl_Structure*> anElements; |
186 | |
187 | const OpenGl_LayerList& aList = myView->LayerList(); |
188 | |
189 | for (OpenGl_SequenceOfLayers::Iterator anLayerIt (aList.Layers()); anLayerIt.More(); anLayerIt.Next()) |
190 | { |
191 | const OpenGl_PriorityList& aPriorityList = anLayerIt.Value(); |
192 | |
193 | if (aPriorityList.NbStructures() == 0) |
194 | continue; |
195 | |
196 | const OpenGl_ArrayOfStructure& aStructArray = aPriorityList.ArrayOfStructures(); |
197 | |
68333c8f |
198 | for (Standard_Integer anIndex = 0; anIndex < aStructArray.Length(); ++anIndex) |
e276548b |
199 | { |
200 | OpenGl_SequenceOfStructure::Iterator aStructIt; |
201 | |
202 | for (aStructIt.Init (aStructArray (anIndex)); aStructIt.More(); aStructIt.Next()) |
203 | { |
204 | const OpenGl_Structure* aStructure = aStructIt.Value(); |
205 | |
206 | if (theCheck) |
207 | { |
208 | if (CheckRaytraceStructure (aStructure)) |
209 | { |
210 | return UpdateRaytraceGeometry (Standard_False); |
211 | } |
212 | } |
213 | else |
214 | { |
215 | if (!aStructure->IsRaytracable()) |
216 | continue; |
217 | |
218 | if (aStructure->Transformation()->mat != NULL) |
219 | { |
220 | if (aTransform == NULL) |
265d4508 |
221 | aTransform = new Standard_ShortReal[16]; |
e276548b |
222 | |
68333c8f |
223 | for (Standard_Integer i = 0; i < 4; ++i) |
224 | for (Standard_Integer j = 0; j < 4; ++j) |
e276548b |
225 | { |
226 | aTransform[j * 4 + i] = aStructure->Transformation()->mat[i][j]; |
227 | } |
228 | } |
229 | |
230 | AddRaytraceStructure (aStructure, aTransform, anElements); |
231 | } |
232 | } |
233 | } |
234 | } |
235 | |
236 | if (!theCheck) |
237 | { |
238 | // Actualize the hash map of structures -- remove out-of-date records |
239 | std::map<const OpenGl_Structure*, Standard_Size>::iterator anIter = myStructureStates.begin(); |
240 | |
241 | while (anIter != myStructureStates.end()) |
242 | { |
243 | if (anElements.find (anIter->first) == anElements.end()) |
244 | { |
245 | myStructureStates.erase (anIter++); |
246 | } |
247 | else |
248 | { |
249 | ++anIter; |
250 | } |
251 | } |
252 | |
253 | // Actualize OpenGL layer list state |
254 | myLayersModificationStatus = myView->LayerList().ModificationState(); |
255 | |
265d4508 |
256 | // Rebuild bottom-level and high-level BVHs |
257 | myRaytraceGeometry.ProcessAcceleration(); |
e276548b |
258 | |
265d4508 |
259 | const Standard_ShortReal aMinRadius = Max (fabs (myRaytraceGeometry.Box().CornerMin().x()), Max ( |
260 | fabs (myRaytraceGeometry.Box().CornerMin().y()), fabs (myRaytraceGeometry.Box().CornerMin().z()))); |
261 | const Standard_ShortReal aMaxRadius = Max (fabs (myRaytraceGeometry.Box().CornerMax().x()), Max ( |
262 | fabs (myRaytraceGeometry.Box().CornerMax().y()), fabs (myRaytraceGeometry.Box().CornerMax().z()))); |
e276548b |
263 | |
265d4508 |
264 | myRaytraceSceneRadius = 2.f /* scale factor */ * Max (aMinRadius, aMaxRadius); |
e276548b |
265 | |
265d4508 |
266 | myRaytraceSceneEpsilon = Max (1e-4f, |
267 | myRaytraceGeometry.Box().Size().Length() * 1e-4f); |
e276548b |
268 | |
269 | return WriteRaytraceSceneToDevice(); |
270 | } |
271 | |
272 | delete [] aTransform; |
273 | |
274 | return Standard_True; |
275 | } |
276 | |
277 | // ======================================================================= |
278 | // function : CheckRaytraceStructure |
265d4508 |
279 | // purpose : Checks to see if the structure is modified |
e276548b |
280 | // ======================================================================= |
281 | Standard_Boolean OpenGl_Workspace::CheckRaytraceStructure (const OpenGl_Structure* theStructure) |
282 | { |
283 | if (!theStructure->IsRaytracable()) |
284 | { |
285 | // Checks to see if all ray-tracable elements were |
286 | // removed from the structure |
287 | if (theStructure->ModificationState() > 0) |
288 | { |
289 | theStructure->ResetModificationState(); |
290 | return Standard_True; |
291 | } |
292 | |
293 | return Standard_False; |
294 | } |
295 | |
296 | std::map<const OpenGl_Structure*, Standard_Size>::iterator aStructState = myStructureStates.find (theStructure); |
297 | |
298 | if (aStructState != myStructureStates.end()) |
299 | return aStructState->second != theStructure->ModificationState(); |
300 | |
301 | return Standard_True; |
302 | } |
303 | |
304 | // ======================================================================= |
305 | // function : CreateMaterial |
306 | // purpose : Creates ray-tracing material properties |
307 | // ======================================================================= |
265d4508 |
308 | void CreateMaterial (const OPENGL_SURF_PROP& theProp, OpenGl_RaytraceMaterial& theMaterial) |
e276548b |
309 | { |
64c759f8 |
310 | const float* aSrcAmb = theProp.isphysic ? theProp.ambcol.rgb : theProp.matcol.rgb; |
265d4508 |
311 | theMaterial.Ambient = BVH_Vec4f (aSrcAmb[0] * theProp.amb, |
312 | aSrcAmb[1] * theProp.amb, |
313 | aSrcAmb[2] * theProp.amb, |
314 | 1.0f); |
64c759f8 |
315 | |
316 | const float* aSrcDif = theProp.isphysic ? theProp.difcol.rgb : theProp.matcol.rgb; |
265d4508 |
317 | theMaterial.Diffuse = BVH_Vec4f (aSrcDif[0] * theProp.diff, |
318 | aSrcDif[1] * theProp.diff, |
319 | aSrcDif[2] * theProp.diff, |
320 | 1.0f); |
64c759f8 |
321 | |
322 | const float aDefSpecCol[4] = {1.0f, 1.0f, 1.0f, 1.0f}; |
323 | const float* aSrcSpe = theProp.isphysic ? theProp.speccol.rgb : aDefSpecCol; |
265d4508 |
324 | theMaterial.Specular = BVH_Vec4f (aSrcSpe[0] * theProp.spec, |
325 | aSrcSpe[1] * theProp.spec, |
326 | aSrcSpe[2] * theProp.spec, |
327 | theProp.shine); |
e276548b |
328 | |
64c759f8 |
329 | const float* aSrcEms = theProp.isphysic ? theProp.emscol.rgb : theProp.matcol.rgb; |
265d4508 |
330 | theMaterial.Emission = BVH_Vec4f (aSrcEms[0] * theProp.emsv, |
331 | aSrcEms[1] * theProp.emsv, |
332 | aSrcEms[2] * theProp.emsv, |
333 | 1.0f); |
e276548b |
334 | |
335 | // Note: Here we use sub-linear transparency function |
336 | // to produce realistic-looking transparency effect |
265d4508 |
337 | theMaterial.Transparency = BVH_Vec4f (powf (theProp.trans, 0.75f), |
338 | 1.f - theProp.trans, |
339 | 1.f, |
340 | 1.f); |
e276548b |
341 | |
342 | const float aMaxRefl = Max (theMaterial.Diffuse.x() + theMaterial.Specular.x(), |
343 | Max (theMaterial.Diffuse.y() + theMaterial.Specular.y(), |
344 | theMaterial.Diffuse.z() + theMaterial.Specular.z())); |
345 | |
346 | const float aReflectionScale = 0.75f / aMaxRefl; |
347 | |
265d4508 |
348 | theMaterial.Reflection = BVH_Vec4f (theProp.speccol.rgb[0] * theProp.spec, |
349 | theProp.speccol.rgb[1] * theProp.spec, |
350 | theProp.speccol.rgb[2] * theProp.spec, |
351 | 0.f) * aReflectionScale; |
e276548b |
352 | } |
353 | |
354 | // ======================================================================= |
355 | // function : AddRaytraceStructure |
356 | // purpose : Adds OpenGL structure to ray-traced scene geometry |
357 | // ======================================================================= |
265d4508 |
358 | Standard_Boolean OpenGl_Workspace::AddRaytraceStructure (const OpenGl_Structure* theStructure, |
359 | const Standard_ShortReal* theTransform, std::set<const OpenGl_Structure*>& theElements) |
e276548b |
360 | { |
e276548b |
361 | theElements.insert (theStructure); |
362 | |
363 | if (!theStructure->IsVisible()) |
364 | { |
365 | myStructureStates[theStructure] = theStructure->ModificationState(); |
366 | return Standard_True; |
367 | } |
368 | |
369 | // Get structure material |
68333c8f |
370 | Standard_Integer aStructMatID = -1; |
e276548b |
371 | |
372 | if (theStructure->AspectFace() != NULL) |
373 | { |
265d4508 |
374 | aStructMatID = static_cast<Standard_Integer> (myRaytraceGeometry.Materials.size()); |
e276548b |
375 | |
376 | OpenGl_RaytraceMaterial aStructMaterial; |
377 | CreateMaterial (theStructure->AspectFace()->IntFront(), aStructMaterial); |
378 | |
265d4508 |
379 | myRaytraceGeometry.Materials.push_back (aStructMaterial); |
e276548b |
380 | } |
381 | |
b64d84be |
382 | for (OpenGl_Structure::GroupIterator aGroupIter (theStructure->DrawGroups()); aGroupIter.More(); aGroupIter.Next()) |
e276548b |
383 | { |
384 | // Get group material |
68333c8f |
385 | Standard_Integer aGroupMatID = -1; |
b64d84be |
386 | if (aGroupIter.Value()->AspectFace() != NULL) |
e276548b |
387 | { |
265d4508 |
388 | aGroupMatID = static_cast<Standard_Integer> (myRaytraceGeometry.Materials.size()); |
e276548b |
389 | |
390 | OpenGl_RaytraceMaterial aGroupMaterial; |
b64d84be |
391 | CreateMaterial (aGroupIter.Value()->AspectFace()->IntFront(), aGroupMaterial); |
e276548b |
392 | |
265d4508 |
393 | myRaytraceGeometry.Materials.push_back (aGroupMaterial); |
e276548b |
394 | } |
395 | |
68333c8f |
396 | Standard_Integer aMatID = aGroupMatID < 0 ? aStructMatID : aGroupMatID; |
e276548b |
397 | |
265d4508 |
398 | if (aMatID < 0) |
e276548b |
399 | { |
265d4508 |
400 | aMatID = static_cast<Standard_Integer> (myRaytraceGeometry.Materials.size()); |
e276548b |
401 | |
265d4508 |
402 | myRaytraceGeometry.Materials.push_back (OpenGl_RaytraceMaterial()); |
e276548b |
403 | } |
404 | |
265d4508 |
405 | // Add OpenGL elements from group (extract primitives arrays and aspects) |
b64d84be |
406 | for (const OpenGl_ElementNode* aNode = aGroupIter.Value()->FirstNode(); aNode != NULL; aNode = aNode->next) |
e276548b |
407 | { |
5322131b |
408 | OpenGl_AspectFace* anAspect = dynamic_cast<OpenGl_AspectFace*> (aNode->elem); |
409 | if (anAspect != NULL) |
e276548b |
410 | { |
5322131b |
411 | aMatID = static_cast<Standard_Integer> (myRaytraceGeometry.Materials.size()); |
e276548b |
412 | |
5322131b |
413 | OpenGl_RaytraceMaterial aMaterial; |
414 | CreateMaterial (anAspect->IntFront(), aMaterial); |
e276548b |
415 | |
5322131b |
416 | myRaytraceGeometry.Materials.push_back (aMaterial); |
e276548b |
417 | } |
5322131b |
418 | else |
e276548b |
419 | { |
420 | OpenGl_PrimitiveArray* aPrimArray = dynamic_cast<OpenGl_PrimitiveArray*> (aNode->elem); |
e276548b |
421 | if (aPrimArray != NULL) |
422 | { |
265d4508 |
423 | NCollection_Handle<BVH_Object<Standard_ShortReal, 4> > aSet = |
424 | AddRaytracePrimitiveArray (aPrimArray->PArray(), aMatID, theTransform); |
425 | |
426 | if (!aSet.IsNull()) |
427 | myRaytraceGeometry.Objects().Append (aSet); |
e276548b |
428 | } |
429 | } |
430 | } |
e276548b |
431 | } |
432 | |
265d4508 |
433 | Standard_ShortReal* aTransform (NULL); |
e276548b |
434 | |
435 | // Process all connected OpenGL structures |
265d4508 |
436 | for (OpenGl_ListOfStructure::Iterator anIts (theStructure->ConnectedStructures()); anIts.More(); anIts.Next()) |
e276548b |
437 | { |
438 | if (anIts.Value()->Transformation()->mat != NULL) |
439 | { |
265d4508 |
440 | Standard_ShortReal* aTransform = new Standard_ShortReal[16]; |
e276548b |
441 | |
68333c8f |
442 | for (Standard_Integer i = 0; i < 4; ++i) |
443 | for (Standard_Integer j = 0; j < 4; ++j) |
e276548b |
444 | { |
445 | aTransform[j * 4 + i] = |
446 | anIts.Value()->Transformation()->mat[i][j]; |
447 | } |
448 | } |
449 | |
450 | if (anIts.Value()->IsRaytracable()) |
451 | AddRaytraceStructure (anIts.Value(), aTransform != NULL ? aTransform : theTransform, theElements); |
e276548b |
452 | } |
453 | |
454 | delete[] aTransform; |
455 | |
456 | myStructureStates[theStructure] = theStructure->ModificationState(); |
457 | |
458 | return Standard_True; |
459 | } |
460 | |
461 | // ======================================================================= |
462 | // function : AddRaytracePrimitiveArray |
463 | // purpose : Adds OpenGL primitive array to ray-traced scene geometry |
464 | // ======================================================================= |
265d4508 |
465 | OpenGl_TriangleSet* OpenGl_Workspace::AddRaytracePrimitiveArray ( |
466 | const CALL_DEF_PARRAY* theArray, Standard_Integer theMatID, const Standard_ShortReal* theTransform) |
e276548b |
467 | { |
468 | if (theArray->type != TelPolygonsArrayType && |
469 | theArray->type != TelTrianglesArrayType && |
470 | theArray->type != TelQuadranglesArrayType && |
471 | theArray->type != TelTriangleFansArrayType && |
472 | theArray->type != TelTriangleStripsArrayType && |
473 | theArray->type != TelQuadrangleStripsArrayType) |
474 | { |
265d4508 |
475 | return NULL; |
e276548b |
476 | } |
477 | |
478 | if (theArray->vertices == NULL) |
265d4508 |
479 | return NULL; |
e276548b |
480 | |
481 | #ifdef RAY_TRACE_PRINT_INFO |
482 | switch (theArray->type) |
483 | { |
484 | case TelPolygonsArrayType: |
265d4508 |
485 | std::cout << "\tAdding TelPolygonsArrayType" << std::endl; break; |
e276548b |
486 | case TelTrianglesArrayType: |
265d4508 |
487 | std::cout << "\tAdding TelTrianglesArrayType" << std::endl; break; |
e276548b |
488 | case TelQuadranglesArrayType: |
265d4508 |
489 | std::cout << "\tAdding TelQuadranglesArrayType" << std::endl; break; |
e276548b |
490 | case TelTriangleFansArrayType: |
265d4508 |
491 | std::cout << "\tAdding TelTriangleFansArrayType" << std::endl; break; |
e276548b |
492 | case TelTriangleStripsArrayType: |
265d4508 |
493 | std::cout << "\tAdding TelTriangleStripsArrayType" << std::endl; break; |
e276548b |
494 | case TelQuadrangleStripsArrayType: |
265d4508 |
495 | std::cout << "\tAdding TelQuadrangleStripsArrayType" << std::endl; break; |
e276548b |
496 | } |
497 | #endif |
498 | |
265d4508 |
499 | OpenGl_TriangleSet* aSet = new OpenGl_TriangleSet; |
e276548b |
500 | |
e276548b |
501 | { |
265d4508 |
502 | aSet->Vertices.reserve (theArray->num_vertexs); |
e276548b |
503 | |
265d4508 |
504 | for (Standard_Integer aVert = 0; aVert < theArray->num_vertexs; ++aVert) |
505 | { |
506 | BVH_Vec4f aVertex (theArray->vertices[aVert].xyz[0], |
507 | theArray->vertices[aVert].xyz[1], |
508 | theArray->vertices[aVert].xyz[2], |
509 | 1.f); |
510 | if (theTransform) |
511 | aVertex = MatVecMult (theTransform, aVertex); |
e276548b |
512 | |
265d4508 |
513 | aSet->Vertices.push_back (aVertex); |
514 | } |
e276548b |
515 | |
265d4508 |
516 | aSet->Normals.reserve (theArray->num_vertexs); |
e276548b |
517 | |
265d4508 |
518 | for (Standard_Integer aNorm = 0; aNorm < theArray->num_vertexs; ++aNorm) |
519 | { |
520 | BVH_Vec4f aNormal; |
e276548b |
521 | |
265d4508 |
522 | // Note: In case of absence of normals, the |
523 | // renderer uses generated geometric normals |
e276548b |
524 | |
265d4508 |
525 | if (theArray->vnormals != NULL) |
526 | { |
527 | aNormal = BVH_Vec4f (theArray->vnormals[aNorm].xyz[0], |
528 | theArray->vnormals[aNorm].xyz[1], |
529 | theArray->vnormals[aNorm].xyz[2], |
530 | 0.f); |
e276548b |
531 | |
265d4508 |
532 | if (theTransform) |
533 | aNormal = MatVecMult (theTransform, aNormal); |
534 | } |
e276548b |
535 | |
265d4508 |
536 | aSet->Normals.push_back (aNormal); |
e276548b |
537 | } |
538 | |
265d4508 |
539 | if (theArray->num_bounds > 0) |
540 | { |
541 | #ifdef RAY_TRACE_PRINT_INFO |
542 | std::cout << "\tNumber of bounds = " << theArray->num_bounds << std::endl; |
543 | #endif |
e276548b |
544 | |
265d4508 |
545 | Standard_Integer aBoundStart = 0; |
e276548b |
546 | |
265d4508 |
547 | for (Standard_Integer aBound = 0; aBound < theArray->num_bounds; ++aBound) |
548 | { |
549 | const Standard_Integer aVertNum = theArray->bounds[aBound]; |
e276548b |
550 | |
265d4508 |
551 | #ifdef RAY_TRACE_PRINT_INFO |
552 | std::cout << "\tAdding indices from bound " << aBound << ": " << |
553 | aBoundStart << " .. " << aVertNum << std::endl; |
554 | #endif |
555 | |
556 | if (!AddRaytraceVertexIndices (aSet, theArray, aBoundStart, aVertNum, theMatID)) |
557 | { |
558 | delete aSet; |
559 | return NULL; |
560 | } |
561 | |
562 | aBoundStart += aVertNum; |
563 | } |
564 | } |
565 | else |
e276548b |
566 | { |
265d4508 |
567 | const Standard_Integer aVertNum = theArray->num_edges > 0 ? theArray->num_edges : theArray->num_vertexs; |
e276548b |
568 | |
265d4508 |
569 | #ifdef RAY_TRACE_PRINT_INFO |
570 | std::cout << "\tAdding indices from array: " << aVertNum << std::endl; |
571 | #endif |
e276548b |
572 | |
265d4508 |
573 | if (!AddRaytraceVertexIndices (aSet, theArray, 0, aVertNum, theMatID)) |
e276548b |
574 | { |
265d4508 |
575 | delete aSet; |
576 | return NULL; |
e276548b |
577 | } |
e276548b |
578 | } |
579 | } |
e276548b |
580 | |
265d4508 |
581 | if (aSet->Size() != 0) |
582 | aSet->MarkDirty(); |
e276548b |
583 | |
265d4508 |
584 | return aSet; |
e276548b |
585 | } |
586 | |
587 | // ======================================================================= |
588 | // function : AddRaytraceVertexIndices |
589 | // purpose : Adds vertex indices to ray-traced scene geometry |
590 | // ======================================================================= |
265d4508 |
591 | Standard_Boolean OpenGl_Workspace::AddRaytraceVertexIndices (OpenGl_TriangleSet* theSet, |
592 | const CALL_DEF_PARRAY* theArray, Standard_Integer theOffset, Standard_Integer theCount, Standard_Integer theMatID) |
e276548b |
593 | { |
e276548b |
594 | switch (theArray->type) |
595 | { |
265d4508 |
596 | case TelTrianglesArrayType: |
597 | return AddRaytraceTriangleArray (theSet, theArray, theOffset, theCount, theMatID); |
598 | |
599 | case TelQuadranglesArrayType: |
600 | return AddRaytraceQuadrangleArray (theSet, theArray, theOffset, theCount, theMatID); |
601 | |
602 | case TelTriangleFansArrayType: |
603 | return AddRaytraceTriangleFanArray (theSet, theArray, theOffset, theCount, theMatID); |
604 | |
5322131b |
605 | case TelTriangleStripsArrayType: |
265d4508 |
606 | return AddRaytraceTriangleStripArray (theSet, theArray, theOffset, theCount, theMatID); |
607 | |
608 | case TelQuadrangleStripsArrayType: |
609 | return AddRaytraceQuadrangleStripArray (theSet, theArray, theOffset, theCount, theMatID); |
610 | |
611 | default: |
612 | return AddRaytracePolygonArray (theSet, theArray, theOffset, theCount, theMatID); |
e276548b |
613 | } |
e276548b |
614 | } |
615 | |
616 | // ======================================================================= |
617 | // function : AddRaytraceTriangleArray |
618 | // purpose : Adds OpenGL triangle array to ray-traced scene geometry |
619 | // ======================================================================= |
265d4508 |
620 | Standard_Boolean OpenGl_Workspace::AddRaytraceTriangleArray (OpenGl_TriangleSet* theSet, |
621 | const CALL_DEF_PARRAY* theArray, Standard_Integer theOffset, Standard_Integer theCount, Standard_Integer theMatID) |
e276548b |
622 | { |
265d4508 |
623 | if (theCount < 3) |
e276548b |
624 | return Standard_True; |
625 | |
265d4508 |
626 | theSet->Elements.reserve (theSet->Elements.size() + theCount / 3); |
627 | |
e276548b |
628 | if (theArray->num_edges > 0) |
629 | { |
265d4508 |
630 | for (Standard_Integer aVert = theOffset; aVert < theOffset + theCount - 2; aVert += 3) |
e276548b |
631 | { |
265d4508 |
632 | theSet->Elements.push_back (BVH_Vec4i (theArray->edges[aVert + 0], |
633 | theArray->edges[aVert + 1], |
634 | theArray->edges[aVert + 2], |
635 | theMatID)); |
e276548b |
636 | } |
637 | } |
638 | else |
639 | { |
265d4508 |
640 | for (Standard_Integer aVert = theOffset; aVert < theOffset + theCount - 2; aVert += 3) |
e276548b |
641 | { |
265d4508 |
642 | theSet->Elements.push_back (BVH_Vec4i (aVert + 0, |
643 | aVert + 1, |
644 | aVert + 2, |
645 | theMatID)); |
e276548b |
646 | } |
647 | } |
648 | |
649 | return Standard_True; |
650 | } |
651 | |
652 | // ======================================================================= |
653 | // function : AddRaytraceTriangleFanArray |
654 | // purpose : Adds OpenGL triangle fan array to ray-traced scene geometry |
655 | // ======================================================================= |
265d4508 |
656 | Standard_Boolean OpenGl_Workspace::AddRaytraceTriangleFanArray (OpenGl_TriangleSet* theSet, |
657 | const CALL_DEF_PARRAY* theArray, Standard_Integer theOffset, Standard_Integer theCount, Standard_Integer theMatID) |
e276548b |
658 | { |
265d4508 |
659 | if (theCount < 3) |
e276548b |
660 | return Standard_True; |
661 | |
265d4508 |
662 | theSet->Elements.reserve (theSet->Elements.size() + theCount - 2); |
663 | |
e276548b |
664 | if (theArray->num_edges > 0) |
665 | { |
265d4508 |
666 | for (Standard_Integer aVert = theOffset; aVert < theOffset + theCount - 2; ++aVert) |
e276548b |
667 | { |
265d4508 |
668 | theSet->Elements.push_back (BVH_Vec4i (theArray->edges[theOffset], |
669 | theArray->edges[aVert + 1], |
670 | theArray->edges[aVert + 2], |
671 | theMatID)); |
e276548b |
672 | } |
673 | } |
674 | else |
675 | { |
265d4508 |
676 | for (Standard_Integer aVert = theOffset; aVert < theOffset + theCount - 2; ++aVert) |
e276548b |
677 | { |
265d4508 |
678 | theSet->Elements.push_back (BVH_Vec4i (theOffset, |
679 | aVert + 1, |
680 | aVert + 2, |
681 | theMatID)); |
e276548b |
682 | } |
683 | } |
684 | |
685 | return Standard_True; |
686 | } |
687 | |
688 | // ======================================================================= |
689 | // function : AddRaytraceTriangleStripArray |
690 | // purpose : Adds OpenGL triangle strip array to ray-traced scene geometry |
691 | // ======================================================================= |
265d4508 |
692 | Standard_Boolean OpenGl_Workspace::AddRaytraceTriangleStripArray (OpenGl_TriangleSet* theSet, |
693 | const CALL_DEF_PARRAY* theArray, Standard_Integer theOffset, Standard_Integer theCount, Standard_Integer theMatID) |
e276548b |
694 | { |
265d4508 |
695 | if (theCount < 3) |
e276548b |
696 | return Standard_True; |
697 | |
265d4508 |
698 | theSet->Elements.reserve (theSet->Elements.size() + theCount - 2); |
699 | |
e276548b |
700 | if (theArray->num_edges > 0) |
701 | { |
265d4508 |
702 | for (Standard_Integer aVert = theOffset, aCW = 0; aVert < theOffset + theCount - 2; ++aVert, aCW = (aCW + 1) % 2) |
e276548b |
703 | { |
265d4508 |
704 | theSet->Elements.push_back (BVH_Vec4i (theArray->edges[aVert + aCW ? 1 : 0], |
705 | theArray->edges[aVert + aCW ? 0 : 1], |
706 | theArray->edges[aVert + 2], |
707 | theMatID)); |
e276548b |
708 | } |
709 | } |
710 | else |
711 | { |
265d4508 |
712 | for (Standard_Integer aVert = theOffset, aCW = 0; aVert < theOffset + theCount - 2; ++aVert, aCW = (aCW + 1) % 2) |
e276548b |
713 | { |
265d4508 |
714 | theSet->Elements.push_back (BVH_Vec4i (aVert + aCW ? 1 : 0, |
715 | aVert + aCW ? 0 : 1, |
716 | aVert + 2, |
717 | theMatID)); |
e276548b |
718 | } |
719 | } |
720 | |
721 | return Standard_True; |
722 | } |
723 | |
724 | // ======================================================================= |
725 | // function : AddRaytraceQuadrangleArray |
726 | // purpose : Adds OpenGL quad array to ray-traced scene geometry |
727 | // ======================================================================= |
265d4508 |
728 | Standard_Boolean OpenGl_Workspace::AddRaytraceQuadrangleArray (OpenGl_TriangleSet* theSet, |
729 | const CALL_DEF_PARRAY* theArray, Standard_Integer theOffset, Standard_Integer theCount, Standard_Integer theMatID) |
e276548b |
730 | { |
265d4508 |
731 | if (theCount < 4) |
e276548b |
732 | return Standard_True; |
733 | |
265d4508 |
734 | theSet->Elements.reserve (theSet->Elements.size() + theCount / 2); |
735 | |
e276548b |
736 | if (theArray->num_edges > 0) |
737 | { |
265d4508 |
738 | for (Standard_Integer aVert = theOffset; aVert < theOffset + theCount - 3; aVert += 4) |
e276548b |
739 | { |
265d4508 |
740 | theSet->Elements.push_back (BVH_Vec4i (theArray->edges[aVert + 0], |
741 | theArray->edges[aVert + 1], |
742 | theArray->edges[aVert + 2], |
743 | theMatID)); |
744 | |
745 | theSet->Elements.push_back (BVH_Vec4i (theArray->edges[aVert + 0], |
746 | theArray->edges[aVert + 2], |
747 | theArray->edges[aVert + 3], |
748 | theMatID)); |
e276548b |
749 | } |
750 | } |
751 | else |
752 | { |
265d4508 |
753 | for (Standard_Integer aVert = theOffset; aVert < theOffset + theCount - 3; aVert += 4) |
e276548b |
754 | { |
265d4508 |
755 | theSet->Elements.push_back (BVH_Vec4i (aVert + 0, |
756 | aVert + 1, |
757 | aVert + 2, |
758 | theMatID)); |
759 | |
760 | theSet->Elements.push_back (BVH_Vec4i (aVert + 0, |
761 | aVert + 2, |
762 | aVert + 3, |
763 | theMatID)); |
e276548b |
764 | } |
765 | } |
766 | |
767 | return Standard_True; |
768 | } |
769 | |
770 | // ======================================================================= |
771 | // function : AddRaytraceQuadrangleStripArray |
772 | // purpose : Adds OpenGL quad strip array to ray-traced scene geometry |
773 | // ======================================================================= |
265d4508 |
774 | Standard_Boolean OpenGl_Workspace::AddRaytraceQuadrangleStripArray (OpenGl_TriangleSet* theSet, |
775 | const CALL_DEF_PARRAY* theArray, Standard_Integer theOffset, Standard_Integer theCount, Standard_Integer theMatID) |
e276548b |
776 | { |
265d4508 |
777 | if (theCount < 4) |
e276548b |
778 | return Standard_True; |
779 | |
265d4508 |
780 | theSet->Elements.reserve (theSet->Elements.size() + 2 * theCount - 6); |
781 | |
e276548b |
782 | if (theArray->num_edges > 0) |
783 | { |
265d4508 |
784 | for (Standard_Integer aVert = theOffset; aVert < theOffset + theCount - 3; aVert += 2) |
e276548b |
785 | { |
265d4508 |
786 | theSet->Elements.push_back (BVH_Vec4i (theArray->edges[aVert + 0], |
787 | theArray->edges[aVert + 1], |
788 | theArray->edges[aVert + 2], |
789 | theMatID)); |
790 | |
791 | theSet->Elements.push_back (BVH_Vec4i (theArray->edges[aVert + 1], |
792 | theArray->edges[aVert + 3], |
793 | theArray->edges[aVert + 2], |
794 | theMatID)); |
e276548b |
795 | } |
796 | } |
797 | else |
798 | { |
265d4508 |
799 | for (Standard_Integer aVert = theOffset; aVert < theOffset + theCount - 3; aVert += 2) |
e276548b |
800 | { |
265d4508 |
801 | theSet->Elements.push_back (BVH_Vec4i (aVert + 0, |
802 | aVert + 1, |
803 | aVert + 2, |
804 | theMatID)); |
805 | |
806 | theSet->Elements.push_back (BVH_Vec4i (aVert + 1, |
807 | aVert + 3, |
808 | aVert + 2, |
809 | theMatID)); |
e276548b |
810 | } |
811 | } |
812 | |
813 | return Standard_True; |
814 | } |
815 | |
816 | // ======================================================================= |
817 | // function : AddRaytracePolygonArray |
818 | // purpose : Adds OpenGL polygon array to ray-traced scene geometry |
819 | // ======================================================================= |
265d4508 |
820 | Standard_Boolean OpenGl_Workspace::AddRaytracePolygonArray (OpenGl_TriangleSet* theSet, |
821 | const CALL_DEF_PARRAY* theArray, Standard_Integer theOffset, Standard_Integer theCount, Standard_Integer theMatID) |
e276548b |
822 | { |
265d4508 |
823 | if (theCount < 3) |
e276548b |
824 | return Standard_True; |
825 | |
265d4508 |
826 | theSet->Elements.reserve (theSet->Elements.size() + theCount - 2); |
827 | |
828 | if (theArray->num_edges > 0) |
e276548b |
829 | { |
265d4508 |
830 | for (Standard_Integer aVert = theOffset; aVert < theOffset + theCount - 2; ++aVert) |
e276548b |
831 | { |
265d4508 |
832 | theSet->Elements.push_back (BVH_Vec4i (theArray->edges[theOffset], |
833 | theArray->edges[aVert + 1], |
834 | theArray->edges[aVert + 2], |
835 | theMatID)); |
e276548b |
836 | } |
837 | } |
838 | else |
839 | { |
265d4508 |
840 | for (Standard_Integer aVert = theOffset; aVert < theOffset + theCount - 2; ++aVert) |
e276548b |
841 | { |
265d4508 |
842 | theSet->Elements.push_back (BVH_Vec4i (theOffset, |
843 | aVert + 1, |
844 | aVert + 2, |
845 | theMatID)); |
e276548b |
846 | } |
847 | } |
848 | |
849 | return Standard_True; |
850 | } |
851 | |
852 | // ======================================================================= |
853 | // function : UpdateRaytraceLightSources |
854 | // purpose : Updates 3D scene light sources for ray-tracing |
855 | // ======================================================================= |
856 | Standard_Boolean OpenGl_Workspace::UpdateRaytraceLightSources (const GLdouble theInvModelView[16]) |
857 | { |
265d4508 |
858 | myRaytraceGeometry.Sources.clear(); |
e276548b |
859 | |
265d4508 |
860 | myRaytraceGeometry.GlobalAmbient = BVH_Vec4f (0.0f, 0.0f, 0.0f, 0.0f); |
861 | |
862 | for (OpenGl_ListOfLight::Iterator anItl (myView->LightList()); anItl.More(); anItl.Next()) |
e276548b |
863 | { |
12381341 |
864 | const OpenGl_Light& aLight = anItl.Value(); |
265d4508 |
865 | |
12381341 |
866 | if (aLight.Type == Visual3d_TOLS_AMBIENT) |
e276548b |
867 | { |
265d4508 |
868 | myRaytraceGeometry.GlobalAmbient += BVH_Vec4f (aLight.Color.r(), |
869 | aLight.Color.g(), |
870 | aLight.Color.b(), |
871 | 0.0f); |
e276548b |
872 | continue; |
873 | } |
874 | |
265d4508 |
875 | BVH_Vec4f aDiffuse (aLight.Color.r(), |
876 | aLight.Color.g(), |
877 | aLight.Color.b(), |
878 | 1.0f); |
879 | |
880 | BVH_Vec4f aPosition (-aLight.Direction.x(), |
881 | -aLight.Direction.y(), |
882 | -aLight.Direction.z(), |
883 | 0.0f); |
884 | |
12381341 |
885 | if (aLight.Type != Visual3d_TOLS_DIRECTIONAL) |
e276548b |
886 | { |
265d4508 |
887 | aPosition = BVH_Vec4f (aLight.Position.x(), |
888 | aLight.Position.y(), |
889 | aLight.Position.z(), |
890 | 1.0f); |
e276548b |
891 | } |
265d4508 |
892 | |
12381341 |
893 | if (aLight.IsHeadlight) |
e276548b |
894 | aPosition = MatVecMult (theInvModelView, aPosition); |
e276548b |
895 | |
265d4508 |
896 | myRaytraceGeometry.Sources.push_back (OpenGl_RaytraceLight (aDiffuse, aPosition)); |
e276548b |
897 | } |
898 | |
899 | cl_int anError = CL_SUCCESS; |
900 | |
901 | if (myRaytraceLightSourceBuffer != NULL) |
902 | clReleaseMemObject (myRaytraceLightSourceBuffer); |
903 | |
265d4508 |
904 | Standard_Integer aLightBufferSize = myRaytraceGeometry.Sources.size() != 0 ? |
905 | static_cast<Standard_Integer> (myRaytraceGeometry.Sources.size()) : 1; |
e276548b |
906 | |
907 | myRaytraceLightSourceBuffer = clCreateBuffer (myComputeContext, CL_MEM_READ_ONLY, |
265d4508 |
908 | aLightBufferSize * sizeof(OpenGl_RaytraceLight), NULL, &anError); |
e276548b |
909 | |
265d4508 |
910 | if (myRaytraceGeometry.Sources.size() != 0) |
e276548b |
911 | { |
265d4508 |
912 | const void* aDataPtr = myRaytraceGeometry.Sources.front().Packed(); |
913 | |
914 | anError |= clEnqueueWriteBuffer (myComputeQueue, myRaytraceLightSourceBuffer, CL_TRUE, 0, |
915 | aLightBufferSize * sizeof(OpenGl_RaytraceLight), aDataPtr, 0, NULL, NULL); |
e276548b |
916 | } |
917 | |
918 | #ifdef RAY_TRACE_PRINT_INFO |
919 | if (anError != CL_SUCCESS) |
920 | { |
265d4508 |
921 | std::cout << "Error! Failed to set light sources"; |
e276548b |
922 | |
923 | return Standard_False; |
924 | } |
925 | #endif |
926 | |
927 | return Standard_True; |
928 | } |
929 | |
930 | // ======================================================================= |
931 | // function : CheckOpenCL |
932 | // purpose : Checks OpenCL dynamic library availability |
933 | // ======================================================================= |
934 | Standard_Boolean CheckOpenCL() |
935 | { |
936 | #if defined ( _WIN32 ) |
937 | |
938 | __try |
939 | { |
940 | cl_uint aNbPlatforms; |
941 | clGetPlatformIDs (0, NULL, &aNbPlatforms); |
942 | } |
943 | __except (EXCEPTION_EXECUTE_HANDLER) |
944 | { |
945 | return Standard_False; |
946 | } |
947 | |
948 | #endif |
949 | |
950 | return Standard_True; |
951 | } |
952 | |
953 | // ======================================================================= |
954 | // function : InitOpenCL |
955 | // purpose : Initializes OpenCL objects |
956 | // ======================================================================= |
957 | Standard_Boolean OpenGl_Workspace::InitOpenCL() |
958 | { |
959 | if (myComputeInitStatus != OpenGl_CLIS_NONE) |
960 | { |
961 | return myComputeInitStatus == OpenGl_CLIS_INIT; |
962 | } |
963 | |
964 | if (!CheckOpenCL()) |
965 | { |
966 | myComputeInitStatus = OpenGl_CLIS_FAIL; // fail to load OpenCL library |
967 | myGlContext->PushMessage (GL_DEBUG_SOURCE_APPLICATION_ARB, |
968 | GL_DEBUG_TYPE_ERROR_ARB, |
969 | 0, |
970 | GL_DEBUG_SEVERITY_HIGH_ARB, |
971 | "Failed to load OpenCL dynamic library!"); |
972 | return Standard_False; |
973 | } |
974 | |
975 | // Obtain the list of platforms available |
976 | cl_uint aNbPlatforms = 0; |
977 | cl_int anError = clGetPlatformIDs (0, NULL, &aNbPlatforms); |
978 | cl_platform_id* aPlatforms = (cl_platform_id* )alloca (aNbPlatforms * sizeof(cl_platform_id)); |
979 | anError |= clGetPlatformIDs (aNbPlatforms, aPlatforms, NULL); |
980 | if (anError != CL_SUCCESS |
981 | || aNbPlatforms == 0) |
982 | { |
983 | myComputeInitStatus = OpenGl_CLIS_FAIL; |
984 | myGlContext->PushMessage (GL_DEBUG_SOURCE_APPLICATION_ARB, |
985 | GL_DEBUG_TYPE_ERROR_ARB, |
986 | 0, |
987 | GL_DEBUG_SEVERITY_HIGH_ARB, |
988 | "No any OpenCL platform installed!"); |
989 | return Standard_False; |
990 | } |
991 | |
992 | // Note: We try to find NVIDIA or AMD platforms with GPU devices! |
993 | cl_platform_id aPrefPlatform = NULL; |
994 | for (cl_uint aPlatIter = 0; aPlatIter < aNbPlatforms; ++aPlatIter) |
995 | { |
996 | char aName[256]; |
997 | anError = clGetPlatformInfo (aPlatforms[aPlatIter], CL_PLATFORM_NAME, |
998 | sizeof(aName), aName, NULL); |
999 | if (anError != CL_SUCCESS) |
1000 | { |
1001 | continue; |
1002 | } |
1003 | |
1004 | if (strncmp (aName, "NVIDIA", strlen ("NVIDIA")) == 0) |
1005 | { |
1006 | aPrefPlatform = aPlatforms[aPlatIter]; |
1007 | |
1008 | // Use optimizations for NVIDIA GPUs |
1009 | myIsAmdComputePlatform = Standard_False; |
1010 | } |
1011 | else if (strncmp (aName, "AMD", strlen ("AMD")) == 0) |
1012 | { |
1013 | aPrefPlatform = (aPrefPlatform == NULL) |
1014 | ? aPlatforms[aPlatIter] |
1015 | : aPrefPlatform; |
1016 | |
1017 | // Use optimizations for ATI/AMD platform |
1018 | myIsAmdComputePlatform = Standard_True; |
1019 | } |
1020 | } |
1021 | |
1022 | if (aPrefPlatform == NULL) |
1023 | { |
1024 | aPrefPlatform = aPlatforms[0]; |
1025 | } |
1026 | |
1027 | // Obtain the list of devices available in the selected platform |
1028 | cl_uint aNbDevices = 0; |
1029 | anError = clGetDeviceIDs (aPrefPlatform, CL_DEVICE_TYPE_GPU, |
1030 | 0, NULL, &aNbDevices); |
1031 | |
1032 | cl_device_id* aDevices = (cl_device_id* )alloca (aNbDevices * sizeof(cl_device_id)); |
1033 | anError |= clGetDeviceIDs (aPrefPlatform, CL_DEVICE_TYPE_GPU, |
1034 | aNbDevices, aDevices, NULL); |
1035 | if (anError != CL_SUCCESS) |
1036 | { |
1037 | myComputeInitStatus = OpenGl_CLIS_FAIL; |
1038 | myGlContext->PushMessage (GL_DEBUG_SOURCE_APPLICATION_ARB, |
1039 | GL_DEBUG_TYPE_ERROR_ARB, |
1040 | 0, |
1041 | GL_DEBUG_SEVERITY_HIGH_ARB, |
1042 | "Failed to get OpenCL GPU device!"); |
1043 | return Standard_False; |
1044 | } |
1045 | |
1046 | // Note: Simply get first available GPU |
1047 | cl_device_id aDevice = aDevices[0]; |
1048 | |
1049 | // detect old contexts |
1050 | char aVerClStr[256]; |
1051 | clGetDeviceInfo (aDevice, CL_DEVICE_VERSION, |
1052 | sizeof(aVerClStr), aVerClStr, NULL); |
1053 | aVerClStr[strlen ("OpenCL 1.0")] = '\0'; |
1054 | const bool isVer10 = strncmp (aVerClStr, "OpenCL 1.0", strlen ("OpenCL 1.0")) == 0; |
1055 | |
1056 | // Create OpenCL context |
1057 | cl_context_properties aCtxProp[] = |
1058 | { |
1059 | #if defined(__APPLE__) && !defined(MACOSX_USE_GLX) |
1060 | CL_CONTEXT_PROPERTY_USE_CGL_SHAREGROUP_APPLE, |
1061 | (cl_context_properties )CGLGetShareGroup (CGLGetCurrentContext()), |
1062 | #elif defined(_WIN32) |
1063 | CL_CONTEXT_PLATFORM, (cl_context_properties )aPrefPlatform, |
1064 | CL_GL_CONTEXT_KHR, (cl_context_properties )wglGetCurrentContext(), |
1065 | CL_WGL_HDC_KHR, (cl_context_properties )wglGetCurrentDC(), |
1066 | #else |
1067 | CL_GL_CONTEXT_KHR, (cl_context_properties )glXGetCurrentContext(), |
1068 | CL_GLX_DISPLAY_KHR, (cl_context_properties )glXGetCurrentDisplay(), |
1069 | CL_CONTEXT_PLATFORM, (cl_context_properties )aPrefPlatform, |
1070 | #endif |
1071 | 0 |
1072 | }; |
1073 | |
1074 | myComputeContext = clCreateContext (aCtxProp, |
1075 | #if defined(__APPLE__) && !defined(MACOSX_USE_GLX) |
1076 | 0, NULL, // device will be taken from GL context |
1077 | #else |
1078 | 1, &aDevice, |
1079 | #endif |
1080 | NULL, NULL, &anError); |
1081 | if (anError != CL_SUCCESS) |
1082 | { |
1083 | myComputeInitStatus = OpenGl_CLIS_FAIL; |
1084 | myGlContext->PushMessage (GL_DEBUG_SOURCE_APPLICATION_ARB, |
1085 | GL_DEBUG_TYPE_ERROR_ARB, |
1086 | 0, |
1087 | GL_DEBUG_SEVERITY_HIGH_ARB, |
1088 | "Failed to initialize OpenCL context!"); |
1089 | return Standard_False; |
1090 | } |
1091 | |
1092 | // Create OpenCL program |
1093 | const char* aSources[] = |
1094 | { |
1095 | isVer10 ? "#define M_PI_F ( float )( 3.14159265359f )\n" : "", |
1096 | THE_RAY_TRACE_OPENCL_SOURCE |
1097 | }; |
1098 | myRaytraceProgram = clCreateProgramWithSource (myComputeContext, 2, |
1099 | aSources, NULL, &anError); |
1100 | if (anError != CL_SUCCESS) |
1101 | { |
1102 | myComputeInitStatus = OpenGl_CLIS_FAIL; |
1103 | myGlContext->PushMessage (GL_DEBUG_SOURCE_APPLICATION_ARB, |
1104 | GL_DEBUG_TYPE_ERROR_ARB, |
1105 | 0, |
1106 | GL_DEBUG_SEVERITY_HIGH_ARB, |
1107 | "Failed to create OpenCL ray-tracing program!"); |
1108 | return Standard_False; |
1109 | } |
1110 | |
1111 | anError = clBuildProgram (myRaytraceProgram, 0, |
1112 | NULL, NULL, NULL, NULL); |
1113 | { |
1114 | // Fetch build log |
1115 | size_t aLogLen = 0; |
1116 | cl_int aResult = clGetProgramBuildInfo (myRaytraceProgram, aDevice, |
1117 | CL_PROGRAM_BUILD_LOG, 0, NULL, &aLogLen); |
1118 | |
1119 | char* aBuildLog = (char* )alloca (aLogLen); |
1120 | aResult |= clGetProgramBuildInfo (myRaytraceProgram, aDevice, |
1121 | CL_PROGRAM_BUILD_LOG, aLogLen, aBuildLog, NULL); |
1122 | if (aResult == CL_SUCCESS) |
1123 | { |
1124 | if (anError != CL_SUCCESS) |
1125 | { |
1126 | myGlContext->PushMessage (GL_DEBUG_SOURCE_APPLICATION_ARB, |
1127 | GL_DEBUG_TYPE_ERROR_ARB, |
1128 | 0, |
1129 | GL_DEBUG_SEVERITY_HIGH_ARB, |
1130 | aBuildLog); |
1131 | } |
1132 | else |
1133 | { |
1134 | #ifdef RAY_TRACE_PRINT_INFO |
1135 | std::cout << aBuildLog << std::endl; |
1136 | #endif |
1137 | } |
1138 | } |
1139 | } |
1140 | |
1141 | if (anError != CL_SUCCESS) |
1142 | { |
1143 | return Standard_False; |
1144 | } |
1145 | |
1146 | // Create OpenCL ray tracing kernels |
265d4508 |
1147 | myRaytraceRenderKernel = clCreateKernel (myRaytraceProgram, "RaytraceRender", &anError); |
e276548b |
1148 | if (anError != CL_SUCCESS) |
1149 | { |
1150 | myComputeInitStatus = OpenGl_CLIS_FAIL; |
1151 | myGlContext->PushMessage (GL_DEBUG_SOURCE_APPLICATION_ARB, |
1152 | GL_DEBUG_TYPE_ERROR_ARB, |
1153 | 0, |
1154 | GL_DEBUG_SEVERITY_HIGH_ARB, |
1155 | "Failed to create OpenCL ray-tracing kernel!"); |
1156 | return Standard_False; |
1157 | } |
1158 | |
265d4508 |
1159 | myRaytraceSmoothKernel = clCreateKernel (myRaytraceProgram, "RaytraceSmooth", &anError); |
e276548b |
1160 | if (anError != CL_SUCCESS) |
1161 | { |
1162 | myComputeInitStatus = OpenGl_CLIS_FAIL; |
1163 | myGlContext->PushMessage (GL_DEBUG_SOURCE_APPLICATION_ARB, |
1164 | GL_DEBUG_TYPE_ERROR_ARB, |
1165 | 0, |
1166 | GL_DEBUG_SEVERITY_HIGH_ARB, |
1167 | "Failed to create OpenCL ray-tracing kernel!"); |
1168 | return Standard_False; |
1169 | } |
1170 | |
1171 | // Create OpenCL command queue |
1172 | // Note: For profiling set CL_QUEUE_PROFILING_ENABLE |
1173 | cl_command_queue_properties aProps = CL_QUEUE_PROFILING_ENABLE; |
1174 | |
265d4508 |
1175 | myComputeQueue = clCreateCommandQueue (myComputeContext, aDevice, aProps, &anError); |
e276548b |
1176 | if (anError != CL_SUCCESS) |
1177 | { |
1178 | myComputeInitStatus = OpenGl_CLIS_FAIL; |
1179 | myGlContext->PushMessage (GL_DEBUG_SOURCE_APPLICATION_ARB, |
1180 | GL_DEBUG_TYPE_ERROR_ARB, |
1181 | 0, |
1182 | GL_DEBUG_SEVERITY_HIGH_ARB, |
1183 | "Failed to create OpenCL command queue!"); |
1184 | |
1185 | return Standard_False; |
1186 | } |
1187 | |
1188 | myComputeInitStatus = OpenGl_CLIS_INIT; // initialized in normal way |
1189 | return Standard_True; |
1190 | } |
1191 | |
1192 | // ======================================================================= |
1193 | // function : GetOpenClDeviceInfo |
1194 | // purpose : Returns information about device used for computations |
1195 | // ======================================================================= |
1196 | Standard_Boolean OpenGl_Workspace::GetOpenClDeviceInfo (NCollection_DataMap<TCollection_AsciiString, |
1197 | TCollection_AsciiString>& theInfo) const |
1198 | { |
1199 | theInfo.Clear(); |
1200 | if (myComputeContext == NULL) |
1201 | { |
1202 | return Standard_False; |
1203 | } |
1204 | |
1205 | size_t aDevicesSize = 0; |
1206 | cl_int anError = clGetContextInfo (myComputeContext, CL_CONTEXT_DEVICES, 0, NULL, &aDevicesSize); |
1207 | cl_device_id* aDevices = (cl_device_id* )alloca (aDevicesSize); |
1208 | anError |= clGetContextInfo (myComputeContext, CL_CONTEXT_DEVICES, aDevicesSize, aDevices, NULL); |
1209 | if (anError != CL_SUCCESS) |
1210 | { |
1211 | return Standard_False; |
1212 | } |
1213 | |
1214 | char aDeviceName[256]; |
1215 | anError |= clGetDeviceInfo (aDevices[0], CL_DEVICE_NAME, sizeof(aDeviceName), aDeviceName, NULL); |
1216 | theInfo.Bind ("Name", aDeviceName); |
1217 | |
1218 | char aDeviceVendor[256]; |
1219 | anError |= clGetDeviceInfo (aDevices[0], CL_DEVICE_VENDOR, sizeof(aDeviceVendor), aDeviceVendor, NULL); |
1220 | theInfo.Bind ("Vendor", aDeviceVendor); |
1221 | |
1222 | cl_device_type aDeviceType; |
1223 | anError |= clGetDeviceInfo (aDevices[0], CL_DEVICE_TYPE, sizeof(aDeviceType), &aDeviceType, NULL); |
1224 | theInfo.Bind ("Type", aDeviceType == CL_DEVICE_TYPE_GPU ? "GPU" : "CPU"); |
1225 | return Standard_True; |
1226 | } |
1227 | |
1228 | // ======================================================================= |
1229 | // function : ReleaseOpenCL |
1230 | // purpose : Releases resources of OpenCL objects |
1231 | // ======================================================================= |
1232 | void OpenGl_Workspace::ReleaseOpenCL() |
1233 | { |
1234 | clReleaseKernel (myRaytraceRenderKernel); |
1235 | clReleaseKernel (myRaytraceSmoothKernel); |
1236 | |
1237 | clReleaseProgram (myRaytraceProgram); |
265d4508 |
1238 | clReleaseCommandQueue (myComputeQueue); |
5322131b |
1239 | |
e276548b |
1240 | clReleaseMemObject (myRaytraceOutputImage); |
1241 | clReleaseMemObject (myRaytraceEnvironment); |
68333c8f |
1242 | clReleaseMemObject (myRaytraceOutputImageAA); |
e276548b |
1243 | |
e276548b |
1244 | clReleaseMemObject (myRaytraceMaterialBuffer); |
1245 | clReleaseMemObject (myRaytraceLightSourceBuffer); |
1246 | |
265d4508 |
1247 | clReleaseMemObject (mySceneNodeInfoBuffer); |
1248 | clReleaseMemObject (mySceneMinPointBuffer); |
1249 | clReleaseMemObject (mySceneMaxPointBuffer); |
1250 | |
1251 | clReleaseMemObject (myObjectNodeInfoBuffer); |
1252 | clReleaseMemObject (myObjectMinPointBuffer); |
1253 | clReleaseMemObject (myObjectMaxPointBuffer); |
1254 | |
1255 | clReleaseMemObject (myGeometryVertexBuffer); |
1256 | clReleaseMemObject (myGeometryNormalBuffer); |
1257 | clReleaseMemObject (myGeometryTriangBuffer); |
e276548b |
1258 | |
1259 | clReleaseContext (myComputeContext); |
1260 | |
68333c8f |
1261 | if (!myGlContext.IsNull()) |
1262 | { |
1263 | if (!myRaytraceOutputTexture.IsNull()) |
1264 | myGlContext->DelayedRelease (myRaytraceOutputTexture); |
1265 | myRaytraceOutputTexture.Nullify(); |
1266 | |
1267 | if (!myRaytraceOutputTextureAA.IsNull()) |
1268 | myGlContext->DelayedRelease (myRaytraceOutputTextureAA); |
1269 | myRaytraceOutputTextureAA.Nullify(); |
1270 | } |
e276548b |
1271 | } |
1272 | |
1273 | // ======================================================================= |
1274 | // function : ResizeRaytraceOutputBuffer |
1275 | // purpose : Resizes OpenCL output image |
1276 | // ======================================================================= |
1277 | Standard_Boolean OpenGl_Workspace::ResizeRaytraceOutputBuffer (const cl_int theSizeX, |
1278 | const cl_int theSizeY) |
1279 | { |
1280 | if (myComputeContext == NULL) |
1281 | { |
1282 | return Standard_False; |
1283 | } |
1284 | |
68333c8f |
1285 | if (!myRaytraceOutputTexture.IsNull()) |
e276548b |
1286 | { |
68333c8f |
1287 | Standard_Boolean toResize = myRaytraceOutputTexture->SizeX() != theSizeX || |
1288 | myRaytraceOutputTexture->SizeY() != theSizeY; |
e276548b |
1289 | |
68333c8f |
1290 | if (!toResize) |
1291 | return Standard_True; |
e276548b |
1292 | |
68333c8f |
1293 | if (!myGlContext.IsNull()) |
e276548b |
1294 | { |
68333c8f |
1295 | if (!myRaytraceOutputTexture.IsNull()) |
1296 | myGlContext->DelayedRelease (myRaytraceOutputTexture); |
1297 | if (!myRaytraceOutputTextureAA.IsNull()) |
1298 | myGlContext->DelayedRelease (myRaytraceOutputTextureAA); |
e276548b |
1299 | } |
1300 | } |
5322131b |
1301 | |
68333c8f |
1302 | myRaytraceOutputTexture = new OpenGl_Texture(); |
5322131b |
1303 | |
68333c8f |
1304 | myRaytraceOutputTexture->Create (myGlContext); |
1305 | myRaytraceOutputTexture->InitRectangle (myGlContext, |
1306 | theSizeX, theSizeY, OpenGl_TextureFormat::Create<GLfloat, 4>()); |
5322131b |
1307 | |
68333c8f |
1308 | myRaytraceOutputTextureAA = new OpenGl_Texture(); |
5322131b |
1309 | |
68333c8f |
1310 | myRaytraceOutputTextureAA->Create (myGlContext); |
1311 | myRaytraceOutputTextureAA->InitRectangle (myGlContext, |
1312 | theSizeX, theSizeY, OpenGl_TextureFormat::Create<GLfloat, 4>()); |
e276548b |
1313 | |
68333c8f |
1314 | if (myRaytraceOutputImage != NULL) |
1315 | clReleaseMemObject (myRaytraceOutputImage); |
e276548b |
1316 | |
68333c8f |
1317 | if (myRaytraceOutputImageAA != NULL) |
1318 | clReleaseMemObject (myRaytraceOutputImageAA); |
e276548b |
1319 | |
1320 | cl_int anError = CL_SUCCESS; |
1321 | |
68333c8f |
1322 | myRaytraceOutputImage = clCreateFromGLTexture2D (myComputeContext, |
1323 | CL_MEM_READ_WRITE, GL_TEXTURE_RECTANGLE, 0, myRaytraceOutputTexture->TextureId(), &anError); |
e276548b |
1324 | |
e276548b |
1325 | if (anError != CL_SUCCESS) |
1326 | { |
1327 | #ifdef RAY_TRACE_PRINT_INFO |
1328 | std::cout << "Error! Failed to create output image!" << std::endl; |
1329 | #endif |
1330 | return Standard_False; |
1331 | } |
1332 | |
68333c8f |
1333 | myRaytraceOutputImageAA = clCreateFromGLTexture2D (myComputeContext, |
1334 | CL_MEM_READ_WRITE, GL_TEXTURE_RECTANGLE, 0, myRaytraceOutputTextureAA->TextureId(), &anError); |
1335 | |
e276548b |
1336 | if (anError != CL_SUCCESS) |
1337 | { |
1338 | #ifdef RAY_TRACE_PRINT_INFO |
1339 | std::cout << "Error! Failed to create anti-aliased output image!" << std::endl; |
1340 | #endif |
1341 | return Standard_False; |
1342 | } |
1343 | |
1344 | return Standard_True; |
1345 | } |
1346 | |
1347 | // ======================================================================= |
1348 | // function : WriteRaytraceSceneToDevice |
265d4508 |
1349 | // purpose : Writes scene geometry to OpenCL device |
e276548b |
1350 | // ======================================================================= |
1351 | Standard_Boolean OpenGl_Workspace::WriteRaytraceSceneToDevice() |
1352 | { |
1353 | if (myComputeContext == NULL) |
1354 | return Standard_False; |
1355 | |
265d4508 |
1356 | cl_int anErrorRes = CL_SUCCESS; |
1357 | |
1358 | if (mySceneNodeInfoBuffer != NULL) |
1359 | anErrorRes |= clReleaseMemObject (mySceneNodeInfoBuffer); |
1360 | |
1361 | if (mySceneMinPointBuffer != NULL) |
1362 | anErrorRes |= clReleaseMemObject (mySceneMinPointBuffer); |
1363 | |
1364 | if (mySceneMaxPointBuffer != NULL) |
1365 | anErrorRes |= clReleaseMemObject (mySceneMaxPointBuffer); |
e276548b |
1366 | |
265d4508 |
1367 | if (myObjectNodeInfoBuffer != NULL) |
1368 | anErrorRes |= clReleaseMemObject (myObjectNodeInfoBuffer); |
e276548b |
1369 | |
265d4508 |
1370 | if (myObjectMinPointBuffer != NULL) |
1371 | anErrorRes |= clReleaseMemObject (myObjectMinPointBuffer); |
e276548b |
1372 | |
265d4508 |
1373 | if (myObjectMaxPointBuffer != NULL) |
1374 | anErrorRes |= clReleaseMemObject (myObjectMaxPointBuffer); |
e276548b |
1375 | |
265d4508 |
1376 | if (myGeometryVertexBuffer != NULL) |
1377 | anErrorRes |= clReleaseMemObject (myGeometryVertexBuffer); |
e276548b |
1378 | |
265d4508 |
1379 | if (myGeometryNormalBuffer != NULL) |
1380 | anErrorRes |= clReleaseMemObject (myGeometryNormalBuffer); |
e276548b |
1381 | |
265d4508 |
1382 | if (myGeometryTriangBuffer != NULL) |
1383 | anErrorRes |= clReleaseMemObject (myGeometryTriangBuffer); |
e276548b |
1384 | |
1385 | if (myRaytraceMaterialBuffer != NULL) |
265d4508 |
1386 | anErrorRes |= clReleaseMemObject (myRaytraceMaterialBuffer); |
e276548b |
1387 | |
265d4508 |
1388 | if (anErrorRes != CL_SUCCESS) |
e276548b |
1389 | { |
1390 | #ifdef RAY_TRACE_PRINT_INFO |
265d4508 |
1391 | std::cout << "Error! Failed to release OpenCL buffers" << std::endl; |
e276548b |
1392 | #endif |
1393 | return Standard_False; |
1394 | } |
1395 | |
265d4508 |
1396 | ///////////////////////////////////////////////////////////////////////////// |
1397 | // Create material buffer |
1398 | |
1399 | const size_t aMaterialBufferSize = |
1400 | myRaytraceGeometry.Materials.size() != 0 ? myRaytraceGeometry.Materials.size() : 1; |
1401 | |
1402 | myRaytraceMaterialBuffer = clCreateBuffer (myComputeContext, |
1403 | CL_MEM_READ_ONLY, aMaterialBufferSize * sizeof(OpenGl_RaytraceMaterial), NULL, &anErrorRes); |
1404 | |
1405 | if (anErrorRes != CL_SUCCESS) |
e276548b |
1406 | { |
1407 | #ifdef RAY_TRACE_PRINT_INFO |
265d4508 |
1408 | std::cout << "Error! Failed to create OpenCL material buffer" << std::endl; |
e276548b |
1409 | #endif |
1410 | return Standard_False; |
1411 | } |
1412 | |
265d4508 |
1413 | ///////////////////////////////////////////////////////////////////////////// |
1414 | // Create BVHs buffers |
1415 | |
1416 | cl_int anErrorTmp = CL_SUCCESS; |
1417 | |
1418 | const NCollection_Handle<BVH_Tree<Standard_ShortReal, 4> >& aBVH = myRaytraceGeometry.BVH(); |
5322131b |
1419 | |
265d4508 |
1420 | const size_t aSceneMinPointBufferSize = |
1421 | aBVH->MinPointBuffer().size() != 0 ? aBVH->MinPointBuffer().size() : 1; |
1422 | |
1423 | mySceneMinPointBuffer = clCreateBuffer (myComputeContext, |
1424 | CL_MEM_READ_ONLY, aSceneMinPointBufferSize * sizeof(cl_float4), NULL, &anErrorTmp); |
1425 | anErrorRes |= anErrorTmp; |
1426 | |
1427 | const size_t aSceneMaxPointBufferSize = |
1428 | aBVH->MaxPointBuffer().size() != 0 ? aBVH->MaxPointBuffer().size() : 1; |
1429 | |
1430 | mySceneMaxPointBuffer = clCreateBuffer (myComputeContext, |
1431 | CL_MEM_READ_ONLY, aSceneMaxPointBufferSize * sizeof(cl_float4), NULL, &anErrorTmp); |
1432 | anErrorRes |= anErrorTmp; |
1433 | |
1434 | const size_t aSceneNodeInfoBufferSize = |
1435 | aBVH->NodeInfoBuffer().size() != 0 ? aBVH->NodeInfoBuffer().size() : 1; |
1436 | |
1437 | mySceneNodeInfoBuffer = clCreateBuffer (myComputeContext, |
1438 | CL_MEM_READ_ONLY, aSceneNodeInfoBufferSize * sizeof(cl_int4), NULL, &anErrorTmp); |
1439 | anErrorRes |= anErrorTmp; |
1440 | |
1441 | if (anErrorRes != CL_SUCCESS) |
e276548b |
1442 | { |
1443 | #ifdef RAY_TRACE_PRINT_INFO |
265d4508 |
1444 | std::cout << "Error! Failed to create OpenCL buffers for high-level scene BVH" << std::endl; |
e276548b |
1445 | #endif |
1446 | return Standard_False; |
1447 | } |
1448 | |
265d4508 |
1449 | Standard_Integer aTotalVerticesNb = 0; |
1450 | Standard_Integer aTotalElementsNb = 0; |
1451 | Standard_Integer aTotalBVHNodesNb = 0; |
1452 | |
1453 | for (Standard_Integer anElemIndex = 0; anElemIndex < myRaytraceGeometry.Size(); ++anElemIndex) |
1454 | { |
1455 | OpenGl_TriangleSet* aTriangleSet = dynamic_cast<OpenGl_TriangleSet*> ( |
1456 | myRaytraceGeometry.Objects().ChangeValue (anElemIndex).operator->()); |
1457 | |
1458 | Standard_ASSERT_RETURN (aTriangleSet != NULL, |
1459 | "Error! Failed to get triangulation of OpenGL element", Standard_False); |
1460 | |
d5f74e42 |
1461 | aTotalVerticesNb += (int)aTriangleSet->Vertices.size(); |
1462 | aTotalElementsNb += (int)aTriangleSet->Elements.size(); |
265d4508 |
1463 | |
1464 | Standard_ASSERT_RETURN (!aTriangleSet->BVH().IsNull(), |
1465 | "Error! Failed to get bottom-level BVH of OpenGL element", Standard_False); |
1466 | |
d5f74e42 |
1467 | aTotalBVHNodesNb += (int)aTriangleSet->BVH()->NodeInfoBuffer().size(); |
265d4508 |
1468 | } |
5322131b |
1469 | |
265d4508 |
1470 | aTotalBVHNodesNb = aTotalBVHNodesNb > 0 ? aTotalBVHNodesNb : 1; |
1471 | |
1472 | myObjectNodeInfoBuffer = clCreateBuffer (myComputeContext, |
1473 | CL_MEM_READ_ONLY, aTotalBVHNodesNb * sizeof(cl_int4), NULL, &anErrorTmp); |
1474 | anErrorRes |= anErrorTmp; |
1475 | |
1476 | myObjectMinPointBuffer = clCreateBuffer (myComputeContext, |
1477 | CL_MEM_READ_ONLY, aTotalBVHNodesNb * sizeof(cl_float4), NULL, &anErrorTmp); |
1478 | anErrorRes |= anErrorTmp; |
1479 | |
1480 | myObjectMaxPointBuffer = clCreateBuffer (myComputeContext, |
1481 | CL_MEM_READ_ONLY, aTotalBVHNodesNb * sizeof(cl_float4), NULL, &anErrorTmp); |
1482 | anErrorRes |= anErrorTmp; |
1483 | |
1484 | if (anErrorRes != CL_SUCCESS) |
e276548b |
1485 | { |
1486 | #ifdef RAY_TRACE_PRINT_INFO |
265d4508 |
1487 | std::cout << "Error! Failed to create OpenCL buffers for bottom-level scene BVHs" << std::endl; |
e276548b |
1488 | #endif |
1489 | return Standard_False; |
1490 | } |
1491 | |
265d4508 |
1492 | ///////////////////////////////////////////////////////////////////////////// |
1493 | // Create geometry buffers |
1494 | |
1495 | aTotalVerticesNb = aTotalVerticesNb > 0 ? aTotalVerticesNb : 1; |
1496 | |
1497 | myGeometryVertexBuffer = clCreateBuffer (myComputeContext, |
1498 | CL_MEM_READ_ONLY, aTotalVerticesNb * sizeof(cl_float4), NULL, &anErrorTmp); |
1499 | anErrorRes |= anErrorTmp; |
1500 | |
1501 | myGeometryNormalBuffer = clCreateBuffer (myComputeContext, |
1502 | CL_MEM_READ_ONLY, aTotalVerticesNb * sizeof(cl_float4), NULL, &anErrorTmp); |
1503 | anErrorRes |= anErrorTmp; |
1504 | |
1505 | aTotalElementsNb = aTotalElementsNb > 0 ? aTotalElementsNb : 1; |
1506 | |
1507 | myGeometryTriangBuffer = clCreateBuffer (myComputeContext, |
1508 | CL_MEM_READ_ONLY, aTotalElementsNb * sizeof(cl_int4), NULL, &anErrorTmp); |
1509 | anErrorRes |= anErrorTmp; |
1510 | |
1511 | if (anErrorRes != CL_SUCCESS) |
1512 | { |
1513 | #ifdef RAY_TRACE_PRINT_INFO |
1514 | std::cout << "Error! Failed to create OpenCL geometry buffers" << std::endl; |
1515 | #endif |
1516 | return Standard_False; |
1517 | } |
1518 | |
1519 | ///////////////////////////////////////////////////////////////////////////// |
1520 | // Write BVH and geometry buffers |
1521 | |
1522 | if (aBVH->NodeInfoBuffer().size() != 0) |
1523 | { |
1524 | anErrorRes |= clEnqueueWriteBuffer (myComputeQueue, mySceneNodeInfoBuffer, CL_FALSE, 0, |
1525 | aSceneNodeInfoBufferSize * sizeof(cl_int4), &aBVH->NodeInfoBuffer().front(), 0, NULL, NULL); |
1526 | |
1527 | anErrorRes |= clEnqueueWriteBuffer (myComputeQueue, mySceneMinPointBuffer, CL_FALSE, 0, |
1528 | aSceneMinPointBufferSize * sizeof(cl_float4), &aBVH->MinPointBuffer().front(), 0, NULL, NULL); |
1529 | |
1530 | anErrorRes |= clEnqueueWriteBuffer (myComputeQueue, mySceneMaxPointBuffer, CL_FALSE, 0, |
1531 | aSceneMaxPointBufferSize * sizeof(cl_float4), &aBVH->MaxPointBuffer().front(), 0, NULL, NULL); |
1532 | |
1533 | anErrorRes |= clFinish (myComputeQueue); |
1534 | |
1535 | if (anErrorRes != CL_SUCCESS) |
e276548b |
1536 | { |
265d4508 |
1537 | #ifdef RAY_TRACE_PRINT_INFO |
1538 | std::cout << "Error! Failed to write OpenCL buffers for high-level scene BVH" << std::endl; |
1539 | #endif |
e276548b |
1540 | return Standard_False; |
1541 | } |
e276548b |
1542 | |
265d4508 |
1543 | for (Standard_Integer aNodeIdx = 0; aNodeIdx < aBVH->Length(); ++aNodeIdx) |
e276548b |
1544 | { |
265d4508 |
1545 | if (!aBVH->IsOuter (aNodeIdx)) |
1546 | continue; |
1547 | |
1548 | OpenGl_TriangleSet* aTriangleSet = myRaytraceGeometry.TriangleSet (aNodeIdx); |
1549 | |
1550 | Standard_ASSERT_RETURN (aTriangleSet != NULL, |
1551 | "Error! Failed to get triangulation of OpenGL element", Standard_False); |
1552 | |
1553 | const size_t aBVHBuffserSize = |
1554 | aTriangleSet->BVH()->NodeInfoBuffer().size() != 0 ? aTriangleSet->BVH()->NodeInfoBuffer().size() : 1; |
1555 | |
1556 | const Standard_Integer aBVHOffset = myRaytraceGeometry.AccelerationOffset (aNodeIdx); |
1557 | |
1558 | Standard_ASSERT_RETURN (aBVHOffset != OpenGl_RaytraceGeometry::INVALID_OFFSET, |
1559 | "Error! Failed to get offset for bottom-level BVH", Standard_False); |
1560 | |
1561 | anErrorRes |= clEnqueueWriteBuffer (myComputeQueue, myObjectNodeInfoBuffer, CL_FALSE, aBVHOffset * sizeof(cl_int4), |
1562 | aBVHBuffserSize * sizeof(cl_int4), &aTriangleSet->BVH()->NodeInfoBuffer().front(), 0, NULL, NULL); |
1563 | |
1564 | anErrorRes |= clEnqueueWriteBuffer (myComputeQueue, myObjectMinPointBuffer, CL_FALSE, aBVHOffset * sizeof(cl_float4), |
1565 | aBVHBuffserSize * sizeof(cl_float4), &aTriangleSet->BVH()->MinPointBuffer().front(), 0, NULL, NULL); |
1566 | |
1567 | anErrorRes |= clEnqueueWriteBuffer (myComputeQueue, myObjectMaxPointBuffer, CL_FALSE, aBVHOffset * sizeof(cl_float4), |
1568 | aBVHBuffserSize * sizeof(cl_float4), &aTriangleSet->BVH()->MaxPointBuffer().front(), 0, NULL, NULL); |
1569 | |
1570 | anErrorRes |= clFinish (myComputeQueue); |
1571 | |
1572 | if (anErrorRes != CL_SUCCESS) |
1573 | { |
1574 | #ifdef RAY_TRACE_PRINT_INFO |
1575 | std::cout << "Error! Failed to write OpenCL buffers for bottom-level scene BVHs" << std::endl; |
1576 | #endif |
1577 | return Standard_False; |
1578 | } |
1579 | |
1580 | const Standard_Integer aVerticesOffset = myRaytraceGeometry.VerticesOffset (aNodeIdx); |
1581 | |
1582 | Standard_ASSERT_RETURN (aVerticesOffset != OpenGl_RaytraceGeometry::INVALID_OFFSET, |
1583 | "Error! Failed to get offset for triangulation vertices of OpenGL element", Standard_False); |
1584 | |
1585 | anErrorRes |= clEnqueueWriteBuffer (myComputeQueue, myGeometryVertexBuffer, CL_FALSE, aVerticesOffset * sizeof(cl_float4), |
1586 | aTriangleSet->Vertices.size() * sizeof(cl_float4), &aTriangleSet->Vertices.front(), 0, NULL, NULL); |
1587 | |
1588 | anErrorRes |= clEnqueueWriteBuffer (myComputeQueue, myGeometryNormalBuffer, CL_FALSE, aVerticesOffset * sizeof(cl_float4), |
1589 | aTriangleSet->Normals.size() * sizeof(cl_float4), &aTriangleSet->Normals.front(), 0, NULL, NULL); |
1590 | |
1591 | const Standard_Integer anElementsOffset = myRaytraceGeometry.ElementsOffset (aNodeIdx); |
1592 | |
1593 | Standard_ASSERT_RETURN (anElementsOffset != OpenGl_RaytraceGeometry::INVALID_OFFSET, |
1594 | "Error! Failed to get offset for triangulation elements of OpenGL element", Standard_False); |
1595 | |
1596 | anErrorRes |= clEnqueueWriteBuffer (myComputeQueue, myGeometryTriangBuffer, CL_FALSE, anElementsOffset * sizeof(cl_int4), |
1597 | aTriangleSet->Elements.size() * sizeof(cl_int4), &aTriangleSet->Elements.front(), 0, NULL, NULL); |
1598 | |
1599 | anErrorRes |= clFinish (myComputeQueue); |
1600 | |
1601 | if (anErrorRes != CL_SUCCESS) |
1602 | { |
1603 | #ifdef RAY_TRACE_PRINT_INFO |
1604 | std::cout << "Error! Failed to write OpenCL triangulation buffers for OpenGL element" << std::endl; |
1605 | #endif |
1606 | return Standard_False; |
1607 | } |
e276548b |
1608 | } |
1609 | } |
1610 | |
265d4508 |
1611 | ///////////////////////////////////////////////////////////////////////////// |
1612 | // Write material buffer |
1613 | |
1614 | if (myRaytraceGeometry.Materials.size() != 0) |
e276548b |
1615 | { |
265d4508 |
1616 | const void* aDataPtr = myRaytraceGeometry.Materials.front().Packed(); |
e276548b |
1617 | |
265d4508 |
1618 | anErrorRes |= clEnqueueWriteBuffer (myComputeQueue, myRaytraceMaterialBuffer, |
1619 | CL_FALSE, 0, aMaterialBufferSize * sizeof(OpenGl_RaytraceMaterial), aDataPtr, 0, NULL, NULL); |
1620 | |
1621 | if (anErrorRes != CL_SUCCESS) |
e276548b |
1622 | { |
1623 | #ifdef RAY_TRACE_PRINT_INFO |
265d4508 |
1624 | std::cout << "Error! Failed to write OpenCL material buffer" << std::endl; |
e276548b |
1625 | #endif |
1626 | return Standard_False; |
1627 | } |
1628 | } |
1629 | |
265d4508 |
1630 | anErrorRes |= clFinish (myComputeQueue); |
1631 | |
1632 | if (anErrorRes == CL_SUCCESS) |
1633 | { |
1634 | myIsRaytraceDataValid = myRaytraceGeometry.Objects().Size() != 0; |
1635 | } |
e276548b |
1636 | #ifdef RAY_TRACE_PRINT_INFO |
265d4508 |
1637 | else |
1638 | { |
1639 | std::cout << "Error! Failed to set scene data buffers" << std::endl; |
1640 | } |
e276548b |
1641 | #endif |
1642 | |
e276548b |
1643 | #ifdef RAY_TRACE_PRINT_INFO |
5322131b |
1644 | |
265d4508 |
1645 | Standard_ShortReal aMemUsed = 0.f; |
5322131b |
1646 | |
265d4508 |
1647 | for (Standard_Integer anElemIdx = 0; anElemIdx < myRaytraceGeometry.Size(); ++anElemIdx) |
1648 | { |
1649 | OpenGl_TriangleSet* aTriangleSet = dynamic_cast<OpenGl_TriangleSet*> ( |
1650 | myRaytraceGeometry.Objects().ChangeValue (anElemIdx).operator->()); |
e276548b |
1651 | |
265d4508 |
1652 | aMemUsed += static_cast<Standard_ShortReal> ( |
1653 | aTriangleSet->Vertices.size() * sizeof (BVH_Vec4f)); |
1654 | aMemUsed += static_cast<Standard_ShortReal> ( |
1655 | aTriangleSet->Normals.size() * sizeof (BVH_Vec4f)); |
1656 | aMemUsed += static_cast<Standard_ShortReal> ( |
1657 | aTriangleSet->Elements.size() * sizeof (BVH_Vec4i)); |
e276548b |
1658 | |
265d4508 |
1659 | aMemUsed += static_cast<Standard_ShortReal> ( |
1660 | aTriangleSet->BVH()->NodeInfoBuffer().size() * sizeof (BVH_Vec4i)); |
1661 | aMemUsed += static_cast<Standard_ShortReal> ( |
1662 | aTriangleSet->BVH()->MinPointBuffer().size() * sizeof (BVH_Vec4f)); |
1663 | aMemUsed += static_cast<Standard_ShortReal> ( |
1664 | aTriangleSet->BVH()->MaxPointBuffer().size() * sizeof (BVH_Vec4f)); |
1665 | } |
5322131b |
1666 | |
265d4508 |
1667 | aMemUsed += static_cast<Standard_ShortReal> ( |
1668 | myRaytraceGeometry.BVH()->NodeInfoBuffer().size() * sizeof (BVH_Vec4i)); |
1669 | aMemUsed += static_cast<Standard_ShortReal> ( |
1670 | myRaytraceGeometry.BVH()->MinPointBuffer().size() * sizeof (BVH_Vec4f)); |
1671 | aMemUsed += static_cast<Standard_ShortReal> ( |
1672 | myRaytraceGeometry.BVH()->MaxPointBuffer().size() * sizeof (BVH_Vec4f)); |
e276548b |
1673 | |
265d4508 |
1674 | std::cout << "GPU Memory Used (MB): ~" << aMemUsed / 1048576 << std::endl; |
e276548b |
1675 | |
1676 | #endif |
1677 | |
265d4508 |
1678 | return (CL_SUCCESS == anErrorRes); |
e276548b |
1679 | } |
1680 | |
265d4508 |
1681 | // Use it to estimate the optimal size of OpenCL work group |
1682 | // #define OPENCL_GROUP_SIZE_TEST |
e276548b |
1683 | |
1684 | // ======================================================================= |
1685 | // function : RunRaytraceOpenCLKernels |
1686 | // purpose : Runs OpenCL ray-tracing kernels |
1687 | // ======================================================================= |
265d4508 |
1688 | Standard_Boolean OpenGl_Workspace::RunRaytraceOpenCLKernels (const Graphic3d_CView& theCView, |
1689 | const Standard_ShortReal theOrigins[16], |
1690 | const Standard_ShortReal theDirects[16], |
1691 | const Standard_Integer theSizeX, |
1692 | const Standard_Integer theSizeY) |
e276548b |
1693 | { |
265d4508 |
1694 | if (myRaytraceRenderKernel == NULL || myComputeQueue == NULL) |
e276548b |
1695 | return Standard_False; |
1696 | |
265d4508 |
1697 | //////////////////////////////////////////////////////////////////////// |
e276548b |
1698 | // Set kernel arguments |
1699 | |
1700 | cl_uint anIndex = 0; |
1701 | cl_int anError = 0; |
1702 | |
d5f74e42 |
1703 | cl_int aLightSourceBufferSize = (cl_int)myRaytraceGeometry.Sources.size(); |
265d4508 |
1704 | |
1705 | anError |= clSetKernelArg ( |
1706 | myRaytraceRenderKernel, anIndex++, sizeof(cl_int), &theSizeX); |
1707 | anError |= clSetKernelArg ( |
1708 | myRaytraceRenderKernel, anIndex++, sizeof(cl_int), &theSizeY); |
1709 | anError |= clSetKernelArg ( |
1710 | myRaytraceRenderKernel, anIndex++, sizeof(cl_float16), theOrigins); |
1711 | anError |= clSetKernelArg ( |
1712 | myRaytraceRenderKernel, anIndex++, sizeof(cl_float16), theDirects); |
1713 | anError |= clSetKernelArg ( |
1714 | myRaytraceRenderKernel, anIndex++, sizeof(cl_mem), &myRaytraceEnvironment); |
1715 | anError |= clSetKernelArg ( |
1716 | myRaytraceRenderKernel, anIndex++, sizeof(cl_mem), &myRaytraceOutputImage); |
1717 | anError |= clSetKernelArg ( |
1718 | myRaytraceRenderKernel, anIndex++, sizeof(cl_mem), &mySceneNodeInfoBuffer); |
1719 | anError |= clSetKernelArg ( |
1720 | myRaytraceRenderKernel, anIndex++, sizeof(cl_mem), &mySceneMinPointBuffer); |
1721 | anError |= clSetKernelArg ( |
1722 | myRaytraceRenderKernel, anIndex++, sizeof(cl_mem), &mySceneMaxPointBuffer); |
1723 | anError |= clSetKernelArg ( |
1724 | myRaytraceRenderKernel, anIndex++, sizeof(cl_mem), &myObjectNodeInfoBuffer); |
1725 | anError |= clSetKernelArg ( |
1726 | myRaytraceRenderKernel, anIndex++, sizeof(cl_mem), &myObjectMinPointBuffer); |
1727 | anError |= clSetKernelArg ( |
1728 | myRaytraceRenderKernel, anIndex++, sizeof(cl_mem), &myObjectMaxPointBuffer); |
1729 | anError |= clSetKernelArg ( |
1730 | myRaytraceRenderKernel, anIndex++, sizeof(cl_mem), &myGeometryTriangBuffer); |
1731 | anError |= clSetKernelArg ( |
1732 | myRaytraceRenderKernel, anIndex++, sizeof(cl_mem), &myGeometryVertexBuffer); |
1733 | anError |= clSetKernelArg ( |
1734 | myRaytraceRenderKernel, anIndex++, sizeof(cl_mem), &myGeometryNormalBuffer); |
1735 | anError |= clSetKernelArg ( |
1736 | myRaytraceRenderKernel, anIndex++, sizeof(cl_mem), &myRaytraceLightSourceBuffer); |
1737 | anError |= clSetKernelArg ( |
1738 | myRaytraceRenderKernel, anIndex++, sizeof(cl_mem), &myRaytraceMaterialBuffer); |
1739 | anError |= clSetKernelArg ( |
1740 | myRaytraceRenderKernel, anIndex++, sizeof(cl_float4), &myRaytraceGeometry.GlobalAmbient); |
1741 | anError |= clSetKernelArg ( |
1742 | myRaytraceRenderKernel, anIndex++, sizeof(cl_int), &aLightSourceBufferSize); |
1743 | anError |= clSetKernelArg ( |
1744 | myRaytraceRenderKernel, anIndex++, sizeof(cl_int), &theCView.IsShadowsEnabled); |
1745 | anError |= clSetKernelArg ( |
1746 | myRaytraceRenderKernel, anIndex++, sizeof(cl_int), &theCView.IsReflectionsEnabled); |
1747 | anError |= clSetKernelArg ( |
1748 | myRaytraceRenderKernel, anIndex++, sizeof(cl_float), &myRaytraceSceneEpsilon); |
1749 | anError |= clSetKernelArg ( |
1750 | myRaytraceRenderKernel, anIndex++, sizeof(cl_float), &myRaytraceSceneRadius); |
1751 | |
e276548b |
1752 | if (anError != CL_SUCCESS) |
1753 | { |
265d4508 |
1754 | const TCollection_ExtendedString aMessage = "Error! Failed to set arguments of ray-tracing kernel!"; |
5322131b |
1755 | |
e276548b |
1756 | myGlContext->PushMessage (GL_DEBUG_SOURCE_APPLICATION_ARB, |
265d4508 |
1757 | GL_DEBUG_TYPE_ERROR_ARB, 0, GL_DEBUG_SEVERITY_HIGH_ARB, aMessage); |
1758 | |
e276548b |
1759 | return Standard_False; |
1760 | } |
1761 | |
265d4508 |
1762 | // Second-pass 'smoothing' kernel runs only if anti-aliasing is enabled |
e276548b |
1763 | if (theCView.IsAntialiasingEnabled) |
1764 | { |
1765 | anIndex = 0; |
265d4508 |
1766 | |
1767 | anError |= clSetKernelArg ( |
1768 | myRaytraceSmoothKernel, anIndex++, sizeof(cl_int), &theSizeX); |
1769 | anError |= clSetKernelArg ( |
1770 | myRaytraceSmoothKernel, anIndex++, sizeof(cl_int), &theSizeY); |
1771 | anError |= clSetKernelArg ( |
1772 | myRaytraceSmoothKernel, anIndex++, sizeof(cl_float16), theOrigins); |
1773 | anError |= clSetKernelArg ( |
1774 | myRaytraceSmoothKernel, anIndex++, sizeof(cl_float16), theDirects); |
1775 | anError |= clSetKernelArg ( |
1776 | myRaytraceSmoothKernel, anIndex++, sizeof(cl_mem), &myRaytraceOutputImage); |
1777 | anError |= clSetKernelArg ( |
1778 | myRaytraceSmoothKernel, anIndex++, sizeof(cl_mem), &myRaytraceEnvironment); |
1779 | anError |= clSetKernelArg ( |
1780 | myRaytraceSmoothKernel, anIndex++, sizeof(cl_mem), &myRaytraceOutputImageAA); |
1781 | anError |= clSetKernelArg ( |
1782 | myRaytraceSmoothKernel, anIndex++, sizeof(cl_mem), &mySceneNodeInfoBuffer); |
1783 | anError |= clSetKernelArg ( |
1784 | myRaytraceSmoothKernel, anIndex++, sizeof(cl_mem), &mySceneMinPointBuffer); |
1785 | anError |= clSetKernelArg ( |
1786 | myRaytraceSmoothKernel, anIndex++, sizeof(cl_mem), &mySceneMaxPointBuffer); |
1787 | anError |= clSetKernelArg ( |
1788 | myRaytraceSmoothKernel, anIndex++, sizeof(cl_mem), &myObjectNodeInfoBuffer); |
1789 | anError |= clSetKernelArg ( |
1790 | myRaytraceSmoothKernel, anIndex++, sizeof(cl_mem), &myObjectMinPointBuffer); |
1791 | anError |= clSetKernelArg ( |
1792 | myRaytraceSmoothKernel, anIndex++, sizeof(cl_mem), &myObjectMaxPointBuffer); |
1793 | anError |= clSetKernelArg ( |
1794 | myRaytraceSmoothKernel, anIndex++, sizeof(cl_mem), &myGeometryTriangBuffer); |
1795 | anError |= clSetKernelArg ( |
1796 | myRaytraceSmoothKernel, anIndex++, sizeof(cl_mem), &myGeometryVertexBuffer); |
1797 | anError |= clSetKernelArg ( |
1798 | myRaytraceSmoothKernel, anIndex++, sizeof(cl_mem), &myGeometryNormalBuffer); |
1799 | anError |= clSetKernelArg ( |
1800 | myRaytraceSmoothKernel, anIndex++, sizeof(cl_mem), &myRaytraceLightSourceBuffer); |
1801 | anError |= clSetKernelArg ( |
1802 | myRaytraceSmoothKernel, anIndex++, sizeof(cl_mem), &myRaytraceMaterialBuffer); |
1803 | anError |= clSetKernelArg ( |
1804 | myRaytraceSmoothKernel, anIndex++, sizeof(cl_float4), &myRaytraceGeometry.GlobalAmbient); |
1805 | anError |= clSetKernelArg ( |
1806 | myRaytraceSmoothKernel, anIndex++, sizeof(cl_int), &aLightSourceBufferSize); |
1807 | anError |= clSetKernelArg ( |
1808 | myRaytraceSmoothKernel, anIndex++, sizeof(cl_int), &theCView.IsShadowsEnabled); |
1809 | anError |= clSetKernelArg ( |
1810 | myRaytraceSmoothKernel, anIndex++, sizeof(cl_int), &theCView.IsReflectionsEnabled); |
1811 | anError |= clSetKernelArg ( |
1812 | myRaytraceSmoothKernel, anIndex++, sizeof(cl_float), &myRaytraceSceneEpsilon); |
1813 | anError |= clSetKernelArg ( |
1814 | myRaytraceSmoothKernel, anIndex++, sizeof(cl_float), &myRaytraceSceneRadius); |
1815 | |
e276548b |
1816 | if (anError != CL_SUCCESS) |
1817 | { |
265d4508 |
1818 | const TCollection_ExtendedString aMessage = "Error! Failed to set arguments of smoothing kernel!"; |
1819 | |
e276548b |
1820 | myGlContext->PushMessage (GL_DEBUG_SOURCE_APPLICATION_ARB, |
265d4508 |
1821 | GL_DEBUG_TYPE_ERROR_ARB, 0, GL_DEBUG_SEVERITY_HIGH_ARB, aMessage); |
1822 | |
e276548b |
1823 | return Standard_False; |
1824 | } |
1825 | } |
1826 | |
265d4508 |
1827 | //////////////////////////////////////////////////////////////////////// |
e276548b |
1828 | // Set work size |
265d4508 |
1829 | |
1830 | size_t aLocWorkSize[] = { myIsAmdComputePlatform ? 2 : 4, 32 }; |
e276548b |
1831 | |
1832 | #ifdef OPENCL_GROUP_SIZE_TEST |
265d4508 |
1833 | for (size_t aLocX = 2; aLocX <= 32; aLocX <<= 1) |
1834 | for (size_t aLocY = 2; aLocY <= 32; aLocY <<= 1) |
e276548b |
1835 | { |
265d4508 |
1836 | aLocWorkSize[0] = aLocX; |
1837 | aLocWorkSize[1] = aLocY; |
e276548b |
1838 | #endif |
1839 | |
1840 | size_t aWorkSizeX = theSizeX; |
265d4508 |
1841 | if (aWorkSizeX % aLocWorkSize[0] != 0) |
1842 | aWorkSizeX += aLocWorkSize[0] - aWorkSizeX % aLocWorkSize[0]; |
e276548b |
1843 | |
1844 | size_t aWokrSizeY = theSizeY; |
265d4508 |
1845 | if (aWokrSizeY % aLocWorkSize[1] != 0 ) |
1846 | aWokrSizeY += aLocWorkSize[1] - aWokrSizeY % aLocWorkSize[1]; |
1847 | |
1848 | size_t aTotWorkSize[] = { aWorkSizeX, aWokrSizeY }; |
e276548b |
1849 | |
265d4508 |
1850 | cl_event anEvent = NULL, anEventSmooth = NULL; |
1851 | |
1852 | anError = clEnqueueNDRangeKernel (myComputeQueue, |
1853 | myRaytraceRenderKernel, 2, NULL, aTotWorkSize, aLocWorkSize, 0, NULL, &anEvent); |
e276548b |
1854 | |
e276548b |
1855 | if (anError != CL_SUCCESS) |
1856 | { |
265d4508 |
1857 | const TCollection_ExtendedString aMessage = "Error! Failed to execute the ray-tracing kernel!"; |
1858 | |
e276548b |
1859 | myGlContext->PushMessage (GL_DEBUG_SOURCE_APPLICATION_ARB, |
265d4508 |
1860 | GL_DEBUG_TYPE_ERROR_ARB, 0, GL_DEBUG_SEVERITY_HIGH_ARB, aMessage); |
1861 | |
e276548b |
1862 | return Standard_False; |
1863 | } |
265d4508 |
1864 | |
e276548b |
1865 | clWaitForEvents (1, &anEvent); |
1866 | |
1867 | if (theCView.IsAntialiasingEnabled) |
1868 | { |
265d4508 |
1869 | size_t aLocWorkSizeSmooth[] = { myIsAmdComputePlatform ? 8 : 4, |
1870 | myIsAmdComputePlatform ? 8 : 32 }; |
e276548b |
1871 | |
1872 | #ifdef OPENCL_GROUP_SIZE_TEST |
265d4508 |
1873 | aLocWorkSizeSmooth[0] = aLocX; |
1874 | aLocWorkSizeSmooth[1] = aLocY; |
e276548b |
1875 | #endif |
1876 | |
1877 | aWorkSizeX = theSizeX; |
265d4508 |
1878 | if (aWorkSizeX % aLocWorkSizeSmooth[0] != 0) |
1879 | aWorkSizeX += aLocWorkSizeSmooth[0] - aWorkSizeX % aLocWorkSizeSmooth[0]; |
e276548b |
1880 | |
1881 | size_t aWokrSizeY = theSizeY; |
265d4508 |
1882 | if (aWokrSizeY % aLocWorkSizeSmooth[1] != 0 ) |
1883 | aWokrSizeY += aLocWorkSizeSmooth[1] - aWokrSizeY % aLocWorkSizeSmooth[1]; |
1884 | |
1885 | size_t aTotWorkSizeSmooth [] = { aWorkSizeX, aWokrSizeY }; |
1886 | |
1887 | anError = clEnqueueNDRangeKernel (myComputeQueue, myRaytraceSmoothKernel, |
1888 | 2, NULL, aTotWorkSizeSmooth, aLocWorkSizeSmooth, 0, NULL, &anEventSmooth); |
e276548b |
1889 | |
e276548b |
1890 | clWaitForEvents (1, &anEventSmooth); |
1891 | |
1892 | if (anError != CL_SUCCESS) |
1893 | { |
265d4508 |
1894 | const TCollection_ExtendedString aMessage = "Error! Failed to execute the smoothing kernel"; |
1895 | |
e276548b |
1896 | myGlContext->PushMessage (GL_DEBUG_SOURCE_APPLICATION_ARB, |
265d4508 |
1897 | GL_DEBUG_TYPE_ERROR_ARB, 0, GL_DEBUG_SEVERITY_HIGH_ARB, aMessage); |
1898 | |
e276548b |
1899 | return Standard_False; |
1900 | } |
1901 | } |
1902 | |
265d4508 |
1903 | #if defined (RAY_TRACE_PRINT_INFO) || defined (OPENCL_GROUP_SIZE_TEST) |
1904 | |
1905 | static cl_ulong ttt1 = 10000000000; |
1906 | static cl_ulong ttt2 = 10000000000; |
1907 | |
1908 | cl_ulong aBegTime = 0; |
1909 | cl_ulong aEndTime = 0; |
e276548b |
1910 | |
265d4508 |
1911 | clGetEventProfilingInfo (anEvent, |
1912 | CL_PROFILING_COMMAND_START, sizeof(aBegTime), &aBegTime, NULL); |
1913 | clGetEventProfilingInfo (anEvent, |
1914 | CL_PROFILING_COMMAND_END, sizeof(aEndTime), &aEndTime, NULL); |
e276548b |
1915 | |
265d4508 |
1916 | ttt1 = aEndTime - aBegTime < ttt1 ? aEndTime - aBegTime : ttt1; |
1917 | |
1918 | std::cout << "\tRender time (ms): " << ttt1 / 1e6f << std::endl; |
e276548b |
1919 | |
1920 | if (theCView.IsAntialiasingEnabled) |
1921 | { |
265d4508 |
1922 | clGetEventProfilingInfo (anEventSmooth, |
1923 | CL_PROFILING_COMMAND_START, sizeof(aBegTime), &aBegTime, NULL); |
1924 | clGetEventProfilingInfo (anEventSmooth, |
1925 | CL_PROFILING_COMMAND_END, sizeof(aEndTime), &aEndTime, NULL); |
1926 | |
1927 | ttt2 = aEndTime - aBegTime < ttt2 ? aEndTime - aBegTime : ttt2; |
1928 | |
1929 | std::cout << "\tSmooth time (ms): " << ttt2 / 1e6f << std::endl; |
e276548b |
1930 | } |
265d4508 |
1931 | |
e276548b |
1932 | #endif |
1933 | |
1934 | if (anEvent != NULL) |
1935 | clReleaseEvent (anEvent); |
1936 | |
1937 | if (anEventSmooth != NULL) |
1938 | clReleaseEvent (anEventSmooth); |
265d4508 |
1939 | |
1940 | #ifdef OPENCL_GROUP_SIZE_TEST |
e276548b |
1941 | } |
265d4508 |
1942 | #endif |
e276548b |
1943 | |
1944 | return Standard_True; |
1945 | } |
1946 | |
1947 | // ======================================================================= |
1948 | // function : ComputeInverseMatrix |
1949 | // purpose : Computes inversion of 4x4 floating-point matrix |
1950 | // ======================================================================= |
1951 | template <typename T> |
1952 | void ComputeInverseMatrix (const T m[16], T inv[16]) |
1953 | { |
1954 | inv[ 0] = m[ 5] * (m[10] * m[15] - m[11] * m[14]) - |
1955 | m[ 9] * (m[ 6] * m[15] - m[ 7] * m[14]) - |
1956 | m[13] * (m[ 7] * m[10] - m[ 6] * m[11]); |
1957 | |
1958 | inv[ 1] = m[ 1] * (m[11] * m[14] - m[10] * m[15]) - |
1959 | m[ 9] * (m[ 3] * m[14] - m[ 2] * m[15]) - |
1960 | m[13] * (m[ 2] * m[11] - m[ 3] * m[10]); |
1961 | |
1962 | inv[ 2] = m[ 1] * (m[ 6] * m[15] - m[ 7] * m[14]) - |
1963 | m[ 5] * (m[ 2] * m[15] - m[ 3] * m[14]) - |
1964 | m[13] * (m[ 3] * m[ 6] - m[ 2] * m[ 7]); |
1965 | |
1966 | inv[ 3] = m[ 1] * (m[ 7] * m[10] - m[ 6] * m[11]) - |
1967 | m[ 5] * (m[ 3] * m[10] - m[ 2] * m[11]) - |
1968 | m[ 9] * (m[ 2] * m[ 7] - m[ 3] * m[ 6]); |
1969 | |
1970 | inv[ 4] = m[ 4] * (m[11] * m[14] - m[10] * m[15]) - |
1971 | m[ 8] * (m[ 7] * m[14] - m[ 6] * m[15]) - |
1972 | m[12] * (m[ 6] * m[11] - m[ 7] * m[10]); |
1973 | |
1974 | inv[ 5] = m[ 0] * (m[10] * m[15] - m[11] * m[14]) - |
1975 | m[ 8] * (m[ 2] * m[15] - m[ 3] * m[14]) - |
1976 | m[12] * (m[ 3] * m[10] - m[ 2] * m[11]); |
1977 | |
1978 | inv[ 6] = m[ 0] * (m[ 7] * m[14] - m[ 6] * m[15]) - |
1979 | m[ 4] * (m[ 3] * m[14] - m[ 2] * m[15]) - |
1980 | m[12] * (m[ 2] * m[ 7] - m[ 3] * m[ 6]); |
1981 | |
1982 | inv[ 7] = m[ 0] * (m[ 6] * m[11] - m[ 7] * m[10]) - |
1983 | m[ 4] * (m[ 2] * m[11] - m[ 3] * m[10]) - |
1984 | m[ 8] * (m[ 3] * m[ 6] - m[ 2] * m[ 7]); |
1985 | |
1986 | inv[ 8] = m[ 4] * (m[ 9] * m[15] - m[11] * m[13]) - |
1987 | m[ 8] * (m[ 5] * m[15] - m[ 7] * m[13]) - |
1988 | m[12] * (m[ 7] * m[ 9] - m[ 5] * m[11]); |
1989 | |
1990 | inv[ 9] = m[ 0] * (m[11] * m[13] - m[ 9] * m[15]) - |
1991 | m[ 8] * (m[ 3] * m[13] - m[ 1] * m[15]) - |
1992 | m[12] * (m[ 1] * m[11] - m[ 3] * m[ 9]); |
1993 | |
1994 | inv[10] = m[ 0] * (m[ 5] * m[15] - m[ 7] * m[13]) - |
1995 | m[ 4] * (m[ 1] * m[15] - m[ 3] * m[13]) - |
1996 | m[12] * (m[ 3] * m[ 5] - m[ 1] * m[ 7]); |
1997 | |
1998 | inv[11] = m[ 0] * (m[ 7] * m[ 9] - m[ 5] * m[11]) - |
1999 | m[ 4] * (m[ 3] * m[ 9] - m[ 1] * m[11]) - |
2000 | m[ 8] * (m[ 1] * m[ 7] - m[ 3] * m[ 5]); |
2001 | |
2002 | inv[12] = m[ 4] * (m[10] * m[13] - m[ 9] * m[14]) - |
2003 | m[ 8] * (m[ 6] * m[13] - m[ 5] * m[14]) - |
2004 | m[12] * (m[ 5] * m[10] - m[ 6] * m[ 9]); |
2005 | |
2006 | inv[13] = m[ 0] * (m[ 9] * m[14] - m[10] * m[13]) - |
2007 | m[ 8] * (m[ 1] * m[14] - m[ 2] * m[13]) - |
2008 | m[12] * (m[ 2] * m[ 9] - m[ 1] * m[10]); |
2009 | |
2010 | inv[14] = m[ 0] * (m[ 6] * m[13] - m[ 5] * m[14]) - |
2011 | m[ 4] * (m[ 2] * m[13] - m[ 1] * m[14]) - |
2012 | m[12] * (m[ 1] * m[ 6] - m[ 2] * m[ 5]); |
2013 | |
2014 | inv[15] = m[ 0] * (m[ 5] * m[10] - m[ 6] * m[ 9]) - |
2015 | m[ 4] * (m[ 1] * m[10] - m[ 2] * m[ 9]) - |
2016 | m[ 8] * (m[ 2] * m[ 5] - m[ 1] * m[ 6]); |
2017 | |
2018 | T det = m[0] * inv[ 0] + |
2019 | m[1] * inv[ 4] + |
2020 | m[2] * inv[ 8] + |
2021 | m[3] * inv[12]; |
2022 | |
2023 | if (det == T (0.0)) return; |
2024 | |
2025 | det = T (1.0) / det; |
2026 | |
68333c8f |
2027 | for (Standard_Integer i = 0; i < 16; ++i) |
e276548b |
2028 | inv[i] *= det; |
2029 | } |
2030 | |
2031 | // ======================================================================= |
2032 | // function : GenerateCornerRays |
2033 | // purpose : Generates primary rays for corners of screen quad |
2034 | // ======================================================================= |
2035 | void GenerateCornerRays (const GLdouble theInvModelProj[16], |
68333c8f |
2036 | cl_float theOrigins[16], |
2037 | cl_float theDirects[16]) |
e276548b |
2038 | { |
68333c8f |
2039 | Standard_Integer aOriginIndex = 0; |
2040 | Standard_Integer aDirectIndex = 0; |
e276548b |
2041 | |
68333c8f |
2042 | for (Standard_Integer y = -1; y <= 1; y += 2) |
e276548b |
2043 | { |
68333c8f |
2044 | for (Standard_Integer x = -1; x <= 1; x += 2) |
e276548b |
2045 | { |
265d4508 |
2046 | BVH_Vec4f aOrigin (float(x), |
2047 | float(y), |
2048 | -1.f, |
2049 | 1.f); |
e276548b |
2050 | |
2051 | aOrigin = MatVecMult (theInvModelProj, aOrigin); |
b5ac8292 |
2052 | aOrigin.x() = aOrigin.x() / aOrigin.w(); |
2053 | aOrigin.y() = aOrigin.y() / aOrigin.w(); |
2054 | aOrigin.z() = aOrigin.z() / aOrigin.w(); |
2055 | aOrigin.w() = 1.f; |
e276548b |
2056 | |
265d4508 |
2057 | BVH_Vec4f aDirect (float(x), |
2058 | float(y), |
2059 | 1.f, |
2060 | 1.f); |
e276548b |
2061 | |
b5ac8292 |
2062 | aDirect = MatVecMult (theInvModelProj, aDirect); |
2063 | aDirect.x() = aDirect.x() / aDirect.w(); |
2064 | aDirect.y() = aDirect.y() / aDirect.w(); |
2065 | aDirect.z() = aDirect.z() / aDirect.w(); |
2066 | aDirect.w() = 1.f; |
2067 | |
2068 | aDirect = aDirect - aOrigin; |
e276548b |
2069 | |
2070 | GLdouble aInvLen = 1.f / sqrt (aDirect.x() * aDirect.x() + |
2071 | aDirect.y() * aDirect.y() + |
2072 | aDirect.z() * aDirect.z()); |
2073 | |
2074 | theOrigins [aOriginIndex++] = static_cast<GLfloat> (aOrigin.x()); |
2075 | theOrigins [aOriginIndex++] = static_cast<GLfloat> (aOrigin.y()); |
2076 | theOrigins [aOriginIndex++] = static_cast<GLfloat> (aOrigin.z()); |
2077 | theOrigins [aOriginIndex++] = 1.f; |
2078 | |
2079 | theDirects [aDirectIndex++] = static_cast<GLfloat> (aDirect.x() * aInvLen); |
2080 | theDirects [aDirectIndex++] = static_cast<GLfloat> (aDirect.y() * aInvLen); |
2081 | theDirects [aDirectIndex++] = static_cast<GLfloat> (aDirect.z() * aInvLen); |
2082 | theDirects [aDirectIndex++] = 0.f; |
2083 | } |
2084 | } |
2085 | } |
2086 | |
2087 | // ======================================================================= |
2088 | // function : Raytrace |
2089 | // purpose : Redraws the window using OpenCL ray tracing |
2090 | // ======================================================================= |
2091 | Standard_Boolean OpenGl_Workspace::Raytrace (const Graphic3d_CView& theCView, |
68333c8f |
2092 | const Standard_Integer theSizeX, |
2093 | const Standard_Integer theSizeY, |
e276548b |
2094 | const Tint theToSwap) |
2095 | { |
2096 | if (!InitOpenCL()) |
2097 | return Standard_False; |
2098 | |
2099 | if (!ResizeRaytraceOutputBuffer (theSizeX, theSizeY)) |
2100 | return Standard_False; |
2101 | |
2102 | if (!UpdateRaytraceEnvironmentMap()) |
2103 | return Standard_False; |
2104 | |
2105 | if (!UpdateRaytraceGeometry (Standard_True)) |
2106 | return Standard_False; |
2107 | |
2108 | // Get model-view and projection matrices |
2109 | TColStd_Array2OfReal theOrientation (0, 3, 0, 3); |
2110 | TColStd_Array2OfReal theViewMapping (0, 3, 0, 3); |
2111 | |
b5ac8292 |
2112 | myView->GetMatrices (theOrientation, theViewMapping); |
e276548b |
2113 | |
2114 | GLdouble aOrientationMatrix[16]; |
2115 | GLdouble aViewMappingMatrix[16]; |
2116 | GLdouble aOrientationInvers[16]; |
2117 | |
68333c8f |
2118 | for (Standard_Integer j = 0; j < 4; ++j) |
2119 | for (Standard_Integer i = 0; i < 4; ++i) |
e276548b |
2120 | { |
2121 | aOrientationMatrix [4 * j + i] = theOrientation (i, j); |
2122 | aViewMappingMatrix [4 * j + i] = theViewMapping (i, j); |
2123 | } |
2124 | |
2125 | ComputeInverseMatrix (aOrientationMatrix, aOrientationInvers); |
2126 | |
2127 | if (!UpdateRaytraceLightSources (aOrientationInvers)) |
2128 | return Standard_False; |
2129 | |
2130 | // Generate primary rays for corners of the screen quad |
2131 | glMatrixMode (GL_MODELVIEW); |
2132 | |
2133 | glLoadMatrixd (aViewMappingMatrix); |
2134 | glMultMatrixd (aOrientationMatrix); |
2135 | |
2136 | GLdouble aModelProject[16]; |
2137 | GLdouble aInvModelProj[16]; |
2138 | |
2139 | glGetDoublev (GL_MODELVIEW_MATRIX, aModelProject); |
2140 | |
2141 | ComputeInverseMatrix (aModelProject, aInvModelProj); |
2142 | |
2143 | GLfloat aOrigins[16]; |
2144 | GLfloat aDirects[16]; |
2145 | |
2146 | GenerateCornerRays (aInvModelProj, |
2147 | aOrigins, |
2148 | aDirects); |
2149 | |
2150 | // Compute ray-traced image using OpenCL kernel |
68333c8f |
2151 | cl_mem anImages[] = { myRaytraceOutputImage, myRaytraceOutputImageAA }; |
265d4508 |
2152 | cl_int anError = clEnqueueAcquireGLObjects (myComputeQueue, |
e276548b |
2153 | 2, anImages, |
2154 | 0, NULL, NULL); |
265d4508 |
2155 | clFinish (myComputeQueue); |
e276548b |
2156 | |
2157 | if (myIsRaytraceDataValid) |
2158 | { |
2159 | RunRaytraceOpenCLKernels (theCView, |
2160 | aOrigins, |
2161 | aDirects, |
2162 | theSizeX, |
2163 | theSizeY); |
2164 | } |
2165 | |
265d4508 |
2166 | anError |= clEnqueueReleaseGLObjects (myComputeQueue, |
e276548b |
2167 | 2, anImages, |
2168 | 0, NULL, NULL); |
265d4508 |
2169 | clFinish (myComputeQueue); |
e276548b |
2170 | |
2171 | // Draw background |
2172 | glPushAttrib (GL_ENABLE_BIT | |
2173 | GL_CURRENT_BIT | |
2174 | GL_COLOR_BUFFER_BIT | |
2175 | GL_DEPTH_BUFFER_BIT); |
2176 | |
2177 | glDisable (GL_DEPTH_TEST); |
2178 | |
2179 | if (NamedStatus & OPENGL_NS_WHITEBACK) |
2180 | { |
2181 | glClearColor (1.0f, 1.0f, 1.0f, 1.0f); |
2182 | } |
2183 | else |
2184 | { |
2185 | glClearColor (myBgColor.rgb[0], |
2186 | myBgColor.rgb[1], |
2187 | myBgColor.rgb[2], |
2188 | 1.0f); |
2189 | } |
2190 | |
2191 | glClear (GL_COLOR_BUFFER_BIT); |
2192 | |
2193 | Handle(OpenGl_Workspace) aWorkspace (this); |
2194 | myView->DrawBackground (aWorkspace); |
2195 | |
2196 | // Draw dummy quad to show result image |
2197 | glEnable (GL_COLOR_MATERIAL); |
2198 | glEnable (GL_BLEND); |
2199 | |
2200 | glDisable (GL_DEPTH_TEST); |
2201 | |
2202 | glBlendFunc (GL_ONE, GL_SRC_ALPHA); |
2203 | |
2204 | glEnable (GL_TEXTURE_RECTANGLE); |
2205 | |
2206 | glMatrixMode (GL_PROJECTION); |
2207 | glLoadIdentity(); |
2208 | |
2209 | glMatrixMode (GL_MODELVIEW); |
2210 | glLoadIdentity(); |
2211 | |
2212 | glColor3f (1.0f, 1.0f, 1.0f); |
2213 | |
68333c8f |
2214 | if (!theCView.IsAntialiasingEnabled) |
2215 | myRaytraceOutputTexture->Bind (myGlContext); |
2216 | else |
2217 | myRaytraceOutputTextureAA->Bind (myGlContext); |
e276548b |
2218 | |
2219 | if (myIsRaytraceDataValid) |
2220 | { |
2221 | glBegin (GL_QUADS); |
2222 | { |
2223 | glTexCoord2i ( 0, 0); glVertex2f (-1.f, -1.f); |
2224 | glTexCoord2i ( 0, theSizeY); glVertex2f (-1.f, 1.f); |
2225 | glTexCoord2i (theSizeX, theSizeY); glVertex2f ( 1.f, 1.f); |
2226 | glTexCoord2i (theSizeX, 0); glVertex2f ( 1.f, -1.f); |
2227 | } |
2228 | glEnd(); |
2229 | } |
2230 | |
2231 | glPopAttrib(); |
2232 | |
2233 | // Swap the buffers |
2234 | if (theToSwap) |
2235 | { |
2236 | GetGlContext()->SwapBuffers(); |
2237 | myBackBufferRestored = Standard_False; |
2238 | } |
2239 | else |
2240 | glFlush(); |
2241 | |
2242 | return Standard_True; |
2243 | } |
2244 | |
2245 | #endif |