При захвате кадра графического процессора из моего приложения я вижу следующие сообщения в разделе Статистика конвейера -> Примечания:
Ошибка предварительной загрузки буфера
Убедитесь, что ваш размер данных кратен4 байта и выровнены по 4 байта и попробуйте использовать простой шаблон доступа.Для постоянных буферов попробуйте использовать фиксированный размер буфера.
pointLightBufferCenterAndRadius не может быть повышен - lightCuller.metal: light_culler
Вот моя инициализация буфера:
const int MaxLights = 2048;
pointLightCenterAndRadiusBuffer = [GfxDevice::GetMetalDevice() newBufferWithLength:MaxLights * sizeof( Vec4 )
options:MTLResourceCPUCacheModeDefaultCache];
pointLightCenterAndRadiusBuffer.label = @"pointLightCenterAndRadiusBuffer";
Вот соответствующая часть моего шейдера:
kernel void light_culler(texture2d<float, access::read> depthNormalsTexture [[texture(0)]],
constant Uniforms& uniforms [[ buffer(0) ]],
constant float4* pointLightBufferCenterAndRadius [[ buffer(1) ]],
device uint* perTileLightIndexBufferOut [[ buffer(2) ]],
constant float4* spotLightBufferCenterAndRadius [[ buffer(3) ]],
uint2 gid [[thread_position_in_grid]],
uint2 tid [[thread_position_in_threadgroup]],
uint2 dtid [[threadgroup_position_in_grid]])
{
threadgroup uint ldsLightIdx[ MAX_NUM_LIGHTS_PER_TILE ];
threadgroup atomic_uint ldsZMax;
threadgroup atomic_uint ldsZMin;
threadgroup atomic_uint ldsLightIdxCounter;
uint2 globalIdx = gid;
uint2 localIdx = tid;
uint2 groupIdx = dtid;
uint localIdxFlattened = localIdx.x + localIdx.y * TILE_RES;
uint tileIdxFlattened = groupIdx.x + groupIdx.y * GetNumTilesX( uniforms.windowWidth );
if (localIdxFlattened == 0)
{
atomic_store_explicit( &ldsZMin, 0x7f7fffff, memory_order_relaxed ); // FLT_MAX as uint
atomic_store_explicit( &ldsZMax, 0, memory_order_relaxed );
atomic_store_explicit( &ldsLightIdxCounter, 0, memory_order_relaxed );
}
float4 frustumEqn[ 4 ];
{
uint pxm = TILE_RES * groupIdx.x;
uint pym = TILE_RES * groupIdx.y;
uint pxp = TILE_RES * (groupIdx.x + 1);
uint pyp = TILE_RES * (groupIdx.y + 1);
float winWidth = float( TILE_RES * GetNumTilesX( uniforms.windowWidth ) );
float winHeight = float( TILE_RES * GetNumTilesY( uniforms.windowHeight) );
float4 v0 = float4( pxm / winWidth * 2.0f - 1.0f, (winHeight - pym) / winHeight * 2.0f - 1.0f, 1.0f, 1.0f );
float4 v1 = float4( pxp / winWidth * 2.0f - 1.0f, (winHeight - pym) / winHeight * 2.0f - 1.0f, 1.0f, 1.0f );
float4 v2 = float4( pxp / winWidth * 2.0f - 1.0f, (winHeight - pyp) / winHeight * 2.0f - 1.0f, 1.0f, 1.0f );
float4 v3 = float4( pxm / winWidth * 2.0f - 1.0f, (winHeight - pyp) / winHeight * 2.0f - 1.0f, 1.0f, 1.0f );
float4 frustum[ 4 ];
frustum[ 0 ] = ConvertClipToView( v0, uniforms.clipToView );
frustum[ 1 ] = ConvertClipToView( v1, uniforms.clipToView );
frustum[ 2 ] = ConvertClipToView( v2, uniforms.clipToView );
frustum[ 3 ] = ConvertClipToView( v3, uniforms.clipToView );
for (uint i = 0; i < 4; ++i)
{
frustumEqn[ i ] = CreatePlaneEquation( frustum[ i ], frustum[ (i + 1) & 3 ] );
}
}
threadgroup_barrier( mem_flags::mem_threadgroup );
float minZ = FLT_MAX;
float maxZ = 0.0f;
float depth = depthNormalsTexture.read( globalIdx.xy ).x;
uint z = as_type< uint >( depth );
if (depth != 0.0f)
{
/*uint i =*/ atomic_fetch_min_explicit( &ldsZMin, z, memory_order::memory_order_relaxed );
/*uint j =*/ atomic_fetch_max_explicit( &ldsZMax, z, memory_order::memory_order_relaxed );
}
threadgroup_barrier( mem_flags::mem_threadgroup );
uint zMin = atomic_load_explicit( &ldsZMin, memory_order::memory_order_relaxed );
uint zMax = atomic_load_explicit( &ldsZMax, memory_order::memory_order_relaxed );
minZ = as_type< float >( zMax );
maxZ = as_type< float >( zMin );
int numPointLights = uniforms.numLights & 0xFFFFu;
for (int i = 0; i < numPointLights; i += NUM_THREADS_PER_TILE)
{
int il = localIdxFlattened + i;
if (il < numPointLights)
{
float4 center = pointLightBufferCenterAndRadius[ il ];
float radius = center.w;
center.xyz = (uniforms.localToView * float4( center.xyz, 1.0f ) ).xyz;
if (-center.z + minZ < radius && center.z - maxZ < radius)
{
if ((GetSignedDistanceFromPlane( center, frustumEqn[ 0 ] ) < radius) &&
(GetSignedDistanceFromPlane( center, frustumEqn[ 1 ] ) < radius) &&
(GetSignedDistanceFromPlane( center, frustumEqn[ 2 ] ) < radius) &&
(GetSignedDistanceFromPlane( center, frustumEqn[ 3 ] ) < radius))
{
// do a thread-safe increment of the list counter
// and put the index of this light into the list
int dstIdx = atomic_fetch_add_explicit( &ldsLightIdxCounter, 1, memory_order::memory_order_relaxed );
ldsLightIdx[ dstIdx ] = il;
}
}
}
}
threadgroup_barrier( mem_flags::mem_threadgroup );
int numPointLightsInThisTile = atomic_load_explicit( &ldsLightIdxCounter, memory_order::memory_order_relaxed );
// Spot lights.
int numSpotLights = (uniforms.numLights & 0xFFFF0000u) >> 16;
for (int i = 0; i < numSpotLights; i += NUM_THREADS_PER_TILE)
{
int il = localIdxFlattened + i;
if (il < numSpotLights)
{
float4 center = spotLightBufferCenterAndRadius[ il ];
float radius = center.w * 5.0f; // FIXME: Multiply was added, but more clever culling should be done instead.
center.xyz = (uniforms.localToView * float4( center.xyz, 1.0f )).xyz;
if (-center.z + minZ < radius && center.z - maxZ < radius)
{
if ((GetSignedDistanceFromPlane( center, frustumEqn[ 0 ] ) < radius) &&
(GetSignedDistanceFromPlane( center, frustumEqn[ 1 ] ) < radius) &&
(GetSignedDistanceFromPlane( center, frustumEqn[ 2 ] ) < radius) &&
(GetSignedDistanceFromPlane( center, frustumEqn[ 3 ] ) < radius))
{
int dstIdx = atomic_fetch_add_explicit( &ldsLightIdxCounter, 1, memory_order::memory_order_relaxed );
ldsLightIdx[ dstIdx ] = il;
}
}
}
}
threadgroup_barrier( mem_flags::mem_threadgroup );
{ // write back
int startOffset = uniforms.maxNumLightsPerTile * tileIdxFlattened;
for (int i = localIdxFlattened; i < numPointLightsInThisTile; i += NUM_THREADS_PER_TILE)
{
// per-tile list of light indices
perTileLightIndexBufferOut[ startOffset + i ] = ldsLightIdx[ i ];
}
int jMax = atomic_load_explicit( &ldsLightIdxCounter, memory_order::memory_order_relaxed );
for (int j = localIdxFlattened + numPointLightsInThisTile; j < jMax; j += NUM_THREADS_PER_TILE)
{
// per-tile list of light indices
perTileLightIndexBufferOut[ startOffset + j + 1 ] = ldsLightIdx[ j ];
}
if (localIdxFlattened == 0)
{
perTileLightIndexBufferOut[ startOffset + numPointLightsInThisTile ] = LIGHT_INDEX_BUFFER_SENTINEL;
int offs = atomic_load_explicit( &ldsLightIdxCounter, memory_order::memory_order_relaxed );
perTileLightIndexBufferOut[ startOffset + offs + 1 ] = LIGHT_INDEX_BUFFER_SENTINEL;
}
}
}
Я отлаживаю приложение на iPad Pro 10.5 "с использованием iOS 11.4 и Xcode 9.4. Как исправить предупреждение?
Я также попытался изменить тип буфера с constant float4* на constant PointLight& pointLightBufferCenterAndRadius, где PointLight - это struct PointLight { float4 d[ 2048 ]; }, как предложено в Apple Metal WWDC talk .