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