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