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