From a72727a1734ab525766f1738aeea54194d76e7c5 Mon Sep 17 00:00:00 2001 From: Philip Rebohle Date: Thu, 28 Dec 2017 16:03:17 +0100 Subject: [PATCH] [dxbc] Added support for structured and raw buffers --- src/dxbc/dxbc_compiler.cpp | 429 ++++++++++++++++++++++++++++++++++--- src/dxbc/dxbc_compiler.h | 69 +++++- src/dxbc/dxbc_decoder.h | 32 ++- src/dxbc/dxbc_defs.cpp | 192 ++++++++++++++--- src/dxbc/dxbc_defs.h | 3 + src/dxbc/dxbc_enums.h | 7 + src/spirv/spirv_module.cpp | 34 +++ src/spirv/spirv_module.h | 12 ++ 8 files changed, 713 insertions(+), 65 deletions(-) diff --git a/src/dxbc/dxbc_compiler.cpp b/src/dxbc/dxbc_compiler.cpp index 4c3f44bbb..484c9aeb7 100644 --- a/src/dxbc/dxbc_compiler.cpp +++ b/src/dxbc/dxbc_compiler.cpp @@ -66,6 +66,15 @@ namespace dxvk { case DxbcInstClass::GeometryEmit: return this->emitGeometryEmit(ins); + case DxbcInstClass::Atomic: + return this->emitAtomic(ins); + + case DxbcInstClass::BufferLoad: + return this->emitBufferLoad(ins); + + case DxbcInstClass::BufferStore: + return this->emitBufferStore(ins); + case DxbcInstClass::TextureQuery: return this->emitTextureQuery(ins); @@ -180,9 +189,20 @@ namespace dxvk { case DxbcOpcode::DclSampler: return this->emitDclSampler(ins); +// case DxbcOpcode::DclUavTyped: case DxbcOpcode::DclResource: - return this->emitDclResource(ins); + return this->emitDclResourceTyped(ins); + + case DxbcOpcode::DclUavRaw: + case DxbcOpcode::DclResourceRaw: + case DxbcOpcode::DclUavStructured: + case DxbcOpcode::DclResourceStructured: + return this->emitDclResourceRawStructured(ins); + case DxbcOpcode::DclThreadGroupSharedMemoryRaw: + case DxbcOpcode::DclThreadGroupSharedMemoryStructured: + return this->emitDclThreadGroupSharedMemory(ins); + case DxbcOpcode::DclGsInputPrimitive: return this->emitDclGsInputPrimitive(ins); @@ -191,6 +211,9 @@ namespace dxvk { case DxbcOpcode::DclMaxOutputVertexCount: return this->emitDclMaxOutputVertexCount(ins); + + case DxbcOpcode::DclThreadGroup: + return this->emitDclThreadGroup(ins); default: Logger::warn( @@ -471,7 +494,6 @@ namespace dxvk { void DxbcCompiler::emitDclSampler(const DxbcShaderInstruction& ins) { // dclSampler takes one operand: // (dst0) The sampler register to declare - // TODO implement sampler mode (default / comparison / mono) const uint32_t samplerId = ins.dst[0].idx[0].offset; // The sampler type is opaque, but we still have to @@ -504,10 +526,10 @@ namespace dxvk { } - void DxbcCompiler::emitDclResource(const DxbcShaderInstruction& ins) { + void DxbcCompiler::emitDclResourceTyped(const DxbcShaderInstruction& ins) { // dclResource takes two operands: - // (dst0) The resource register ID - // (imm0) The resource return type + // (dst0) The resource register ID + // (imm0) The resource return type const uint32_t registerId = ins.dst[0].idx[0].offset; // Defines the type of the resource (texture2D, ...) @@ -588,12 +610,17 @@ namespace dxvk { m_module.setDebugName(varId, str::format("t", registerId).c_str()); - m_textures.at(registerId).imageInfo = typeInfo; - m_textures.at(registerId).varId = varId; - m_textures.at(registerId).sampledType = sampledType; - m_textures.at(registerId).sampledTypeId = sampledTypeId; - m_textures.at(registerId).colorTypeId = colorTypeId; - m_textures.at(registerId).depthTypeId = depthTypeId; + DxbcShaderResource res; + res.type = DxbcResourceType::Typed; + res.imageInfo = typeInfo; + res.varId = varId; + res.sampledType = sampledType; + res.sampledTypeId = sampledTypeId; + res.colorTypeId = colorTypeId; + res.depthTypeId = depthTypeId; + res.structStride = 0; + + m_textures.at(registerId) = res; // Compute the DXVK binding slot index for the resource. // D3D11 needs to bind the actual resource to this slot. @@ -613,6 +640,105 @@ namespace dxvk { } + void DxbcCompiler::emitDclResourceRawStructured(const DxbcShaderInstruction& ins) { + // dcl_resource_raw and dcl_uav_raw take one argument: + // (dst0) The resource register ID + // dcl_resource_structured and dcl_uav_structured take two arguments: + // (dst0) The resource register ID + // (imm0) Structure stride, in bytes + const uint32_t registerId = ins.dst[0].idx[0].offset; + + const bool isUav = ins.op == DxbcOpcode::DclUavRaw + || ins.op == DxbcOpcode::DclUavStructured; + + const bool isStructured = ins.op == DxbcOpcode::DclUavStructured + || ins.op == DxbcOpcode::DclResourceStructured; + + // Structured and raw buffers are represented as + // texel buffers consisting of 32-bit integers. + m_module.enableCapability(spv::CapabilityImageBuffer); + + const DxbcScalarType sampledType = DxbcScalarType::Uint32; + const uint32_t sampledTypeId = getScalarTypeId(sampledType); + + const DxbcImageInfo typeInfo = { spv::DimBuffer, 0, 0, isUav ? 2u : 1u }; + + // Declare the resource type + const uint32_t resTypeId = m_module.defImageType(sampledTypeId, + typeInfo.dim, 0, typeInfo.array, typeInfo.ms, typeInfo.sampled, + spv::ImageFormatR32ui); + + const uint32_t varId = m_module.newVar( + m_module.defPointerType(resTypeId, spv::StorageClassUniformConstant), + spv::StorageClassUniformConstant); + + m_module.setDebugName(varId, + str::format(isUav ? "u" : "t", registerId).c_str()); + + // Write back resource info + const DxbcResourceType resType = isStructured + ? DxbcResourceType::Structured + : DxbcResourceType::Raw; + + const uint32_t resStride = isStructured + ? ins.imm[0].u32 + : 0; + + if (isUav) { + DxbcUav uav; + uav.type = resType; + uav.imageInfo = typeInfo; + uav.varId = varId; + uav.sampledType = sampledType; + uav.sampledTypeId = sampledTypeId; + uav.imageTypeId = resTypeId; + uav.structStride = resStride; + m_uavs.at(registerId) = uav; + } else { + DxbcShaderResource res; + res.type = resType; + res.imageInfo = typeInfo; + res.varId = varId; + res.sampledType = sampledType; + res.sampledTypeId = sampledTypeId; + res.colorTypeId = resTypeId; + res.depthTypeId = resTypeId; + res.structStride = resStride; + m_textures.at(registerId) = res; + } + + // Compute the DXVK binding slot index for the resource. + const uint32_t bindingId = computeResourceSlotId( + m_version.type(), isUav + ? DxbcBindingType::UnorderedAccessView + : DxbcBindingType::ShaderResource, + registerId); + + m_module.decorateDescriptorSet(varId, 0); + m_module.decorateBinding(varId, bindingId); + + // Store descriptor info for the shader interface + DxvkResourceSlot resource; + resource.slot = bindingId; + resource.type = isUav + ? VK_DESCRIPTOR_TYPE_STORAGE_TEXEL_BUFFER + : VK_DESCRIPTOR_TYPE_UNIFORM_TEXEL_BUFFER; + m_resourceSlots.push_back(resource); + } + + + void DxbcCompiler::emitDclThreadGroupSharedMemory(const DxbcShaderInstruction& ins) { + // dcl_tgsm_raw takes two arguments: + // (dst0) The resource register ID + // (imm0) Block size, in DWORDs + // dcl_tgsm_structured takes three arguments: + // (dst0) The resource register ID + // (imm0) Structure stride, in bytes + // (imm0) Structure count + Logger::err("DxbcCompiler: emitDclThreadGroupSharedMemory not implemented"); + } + + void DxbcCompiler::emitDclGsInputPrimitive(const DxbcShaderInstruction& ins) { // The input primitive type is stored within in the // control bits of the opcode token. In SPIR-V, we @@ -665,6 +791,16 @@ namespace dxvk { } + void DxbcCompiler::emitDclThreadGroup(const DxbcShaderInstruction& ins) { + // dcl_thread_group has three operands: + // (imm0) Number of threads in X dimension + // (imm1) Number of threads in Y dimension + // (imm2) Number of threads in Z dimension + m_module.setLocalSize(m_entryPointId, + ins.imm[0].u32, ins.imm[1].u32, ins.imm[2].u32); + } + + void DxbcCompiler::emitDclImmediateConstantBuffer(const DxbcShaderInstruction& ins) { if (m_immConstBuf != 0) throw DxvkError("DxbcCompiler: Immediate constant buffer already declared"); @@ -1318,11 +1454,55 @@ namespace dxvk { } + void DxbcCompiler::emitAtomic(const DxbcShaderInstruction& ins) { + Logger::err("DxbcCompiler: emitAtomic not implemented"); + } + + + void DxbcCompiler::emitBufferLoad(const DxbcShaderInstruction& ins) { + // ld_raw takes three arguments: + // (dst0) Destination register + // (src0) Byte offset + // (src1) Source register + // ld_structured takes four arguments: + // (dst0) Destination register + // (src0) Structure index + // (src1) Byte offset + // (src2) Source register + const bool isStructured = ins.op == DxbcOpcode::LdStructured; + + // Source register. The exact way we access + // the data depends on the register type. + const DxbcRegister& dstReg = ins.dst[0]; + const DxbcRegister& srcReg = isStructured ? ins.src[2] : ins.src[1]; + + // Retrieve common info about the buffer + const DxbcBufferInfo bufferInfo = getBufferInfo(srcReg); + + // Compute element index + const DxbcRegisterValue elementIndex = isStructured + ? emitCalcBufferIndexStructured( + emitRegisterLoad(ins.src[0], DxbcRegMask(true, false, false, false)), + emitRegisterLoad(ins.src[1], DxbcRegMask(true, false, false, false)), + bufferInfo.stride) + : emitCalcBufferIndexRaw( + emitRegisterLoad(ins.src[0], DxbcRegMask(true, false, false, false))); + + emitRegisterStore(dstReg, + emitRawBufferLoad(srcReg, elementIndex, dstReg.mask)); + } + + + void DxbcCompiler::emitBufferStore(const DxbcShaderInstruction& ins) { + Logger::err("DxbcCompiler: emitBufferStore not implemented"); + } + + void DxbcCompiler::emitTextureQuery(const DxbcShaderInstruction& ins) { // resinfo has three operands: - // (dst0) The destination register - // (src0) Resource LOD to query - // (src1) Resource to query + // (dst0) The destination register + // (src0) Resource LOD to query + // (src1) Resource to query const DxbcResinfoType resinfoType = ins.controls.resinfoType; if (ins.src[1].type != DxbcOperandType::Resource) { @@ -1449,9 +1629,9 @@ namespace dxvk { void DxbcCompiler::emitTextureFetch(const DxbcShaderInstruction& ins) { // ld has three operands: - // (dst0) The destination register - // (src0) Source address - // (src1) Source texture + // (dst0) The destination register + // (src0) Source address + // (src1) Source texture const uint32_t textureId = ins.src[1].idx[0].offset; // Image type, which stores the image dimensions etc. @@ -1526,10 +1706,10 @@ namespace dxvk { // TODO support remaining sample ops // All sample instructions have at least these operands: - // (dst0) The destination register - // (src0) Texture coordinates - // (src1) The texture itself - // (src2) The sampler object + // (dst0) The destination register + // (src0) Texture coordinates + // (src1) The texture itself + // (src2) The sampler object const DxbcRegister& texCoordReg = ins.src[0]; const DxbcRegister& textureReg = ins.src[1]; const DxbcRegister& samplerReg = ins.src[2]; @@ -1589,7 +1769,7 @@ namespace dxvk { : DxbcRegisterValue(); // Determine the sampled image type based on the opcode. - // FIXME while this is in line what the officla glsl compiler + // FIXME while this is in line what the offical glsl compiler // does, this might actually violate the SPIR-V specification. const uint32_t sampledImageType = isDepthCompare ? m_module.defSampledImageType(m_textures.at(textureId).depthTypeId) @@ -2299,6 +2479,126 @@ namespace dxvk { } + DxbcRegisterValue DxbcCompiler::emitRawBufferLoad( + const DxbcRegister& operand, + DxbcRegisterValue elementIndex, + DxbcRegMask writeMask) { + const DxbcBufferInfo bufferInfo = getBufferInfo(operand); + + // Shared memory is the only type of buffer that + // is not accessed through a texel buffer view + const bool isTgsm = operand.type == DxbcOperandType::ThreadGroupSharedMemory; + + const uint32_t bufferId = isTgsm + ? 0 : m_module.opLoad(bufferInfo.typeId, bufferInfo.varId); + + // Since all data is represented as a sequence of 32-bit + // integers, we have to load each component individually. + std::array componentIds = { 0, 0, 0, 0 }; + std::array swizzleIds = { 0, 0, 0, 0 }; + + uint32_t componentIndex = 0; + + const uint32_t vectorTypeId = getVectorTypeId({ DxbcScalarType::Uint32, 4 }); + const uint32_t scalarTypeId = getVectorTypeId({ DxbcScalarType::Uint32, 1 }); + + for (uint32_t i = 0; i < 4; i++) { + // We'll apply both the write mask and the source operand swizzle + // immediately. Unused components are not loaded, and the scalar + // IDs are written to the array in the order they are requested. + if (writeMask[i]) { + const uint32_t swizzleIndex = operand.swizzle[i]; + + if (componentIds[swizzleIndex] == 0) { + // Add the component offset to the element index + const uint32_t elementIndexAdjusted = swizzleIndex != 0 + ? m_module.opIAdd(getVectorTypeId(elementIndex.type), + elementIndex.id, m_module.consti32(swizzleIndex)) + : elementIndex.id; + + // Load requested component from the buffer + componentIds[swizzleIndex] = [&] { + const uint32_t zero = 0; + + switch (operand.type) { + case DxbcOperandType::Resource: + return m_module.opCompositeExtract(scalarTypeId, + m_module.opImageFetch(vectorTypeId, + bufferId, elementIndexAdjusted, + SpirvImageOperands()), 1, &zero); + + case DxbcOperandType::UnorderedAccessView: + return m_module.opImageRead( + scalarTypeId, bufferId, elementIndexAdjusted, + SpirvImageOperands()); + + default: + throw DxvkError("DxbcCompiler: Invalid operand type for strucured/raw load"); + } + }(); + } + + // Append current component to the list of scalar IDs. + // These will be used to construct the resulting vector. + swizzleIds[componentIndex++] = componentIds[swizzleIndex]; + } + } + + // Create result vector + DxbcRegisterValue result; + result.type.ctype = DxbcScalarType::Uint32; + result.type.ccount = writeMask.setCount(); + + result.id = result.type.ccount > 1 + ? m_module.opCompositeConstruct(getVectorTypeId(result.type), + result.type.ccount, swizzleIds.data()) + : swizzleIds[0]; + return result; + } + + + void DxbcCompiler::emitRawBufferStore( + const DxbcRegister& operand, + DxbcRegisterValue elementIndex, + DxbcRegisterValue value) { + const DxbcBufferInfo bufferInfo = getBufferInfo(operand); + // TODO implement + } + + + DxbcRegisterValue DxbcCompiler::emitCalcBufferIndexStructured( + DxbcRegisterValue structId, + DxbcRegisterValue structOffset, + uint32_t structStride) { + DxbcRegisterValue result; + result.type.ctype = DxbcScalarType::Sint32; + result.type.ccount = 1; + + const uint32_t typeId = getVectorTypeId(result.type); + + result.id = m_module.opShiftRightArithmetic(typeId, + m_module.opIAdd(typeId, + m_module.opIMul(typeId, structId.id, + m_module.consti32(structStride)), + structOffset.id), + m_module.consti32(2)); + return result; + } + + + DxbcRegisterValue DxbcCompiler::emitCalcBufferIndexRaw( + DxbcRegisterValue byteOffset) { + DxbcRegisterValue result; + result.type.ctype = DxbcScalarType::Sint32; + result.type.ccount = 1; + result.id = m_module.opShiftRightArithmetic( + getVectorTypeId(result.type), + byteOffset.id, + m_module.consti32(2)); + return result; + } + + DxbcRegisterValue DxbcCompiler::emitIndexLoad( DxbcRegIndex index) { if (index.relReg != nullptr) { @@ -2448,8 +2748,9 @@ namespace dxvk { const DxbcRegisterValue value = [&] { switch (m_version.type()) { - case DxbcProgramType::VertexShader: return emitVsSystemValueLoad(map.sv, map.regMask); - case DxbcProgramType::PixelShader: return emitPsSystemValueLoad(map.sv, map.regMask); + case DxbcProgramType::VertexShader: return emitVsSystemValueLoad(map.sv, map.regMask); + case DxbcProgramType::PixelShader: return emitPsSystemValueLoad(map.sv, map.regMask); + case DxbcProgramType::ComputeShader: return emitCsSystemValueLoad(map.sv, map.regMask); default: throw DxvkError(str::format("DxbcCompiler: Unexpected stage: ", m_version.type())); } }(); @@ -2623,6 +2924,17 @@ namespace dxvk { } + DxbcRegisterValue DxbcCompiler::emitCsSystemValueLoad( + DxbcSystemValue sv, + DxbcRegMask mask) { + switch (sv) { + default: + throw DxvkError(str::format( + "DxbcCompiler: Unhandled CS SV input: ", sv)); + } + } + + void DxbcCompiler::emitVsSystemValueStore( DxbcSystemValue sv, DxbcRegMask mask, @@ -2711,7 +3023,42 @@ namespace dxvk { void DxbcCompiler::emitCsInitBuiltins() { - // TODO implement + m_cs.builtinGlobalInvocationId = emitNewBuiltinVariable({ + { DxbcScalarType::Uint32, 3, 0 }, + spv::StorageClassInput }, + spv::BuiltInGlobalInvocationId, + "cs_global_invocation_id"); + + m_cs.builtinLocalInvocationId = emitNewBuiltinVariable({ + { DxbcScalarType::Uint32, 3, 0 }, + spv::StorageClassInput }, + spv::BuiltInLocalInvocationId, + "cs_local_invocation_id"); + + // FIXME Vulkan might not support this? not documented + m_cs.builtinLocalInvocationIndex = emitNewBuiltinVariable({ + { DxbcScalarType::Uint32, 1, 0 }, + spv::StorageClassInput }, + spv::BuiltInLocalInvocationIndex, + "cs_local_invocation_index"); + + m_cs.builtinWorkgroupId = emitNewBuiltinVariable({ + { DxbcScalarType::Uint32, 3, 0 }, + spv::StorageClassInput }, + spv::BuiltInWorkgroupId, + "cs_workgroup_id"); + + m_cs.builtinWorkgroupSize = emitNewBuiltinVariable({ + { DxbcScalarType::Uint32, 3, 0 }, + spv::StorageClassInput }, + spv::BuiltInWorkgroupSize, + "cs_workgroup_size"); + + m_cs.builtinWorkgroupCount = emitNewBuiltinVariable({ + { DxbcScalarType::Uint32, 3, 0 }, + spv::StorageClassInput }, + spv::BuiltInNumWorkgroups, + "cs_workgroup_count"); } @@ -2966,6 +3313,38 @@ namespace dxvk { } + DxbcBufferInfo DxbcCompiler::getBufferInfo(const DxbcRegister& reg) { + const uint32_t registerId = reg.idx[0].offset; + + switch (reg.type) { + case DxbcOperandType::Resource: { + DxbcBufferInfo result; + result.type = m_textures.at(registerId).type; + result.typeId = m_textures.at(registerId).colorTypeId; + result.varId = m_textures.at(registerId).varId; + result.stride = m_textures.at(registerId).structStride; + return result; + } break; + + case DxbcOperandType::UnorderedAccessView: { + DxbcBufferInfo result; + result.type = m_uavs.at(registerId).type; + result.typeId = m_uavs.at(registerId).imageTypeId; + result.varId = m_uavs.at(registerId).varId; + result.stride = m_uavs.at(registerId).structStride; + return result; + } break; + + // TODO implement +// case DxbcOperandType::ThreadGroupSharedMemory: { +// } break; + + default: + throw DxvkError(str::format("DxbcCompiler: Invalid operand type for buffer: ", reg.type)); + } + } + + uint32_t DxbcCompiler::getScalarTypeId(DxbcScalarType type) { switch (type) { case DxbcScalarType::Uint32: return m_module.defIntType(32, 0); diff --git a/src/dxbc/dxbc_compiler.h b/src/dxbc/dxbc_compiler.h index fc25eba86..186fdb126 100644 --- a/src/dxbc/dxbc_compiler.h +++ b/src/dxbc/dxbc_compiler.h @@ -128,6 +128,13 @@ namespace dxvk { */ struct DxbcCompilerCsPart { uint32_t functionId = 0; + + uint32_t builtinGlobalInvocationId = 0; + uint32_t builtinLocalInvocationId = 0; + uint32_t builtinLocalInvocationIndex = 0; + uint32_t builtinWorkgroupId = 0; + uint32_t builtinWorkgroupSize = 0; + uint32_t builtinWorkgroupCount = 0; }; @@ -162,6 +169,14 @@ namespace dxvk { }; + struct DxbcBufferInfo { + DxbcResourceType type; + uint32_t typeId; + uint32_t varId; + uint32_t stride; + }; + + /** * \brief DXBC to SPIR-V shader compiler * @@ -230,6 +245,7 @@ namespace dxvk { std::array m_constantBuffers; std::array m_samplers; std::array m_textures; + std::array m_uavs; /////////////////////////////////////////////// // Control flow information. Stores labels for @@ -302,7 +318,13 @@ namespace dxvk { void emitDclSampler( const DxbcShaderInstruction& ins); - void emitDclResource( + void emitDclResourceTyped( + const DxbcShaderInstruction& ins); + + void emitDclResourceRawStructured( + const DxbcShaderInstruction& ins); + + void emitDclThreadGroupSharedMemory( const DxbcShaderInstruction& ins); void emitDclGsInputPrimitive( @@ -314,6 +336,9 @@ namespace dxvk { void emitDclMaxOutputVertexCount( const DxbcShaderInstruction& ins); + void emitDclThreadGroup( + const DxbcShaderInstruction& ins); + //////////////////////// // Custom data handlers void emitDclImmediateConstantBuffer( @@ -354,6 +379,15 @@ namespace dxvk { void emitGeometryEmit( const DxbcShaderInstruction& ins); + void emitAtomic( + const DxbcShaderInstruction& ins); + + void emitBufferLoad( + const DxbcShaderInstruction& ins); + + void emitBufferStore( + const DxbcShaderInstruction& ins); + void emitTextureQuery( const DxbcShaderInstruction& ins); @@ -457,6 +491,28 @@ namespace dxvk { DxbcRegisterPointer emitGetOperandPtr( const DxbcRegister& operand); + /////////////////////////////// + // Resource load/store methods + DxbcRegisterValue emitRawBufferLoad( + const DxbcRegister& operand, + DxbcRegisterValue elementIndex, + DxbcRegMask writeMask); + + void emitRawBufferStore( + const DxbcRegister& operand, + DxbcRegisterValue elementIndex, + DxbcRegisterValue value); + + //////////////////////////////////// + // Buffer index calculation methods + DxbcRegisterValue emitCalcBufferIndexStructured( + DxbcRegisterValue structId, + DxbcRegisterValue structOffset, + uint32_t structStride); + + DxbcRegisterValue emitCalcBufferIndexRaw( + DxbcRegisterValue byteOffset); + ////////////////////////////// // Operand load/store methods DxbcRegisterValue emitIndexLoad( @@ -500,6 +556,10 @@ namespace dxvk { DxbcSystemValue sv, DxbcRegMask mask); + DxbcRegisterValue emitCsSystemValueLoad( + DxbcSystemValue sv, + DxbcRegMask mask); + /////////////////////////////////////////// // System value store methods (per shader) void emitVsSystemValueStore( @@ -552,10 +612,13 @@ namespace dxvk { spv::BuiltIn builtIn, const char* name); - ///////////////////////////////////// - // Control flow block search methods + //////////////// + // Misc methods DxbcCfgBlock* cfgFindLoopBlock(); + DxbcBufferInfo getBufferInfo( + const DxbcRegister& reg); + /////////////////////////// // Type definition methods uint32_t getScalarTypeId( diff --git a/src/dxbc/dxbc_decoder.h b/src/dxbc/dxbc_decoder.h index a79b4d970..6b2ac7ced 100644 --- a/src/dxbc/dxbc_decoder.h +++ b/src/dxbc/dxbc_decoder.h @@ -69,14 +69,34 @@ namespace dxvk { * and associated type IDs. */ struct DxbcShaderResource { - DxbcImageInfo imageInfo; - uint32_t varId = 0; - DxbcScalarType sampledType = DxbcScalarType::Float32; - uint32_t sampledTypeId = 0; - uint32_t colorTypeId = 0; - uint32_t depthTypeId = 0; + DxbcResourceType type = DxbcResourceType::Typed; + DxbcImageInfo imageInfo; + uint32_t varId = 0; + DxbcScalarType sampledType = DxbcScalarType::Float32; + uint32_t sampledTypeId = 0; + uint32_t colorTypeId = 0; + uint32_t depthTypeId = 0; + uint32_t structStride = 0; }; + + /** + * \brief Unordered access binding + * + * Stores a resource variable that is provided + * by a UAV, as well as associated type IDs. + */ + struct DxbcUav { + DxbcResourceType type = DxbcResourceType::Typed; + DxbcImageInfo imageInfo; + uint32_t varId = 0; + DxbcScalarType sampledType = DxbcScalarType::Float32; + uint32_t sampledTypeId = 0; + uint32_t imageTypeId = 0; + uint32_t structStride = 0; + }; + + /** * \brief Component swizzle * diff --git a/src/dxbc/dxbc_defs.cpp b/src/dxbc/dxbc_defs.cpp index 7c4e82b03..8d92c406b 100644 --- a/src/dxbc/dxbc_defs.cpp +++ b/src/dxbc/dxbc_defs.cpp @@ -611,75 +611,205 @@ namespace dxvk { /* DclHsJoinPhaseInstanceCount */ { }, /* DclThreadGroup */ - { }, + { 3, DxbcInstClass::Declaration, { + { DxbcOperandKind::Imm32, DxbcScalarType::Uint32 }, + { DxbcOperandKind::Imm32, DxbcScalarType::Uint32 }, + { DxbcOperandKind::Imm32, DxbcScalarType::Uint32 }, + } }, /* DclUavTyped */ - { }, + { 2, DxbcInstClass::Declaration, { + { DxbcOperandKind::DstReg, DxbcScalarType::Float32 }, + { DxbcOperandKind::Imm32, DxbcScalarType::Uint32 }, + } }, /* DclUavRaw */ - { }, + { 1, DxbcInstClass::Declaration, { + { DxbcOperandKind::DstReg, DxbcScalarType::Float32 }, + } }, /* DclUavStructured */ - { }, + { 2, DxbcInstClass::Declaration, { + { DxbcOperandKind::DstReg, DxbcScalarType::Float32 }, + { DxbcOperandKind::Imm32, DxbcScalarType::Uint32 }, + } }, /* DclThreadGroupSharedMemoryRaw */ - { }, + { 2, DxbcInstClass::Declaration, { + { DxbcOperandKind::DstReg, DxbcScalarType::Float32 }, + { DxbcOperandKind::Imm32, DxbcScalarType::Uint32 }, + } }, /* DclThreadGroupSharedMemoryStructured */ - { }, + { 3, DxbcInstClass::Declaration, { + { DxbcOperandKind::DstReg, DxbcScalarType::Float32 }, + { DxbcOperandKind::Imm32, DxbcScalarType::Uint32 }, + { DxbcOperandKind::Imm32, DxbcScalarType::Uint32 }, + } }, /* DclResourceRaw */ - { }, + { 1, DxbcInstClass::Declaration, { + { DxbcOperandKind::DstReg, DxbcScalarType::Float32 }, + } }, /* DclResourceStructured */ - { }, + { 2, DxbcInstClass::Declaration, { + { DxbcOperandKind::DstReg, DxbcScalarType::Float32 }, + { DxbcOperandKind::Imm32, DxbcScalarType::Uint32 }, + } }, /* LdUavTyped */ { }, /* StoreUavTyped */ { }, /* LdRaw */ - { }, + { 3, DxbcInstClass::BufferLoad, { + { DxbcOperandKind::DstReg, DxbcScalarType::Uint32 }, + { DxbcOperandKind::SrcReg, DxbcScalarType::Sint32 }, + { DxbcOperandKind::SrcReg, DxbcScalarType::Uint32 }, + } }, /* StoreRaw */ - { }, + { 3, DxbcInstClass::BufferStore, { + { DxbcOperandKind::DstReg, DxbcScalarType::Uint32 }, + { DxbcOperandKind::SrcReg, DxbcScalarType::Sint32 }, + { DxbcOperandKind::SrcReg, DxbcScalarType::Uint32 }, + } }, /* LdStructured */ - { }, + { 4, DxbcInstClass::BufferLoad, { + { DxbcOperandKind::DstReg, DxbcScalarType::Uint32 }, + { DxbcOperandKind::SrcReg, DxbcScalarType::Sint32 }, + { DxbcOperandKind::SrcReg, DxbcScalarType::Sint32 }, + { DxbcOperandKind::SrcReg, DxbcScalarType::Uint32 }, + } }, /* StoreStructured */ - { }, + { 4, DxbcInstClass::BufferStore, { + { DxbcOperandKind::DstReg, DxbcScalarType::Uint32 }, + { DxbcOperandKind::SrcReg, DxbcScalarType::Sint32 }, + { DxbcOperandKind::SrcReg, DxbcScalarType::Sint32 }, + { DxbcOperandKind::SrcReg, DxbcScalarType::Uint32 }, + } }, /* AtomicAnd */ - { }, + { 3, DxbcInstClass::Atomic, { + { DxbcOperandKind::DstReg, DxbcScalarType::Uint32 }, + { DxbcOperandKind::SrcReg, DxbcScalarType::Sint32 }, + { DxbcOperandKind::SrcReg, DxbcScalarType::Uint32 }, + } }, /* AtomicOr */ - { }, + { 3, DxbcInstClass::Atomic, { + { DxbcOperandKind::DstReg, DxbcScalarType::Uint32 }, + { DxbcOperandKind::SrcReg, DxbcScalarType::Sint32 }, + { DxbcOperandKind::SrcReg, DxbcScalarType::Uint32 }, + } }, /* AtomicXor */ - { }, + { 3, DxbcInstClass::Atomic, { + { DxbcOperandKind::DstReg, DxbcScalarType::Uint32 }, + { DxbcOperandKind::SrcReg, DxbcScalarType::Sint32 }, + { DxbcOperandKind::SrcReg, DxbcScalarType::Uint32 }, + } }, /* AtomicCmpStore */ - { }, + { 4, DxbcInstClass::Atomic, { + { DxbcOperandKind::DstReg, DxbcScalarType::Uint32 }, + { DxbcOperandKind::SrcReg, DxbcScalarType::Sint32 }, + { DxbcOperandKind::SrcReg, DxbcScalarType::Uint32 }, + { DxbcOperandKind::SrcReg, DxbcScalarType::Uint32 }, + } }, /* AtomicIAdd */ - { }, + { 3, DxbcInstClass::Atomic, { + { DxbcOperandKind::DstReg, DxbcScalarType::Uint32 }, + { DxbcOperandKind::SrcReg, DxbcScalarType::Sint32 }, + { DxbcOperandKind::SrcReg, DxbcScalarType::Uint32 }, + } }, /* AtomicIMax */ - { }, + { 3, DxbcInstClass::Atomic, { + { DxbcOperandKind::DstReg, DxbcScalarType::Sint32 }, + { DxbcOperandKind::SrcReg, DxbcScalarType::Sint32 }, + { DxbcOperandKind::SrcReg, DxbcScalarType::Sint32 }, + } }, /* AtomicIMin */ - { }, + { 3, DxbcInstClass::Atomic, { + { DxbcOperandKind::DstReg, DxbcScalarType::Sint32 }, + { DxbcOperandKind::SrcReg, DxbcScalarType::Sint32 }, + { DxbcOperandKind::SrcReg, DxbcScalarType::Sint32 }, + } }, /* AtomicUMax */ - { }, + { 3, DxbcInstClass::Atomic, { + { DxbcOperandKind::DstReg, DxbcScalarType::Uint32 }, + { DxbcOperandKind::SrcReg, DxbcScalarType::Sint32 }, + { DxbcOperandKind::SrcReg, DxbcScalarType::Uint32 }, + } }, /* AtomicUMin */ - { }, + { 3, DxbcInstClass::Atomic, { + { DxbcOperandKind::DstReg, DxbcScalarType::Uint32 }, + { DxbcOperandKind::SrcReg, DxbcScalarType::Sint32 }, + { DxbcOperandKind::SrcReg, DxbcScalarType::Uint32 }, + } }, /* ImmAtomicAlloc */ { }, /* ImmAtomicConsume */ { }, /* ImmAtomicIAdd */ - { }, + { 4, DxbcInstClass::Atomic, { + { DxbcOperandKind::DstReg, DxbcScalarType::Uint32 }, + { DxbcOperandKind::DstReg, DxbcScalarType::Uint32 }, + { DxbcOperandKind::SrcReg, DxbcScalarType::Sint32 }, + { DxbcOperandKind::SrcReg, DxbcScalarType::Uint32 }, + } }, /* ImmAtomicAnd */ - { }, + { 4, DxbcInstClass::Atomic, { + { DxbcOperandKind::DstReg, DxbcScalarType::Uint32 }, + { DxbcOperandKind::DstReg, DxbcScalarType::Uint32 }, + { DxbcOperandKind::SrcReg, DxbcScalarType::Sint32 }, + { DxbcOperandKind::SrcReg, DxbcScalarType::Uint32 }, + } }, /* ImmAtomicOr */ - { }, + { 4, DxbcInstClass::Atomic, { + { DxbcOperandKind::DstReg, DxbcScalarType::Uint32 }, + { DxbcOperandKind::DstReg, DxbcScalarType::Uint32 }, + { DxbcOperandKind::SrcReg, DxbcScalarType::Sint32 }, + { DxbcOperandKind::SrcReg, DxbcScalarType::Uint32 }, + } }, /* ImmAtomicXor */ - { }, + { 4, DxbcInstClass::Atomic, { + { DxbcOperandKind::DstReg, DxbcScalarType::Uint32 }, + { DxbcOperandKind::DstReg, DxbcScalarType::Uint32 }, + { DxbcOperandKind::SrcReg, DxbcScalarType::Sint32 }, + { DxbcOperandKind::SrcReg, DxbcScalarType::Uint32 }, + } }, /* ImmAtomicExch */ - { }, + { 4, DxbcInstClass::Atomic, { + { DxbcOperandKind::DstReg, DxbcScalarType::Uint32 }, + { DxbcOperandKind::DstReg, DxbcScalarType::Uint32 }, + { DxbcOperandKind::SrcReg, DxbcScalarType::Sint32 }, + { DxbcOperandKind::SrcReg, DxbcScalarType::Uint32 }, + } }, /* ImmAtomicCmpExch */ - { }, + { 5, DxbcInstClass::Atomic, { + { DxbcOperandKind::DstReg, DxbcScalarType::Uint32 }, + { DxbcOperandKind::DstReg, DxbcScalarType::Uint32 }, + { DxbcOperandKind::SrcReg, DxbcScalarType::Sint32 }, + { DxbcOperandKind::SrcReg, DxbcScalarType::Uint32 }, + { DxbcOperandKind::SrcReg, DxbcScalarType::Uint32 }, + } }, /* ImmAtomicImax */ - { }, + { 4, DxbcInstClass::Atomic, { + { DxbcOperandKind::DstReg, DxbcScalarType::Sint32 }, + { DxbcOperandKind::DstReg, DxbcScalarType::Sint32 }, + { DxbcOperandKind::SrcReg, DxbcScalarType::Sint32 }, + { DxbcOperandKind::SrcReg, DxbcScalarType::Sint32 }, + } }, /* ImmAtomicImin */ - { }, + { 4, DxbcInstClass::Atomic, { + { DxbcOperandKind::DstReg, DxbcScalarType::Sint32 }, + { DxbcOperandKind::DstReg, DxbcScalarType::Sint32 }, + { DxbcOperandKind::SrcReg, DxbcScalarType::Sint32 }, + { DxbcOperandKind::SrcReg, DxbcScalarType::Sint32 }, + } }, /* ImmAtomicUmax */ - { }, + { 4, DxbcInstClass::Atomic, { + { DxbcOperandKind::DstReg, DxbcScalarType::Uint32 }, + { DxbcOperandKind::DstReg, DxbcScalarType::Uint32 }, + { DxbcOperandKind::SrcReg, DxbcScalarType::Sint32 }, + { DxbcOperandKind::SrcReg, DxbcScalarType::Uint32 }, + } }, /* ImmAtomicUmin */ - { }, + { 4, DxbcInstClass::Atomic, { + { DxbcOperandKind::DstReg, DxbcScalarType::Uint32 }, + { DxbcOperandKind::DstReg, DxbcScalarType::Uint32 }, + { DxbcOperandKind::SrcReg, DxbcScalarType::Sint32 }, + { DxbcOperandKind::SrcReg, DxbcScalarType::Uint32 }, + } }, /* Sync */ { }, /* DAdd */ diff --git a/src/dxbc/dxbc_defs.h b/src/dxbc/dxbc_defs.h index 2080d9351..70f5e6258 100644 --- a/src/dxbc/dxbc_defs.h +++ b/src/dxbc/dxbc_defs.h @@ -32,6 +32,9 @@ namespace dxvk { CustomData, ///< Immediate constant buffer ControlFlow, ///< Control flow instructions GeometryEmit, ///< Special geometry shader instructions + Atomic, ///< Atomic operations + BufferLoad, ///< Structured or raw buffer load + BufferStore, ///< Structured or raw buffer store TextureQuery, ///< Texture query instruction TextureFetch, ///< Texture fetch instruction TextureSample, ///< Texture sampling instruction diff --git a/src/dxbc/dxbc_enums.h b/src/dxbc/dxbc_enums.h index 02aee0375..fc9f3648f 100644 --- a/src/dxbc/dxbc_enums.h +++ b/src/dxbc/dxbc_enums.h @@ -560,4 +560,11 @@ namespace dxvk { ImmConstBuf = 3, }; + + enum class DxbcResourceType : uint32_t { + Typed = 0, + Raw = 1, + Structured = 2, + }; + } \ No newline at end of file diff --git a/src/spirv/spirv_module.cpp b/src/spirv/spirv_module.cpp index 29868ae68..59e785853 100644 --- a/src/spirv/spirv_module.cpp +++ b/src/spirv/spirv_module.cpp @@ -1680,6 +1680,40 @@ namespace dxvk { } + uint32_t SpirvModule::opImageRead( + uint32_t resultType, + uint32_t image, + uint32_t coordinates, + const SpirvImageOperands& operands) { + uint32_t resultId = this->allocateId(); + + m_code.putIns (spv::OpImageRead, + 5 + getImageOperandWordCount(operands)); + m_code.putWord(resultType); + m_code.putWord(resultId); + m_code.putWord(image); + m_code.putWord(coordinates); + + putImageOperands(operands); + return resultId; + } + + + void SpirvModule::opImageWrite( + uint32_t image, + uint32_t coordinates, + uint32_t texel, + const SpirvImageOperands& operands) { + m_code.putIns (spv::OpImageWrite, + 4 + getImageOperandWordCount(operands)); + m_code.putWord(image); + m_code.putWord(coordinates); + m_code.putWord(texel); + + putImageOperands(operands); + } + + uint32_t SpirvModule::opSampledImage( uint32_t resultType, uint32_t image, diff --git a/src/spirv/spirv_module.h b/src/spirv/spirv_module.h index aa42798c1..cde50109e 100644 --- a/src/spirv/spirv_module.h +++ b/src/spirv/spirv_module.h @@ -581,6 +581,18 @@ namespace dxvk { uint32_t pointerId, uint32_t valueId); + uint32_t opImageRead( + uint32_t resultType, + uint32_t image, + uint32_t coordinates, + const SpirvImageOperands& operands); + + void opImageWrite( + uint32_t image, + uint32_t coordinates, + uint32_t texel, + const SpirvImageOperands& operands); + uint32_t opSampledImage( uint32_t resultType, uint32_t image,