SYCL fix:

remembering V coordinate of custom primitives in variable

Author: stefanatwork

kernels/common/geometry.h

@@ -220,12 +220,6 @@ namespace embree
                    MTY_FLAT_HERMITE_CURVE | MTY_ROUND_HERMITE_CURVE | MTY_ORIENTED_HERMITE_CURVE |
                    MTY_FLAT_CATMULL_ROM_CURVE | MTY_ROUND_CATMULL_ROM_CURVE | MTY_ORIENTED_CATMULL_ROM_CURVE,
 
-      MTY_FLAT_CURVES = MTY_FLAT_LINEAR_CURVE | MTY_ORIENTED_LINEAR_CURVE |
-                        MTY_FLAT_BEZIER_CURVE | MTY_ORIENTED_BEZIER_CURVE |
-                        MTY_FLAT_BSPLINE_CURVE | MTY_ORIENTED_BSPLINE_CURVE |
-                        MTY_FLAT_HERMITE_CURVE | MTY_ORIENTED_HERMITE_CURVE |
-                        MTY_FLAT_CATMULL_ROM_CURVE | MTY_ORIENTED_CATMULL_ROM_CURVE,
-
       MTY_SPHERE_POINT = 1ul << GTY_SPHERE_POINT,
       MTY_DISC_POINT = 1ul << GTY_DISC_POINT,
       MTY_ORIENTED_DISC_POINT = 1ul << GTY_ORIENTED_DISC_POINT,

kernels/sycl/rthwif_embree.cpp

@@ -584,21 +584,19 @@ __forceinline bool invokeTriangleIntersectionFilter(intel_ray_query_t& query, Ge
   return ishit;
 }
 
-__forceinline bool commit_potential_hit(intel_ray_query_t& query, RayHit& ray, bool scale_v) {
-    /* Xe3 may store UVs in UNORM format, so we can't use negative V coordinates here
-      Leaving the CPU code path as is, GPUs will apply the scale later. */
-  float v = scale_v ? madd(ray.v, 0.5f, 0.5f) : ray.v;
-  intel_ray_query_commit_potential_hit_override (query, ray.tfar, float2(ray.u, v));
+__forceinline bool commit_potential_hit(intel_ray_query_t& query, RayHit& ray, float &ray_v) {
+  ray_v = ray.v;
+  intel_ray_query_commit_potential_hit_override (query, ray.tfar, float2(ray.u, ray.v));
   return false;
 }
 
-__forceinline bool commit_potential_hit(intel_ray_query_t& query, Ray& ray, bool) {
+__forceinline bool commit_potential_hit(intel_ray_query_t& query, Ray& ray, float &) {
   intel_ray_query_commit_potential_hit_override (query, ray.tfar, float2(0.0f, 0.0f));
   return true;
 }
 
 template<typename Ray>
-__forceinline void trav_loop(intel_ray_query_t& query, Ray& ray, Scene* scene, sycl::private_ptr<RayQueryContext> context, const RTCFeatureFlags feature_mask)
+__forceinline void trav_loop(intel_ray_query_t& query, Ray& ray, Scene* scene, sycl::private_ptr<RayQueryContext> context, const RTCFeatureFlags feature_mask, float &ray_v)
 {
   Scenes scenes(scene);
   while (!intel_is_traversal_done(query))
@@ -647,8 +645,7 @@ __forceinline void trav_loop(intel_ray_query_t& query, Ray& ray, Scene* scene, s
       if (candidate == intel_candidate_type_procedural)
       {
         if (intersect_primitive(query,ray,scenes,geom,context,geomID,primID,feature_mask)) {
-          bool scale_v = geom->getTypeMask() & Geometry::MTY_FLAT_CURVES;
-          if (commit_potential_hit (query, ray, scale_v))
+          if (commit_potential_hit (query, ray, ray_v))
             break; // shadow rays break at first hit
         }
       }
@@ -726,8 +723,9 @@ SYCL_EXTERNAL __attribute__((always_inline)) void rtcIntersectRTHW(sycl::global_
   intel_ray_query_start_traversal(query);
   intel_ray_query_sync(query);
 
+  float hit_v = 0.0f;
   if (args->feature_mask & TRAV_LOOP_FEATURES) {
-    trav_loop(query,ray,scene,&context,args->feature_mask);
+    trav_loop(query,ray,scene,&context,args->feature_mask, hit_v);
   }
 
   bool valid = intel_has_committed_hit(query);
@@ -739,27 +737,22 @@ SYCL_EXTERNAL __attribute__((always_inline)) void rtcIntersectRTHW(sycl::global_
     unsigned int geomID = intel_get_hit_geometry_id(query, intel_hit_type_committed_hit);
 
     unsigned int primID = ray.primID;
-    if (intel_get_hit_candidate(query, intel_hit_type_committed_hit) == intel_candidate_type_triangle)
+    if (intel_get_hit_candidate(query, intel_hit_type_committed_hit) == intel_candidate_type_triangle) {
       primID = intel_get_hit_triangle_primitive_id(query, intel_hit_type_committed_hit);
+    }
+    else {
+      /* Xe3 hardware stores UV as UNORM. Since ribbons use a range of -1.0 to 1.0
+        the V coordinate needs to be passed in a separate variable. */
+      if ((args->feature_mask & RTC_FEATURE_FLAG_FLAT_CURVES) ||
+          (args->feature_mask & RTC_FEATURE_FLAG_NORMAL_ORIENTED_CURVES))
+        uv.y() = hit_v;
+    }
 
     rayhit_i->ray.tfar = t;
     rayhit_i->hit.geomID = geomID;
     rayhit_i->hit.primID = primID;
     rayhit_i->hit.u = uv.x();
     rayhit_i->hit.v = uv.y();
-
-    if (RTC_FEATURE_FLAG_CURVES & TRAV_LOOP_FEATURES) {
-      if (intel_get_hit_candidate(query, intel_hit_type_committed_hit) == intel_candidate_type_procedural)
-      {
-        Geometry* geom = scene->get(geomID);
-        if (geom->getTypeMask() & Geometry::MTY_CURVES &&
-           (geom->getCurveType() == Geometry::GTY_SUBTYPE_FLAT_CURVE || geom->getCurveType() == Geometry::GTY_SUBTYPE_ORIENTED_CURVE))
-        {
-          /* for curves we need to convert the v coordinates from [0,1] to [-1,1] */
-          rayhit_i->hit.v = rayhit_i->hit.v * 2.0f - 1.0f;
-        }
-      }
-    }
 
 #if RTC_MAX_INSTANCE_LEVEL_COUNT > 1
     for (uint32_t l=0; l<RTC_MAX_INSTANCE_LEVEL_COUNT; l++) {
@@ -839,8 +832,9 @@ SYCL_EXTERNAL __attribute__((always_inline)) void rtcOccludedRTHW(sycl::global_p
   intel_ray_query_start_traversal(query);
   intel_ray_query_sync(query);
 
+  float v_unused MAYBE_UNUSED = 0.0f;
   if (args->feature_mask & TRAV_LOOP_FEATURES) {
-    trav_loop(query,ray,scene,&context,args->feature_mask);
+    trav_loop(query,ray,scene,&context,args->feature_mask, v_unused);
   }
 
   if (intel_has_committed_hit(query))