video_core: Use texture buffers for untyped format load/store

Author: raphaelthegreat

src/core/libraries/avplayer/avplayer.cpp

@@ -120,6 +120,7 @@ bool PS4_SYSV_ABI sceAvPlayerGetVideoDataEx(SceAvPlayerHandle handle,
 }
 
 SceAvPlayerHandle PS4_SYSV_ABI sceAvPlayerInit(SceAvPlayerInitData* data) {
+    return nullptr;
     LOG_TRACE(Lib_AvPlayer, "called");
     if (data == nullptr) {
         return nullptr;
@@ -325,4 +326,4 @@ void RegisterlibSceAvPlayer(Core::Loader::SymbolsResolver* sym) {
     LIB_FUNCTION("yN7Jhuv8g24", "libSceAvPlayer", 1, "libSceAvPlayer", 1, 0, sceAvPlayerVprintf);
 };
 
-} // namespace Libraries::AvPlayer
+} // namespace Libraries::AvPlayer

src/core/libraries/kernel/thread_management.cpp

@@ -1066,7 +1066,16 @@ ScePthread PThreadPool::Create() {
         }
     }
 
+#ifdef _WIN64
     auto* ret = new PthreadInternal{};
+#else
+    // TODO: Linux specific hack
+    static u8* hint_address = reinterpret_cast<u8*>(0x7FFFFC000ULL);
+    auto* ret = reinterpret_cast<PthreadInternal*>(
+        mmap(hint_address, sizeof(PthreadInternal), PROT_READ | PROT_WRITE,
+             MAP_PRIVATE | MAP_ANONYMOUS | MAP_FIXED, -1, 0));
+    hint_address += Common::AlignUp(sizeof(PthreadInternal), 4_KB);
+#endif
     ret->is_free = false;
     ret->is_detached = false;
     ret->is_almost_done = false;

src/shader_recompiler/backend/spirv/emit_spirv.cpp

@@ -189,6 +189,9 @@ void DefineEntryPoint(const IR::Program& program, EmitContext& ctx, Id main) {
         ctx.AddCapability(spv::Capability::StorageImageExtendedFormats);
         ctx.AddCapability(spv::Capability::StorageImageWriteWithoutFormat);
     }
+    if (info.has_texel_buffers) {
+        ctx.AddCapability(spv::Capability::SampledBuffer);
+    }
     switch (program.info.stage) {
     case Stage::Compute: {
         const std::array<u32, 3> workgroup_size{program.info.workgroup_size};

src/shader_recompiler/backend/spirv/emit_spirv_context_get_set.cpp

@@ -262,171 +262,15 @@ Id EmitLoadBufferF32x4(EmitContext& ctx, IR::Inst*, u32 handle, Id address) {
     return EmitLoadBufferF32xN<4>(ctx, handle, address);
 }
 
-static bool IsSignedInteger(AmdGpu::NumberFormat format) {
-    switch (format) {
-    case AmdGpu::NumberFormat::Unorm:
-    case AmdGpu::NumberFormat::Uscaled:
-    case AmdGpu::NumberFormat::Uint:
-        return false;
-    case AmdGpu::NumberFormat::Snorm:
-    case AmdGpu::NumberFormat::Sscaled:
-    case AmdGpu::NumberFormat::Sint:
-    case AmdGpu::NumberFormat::SnormNz:
-        return true;
-    case AmdGpu::NumberFormat::Float:
-    default:
-        UNREACHABLE();
-    }
-}
-
-static u32 UXBitsMax(u32 bit_width) {
-    return (1u << bit_width) - 1u;
-}
-
-static u32 SXBitsMax(u32 bit_width) {
-    return (1u << (bit_width - 1u)) - 1u;
-}
-
-static Id ConvertValue(EmitContext& ctx, Id value, AmdGpu::NumberFormat format, u32 bit_width) {
-    switch (format) {
-    case AmdGpu::NumberFormat::Unorm:
-        return ctx.OpFDiv(ctx.F32[1], value, ctx.ConstF32(float(UXBitsMax(bit_width))));
-    case AmdGpu::NumberFormat::Snorm:
-        return ctx.OpFDiv(ctx.F32[1], value, ctx.ConstF32(float(SXBitsMax(bit_width))));
-    case AmdGpu::NumberFormat::SnormNz:
-        // (x * 2 + 1) / (Format::SMAX * 2)
-        value = ctx.OpFMul(ctx.F32[1], value, ctx.ConstF32(2.f));
-        value = ctx.OpFAdd(ctx.F32[1], value, ctx.ConstF32(1.f));
-        return ctx.OpFDiv(ctx.F32[1], value, ctx.ConstF32(float(SXBitsMax(bit_width) * 2)));
-    case AmdGpu::NumberFormat::Uscaled:
-    case AmdGpu::NumberFormat::Sscaled:
-    case AmdGpu::NumberFormat::Uint:
-    case AmdGpu::NumberFormat::Sint:
-    case AmdGpu::NumberFormat::Float:
-        return value;
-    default:
-        UNREACHABLE_MSG("Unsupported number format for conversion: {}",
-                        magic_enum::enum_name(format));
-    }
-}
-
-static Id ComponentOffset(EmitContext& ctx, Id address, u32 stride, u32 bit_offset) {
-    Id comp_offset = ctx.ConstU32(bit_offset);
-    if (stride < 4) {
-        // comp_offset += (address % 4) * 8;
-        const Id byte_offset = ctx.OpUMod(ctx.U32[1], address, ctx.ConstU32(4u));
-        const Id bit_offset = ctx.OpShiftLeftLogical(ctx.U32[1], byte_offset, ctx.ConstU32(3u));
-        comp_offset = ctx.OpIAdd(ctx.U32[1], comp_offset, bit_offset);
-    }
-    return comp_offset;
-}
-
-static Id GetBufferFormatValue(EmitContext& ctx, u32 handle, Id address, u32 comp) {
-    auto& buffer = ctx.buffers[handle];
-    const auto format = buffer.dfmt;
-    switch (format) {
-    case AmdGpu::DataFormat::FormatInvalid:
-        return ctx.f32_zero_value;
-    case AmdGpu::DataFormat::Format8:
-    case AmdGpu::DataFormat::Format16:
-    case AmdGpu::DataFormat::Format32:
-    case AmdGpu::DataFormat::Format8_8:
-    case AmdGpu::DataFormat::Format16_16:
-    case AmdGpu::DataFormat::Format10_11_11:
-    case AmdGpu::DataFormat::Format11_11_10:
-    case AmdGpu::DataFormat::Format10_10_10_2:
-    case AmdGpu::DataFormat::Format2_10_10_10:
-    case AmdGpu::DataFormat::Format8_8_8_8:
-    case AmdGpu::DataFormat::Format32_32:
-    case AmdGpu::DataFormat::Format16_16_16_16:
-    case AmdGpu::DataFormat::Format32_32_32:
-    case AmdGpu::DataFormat::Format32_32_32_32: {
-        const u32 num_components = AmdGpu::NumComponents(format);
-        if (comp >= num_components) {
-            return ctx.f32_zero_value;
-        }
-
-        // uint index = address / 4;
-        Id index = ctx.OpShiftRightLogical(ctx.U32[1], address, ctx.ConstU32(2u));
-        const u32 stride = buffer.stride;
-        if (stride > 4) {
-            const u32 index_offset = u32(AmdGpu::ComponentOffset(format, comp) / 32);
-            if (index_offset > 0) {
-                // index += index_offset;
-                index = ctx.OpIAdd(ctx.U32[1], index, ctx.ConstU32(index_offset));
-            }
-        }
-        const Id ptr = ctx.OpAccessChain(buffer.pointer_type, buffer.id, ctx.u32_zero_value, index);
-
-        const u32 bit_offset = AmdGpu::ComponentOffset(format, comp) % 32;
-        const u32 bit_width = AmdGpu::ComponentBits(format, comp);
-        const auto num_format = buffer.nfmt;
-        if (num_format == AmdGpu::NumberFormat::Float) {
-            if (bit_width == 32) {
-                return ctx.OpLoad(ctx.F32[1], ptr);
-            } else if (bit_width == 16) {
-                const Id comp_offset = ComponentOffset(ctx, address, stride, bit_offset);
-                Id value = ctx.OpLoad(ctx.U32[1], ptr);
-                value =
-                    ctx.OpBitFieldSExtract(ctx.S32[1], value, comp_offset, ctx.ConstU32(bit_width));
-                value = ctx.OpSConvert(ctx.U16, value);
-                value = ctx.OpBitcast(ctx.F16[1], value);
-                return ctx.OpFConvert(ctx.F32[1], value);
-            } else {
-                UNREACHABLE_MSG("Invalid float bit width {}", bit_width);
-            }
-        } else {
-            Id value = ctx.OpLoad(ctx.U32[1], ptr);
-            const bool is_signed = IsSignedInteger(num_format);
-            if (bit_width < 32) {
-                const Id comp_offset = ComponentOffset(ctx, address, stride, bit_offset);
-                if (is_signed) {
-                    value = ctx.OpBitFieldSExtract(ctx.S32[1], value, comp_offset,
-                                                   ctx.ConstU32(bit_width));
-                } else {
-                    value = ctx.OpBitFieldUExtract(ctx.U32[1], value, comp_offset,
-                                                   ctx.ConstU32(bit_width));
-                }
-            }
-            value = ctx.OpBitcast(ctx.F32[1], value);
-            return ConvertValue(ctx, value, num_format, bit_width);
-        }
-        break;
-    }
-    default:
-        UNREACHABLE_MSG("Invalid format for conversion: {}", magic_enum::enum_name(format));
-    }
-}
-
-template <u32 N>
-static Id EmitLoadBufferFormatF32xN(EmitContext& ctx, IR::Inst* inst, u32 handle, Id address) {
-    auto& buffer = ctx.buffers[handle];
-    address = ctx.OpIAdd(ctx.U32[1], address, buffer.offset);
-    if constexpr (N == 1) {
-        return GetBufferFormatValue(ctx, handle, address, 0);
-    } else {
-        boost::container::static_vector<Id, N> ids;
-        for (u32 i = 0; i < N; i++) {
-            ids.push_back(GetBufferFormatValue(ctx, handle, address, i));
-        }
-        return ctx.OpCompositeConstruct(ctx.F32[N], ids);
-    }
-}
-
 Id EmitLoadBufferFormatF32(EmitContext& ctx, IR::Inst* inst, u32 handle, Id address) {
-    return EmitLoadBufferFormatF32xN<1>(ctx, inst, handle, address);
-}
-
-Id EmitLoadBufferFormatF32x2(EmitContext& ctx, IR::Inst* inst, u32 handle, Id address) {
-    return EmitLoadBufferFormatF32xN<2>(ctx, inst, handle, address);
-}
-
-Id EmitLoadBufferFormatF32x3(EmitContext& ctx, IR::Inst* inst, u32 handle, Id address) {
-    return EmitLoadBufferFormatF32xN<3>(ctx, inst, handle, address);
-}
-
-Id EmitLoadBufferFormatF32x4(EmitContext& ctx, IR::Inst* inst, u32 handle, Id address) {
-    return EmitLoadBufferFormatF32xN<4>(ctx, inst, handle, address);
+    const auto& buffer = ctx.texture_buffers[handle];
+    const Id tex_buffer = ctx.OpLoad(buffer.image_type, buffer.id);
+    const Id coord = ctx.OpIAdd(ctx.U32[1], address, buffer.coord_offset);
+    Id texel = ctx.OpImageFetch(buffer.result_type, tex_buffer, coord);
+    if (buffer.is_integer) {
+        texel = ctx.OpBitcast(ctx.F32[4], texel);
+    }
+    return texel;
 }
 
 template <u32 N>
@@ -467,6 +311,7 @@ void EmitStoreBufferU32(EmitContext& ctx, IR::Inst* inst, u32 handle, Id address
     EmitStoreBufferF32xN<1>(ctx, handle, address, value);
 }
 
+<<<<<<< HEAD
 static Id ConvertF32ToFormat(EmitContext& ctx, Id value, AmdGpu::NumberFormat format,
                              u32 bit_width) {
     switch (format) {
@@ -541,23 +386,16 @@ static void EmitStoreBufferFormatF32xN(EmitContext& ctx, u32 handle, Id address,
     }
 }
 
+=======
+>>>>>>> 8b824588 (video_core: Use texture buffers for untyped format load/store)
 void EmitStoreBufferFormatF32(EmitContext& ctx, IR::Inst* inst, u32 handle, Id address, Id value) {
-    EmitStoreBufferFormatF32xN<1>(ctx, handle, address, value);
-}
-
-void EmitStoreBufferFormatF32x2(EmitContext& ctx, IR::Inst* inst, u32 handle, Id address,
-                                Id value) {
-    EmitStoreBufferFormatF32xN<2>(ctx, handle, address, value);
-}
-
-void EmitStoreBufferFormatF32x3(EmitContext& ctx, IR::Inst* inst, u32 handle, Id address,
-                                Id value) {
-    EmitStoreBufferFormatF32xN<3>(ctx, handle, address, value);
-}
-
-void EmitStoreBufferFormatF32x4(EmitContext& ctx, IR::Inst* inst, u32 handle, Id address,
-                                Id value) {
-    EmitStoreBufferFormatF32xN<4>(ctx, handle, address, value);
+    const auto& buffer = ctx.texture_buffers[handle];
+    const Id tex_buffer = ctx.OpLoad(buffer.image_type, buffer.id);
+    const Id coord = ctx.OpIAdd(ctx.U32[1], address, buffer.coord_offset);
+    if (buffer.is_integer) {
+        value = ctx.OpBitcast(ctx.U32[4], value);
+    }
+    ctx.OpImageWrite(tex_buffer, coord, value);
 }
 
 } // namespace Shader::Backend::SPIRV