mirror of
https://github.com/PixarAnimationStudios/OpenSubdiv
synced 2024-11-09 13:50:05 +00:00
Merge pull request #1132 from davidgyu/mtl_ptexviewer_fix
Metal PtexViewer fixes and improvements
This commit is contained in:
Barry Fowler
GitHub
commit
fdd08b96f1
@ -46,7 +46,17 @@ const char* MTLPtexMipmapTexture::GetShaderSource() {
|
||||
MTLPtexMipmapTexture * MTLPtexMipmapTexture::Create(Osd::MTLContext *deviceContext, PtexTexture *reader, int maxLevels) {
|
||||
const auto maxNumPages = 2048;
|
||||
|
||||
PtexMipmapTextureLoader loader(reader, maxNumPages, maxLevels);
|
||||
size_t targetMemory = 0;
|
||||
|
||||
// Read the ptex data and pack the texels
|
||||
bool padAlpha = reader->numChannels()==3;
|
||||
|
||||
PtexMipmapTextureLoader loader(reader,
|
||||
maxNumPages,
|
||||
maxLevels,
|
||||
targetMemory,
|
||||
true/*seemlessMipmap*/,
|
||||
padAlpha);
|
||||
const auto numFaces = loader.GetNumFaces();
|
||||
|
||||
const auto layoutBuffer = [deviceContext->device newBufferWithBytes:loader.GetLayoutBuffer() length:numFaces * 6 * sizeof(short) options:Osd::MTLDefaultStorageMode];
|
||||
@ -54,7 +64,7 @@ MTLPtexMipmapTexture * MTLPtexMipmapTexture::Create(Osd::MTLContext *deviceConte
|
||||
const auto textureDescriptor = [MTLTextureDescriptor new];
|
||||
int bpp = 0;
|
||||
textureDescriptor.pixelFormat = [&]() {
|
||||
const auto numChannels = reader->numChannels();
|
||||
const auto numChannels = reader->numChannels() + padAlpha;
|
||||
switch(reader->dataType()) {
|
||||
case Ptex::dt_uint16:
|
||||
bpp = sizeof(short) * numChannels;
|
||||
|
||||
@ -90,12 +90,17 @@ enum {
|
||||
|
||||
-(void)keyDown:(NSEvent *)event {
|
||||
const auto key = [event.charactersIgnoringModifiers characterAtIndex:0];
|
||||
if(key == '=') {
|
||||
_controller.osdRenderer.tessellationLevel = std::min(_controller.osdRenderer.tessellationLevel + 1, 16.0f);
|
||||
if (hud.KeyDown(key)) {
|
||||
return;
|
||||
} else if(key == '=') {
|
||||
_controller.osdRenderer.tessellationLevel = std::min(_controller.osdRenderer.tessellationLevel + 1, 16);
|
||||
} else if (key == '-') {
|
||||
_controller.osdRenderer.tessellationLevel = std::max(_controller.osdRenderer.tessellationLevel - 1, 0.0f);
|
||||
} else if(!hud.KeyDown(key))
|
||||
_controller.osdRenderer.tessellationLevel = std::max(_controller.osdRenderer.tessellationLevel - 1, 0);
|
||||
} else if (key == 'f') {
|
||||
[_controller.osdRenderer fitFrame];
|
||||
} else {
|
||||
[super keyDown:event];
|
||||
}
|
||||
}
|
||||
|
||||
-(void)scrollWheel:(NSEvent *)event {
|
||||
@ -184,6 +189,9 @@ enum {
|
||||
case kHUD_CB_FRACTIONAL_SPACING:
|
||||
self.osdRenderer.useFractionalTessellation = checked;
|
||||
break;
|
||||
case kHUD_CB_SEAMLESS_MIPMAP:
|
||||
self.osdRenderer.useSeamlessMipmap = checked;
|
||||
break;
|
||||
default:
|
||||
assert("Unknown checkbox ID" && 0);
|
||||
}
|
||||
@ -397,8 +405,6 @@ enum {
|
||||
|
||||
auto commandBuffer = [_commandQueue commandBuffer];
|
||||
|
||||
_osdRenderer.displacementScale = 3;
|
||||
|
||||
double avg = 0;
|
||||
for(int i = 0; i < FRAME_HISTORY; i++)
|
||||
avg += _frameBeginTimestamp[i];
|
||||
@ -408,7 +414,7 @@ enum {
|
||||
|
||||
auto& hud = self.view->hud;
|
||||
if(hud.IsVisible()) {
|
||||
hud.DrawString(10, -120, "Tess level : %f", _osdRenderer.tessellationLevel);
|
||||
hud.DrawString(10, -120, "Tess level : %d", _osdRenderer.tessellationLevel);
|
||||
hud.DrawString(10, -20, "FPS = %3.1f", 1.0 / avg);
|
||||
|
||||
//Disable Culling & Force Fill mode when drawing the UI
|
||||
|
||||
@ -1,111 +0,0 @@
|
||||
//
|
||||
// Copyright 2013 Pixar
|
||||
//
|
||||
// Licensed under the Apache License, Version 2.0 (the "Apache License")
|
||||
// with the following modification; you may not use this file except in
|
||||
// compliance with the Apache License and the following modification to it:
|
||||
// Section 6. Trademarks. is deleted and replaced with:
|
||||
//
|
||||
// 6. Trademarks. This License does not grant permission to use the trade
|
||||
// names, trademarks, service marks, or product names of the Licensor
|
||||
// and its affiliates, except as required to comply with Section 4(c) of
|
||||
// the License and to reproduce the content of the NOTICE file.
|
||||
//
|
||||
// You may obtain a copy of the Apache License at
|
||||
//
|
||||
// http://www.apache.org/licenses/LICENSE-2.0
|
||||
//
|
||||
// Unless required by applicable law or agreed to in writing, software
|
||||
// distributed under the Apache License with the above modification is
|
||||
// distributed on an "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
|
||||
// KIND, either express or implied. See the Apache License for the specific
|
||||
// language governing permissions and limitations under the Apache License.
|
||||
//
|
||||
|
||||
#include "../../regression/common/shape_utils.h"
|
||||
#include "../../regression/shapes/all.h"
|
||||
|
||||
|
||||
static std::vector<ShapeDesc> g_defaultShapes;
|
||||
|
||||
//------------------------------------------------------------------------------
|
||||
static void initShapes() {
|
||||
// g_defaultShapes.push_back(ShapeDesc("temple", Viewer::shapes::temple, kCatmark));
|
||||
g_defaultShapes.push_back(ShapeDesc("catmark_cube_corner0", catmark_cube_corner0, kCatmark));
|
||||
g_defaultShapes.push_back(ShapeDesc("catmark_cube_corner1", catmark_cube_corner1, kCatmark));
|
||||
g_defaultShapes.push_back(ShapeDesc("catmark_cube_corner2", catmark_cube_corner2, kCatmark));
|
||||
g_defaultShapes.push_back(ShapeDesc("catmark_cube_corner3", catmark_cube_corner3, kCatmark));
|
||||
g_defaultShapes.push_back(ShapeDesc("catmark_cube_corner4", catmark_cube_corner4, kCatmark));
|
||||
g_defaultShapes.push_back(ShapeDesc("catmark_cube_creases0", catmark_cube_creases0, kCatmark));
|
||||
g_defaultShapes.push_back(ShapeDesc("catmark_cube_creases1", catmark_cube_creases1, kCatmark));
|
||||
g_defaultShapes.push_back(ShapeDesc("catmark_cube_creases2", catmark_cube_creases2, kCatmark));
|
||||
g_defaultShapes.push_back(ShapeDesc("catmark_cube", catmark_cube, kCatmark));
|
||||
g_defaultShapes.push_back(ShapeDesc("catmark_dart_edgecorner", catmark_dart_edgecorner, kCatmark));
|
||||
g_defaultShapes.push_back(ShapeDesc("catmark_dart_edgeonly", catmark_dart_edgeonly, kCatmark));
|
||||
g_defaultShapes.push_back(ShapeDesc("catmark_edgecorner", catmark_edgecorner, kCatmark));
|
||||
g_defaultShapes.push_back(ShapeDesc("catmark_edgeonly", catmark_edgeonly, kCatmark));
|
||||
g_defaultShapes.push_back(ShapeDesc("catmark_chaikin0", catmark_chaikin0, kCatmark));
|
||||
g_defaultShapes.push_back(ShapeDesc("catmark_chaikin1", catmark_chaikin1, kCatmark));
|
||||
g_defaultShapes.push_back(ShapeDesc("catmark_chaikin2", catmark_chaikin2, kCatmark));
|
||||
g_defaultShapes.push_back(ShapeDesc("catmark_fan", catmark_fan, kCatmark));
|
||||
g_defaultShapes.push_back(ShapeDesc("catmark_flap", catmark_flap, kCatmark));
|
||||
g_defaultShapes.push_back(ShapeDesc("catmark_flap2", catmark_flap2, kCatmark));
|
||||
g_defaultShapes.push_back(ShapeDesc("catmark_fvar_bound0", catmark_fvar_bound0, kCatmark));
|
||||
g_defaultShapes.push_back(ShapeDesc("catmark_fvar_bound1", catmark_fvar_bound1, kCatmark));
|
||||
g_defaultShapes.push_back(ShapeDesc("catmark_fvar_bound2", catmark_fvar_bound2, kCatmark));
|
||||
g_defaultShapes.push_back(ShapeDesc("catmark_gregory_test0", catmark_gregory_test0, kCatmark));
|
||||
g_defaultShapes.push_back(ShapeDesc("catmark_gregory_test1", catmark_gregory_test1, kCatmark));
|
||||
g_defaultShapes.push_back(ShapeDesc("catmark_gregory_test2", catmark_gregory_test2, kCatmark));
|
||||
g_defaultShapes.push_back(ShapeDesc("catmark_gregory_test3", catmark_gregory_test3, kCatmark));
|
||||
g_defaultShapes.push_back(ShapeDesc("catmark_gregory_test4", catmark_gregory_test4, kCatmark));
|
||||
g_defaultShapes.push_back(ShapeDesc("catmark_gregory_test5", catmark_gregory_test5, kCatmark));
|
||||
g_defaultShapes.push_back(ShapeDesc("catmark_gregory_test6", catmark_gregory_test6, kCatmark));
|
||||
g_defaultShapes.push_back(ShapeDesc("catmark_gregory_test7", catmark_gregory_test7, kCatmark));
|
||||
g_defaultShapes.push_back(ShapeDesc("catmark_hole_test1", catmark_hole_test1, kCatmark));
|
||||
g_defaultShapes.push_back(ShapeDesc("catmark_hole_test2", catmark_hole_test2, kCatmark));
|
||||
g_defaultShapes.push_back(ShapeDesc("catmark_hole_test3", catmark_hole_test3, kCatmark));
|
||||
g_defaultShapes.push_back(ShapeDesc("catmark_hole_test4", catmark_hole_test4, kCatmark));
|
||||
g_defaultShapes.push_back(ShapeDesc("catmark_lefthanded", catmark_lefthanded, kCatmark, true /*isLeftHanded*/));
|
||||
g_defaultShapes.push_back(ShapeDesc("catmark_righthanded", catmark_righthanded, kCatmark));
|
||||
// g_defaultShapes.push_back(ShapeDesc("catmark_pole8", catmark_pole8, kCatmark));
|
||||
// g_defaultShapes.push_back(ShapeDesc("catmark_pole64", catmark_pole64, kCatmark));
|
||||
// g_defaultShapes.push_back(ShapeDesc("catmark_nonman_quadpole8", catmark_nonman_quadpole8, kCatmark));
|
||||
// g_defaultShapes.push_back(ShapeDesc("catmark_nonman_quadpole64", catmark_nonman_quadpole64, kCatmark));
|
||||
// g_defaultShapes.push_back(ShapeDesc("catmark_nonman_quadpole360", catmark_nonman_quadpole360, kCatmark));
|
||||
g_defaultShapes.push_back(ShapeDesc("catmark_pyramid_creases0", catmark_pyramid_creases0, kCatmark));
|
||||
g_defaultShapes.push_back(ShapeDesc("catmark_pyramid_creases1", catmark_pyramid_creases1, kCatmark));
|
||||
g_defaultShapes.push_back(ShapeDesc("catmark_pyramid", catmark_pyramid, kCatmark));
|
||||
g_defaultShapes.push_back(ShapeDesc("catmark_tent_creases0", catmark_tent_creases0, kCatmark));
|
||||
g_defaultShapes.push_back(ShapeDesc("catmark_tent_creases1", catmark_tent_creases1, kCatmark));
|
||||
g_defaultShapes.push_back(ShapeDesc("catmark_tent", catmark_tent, kCatmark));
|
||||
g_defaultShapes.push_back(ShapeDesc("catmark_torus", catmark_torus, kCatmark));
|
||||
g_defaultShapes.push_back(ShapeDesc("catmark_torus_creases0", catmark_torus_creases0, kCatmark));
|
||||
g_defaultShapes.push_back(ShapeDesc("catmark_single_crease", catmark_single_crease, kCatmark));
|
||||
g_defaultShapes.push_back(ShapeDesc("catmark_smoothtris0", catmark_smoothtris0, kCatmark));
|
||||
g_defaultShapes.push_back(ShapeDesc("catmark_smoothtris1", catmark_smoothtris1, kCatmark));
|
||||
// // g_defaultShapes.push_back( ShapeDesc("catmark_square_hedit0", catmark_square_hedit0, kCatmark ) );
|
||||
// // g_defaultShapes.push_back( ShapeDesc("catmark_square_hedit1", catmark_square_hedit1, kCatmark ) );
|
||||
// // g_defaultShapes.push_back( ShapeDesc("catmark_square_hedit2", catmark_square_hedit2, kCatmark ) );
|
||||
// // g_defaultShapes.push_back( ShapeDesc("catmark_square_hedit3", catmark_square_hedit3, kCatmark ) );
|
||||
g_defaultShapes.push_back(ShapeDesc("catmark_bishop", catmark_bishop, kCatmark));
|
||||
g_defaultShapes.push_back(ShapeDesc("catmark_car", catmark_car, kCatmark));
|
||||
g_defaultShapes.push_back(ShapeDesc("catmark_helmet", catmark_helmet, kCatmark));
|
||||
g_defaultShapes.push_back(ShapeDesc("catmark_pawn", catmark_pawn, kCatmark));
|
||||
g_defaultShapes.push_back(ShapeDesc("catmark_rook", catmark_rook, kCatmark));
|
||||
//
|
||||
// g_defaultShapes.push_back(ShapeDesc("bilinear_cube", bilinear_cube, kBilinear));
|
||||
//
|
||||
// g_defaultShapes.push_back(ShapeDesc("loop_cube_creases0", loop_cube_creases0, kLoop));
|
||||
// g_defaultShapes.push_back(ShapeDesc("loop_cube_creases1", loop_cube_creases1, kLoop));
|
||||
// g_defaultShapes.push_back(ShapeDesc("loop_cube", loop_cube, kLoop));
|
||||
// g_defaultShapes.push_back(ShapeDesc("loop_icosahedron", loop_icosahedron, kLoop));
|
||||
// g_defaultShapes.push_back(ShapeDesc("loop_saddle_edgecorner", loop_saddle_edgecorner, kLoop));
|
||||
// g_defaultShapes.push_back(ShapeDesc("loop_saddle_edgeonly", loop_saddle_edgeonly, kLoop));
|
||||
// g_defaultShapes.push_back(ShapeDesc("loop_triangle_edgecorner", loop_triangle_edgecorner, kLoop));
|
||||
// g_defaultShapes.push_back(ShapeDesc("loop_triangle_edgeonly", loop_triangle_edgeonly, kLoop));
|
||||
// g_defaultShapes.push_back(ShapeDesc("loop_chaikin0", loop_chaikin0, kLoop));
|
||||
// g_defaultShapes.push_back(ShapeDesc("loop_chaikin1", loop_chaikin1, kLoop));
|
||||
// g_defaultShapes.push_back(ShapeDesc("loop_pole8", loop_pole8, kLoop));
|
||||
// g_defaultShapes.push_back(ShapeDesc("loop_pole64", loop_pole64, kLoop));
|
||||
// g_defaultShapes.push_back(ShapeDesc("loop_pole360", loop_pole360, kLoop));
|
||||
}
|
||||
@ -63,13 +63,12 @@ typedef struct {
|
||||
|
||||
-(id<MTLRenderCommandEncoder>)drawFrame:(id<MTLCommandBuffer>)commandBuffer;
|
||||
|
||||
-(void)fitFrame;
|
||||
|
||||
@property (readonly, nonatomic) id<OSDRendererDelegate> delegate;
|
||||
|
||||
@property (nonatomic) unsigned refinementLevel;
|
||||
@property (nonatomic) float tessellationLevel;
|
||||
|
||||
@property (readonly, nonatomic) NSArray<NSString*>* loadedModels;
|
||||
@property (nonatomic) NSString* currentModel;
|
||||
@property (nonatomic) int tessellationLevel;
|
||||
|
||||
@property (readonly, nonatomic) Camera* camera;
|
||||
|
||||
@ -80,7 +79,8 @@ typedef struct {
|
||||
@property (nonatomic) bool usePatchIndexBuffer;
|
||||
@property (nonatomic) bool usePatchBackfaceCulling;
|
||||
@property (nonatomic) bool usePatchClipCulling;
|
||||
@property (nonatomic) bool useSingleCrease;
|
||||
@property (nonatomic) bool useSmoothCornerPatch;
|
||||
@property (nonatomic) bool useSingleCreasePatch;
|
||||
@property (nonatomic) bool useInfinitelySharpPatch;
|
||||
@property (nonatomic) bool useStageIn;
|
||||
@property (nonatomic) bool usePrimitiveBackfaceCulling;
|
||||
@ -91,6 +91,7 @@ typedef struct {
|
||||
@property (nonatomic) bool displayControlMeshVertices;
|
||||
@property (nonatomic) bool displaySpecular;
|
||||
@property (nonatomic) bool displayOcclusion;
|
||||
@property (nonatomic) bool yup;
|
||||
@property (nonatomic) float mipmapBias;
|
||||
@property (nonatomic) float displacementScale;
|
||||
|
||||
|
||||
@ -1,6 +1,6 @@
|
||||
#line 0 "examples/mtlPtexViewer/mtlPtexViewer.metal"
|
||||
//
|
||||
// Copyright 2013 Pixar
|
||||
// Copyright 2013-2019 Pixar
|
||||
//
|
||||
// Licensed under the Apache License, Version 2.0 (the "Apache License")
|
||||
// with the following modification; you may not use this file except in
|
||||
@ -26,6 +26,10 @@
|
||||
#include <metal_stdlib>
|
||||
using namespace metal;
|
||||
|
||||
#if OSD_IS_ADAPTIVE
|
||||
static_assert(!OSD_ENABLE_SCREENSPACE_TESSELLATION || !USE_PTVS_FACTORS, "USE_PTVS_FACTORS cannot be enabled if OSD_ENABLE_SCREENSPACE_TESSELLATION is enabled");
|
||||
#endif
|
||||
|
||||
struct Config {
|
||||
float displacementScale;
|
||||
float mipmapBias;
|
||||
@ -64,11 +68,11 @@ const constant float4 patchColors[] = {
|
||||
float4(0.0f, 0.8f, 0.75f, 1.0f), // boundary pattern 4
|
||||
|
||||
float4(0.0f, 1.0f, 0.0f, 1.0f), // corner
|
||||
float4(0.25f, 0.25f, 0.25f, 1.0f), // corner pattern 0
|
||||
float4(0.25f, 0.25f, 0.25f, 1.0f), // corner pattern 1
|
||||
float4(0.25f, 0.25f, 0.25f, 1.0f), // corner pattern 2
|
||||
float4(0.25f, 0.25f, 0.25f, 1.0f), // corner pattern 3
|
||||
float4(0.25f, 0.25f, 0.25f, 1.0f), // corner pattern 4
|
||||
float4(0.5f, 1.0f, 0.5f, 1.0f), // corner pattern 0
|
||||
float4(0.5f, 1.0f, 0.5f, 1.0f), // corner pattern 1
|
||||
float4(0.5f, 1.0f, 0.5f, 1.0f), // corner pattern 2
|
||||
float4(0.5f, 1.0f, 0.5f, 1.0f), // corner pattern 3
|
||||
float4(0.5f, 1.0f, 0.5f, 1.0f), // corner pattern 4
|
||||
|
||||
float4(1.0f, 1.0f, 0.0f, 1.0f), // gregory
|
||||
float4(1.0f, 1.0f, 0.0f, 1.0f), // gregory
|
||||
@ -95,33 +99,32 @@ const constant float4 patchColors[] = {
|
||||
float4
|
||||
getAdaptivePatchColor(int3 patchParam, float sharpness)
|
||||
{
|
||||
int pattern = popcount(OsdGetPatchTransitionMask(patchParam));
|
||||
int edgeCount = popcount(OsdGetPatchBoundaryMask(patchParam));
|
||||
|
||||
int patchType = 0;
|
||||
#if OSD_PATCH_ENABLE_SINGLE_CREASE
|
||||
if (sharpness > 0) {
|
||||
pattern = 1;
|
||||
}
|
||||
#endif
|
||||
|
||||
int edgeCount = popcount(OsdGetPatchBoundaryMask(patchParam));
|
||||
if (edgeCount == 1) {
|
||||
patchType = 2; // BOUNDARY
|
||||
}
|
||||
if (edgeCount == 2) {
|
||||
if (edgeCount > 1) {
|
||||
patchType = 3; // CORNER
|
||||
}
|
||||
|
||||
// XXX: it looks like edgeCount != 0 for some gregory boundary patches.
|
||||
// there might be a bug somewhere...
|
||||
#if OSD_PATCH_GREGORY
|
||||
#if OSD_PATCH_ENABLE_SINGLE_CREASE
|
||||
if (sharpness > 0) {
|
||||
patchType = 1;
|
||||
}
|
||||
#elif OSD_PATCH_GREGORY
|
||||
patchType = 4;
|
||||
#elif OSD_PATCH_GREGORY_BOUNDARY
|
||||
patchType = 5;
|
||||
#elif OSD_PATCH_GREGORY_BASIS
|
||||
patchType = 6;
|
||||
#elif OSD_PATCH_GREGORY_TRIANGLE
|
||||
patchType = 6;
|
||||
#endif
|
||||
|
||||
int pattern = popcount(OsdGetPatchTransitionMask(patchParam));
|
||||
|
||||
return patchColors[6*patchType + pattern];
|
||||
}
|
||||
|
||||
@ -248,10 +251,9 @@ struct FragmentInput
|
||||
|
||||
#if OSD_IS_ADAPTIVE
|
||||
#if USE_STAGE_IN
|
||||
#if OSD_PATCH_REGULAR
|
||||
#if OSD_PATCH_REGULAR || OSD_PATCH_BOX_SPLINE_TRIANGLE
|
||||
struct ControlPoint
|
||||
{
|
||||
|
||||
float3 P [[attribute(0)]];
|
||||
#if OSD_PATCH_ENABLE_SINGLE_CREASE
|
||||
float3 P1 [[attribute(1)]];
|
||||
@ -261,6 +263,21 @@ struct ControlPoint
|
||||
#endif
|
||||
#endif
|
||||
};
|
||||
#elif OSD_PATCH_GREGORY || OSD_PATCH_GREGORY_BOUNDARY
|
||||
struct ControlPoint
|
||||
{
|
||||
float3 P [[attribute(0)]];
|
||||
float3 Ep [[attribute(1)]];
|
||||
float3 Em [[attribute(2)]];
|
||||
float3 Fp [[attribute(3)]];
|
||||
float3 Fm [[attribute(4)]];
|
||||
};
|
||||
#elif OSD_PATCH_GREGORY_BASIS || OSD_PATCH_GREGORY_TRIANGLE
|
||||
struct ControlPoint
|
||||
{
|
||||
float3 position [[attribute(0)]];
|
||||
};
|
||||
#endif
|
||||
|
||||
struct PatchInput
|
||||
{
|
||||
@ -271,77 +288,67 @@ struct PatchInput
|
||||
#endif
|
||||
int3 patchParam [[attribute(10)]];
|
||||
};
|
||||
#elif OSD_PATCH_GREGORY || OSD_PATCH_GREGORY_BOUNDARY
|
||||
struct ControlPoint
|
||||
{
|
||||
|
||||
float3 P [[attribute(0)]];
|
||||
float3 Ep [[attribute(1)]];
|
||||
float3 Em [[attribute(2)]];
|
||||
float3 Fp [[attribute(3)]];
|
||||
float3 Fm [[attribute(4)]];
|
||||
};
|
||||
|
||||
struct PatchInput
|
||||
{
|
||||
patch_control_point<ControlPoint> cv;
|
||||
int3 patchParam [[attribute(10)]];
|
||||
};
|
||||
#elif OSD_PATCH_GREGORY_BASIS
|
||||
struct ControlPoint
|
||||
{
|
||||
float3 position [[attribute(0)]];
|
||||
};
|
||||
|
||||
struct PatchInput
|
||||
{
|
||||
patch_control_point<ControlPoint> cv;
|
||||
int3 patchParam [[attribute(10)]];
|
||||
};
|
||||
#endif
|
||||
|
||||
|
||||
#if OSD_PATCH_REGULAR || OSD_PATCH_GREGORY_BASIS || OSD_PATCH_GREGORY || OSD_PATCH_GREGORY_BOUNDARY
|
||||
typedef MTLQuadTessellationFactorsHalf PatchTessFactors;
|
||||
#elif OSD_PATCH_BOX_SPLINE_TRIANGLE || OSD_PATCH_GREGORY_TRIANGLE
|
||||
typedef MTLTriangleTessellationFactorsHalf PatchTessFactors;
|
||||
#endif
|
||||
|
||||
|
||||
kernel void compute_main(
|
||||
const constant PerFrameConstants& frameConsts [[buffer(FRAME_CONST_BUFFER_INDEX)]],
|
||||
unsigned thread_position_in_grid [[thread_position_in_grid]],
|
||||
unsigned thread_position_in_threadgroup [[thread_position_in_threadgroup]],
|
||||
unsigned threadgroup_position_in_grid [[threadgroup_position_in_grid]],
|
||||
OsdPatchParamBufferSet osdBuffers,
|
||||
device MTLQuadTessellationFactorsHalf* quadTessellationFactors [[buffer(QUAD_TESSFACTORS_INDEX)]]
|
||||
device PatchTessFactors* patchTessellationFactors [[buffer(PATCH_TESSFACTORS_INDEX)]]
|
||||
#if OSD_USE_PATCH_INDEX_BUFFER
|
||||
,device unsigned* patchIndex [[buffer(OSD_PATCH_INDEX_BUFFER_INDEX)]]
|
||||
,device MTLDrawPatchIndirectArguments* drawIndirectCommands [[buffer(OSD_DRAWINDIRECT_BUFFER_INDEX)]]
|
||||
#endif
|
||||
)
|
||||
{
|
||||
//----------------------------------------------------------
|
||||
// OSD Kernel Setup
|
||||
//----------------------------------------------------------
|
||||
|
||||
#define PATCHES_PER_THREADGROUP (THREADS_PER_THREADGROUP / THREADS_PER_PATCH)
|
||||
int const primitiveID = thread_position_in_grid / THREADS_PER_PATCH;
|
||||
int const primitiveIDInTG = thread_position_in_threadgroup / THREADS_PER_PATCH;
|
||||
int const vertexIndex = threadgroup_position_in_grid * PATCHES_PER_THREADGROUP * CONTROL_POINTS_PER_PATCH +
|
||||
thread_position_in_threadgroup * CONTROL_POINTS_PER_THREAD;
|
||||
int const vertexIndexInTG = thread_position_in_threadgroup * CONTROL_POINTS_PER_THREAD;
|
||||
int const invocationID = (thread_position_in_threadgroup * VERTEX_CONTROL_POINTS_PER_THREAD) % (THREADS_PER_PATCH*VERTEX_CONTROL_POINTS_PER_THREAD);
|
||||
|
||||
//Contains the shared patchParam value used by all threads that act upon a single patch
|
||||
//the .z (sharpness) field is set to -1 (NAN) if that patch should be culled to signal other threads to return.
|
||||
threadgroup int3 patchParam[PATCHES_PER_THREADGROUP];
|
||||
threadgroup PatchVertexType patchVertices[PATCHES_PER_THREADGROUP * CONTROL_POINTS_PER_PATCH];
|
||||
|
||||
const auto real_threadgroup = thread_position_in_grid / REAL_THREADGROUP_DIVISOR;
|
||||
const auto subthreadgroup_in_threadgroup = thread_position_in_threadgroup / REAL_THREADGROUP_DIVISOR;
|
||||
const auto real_thread_in_threadgroup = thread_position_in_threadgroup & (REAL_THREADGROUP_DIVISOR - 1);
|
||||
|
||||
#if NEEDS_BARRIER
|
||||
const auto validThread = thread_position_in_grid * CONTROL_POINTS_PER_THREAD < osdBuffers.kernelExecutionLimit;
|
||||
#else
|
||||
const auto validThread = true;
|
||||
if(thread_position_in_grid * CONTROL_POINTS_PER_THREAD >= osdBuffers.kernelExecutionLimit)
|
||||
return;
|
||||
#endif
|
||||
|
||||
if(validThread)
|
||||
//----------------------------------------------------------
|
||||
// OSD Vertex Transform
|
||||
//----------------------------------------------------------
|
||||
{
|
||||
patchParam[subthreadgroup_in_threadgroup] = OsdGetPatchParam(real_threadgroup, osdBuffers.patchParamBuffer);
|
||||
patchParam[primitiveIDInTG] = OsdGetPatchParam(primitiveID, osdBuffers.patchParamBuffer);
|
||||
|
||||
for(unsigned threadOffset = 0; threadOffset < CONTROL_POINTS_PER_THREAD; threadOffset++)
|
||||
for (unsigned threadOffset = 0; threadOffset < CONTROL_POINTS_PER_THREAD; ++threadOffset)
|
||||
{
|
||||
const auto vertexId = osdBuffers.indexBuffer[(thread_position_in_grid * CONTROL_POINTS_PER_THREAD + threadOffset) * IndexLookupStride];
|
||||
const auto v = osdBuffers.vertexBuffer[vertexId];
|
||||
if (vertexIndexInTG + threadOffset < PATCHES_PER_THREADGROUP * CONTROL_POINTS_PER_PATCH)
|
||||
{
|
||||
const auto vertexId = osdBuffers.indexBuffer[(vertexIndex + threadOffset)];
|
||||
const auto v = osdBuffers.vertexBuffer[vertexId];
|
||||
|
||||
threadgroup auto& patchVertex = patchVertices[thread_position_in_threadgroup * CONTROL_POINTS_PER_THREAD + threadOffset];
|
||||
OsdComputePerVertex(float4(v.position,1), patchVertex, vertexId, frameConsts.ModelViewProjectionMatrix, osdBuffers);
|
||||
//User per vertex goes here, modifying 'patchVertex'
|
||||
threadgroup auto& patchVertex = patchVertices[vertexIndexInTG + threadOffset];
|
||||
|
||||
//----------------------------------------------------------
|
||||
// User Vertex Transform
|
||||
//----------------------------------------------------------
|
||||
|
||||
OsdComputePerVertex(float4(v.position,1), patchVertex, vertexId, frameConsts.ModelViewProjectionMatrix, osdBuffers);
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
@ -349,27 +356,31 @@ kernel void compute_main(
|
||||
threadgroup_barrier(mem_flags::mem_threadgroup);
|
||||
#endif
|
||||
|
||||
if(validThread)
|
||||
//----------------------------------------------------------
|
||||
// OSD Patch Cull
|
||||
//----------------------------------------------------------
|
||||
{
|
||||
#if PATCHES_PER_THREADGROUP > 1
|
||||
auto patch = patchVertices + subthreadgroup_in_threadgroup * CONTROL_POINTS_PER_THREAD * CONTROL_POINTS_PER_PATCH;
|
||||
#else
|
||||
//Small optimization for the '1 patch per threadgroup' case
|
||||
auto patch = patchVertices;
|
||||
#endif
|
||||
auto patch = patchVertices + primitiveIDInTG * CONTROL_POINTS_PER_PATCH;
|
||||
|
||||
if(!OsdCullPerPatchVertex(patch, frameConsts.ModelViewMatrix))
|
||||
if (!OsdCullPerPatchVertex(patch, frameConsts.ModelViewMatrix))
|
||||
{
|
||||
#if !OSD_USE_PATCH_INDEX_BUFFER
|
||||
quadTessellationFactors[real_threadgroup].edgeTessellationFactor[0] = 0.0h;
|
||||
quadTessellationFactors[real_threadgroup].edgeTessellationFactor[1] = 0.0h;
|
||||
quadTessellationFactors[real_threadgroup].edgeTessellationFactor[2] = 0.0h;
|
||||
quadTessellationFactors[real_threadgroup].edgeTessellationFactor[3] = 0.0h;
|
||||
quadTessellationFactors[real_threadgroup].insideTessellationFactor[0] = 0.0h;
|
||||
quadTessellationFactors[real_threadgroup].insideTessellationFactor[1] = 0.0h;
|
||||
#if OSD_PATCH_REGULAR || OSD_PATCH_GREGORY_BASIS || OSD_PATCH_GREGORY || OSD_PATCH_GREGORY_BOUNDARY
|
||||
patchTessellationFactors[primitiveID].edgeTessellationFactor[0] = 0.0h;
|
||||
patchTessellationFactors[primitiveID].edgeTessellationFactor[1] = 0.0h;
|
||||
patchTessellationFactors[primitiveID].edgeTessellationFactor[2] = 0.0h;
|
||||
patchTessellationFactors[primitiveID].edgeTessellationFactor[3] = 0.0h;
|
||||
patchTessellationFactors[primitiveID].insideTessellationFactor[0] = 0.0h;
|
||||
patchTessellationFactors[primitiveID].insideTessellationFactor[1] = 0.0h;
|
||||
#elif OSD_PATCH_BOX_SPLINE_TRIANGLE || OSD_PATCH_GREGORY_TRIANGLE
|
||||
patchTessellationFactors[primitiveID].edgeTessellationFactor[0] = 0.0h;
|
||||
patchTessellationFactors[primitiveID].edgeTessellationFactor[1] = 0.0h;
|
||||
patchTessellationFactors[primitiveID].edgeTessellationFactor[2] = 0.0h;
|
||||
patchTessellationFactors[primitiveID].insideTessellationFactor = 0.0h;
|
||||
#endif
|
||||
#endif
|
||||
|
||||
patchParam[subthreadgroup_in_threadgroup].z = -1;
|
||||
patchParam[primitiveIDInTG].z = -1;
|
||||
#if !NEEDS_BARRIER
|
||||
return;
|
||||
#endif
|
||||
@ -380,18 +391,24 @@ kernel void compute_main(
|
||||
threadgroup_barrier(mem_flags::mem_threadgroup);
|
||||
#endif
|
||||
|
||||
if(validThread && patchParam[subthreadgroup_in_threadgroup].z != -1)
|
||||
//----------------------------------------------------------
|
||||
// OSD Patch Compute
|
||||
//----------------------------------------------------------
|
||||
if (patchParam[primitiveIDInTG].z != -1)
|
||||
{
|
||||
for(unsigned threadOffset = 0; threadOffset < CONTROL_POINTS_PER_THREAD; threadOffset++)
|
||||
for (unsigned threadOffset = 0; threadOffset < VERTEX_CONTROL_POINTS_PER_THREAD; ++threadOffset)
|
||||
{
|
||||
OsdComputePerPatchVertex(
|
||||
patchParam[subthreadgroup_in_threadgroup],
|
||||
real_thread_in_threadgroup * CONTROL_POINTS_PER_THREAD + threadOffset,
|
||||
real_threadgroup,
|
||||
thread_position_in_grid * CONTROL_POINTS_PER_THREAD + threadOffset,
|
||||
patchVertices + subthreadgroup_in_threadgroup * CONTROL_POINTS_PER_PATCH,
|
||||
osdBuffers
|
||||
);
|
||||
if (invocationID + threadOffset < VERTEX_CONTROL_POINTS_PER_PATCH)
|
||||
{
|
||||
OsdComputePerPatchVertex(
|
||||
patchParam[primitiveIDInTG],
|
||||
invocationID + threadOffset,
|
||||
primitiveID,
|
||||
invocationID + threadOffset + primitiveID * VERTEX_CONTROL_POINTS_PER_PATCH,
|
||||
patchVertices + primitiveIDInTG * CONTROL_POINTS_PER_PATCH,
|
||||
osdBuffers
|
||||
);
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
@ -399,26 +416,29 @@ kernel void compute_main(
|
||||
threadgroup_barrier(mem_flags::mem_device_and_threadgroup);
|
||||
#endif
|
||||
|
||||
if(validThread && real_thread_in_threadgroup == 0)
|
||||
//----------------------------------------------------------
|
||||
// OSD Tessellation Factors
|
||||
//----------------------------------------------------------
|
||||
if (invocationID == 0)
|
||||
{
|
||||
|
||||
#if OSD_USE_PATCH_INDEX_BUFFER
|
||||
const auto patchId = atomic_fetch_add_explicit((device atomic_uint*)&drawIndirectCommands->patchCount, 1, memory_order_relaxed);
|
||||
patchIndex[patchId] = real_threadgroup;
|
||||
patchIndex[patchId] = primitiveID;
|
||||
#else
|
||||
const auto patchId = real_threadgroup;
|
||||
const auto patchId = primitiveID;
|
||||
#endif
|
||||
|
||||
OsdComputePerPatchFactors(
|
||||
patchParam[subthreadgroup_in_threadgroup],
|
||||
patchParam[primitiveIDInTG],
|
||||
frameConsts.TessLevel,
|
||||
real_threadgroup,
|
||||
primitiveID,
|
||||
frameConsts.ProjectionMatrix,
|
||||
frameConsts.ModelViewMatrix,
|
||||
osdBuffers,
|
||||
patchVertices + subthreadgroup_in_threadgroup * CONTROL_POINTS_PER_PATCH,
|
||||
quadTessellationFactors[patchId]
|
||||
);
|
||||
patchVertices + primitiveIDInTG * CONTROL_POINTS_PER_PATCH,
|
||||
patchTessellationFactors[patchId]
|
||||
);
|
||||
}
|
||||
}
|
||||
|
||||
@ -607,14 +627,14 @@ fragment float4 fragment_main(
|
||||
float3 normal = perturbNormalFromDisplacement(input.position.xyz,
|
||||
input.normal,
|
||||
input.patchCoord,
|
||||
config.mipmapBias,
|
||||
mipmapBias,
|
||||
textureDisplace_Data,
|
||||
textureDisplace_Packing,
|
||||
config.displacementScale);
|
||||
displacementScale);
|
||||
#elif NORMAL_BIQUADRATIC || NORMAL_BIQUADRATIC_WG
|
||||
float4 du, dv;
|
||||
float4 disp = PtexMipmapLookupQuadratic(du, dv, input.patchCoord,
|
||||
config.mipmapBias,
|
||||
mipmapBias,
|
||||
textureDisplace_Data,
|
||||
textureDisplace_Packing);
|
||||
|
||||
@ -655,15 +675,15 @@ fragment float4 fragment_main(
|
||||
textureImage_Packing);
|
||||
#elif COLOR_PATCHTYPE
|
||||
float4 texColor = lighting(float4(input.patchColor), input.position.xyz, normal, 0, lightSource);
|
||||
outColor = texColor;
|
||||
outColor = max(texColor, shade);
|
||||
return outColor;
|
||||
#elif COLOR_PATCHCOORD
|
||||
float4 texColor = lighting(input.patchCoord, input.position.xyz, normal, 0, lightSource);
|
||||
outColor = texColor;
|
||||
outColor = max(texColor, shade);
|
||||
return outColor;
|
||||
#elif COLOR_NORMAL
|
||||
float4 texColor = float4(normal.x, normal.y, normal.z, 1);
|
||||
outColor = texColor;
|
||||
outColor = max(texColor, shade);
|
||||
return outColor;
|
||||
#else // COLOR_NONE
|
||||
float4 texColor = float4(0.5, 0.5, 0.5, 1);
|
||||
@ -672,7 +692,7 @@ fragment float4 fragment_main(
|
||||
// ------------ occlusion ---------------
|
||||
|
||||
#if USE_PTEX_OCCLUSION
|
||||
float occ = PtexMipmapLookup(input.patchCoord, config.mipmapBias,
|
||||
float occ = PtexMipmapLookup(input.patchCoord, mipmapBias,
|
||||
textureOcclusion_Data,
|
||||
textureOcclusion_Packing).x;
|
||||
#else
|
||||
@ -682,7 +702,7 @@ fragment float4 fragment_main(
|
||||
// ------------ specular ---------------
|
||||
|
||||
#if USE_PTEX_SPECULAR
|
||||
float specular = PtexMipmapLookup(input.patchCoord, config.mipmapBias,
|
||||
float specular = PtexMipmapLookup(input.patchCoord, mipmapBias,
|
||||
textureSpecular_Data,
|
||||
textureSpecular_Packing).x;
|
||||
#else
|
||||
|
||||
File diff suppressed because it is too large
Load Diff
Loading…
Reference in New Issue
Block a user