FernandoS27

// Copyright 2019 yuzu Emulator Project
// Licensed under GPLv2 or any later version
// Refer to the license.txt file included.

#pragma once

#include <unordered_map>
#include "common/common_types.h"
#include "common/hash.h"
#include "video_core/engines/const_buffer_engine_interface.h"

FernandoS27

struct BranchIndirectInfo {
    u32 buffer{};
    u32 offset{};
    u32 entries{};
};

std::optional<BranchIndirectInfo> TrackBranchIndirectInfo(const CFGRebuildState& state,
                                                          u32 start_address, u32 current_position) {
    const u32 shader_start = state.start;
    u32 pos = current_position;

FernandoS27

#version 430 core
#extension GL_ARB_separate_shader_objects : enable

#define EMULATION_UBO_BINDING 2
#define CBUF_BINDING_1 3
#define CBUF_BINDING_3 4
#define CBUF_BINDING_4 5
#define SAMPLER_BINDING_0 0
#define SAMPLER_BINDING_1 1
#define SAMPLER_BINDING_2 2

FernandoS27

        /*0538*/               @P0 BRK ;
        /*0548*/                   IADD32I R5, R5, -0x1 ;
        /*0550*/                   IMNMX.U32 R5, R5, 0x2, PT ;
        /*0558*/                   SHL R5, R5, 0x2 ;
        /*0568*/                   LDC R5, c[0x1][R5+0xc0] ;
        /*0570*/                   BRX R5 -0x578 ;
        /*0578*/         {         MOV R6, R1 ;
        /*0588*/                   BRK         }
        /*0590*/         {         MOV R6, R2 ;
        /*0598*/                   BRK         }

FernandoS27

void SwizzleSubrect(u32 subrect_width, u32 subrect_height, u32 source_pitch, u32 swizzled_width,
                    u32 bytes_per_pixel, VAddr swizzled_data, VAddr unswizzled_data,
                    u32 block_height) {
    const u32 image_width_in_gobs{(swizzled_width * bytes_per_pixel + 63) / 64};
    for (u32 line = 0; line < subrect_height; ++line) {
        const u32 gob_address_y =
            (line / (8 * block_height)) * 512 * block_height * image_width_in_gobs +
            (line % (8 * block_height) / 8) * 512;
        auto& table = legacy_swizzle_table[line % 8];
        for (u32 x = 0; x < subrect_width; ++x) {

FernandoS27

// Example program
#include <iostream>
#include <string>
#include <cmath>
#include <utility>
#include <array>
#include <cstring>

typedef unsigned char u8;
typedef unsigned int u32;

FernandoS27

#version 430 core
#extension GL_ARB_separate_shader_objects : enable

#define MAX_CONSTBUFFER_ELEMENTS 4096
bool exec_fragment();

in vec4 position;
layout(location = 0) out vec4 FragColor0;
layout(location = 1) out vec4 FragColor1;
layout(location = 2) out vec4 FragColor2;

FernandoS27

bool IsSchedInstruction(u32 offset, u32 main_offset) {
    // sched instructions appear once every 4 instructions.
    static constexpr size_t SchedPeriod = 4;
    u32 absolute_offset = offset - main_offset;

    return (absolute_offset % SchedPeriod) == 0;
}

void ShaderDumper::dump() {
    FileUtil::IOFile sFile;

FernandoS27

unsigned int approx(float f) {
    float ax = fabs(f);
    unsigned int extra = 1;
    if (ax >= 128.f) {
        ax = 129.f;
        extra = 0;
    }
    double w;
    float ax2 = modf(ax, &w);
    unsigned int sign = (f < 0.0f) << 31;

FernandoS27

// Kernel definition
__global__ void Test(float* A, float* B, float* C)
{
    int i = threadIdx.x;
    C[i] = A[i] + B[i];
}