diff options
| -rw-r--r-- | source/slang/slang-reflection-api.cpp | 304 | ||||
| -rw-r--r-- | tools/gfx/cuda/render-cuda.cpp | 77 | ||||
| -rw-r--r-- | tools/gfx/render-graphics-common.cpp | 2 |
3 files changed, 246 insertions, 137 deletions
diff --git a/source/slang/slang-reflection-api.cpp b/source/slang/slang-reflection-api.cpp index dcfed07b0..1cfb8fecc 100644 --- a/source/slang/slang-reflection-api.cpp +++ b/source/slang/slang-reflection-api.cpp @@ -1172,6 +1172,35 @@ namespace Slang } } + SlangBindingType _calcBindingType( + LayoutResourceKind kind) + { + switch( kind ) + { + default: + return SLANG_BINDING_TYPE_UNKNOWN; + + // Some cases of `LayoutResourceKind` can be mapped + // directly to a `BindingType` because there is only + // one case of types that have that resource kind. + + #define CASE(FROM, TO) \ + case LayoutResourceKind::FROM: return SLANG_BINDING_TYPE_##TO + + CASE(ConstantBuffer, CONSTANT_BUFFER); + CASE(SamplerState, SAMPLER); + CASE(VaryingInput, VARYING_INPUT); + CASE(VaryingOutput, VARYING_OUTPUT); + CASE(ExistentialObjectParam, EXISTENTIAL_VALUE); + CASE(PushConstantBuffer, PUSH_CONSTANT); + CASE(Uniform, INLINE_UNIFORM_DATA); + // TODO: register space + + #undef CASE + } + } + + SlangBindingType _calcBindingType( Slang::TypeLayout* typeLayout, @@ -1193,29 +1222,7 @@ namespace Slang // multiple different kinds of binding, depending on where/how // it is used (e.g., as a varying parameter, a root constant, etc.). // - switch( kind ) - { - default: - return SLANG_BINDING_TYPE_UNKNOWN; - - // Some cases of `LayoutResourceKind` can be mapped - // directly to a `BindingType` because there is only - // one case of types that have that resource kind. - - #define CASE(FROM, TO) \ - case LayoutResourceKind::FROM: return SLANG_BINDING_TYPE_##TO - - CASE(ConstantBuffer, CONSTANT_BUFFER); - CASE(SamplerState, SAMPLER); - CASE(VaryingInput, VARYING_INPUT); - CASE(VaryingOutput, VARYING_OUTPUT); - CASE(ExistentialObjectParam, EXISTENTIAL_VALUE); - CASE(PushConstantBuffer, PUSH_CONSTANT); - CASE(Uniform, INLINE_UNIFORM_DATA); - // TODO: register space - - #undef CASE - } + return _calcBindingType(kind); } static DeclRefType* asInterfaceType(Type* type) @@ -1538,147 +1545,186 @@ namespace Slang } else if(asInterfaceType(typeLayout->type)) { - // An `interface` type should introduce a sub-object range, - // with no concrete descriptor ranges to store its value - // (since we don't know until runtime what type of - // value will be plugged in). + // An `interface` type should introduce a binding range and a matching + // sub-object range. + // + // We currently do *not* allocate any descriptor ranges to represent + // an interface-type field, since the only direct storage required + // is all uniform/ordinary data. // - - LayoutResourceKind kind = LayoutResourceKind::ExistentialObjectParam; - auto count = multiplier; - auto spaceOffset = _calcSpaceOffset(path, kind); - - Int descriptorSetIndex = _findOrAddDescriptorSet(spaceOffset); - auto descriptorSet = m_extendedInfo->m_descriptorSets[descriptorSetIndex]; - TypeLayout::ExtendedInfo::BindingRangeInfo bindingRange; bindingRange.leafTypeLayout = typeLayout; bindingRange.bindingType = SLANG_BINDING_TYPE_EXISTENTIAL_VALUE; bindingRange.count = multiplier; - bindingRange.descriptorSetIndex = descriptorSetIndex; - bindingRange.firstDescriptorRangeIndex = descriptorSet->descriptorRanges.getCount(); - bindingRange.descriptorRangeCount = 1; + bindingRange.descriptorSetIndex = 0; + bindingRange.descriptorRangeCount = 0; + bindingRange.firstDescriptorRangeIndex = 0; TypeLayout::ExtendedInfo::SubObjectRangeInfo subObjectRange; subObjectRange.bindingRangeIndex = m_extendedInfo->m_bindingRanges.getCount(); + // TODO: if we have "pending" layout information that tells us where + // data for the sub-object range has been allocated, then we need + // a way to reference that data here. + m_extendedInfo->m_bindingRanges.add(bindingRange); m_extendedInfo->m_subObjectRanges.add(subObjectRange); } - // TODO: We need to handle `interface` types here, because they are - // another case that introduces a "sub-object" for the purposes of - // application-side allocation. - // - // TODO: There are a few cases of "leaf" fields that might - // still result in multiple descriptors (or at least multiple - // `LayoutResourceKind`s) depending on the target. - // - // For eample, combined texture-sampler types should be treated - // as "leaf" fields for this code (since a portable engine would - // need to abstract over them), but would map to two descriptors - // on targets that don't actually support combining them. else { - Int resourceKindCount = typeLayout->resourceInfos.getCount(); - if(resourceKindCount == 0) + // Here we have the catch-all case that handles "leaf" fields + // that should never introduce a sub-object range, but might + // need to introduce a binding range and descriptor ranges. + // + // First, we want to determine what type of binding this + // leaf field should map to, if any. We being by querying + // the type itself, since there are many distinct descriptor + // types for textures/buffers that can only be determined + // by type, rather than by a `LayoutResourceKind`. + // + auto bindingType = _calcResourceBindingType(typeLayout); + + // It is possible that the type alone isn't enough to tell + // us a specific binding type, at which point we need to + // start looking at the actual resources the type layout + // consumes. + // + if(bindingType == SLANG_BINDING_TYPE_UNKNOWN) { - // This is a field that consumes no resources, and as - // such does not need a binding or descriptor ranges - // allocated for it. + // We will search through all the resource kinds that + // the type layout consumes, to see if we can find + // one that indicates a binding type we actually + // want to reflect. // - return; - } - else - { - // `resourceKindCount` should be 1 in most cases. - // However if this field is a buffer of existential type, - // the resourceInfo array will contain additional ExistentialTypeParam - // and ExistentialObjectParam entries. These entries doesn't affect the - // logic here, so we only need to care about the first entry. - auto& resInfo = typeLayout->resourceInfos[0]; - LayoutResourceKind kind = resInfo.kind; - - auto bindingType = _calcBindingType(typeLayout, kind); - if(bindingType == SLANG_BINDING_TYPE_INLINE_UNIFORM_DATA) + for( auto resInfo : typeLayout->resourceInfos ) { - // We do not consider uniform resource usage - // in the ranges we compute. + auto kind = resInfo.kind; + if(kind == LayoutResourceKind::Uniform) + continue; + + auto kindBindingType = _calcBindingType(kind); + if(kindBindingType == SLANG_BINDING_TYPE_UNKNOWN) + continue; + + // If we find a relevant binding type based on + // one of the resource kinds that are consumed, + // then we immediately stop the search and use + // the first one found (whether or not later + // entries might also provide something relevant). // - // TODO: We may need to revise that rule for types that - // represent resources, even when one or more targets - // map those resource types to ordinary/uniform data. - // - return; + bindingType = kindBindingType; + break; } + } - // This leaf field will map to a single binding range and, - // if it is appropriate, a single descriptor range. - // - auto count = resInfo.count * multiplier; - auto indexOffset = _calcIndexOffset(path, kind); - auto spaceOffset = _calcSpaceOffset(path, kind); + // After we've tried to determine a binding type, if + // we have nothing to go on then we don't want to add + // a binding range. + // + if(bindingType == SLANG_BINDING_TYPE_UNKNOWN) + return; + + // We now know that the leaf field will map to a single binding range, + // and zero or more descriptor ranges. + // + TypeLayout::ExtendedInfo::BindingRangeInfo bindingRange; + bindingRange.leafTypeLayout = typeLayout; + bindingRange.bindingType = bindingType; + bindingRange.count = multiplier; + bindingRange.descriptorSetIndex = 0; + bindingRange.firstDescriptorRangeIndex = 0; + bindingRange.descriptorRangeCount = 0; - Int descriptorSetIndex = -1; - Int firstDescriptorIndex = 0; - RefPtr<TypeLayout::ExtendedInfo::DescriptorSetInfo> descriptorSet; + // We will associate the binding range with a specific descriptor + // set on demand *if* we discover that it shold contain any + // descriptor ranges. + // + RefPtr<TypeLayout::ExtendedInfo::DescriptorSetInfo> descriptorSet; + + + // We will add a descriptor range for each relevant resource kind + // that the type layout consumes. + // + for(auto resInfo : typeLayout->resourceInfos) + { + auto kind = resInfo.kind; switch( kind ) { + default: + break; + + + // There are many resource kinds that we do not want + // to expose as descriptor ranges simply because they + // do not actually allocate descriptors on our target + // APIs. + // + // Notably included here are uniform/ordinary data and + // varying input/output (including the ray-tracing cases). + // + // It is worth noting that we *do* allow root/push-constant + // ranges to be reflected as "descriptor" ranges here, + // despite the fact that they are not descriptor-bound + // under D3D12/Vulkan. + // + // In practice, even with us filtering out some cases here, + // an application/renderer layer will need to filter/translate + // or descriptor ranges into API-specific ones, and a one-to-one + // mapping should not be assumed. + // + // TODO: Make some clear decisions about what should and should + // not appear here. + // + case LayoutResourceKind::Uniform: case LayoutResourceKind::RegisterSpace: case LayoutResourceKind::VaryingInput: case LayoutResourceKind::VaryingOutput: case LayoutResourceKind::HitAttributes: case LayoutResourceKind::RayPayload: - // Resource kinds that represent "varying" input/output - // do not manifest as entries in API descriptor tables. - // - // TODO: Neither do root constants, if we are being - // precise. This API really needs to carefully match - // the semantics of the target platform/API in terms - // of what things are descriptor-bound and which are - // not, so that a user can easily allocate the platform-specific - // descriptor sets using this info. - // - // (That said, we are purposefully *not* breaking apart - // samplers and SRV/UAV/CBV stuff for our D3D reflection - // of descriptor sets. It seems like the policy here - // really requires careful thought) - // - // TODO: Maybe the best answer is to leave decomposition - // of stuff into descriptor sets up to the application - // layer? This is especially true if a common case would - // be an application that doesn't support arbitrary manual - // binding of parameters to register/spaces. - // - break; + continue; + } - default: - { - TypeLayout::ExtendedInfo::DescriptorRangeInfo descriptorRange; - descriptorRange.kind = kind; - descriptorRange.bindingType = bindingType; - descriptorRange.count = count; - descriptorRange.indexOffset = indexOffset; + // We will prefer to use a binding type derived from the specific + // resource kind, but will fall back to information from the + // type layout when that is not available. + // + // TODO: This logic probably needs a bit more work to handle + // the case of a combined texture-sampler field that is being + // compiled for an API with separate textures and samplers. + // + auto kindBindingType = _calcBindingType(kind); + if( kindBindingType == SLANG_BINDING_TYPE_UNKNOWN ) + { + kindBindingType = bindingType; + } - descriptorSetIndex = _findOrAddDescriptorSet(spaceOffset); - descriptorSet = m_extendedInfo->m_descriptorSets[descriptorSetIndex]; + // We now expect to allocate a descriptor range for this + // `resInfo` representing resouce usage. + // + auto count = resInfo.count * multiplier; + auto indexOffset = _calcIndexOffset(path, kind); + auto spaceOffset = _calcSpaceOffset(path, kind); - firstDescriptorIndex = descriptorSet->descriptorRanges.getCount(); - descriptorSet->descriptorRanges.add(descriptorRange); - } - break; - } + TypeLayout::ExtendedInfo::DescriptorRangeInfo descriptorRange; + descriptorRange.kind = kind; + descriptorRange.bindingType = kindBindingType; + descriptorRange.count = count; + descriptorRange.indexOffset = indexOffset; + if(!descriptorSet) + { + Int descriptorSetIndex = _findOrAddDescriptorSet(spaceOffset); + descriptorSet = m_extendedInfo->m_descriptorSets[descriptorSetIndex]; - TypeLayout::ExtendedInfo::BindingRangeInfo bindingRange; - bindingRange.leafTypeLayout = typeLayout; - bindingRange.bindingType = _calcBindingType(typeLayout, kind); - bindingRange.count = count; - bindingRange.descriptorSetIndex = descriptorSetIndex; - bindingRange.firstDescriptorRangeIndex = firstDescriptorIndex; - bindingRange.descriptorRangeCount = 1; + bindingRange.descriptorSetIndex = descriptorSetIndex; + bindingRange.firstDescriptorRangeIndex = descriptorSet->descriptorRanges.getCount(); + } - m_extendedInfo->m_bindingRanges.add(bindingRange); + descriptorSet->descriptorRanges.add(descriptorRange); + bindingRange.descriptorRangeCount++; } + + m_extendedInfo->m_bindingRanges.add(bindingRange); } } }; diff --git a/tools/gfx/cuda/render-cuda.cpp b/tools/gfx/cuda/render-cuda.cpp index 4f87bdfc9..906d553cb 100644 --- a/tools/gfx/cuda/render-cuda.cpp +++ b/tools/gfx/cuda/render-cuda.cpp @@ -250,6 +250,20 @@ public: slang::BindingType bindingType; Index count; Index baseIndex; // Flat index for sub-ojects + + // TODO: The `uniformOffset` field should be removed, + // since it cannot be supported by the Slang reflection + // API once we fix some design issues. + // + // It is only being used today for pre-allocation of sub-objects + // for constant buffers and parameter blocks (which should be + // deprecated/removed anyway). + // + // Note: We would need to bring this field back, plus + // a lot of other complexity, if we ever want to support + // setting of resources/buffers directly by a binding + // range index and array index. + // Index uniformOffset; // Uniform offset for a resource typed field. }; @@ -310,6 +324,14 @@ public: SlangInt rangeIndexInDescriptorSet = m_elementTypeLayout->getBindingRangeFirstDescriptorRangeIndex(r); + // TODO: This logic assumes that for any binding range that might consume + // multiple kinds of resources, the descriptor range for its uniform + // usage will be the first one in the range. + // + // We need to decide whether that assumption is one we intend to support + // applications making, or whether they should be forced to perform a + // linear search over the descriptor ranges for a specific binding range. + // auto uniformOffset = m_elementTypeLayout->getDescriptorSetDescriptorRangeIndexOffset( descriptorSetIndex, rangeIndexInDescriptorSet); @@ -464,6 +486,17 @@ public: return SLANG_OK; } virtual SLANG_NO_THROW Result SLANG_MCALL + setDeviceData(size_t offset, void* data, size_t size) + { + size = Math::Min(size, bufferResource->getDesc()->sizeInBytes - offset); + SLANG_CUDA_RETURN_ON_FAIL(cudaMemcpy( + (uint8_t*)bufferResource->m_cudaMemory + offset, + data, + size, + cudaMemcpyHostToDevice)); + return SLANG_OK; + } + virtual SLANG_NO_THROW Result SLANG_MCALL getObject(ShaderOffset const& offset, IShaderObject** object) { auto subObjectIndex = @@ -480,17 +513,34 @@ public: virtual SLANG_NO_THROW Result SLANG_MCALL setObject(ShaderOffset const& offset, IShaderObject* object) { - auto subObjectIndex = - getLayout()->m_bindingRanges[offset.bindingRangeIndex].baseIndex + offset.bindingArrayIndex; - SLANG_ASSERT( - offset.uniformOffset == - getLayout()->m_bindingRanges[offset.bindingRangeIndex].uniformOffset + - offset.bindingArrayIndex * sizeof(void*)); + auto layout = getLayout(); + SLANG_ASSERT(offset.bindingRangeIndex >= 0); + SLANG_ASSERT(offset.bindingRangeIndex < layout->m_bindingRanges.getCount()); + auto& bindingRange = layout->m_bindingRanges[offset.bindingRangeIndex]; + + auto subObjectIndex = bindingRange.baseIndex + offset.bindingArrayIndex; auto cudaObject = dynamic_cast<CUDAShaderObject*>(object); if (subObjectIndex >= objects.getCount()) objects.setCount(subObjectIndex + 1); objects[subObjectIndex] = cudaObject; - return setData(offset, &cudaObject->bufferResource->m_cudaMemory, sizeof(void*)); + + switch( bindingRange.bindingType ) + { + default: + SLANG_RETURN_ON_FAIL(setData(offset, &cudaObject->bufferResource->m_cudaMemory, sizeof(void*))); + break; + + case slang::BindingType::ExistentialValue: + // TODO: handle the "does it fit" logic + { + auto valueSize = cudaObject->m_layout->getElementTypeLayout()->getSize(); + auto valueOffset = 16; + SLANG_RETURN_ON_FAIL(setDeviceData(offset.uniformOffset + valueOffset, cudaObject->getBuffer(), valueSize)); + } + break; + } + + return SLANG_OK; } virtual SLANG_NO_THROW Result SLANG_MCALL setResource(ShaderOffset const& offset, IResourceView* resourceView) @@ -615,6 +665,19 @@ public: return SLANG_OK; } + virtual SLANG_NO_THROW Result SLANG_MCALL + setDeviceData(size_t offset, void* data, size_t size) + { + size = Math::Min(size, uniformBufferSize - offset); + SLANG_CUDA_RETURN_ON_FAIL(cudaMemcpy( + (uint8_t*)hostBuffer + offset, + data, + size, + cudaMemcpyDeviceToHost)); + return SLANG_OK; + } + + virtual SLANG_NO_THROW void* SLANG_MCALL getBuffer() override { return hostBuffer; diff --git a/tools/gfx/render-graphics-common.cpp b/tools/gfx/render-graphics-common.cpp index 5f083538d..9e95df95c 100644 --- a/tools/gfx/render-graphics-common.cpp +++ b/tools/gfx/render-graphics-common.cpp @@ -153,6 +153,7 @@ public: switch (slangBindingType) { case slang::BindingType::ExistentialValue: + case slang::BindingType::InlineUniformData: continue; default: break; @@ -172,7 +173,6 @@ public: auto category = typeLayout->getDescriptorSetDescriptorRangeCategory(s, r); descriptorRangeDesc.binding += varLayout->getOffset(category); } - descriptorSetInfo->slotRangeDescs.add(descriptorRangeDesc); } } |