Merge pull request #966 from Autodesk/krickw/MAYA-106937/prototype_co…

…mpute_normals_and_OSD_on_the_gpu

MAYA-106937 Prototype support for GPU Normals and OSD using glsl.

cmake/modules/FindMaya.cmake

@@ -207,7 +207,8 @@ foreach(MAYA_LIB
     IMFbase
     tbb
     cg
-    cgGL)
+    cgGL
+    clew)
 
     find_library(MAYA_${MAYA_LIB}_LIBRARY
             ${MAYA_LIB}

lib/mayaUsd/render/CMakeLists.txt

@@ -1,4 +1,5 @@
 add_subdirectory(px_vp20)
 add_subdirectory(pxrUsdMayaGL)
+add_subDirectory(vp2ComputeShaders)
 add_subdirectory(vp2RenderDelegate)
 add_subdirectory(vp2ShaderFragments)

lib/mayaUsd/render/vp2ComputeShaders/CMakeLists.txt

@@ -0,0 +1,13 @@
+# -----------------------------------------------------------------------------
+# install
+# -----------------------------------------------------------------------------
+list(APPEND MAYAUSD_COMPUTESHADERS
+    computeNormals.glsl
+    computeNormals.cl
+    plugInfo.json
+)
+
+install(FILES ${MAYAUSD_COMPUTESHADERS}
+    DESTINATION ${CMAKE_INSTALL_PREFIX}/lib/usd/mayaUSD_ComputeShaders/resources
+)
+

lib/mayaUsd/render/vp2ComputeShaders/computeNormals.cl

@@ -0,0 +1,34 @@
+__kernel void computeNormals(
+    __global const float* positions,
+    const unsigned int    vertexCount, // number of posisions and normals
+                                       /*  Adjacency buffer is two distinct parts.
+                                           First, two ints per vertex the offset and the valence. The valence is the number of adjacent
+                                          vertices. The offset is the offset into the adjacency buffer to find the vertex ids of the
+                                          adjacent vertices.                                    Next, a list of vertex ids of the
+                                          adjacent vertices for each vertex, found                                    using the information from the first part of the
+                                          buffer.
+                                       */
+    __global const int* adjacency,
+    __global float*     normals)
+{
+    unsigned int vertexId = get_global_id(0);
+    if (vertexId >= vertexCount)
+        return;
+
+    unsigned int  offsetIdx = vertexId * 2;
+    int           offset = adjacency[offsetIdx];
+    int           valence = adjacency[offsetIdx + 1];
+    __global int* currAdj = &adjacency[offset];
+
+    const float3 currVertex = vload3(vertexId, positions);
+    float3       accumulatedNormal = (float3)(0.0f, 0.0f, 0.0f);
+
+    for (int neighbour = 0; neighbour < valence; neighbour++) {
+        float3 prevVertex = vload3(*currAdj++, positions);
+        float3 nextVertex = vload3(*currAdj++, positions);
+        accumulatedNormal += cross(nextVertex - currVertex, prevVertex - currVertex);
+    }
+
+    accumulatedNormal = normalize(accumulatedNormal);
+    vstore3(accumulatedNormal, vertexId, normals);
+}

lib/mayaUsd/render/vp2ComputeShaders/computeNormals.glsl

@@ -0,0 +1,69 @@
+#version 430
+
+layout( std140, binding=0 ) uniform Values
+{
+	uint VertexCount;
+};
+
+// This is a float3 but for buffer layout to be correct use float
+layout( std430, binding=1 ) buffer Pos
+{
+	float Positions[ ];
+};
+
+layout( std430, binding=2 ) buffer Adj
+{
+	int Adjacency[ ];
+};
+
+layout( std430, binding=3 ) buffer RtoS
+{
+	int RenderingToScene[ ];
+};
+
+layout( std430, binding=4 ) buffer StoR
+{
+	int SceneToRendering[ ];
+};
+
+// This is a float3 but for buffer layout to be correct use float
+layout( std430, binding=5 ) buffer Norm
+{
+	float Normals[ ];
+};
+
+layout( local_size_x = 256, local_size_y = 1, local_size_z = 1) in;
+
+void main() {
+	uint renderingVertexId = gl_GlobalInvocationID.x;
+	uint renderingVertexOffset = renderingVertexId *3;
+
+	if (renderingVertexId < VertexCount)
+	{
+		uint sceneVertexId = RenderingToScene[renderingVertexId];
+
+		uint adjOffsetIdx = sceneVertexId*2;
+		int offset = Adjacency[adjOffsetIdx];
+		int valence = Adjacency[adjOffsetIdx+1];
+
+		vec3 currVertex = vec3(Positions[renderingVertexOffset], Positions[renderingVertexOffset+1], Positions[renderingVertexOffset+2]);
+		vec3 accumulatedNormal = vec3(0.0, 0.0, 0.0);
+
+		for (int neighbour=0; neighbour<valence; neighbour++)
+		{
+			int prevVertexOffset = SceneToRendering[Adjacency[offset++]] * 3;
+			vec3 prevVertex = vec3(Positions[prevVertexOffset], Positions[prevVertexOffset+1], Positions[prevVertexOffset+2]);
+			int nextVertexOffset = SceneToRendering[Adjacency[offset++]] * 3;
+			vec3 nextVertex = vec3(Positions[nextVertexOffset], Positions[nextVertexOffset+1], Positions[nextVertexOffset+2]);
+			accumulatedNormal += cross(nextVertex - currVertex, prevVertex - currVertex);
+		}
+
+		accumulatedNormal = normalize(accumulatedNormal);
+
+		Normals[renderingVertexOffset] = accumulatedNormal.x;
+		Normals[renderingVertexOffset+1] = accumulatedNormal.y;
+		Normals[renderingVertexOffset+2] = accumulatedNormal.z;
+	}
+}
+
+

lib/mayaUsd/render/vp2ComputeShaders/plugInfo.json

@@ -0,0 +1,11 @@
+{
+    "Plugins": [
+        {
+            "Name": "mayaUsd_ComputeShaders", 
+            "Info": {}, 
+            "ResourcePath": "../../usd/mayaUsd_ComputeShaders/resources", 
+            "Root": "..", 
+            "Type": "resource"
+        }
+    ]
+}

lib/mayaUsd/render/vp2RenderDelegate/CMakeLists.txt

@@ -10,6 +10,7 @@ target_sources(${PROJECT_NAME}
         instancer.cpp
         material.cpp
         mesh.cpp
+        meshViewportCompute.cpp
         proxyRenderDelegate.cpp
         render_delegate.cpp
         render_param.cpp