diff --git a/kernels/common/buffer.h b/kernels/common/buffer.h index c761bf7a98..340e962520 100644 --- a/kernels/common/buffer.h +++ b/kernels/common/buffer.h @@ -147,7 +147,7 @@ namespace embree public: /*! Buffer construction */ RawBufferView() - : ptr_ofs(nullptr), stride(0), num(0), format(RTC_FORMAT_UNDEFINED), modCounter(1), modified(true), userData(0) {} + : ptr_ofs(nullptr), dptr_ofs(nullptr), stride(0), num(0), format(RTC_FORMAT_UNDEFINED), modCounter(1), modified(true), userData(0) {} public: /*! sets the buffer view */ diff --git a/kernels/common/scene.cpp b/kernels/common/scene.cpp index 7373a86c2c..4c987b1aed 100644 --- a/kernels/common/scene.cpp +++ b/kernels/common/scene.cpp @@ -49,6 +49,7 @@ namespace embree maxTimeSegments(0), geometries_device(nullptr), geometries_data_device(nullptr), + num_geometries_device(0), taskGroup(new TaskGroup()), progressInterface(this), progress_monitor_function(nullptr), progress_monitor_ptr(nullptr), progress_monitor_counter(0) { @@ -71,6 +72,15 @@ namespace embree Scene::~Scene() noexcept { +#if defined(EMBREE_SYCL_SUPPORT) + if (geometries_data_device) { + device->free(geometries_data_device); + } + if (geometries_device) { + device->free(geometries_device); + } +#endif + device->refDec(); } @@ -922,7 +932,7 @@ namespace embree } catch (...) { accels_clear(); - Lock lock(taskGroup->schedulerMutex);, *global_gpu_queue + Lock lock(taskGroup->schedulerMutex); taskGroup->scheduler = nullptr; throw; } diff --git a/kernels/common/scene_grid_mesh.cpp b/kernels/common/scene_grid_mesh.cpp index 5a5a573688..2dead48268 100644 --- a/kernels/common/scene_grid_mesh.cpp +++ b/kernels/common/scene_grid_mesh.cpp @@ -204,6 +204,8 @@ namespace embree interpolate_impl<4>(args); } +#if defined(EMBREE_SYCL_SUPPORT) + size_t GridMesh::getGeometryDataDeviceByteSize() const { size_t byte_size = sizeof(GridMesh); byte_size += numTimeSteps * sizeof(BufferView); @@ -213,8 +215,6 @@ namespace embree return 16 * ((byte_size + 15) / 16); } -#if defined(EMBREE_SYCL_SUPPORT) - void GridMesh::convertToDeviceRepresentation(size_t offset, char* data_host, char* data_device) const { GridMesh* mesh = (GridMesh*)(data_host + offset); std::memcpy(data_host + offset, (void*)this, sizeof(GridMesh)); diff --git a/kernels/common/scene_grid_mesh.h b/kernels/common/scene_grid_mesh.h index 8d3eef97e4..b98ec6948e 100644 --- a/kernels/common/scene_grid_mesh.h +++ b/kernels/common/scene_grid_mesh.h @@ -55,9 +55,14 @@ namespace embree void commit(); bool verify(); void interpolate(const RTCInterpolateArguments* const args); + +#if defined(EMBREE_SYCL_SUPPORT) + size_t getGeometryDataDeviceByteSize() const; void convertToDeviceRepresentation(size_t offset, char* data_host, char* data_device) const; +#endif + template void interpolate_impl(const RTCInterpolateArguments* const args) { diff --git a/kernels/sycl/scene_sycl.cpp b/kernels/sycl/scene_sycl.cpp index d3cc9a117d..9bc89da8d1 100644 --- a/kernels/sycl/scene_sycl.cpp +++ b/kernels/sycl/scene_sycl.cpp @@ -163,6 +163,9 @@ void Scene::syncWithDevice(sycl::queue* queue_in) Geometry** geometries_host = (Geometry**)device->malloc(sizeof(Geometry*)*geometries.size(), 16, EmbreeMemoryType::UNKNOWN); char* geometries_data_host = (char*)device->malloc(geometry_data_byte_size, 16, EmbreeMemoryType::UNKNOWN); + std::memset(geometries_host, 0, sizeof(Geometry*) * geometries.size()); + std::memset(geometries_data_host, 0, geometry_data_byte_size); + for (size_t i = 0; i < geometries.size(); ++i) { geometries[i]->convertToDeviceRepresentation(offsets[i], geometries_data_host, geometries_data_device); geometries_host[i] = (Geometry*)(geometries_data_device + offsets[i]); diff --git a/tutorials/CMakeLists.txt b/tutorials/CMakeLists.txt index 0cf51a943f..28e5bec4eb 100644 --- a/tutorials/CMakeLists.txt +++ b/tutorials/CMakeLists.txt @@ -113,6 +113,7 @@ ADD_SUBDIRECTORY(next_hit) ADD_SUBDIRECTORY(multiscene_geometry) ADD_SUBDIRECTORY(ray_mask) ADD_SUBDIRECTORY(forest) +ADD_SUBDIRECTORY(debug_device_memory) ADD_SUBDIRECTORY(embree_tests) diff --git a/tutorials/curve_geometry/curve_geometry_device.cpp b/tutorials/curve_geometry/curve_geometry_device.cpp index a5b7ce7e89..d6822ae916 100644 --- a/tutorials/curve_geometry/curve_geometry_device.cpp +++ b/tutorials/curve_geometry/curve_geometry_device.cpp @@ -190,7 +190,7 @@ void renderPixelStandard(const TutorialData& data, RTCOccludedArguments sargs; rtcInitOccludedArguments(&sargs); sargs.feature_mask = (RTCFeatureFlags) (FEATURE_MASK); - //rtcOccluded1(data.g_scene,RTCRay_(shadow),&sargs); + rtcOccluded1(data.g_scene,RTCRay_(shadow),&sargs); RayStats_addShadowRay(stats); /* add light contribution */ diff --git a/tutorials/debug_device_memory/CMakeLists.txt b/tutorials/debug_device_memory/CMakeLists.txt new file mode 100644 index 0000000000..5332594037 --- /dev/null +++ b/tutorials/debug_device_memory/CMakeLists.txt @@ -0,0 +1,11 @@ +## Copyright 2009-2021 Intel Corporation +## SPDX-License-Identifier: Apache-2.0 + +IF (EMBREE_GEOMETRY_TRIANGLE) + +INCLUDE(tutorial) +ADD_TUTORIAL(debug_device_memory) +ADD_TUTORIAL_ISPC(debug_device_memory) +ADD_TUTORIAL_SYCL(debug_device_memory) + +ENDIF() \ No newline at end of file diff --git a/tutorials/debug_device_memory/debug_device_memory.cpp b/tutorials/debug_device_memory/debug_device_memory.cpp new file mode 100644 index 0000000000..307b5a053a --- /dev/null +++ b/tutorials/debug_device_memory/debug_device_memory.cpp @@ -0,0 +1,39 @@ +// Copyright 2009-2021 Intel Corporation +// SPDX-License-Identifier: Apache-2.0 + +#include "../common/tutorial/tutorial.h" +#include "../common/tutorial/benchmark_render.h" + +#if defined(EMBREE_SYCL_TUTORIAL) +# define NAME "debug_device_memory_sycl" +# define FEATURES FEATURE_RTCORE | FEATURE_SYCL +#else +# define NAME "debug_device_memory" +# define FEATURES FEATURE_RTCORE +#endif + +namespace embree +{ + struct Tutorial : public TutorialApplication + { + Tutorial() + : TutorialApplication(NAME,FEATURES) + { + /* set default camera */ + camera.from = Vec3fa(1.5f,1.5f,-1.5f); + camera.to = Vec3fa(0.0f,0.0f,0.0f); + } + }; + +} + +int main(int argc, char** argv) { + try { + if (embree::TutorialBenchmark::benchmark(argc, argv)) { + return embree::TutorialBenchmark(embree::renderBenchFunc).main(argc, argv, "debug_device_memory"); + } + return embree::Tutorial().main(argc,argv); + } catch (std::exception& e ) { + std::cerr << "Exception caught: " << e.what() << std::endl; + } +} diff --git a/tutorials/debug_device_memory/debug_device_memory_device.cpp b/tutorials/debug_device_memory/debug_device_memory_device.cpp new file mode 100644 index 0000000000..ac66c9fd2d --- /dev/null +++ b/tutorials/debug_device_memory/debug_device_memory_device.cpp @@ -0,0 +1,455 @@ +// Copyright 2009-2021 Intel Corporation +// SPDX-License-Identifier: Apache-2.0 + +#include "debug_device_memory_device.h" + +#include + +namespace embree { + +AffineSpace3fa axfm0; +AffineSpace3fa axfm1; + +/* all features required by this tutorial */ +#define FEATURE_MASK \ + RTC_FEATURE_FLAG_TRIANGLE | \ + RTC_FEATURE_FLAG_INSTANCE | \ + RTC_FEATURE_FLAG_MOTION_BLUR + +RTCScene g_scene = nullptr; +RTCScene g_scene1 = nullptr; +TutorialData data; + +/* adds a cube to the scene */ +unsigned int addCube (RTCScene scene_i) +{ + /* create a triangulated cube with 12 triangles and 8 vertices */ + RTCGeometry mesh = rtcNewGeometry(g_device, RTC_GEOMETRY_TYPE_TRIANGLE); + + rtcSetGeometryTimeStepCount(mesh, 2); + + /* create face and vertex color arrays */ + data.face_colors = (Vec3fa*) alignedUSMMalloc((12)*sizeof(Vec3fa),16); + data.vertex_colors = (Vec3fa*) alignedUSMMalloc((8)*sizeof(Vec3fa),16); + + /* set vertices and vertex colors */ + Vertex* vertices_device; + Vertex* vertices1_device; + Vertex* vertices = (Vertex*) rtcSetNewGeometryBufferXPU(mesh,RTC_BUFFER_TYPE_VERTEX,0,RTC_FORMAT_FLOAT3,sizeof(Vertex),8,(void**)&vertices_device); + Vertex* vertices1 = (Vertex*) rtcSetNewGeometryBufferXPU(mesh,RTC_BUFFER_TYPE_VERTEX,1,RTC_FORMAT_FLOAT3,sizeof(Vertex),8,(void**)&vertices1_device); + +#if defined(EMBREE_SYCL_TUTORIAL) + sycl::usm::alloc allocType; + allocType = sycl::get_pointer_type(vertices, global_gpu_queue->get_context()); + if (allocType != sycl::usm::alloc::host) + std::cout << "cube vertices have wrong alloc type!" <get_context()); + if (allocType != sycl::usm::alloc::device) + std::cout << "cube vertices have wrong alloc type!" <memcpy(vertices_device, vertices, 8 * sizeof(Vertex)); + global_gpu_queue->memcpy(vertices1_device, vertices1, 8 * sizeof(Vertex)); + global_gpu_queue->memcpy(triangles_device, triangles, 12 * sizeof(Triangle)); + global_gpu_queue->wait_and_throw(); +#endif + + rtcCommitGeometry(mesh); + unsigned int geomID = rtcAttachGeometry(scene_i,mesh); + rtcReleaseGeometry(mesh); + return geomID; +} + +unsigned int addCubeShared (RTCScene scene_i) +{ + /* create a triangulated cube with 12 triangles and 8 vertices */ + RTCGeometry mesh = rtcNewGeometry(g_device, RTC_GEOMETRY_TYPE_TRIANGLE); + + rtcSetGeometryTimeStepCount(mesh, 2); + + /* create face and vertex color arrays */ + data.face_colors = (Vec3fa*) alignedUSMMalloc((12)*sizeof(Vec3fa),16); + data.vertex_colors = (Vec3fa*) alignedUSMMalloc((8)*sizeof(Vec3fa),16); + + Vertex* vertices = (Vertex*) alignedMalloc(8*sizeof(Vertex), 16); + Vertex* vertices1 = (Vertex*) alignedMalloc(8*sizeof(Vertex), 16); + Triangle* triangles = (Triangle*) alignedMalloc(12*sizeof(Triangle), 16); + +#if defined(EMBREE_SYCL_TUTORIAL) + Vertex* vertices_device = sycl::aligned_alloc_device(16, 8, *global_gpu_queue); + Vertex* vertices1_device = sycl::aligned_alloc_device(16, 8, *global_gpu_queue); + Triangle* triangles_device = sycl::aligned_alloc_device(16, 12, *global_gpu_queue); +#else + Vertex* vertices_device = nullptr; + Vertex* vertices1_device = nullptr; + Triangle* triangles_device = nullptr; +#endif + + rtcSetSharedGeometryBufferXPU(mesh,RTC_BUFFER_TYPE_VERTEX,0,RTC_FORMAT_FLOAT3,vertices,vertices_device, 0, sizeof(Vertex),8); + rtcSetSharedGeometryBufferXPU(mesh,RTC_BUFFER_TYPE_VERTEX,1,RTC_FORMAT_FLOAT3,vertices1,vertices1_device, 0, sizeof(Vertex),8); + + data.vertex_colors[0] = Vec3fa(0,0,0); + data.vertex_colors[1] = Vec3fa(0,0,1); + data.vertex_colors[2] = Vec3fa(0,1,0); + data.vertex_colors[3] = Vec3fa(0,1,1); + data.vertex_colors[4] = Vec3fa(1,0,0); + data.vertex_colors[5] = Vec3fa(1,0,1); + data.vertex_colors[6] = Vec3fa(1,1,0); + data.vertex_colors[7] = Vec3fa(1,1,1); + + vertices[0].x = -1; vertices[0].y = -1; vertices[0].z = -1; + vertices[1].x = -1; vertices[1].y = -1; vertices[1].z = +1; + vertices[2].x = -1; vertices[2].y = +1; vertices[2].z = -1; + vertices[3].x = -1; vertices[3].y = +1; vertices[3].z = +1; + vertices[4].x = +1; vertices[4].y = -1; vertices[4].z = -1; + vertices[5].x = +1; vertices[5].y = -1; vertices[5].z = +1; + vertices[6].x = +1; vertices[6].y = +1; vertices[6].z = -1; + vertices[7].x = +1; vertices[7].y = +1; vertices[7].z = +1; + + vertices1[0].x = -1 + 0.5f; vertices1[0].y = -1 + 0.5f; vertices1[0].z = -1 + 0.5f; + vertices1[1].x = -1 + 0.5f; vertices1[1].y = -1 + 0.5f; vertices1[1].z = +1 + 0.5f; + vertices1[2].x = -1 + 0.5f; vertices1[2].y = +1 + 0.5f; vertices1[2].z = -1 + 0.5f; + vertices1[3].x = -1 + 0.5f; vertices1[3].y = +1 + 0.5f; vertices1[3].z = +1 + 0.5f; + vertices1[4].x = +1 + 0.5f; vertices1[4].y = -1 + 0.5f; vertices1[4].z = -1 + 0.5f; + vertices1[5].x = +1 + 0.5f; vertices1[5].y = -1 + 0.5f; vertices1[5].z = +1 + 0.5f; + vertices1[6].x = +1 + 0.5f; vertices1[6].y = +1 + 0.5f; vertices1[6].z = -1 + 0.5f; + vertices1[7].x = +1 + 0.5f; vertices1[7].y = +1 + 0.5f; vertices1[7].z = +1 + 0.5f; + + /* set triangles and face colors */ + int tri = 0; + rtcSetSharedGeometryBufferXPU(mesh,RTC_BUFFER_TYPE_INDEX,0,RTC_FORMAT_UINT3,triangles,triangles_device,0,sizeof(Triangle),12); + + // left side + data.face_colors[tri] = Vec3fa(1,0,0); triangles[tri].v0 = 0; triangles[tri].v1 = 1; triangles[tri].v2 = 2; tri++; + data.face_colors[tri] = Vec3fa(1,0,0); triangles[tri].v0 = 1; triangles[tri].v1 = 3; triangles[tri].v2 = 2; tri++; + + // right side + data.face_colors[tri] = Vec3fa(0,1,0); triangles[tri].v0 = 4; triangles[tri].v1 = 6; triangles[tri].v2 = 5; tri++; + data.face_colors[tri] = Vec3fa(0,1,0); triangles[tri].v0 = 5; triangles[tri].v1 = 6; triangles[tri].v2 = 7; tri++; + + // bottom side + data.face_colors[tri] = Vec3fa(0.5f); triangles[tri].v0 = 0; triangles[tri].v1 = 4; triangles[tri].v2 = 1; tri++; + data.face_colors[tri] = Vec3fa(0.5f); triangles[tri].v0 = 1; triangles[tri].v1 = 4; triangles[tri].v2 = 5; tri++; + + // top side + data.face_colors[tri] = Vec3fa(1.0f); triangles[tri].v0 = 2; triangles[tri].v1 = 3; triangles[tri].v2 = 6; tri++; + data.face_colors[tri] = Vec3fa(1.0f); triangles[tri].v0 = 3; triangles[tri].v1 = 7; triangles[tri].v2 = 6; tri++; + + // front side + data.face_colors[tri] = Vec3fa(0,0,1); triangles[tri].v0 = 0; triangles[tri].v1 = 2; triangles[tri].v2 = 4; tri++; + data.face_colors[tri] = Vec3fa(0,0,1); triangles[tri].v0 = 2; triangles[tri].v1 = 6; triangles[tri].v2 = 4; tri++; + + // back side + data.face_colors[tri] = Vec3fa(1,1,0); triangles[tri].v0 = 1; triangles[tri].v1 = 5; triangles[tri].v2 = 3; tri++; + data.face_colors[tri] = Vec3fa(1,1,0); triangles[tri].v0 = 3; triangles[tri].v1 = 5; triangles[tri].v2 = 7; tri++; + + rtcSetGeometryVertexAttributeCount(mesh,1); + rtcSetSharedGeometryBuffer(mesh,RTC_BUFFER_TYPE_VERTEX_ATTRIBUTE,0,RTC_FORMAT_FLOAT3,data.vertex_colors,0,sizeof(Vec3fa),8); + +#if defined(EMBREE_SYCL_TUTORIAL) + global_gpu_queue->memcpy(vertices_device, vertices, 8 * sizeof(Vertex)); + global_gpu_queue->memcpy(vertices1_device, vertices1, 8 * sizeof(Vertex)); + global_gpu_queue->memcpy(triangles_device, triangles, 12 * sizeof(Triangle)); +#endif + + rtcCommitGeometry(mesh); + unsigned int geomID = rtcAttachGeometry(scene_i,mesh); + rtcReleaseGeometry(mesh); + return geomID; +} + +/* adds a ground plane to the scene */ +unsigned int addGroundPlane (RTCScene scene_i) +{ + /* create a triangulated plane with 2 triangles and 4 vertices */ + RTCGeometry mesh = rtcNewGeometry (g_device, RTC_GEOMETRY_TYPE_TRIANGLE); + + /* set vertices */ + Vertex* vertices = (Vertex*) rtcSetNewGeometryBuffer(mesh,RTC_BUFFER_TYPE_VERTEX,0,RTC_FORMAT_FLOAT3,sizeof(Vertex),4); + vertices[0].x = -10; vertices[0].y = -2; vertices[0].z = -10; + vertices[1].x = -10; vertices[1].y = -2; vertices[1].z = +10; + vertices[2].x = +10; vertices[2].y = -2; vertices[2].z = -10; + vertices[3].x = +10; vertices[3].y = -2; vertices[3].z = +10; + + /* set triangles */ + Triangle* triangles = (Triangle*) rtcSetNewGeometryBuffer(mesh,RTC_BUFFER_TYPE_INDEX,0,RTC_FORMAT_UINT3,sizeof(Triangle),2); + triangles[0].v0 = 0; triangles[0].v1 = 1; triangles[0].v2 = 2; + triangles[1].v0 = 1; triangles[1].v1 = 3; triangles[1].v2 = 2; + + rtcCommitGeometry(mesh); + unsigned int geomID = rtcAttachGeometry(scene_i,mesh); + rtcReleaseGeometry(mesh); + return geomID; +} + +/* called by the C++ code for initialization */ +extern "C" void device_init (char* cfg) +{ + /* create scene */ + TutorialData_Constructor(&data); + g_scene = data.g_scene = rtcNewScene(g_device); + g_scene1 = rtcNewScene(g_device); + rtcSetSceneFlags(data.g_scene, RTC_SCENE_FLAG_PREFETCH_USM_SHARED_ON_GPU); + + /* add cube */ + addCube(g_scene1); + + RTCGeometry inst = rtcNewGeometry(g_device, RTC_GEOMETRY_TYPE_INSTANCE); + + rtcSetGeometryTimeStepCount(inst, 2); + rtcSetGeometryInstancedScene(inst, g_scene1); + + LinearSpace3fa xfm = one; + axfm0 = AffineSpace3fa(xfm,Vec3fa(0.f, 0.f, 0.f)); + axfm1 = AffineSpace3fa(xfm,Vec3fa(3.f, 0.f, 0.f)); + rtcSetGeometryTransform(inst,0,RTC_FORMAT_FLOAT4X4_COLUMN_MAJOR,(float*)&(axfm0.l.vx.x)); + rtcSetGeometryTransform(inst,1,RTC_FORMAT_FLOAT4X4_COLUMN_MAJOR,(float*)&(axfm1.l.vx.x)); + + rtcAttachGeometry(data.g_scene,inst); + rtcReleaseGeometry(inst); + rtcCommitGeometry(inst); + + /* add ground plane */ + addGroundPlane(data.g_scene); + + /* commit changes to scene */ +#if defined(EMBREE_SYCL_SUPPORT) && defined(SYCL_LANGUAGE_VERSION) + rtcCommitSceneWithQueue (g_scene1, *global_gpu_queue); + rtcCommitSceneWithQueue (data.g_scene, *global_gpu_queue); +#else + rtcCommitScene (g_scene1); + rtcCommitScene (data.g_scene); +#endif +} + +static inline uint32_t doodle(uint32_t x) +{ + x ^= x << 13; + x ^= x >> 17; + x ^= x << 5; + return x; +} + +static inline float doodlef(uint32_t x) +{ + return ((float)doodle(x)) / (float)(uint32_t(-1)); +} + +/* task that renders a single screen tile */ +void renderPixelStandard(const TutorialData& data, + int x, int y, + int* pixels, + const unsigned int width, + const unsigned int height, + const float time, + const ISPCCamera& camera, RayStats& stats) +{ + /* initialize ray */ + uint32_t state = doodle(x + y * width); + state = doodle(state); + float t = doodlef(state); + Ray ray(Vec3fa(camera.xfm.p), Vec3fa(normalize(x*camera.xfm.l.vx + y*camera.xfm.l.vy + camera.xfm.l.vz)), 0.0f, inf, t); + + /* intersect ray with scene */ + RTCIntersectArguments iargs; + rtcInitIntersectArguments(&iargs); + iargs.feature_mask = (RTCFeatureFlags) (FEATURE_MASK); + rtcIntersect1(data.g_scene,RTCRayHit_(ray),&iargs); + RayStats_addRay(stats); + + /* shade pixels */ + Vec3fa color = Vec3fa(0.0f); + if (ray.geomID != RTC_INVALID_GEOMETRY_ID || ray.instID[0] != RTC_INVALID_GEOMETRY_ID) + { +#if 1 + Vec3fa diffuse = data.face_colors[ray.primID]; + color = color + diffuse*0.5f; + Vec3fa lightDir = normalize(Vec3fa(-1,-1,-1)); + + /* initialize shadow ray */ + Ray shadow(ray.org + ray.tfar*ray.dir, neg(lightDir), 0.001f, inf, 1.f - t); + + /* trace shadow ray */ + RTCOccludedArguments sargs; + rtcInitOccludedArguments(&sargs); + sargs.feature_mask = (RTCFeatureFlags) (FEATURE_MASK); + rtcOccluded1(data.g_scene,RTCRay_(shadow),&sargs); + RayStats_addShadowRay(stats); + + /* add light contribution */ + if (shadow.tfar >= 0.0f) + color = color + diffuse*clamp(-dot(lightDir,normalize(ray.Ng)),0.0f,1.0f); +#else + +#if 0 + if (ray.geomID == 0) + color = Vec3fa(0.f, 0.f, 1.f); + else if (ray.geomID == 1) + color = Vec3fa(0.f, 1.f, 0.f); + else if (ray.geomID == 2) + color = Vec3fa(1.f, 0.f, 0.f); + else + color = Vec3fa(1.f); +#endif + +#if 0 + if (ray.primID == 0) + color = Vec3fa(0.f, 0.f, 1.f); + else if (ray.primID == 1) + color = Vec3fa(0.f, 1.f, 0.f); + else if (ray.primID == 2) + color = Vec3fa(1.f, 0.f, 0.f); + else + color = Vec3fa(1.f); +#endif + +#if 0 + color = Vec3fa(ray.u, ray.v, 0.f); +#endif + + color = Vec3fa(0.f, 1.f, 0.f); + +#endif + + } + + /* write color to framebuffer */ + unsigned int r = (unsigned int) (255.0f * clamp(color.x,0.0f,1.0f)); + unsigned int g = (unsigned int) (255.0f * clamp(color.y,0.0f,1.0f)); + unsigned int b = (unsigned int) (255.0f * clamp(color.z,0.0f,1.0f)); + pixels[y*width+x] = (b << 16) + (g << 8) + r; +} + +/* task that renders a single screen tile */ +void renderTileTask (int taskIndex, int threadIndex, int* pixels, + const unsigned int width, + const unsigned int height, + const float time, + const ISPCCamera& camera, + const int numTilesX, + const int numTilesY) +{ + const unsigned int tileY = taskIndex / numTilesX; + const unsigned int tileX = taskIndex - tileY * numTilesX; + const unsigned int x0 = tileX * TILE_SIZE_X; + const unsigned int x1 = min(x0+TILE_SIZE_X,width); + const unsigned int y0 = tileY * TILE_SIZE_Y; + const unsigned int y1 = min(y0+TILE_SIZE_Y,height); + + for (unsigned int y=y0; ysubmit([=](sycl::handler& cgh){ + const sycl::nd_range<2> nd_range = make_nd_range(height,width); + cgh.parallel_for(nd_range,[=](sycl::nd_item<2> item) { + const unsigned int x = item.get_global_id(1); if (x >= width ) return; + const unsigned int y = item.get_global_id(0); if (y >= height) return; + RayStats stats; + renderPixelStandard(ldata,x,y,pixels,width,height,time,camera,stats); + }); + }); + global_gpu_queue->wait_and_throw(); + + const auto t0 = event.template get_profiling_info(); + const auto t1 = event.template get_profiling_info(); + const double dt = (t1-t0)*1E-9; + ((ISPCCamera*)&camera)->render_time = dt; +#else + const int numTilesX = (width +TILE_SIZE_X-1)/TILE_SIZE_X; + const int numTilesY = (height+TILE_SIZE_Y-1)/TILE_SIZE_Y; + parallel_for(size_t(0),size_t(numTilesX*numTilesY),[&](const range& range) { + const int threadIndex = (int)TaskScheduler::threadIndex(); + for (size_t i=range.begin(); ig_scene = nullptr; + This->face_colors = nullptr; + This->vertex_colors = nullptr; +} + +inline void TutorialData_Destructor(TutorialData* This) +{ + rtcReleaseScene (This->g_scene); This->g_scene = nullptr; + alignedUSMFree(This->face_colors); This->face_colors = nullptr; + alignedUSMFree(This->vertex_colors); This->vertex_colors = nullptr; +} + +} // namespace embree diff --git a/tutorials/debug_device_memory/debug_device_memory_device.ispc b/tutorials/debug_device_memory/debug_device_memory_device.ispc new file mode 100644 index 0000000000..3f15d4dfca --- /dev/null +++ b/tutorials/debug_device_memory/debug_device_memory_device.ispc @@ -0,0 +1,227 @@ +// Copyright 2009-2021 Intel Corporation +// SPDX-License-Identifier: Apache-2.0 + +#include "debug_device_memory_device.isph" + +/* all features required by this tutorial */ +#define FEATURE_MASK \ + RTC_FEATURE_FLAG_TRIANGLE + +RTCScene g_scene = NULL; +uniform TutorialData data; + +/* adds a cube to the scene */ +uniform unsigned int addCube (RTCScene scene_i) +{ + /* create a triangulated cube with 12 triangles and 8 vertices */ + RTCGeometry mesh = rtcNewGeometry(g_device, RTC_GEOMETRY_TYPE_TRIANGLE); + + /* create face and vertex color arrays */ + data.face_colors = uniform new uniform Vec3f[12]; + data.vertex_colors = uniform new uniform Vec3fa[8]; + + /* set vertices and vertex colors */ + uniform Vertex* uniform vertices = (uniform Vertex* uniform) rtcSetNewGeometryBuffer(mesh,RTC_BUFFER_TYPE_VERTEX,0,RTC_FORMAT_FLOAT3,sizeof(uniform Vertex),8); + data.vertex_colors[0] = make_Vec3fa(0,0,0); vertices[0].x = -1; vertices[0].y = -1; vertices[0].z = -1; + data.vertex_colors[1] = make_Vec3fa(0,0,1); vertices[1].x = -1; vertices[1].y = -1; vertices[1].z = +1; + data.vertex_colors[2] = make_Vec3fa(0,1,0); vertices[2].x = -1; vertices[2].y = +1; vertices[2].z = -1; + data.vertex_colors[3] = make_Vec3fa(0,1,1); vertices[3].x = -1; vertices[3].y = +1; vertices[3].z = +1; + data.vertex_colors[4] = make_Vec3fa(1,0,0); vertices[4].x = +1; vertices[4].y = -1; vertices[4].z = -1; + data.vertex_colors[5] = make_Vec3fa(1,0,1); vertices[5].x = +1; vertices[5].y = -1; vertices[5].z = +1; + data.vertex_colors[6] = make_Vec3fa(1,1,0); vertices[6].x = +1; vertices[6].y = +1; vertices[6].z = -1; + data.vertex_colors[7] = make_Vec3fa(1,1,1); vertices[7].x = +1; vertices[7].y = +1; vertices[7].z = +1; + + /* set triangles and face colors */ + uniform int tri = 0; + uniform Triangle* uniform triangles = (uniform Triangle* uniform) rtcSetNewGeometryBuffer(mesh,RTC_BUFFER_TYPE_INDEX,0,RTC_FORMAT_UINT3,sizeof(uniform Triangle),12); + + // left side + data.face_colors[tri] = make_Vec3f(1,0,0); triangles[tri].v0 = 0; triangles[tri].v1 = 1; triangles[tri].v2 = 2; tri++; + data.face_colors[tri] = make_Vec3f(1,0,0); triangles[tri].v0 = 1; triangles[tri].v1 = 3; triangles[tri].v2 = 2; tri++; + + // right side + data.face_colors[tri] = make_Vec3f(0,1,0); triangles[tri].v0 = 4; triangles[tri].v1 = 6; triangles[tri].v2 = 5; tri++; + data.face_colors[tri] = make_Vec3f(0,1,0); triangles[tri].v0 = 5; triangles[tri].v1 = 6; triangles[tri].v2 = 7; tri++; + + // bottom side + data.face_colors[tri] = make_Vec3f(0.5f); triangles[tri].v0 = 0; triangles[tri].v1 = 4; triangles[tri].v2 = 1; tri++; + data.face_colors[tri] = make_Vec3f(0.5f); triangles[tri].v0 = 1; triangles[tri].v1 = 4; triangles[tri].v2 = 5; tri++; + + // top side + data.face_colors[tri] = make_Vec3f(1.0f); triangles[tri].v0 = 2; triangles[tri].v1 = 3; triangles[tri].v2 = 6; tri++; + data.face_colors[tri] = make_Vec3f(1.0f); triangles[tri].v0 = 3; triangles[tri].v1 = 7; triangles[tri].v2 = 6; tri++; + + // front side + data.face_colors[tri] = make_Vec3f(0,0,1); triangles[tri].v0 = 0; triangles[tri].v1 = 2; triangles[tri].v2 = 4; tri++; + data.face_colors[tri] = make_Vec3f(0,0,1); triangles[tri].v0 = 2; triangles[tri].v1 = 6; triangles[tri].v2 = 4; tri++; + + // back side + data.face_colors[tri] = make_Vec3f(1,1,0); triangles[tri].v0 = 1; triangles[tri].v1 = 5; triangles[tri].v2 = 3; tri++; + data.face_colors[tri] = make_Vec3f(1,1,0); triangles[tri].v0 = 3; triangles[tri].v1 = 5; triangles[tri].v2 = 7; tri++; + + rtcSetGeometryVertexAttributeCount(mesh,1); + rtcSetSharedGeometryBuffer(mesh,RTC_BUFFER_TYPE_VERTEX_ATTRIBUTE,0,RTC_FORMAT_FLOAT3,data.vertex_colors,0,sizeof(uniform Vec3fa),8); + + rtcCommitGeometry(mesh); + uniform unsigned int geomID = rtcAttachGeometry(scene_i,mesh); + rtcReleaseGeometry(mesh); + return geomID; +} + +/* adds a ground plane to the scene */ +uniform unsigned int addGroundPlane (RTCScene scene_i) +{ + /* create a triangulated plane with 2 triangles and 4 vertices */ + RTCGeometry mesh = rtcNewGeometry (g_device, RTC_GEOMETRY_TYPE_TRIANGLE); + + /* set vertices */ + uniform Vertex* uniform vertices = (uniform Vertex* uniform) rtcSetNewGeometryBuffer(mesh,RTC_BUFFER_TYPE_VERTEX,0,RTC_FORMAT_FLOAT3,sizeof(uniform Vertex),4); + vertices[0].x = -10; vertices[0].y = -2; vertices[0].z = -10; + vertices[1].x = -10; vertices[1].y = -2; vertices[1].z = +10; + vertices[2].x = +10; vertices[2].y = -2; vertices[2].z = -10; + vertices[3].x = +10; vertices[3].y = -2; vertices[3].z = +10; + + /* set triangles */ + uniform Triangle* uniform triangles = (uniform Triangle* uniform) rtcSetNewGeometryBuffer(mesh,RTC_BUFFER_TYPE_INDEX,0,RTC_FORMAT_UINT3,sizeof(uniform Triangle),2); + triangles[0].v0 = 0; triangles[0].v1 = 1; triangles[0].v2 = 2; + triangles[1].v0 = 1; triangles[1].v1 = 3; triangles[1].v2 = 2; + + rtcCommitGeometry(mesh); + uniform unsigned int geomID = rtcAttachGeometry(scene_i,mesh); + rtcReleaseGeometry(mesh); + return geomID; +} + +/* called by the C++ code for initialization */ +export void device_init (uniform int8* uniform cfg) +{ + /* create scene */ + TutorialData_Constructor(&data); + g_scene = data.g_scene = rtcNewScene(g_device); + + /* add cube */ + addCube(data.g_scene); + + /* add ground plane */ + addGroundPlane(data.g_scene); + + /* commit changes to scene */ + rtcCommitScene (data.g_scene); +} + +/* task that renders a single screen tile */ +void renderPixelStandard(const uniform TutorialData& data, + int x, int y, + uniform int* uniform pixels, + const uniform unsigned int width, + const uniform unsigned int height, + const float time, + const uniform ISPCCamera& camera, uniform RayStats& stats) +{ + /* initialize ray */ + Ray ray = make_Ray(make_Vec3f(camera.xfm.p), make_Vec3f(normalize(x*camera.xfm.l.vx + y*camera.xfm.l.vy + camera.xfm.l.vz)), 0.0f, inf); + + /* intersect ray with scene */ + uniform RTCIntersectArguments iargs; + rtcInitIntersectArguments(&iargs); + iargs.feature_mask = (uniform RTCFeatureFlags) (FEATURE_MASK); + rtcIntersectV(data.g_scene,RTCRayHit_(ray),&iargs); + RayStats_addRay(stats); + + /* shade pixels */ + Vec3f color = make_Vec3f(0.0f); + if (ray.geomID != RTC_INVALID_GEOMETRY_ID) + { + Vec3f diffuse = data.face_colors[ray.primID]; + color = color + diffuse*0.5f; + Vec3f lightDir = normalize(make_Vec3f(-1,-1,-1)); + + /* initialize shadow ray */ + Ray shadow = make_Ray(ray.org + ray.tfar*ray.dir, neg(lightDir), 0.001f, inf, 0.0f); + + /* trace shadow ray */ + uniform RTCOccludedArguments sargs; + rtcInitOccludedArguments(&sargs); + sargs.feature_mask = (uniform RTCFeatureFlags) (FEATURE_MASK); + rtcOccludedV(data.g_scene,RTCRay_(shadow),&sargs); + RayStats_addShadowRay(stats); + + /* add light contribution */ + if (shadow.tfar >= 0.0f) + color = color + diffuse*clamp(-dot(lightDir,normalize(ray.Ng)),0.0f,1.0f); + } + + /* write color to framebuffer */ + unsigned int r = (unsigned int) (255.0f * clamp(color.x,0.0f,1.0f)); + unsigned int g = (unsigned int) (255.0f * clamp(color.y,0.0f,1.0f)); + unsigned int b = (unsigned int) (255.0f * clamp(color.z,0.0f,1.0f)); + pixels[y*width+x] = (b << 16) + (g << 8) + r; +} + +/* task that renders a single screen tile */ +task void renderTileTask(uniform int* uniform pixels, + const uniform unsigned int width, + const uniform unsigned int height, + const uniform float time, + const uniform ISPCCamera& camera, + const uniform int numTilesX, + const uniform int numTilesY) +{ + const uniform unsigned int tileY = taskIndex / numTilesX; + const uniform unsigned int tileX = taskIndex - tileY * numTilesX; + const uniform unsigned int x0 = tileX * TILE_SIZE_X; + const uniform unsigned int x1 = min(x0+TILE_SIZE_X,width); + const uniform unsigned int y0 = tileY * TILE_SIZE_Y; + const uniform unsigned int y1 = min(y0+TILE_SIZE_Y,height); + + foreach_tiled (y = y0 ... y1, x = x0 ... x1) + { + renderPixelStandard(data,x,y,pixels,width,height,time,camera,g_stats[threadIndex]); + } +} + +/* called by the C++ code to render */ +export void renderFrameStandard (uniform int* uniform pixels, + const uniform unsigned int width, + const uniform unsigned int height, + const uniform float time, + const uniform ISPCCamera& camera) +{ +#if defined(EMBREE_SYCL_TUTORIAL) && !defined(EMBREE_SYCL_RT_SIMULATION) + TutorialData ldata = data; + sycl::event event = global_gpu_queue->submit([=](sycl::handler& cgh){ + const sycl::nd_range<2> nd_range = make_nd_range(height,width); + cgh.parallel_for(nd_range,[=](sycl::nd_item<2> item) { + const unsigned int x = item.get_global_id(1); if (x >= width ) return; + const unsigned int y = item.get_global_id(0); if (y >= height) return; + RayStats stats; + renderPixelStandard(ldata,x,y,pixels,width,height,time,camera,stats); + }); + }); + global_gpu_queue->wait_and_throw(); + + const auto t0 = event.template get_profiling_info(); + const auto t1 = event.template get_profiling_info(); + const double dt = (t1-t0)*1E-9; + ((ISPCCamera*)&camera)->render_time = dt; +#else + const uniform int numTilesX = (width +TILE_SIZE_X-1)/TILE_SIZE_X; + const uniform int numTilesY = (height+TILE_SIZE_Y-1)/TILE_SIZE_Y; + launch[numTilesX*numTilesY] renderTileTask(pixels,width,height,time,camera,numTilesX,numTilesY); sync; +#endif +} + +/* called by the C++ code to render */ +export void device_render (uniform int* uniform pixels, + const uniform unsigned int width, + const uniform unsigned int height, + const uniform float time, + const uniform ISPCCamera& camera) +{ +} + +/* called by the C++ code for cleanup */ +export void device_cleanup () +{ + TutorialData_Destructor(&data); +} diff --git a/tutorials/debug_device_memory/debug_device_memory_device.isph b/tutorials/debug_device_memory/debug_device_memory_device.isph new file mode 100644 index 0000000000..6e3b1ef3b1 --- /dev/null +++ b/tutorials/debug_device_memory/debug_device_memory_device.isph @@ -0,0 +1,25 @@ +// Copyright 2009-2021 Intel Corporation +// SPDX-License-Identifier: Apache-2.0 + +#include "../common/tutorial/tutorial_device.isph" + +struct TutorialData +{ + RTCScene g_scene; + uniform Vec3f* uniform face_colors; + uniform Vec3fa* uniform vertex_colors; +}; + +inline void TutorialData_Constructor(uniform TutorialData* uniform This) +{ + This->g_scene = NULL; + This->face_colors = NULL; + This->vertex_colors = NULL; +} + +inline void TutorialData_Destructor(uniform TutorialData* uniform This) +{ + rtcReleaseScene (This->g_scene); This->g_scene = NULL; + delete[] This->face_colors; This->face_colors = NULL; + delete[] This->vertex_colors; This->vertex_colors = NULL; +} diff --git a/tutorials/point_geometry/point_geometry_device.cpp b/tutorials/point_geometry/point_geometry_device.cpp index 906fd2ba0f..f207617749 100644 --- a/tutorials/point_geometry/point_geometry_device.cpp +++ b/tutorials/point_geometry/point_geometry_device.cpp @@ -27,9 +27,7 @@ void addPoints (RTCScene scene, RTCGeometryType gtype, const Vec3fa& pos) #define NORMAL RandomSampler_get1D(rng) * 2.f - 1.f RTCGeometry geom = rtcNewGeometry (g_device, gtype); - rtcSetGeometryTimeStepCount(geom, 2); Vec4f* point_vertices = (Vec4f*)rtcSetNewGeometryBuffer(geom,RTC_BUFFER_TYPE_VERTEX, 0, RTC_FORMAT_FLOAT4, sizeof(Vec4f), NUM_POINTS); - Vec4f* point_vertices_1 = (Vec4f*)rtcSetNewGeometryBuffer(geom,RTC_BUFFER_TYPE_VERTEX, 1, RTC_FORMAT_FLOAT4, sizeof(Vec4f), NUM_POINTS); for (int i = 0; i < NUM_POINTS; i++) { @@ -38,7 +36,6 @@ void addPoints (RTCScene scene, RTCGeometryType gtype, const Vec3fa& pos) const float vz = COORD; const float vr = RADIUS; point_vertices[i] = Vec4f(pos.x,pos.y,pos.z,0.0f) + Vec4f(vx, vy, vz, vr); - point_vertices_1[i] = Vec4f(pos.x,pos.y,pos.z,0.0f) + Vec4f(vx, vy, vz, vr) - Vec4f(0.f, 1.f, 0.f, 0.f); const float cr = COLOR; const float cg = COLOR; const float cb = COLOR; @@ -47,14 +44,12 @@ void addPoints (RTCScene scene, RTCGeometryType gtype, const Vec3fa& pos) if (gtype == RTC_GEOMETRY_TYPE_ORIENTED_DISC_POINT) { Vec3fa* point_normals = (Vec3fa*)rtcSetNewGeometryBuffer(geom,RTC_BUFFER_TYPE_NORMAL, 0, RTC_FORMAT_FLOAT3, sizeof(Vec3fa), NUM_POINTS); - Vec3fa* point_normals_1 = (Vec3fa*)rtcSetNewGeometryBuffer(geom,RTC_BUFFER_TYPE_NORMAL, 1, RTC_FORMAT_FLOAT3, sizeof(Vec3fa), NUM_POINTS); for (int i = 0; i < NUM_POINTS; i++) { const float nx = NORMAL; const float ny = NORMAL; const float nz = NORMAL; point_normals[i] = Vec3fa(nx,ny,nz); point_normals[i] = normalize(point_normals[i]); - point_normals_1[i] = point_normals[i]; } } @@ -106,19 +101,6 @@ extern "C" void device_init (char* cfg) rtcCommitScene (g_scene); } -static inline uint32_t doodle(uint32_t x) -{ - x ^= x << 13; - x ^= x >> 17; - x ^= x << 5; - return x; -} - -static inline float doodlef(uint32_t x) -{ - return ((float)doodle(x)) / (float)(uint32_t(-1)); -} - /* task that renders a single screen tile */ void renderPixelStandard(const TutorialData& data, int x, int y, @@ -129,10 +111,7 @@ void renderPixelStandard(const TutorialData& data, const ISPCCamera& camera, RayStats& stats) { /* initialize ray */ - uint32_t state = doodle(x + y * width); - state = doodle(state); - float t = doodlef(state); - Ray ray(Vec3fa(camera.xfm.p), Vec3fa(normalize(x*camera.xfm.l.vx + y*camera.xfm.l.vy + camera.xfm.l.vz)), 0.0f, inf, t); + Ray ray(Vec3fa(camera.xfm.p), Vec3fa(normalize(x*camera.xfm.l.vx + y*camera.xfm.l.vy + camera.xfm.l.vz)), 0.0f, inf); /* intersect ray with scene */ RTCIntersectArguments iargs; diff --git a/tutorials/triangle_geometry/triangle_geometry_device.cpp b/tutorials/triangle_geometry/triangle_geometry_device.cpp index 7784a83837..f7b257ff12 100644 --- a/tutorials/triangle_geometry/triangle_geometry_device.cpp +++ b/tutorials/triangle_geometry/triangle_geometry_device.cpp @@ -3,21 +3,13 @@ #include "triangle_geometry_device.h" -#include - namespace embree { -AffineSpace3fa axfm0; -AffineSpace3fa axfm1; - /* all features required by this tutorial */ #define FEATURE_MASK \ - RTC_FEATURE_FLAG_TRIANGLE | \ - RTC_FEATURE_FLAG_INSTANCE | \ - RTC_FEATURE_FLAG_MOTION_BLUR + RTC_FEATURE_FLAG_TRIANGLE RTCScene g_scene = nullptr; -RTCScene g_scene1 = nullptr; TutorialData data; /* adds a cube to the scene */ @@ -25,60 +17,25 @@ unsigned int addCube (RTCScene scene_i) { /* create a triangulated cube with 12 triangles and 8 vertices */ RTCGeometry mesh = rtcNewGeometry(g_device, RTC_GEOMETRY_TYPE_TRIANGLE); - - rtcSetGeometryTimeStepCount(mesh, 2); /* create face and vertex color arrays */ data.face_colors = (Vec3fa*) alignedUSMMalloc((12)*sizeof(Vec3fa),16); data.vertex_colors = (Vec3fa*) alignedUSMMalloc((8)*sizeof(Vec3fa),16); /* set vertices and vertex colors */ - Vertex* vertices_device; - Vertex* vertices1_device; - Vertex* vertices = (Vertex*) rtcSetNewGeometryBufferXPU(mesh,RTC_BUFFER_TYPE_VERTEX,0,RTC_FORMAT_FLOAT3,sizeof(Vertex),8,(void**)&vertices_device); - Vertex* vertices1 = (Vertex*) rtcSetNewGeometryBufferXPU(mesh,RTC_BUFFER_TYPE_VERTEX,1,RTC_FORMAT_FLOAT3,sizeof(Vertex),8,(void**)&vertices1_device); - -#if defined(EMBREE_SYCL_TUTORIAL) - sycl::usm::alloc allocType; - allocType = sycl::get_pointer_type(vertices, global_gpu_queue->get_context()); - if (allocType != sycl::usm::alloc::host) - std::cout << "cube vertices have wrong alloc type!" <get_context()); - if (allocType != sycl::usm::alloc::device) - std::cout << "cube vertices have wrong alloc type!" <memcpy(vertices_device, vertices, 8 * sizeof(Vertex)); - global_gpu_queue->memcpy(vertices1_device, vertices1, 8 * sizeof(Vertex)); - global_gpu_queue->memcpy(triangles_device, triangles, 12 * sizeof(Triangle)); - global_gpu_queue->wait_and_throw(); -#endif - - rtcCommitGeometry(mesh); - unsigned int geomID = rtcAttachGeometry(scene_i,mesh); - rtcReleaseGeometry(mesh); - return geomID; -} - -unsigned int addCubeShared (RTCScene scene_i) -{ - /* create a triangulated cube with 12 triangles and 8 vertices */ - RTCGeometry mesh = rtcNewGeometry(g_device, RTC_GEOMETRY_TYPE_TRIANGLE); - rtcSetGeometryTimeStepCount(mesh, 2); - - /* create face and vertex color arrays */ - data.face_colors = (Vec3fa*) alignedUSMMalloc((12)*sizeof(Vec3fa),16); - data.vertex_colors = (Vec3fa*) alignedUSMMalloc((8)*sizeof(Vec3fa),16); - - Vertex* vertices = (Vertex*) alignedMalloc(8*sizeof(Vertex), 16); - Vertex* vertices1 = (Vertex*) alignedMalloc(8*sizeof(Vertex), 16); - Triangle* triangles = (Triangle*) alignedMalloc(12*sizeof(Triangle), 16); - -#if defined(EMBREE_SYCL_TUTORIAL) - Vertex* vertices_device = sycl::aligned_alloc_device(16, 8, *global_gpu_queue); - Vertex* vertices1_device = sycl::aligned_alloc_device(16, 8, *global_gpu_queue); - Triangle* triangles_device = sycl::aligned_alloc_device(16, 12, *global_gpu_queue); -#else - Vertex* vertices_device = nullptr; - Vertex* vertices1_device = nullptr; - Triangle* triangles_device = nullptr; -#endif - - rtcSetSharedGeometryBufferXPU(mesh,RTC_BUFFER_TYPE_VERTEX,0,RTC_FORMAT_FLOAT3,vertices,vertices_device, 0, sizeof(Vertex),8); - rtcSetSharedGeometryBufferXPU(mesh,RTC_BUFFER_TYPE_VERTEX,1,RTC_FORMAT_FLOAT3,vertices1,vertices1_device, 0, sizeof(Vertex),8); - - data.vertex_colors[0] = Vec3fa(0,0,0); - data.vertex_colors[1] = Vec3fa(0,0,1); - data.vertex_colors[2] = Vec3fa(0,1,0); - data.vertex_colors[3] = Vec3fa(0,1,1); - data.vertex_colors[4] = Vec3fa(1,0,0); - data.vertex_colors[5] = Vec3fa(1,0,1); - data.vertex_colors[6] = Vec3fa(1,1,0); - data.vertex_colors[7] = Vec3fa(1,1,1); - - vertices[0].x = -1; vertices[0].y = -1; vertices[0].z = -1; - vertices[1].x = -1; vertices[1].y = -1; vertices[1].z = +1; - vertices[2].x = -1; vertices[2].y = +1; vertices[2].z = -1; - vertices[3].x = -1; vertices[3].y = +1; vertices[3].z = +1; - vertices[4].x = +1; vertices[4].y = -1; vertices[4].z = -1; - vertices[5].x = +1; vertices[5].y = -1; vertices[5].z = +1; - vertices[6].x = +1; vertices[6].y = +1; vertices[6].z = -1; - vertices[7].x = +1; vertices[7].y = +1; vertices[7].z = +1; - - vertices1[0].x = -1 + 0.5f; vertices1[0].y = -1 + 0.5f; vertices1[0].z = -1 + 0.5f; - vertices1[1].x = -1 + 0.5f; vertices1[1].y = -1 + 0.5f; vertices1[1].z = +1 + 0.5f; - vertices1[2].x = -1 + 0.5f; vertices1[2].y = +1 + 0.5f; vertices1[2].z = -1 + 0.5f; - vertices1[3].x = -1 + 0.5f; vertices1[3].y = +1 + 0.5f; vertices1[3].z = +1 + 0.5f; - vertices1[4].x = +1 + 0.5f; vertices1[4].y = -1 + 0.5f; vertices1[4].z = -1 + 0.5f; - vertices1[5].x = +1 + 0.5f; vertices1[5].y = -1 + 0.5f; vertices1[5].z = +1 + 0.5f; - vertices1[6].x = +1 + 0.5f; vertices1[6].y = +1 + 0.5f; vertices1[6].z = -1 + 0.5f; - vertices1[7].x = +1 + 0.5f; vertices1[7].y = +1 + 0.5f; vertices1[7].z = +1 + 0.5f; - - /* set triangles and face colors */ - int tri = 0; - rtcSetSharedGeometryBufferXPU(mesh,RTC_BUFFER_TYPE_INDEX,0,RTC_FORMAT_UINT3,triangles,triangles_device,0,sizeof(Triangle),12); - - // left side - data.face_colors[tri] = Vec3fa(1,0,0); triangles[tri].v0 = 0; triangles[tri].v1 = 1; triangles[tri].v2 = 2; tri++; - data.face_colors[tri] = Vec3fa(1,0,0); triangles[tri].v0 = 1; triangles[tri].v1 = 3; triangles[tri].v2 = 2; tri++; - - // right side - data.face_colors[tri] = Vec3fa(0,1,0); triangles[tri].v0 = 4; triangles[tri].v1 = 6; triangles[tri].v2 = 5; tri++; - data.face_colors[tri] = Vec3fa(0,1,0); triangles[tri].v0 = 5; triangles[tri].v1 = 6; triangles[tri].v2 = 7; tri++; - - // bottom side - data.face_colors[tri] = Vec3fa(0.5f); triangles[tri].v0 = 0; triangles[tri].v1 = 4; triangles[tri].v2 = 1; tri++; - data.face_colors[tri] = Vec3fa(0.5f); triangles[tri].v0 = 1; triangles[tri].v1 = 4; triangles[tri].v2 = 5; tri++; - - // top side - data.face_colors[tri] = Vec3fa(1.0f); triangles[tri].v0 = 2; triangles[tri].v1 = 3; triangles[tri].v2 = 6; tri++; - data.face_colors[tri] = Vec3fa(1.0f); triangles[tri].v0 = 3; triangles[tri].v1 = 7; triangles[tri].v2 = 6; tri++; - - // front side - data.face_colors[tri] = Vec3fa(0,0,1); triangles[tri].v0 = 0; triangles[tri].v1 = 2; triangles[tri].v2 = 4; tri++; - data.face_colors[tri] = Vec3fa(0,0,1); triangles[tri].v0 = 2; triangles[tri].v1 = 6; triangles[tri].v2 = 4; tri++; - - // back side - data.face_colors[tri] = Vec3fa(1,1,0); triangles[tri].v0 = 1; triangles[tri].v1 = 5; triangles[tri].v2 = 3; tri++; - data.face_colors[tri] = Vec3fa(1,1,0); triangles[tri].v0 = 3; triangles[tri].v1 = 5; triangles[tri].v2 = 7; tri++; - - rtcSetGeometryVertexAttributeCount(mesh,1); - rtcSetSharedGeometryBuffer(mesh,RTC_BUFFER_TYPE_VERTEX_ATTRIBUTE,0,RTC_FORMAT_FLOAT3,data.vertex_colors,0,sizeof(Vec3fa),8); - -#if defined(EMBREE_SYCL_TUTORIAL) - global_gpu_queue->memcpy(vertices_device, vertices, 8 * sizeof(Vertex)); - global_gpu_queue->memcpy(vertices1_device, vertices1, 8 * sizeof(Vertex)); - global_gpu_queue->memcpy(triangles_device, triangles, 12 * sizeof(Triangle)); -#endif - rtcCommitGeometry(mesh); unsigned int geomID = rtcAttachGeometry(scene_i,mesh); rtcReleaseGeometry(mesh); @@ -235,7 +87,7 @@ unsigned int addGroundPlane (RTCScene scene_i) Triangle* triangles = (Triangle*) rtcSetNewGeometryBuffer(mesh,RTC_BUFFER_TYPE_INDEX,0,RTC_FORMAT_UINT3,sizeof(Triangle),2); triangles[0].v0 = 0; triangles[0].v1 = 1; triangles[0].v2 = 2; triangles[1].v0 = 1; triangles[1].v1 = 3; triangles[1].v2 = 2; - + rtcCommitGeometry(mesh); unsigned int geomID = rtcAttachGeometry(scene_i,mesh); rtcReleaseGeometry(mesh); @@ -248,51 +100,16 @@ extern "C" void device_init (char* cfg) /* create scene */ TutorialData_Constructor(&data); g_scene = data.g_scene = rtcNewScene(g_device); - g_scene1 = rtcNewScene(g_device); rtcSetSceneFlags(data.g_scene, RTC_SCENE_FLAG_PREFETCH_USM_SHARED_ON_GPU); /* add cube */ - addCube(g_scene1); - - RTCGeometry inst = rtcNewGeometry(g_device, RTC_GEOMETRY_TYPE_INSTANCE); - - rtcSetGeometryTimeStepCount(inst, 2); - rtcSetGeometryInstancedScene(inst, g_scene1); - - LinearSpace3fa xfm = one; - axfm0 = AffineSpace3fa(xfm,Vec3fa(0.f, 0.f, 0.f)); - axfm1 = AffineSpace3fa(xfm,Vec3fa(3.f, 0.f, 0.f)); - rtcSetGeometryTransform(inst,0,RTC_FORMAT_FLOAT4X4_COLUMN_MAJOR,(float*)&(axfm0.l.vx.x)); - rtcSetGeometryTransform(inst,1,RTC_FORMAT_FLOAT4X4_COLUMN_MAJOR,(float*)&(axfm1.l.vx.x)); - - rtcAttachGeometry(data.g_scene,inst); - rtcReleaseGeometry(inst); - rtcCommitGeometry(inst); + addCube(data.g_scene); /* add ground plane */ addGroundPlane(data.g_scene); /* commit changes to scene */ -#if defined(EMBREE_SYCL_SUPPORT) && defined(SYCL_LANGUAGE_VERSION) - rtcCommitSceneWithQueue (g_scene1, *global_gpu_queue); - rtcCommitSceneWithQueue (data.g_scene, *global_gpu_queue); -#else - rtcCommitScene (g_scene1); rtcCommitScene (data.g_scene); -#endif -} - -static inline uint32_t doodle(uint32_t x) -{ - x ^= x << 13; - x ^= x >> 17; - x ^= x << 5; - return x; -} - -static inline float doodlef(uint32_t x) -{ - return ((float)doodle(x)) / (float)(uint32_t(-1)); } /* task that renders a single screen tile */ @@ -305,10 +122,7 @@ void renderPixelStandard(const TutorialData& data, const ISPCCamera& camera, RayStats& stats) { /* initialize ray */ - uint32_t state = doodle(x + y * width); - state = doodle(state); - float t = doodlef(state); - Ray ray(Vec3fa(camera.xfm.p), Vec3fa(normalize(x*camera.xfm.l.vx + y*camera.xfm.l.vy + camera.xfm.l.vz)), 0.0f, inf, t); + Ray ray(Vec3fa(camera.xfm.p), Vec3fa(normalize(x*camera.xfm.l.vx + y*camera.xfm.l.vy + camera.xfm.l.vz)), 0.0f, inf); /* intersect ray with scene */ RTCIntersectArguments iargs; @@ -319,15 +133,14 @@ void renderPixelStandard(const TutorialData& data, /* shade pixels */ Vec3fa color = Vec3fa(0.0f); - if (ray.geomID != RTC_INVALID_GEOMETRY_ID || ray.instID[0] != RTC_INVALID_GEOMETRY_ID) + if (ray.geomID != RTC_INVALID_GEOMETRY_ID) { -#if 1 Vec3fa diffuse = data.face_colors[ray.primID]; color = color + diffuse*0.5f; Vec3fa lightDir = normalize(Vec3fa(-1,-1,-1)); /* initialize shadow ray */ - Ray shadow(ray.org + ray.tfar*ray.dir, neg(lightDir), 0.001f, inf, 1.f - t); + Ray shadow(ray.org + ray.tfar*ray.dir, neg(lightDir), 0.001f, inf, 0.0f); /* trace shadow ray */ RTCOccludedArguments sargs; @@ -339,38 +152,6 @@ void renderPixelStandard(const TutorialData& data, /* add light contribution */ if (shadow.tfar >= 0.0f) color = color + diffuse*clamp(-dot(lightDir,normalize(ray.Ng)),0.0f,1.0f); -#else - -#if 0 - if (ray.geomID == 0) - color = Vec3fa(0.f, 0.f, 1.f); - else if (ray.geomID == 1) - color = Vec3fa(0.f, 1.f, 0.f); - else if (ray.geomID == 2) - color = Vec3fa(1.f, 0.f, 0.f); - else - color = Vec3fa(1.f); -#endif - -#if 0 - if (ray.primID == 0) - color = Vec3fa(0.f, 0.f, 1.f); - else if (ray.primID == 1) - color = Vec3fa(0.f, 1.f, 0.f); - else if (ray.primID == 2) - color = Vec3fa(1.f, 0.f, 0.f); - else - color = Vec3fa(1.f); -#endif - -#if 0 - color = Vec3fa(ray.u, ray.v, 0.f); -#endif - - color = Vec3fa(0.f, 1.f, 0.f); - -#endif - } /* write color to framebuffer */