Utilizing GPU compute with Geometries¶
導入部¶
GPUコンピュートは、現在比較的容易にジオメトリデフォメーションの計算のために使用することができます。
InlineDrawingシステムは、既にGPU上にあるジオメトリのレンダリングをサポートするように更新されました。
オブジェクトはまだGPUでサポートされておらず、単純な型の構造体と配列だけはGPUコンピュートカーネルで使用することができます。
ポリゴンメッシュ¶
PolygonMesh は、GPU上へ移動しGPUコンピュートカーネルに渡すことができる PolygonMeshTopology 型の ‘topology’ と呼ばれるメンバーを所有しています。
PolygonMesh をGPUカーネルで使用する前に、そのトポロジーデータをGPUに移動する必要があります。ジオメトリが所有する各 :kl-ref:`GeometryAttribute’ もまた使用する前にGPUに個別に移動する必要があります。
GPUコンピュートカーネルはオブジェクト群をサポートしておらず、ポリゴンメッシュやライン群、ポイント群、または様々なアトリビュート型をGPUカーネルへ渡せないことを意味します。代わりに、これらのオブジェクトに含まれるデータはGPUカーネルに渡すことができます。
Each of the GeometryAttributes owns data that can be passed into GPU Compute kernels. Most attributes have a member called ‘values’ which is an array of the data type supported by the GeometryAttribute.
/*
** Example: Moving a polygonMesh to the GPU
*/
require Geometry;
operator meshDataOnGPU<<<index>>>(io PolygonMeshTopology meshData){
report(meshData);
}
operator entry() {
PolygonMesh mesh = PolygonMesh();
Scalar length = 12.0;
Scalar width = 4.0;
Integer lengthSections = 3;
Integer widthSections = 4;
mesh.addPlane(Xfo(), length, width, lengthSections, widthSections, true, true);
mesh.convertToGPU();
meshDataOnGPU<<<1@true>>>(mesh.topology);
mesh.convertToCPU();
report(mesh.getDesc(true, true));
}
/*
** Output:
{polyData:{data:{data:[0,4,8,56,105,154,12,200,9,153,250,20,296,201,249,346,28,392,297,345,440,36,155,104,489,538,44,251,152,537,586,52,347,248,585,634,60,441,344,633,680,68,539,488,728,777,76,587,536,776,825,84,635,584,824,873,92,681,632,872,920],freeData:{data:[0,0],size:61},freeDataCount:0,freeBlocsListUpperPow2:[]},startIters:[1,6,11,16,21,26,31,36,41,46,51,56],freeIters:{data:[0],size:12},freeItersCount:0},pointData:{data:{data:[0,2147483658,8,49,5,4294967295,4294967295,2147483649,9,5,4294967295,4294967295,4294967295,2147483690,169,10,5,4294967295,4294967295,52,209,50,11,168,4294967295,2147483666,48,89,5,4294967295,4294967295,60,249,90,51,208,4294967295,2147483674,88,129,5,4294967295,4294967295,68,289,130,91,248,4294967295,2147483681,128,5,4294967295,4294967295,4294967295,2147483722,131,288,5,4294967295,4294967295,2147483730,329,170,5,4294967295,4294967295,92,369,210,171,328,4294967295,100,409,250,211,368,4294967295,108,449,290,251,408,4294967295,2147483762,291,448,5,4294967295,4294967295,2147483769,330,5,4294967295,4294967295,4294967295,2147483778,370,331,5,4294967295,4294967295,2147483786,410,371,5,4294967295,4294967295,2147483794,450,411,5,4294967295,4294967295,2147483801,451,5,4294967295,4294967295,4294967295],freeData:{data:[1627790432,410060304,3322953990,29759585],size:121},freeDataCount:38,freeBlocsListUpperPow2:[]},startIters:[7,1,25,37,49,13,19,31,43,55,61,67,73,79,85,91,97,103,109,115],freeIters:{data:[0],size:20},freeItersCount:0},unsharedAttrToPointSlidingArray:[],unsharedAttrToPointSlidingArrayBegin:0}
Mesh: pointCount: 20 polygonCount: 12 nbAttributeVectors: 20
Points (adjacent polygons as 'polygon.neighborPolygonIndex',
borders as '|', closed wing starts as '<<',
unshared attribute as value@attributeIndex ):
0: 1 polygons: |0.1
Attr positions:{x:-2.0,y:+0.0,z:-6.0}
Attr normals:{x:+0.0,y:+1.0,z:+0.0}
Attr uvs0:{x:+0.0,y:+0.0}
1: 2 polygons: |0.0, 1.1
Attr positions:{x:-1.0,y:+0.0,z:-6.0}
Attr normals:{x:+0.0,y:+1.0,z:+0.0}
Attr uvs0:{x:+0.25,y:+0.0}
2: 2 polygons: |1.0, 2.1
Attr positions:{x:+0.0,y:+0.0,z:-6.0}
Attr normals:{x:+0.0,y:+1.0,z:+0.0}
Attr uvs0:{x:+0.5,y:+0.0}
3: 2 polygons: |2.0, 3.1
Attr positions:{x:+1.0,y:+0.0,z:-6.0}
Attr normals:{x:+0.0,y:+1.0,z:+0.0}
Attr uvs0:{x:+0.75,y:+0.0}
4: 1 polygons: |3.0
Attr positions:{x:+2.0,y:+0.0,z:-6.0}
Attr normals:{x:+0.0,y:+1.0,z:+0.0}
Attr uvs0:{x:+1.0,y:+0.0}
5: 2 polygons: |4.1, 0.2
Attr positions:{x:-2.0,y:+0.0,z:-2.0}
Attr normals:{x:+0.0,y:+1.0,z:+0.0}
Attr uvs0:{x:+0.0,y:+0.333313}
6: 4 polygons: <<5.1, 1.2, 0.3, 4.0
Attr positions:{x:-1.0,y:+0.0,z:-2.0}
Attr normals:{x:+0.0,y:+1.0,z:+0.0}
Attr uvs0:{x:+0.25,y:+0.333313}
7: 4 polygons: <<6.1, 2.2, 1.3, 5.0
Attr positions:{x:+0.0,y:+0.0,z:-2.0}
Attr normals:{x:+0.0,y:+1.0,z:+0.0}
Attr uvs0:{x:+0.5,y:+0.333313}
8: 4 polygons: <<7.1, 3.2, 2.3, 6.0
Attr positions:{x:+1.0,y:+0.0,z:-2.0}
Attr normals:{x:+0.0,y:+1.0,z:+0.0}
Attr uvs0:{x:+0.75,y:+0.333313}
9: 2 polygons: |3.3, 7.0
Attr positions:{x:+2.0,y:+0.0,z:-2.0}
Attr normals:{x:+0.0,y:+1.0,z:+0.0}
Attr uvs0:{x:+1.0,y:+0.333313}
10: 2 polygons: |8.1, 4.2
Attr positions:{x:-2.0,y:+0.0,z:+2.0}
Attr normals:{x:+0.0,y:+1.0,z:+0.0}
Attr uvs0:{x:+0.0,y:+0.666626}
11: 4 polygons: <<9.1, 5.2, 4.3, 8.0
Attr positions:{x:-1.0,y:+0.0,z:+2.0}
Attr normals:{x:+0.0,y:+1.0,z:+0.0}
Attr uvs0:{x:+0.25,y:+0.666626}
12: 4 polygons: <<10.1, 6.2, 5.3, 9.0
Attr positions:{x:+0.0,y:+0.0,z:+2.0}
Attr normals:{x:+0.0,y:+1.0,z:+0.0}
Attr uvs0:{x:+0.5,y:+0.666626}
13: 4 polygons: <<11.1, 7.2, 6.3, 10.0
Attr positions:{x:+1.0,y:+0.0,z:+2.0}
Attr normals:{x:+0.0,y:+1.0,z:+0.0}
Attr uvs0:{x:+0.75,y:+0.666626}
14: 2 polygons: |7.3, 11.0
Attr positions:{x:+2.0,y:+0.0,z:+2.0}
Attr normals:{x:+0.0,y:+1.0,z:+0.0}
Attr uvs0:{x:+1.0,y:+0.666626}
15: 1 polygons: |8.2
Attr positions:{x:-2.0,y:+0.0,z:+6.0}
Attr normals:{x:+0.0,y:+1.0,z:+0.0}
Attr uvs0:{x:+0.0,y:+1.0}
16: 2 polygons: |9.2, 8.3
Attr positions:{x:-1.0,y:+0.0,z:+6.0}
Attr normals:{x:+0.0,y:+1.0,z:+0.0}
Attr uvs0:{x:+0.25,y:+1.0}
17: 2 polygons: |10.2, 9.3
Attr positions:{x:+0.0,y:+0.0,z:+6.0}
Attr normals:{x:+0.0,y:+1.0,z:+0.0}
Attr uvs0:{x:+0.5,y:+1.0}
18: 2 polygons: |11.2, 10.3
Attr positions:{x:+1.0,y:+0.0,z:+6.0}
Attr normals:{x:+0.0,y:+1.0,z:+0.0}
Attr uvs0:{x:+0.75,y:+1.0}
19: 1 polygons: |11.3
Attr positions:{x:+2.0,y:+0.0,z:+6.0}
Attr normals:{x:+0.0,y:+1.0,z:+0.0}
Attr uvs0:{x:+1.0,y:+1.0}
Polygons (connected points as 'point.polygonPointIndex', borders as '|'):
0: 4 points: 1.0 |, 0.0 |, 5.1, 6.2
1: 4 points: 2.0 |, 1.1, 6.1, 7.2
2: 4 points: 3.0 |, 2.1, 7.1, 8.2
3: 4 points: 4.0 |, 3.1, 8.1, 9.0 |
4: 4 points: 6.3, 5.0 |, 10.1, 11.2
5: 4 points: 7.3, 6.0, 11.1, 12.2
6: 4 points: 8.3, 7.0, 12.1, 13.2
7: 4 points: 9.1, 8.0, 13.1, 14.0 |
8: 4 points: 11.3, 10.0 |, 15.0 |, 16.1
9: 4 points: 12.3, 11.0, 16.0 |, 17.1
10: 4 points: 13.3, 12.0, 17.0 |, 18.1
11: 4 points: 14.1, 13.0, 18.0 |, 19.0 |
*/
データがGPU上に有れば、 CPUで普段よく使用されるものと同じカーネル群をGPU上でも実行できます。
/*
** Example: Performing a smooth mesh using the CPU and then again using the GPU
*/
require Geometry;
operator noiseOp<<<index>>>(io Vec3 positions[], Scalar height){
positions[index].y = mathRandomFloat32(54775, index) * height;
}
operator smoothMesh<<<index>>>(io PolygonMeshTopology mesh, io Vec3 positions[], Vec3 positionsDoubleBuffer[]) {
//Pseudo-gaussian: center weight = 0.5, neighbor weights sum = 0.5
Vec3 position = positionsDoubleBuffer[ index ];
LocalL16UInt32Array surroundingPoints;
mesh.getPointSurroundingPoints( index, false, surroundingPoints );
UInt32 nbNei = surroundingPoints.size();
if( nbNei ) {
Vec3 neiSum = Vec3(0,0,0);
for( UInt32 i = 0; i < nbNei; ++i ) {
UInt32 neiPt = surroundingPoints.get(i);
neiSum += positionsDoubleBuffer[neiPt];
}
neiSum /= Scalar(nbNei);
position = ( position + neiSum ) * 0.5;
mesh.setPointAttribute(index, positions, position );
}
}
operator entry() {
UInt32 iterations = 40;
Scalar length = 120.0;
Scalar width = 40.0;
Integer lengthSections = 1400;
Integer widthSections = 900;
Scalar height = 10.0;
// first smooth on the CPU.
{
PolygonMesh mesh = PolygonMesh();
mesh.addPlane(Xfo(), length, width, lengthSections, widthSections, true, true);
Ref<Vec3Attribute> positionsAttr = mesh.getAttributes().getPositions();
Vec3 positionsDoubleBuffer[];
positionsDoubleBuffer.resize(positionsAttr.values.size());
noiseOp<<<positionsAttr.values.size()@false>>>(positionsAttr.values, height);
UInt64 start = getCurrentTicks();
for(UInt32 i=0; i<iterations; i++){
smoothMesh<<<mesh.pointCount()@false>>>(mesh.topology, positionsDoubleBuffer, positionsAttr.values);
Vec3 tmp[] = positionsAttr.values;
positionsAttr.values = positionsDoubleBuffer;
positionsAttr.values = tmp;
}
UInt64 end = getCurrentTicks();
report("pointCount: " + mesh.pointCount() + " CPU Time: " + getSecondsBetweenTicks(start, end));
}
// then use the same code to smooth on the GPU.
{
PolygonMesh mesh = PolygonMesh();
mesh.addPlane(Xfo(), length, width, lengthSections, widthSections, true, true);
Ref<Vec3Attribute> positionsAttr = mesh.getAttributes().getPositions();
Vec3 positionsDoubleBuffer[];
positionsDoubleBuffer.resize(positionsAttr.values.size());
mesh.convertToGPU();
positionsAttr.convertToGPU();
positionsDoubleBuffer.convertToGPU();
noiseOp<<<positionsAttr.values.size()@true>>>(positionsAttr.values, height);
UInt64 start = getCurrentTicks();
for(UInt32 i=0; i<iterations; i++){
smoothMesh<<<mesh.pointCount()@true>>>(mesh.topology, positionsDoubleBuffer, positionsAttr.values);
Vec3 tmp[] = positionsAttr.values;
positionsAttr.values = positionsDoubleBuffer;
positionsAttr.values = tmp;
}
UInt64 end = getCurrentTicks();
report("pointCount: " + mesh.pointCount() + " GPU Time: " + getSecondsBetweenTicks(start, end));
}
}
/*
** Output:
pointCount: 1262301 CPU Time: +16.46950416309435
pointCount: 1262301 GPU Time: +16.51505684478995
*/
アトリビュート値の設定¶
通常CPU上でアトリビュート値を設定する際には、PolygonMeshヘルパーメソッドの”setPointAttribute”’を使用することができます。ポリゴンメッシュトポロジー構造は、アトリビュートオブジェクトオブジェクトの代わりにとアトリビュートの値を取得する同様のメソッド群をサポートしています。
/*
** Example: Setting Attribute values using setPointAttribute on the PolygonMeshTopology struct.
*/
require Geometry;
operator randomizeMesh<<<index>>>(io PolygonMeshTopology mesh, io Vec3 positions[]) {
Vec3 position = positionsDoubleBuffer[ index ];
positions[index].y = mathRandomFloat32(54775, index) * height;
mesh.setPointAttribute(index, positions, position );
}
operator entry() {
PolygonMesh mesh();
mesh.addPlane(Xfo(), 120.0, 40.0, 1400, 900, true, true);
Ref<Vec3Attribute> positionsAttr = mesh.getAttributes().getPositions();
mesh.convertToGPU();
positionsAttr.convertToGPU();
randomizeMesh<<<mesh.pointCount()@true>>>(mesh.topology, positionsAttr.values);
}
/*
** Output:
(stdin):9:19: error: 'positionsDoubleBuffer': symbol not found
(stdin):10:3: In assignment:
(stdin):10:24: In binary * expression:
(stdin):10:58: In right-hand side:
(stdin):10:58: error: 'height': symbol not found
(stdin):11:44: error: 'position': symbol not found
*/
スキニング・アトリビュート¶
SkinningAttribute はのGPUコンピュートカーネルに渡される SkinningAttributeData 型、または ‘data’ で呼びだされる構造体メンバーを持っています。
/*
** Example: Using the SkinningAttribute in a GPU compute kernel.
*/
require Geometry;
operator skinMeshPositions<<<index>>>(
io PolygonMeshTopology mesh,
io Vec3 positions[],
io SkinningAttributeData skinningAttr,
Mat44 skinningMatrices[]
){
Vec3 srcPos = positions[index];
LocalL16UInt32Array indices;
LocalL16ScalarArray weights;
skinningAttr.getPairs(index, indices, weights);
Scalar weighSum = 0.0;
Vec3 position(0,0,0);
for( UInt32 i = 0; i < indices.size(); ++i ) {
Scalar boneWeight = weights.get(i);
if( boneWeight == 0.0 )
break;
UInt32 boneId = indices.get(i);
position += (skinningMatrices[boneId] * srcPos) * boneWeight;
weighSum += boneWeight;
}
mesh.setPointAttribute( index, positions, position );
}
operator entry() {
PolygonMesh mesh = PolygonMesh();
mesh.addCuboid(Xfo(), 2.0, 2.0, 2.0);
Ref<Vec3Attribute> positionsAttribute = mesh.getAttributes().positionsAttribute;
Ref<SkinningAttribute> skinningAttribute = mesh.getOrCreateAttribute("skinningData", SkinningAttribute);
// Generate a random set of id/weight pairs per vertex in the mesh.
UInt32 numJoints = 5;
UInt32 maxNumJointerPerVertex = 3;
UInt32 seed = 8516;
UInt32 offset = 0;
for(Integer i=0; i<skinningAttribute.size; i++){
UInt16 numItems = mathRandomFloat32(seed, ++offset) * maxNumJointerPerVertex;
skinningAttribute.setPairCount( i, numItems );
for(Integer j=0; j<numItems; j++){
UInt16 index = mathRandomFloat32(seed, ++offset) * numJoints;
Float32 weight = mathRandomFloat32(seed, ++offset);
skinningAttribute.setPair( i, j, index, weight );
}
}
Mat44 skinningMatrices[];
skinningMatrices.resize(numJoints);
mesh.convertToGPU();
positionsAttribute.convertToGPU();
skinningAttribute.convertToGPU();
skinningMatrices.convertToGPU();
skinMeshPositions<<<mesh.pointCount()@true>>>(
mesh.topology,
positionsAttribute.values,
skinningAttribute.data,
skinningMatrices
);
positionsAttribute.convertToCPU();
skinningAttribute.convertToCPU();
mesh.convertToCPU();
report(mesh.getDesc(true, true));
}
/*
** Output:
(stdin):47:20: error: must use parentheses to call methods
(stdin):64:3: error: must use parentheses to call methods
(stdin):67:5: error: incorrect type (expected 'SkinningAttributeData', actual 'Data')
*/
現在の制限¶
現在のGPUコンピュートのインフラは、以下の制限があります。
- アトリビュートのサイズはGPUコンピュート動作中に変更することはできません。ポリゴンメッシュトポロジーでサポートしているメソッド群だけが値を変更できますが、アトリビュートの共有情報は変更しないでください。
- オブジェクト群はまだGPUでサポートされておらず、単純な型と構造体と配列だけがGPUコンピュートカーネルで使用することができます。
