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