Incorrect SurfaceInteraction for quadric surfaces

Hello,

I noticed an image difference when rendering quadric surfaces between CPU and GPU. I'll attach the scene and command line options at the end of this message.

Look the following two images: Here, there is a single cylinder illuminated by a distant light, rendered by CPU (1st) or GPU (2nd) GPU one looks obviously weird. The point is that rotation (along cylinder's axis, so essentially no effect) applied to the cylinder (look the scene description below). If I comment out the rotation, I get the correct image:

I read the code and it looks that OptiX code for quadric surfaces is weird. Specifically, ray direction passed to an intersection routine for a quadric surface is in local coordinate system in CPU code: https://github.com/mmp/pbrt-v4/blob/05ff05e1ded8299b1de0eb8ee6c11f192d0a64dd/src/pbrt/cpu/primitive.cpp#L115 but it looks to be in the world (render) space in GPU code: https://github.com/mmp/pbrt-v4/blob/05ff05e1ded8299b1de0eb8ee6c11f192d0a64dd/src/pbrt/gpu/optix.cu#L305

I could fix this particular issue by using object-space ray direction instead like the following:

+template <bool fromIS>
 static __device__ inline SurfaceInteraction getQuadricIntersection(
     const QuadricIntersection &si) {
     QuadricRecord &rec = *((QuadricRecord *)optixGetSbtDataPointer());

-    float3 rd = optixGetWorldRayDirection();
-    Vector3f wo = -Vector3f(rd.x, rd.y, rd.z);
+    Vector3f wo;
+    if constexpr (fromIS) {
+        float3 rd = optixGetObjectRayDirection();
+        wo = -Vector3f(rd.x, rd.y, rd.z);
+    }
+    else { // stored as payload in IS
+        wo = -Vector3f(BitsToFloat(optixGetAttribute_4()),
+                       BitsToFloat(optixGetAttribute_5()),
+                       BitsToFloat(optixGetAttribute_6()));
+    }

but I'm not sure if this is the complete/ideal fix. And possibly other non-triangle codes are also wrong (I didn't check though).

Thanks,

The scene I used

Film "gbuffer"
    "integer xresolution" 256 "integer yresolution" 256
    "string filename" "output.exr"

LookAt 3 0 2   0 0 0.5   0 0 1
Camera "perspective" "float fov" 45

Option "wavefront" true
Sampler "halton" "integer pixelsamples" 1
Integrator "path"

WorldBegin

AttributeBegin
LightSource "distant" "point3 from" [1 0 1] "point3 to" [0 0 0] "float scale" 5
AttributeEnd

AttributeBegin

ObjectBegin "wall"
Material "diffuse" "rgb reflectance" [1 1 1]
Shape "cylinder" "float zmin" 0 "float zmax" 1 "float radius" 1
ObjectEnd

AttributeBegin
Rotate 90 0 0 1 # This rotation is the problem for GPU
ObjectInstance "wall"
AttributeEnd

AttributeEnd

The command line options I used: [--gpu] --spp 64 test-v4.pbrt

mmp / pbrt-v4

Incorrect SurfaceInteraction for quadric surfaces #372