code:
$newline never
\<!doctype html>
\<!--[if lt IE 7]> <html class="no-js ie6 oldie" lang="en"> <![endif]-->
\<!--[if IE 7]>    <html class="no-js ie7 oldie" lang="en"> <![endif]-->
\<!--[if IE 8]>    <html class="no-js ie8 oldie" lang="en"> <![endif]-->
\<!--[if gt IE 8]><!-->
<html class="no-js" lang="en"> <!--<![endif]-->
  <head>
    <meta charset="UTF-8">

    <title>#{pageTitle pc}
    <meta name="description" content="">
    <meta name="author" content="">

    <meta name="viewport" content="width=device-width,initial-scale=1">

    ^{pageHead pc}

    \<!--[if lt IE 9]>
    \<script src="http://html5shiv.googlecode.com/svn/trunk/html5.js"></script>
    \<![endif]-->
    <script type="text/javascript" src="https://cdnjs.cloudflare.com/ajax/libs/jquery/2.1.4/jquery.js">
    <script type="text/javascript" src="https://cdnjs.cloudflare.com/ajax/libs/js-cookie/2.0.3/js.cookie.min.js">

    \<!-- Bootstrap-3.3.7 compiled and minified JavaScript -->
    <script src="https://maxcdn.bootstrapcdn.com/bootstrap/3.3.7/js/bootstrap.min.js" integrity="sha384-Tc5IQib027qvyjSMfHjOMaLkfuWVxZxUPnCJA7l2mCWNIpG9mGCD8wGNIcPD7Txa" crossorigin="anonymous">

    <script>
      /* The `defaultCsrfMiddleware` Middleware added in Foundation.hs adds a CSRF token to the request cookies. */
      /* AJAX requests should add that token to a header to be validated by the server. */
      /* See the CSRF documentation in the Yesod.Core.Handler module of the yesod-core package for details. */
      var csrfHeaderName = "#{TE.decodeUtf8 $ CI.foldedCase defaultCsrfHeaderName}";

      var csrfCookieName = "#{TE.decodeUtf8 defaultCsrfCookieName}";
      var csrfToken = Cookies.get(csrfCookieName);

      if (csrfToken) {
      \  $.ajaxPrefilter(function( options, originalOptions, jqXHR ) {
      \      if (!options.crossDomain) {
      \          jqXHR.setRequestHeader(csrfHeaderName, csrfToken);
      \      }
      \  });
      }

    <script>
      document.documentElement.className = document.documentElement.className.replace(/\bno-js\b/,'js');
  <body>
    ^{pageBody pc}

    $maybe analytics <- appAnalytics $ appSettings master
      <script>
        if(!window.location.href.match(/localhost/)){
            (function(i,s,o,g,r,a,m){i['GoogleAnalyticsObject']=r;i[r]=i[r]||function(){
            (i[r].q=i[r].q||[]).push(arguments)},i[r].l=1*new Date();a=s.createElement(o),
            m=s.getElementsByTagName(o)[0];a.async=1;a.src=g;m.parentNode.insertBefore(a,m)
            })(window,document,'script','https://www.google-analytics.com/analytics.js','ga');

            ga('create', '#{analytics}', 'auto');
            ga('send', 'pageview');
        }

repo_name: vaeringjar/yesod

license: unknown


code:
            <button type="button" .navbar-toggle.collapsed data-toggle="collapse" data-target="#navbar" aria-expanded="false" aria-controls="navbar">
                Toggle navigation

                $forall MenuItem label route _ <- navbarLeftFilteredMenuItems

                        #{label}

                $forall MenuItem label route _ <- navbarRightFilteredMenuItems

                        #{label}

    $if not $ Just HomeR == mcurrentRoute

            $forall bc <- parents

                    #{snd bc}

#{title}

    $maybe msg <- mmsg

#{msg}

$if (Just HomeR == mcurrentRoute)
    ^{widget}
$else

                ^{widget}

            #{appCopyright $ appSettings master}


code:
.masthead,
.navbar {
    background-color: rgb(27, 28, 29);
}

.navbar-default .navbar-nav > .active > a {
    background-color: transparent;
    border-bottom: 2px solid white;
}

.navbar-nav {
    padding-bottom: 1em;
}

.masthead {
    margin-top: -21px;
    color: white;
    text-align: center;
    min-height: 500px;
}

.masthead .header {
    max-width: 700px;
    margin: 0 auto;
    font-family: Lato,'Helvetica Neue',Arial,Helvetica,sans-serif;
}

.masthead h1.header {
    margin-top: 1em;
    margin-bottom: 0em;
    font-size: 4.5em;
    line-height: 1.2em;
    font-weight: normal;
 }

.masthead h2 {
    font-size: 1.7em;
    font-weight: normal;
}

.masthead .btn {
    margin: 1em 0;
}

/* Common styles for all types */
.bs-callout {
  padding: 20px;
  margin: 20px 0;
  border: 1px solid #eee;
  border-left-width: 5px;
  border-radius: 3px;
}

.bs-callout p:last-child {
  margin-bottom: 0;
}

.bs-callout-info {
  border-left-color: #1b809e;
}

/* Space things out */
.bs-docs-section {
  margin-bottom: 60px;
}
.bs-docs-section:last-child {
  margin-bottom: 0;
}

#message {
    margin-bottom: 40px;
}


code:
                Yesod—a modern framework for blazing fast websites

                Fast, stable & spiced with great community

                Read the Book

Starting

                        Now that you have a working project you should use the
                        Yesod book to learn more.

                        You can also use this scaffolded site to explore some concepts, and best practices.

                            This page was generated by the <tt>#{handlerName}</tt> handler in
                            <tt>Handler/Home.hs</tt>.

                            The <tt>#{handlerName}</tt> handler is set to generate your
                            site's home screen in Routes file
                            <tt>config/routes

                            We can link to other handlers, like the Profile.
                            Try it out as an anonymous user and see the access denied.
                            Then, try to login with the dummy authentication added
                            while in development.

                            The HTML you are seeing now is actually composed by a number of widgets, #
                            most of them are brought together by the <tt>defaultLayout</tt> function which #
                            is defined in the <tt>Foundation.hs</tt> module, and used by <tt>#{handlerName}</tt>. #
                            All the files for templates and wigdets are in <tt>templates</tt>.

                            A Widget's Html, Css and Javascript are separated in three files with the
                            <tt>.hamlet</tt>, <tt>.lucius</tt> and <tt>.julius</tt> extensions.

                            If you had javascript enabled then you wouldn't be seeing this.

Forms

                        This is an example of a form. Read the
                        Forms chapter #
                        on the yesod book to learn more about them.

                    <form .form-horizontal method=post action=@{HomeR}#forms enctype=#{formEnctype}>
                        ^{formWidget}

                        <button .btn.btn-primary type="submit">
                            Upload it!

                    $maybe (FileForm info con) <- submission
                        Your file type is #{fileContentType info}. You say it has: #{con}

                    $nothing
                        File upload result will be here...

JSON

                        Yesod has JSON support baked-in.
                        The form below makes an AJAX request with Javascript,
                        then updates the page with your submission.
                        (see <tt>Handler/Comment.hs</tt>, <tt>templates/homepage.julius</tt>,
                        and <tt>Handler/Home.hs</tt> for the implementation).

                    <form .form-horizontal ##{commentFormId}>

                            <textarea rows="2" ##{commentTextareaId} placeholder="Your comment here..." required></textarea>

                        <button .btn.btn-primary type=submit>
                           Create comment

                       Your comments will appear here. You can also open the
                       console log to see the raw response from the server.

                       $forall comment <- allComments

#{commentMessage $ entityVal comment}

Testing

                       And last but not least, Testing. In <tt>test/Spec.hs</tt> you will find a #
                       test suite that performs tests on this page.

                       You can run your tests by doing: stack test


code:
document.getElementById(#{toJSON aDomId}).innerHTML = "This text was added by the Javascript part of the homepage widget.";

$(function() {
  $("##{rawJS commentFormId}").submit(function(event) {
    event.preventDefault();

    var message = $("##{rawJS commentTextareaId}").val();
    // (Browsers that enforce the "required" attribute on the textarea won't see this alert)
    if (!message) {
      alert("Please fill out the comment form first.");
      return;
    }

    // Make an AJAX request to the server to create a new comment
    $.ajax({
      url: '@{CommentR}',
      type: 'POST',
      contentType: "application/json",
      data: JSON.stringify({
        message: message,
      }),
      success: function (data) {
        var newNode = $("

");
        newNode.text(data.message);
        console.log(data);
        $("##{rawJS commentListId}").append(newNode);
      },
      error: function (data) {
        console.log("Error creating comment: " + data);
      },
    });

  });
});


code:
h2##{aDomId} {
    color: #990
}

li {
    line-height: 2em;
    font-size: 16px
}

##{commentTextareaId} {
    width: 400px;
    height: 100px;
}


code:
<h1>#{blogPostTitle blogPost}

<article .jumbotron>
    #{blogPostArticle blogPost}


code:
 Latest 50 Posts

    $forall Entity id post <- allPosts

          #{blogPostTitle post}


code:
<h1> Create a post!

<p> Words here...

<form method=post action=@{PostNewR} enctype=#{enctype}>
  ^{widget}
  <button class="btn btn-danger">Create Post!


code:
textarea {
    min-height: 400px;
}


code:
      Access granted!

      This page is protected and access is allowed only for authenticated users.

      Your data is protected with us #{userIdent user}!


code:
{-# LANGUAGE NoImplicitPrelude #-}
{-# LANGUAGE OverloadedStrings #-}
module Handler.CommentSpec (spec) where

import TestImport
import Data.Aeson

spec :: Spec
spec = withApp $ do
    describe "valid request" $ do
        it "gives a 200" $ do
            get HomeR
            statusIs 200

            let message = "My message" :: Text
                body = object [ "message" .= message ]
                encoded = encode body

            request $ do
                setMethod "POST"
                setUrl CommentR
                setRequestBody encoded
                addRequestHeader ("Content-Type", "application/json")

            statusIs 200

            [Entity _id comment] <- runDB $ selectList [CommentMessage ==. message] []
            assertEq "Should have " comment (Comment message Nothing)

    describe "invalid requests" $ do
        it "400s when the JSON body is invalid" $ do
            get HomeR

            let body = object [ "foo" .= ("My message" :: Value) ]

            request $ do
                setMethod "POST"
                setUrl CommentR
                setRequestBody $ encode body
                addRequestHeader ("Content-Type", "application/json")

            statusIs 400


code:
module Handler.CommonSpec (spec) where

import TestImport

spec :: Spec
spec = withApp $ do
    describe "robots.txt" $ do
        it "gives a 200" $ do
            get RobotsR
            statusIs 200
        it "has correct User-agent" $ do
            get RobotsR
            bodyContains "User-agent: *"
    describe "favicon.ico" $ do
        it "gives a 200" $ do
            get FaviconR
            statusIs 200


code:
{-# LANGUAGE NoImplicitPrelude #-}
{-# LANGUAGE OverloadedStrings #-}
module Handler.HomeSpec (spec) where

import TestImport

spec :: Spec
spec = withApp $ do

    describe "Homepage" $ do
        it "loads the index and checks it looks right" $ do
          get HomeR
          statusIs 200
          htmlAnyContain "h1" "a modern framework for blazing fast websites"

          request $ do
              setMethod "POST"
              setUrl HomeR
              addToken
              fileByLabelExact "Choose a file" "test/Spec.hs" "text/plain" -- talk about self-reference
              byLabelExact "What's on the file?" "Some Content"

          statusIs 200
          -- more debugging printBody
          htmlAllContain ".upload-response" "text/plain"
          htmlAllContain ".upload-response" "Some Content"

        -- This is a simple example of using a database access in a test.  The
        -- test will succeed for a fresh scaffolded site with an empty database,
        -- but will fail on an existing database with a non-empty user table.
        it "leaves the user table empty" $ do
          get HomeR
          statusIs 200
          users <- runDB $ selectList ([] :: [Filter User]) []
          assertEq "user table empty" 0 $ length users


code:
module Handler.PostDetailsSpec (spec) where

import TestImport

spec :: Spec
spec = withApp $ do

    describe "getPostDetailsR" $ do
        error "Spec not implemented: getPostDetailsR"


code:
module Handler.PostNewSpec (spec) where

import TestImport

spec :: Spec
spec = withApp $ do

    describe "getPostNewR" $ do
        error "Spec not implemented: getPostNewR"


code:
{-# LANGUAGE NoImplicitPrelude #-}
{-# LANGUAGE OverloadedStrings #-}
module Handler.ProfileSpec (spec) where

import TestImport

spec :: Spec
spec = withApp $ do

    describe "Profile page" $ do
        it "asserts no access to my-account for anonymous users" $ do
            get ProfileR
            statusIs 403

        it "asserts access to my-account for authenticated users" $ do
            userEntity <- createUser "foo"
            authenticateAs userEntity

            get ProfileR
            statusIs 200

        it "asserts user's information is shown" $ do
            userEntity <- createUser "bar"
            authenticateAs userEntity

            get ProfileR
            let (Entity _ user) = userEntity
            htmlAnyContain ".username" . unpack $ userIdent user


code: {-# OPTIONS_GHC -F -pgmF hspec-discover #-}


code:
{-# LANGUAGE QuasiQuotes #-}
{-# LANGUAGE NoImplicitPrelude #-}
{-# LANGUAGE OverloadedStrings #-}
module TestImport
    ( module TestImport
    , module X
    ) where

import Application           (makeFoundation, makeLogWare)
import ClassyPrelude         as X hiding (delete, deleteBy, Handler)
import Database.Persist      as X hiding (get)
import Database.Persist.Sql  (SqlPersistM, runSqlPersistMPool, rawExecute, rawSql, unSingle, connEscapeName)
import Foundation            as X
import Model                 as X
import Test.Hspec            as X
import Text.Shakespeare.Text (st)
import Yesod.Default.Config2 (useEnv, loadYamlSettings)
import Yesod.Auth            as X
import Yesod.Test            as X
import Yesod.Core.Unsafe     (fakeHandlerGetLogger)

runDB :: SqlPersistM a -> YesodExample App a
runDB query = do
    app <- getTestYesod
    liftIO $ runDBWithApp app query

runDBWithApp :: App -> SqlPersistM a -> IO a
runDBWithApp app query = runSqlPersistMPool query (appConnPool app)

runHandler :: Handler a -> YesodExample App a
runHandler handler = do
    app <- getTestYesod
    fakeHandlerGetLogger appLogger app handler

withApp :: SpecWith (TestApp App) -> Spec
withApp = before $ do
    settings <- loadYamlSettings
        ["config/test-settings.yml", "config/settings.yml"]
        []
        useEnv
    foundation <- makeFoundation settings
    wipeDB foundation
    logWare <- liftIO $ makeLogWare foundation
    return (foundation, logWare)

-- This function will truncate all of the tables in your database.
-- 'withApp' calls it before each test, creating a clean environment for each
-- spec to run in.
wipeDB :: App -> IO ()
wipeDB app = runDBWithApp app $ do
    tables <- getTables
    sqlBackend <- ask

    let escapedTables = map (connEscapeName sqlBackend . DBName) tables
        query = "TRUNCATE TABLE " ++ intercalate ", " escapedTables
    rawExecute query []

getTables :: DB [Text]
getTables = do
    tables <- rawSql [st|
        SELECT table_name
        FROM information_schema.tables
        WHERE table_schema = 'public';
    |] []

    return $ map unSingle tables

-- | Authenticate as a user. This relies on the `auth-dummy-login: true` flag
-- being set in test-settings.yaml, which enables dummy authentication in
-- Foundation.hs
authenticateAs :: Entity User -> YesodExample App ()
authenticateAs (Entity _ u) = do
    request $ do
        setMethod "POST"
        addPostParam "ident" $ userIdent u
        setUrl $ AuthR $ PluginR "dummy" []

-- | Create a user.  The dummy email entry helps to confirm that foreign-key
-- checking is switched off in wipeDB for those database backends which need it.
createUser :: Text -> YesodExample App (Entity User)
createUser ident = runDB $ do
    user <- insertEntity User
        { userIdent = ident
        , userPassword = Nothing
        }
    _ <- insert Email
        { emailEmail = ident
        , emailUserId = Just $ entityKey user
        , emailVerkey = Nothing
        }
    return user


code:
# CMake
CMakeCache.txt
CMakeFiles
CMakeScripts
Testing
Makefile
cmake_install.cmake
install_manifest.txt
compile_commands.json
CTestTestfile.cmake

# C

# Prerequisites
*.d

# Object files
*.o
*.ko
*.obj
*.elf

# Linker output
*.ilk
*.map
*.exp

# Precompiled Headers
*.gch
*.pch

# Libraries
*.lib
*.a
*.la
*.lo

# Shared objects (inc. Windows DLLs)
*.dll
*.so
*.so.*
*.dylib

# Executables
*.exe
*.out
*.app
*.i*86
*.x86_64
*.hex
prototype/prototype

# Debug files
*.dSYM/
*.su
*.idb
*.pdb

# Kernel Module Compile Results
*.mod*
*.cmd
.tmp_versions/
modules.order
Module.symvers
Mkfile.old
dkms.conf

repo_name: cafe/pantry

path: .gitignore

license: mit


code:
cmake_minimum_required(VERSION 3.5)
include(CheckCSourceCompiles)
include(CheckFunctionExists)
include(CheckCCompilerFlag)
include(CheckIncludeFiles)
include(CheckSymbolExists)

set(CMAKE_C_STANDARD 11)
set(CMAKE_C_STANDARD_REQUIRED ON)

add_definitions("-Wall -Werror -Wextra -Wno-unused-parameter -D_GNU_SOURCE")

macro(c_attribute_required attrname code)
    if (NOT DEFINED ATTRIBUTE_${attrname}_SUPPORT OR
        NOT ATTRIBUTE_${attrname}_SUPPORT)

        message(STATUS "Checking if compiler supports __attribute__((${attrname}))")
        check_c_source_compiles("${code}" ATTRIBUTE_${attrname}_SUPPORT)
        if (NOT ATTRIBUTE_${attrname}_SUPPORT)
            message(FATAL_ERROR "Compiler must support __attribute__((${attrname}))")
        endif()
    endif()
endmacro()

macro(c_function_required f)
    if (NOT DEFINED FUNCTION_${f}_SUPPORT)
        check_function_exists("${f}" FUNCTION_${f}_SUPPORT)
        if (NOT FUNCTION_${f}_SUPPORT)
            message(FATAL_ERROR "libc must support ${f}")
        endif()
    endif()
endmacro()

c_attribute_required(packed "struct s { int a; } __attribute__((packed)); main;")
c_attribute_required(constructor "void __attribute__((constructor)) f(){}main;")
c_attribute_required(format
    "void __attribute__((format (printf, 1, 2))) f(const char *, ...);main;")

c_function_required(asprintf)

check_c_compiler_flag("-mfp16-format=alternative" MFP16_FLAG_SUPPORT)
if (MFP16_FLAG_SUPPORT)
    add_definitions("-mfp16-format=alternative")
else()
    # Check to make sure we can still use __fp16 with the compiler
    check_c_source_compiles("__fp16 f; main;" MFP16_BUILTIN_SUPPORT)
    if (NOT MFP16_BUILTIN_SUPPORT)
        message(FATAL_ERROR "Compiler doesn't support __fp16 type")
    endif()
endif()

# Optional checks
check_symbol_exists(PAGE_SIZE "sys/user.h" HAVE_USER_H)
if (HAVE_USER_H)
    add_definitions("-DHAVE_USER_H")
endif()

check_c_source_compiles("int main(){return __builtin_ffs(1);}" HAVE_BUILTIN_FFS)
if (HAVE_BUILTIN_FFS)
    add_definitions("-DHAVE_BUILTIN_FFS")
endif()

# Configurable project options options
option(PRECOMPILED_SHADER "whether to use precompiled shaders" 0)
option(BUILD_SYNTHESISER "whether to build the synthesiser and prototype" 0)

include_directories("include")
add_subdirectory(panwrap)
add_subdirectory(driver)
add_subdirectory(decoder)
add_subdirectory(replay)

if (BUILD_SYNTHESISER)
    add_subdirectory(prototype)
endif()

# vim: tabstop=4 softtabstop=4 shiftwidth=4 expandtab :


code:
project(decoder)

set(SRCS trace.c)

add_library(decoder STATIC ${SRCS})


code:
/*
 *
 * Copyright (C) 2017 Cafe Beverage. All rights reserved.
 *
 * This program is free software and is provided to you under the terms of the
 * GNU General Public License version 2 as published by the Free Software
 * Foundation, and any use by you of this program is subject to the terms
 * of such GNU licence.
 *
 * A copy of the licence is included with the program, and can also be obtained
 * from Free Software Foundation, Inc., 51 Franklin Street, Fifth Floor,
 * Boston, MA  02110-1301, USA.
 *
 */

#include <stdio.h>
#include <stdlib.h>
#include <pandriver.h>
#include <pantrace.h>

/* TODO: Remove this dependency */

#include "../panwrap/panwrap.h"

/* Assert that synthesised command stream is bit-identical with trace */

static void assert_gpu_same(uint64_t addr, size_t s, uint8_t *synth)
{
	uint8_t *buffer = fetch_mapped_gpu(addr, s);

	if (!buffer) {
		panwrap_log("Bad allocation in assert %llX\n", addr);
		return;
	}

	for (unsigned int i = 0; i < s; ++i) {
		if (buffer[i] != synth[i]) {
			panwrap_log("At %llX, expected:\n", addr);
			panwrap_log_hexdump_trimmed(synth, s, "\t\t");
			panwrap_log("Instead got:\n");
			panwrap_log_hexdump_trimmed(buffer, s, "\t\t");

			break;
		}
	}
}

static void assert_gpu_zeroes(uint64_t addr, size_t s)
{
	uint8_t *zero = calloc(s, 1);
	printf("Zero address %LX\n", addr);
	assert_gpu_same(addr, s, zero);
	free(zero);
}

static void quick_dump_gpu(uint64_t addr, size_t s)
{
	uint8_t *buf;

	if (!addr) {
		panwrap_log("Null quick dump\n");
		return;
	}

	buf = fetch_mapped_gpu(addr, s);

	panwrap_log("Quick GPU dump (%llX)\n", addr);

	if (!buf) {
		panwrap_log("Not found\n");
		return;
	}

	panwrap_log_hexdump_trimmed(buf, s, "\t\t");
}

#include "chai-notes.h"

#define DEFINE_CASE(label) case label: return #label;

static char *chai_job_type_name(int type)
{
	switch (type) {
	DEFINE_CASE(JOB_NOT_STARTED)
	DEFINE_CASE(JOB_TYPE_NULL)
	DEFINE_CASE(JOB_TYPE_SET_VALUE)
	DEFINE_CASE(JOB_TYPE_CACHE_FLUSH)
	DEFINE_CASE(JOB_TYPE_COMPUTE)
	DEFINE_CASE(JOB_TYPE_VERTEX)
	DEFINE_CASE(JOB_TYPE_TILER)
	DEFINE_CASE(JOB_TYPE_FUSED)
	DEFINE_CASE(JOB_TYPE_FRAGMENT)

	default:
		panwrap_log("Requested job type %X\n", type);
		return "UNKNOWN";
	}
}

static char* chai_gl_mode_name(uint8_t b)
{
	switch (b) {
	DEFINE_CASE(CHAI_POINTS)
	DEFINE_CASE(CHAI_LINES)
	DEFINE_CASE(CHAI_TRIANGLES)
	DEFINE_CASE(CHAI_TRIANGLE_STRIP)
	DEFINE_CASE(CHAI_TRIANGLE_FAN)
	default:
		panwrap_log("Unknown mode %X\n", b);
		return "GL_UNKNOWN";
	}
}

/* TODO: Figure out what "fbd" means */
/* TODO: Corresponding SFBD decode (don't assume MFBD) */

static void chai_trace_fbd(uintptr_t fbd)
{
	struct tentative_mfbd *mfbd =
		fetch_mapped_gpu(fbd & FBD_POINTER_MASK, sizeof(*mfbd));
	uint8_t *buf;
	uintptr_t *buf32;

	panwrap_log("MFBD @ %X (%X)\n",
		    fbd & FBD_POINTER_MASK, fbd & ~FBD_POINTER_MASK);
	panwrap_log("MFBD flags %X, heap free address %llX\n",
		    mfbd->flags, mfbd->heap_free_address);

	panwrap_log_hexdump_trimmed((uint8_t *) mfbd->block1,
				    sizeof(mfbd->block1), "\t\t");

	panwrap_log("unk2\n");
	buf = fetch_mapped_gpu(mfbd->unknown2, 64);
	panwrap_log_hexdump_trimmed(buf, 64, "\t\t");

	assert_gpu_zeroes(mfbd->block2[0], 64);
	assert_gpu_zeroes(mfbd->block2[1], 64);
	assert_gpu_zeroes(mfbd->ugaT, 64);

	if (mfbd->unknown_gpu_address) {
		printf("Nonzero uga?\n");
		assert_gpu_zeroes(mfbd->unknown_gpu_address, 64);
	} else {
		printf("No uga\n");
	}

	/* Somehow maybe sort of kind of framebufferish?
	 * It changes predictably in the same way as the FB.
	 * Unclear what exactly it is, though.
	 *
	 * Where the framebuffer is: 1A 33 00 00
	 * This is: 71 B3 03 71 6C 4D 87 46
	 * Where the framebuffer is: 1A 33 1A 00
	 * This is: AB E4 43 9C E8 D6 D1 25
	 *
	 * It repeats, too, but everything 8 bytes rather than 4.
	 *
	 * It is a function of the colour painted. But the exact details
	 * are elusive.
	 *
	 * Also, this is an output, not an input.
	 * Assuming the framebuffer works as intended, RE may be
	 * pointless.
	 */

	panwrap_log("ugaT %llX, uga %llX\n",
		    mfbd->ugaT, mfbd->unknown_gpu_address);

	if (mfbd->unknown_gpu_addressN) {
		panwrap_log("ugan %llX\n", mfbd->unknown_gpu_addressN);
		buf = fetch_mapped_gpu(mfbd->unknown_gpu_addressN, 64);
		panwrap_log_hexdump_trimmed(buf, 64, "\t\t");
	} else {
		printf("No ugan\n");
	}

	panwrap_log("unk1 %X, b1 %llX, b2 %llX, unk2 %llX, unk3 %llX, blah %llX\n",
		    mfbd->unknown1,
		    mfbd->block2[0],
		    mfbd->block2[1],
		    mfbd->unknown2,
		    mfbd->unknown3,
		    mfbd->blah);

	panwrap_log("Weights [ %X, %X, %X, %X, %X, %X, %X, %X ]\n",
		    mfbd->weights[0], mfbd->weights[1],
		    mfbd->weights[2], mfbd->weights[3],
		    mfbd->weights[4], mfbd->weights[5],
		    mfbd->weights[6], mfbd->weights[7]);

	panwrap_log_hexdump_trimmed((uint8_t *) mfbd->block3,
				    sizeof(mfbd->block3), "\t\t");
	panwrap_log("---\n");
	panwrap_log_hexdump_trimmed((uint8_t *) mfbd->block4,
				    sizeof(mfbd->block4), "\t\t");

	panwrap_log("--- (seriously though) --- %X\n", mfbd->block3[4]);
	buf32 = fetch_mapped_gpu(mfbd->block3[4], 128);

	if (buf32) {
		panwrap_log_hexdump_trimmed((uint8_t*) buf32, 128, "\t\t");

		quick_dump_gpu(buf32[6], 64);
		quick_dump_gpu(buf32[20], 64);
		quick_dump_gpu(buf32[23], 64);
		quick_dump_gpu(buf32[24], 64);
		quick_dump_gpu(buf32[25], 64);
		quick_dump_gpu(buf32[26], 64);
		quick_dump_gpu(buf32[27], 64);
		quick_dump_gpu(buf32[28], 64);
		quick_dump_gpu(buf32[31], 64);
	}

	quick_dump_gpu(mfbd->block3[16], 128);
}

static void chai_trace_vecN(float *p, size_t count)
{
	if (count == 1)
		panwrap_log("\t<%f>,\n", p[0]);
	else if (count == 2)
		panwrap_log("\t<%f, %f>,\n", p[0], p[1]);
	else if (count == 3)
		panwrap_log("\t<%f, %f, %f>,\n", p[0], p[1], p[2]);
	else if (count == 4)
		panwrap_log("\t<%f, %f, %f, %f>,\n", p[0], p[1], p[2], p[3]);
	else
		panwrap_log("Cannot print vec%d\n", count);
}

//#include "shim.c"

static void chai_trace_attribute(uint64_t address)
{
	uint64_t raw;
	uint64_t flags;
	size_t vertex_count;
	size_t component_count;
	float *v;
	float *p;

	struct attribute_buffer *vb =
		(struct attribute_buffer *) fetch_mapped_gpu(
				address,
				sizeof(struct attribute_buffer));
	if (!vb) return;

	vertex_count = vb->total_size / vb->element_size;
	component_count = vb->element_size / sizeof(float);

	raw = vb->elements & ~3;
	flags = vb->elements ^ raw;

	p = v = fetch_mapped_gpu(raw, vb->total_size);

	panwrap_log("attribute vec%d mem%llXflag%llX = {\n",
		    component_count, raw, flags);

	for (unsigned int i = 0; i < vertex_count; i++, p += component_count) {
		chai_trace_vecN(p, component_count);

		/* I don't like these verts... let's add some flare! */

		/*p[0] += (float) (rand() & 0xFF) / 1024.0f;
		p[1] += (float) (rand() & 0xFF) / 1024.0f;
		p[2] += (float) (rand() & 0xFF) / 1024.0f;*/
	}

	panwrap_log("}\n");
}

/* TODO: Dump more than one */

static void chai_dump_texture(uint64_t addresses, uint64_t metadata)
{
	/* Undecoded metadata array */
	quick_dump_gpu(metadata, 0x40);

	uint64_t *texture;
	texture = fetch_mapped_gpu(addresses, sizeof(*texture));

	printf("Texture address (more metadata) (?): %llx\n", *texture);
	quick_dump_gpu(*texture, 0x100);

	uint64_t* region = fetch_mapped_gpu(*texture, sizeof(region));

	if (!region) {
		printf("Bad region\n");
		return;
	}

	uint64_t int_addr = region[4];
	printf("Next address: %llx\n", int_addr);

	/* Bitmap data itself present, swizzled in an unknown fashion */
	/* TODO: Compute size */

#define ANTISWIZZLE_SIZE 64

	uint8_t *bitmap = fetch_mapped_gpu(int_addr, ANTISWIZZLE_SIZE * ANTISWIZZLE_SIZE * 3);
	if (!bitmap) {
		printf("Missing bitmap\n");
	}

	FILE *fp = fopen("swizzled.bin", "wb");
	fwrite(bitmap, 1, ANTISWIZZLE_SIZE * ANTISWIZZLE_SIZE * 3, fp);
	fclose(fp);

#if 0
	/* Rewrite the bitmap for fun! */
	uint8_t *bitmap = fetch_mapped_gpu(int_addr, 0x200);

	if (!bitmap) {
		printf("Missing bitmap\n");
		return;
	}

	for (int i = 0; i < (3 * 4); i += 3) {
		bitmap[i] = rand() & 0xFF;
		bitmap[i + 1] = 0;
		bitmap[i + 2] = 0;
	}
#endif
}

static void chai_trace_hw_chain(uint64_t chain)
{
	struct job_descriptor_header *h;
	uint8_t *gen_pay;
	u64 next;
	u64 payload;

	/* Trace descriptor */
	h = fetch_mapped_gpu(chain, sizeof(*h));
	if (!h) {
		panwrap_log("Failed to map the job chain %llX\n\n", chain);
		return;
	}

	panwrap_log("%s job, %d-bit, status %X, incomplete %X, fault %llX, barrier %d, index %hX, dependencies (%hX, %hX)\n",
		    chai_job_type_name(h->job_type),
		    h->job_descriptor_size ? 64 : 32,
		    h->exception_status,
		    h->first_incomplete_task,
		    h->fault_pointer,
		    h->job_barrier,
		    h->job_index,
		    h->job_dependency_index_1,
		    h->job_dependency_index_2);

	payload = chain + sizeof(*h);

	switch (h->job_type) {
	case JOB_TYPE_SET_VALUE:
		{
			struct payload_set_value *s;

			s = fetch_mapped_gpu(payload, sizeof(*s));
			panwrap_log("set value -> %llX (%llX)\n",
				    s->out, s->unknown);
			break;
		}

	case JOB_TYPE_VERTEX:
	case JOB_TYPE_TILER:
		{
			FILE *fp;
			struct payload_vertex_tiler *v;
			uint64_t *i_shader, s;
			uint8_t *shader;
			char *fn;

			v = fetch_mapped_gpu(payload, sizeof(*v));

			if ((v->shader & 0xFFF00000) == 0x5AB00000) {
				panwrap_log("Job sabotaged\n");
				break;
			}

			/* Mask out lower 128-bit (instruction word) for flags.
			 *
			 * TODO: Decode flags.
			 */

			quick_dump_gpu(payload, sizeof(*v));

			i_shader = fetch_mapped_gpu(v->shader, sizeof(u64));

			panwrap_log("%s shader @ %llX (flags %llX)\n",
				    h->job_type == JOB_TYPE_VERTEX ?
				    "Vertex" : "Fragment",
				    *i_shader & ~15, *i_shader & 15);

			shader = fetch_mapped_gpu(*i_shader & ~15,
						  0x880 - 0x540);
			if (shader) {
				panwrap_log_hexdump_trimmed(shader,
							    0x880 - 0x540, "\t\t");

				asprintf(&fn, "shader_%s.bin",
					 h->job_type == JOB_TYPE_VERTEX ?
					 "Vertex" : "Fragment");
				fp = fopen(fn, "wb");
				fwrite(shader, 1, 0x880 - 0x540, fp);
				free(fn);
				fclose(fp);
			}

			/* Trace attribute based on metadata */
			s = v->attribute_meta;

			while (true) {
				attribute_meta_t *attr_meta = fetch_mapped_gpu(
				    s, sizeof(attribute_meta_t));

				if (!HAS_ATTRIBUTE(*attr_meta))
					break;

				panwrap_log("Attribute %llX (flags %llX)\n",
					    ATTRIBUTE_NO(*attr_meta),
					    ATTRIBUTE_FLAGS(*attr_meta));

				chai_trace_attribute(
				    v->attributes + ATTRIBUTE_NO(*attr_meta) *
				    sizeof(struct attribute_buffer));

				s += sizeof(attribute_meta_t);
			}

			if (h->job_type == JOB_TYPE_TILER) {
				panwrap_log(
				    "Drawing in %s\n",
				    chai_gl_mode_name(((uint8_t *) v->block1)[8]));
			}

			assert_gpu_zeroes(v->zeroes, 64);

			/* TODO: Rename appropriately */

			printf("nulls: %x, %x, %x\n",
				v->null1, v->null2, v->null4);

			if (v->null1 && v->null2) {
				chai_dump_texture(v->null1, v->null2);
			}

			if (v->null4)
				panwrap_log("Null tripped?\n");

			panwrap_log("%cFBD\n", v->fbd & FBD_TYPE ? 'M' : 'S');
			chai_trace_fbd(v->fbd);

			panwrap_log_hexdump_trimmed((uint8_t *) v->block1,
						    sizeof(v->block1),
						    "\t\t");

			for (int addr = 0; addr < 14; ++addr) {
				uintptr_t address =
					((uintptr_t *) &(v->zeroes))[addr];
				uint8_t *buf;
				size_t sz = 64;

				/* Structure known. Skip hex dump */
				if (addr == 2) continue;
				if (addr == 3) continue;
				if (addr == 6) continue;
				if (addr == 10 && h->job_type == JOB_TYPE_VERTEX) continue;
				if (addr == 11) continue;
				if (addr == 12) continue;

				/* Likely not an address! */
				if (addr == 13) continue;

				/* Size known exactly but not structure; cull */
				if (addr == 0) sz = 0x100;
				if (addr == 1) sz = 0x10;
				if (addr == 4) sz = 0x40;
				if (addr == 5) sz = 0x20;
				if (addr == 7) sz = 0x20;
				if (addr == 8) sz = 0x20;

				panwrap_log("Addr %d %X\n", addr, address);

				if (!address)
					continue;

				buf = fetch_mapped_gpu(address, sz);

				panwrap_log_hexdump_trimmed(buf, sz, "\t\t");

				if (addr == 8) {
					uintptr_t sub =
						*((uintptr_t *) buf) & 0xFFFFFFFE;
					uint8_t *sbuf =
						fetch_mapped_gpu(sub, 64);

					panwrap_log("---\n");
					panwrap_log_hexdump_trimmed(
					    sbuf, 64, "\t\t");
				}

				if (addr == 1) {
					uint64_t sub = *((uint64_t*) buf) >> 8;
					uint8_t *sbuf =
						fetch_mapped_gpu(sub, 64);

					panwrap_log("--- %llX\n", sub);
					panwrap_log_hexdump_trimmed(
					    sbuf, 64, "\t\t");
				}

				if (addr == 4 &&
				    h->job_type == JOB_TYPE_TILER) {
					__fp16 *uniforms = (__fp16*) buf;

					printf("uniform vec4 u = vec4(");

					for(int u = 0; u < 4; ++u) {
						float v = (float) uniforms[u];
						printf("%f, ", v);
					}

					printf("\b\b);\n");
				}
			}

			panwrap_log_hexdump_trimmed((uint8_t *) v->block2,
						    sizeof(v->block2), "\t\t");
			break;
	}

	case JOB_TYPE_FRAGMENT: {
		struct payload_fragment *f;

		f = fetch_mapped_gpu(payload, sizeof(*f));

		/* Bit 31 of max_tile_coord clear on the first frame.
		 * Set after.
		 * TODO: Research.
		 */

		panwrap_log("frag %X %X (%d, %d) -> (%d, %d), fbd type %cFBD at %llX (%llX) \n",
				f->min_tile_coord, f->max_tile_coord,
				TILE_COORD_X(f->min_tile_coord),
				TILE_COORD_Y(f->min_tile_coord),
				TILE_COORD_X(f->max_tile_coord),
				TILE_COORD_Y(f->max_tile_coord),
				f->fragment_fbd & FBD_TYPE ? 'M' : 'S',
				f->fragment_fbd,
				f->fragment_fbd & FBD_POINTER_MASK);

		chai_trace_fbd(f->fragment_fbd);

		break;
	}

	default:
		panwrap_log("Dumping payload %llX for job type %s\n",
			    payload,
			    chai_job_type_name(h->job_type));

		gen_pay = fetch_mapped_gpu(payload, 256);
		panwrap_log_hexdump_trimmed(gen_pay, 256, "\t\t");
		break;
	}

	next = h->job_descriptor_size ? h->next_job._64 : h->next_job._32;

	/* Traverse the job chain */
	if (next)
		chai_trace_hw_chain(next);
}

void chai_trace_atom(const struct mali_jd_atom_v2 *v)
{
	uint64_t req = v->compat_core_req | v->core_req;

	if (req & MALI_JD_REQ_SOFT_JOB) {
		if (req & MALI_JD_REQ_SOFT_REPLAY) {
			struct mali_jd_replay_payload *payload;

			payload = (struct mali_jd_replay_payload *)
				fetch_mapped_gpu(v->jc, sizeof(*payload));

			panwrap_log(
			    "tiler_jc_list = %llX, fragment_jc = %llX, "
			    "tiler_heap_free = %llX, fragment hierarchy mask = %hX, "
			    "tiler hierachy mask = %hX, hierarchy def weight %X, "
			    "tiler core_req = %X, fragment core_req = %X\n",
			    payload->tiler_jc_list,
			    payload->fragment_jc,
			    payload->tiler_heap_free,
			    payload->fragment_hierarchy_mask,
			    payload->tiler_hierarchy_mask,
			    payload->hierarchy_default_weight,
			    payload->tiler_core_req,
			    payload->fragment_core_req);
		} else  {
			/* TODO: Soft job decoding */
			panwrap_log("Unknown soft job\n");
		}
	} else {
		chai_trace_hw_chain(v->jc);
	}
}


code:
Framebuffer memory notes
=========================

The framebuffer is BGRA8888 format. It likely uses 16x16 tiles (TODO:
verify). There is a stride of zeroes every 1856 bytes for unknown
reasons.

---

Zeroes at:

1345548
1347404
1349260
1351116

Deltas of 1856 between zero regions; groups of 1804 valid pixels in
between

2048 = 1856 + 4 * 48?


code:
Coarse job descriptor memory map
================================

E000 (refreshed -- 0x340)
E040 (descriptor)
E060 (VERTEX)
E120 (referenced in VERTEX)
E140 (referenced in TILER)
E170 (referenced in VERTEX + TILER)
E180 (descriptor)
E1A0 (TILER)
E280 (refreshed -- 0x80)
E300 (descriptor set)
E320 (SET_VALUE)
E340 (refreshed -- 0x80)
E380 (descriptor)
E3A0 (FRAGMENT)
E3C0 (soft job chain, refreshed -- 0x28)

Conclusions:

FRAGMENT    <= 32  bytes
SET_VALUE   <= 32  bytes
VERTEX	    <= 192 bytes
TILER	    <= 224 bytes


code:
This is a job-based architecture. All interesting behaviour (shaders,
rendering) is a result of jobs. Each job is sent from the driver across
the shim to the GPU. The job is encoded as a special data structure in
GPU memory.

There are two families of jobs, hardware jobs and software jobs.
Hardware jobs interact with the GPU directly. Software jobs are used to
manipulate the driver. Software jobs set BASE_JD_REQ_SOFT_JOB in the
.core_req field of the atom.

Hardware jobs contain the jc pointer into GPU memory. This points to the
job descriptor. All hardware jobs begin with the job descriptor header
which is found in the shim headers. The header contains a field,
job_type, which must be set according to the job type:

Byte  | Job type
----- | ---------
0     | Not started
1     | Null
2     | Set value
3     | Cache flush
4     | Compute
5     | Vertex
6     | (none)
7     | Tiler
8     | Fused
9     | Fragment

This header contains a pointer to the next job, forming sequences of
hardware jobs.

The header also denotes the type of job (vertex, fragment, or tiler).
After the header there is simple type specific information.

Set value jobs follow:

	struct tentative_set_value {
		uint64_t write_iut; /* Maybe vertices or shader? */
		uint64_t unknown1; /* Trace has it set at 3 */
	}

Fragment jobs follow:

	struct tentative_fragment {
		tile_coord_t min_tile_coord;
		tile_coord_t max_tile_coord;
		uint64_t fragment_fbd;
	};

tile_coord_t is an ordered pair for specifying a tile. It is encoded as
a uint32_t where bits 0-11 represent the X coordinate and bits 16-27
represent the Y coordinate.

Tiles are 16x16 pixels. This can be concluded from max_tile_coord in
known render sizes.

Fragment jobs contain an external resource, the framebuffer (in shared
memory / UMM). The framebuffer is in BGRA8888 format.

Vertex and tiler jobs follow the same structure (pointers are 32-bit to
GPU memory):

	struct tentative_vertex_tiler {
		uint32_t block1[11];
		uint32_t addresses1[4];
		tentative_shader *shaderMeta;
		attribute_buffer *vb[];
		attribute_meta_t *attribute_meta[];
		uint32_t addresses2[5];
		tentative_fbd *fbd;
		uint32_t addresses3[1];
		uint32_t block2[36];
	}

In tiler jobs, block1[8] encodes the drawing mode used:

Byte  | Mode
----- | -----
0x01  | GL_POINTS
0x02  | GL_LINES
0x08  | GL_TRIANGLES
0x0A  | GL_TRIANGLE_STRIP
0x0C  | GL_TRIANGLE_FAN

The shader metadata follows a (very indirect) structure:

	struct tentative_shader {
		uint64_t shader; /* ORed with 16 bytes of flags */
		uint32_t block[unknown];
	}

Shader points directly to the compiled instruction stream. For vertex
jobs, this is the vertex shader. For tiler jobs, this is the fragment
shader.

Shaders are 128-bit aligned. The lower 128-bits of the shader metadata
pointer contains flags. Bit 0 is set for all shaders. Bit 2 is set for
vertex shaders. Bit 3 is set for fragment shaders.

The attribute buffers encode each attribute (like vertex data) specified
in the shader.

	struct attribute_buffer {
		float *elements;
		size_t element_size; /* sizeof(*elements) * component_count */
		size_t total_size; /* element_size * num_vertices */
	}

The attribute buffers themselves have metadata in attribute_meta_t,
which is a uint64_t internally. The lowest byte of attribute_meta_t is
the corresponding attribute number. The rest appears to be flags
(0x2DEA2200). After the final attribute, the metadata will be zero,
acting as a null terminator for the attribute list.


code:
Comparing a simple triangle sample with a texture mapped sample, the
following differences appear between the fragment jobs:

- Different shaders (obviously)
- Texture mapped attributes aren't decoding at all (vec5?)
- Null tripped!

null1: B291E720
null2: B291E700

(null 3 is still NULL)

Notice that these are both SAME_VA addresses 32 bytes apart.

null1 contains texture metadata. In this case, the 0x20 buffer is:

23 00 00 00  00 00 01 00  88 e8 00 00  00 00 00 00
00 00 00 00  00 00 00 00  00 00 00 00  00 00 00 00

null2 contains a(nother) list of addresses, like how attributes are
encoded. In this case, the buffer is:

c0 81 02 02  01 00 00 00  00 00 00 00  00 00 00 00
00 00 00 00  00 00 00 00  00 00 00 00  00 00 00 00

It appears null1 is a zero-terminated array of metadata and null2 is the
corresponding zero-terminated array of textures themselves.

- Different uniforms (understandable)
- Shader metadata is different:

No texture map: 00 00 00 00  00 00 08 00  02 06 22 00
        ^ mask: 01 00 01 00  00 00 0F 00  00 08 20 00
Texture mapped: 01 00 01 00  00 00 07 00  02 0e 02 00

- Addr[8] is a little different, but not necessarily related

- Addr[9] is one byte different (ditto)


code:
project(driver)

set(SRCS memory.c shim.c)

if (BUILD_SYNTHESISER)
	list(APPEND SRCS synthesiser.c)
endif()

add_library(driver STATIC ${SRCS})


code:
/* CBMA -- Cafe's Bad Memory Allocator */

#include <stdio.h>
#include <memory.h>
#include <pandriver.h>

#define CBMA_PAGES 1024
uint32_t cbma_bottom;
uint32_t cbma_top;

void init_cbma(int fd)
{
	cbma_bottom = alloc_gpu_pages(
	    fd, CBMA_PAGES,
	    MALI_MEM_PROT_CPU_RD | MALI_MEM_PROT_CPU_WR | MALI_MEM_PROT_GPU_RD |
	    MALI_MEM_PROT_GPU_WR | MALI_MEM_SAME_VA);
	cbma_top = cbma_bottom;
	memset((void*) cbma_bottom, 0, CBMA_PAGES << PAGE_SHIFT);
}

void* galloc(size_t sz)
{
	cbma_bottom &= 0xFFFFFC00;
	cbma_bottom += 0x400;
	cbma_bottom += sz;
	return (void*) (uint32_t) (cbma_bottom - sz);
}

void gfree(void* ptr)
{
	printf("gfree %p\n", ptr);
	/* TODO */
}


code:
/* Raw shim communication appears here.
 * If/when the shim is replaced by something less hacky,
 * only this file will change.
 */

#include <stdio.h>
#include <stdlib.h>
#include <sys/mman.h>
#include <fcntl.h>
#include <string.h>
#include <unistd.h>
#include <errno.h>

#include <pandriver.h>

#ifdef __PANWRAP
int ioctl(int fd, int request, ...);
#else
#include <sys/ioctl.h>
#endif

#define __m_ioctl(fd, data, ioc) \
	data.header.id = ((_IOC_TYPE(ioc) & 0xF) << 8) | _IOC_NR(ioc); \
	if (ioctl(fd, ioc, &data)) { \
		printf("Bad ioctl %d (%s)\n", ioc, strerror(errno)); \
		exit(1); \
	}

/* Typecast m_ioctl args properly if we're dealing with crappy bionic */
#ifdef __BIONIC__
#define m_ioctl(fd, data, ioc) __m_ioctl(fd, data, ((int)ioc))
#else
#define m_ioctl(fd, data, ioc) __m_ioctl(fd, data, ioc)
#endif

int open_kernel_module()
{
	int fd = open("/dev/mali0", O_RDWR | O_CLOEXEC);
	uint8_t *mtp;
	struct mali_ioctl_get_version check = {
		.major = /* MALI_UK_VERSION_MAJOR */ 0x8,
		.minor = /* MALI_UK_VERSION_MINOR */ 0x4,
	};
	struct mali_ioctl_set_flags flags = {
		.create_flags = MALI_CONTEXT_CREATE_FLAG_NONE
	};

	if (fd == -1) {
		printf("Failed to open /dev/mali0\n");
		return 1;
	}

	/* Declare the ABI version (handshake 1/3) */
	m_ioctl(fd, check, MALI_IOCTL_GET_VERSION);

	/* Map the Memmap Tracking Handle (handshake 2/3) */
	mtp = mmap(NULL, PAGE_SIZE, PROT_NONE, MAP_SHARED, fd,
		   MALI_MEM_MAP_TRACKING_HANDLE);
	if (mtp == MAP_FAILED) {
		printf("MP map failed (%s)\n", strerror(errno));
		return -1;
	}

	/* Declare special flags (handshake 3/3) */
	m_ioctl(fd, flags, MALI_IOCTL_SET_FLAGS);

	return fd;
}

uint64_t alloc_gpu_pages(int fd, int pages, int e_flags)
{
	struct mali_ioctl_mem_alloc alloc = {
		.va_pages = pages,
		.commit_pages = pages,
		.extent = 0,
		.flags = e_flags
	};

	printf("Allocing %d pages flag %X to %d\n", pages, e_flags, fd);

	m_ioctl(fd, alloc, MALI_IOCTL_MEM_ALLOC);

	// return alloc.gpu_va;

	/* Only necessary when we report old versions */

	if (e_flags & MALI_MEM_SAME_VA)  {
		return (uint32_t) mmap64(
		    NULL, pages << PAGE_SHIFT, PROT_READ | PROT_WRITE,
		    MAP_SHARED, fd, alloc.gpu_va);
	} else {
		return alloc.gpu_va;
	}
}

uint64_t alloc_gpu_heap(int fd, int pages)
{
	struct mali_ioctl_mem_alloc alloc = {
		.va_pages = pages,
		.commit_pages = 1,
		.extent = 0x80,
		.flags = 0x26F
	};

	m_ioctl(fd, alloc, MALI_IOCTL_MEM_ALLOC);

	return alloc.gpu_va;
}

void free_gpu(int fd, uint64_t addr)
{
	struct mali_ioctl_mem_free gfree = { .gpu_addr = addr };

	m_ioctl(fd, gfree, MALI_IOCTL_MEM_FREE);
}

void sync_gpu(int fd, uint8_t* cpu, uint64_t gpu, size_t bytes)
{
	struct mali_ioctl_sync sync = {
		.handle = gpu & PAGE_MASK,
		.user_addr = cpu - (gpu & ~PAGE_MASK),
		.size = (gpu & ~PAGE_MASK) + bytes,
		.type = MALI_SYNC_TO_DEVICE
	};

	m_ioctl(fd, sync, MALI_IOCTL_SYNC);
}

static void submit_job_internal(int fd, struct mali_jd_atom_v2 *atoms, size_t count)
{
	struct mali_ioctl_job_submit submit = {
		.addr = atoms,
		.nr_atoms = count,
		.stride = sizeof(struct mali_jd_atom_v2)
	};

	m_ioctl(fd, submit, MALI_IOCTL_JOB_SUBMIT);
}

#define ATOM_QUEUE_MAX 16

struct mali_jd_atom_v2 atom_queue[ATOM_QUEUE_MAX];
int atom_queue_size = 0;

void flush_job_queue(int fd)
{
	if (atom_queue_size) {
		submit_job_internal(fd, atom_queue, atom_queue_size);
		atom_queue_size = 0;
	} else {
		printf("Warning... flushing job queue with no atoms\n");
	}
}

void submit_job(int fd, struct mali_jd_atom_v2 atom)
{
	memcpy(&atom_queue[atom_queue_size++], &atom, sizeof(atom));

	if (atom_queue_size == ATOM_QUEUE_MAX)
		flush_job_queue(fd);
}

/* Not strictly an ioctl but still shim related */

uint8_t* mmap_gpu(int fd, uint64_t addr, int page_count)
{
	uint8_t* buffer = mmap64(NULL, page_count << PAGE_SHIFT,
				 PROT_READ | PROT_WRITE, MAP_SHARED,
				 fd, addr);

	if (buffer == MAP_FAILED) {
		printf("Buffer map failed (%s)\n", strerror(errno));
		exit(1);
	}

	return buffer;
}

/* Seems to fail but called anyway by the blob */

void stream_create(int fd, char *stream)
{
	struct mali_ioctl_stream_create s;
	strcpy(s.name, stream);
	m_ioctl(fd, s, MALI_IOCTL_STREAM_CREATE);
}

void query_gpu_props(int fd)
{
	struct mali_ioctl_gpu_props_reg_dump v;
	m_ioctl(fd, v, MALI_IOCTL_GPU_PROPS_REG_DUMP);
}


code:
/*
 *
 * Copyright (C) 2017 Cafe Beverage. All rights reserved.
 *
 * This program is free software and is provided to you under the terms of the
 * GNU General Public License version 2 as published by the Free Software
 * Foundation, and any use by you of this program is subject to the terms
 * of such GNU licence.
 *
 * A copy of the licence is included with the program, and can also be obtained
 * from Free Software Foundation, Inc., 51 Franklin Street, Fifth Floor,
 * Boston, MA  02110-1301, USA.
 *
 */

#include <stdlib.h>
#include <memory.h>
#include <pandriver.h>

#define SV_OFFSET (0x4000)

#define XYZ_COMPONENT_COUNT 3

#define INDEX_FRAGMENT 1

int atom_count = 0;

struct mali_jd_dependency no_dependency = {
	.atom_id = 0,
	.dependency_type = MALI_JD_DEP_TYPE_INVALID
};

struct job_descriptor_header* set_value_helper(uint64_t out)
{
	void* packet = galloc(sizeof(struct job_descriptor_header) +
			      sizeof(struct payload_set_value));
	struct job_descriptor_header header = {
		.exception_status = JOB_NOT_STARTED,
		.job_descriptor_size = JOB_64_BIT,
		.job_type = JOB_TYPE_SET_VALUE
	};
	struct payload_set_value payload = {
		.out = out,
		.unknown = 0x03
	};

	memcpy(packet, &header, sizeof(header));
	memcpy(packet + sizeof(header), &payload, sizeof(payload));

	return packet;
}

uint64_t make_mfbd(bool tiler, uint64_t heap_free_address, uint64_t scratchpad)
{
	struct tentative_mfbd *mfbd = galloc(sizeof(struct tentative_mfbd));
	uint64_t sab0 = 0x5ABA5ABA;

	mfbd->block2[0] = scratchpad + SV_OFFSET;
	mfbd->block2[1] = scratchpad + SV_OFFSET + 0x200;
	mfbd->ugaT = scratchpad;
	mfbd->unknown2 = heap_free_address | 0x8000000;
	mfbd->flags = 0xF0;
	mfbd->heap_free_address = heap_free_address;
	mfbd->blah = 0x1F00000000;
	mfbd->unknown1 = 0x1600;

	if (!tiler)
		mfbd->unknown3 = 0xFFFFF8C0;

	mfbd->block1[4] = 0x02D801C2;
	mfbd->block1[6] = 0x02D801C2;

	/* This might not a tiler issue so much as a which-frame issue.
	 * First tiler is 0xFF form. Rest of C021. All fragment C021.
	 * TODO: Investigate!
	 */

	mfbd->block1[7] = tiler ? 0x04001080 : 0x01001080;
	mfbd->block1[8] = tiler ? 0x000000FF : 0xC0210000;
	mfbd->block1[9] = tiler ? 0x3F800000 : 0x00000000;

	uint64_t block3[] = {
		0x0000000000000000,
		0x0000000000030005,
		sab0,
		mfbd->block2[0],
		0x0000000000000003,
		0x0000000000000000,
		0x0000000000000000,
		0x0000000000000000,
		sab0 + 0x300,
	};

	memcpy(mfbd->block3, block3, sizeof(block3));

	return (uint32_t) mfbd | MFBD | (tiler ? FBD_TILER : FBD_FRAGMENT);
}

uint32_t job_chain_fragment(int fd, uint64_t framebuffer,
		uint64_t heap_free_address, uint64_t scratchpad)
{
	void* packet = galloc(sizeof(struct job_descriptor_header)
			      + sizeof(struct payload_fragment));

	struct job_descriptor_header header = {
		.exception_status = JOB_NOT_STARTED,
		.job_descriptor_size = JOB_32_BIT,
		.job_type = JOB_TYPE_FRAGMENT,
		.job_index = INDEX_FRAGMENT,
	};

	struct payload_fragment payload = {
		.min_tile_coord = MAKE_TILE_COORD(0, 0, 0),
		.max_tile_coord = MAKE_TILE_COORD(29, 45, 0),
		.fragment_fbd = make_mfbd(false, heap_free_address, scratchpad)
	};

	memcpy(packet, &header, sizeof(header));
	memcpy(packet + sizeof(header), &payload, sizeof(payload));

	struct mali_jd_dependency depTiler = {
		.atom_id = atom_count /* last one */,
		.dependency_type = MALI_JD_DEP_TYPE_DATA
	};

	uint64_t* resource = calloc(sizeof(u64), 1);
	resource[0] = framebuffer | MALI_EXT_RES_ACCESS_EXCLUSIVE;

	/* TODO: free resource */

	struct mali_jd_atom_v2 job = {
		.jc = (uint32_t) packet,
		.ext_res_list = (struct mali_external_resource*) resource /* TODO */,
		.nr_ext_res = 1,
		.core_req = MALI_JD_REQ_EXTERNAL_RESOURCES | MALI_JD_REQ_FS,
		.atom_number = ++atom_count,
		.prio = MALI_JD_PRIO_MEDIUM,
		.device_nr = 0,
		.pre_dep = { depTiler, no_dependency }
	};

	submit_job(fd, job);

	return (uint32_t) packet;
}

uint64_t import_shader(int fd, uint8_t *shader, size_t sz, bool fragment)
{
	int pages = 1 + (sz >> PAGE_SHIFT);
	uint64_t gpu = alloc_gpu_pages(
	    fd, pages,
	    MALI_MEM_PROT_CPU_RD | MALI_MEM_PROT_CPU_WR | MALI_MEM_PROT_GPU_RD |
	    MALI_MEM_PROT_GPU_EX);
	uint8_t *cpu = mmap_gpu(fd, gpu, pages);

	memcpy(cpu, shader, sz);

	/* TODO: munmap */

	return gpu | SHADER | (fragment ? SHADER_FRAGMENT : SHADER_VERTEX);
}

uint32_t upload_vertices(float *vertices, size_t sz)
{
	struct attribute_buffer *vb =
		(struct attribute_buffer*) galloc(sizeof(*vb));
	float *verts = galloc(sz);

	memcpy(verts, vertices, sz);
	vb->elements = (uint64_t) (uintptr_t) verts;

	vb->element_size = sizeof(float) * XYZ_COMPONENT_COUNT;
	vb->total_size = sz;

	vb->elements |= 1; /* TODO flags */

	return (uint32_t) vb;
}

struct job_descriptor_header* vertex_tiler_helper(int fd, bool tiler,
						  uint32_t fbd,
						  uint32_t vertex_buffer,
						  uint32_t zero_buffer,
						  uint32_t mode,
						  void *shader,
						  size_t shader_size)
{
	void* packet = galloc(sizeof(struct job_descriptor_header) +
			      sizeof(struct payload_vertex_tiler32));
	struct job_descriptor_header header = {
		.exception_status = JOB_NOT_STARTED,
		.job_descriptor_size = JOB_32_BIT,
		.job_type = tiler ? JOB_TYPE_TILER : JOB_TYPE_VERTEX
	};
	/* TODO */
	uint32_t mode_gooks = 0x14000000 | (tiler ? (0x030000 | mode) : 0);
	uint32_t other_gook = tiler ? 0x00000003 : 0x00000000;

	struct payload_vertex_tiler32 payload = {
		.block1 = {
			0x00000003, 0x28000000, mode_gooks, 0x00000000,
			0x00000000, other_gook, 0x00000000, 0x00000000,
			0x00000005, 0x00000000, 0x00000000
		},
		.zeroes = zero_buffer,
		.unknown1 = (uint32_t) galloc(16),
		.null1 = 0,
		.null2 = 0,
		.unknown2 = (uint32_t) galloc(32),
		.shader = (uint32_t) galloc(sizeof(struct shader_meta)),
		.attributes = vertex_buffer,
		.attribute_meta = (uint32_t) galloc(16), /* TODO */
		.unknown5 = (uint32_t) galloc(32),
		.unknown6 = (uint32_t) galloc(64),
		.nullForVertex = tiler ? (uint32_t) galloc(64) : 0,
		.null4 = 0,
		.fbd = fbd,
		.unknown7 = tiler ? 0 : ((uint32_t) galloc(64) | 1) /* TODO */
	};

	struct shader_meta *s = (struct shader_meta*) payload.shader;
	s->shader = import_shader(fd, shader, shader_size, tiler);

	if (!tiler) {
		uint32_t ni[] = {
			0x43200000, 0x42F00000, 0x3F000000, 0x00000000,
			0x43200000, 0x42F00000, 0x3F000000, 0x00000000
		};

		memcpy((void*) payload.unknown2, ni, sizeof(ni));
	}

	if (tiler) {
		/* Lose precision... on purpose? */
		payload.unknown7 = (uint32_t) s->shader;
	}

	payload.unknown7 = tiler ? 0xDEADBA00 : 0xDEADFA00;

	/* TODO: Decode me! */

	if (tiler) {
		s->unknown1 = 0x0007000000000000;
		s->unknown2 = 0x0000000000020602;
	} else {
		s->unknown1 = 0x0005000100000000;
		s->unknown2 = 0x0000000000420002;
	}

	/* TODO: Generate on the fly (see trace.c) */
	uint32_t *p = (uint32_t*) payload.attribute_meta;
	*p = 0x2DEA2200;

	/* I have *no* idea */

	uint64_t pi[] = {
		0x0000000017E49000, 0x0000000017E49000,
		0x0000000017E49000, 0x0000000017E49000,
		0x00000000179A2200, 0x0000000017E49000,
		0x0000000017E49000
	};

	memcpy((void*) payload.unknown6, pi, sizeof(pi));

	if (tiler) {
		uint32_t ni[] = {
			0xFF800000, 0xFF800000,
			0x7F800000, 0x7F800000,
			0x00000000, 0x3F800000,
			0x00000000, 0x00EF013F,
			0x00000000, 0x0000001F,
			0x02020000, 0x00000001
		};

		memcpy((void*) payload.nullForVertex, ni, sizeof(ni));
	}

	/* Use some magic numbers from the traces */
	uint64_t* unk1 = (uint64_t*) payload.unknown1;
	/* unk1[0] = 0x000000B296271001;
	unk1[1] = 0x000000B296273000; */

	unk1[0] = 0x5a5a5a5a5a5a1001;
	unk1[1] = 0x5a5a5a5a5a5a3000;

	uint32_t writeBuffer = (uint32_t) galloc(64);

	uint64_t* unk5 = (uint64_t*) payload.unknown5;
	unk5[0] = ((uint64_t) (tiler ? 0xDB : 0x7A) << 56) | writeBuffer | 1;
	unk5[1] = 0x0000004000000010;

	if (tiler) {
		uint32_t ni[] = {
			0x00000001, 0x00000000, 0x00070000, 0x00020602,
			0x00000000, 0x00000000, 0x00000000, 0x3712FFFF,
			0x44F0FFFF, 0x0007FF00, 0x0007FF00, 0x00000000,
			0x00000000, 0x00000000, 0x00000000, 0x00000200,
			0x00000000, 0xF0122122, 0x00000000, 0x00000000,
			0x00000000, 0xF0122122, 0x00000000, 0xFF800000,
			0xFF800000, 0x7F800000, 0x7F800000, 0x00000000,
			0x3F800000, 0x00000000, 0xEF013F00, 0x00000000,
			0x0000001F, 0x02020000, 0x00000001, 0x00000000
		};

		memcpy(payload.block2, ni, sizeof(ni));
	} else {
		uint32_t ni[] = {
			0x00000000, 0x0000000C, 0x00000030, 0x2DEA2200,
			0x00000000, 0x00000000, 0x00000000, /* Address to 1 */ 0xCAFEDA01,
			0x57000000, 0x00000010, 0x00000040, 0x17E49000,
			0x00000000, 0x17E49000, 0x00000000, 0x17E49000,
			0x00000000, 0x17E49000, 0x00000000, 0x179A2200,
			0x00000000, 0x17E49000, 0x00000000, 0x17E49000,
			0x00000000, 0x00000000, 0x00000000, 0x43200000,
			0x42F00000, 0x3F000000, 0x00000000, 0x43200000,
			0x42F00000, 0x3F000000, 0x00000000, 0x00000000
		};

		memcpy(payload.block2, ni, sizeof(ni));
	}

	/* Trap tiler job execution */

	if (tiler) {
		payload.shader = 0x5AB00A05;

		/* Hit second */
		//payload.zeroes = 0x5AB01A00;

		payload.unknown1 = 0x5AB02A00;
		payload.unknown2 = 0x5AB03A00;
		payload.attributes = 0x5AB04A00;
		payload.attribute_meta = 0x5AB05A00;
		payload.unknown5 = 0x5AB06A00;
		payload.unknown6 = 0x5AB07A00;
		payload.unknown7 = 0x5AB0DA00;

		/* Hit third */
		//payload.fbd	 = 0x5AB09A00;

		/* Hit first */
		// payload.nullForVertex = 0x5AB08A00;
	}

	memcpy(packet, &header, sizeof(header));
	memcpy(packet + sizeof(header), &payload, sizeof(payload));

	return packet;
}

uint32_t job_chain_vertex_tiler(int fd, float *vertices, size_t vertex_size,
				int mode, void* vertex_shader, size_t vs_sz,
				void *fragment_shader, size_t fs_sz,
				uint64_t heap_free_address, uint64_t scratchpad)
{
	uint32_t vertex_buffer = upload_vertices(vertices, vertex_size);
	uint32_t vertex_fbd = (uint32_t) make_mfbd(true, heap_free_address,
						   scratchpad);
	uint32_t zero_buffer = (uint32_t) alloc_gpu_pages(
	    fd, 0x20,
	    0x3800 | MALI_MEM_PROT_CPU_RD | MALI_MEM_PROT_CPU_WR |
	    MALI_MEM_PROT_GPU_RD);
	struct job_descriptor_header *set =
		set_value_helper(scratchpad + SV_OFFSET);
	struct job_descriptor_header *vertex =
		vertex_tiler_helper(fd, false,
				vertex_fbd, vertex_buffer,
				zero_buffer, mode,
				vertex_shader, vs_sz);
	struct job_descriptor_header *tiler =
		vertex_tiler_helper(fd, true,
				vertex_fbd, vertex_buffer,
				zero_buffer, mode,
				fragment_shader, fs_sz);

	set->next_job._32 = (uint32_t) vertex;
	vertex->next_job._32 = (uint32_t) tiler;

	/* TODO: Determine if these numbers are meaningful */
	set->job_index = 3;
	vertex->job_index = 1;
	tiler->job_index = 2;

	vertex->job_dependency_index_2 = set->job_index;
	tiler->job_dependency_index_1 = vertex->job_index;

	struct mali_jd_atom_v2 job = {
		.jc = (uint32_t) set,
		.ext_res_list = NULL,
		.nr_ext_res = 0,
		.core_req =
			MALI_JD_REQ_CS | MALI_JD_REQ_T | MALI_JD_REQ_CF |
			MALI_JD_REQ_COHERENT_GROUP,
		.atom_number = ++atom_count,
		.prio = MALI_JD_PRIO_MEDIUM,
		.device_nr = 0,
		.pre_dep = { no_dependency, no_dependency }
	};

	submit_job(fd, job);

	return (uint32_t) tiler;
}

void job_chain_replay(int fd, uint32_t tiler_jc, uint32_t fragment_jc,
		      uint64_t heap_free_address, uint64_t framebuffer)
{
	struct mali_jd_replay_payload *payload;
	struct mali_jd_dependency depFragment = {
		.atom_id = atom_count,
		.dependency_type = MALI_JD_DEP_TYPE_DATA
	};
	uint64_t *resource;

	payload = (struct mali_jd_replay_payload*) galloc(sizeof(*payload));

	payload->tiler_jc_list = tiler_jc;
	payload->fragment_jc = fragment_jc;
	payload->tiler_heap_free = heap_free_address;
	payload->fragment_hierarchy_mask = 0;
	payload->tiler_hierarchy_mask = 0;
	payload->hierarchy_default_weight = 0x10000;
	payload->tiler_core_req = MALI_JD_REQ_T | MALI_JD_REQ_COHERENT_GROUP;
	payload->fragment_core_req = MALI_JD_REQ_FS;

	resource = malloc(sizeof(u64) * 1);
	resource[0] = framebuffer | MALI_EXT_RES_ACCESS_EXCLUSIVE;

	struct mali_jd_atom_v2 job = {
		.jc = (uint32_t) payload,
		.ext_res_list = (struct mali_external_resource*)resource,
		.nr_ext_res = 1,
		.core_req =
			MALI_JD_REQ_EXTERNAL_RESOURCES | MALI_JD_REQ_SOFT_REPLAY,
		.atom_number = ++atom_count,
		.prio = MALI_JD_PRIO_LOW,
		.device_nr = 0,
		.pre_dep = { depFragment, no_dependency }
	};

	submit_job(fd, job);
}


code:
/*
 *
 * (C) COPYRIGHT 2010-2016 ARM Limited. All rights reserved.
 * Copyright (C) 2017 Cafe Beverage. All rights reserved.
 *
 * This program is free software and is provided to you under the terms of the
 * GNU General Public License version 2 as published by the Free Software
 * Foundation, and any use by you of this program is subject to the terms
 * of such GNU licence.
 *
 * A copy of the licence is included with the program, and can also be obtained
 * from Free Software Foundation, Inc., 51 Franklin Street, Fifth Floor,
 * Boston, MA  02110-1301, USA.
 *
 */

#ifndef __CHAI_NOTES_H
#define __CHAI_NOTES_H

#include <stddef.h>

#include "jobs.h"

/* See corresponding structures in chai/jobs.h */

struct payload_set_value {
	uint64_t out;
	uint64_t unknown;
};

/* From mali_kbase_replay.c */
/* fbd_address & mask */
#define FBD_POINTER_MASK (~0x3f)

/* set for MFBD; otherwise SFBD */
#define FBD_TYPE 1

#define SFBD (0)
#define MFBD (1)

/* appears to be set for fragment; clear for vertex/tiler */

#define FBD_TILER (0)
#define FBD_FRAGMENT (2)

/* From mali_kbase_10969_workaround.c */
#define X_COORDINATE_MASK 0x00000FFF
#define Y_COORDINATE_MASK 0x0FFF0000

/* The above defines as useful macros */
#define TILE_COORD_X(coord) ((coord) & X_COORDINATE_MASK)
#define TILE_COORD_Y(coord) (((coord) & Y_COORDINATE_MASK) >> 16)

#define MAKE_TILE_COORD(X, Y, flag) ((X) | ((Y) << 16) | (flag))

struct payload_fragment {
	uint32_t min_tile_coord;
	uint32_t max_tile_coord;
	uint64_t fragment_fbd;
};

/* Parts from mali_kbase_replay.c. Mostly from RE. */

struct payload_vertex_tiler {
	uint32_t block1[10];

	uintptr_t null0;
	uintptr_t zeroes;
	uintptr_t unknown1; /* pointer */
	uintptr_t null1;
	uintptr_t null2;
	uintptr_t unknown2; /* pointer */
	uintptr_t shader;   /* struct shader_meta */
	uintptr_t attributes; /* struct attribute_buffer[] */
	uintptr_t attribute_meta; /* attribute_meta[] */
	uintptr_t unknown5; /* pointer */
	uintptr_t unknown6; /* pointer */
	uintptr_t nullForVertex;
	uintptr_t null4;
	uintptr_t fbd;	   /* struct tentative_fbd */
	uintptr_t unknown7; /* pointer */

	uint32_t block2[36];
};

/* FIXME: This is very clearly wrong */

#define SHADER (1 << 0)
#define SHADER_VERTEX (1 << 2)
#define SHADER_FRAGMENT (1 << 3)

struct shader_meta {
	uint64_t shader;
	uint64_t unknown1;
	uint64_t unknown2;
};

struct attribute_buffer {
	uint64_t elements;
	size_t element_size;
	size_t total_size;
};

typedef uint64_t attribute_meta_t;

#define HAS_ATTRIBUTE(meta) (!!meta)
#define ATTRIBUTE_NO(meta) (meta & 0xFF)
#define ATTRIBUTE_FLAGS(meta) (meta & 0xFFFFFFFFFFFFFF00)

/* Synthesis of fragment_job from mali_kbase_reply.c with #defines from
 * errata file */

#define JOB_NOT_STARTED	(0)
#define JOB_TYPE_NULL		(1)
#define JOB_TYPE_SET_VALUE	(2)
#define JOB_TYPE_CACHE_FLUSH	(3)
#define JOB_TYPE_COMPUTE	(4)
#define JOB_TYPE_VERTEX		(5)
#define JOB_TYPE_TILER		(7)
#define JOB_TYPE_FUSED		(8)
#define JOB_TYPE_FRAGMENT	(9)

#define CHAI_POINTS		0x01
#define CHAI_LINES		0x02
#define CHAI_TRIANGLES		0x08
#define CHAI_TRIANGLE_STRIP	0x0A
#define CHAI_TRIANGLE_FAN	0x0C

/* See mali_kbase_replay.c */

#define FBD_HIERARCHY_WEIGHTS 8
#define FBD_HIERARCHY_MASK_MASK 0x1fff

struct tentative_mfbd {
	uint64_t blah;
	uint64_t ugaT;
	uint32_t block1[10];
	uint32_t unknown1;
	uint32_t flags;
	uint64_t block2[2];
	uint64_t heap_free_address;
	uint64_t unknown2;
	uint32_t weights[FBD_HIERARCHY_WEIGHTS];
	uint64_t unknown_gpu_addressN;
	uint32_t block3[22];
	uint64_t unknown_gpu_address;
	uint64_t unknown3;
	uint32_t block4[10];
};

#define JOB_32_BIT 0
#define JOB_64_BIT 1

#endif


code:
/*
 *
 * (C) COPYRIGHT 2014-2016 ARM Limited. All rights reserved.
 *
 * This program is free software and is provided to you under the terms of the
 * GNU General Public License version 2 as published by the Free Software
 * Foundation, and any use by you of this program is subject to the terms
 * of such GNU licence.
 *
 * A copy of the licence is included with the program, and can also be obtained
 * from Free Software Foundation, Inc., 51 Franklin Street, Fifth Floor,
 * Boston, MA  02110-1301, USA.
 *
 */

#ifndef __JOBS_H
#define __JOBS_H

struct job_descriptor_header {
	u32 exception_status;
	u32 first_incomplete_task;
	u64 fault_pointer;
	u8 job_descriptor_size : 1;
	u8 job_type : 7;
	u8 job_barrier : 1;
	u8 _reserved_01 : 1;
	u8 _reserved_1 : 1;
	u8 _reserved_02 : 1;
	u8 _reserved_03 : 1;
	u8 _reserved_2 : 1;
	u8 _reserved_04 : 1;
	u8 _reserved_05 : 1;
	u16 job_index;
	u16 job_dependency_index_1;
	u16 job_dependency_index_2;
	union {
		u64 _64;
		u64 _32;
	} next_job;

} __attribute__((packed));

typedef struct mali_jd_replay_payload {
	/**
	 * Pointer to the first entry in the mali_jd_replay_jc list.  These
	 * will be replayed in @b reverse order (so that extra ones can be added
	 * to the head in future soft jobs without affecting this soft job)
	 */
	u64 tiler_jc_list;

	/**
	 * Pointer to the fragment job chain.
	 */
	u64 fragment_jc;

	/**
	 * Pointer to the tiler heap free FBD field to be modified.
	 */
	u64 tiler_heap_free;

	/**
	 * Hierarchy mask for the replayed fragment jobs. May be zero.
	 */
	u16 fragment_hierarchy_mask;

	/**
	 * Hierarchy mask for the replayed tiler jobs. May be zero.
	 */
	u16 tiler_hierarchy_mask;

	/**
	 * Default weight to be used for hierarchy levels not in the original
	 * mask.
	 */
	u32 hierarchy_default_weight;

	/**
	 * Core requirements for the tiler job chain
	 */
	mali_jd_core_req tiler_core_req;

	/**
	 * Core requirements for the fragment job chain
	 */
	mali_jd_core_req fragment_core_req;

	u8 padding[4];
} mali_jd_replay_payload;

#endif


code:
/*
 * Copyright © 2010 Intel Corporation
 * Copyright © 2010 Francisco Jerez <currojerez@riseup.net>
 *
 * Permission is hereby granted, free of charge, to any person obtaining a
 * copy of this software and associated documentation files (the "Software"),
 * to deal in the Software without restriction, including without limitation
 * the rights to use, copy, modify, merge, publish, distribute, sublicense,
 * and/or sell copies of the Software, and to permit persons to whom the
 * Software is furnished to do so, subject to the following conditions:
 *
 * The above copyright notice and this permission notice (including the next
 * paragraph) shall be included in all copies or substantial portions of the
 * Software.
 *
 * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
 * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
 * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL
 * THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
 * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING
 * FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS
 * IN THE SOFTWARE.
 *
 */

#ifndef _LIST_H_
#define _LIST_H_

/* classic doubly-link circular list */
struct list {
    struct list *next, *prev;
};

#define LIST_HEAD_INIT(name) { &(name), &(name) }

#define LIST_HEAD(name) \
	struct list name = LIST_HEAD_INIT(name)

static void
list_init(struct list *list)
{
    list->next = list->prev = list;
}

static inline void
__list_add(struct list *entry,
	    struct list *prev,
	    struct list *next)
{
    next->prev = entry;
    entry->next = next;
    entry->prev = prev;
    prev->next = entry;
}

static inline void
list_add(struct list *entry, struct list *head)
{
    __list_add(entry, head, head->next);
}

static inline void
__list_del(struct list *prev, struct list *next)
{
    next->prev = prev;
    prev->next = next;
}

static inline void
list_del(struct list *entry)
{
    __list_del(entry->prev, entry->next);
    list_init(entry);
}

static inline int
list_is_empty(struct list *head)
{
    return head->next == head;
}

#ifndef container_of
#define container_of(ptr, type, member) \
    (type *)((char *)(ptr) - (char *) &((type *)0)->member)
#endif

#define list_entry(ptr, type, member) \
    container_of(ptr, type, member)

#define list_first_entry(ptr, type, member) \
    list_entry((ptr)->next, type, member)

#define __container_of(ptr, sample, member)				\
    (void *)((char *)(ptr)						\
	     - ((char *)&(sample)->member - (char *)(sample)))

#define list_for_each_entry(pos, head, member)				\
    for (pos = __container_of((head)->next, pos, member);		\
	 &pos->member != (head);					\
	 pos = __container_of(pos->member.next, pos, member))

#define list_for_each_entry_safe(pos, tmp, head, member)		\
    for (pos = __container_of((head)->next, pos, member),		\
	 tmp = __container_of(pos->member.next, pos, member);		\
	 &pos->member != (head);					\
	 pos = tmp, tmp = __container_of(pos->member.next, tmp, member))

#endif


code:
/*
 * © Copyright 2017 The BiOpenly Community
 *
 * This program is free software and is provided to you under the terms of the
 * GNU General Public License version 2 as published by the Free Software
 * Foundation, and any use by you of this program is subject to the terms
 * of such GNU license.
 *
 * A copy of the licence is included with the program, and can also be obtained
 * from Free Software Foundation, Inc., 51 Franklin Street, Fifth Floor,
 * Boston, MA  02110-1301, USA.
 *
 */

/**
 * Definitions for all of the ioctls for the original open source bifrost GPU
 * kernel driver, written by ARM.
 */

#ifndef __MALI_IOCTL_H__
#define __MALI_IOCTL_H__

#include <panloader-util.h>

#define MALI_GPU_NUM_TEXTURE_FEATURES_REGISTERS 3
#define MALI_GPU_MAX_JOB_SLOTS 16
#define MALI_MAX_COHERENT_GROUPS 16

typedef u8 mali_atom_id;

/**
 * Since these structs are passed to and from the kernel we need to make sure
 * that we get the size of each struct to match exactly what the kernel is
 * expecting. So, when editing this file make sure to add static asserts that
 * check each struct's size against the arg length you see in strace.
 */

enum mali_ioctl_mem_flags {
	/* IN */
	MALI_MEM_PROT_CPU_RD = (1U << 0),      /**< Read access CPU side */
	MALI_MEM_PROT_CPU_WR = (1U << 1),      /**< Write access CPU side */
	MALI_MEM_PROT_GPU_RD = (1U << 2),      /**< Read access GPU side */
	MALI_MEM_PROT_GPU_WR = (1U << 3),      /**< Write access GPU side */
	MALI_MEM_PROT_GPU_EX = (1U << 4),      /**< Execute allowed on the GPU
						    side */

	MALI_MEM_GROW_ON_GPF = (1U << 9),      /**< Grow backing store on GPU
						    Page Fault */

	MALI_MEM_COHERENT_SYSTEM = (1U << 10), /**< Page coherence Outer
						    shareable, if available */
	MALI_MEM_COHERENT_LOCAL = (1U << 11),  /**< Page coherence Inner
						    shareable */
	MALI_MEM_CACHED_CPU = (1U << 12),      /**< Should be cached on the
						    CPU */

	/* IN/OUT */
	MALI_MEM_SAME_VA = (1U << 13), /**< Must have same VA on both the GPU
					    and the CPU */
	/* OUT */
	MALI_MEM_NEED_MMAP = (1U << 14), /**< Must call mmap to acquire a GPU
					     address for the alloc */
	/* IN */
	MALI_MEM_COHERENT_SYSTEM_REQUIRED = (1U << 15), /**< Page coherence
					     Outer shareable, required. */
	MALI_MEM_SECURE = (1U << 16),          /**< Secure memory */
	MALI_MEM_DONT_NEED = (1U << 17),       /**< Not needed physical
						    memory */
	MALI_MEM_IMPORT_SHARED = (1U << 18),   /**< Must use shared CPU/GPU zone
						    (SAME_VA zone) but doesn't
						    require the addresses to
						    be the same */
};

#define MALI_MEM_MAP_TRACKING_HANDLE (0x3000)

enum mali_ioctl_coherency_mode {
	COHERENCY_ACE_LITE = 0,
	COHERENCY_ACE      = 1,
	COHERENCY_NONE     = 31
};

/*
 * Mali Atom priority
 *
 * Only certain priority levels are actually implemented, as specified by the
 * MALI_JD_PRIO_<...> definitions below. It is undefined to use a priority
 * level that is not one of those defined below.
 *
 * Priority levels only affect scheduling between atoms of the same type within
 * a mali context, and only after the atoms have had dependencies resolved.
 * Fragment atoms does not affect non-frament atoms with lower priorities, and
 * the other way around. For example, a low priority atom that has had its
 * dependencies resolved might run before a higher priority atom that has not
 * had its dependencies resolved.
 *
 * The scheduling between mali contexts/processes and between atoms from
 * different mali contexts/processes is unaffected by atom priority.
 *
 * The atoms are scheduled as follows with respect to their priorities:
 * - Let atoms 'X' and 'Y' be for the same job slot who have dependencies
 *   resolved, and atom 'X' has a higher priority than atom 'Y'
 * - If atom 'Y' is currently running on the HW, then it is interrupted to
 *   allow atom 'X' to run soon after
 * - If instead neither atom 'Y' nor atom 'X' are running, then when choosing
 *   the next atom to run, atom 'X' will always be chosen instead of atom 'Y'
 * - Any two atoms that have the same priority could run in any order with
 *   respect to each other. That is, there is no ordering constraint between
 *   atoms of the same priority.
 */
typedef u8 mali_jd_prio;
#define MALI_JD_PRIO_MEDIUM  ((mali_jd_prio)0)
#define MALI_JD_PRIO_HIGH    ((mali_jd_prio)1)
#define MALI_JD_PRIO_LOW     ((mali_jd_prio)2)

/**
 * @brief Job dependency type.
 *
 * A flags field will be inserted into the atom structure to specify whether a
 * dependency is a data or ordering dependency (by putting it before/after
 * 'core_req' in the structure it should be possible to add without changing
 * the structure size).  When the flag is set for a particular dependency to
 * signal that it is an ordering only dependency then errors will not be
 * propagated.
 */
typedef u8 mali_jd_dep_type;
#define MALI_JD_DEP_TYPE_INVALID  (0)       /**< Invalid dependency */
#define MALI_JD_DEP_TYPE_DATA     (1U << 0) /**< Data dependency */
#define MALI_JD_DEP_TYPE_ORDER    (1U << 1) /**< Order dependency */

/**
 * @brief Job chain hardware requirements.
 *
 * A job chain must specify what GPU features it needs to allow the
 * driver to schedule the job correctly.  By not specifying the
 * correct settings can/will cause an early job termination.  Multiple
 * values can be ORed together to specify multiple requirements.
 * Special case is ::MALI_JD_REQ_DEP, which is used to express complex
 * dependencies, and that doesn't execute anything on the hardware.
 */
typedef u32 mali_jd_core_req;

/* Requirements that come from the HW */

/**
 * No requirement, dependency only
 */
#define MALI_JD_REQ_DEP ((mali_jd_core_req)0)

/**
 * Requires fragment shaders
 */
#define MALI_JD_REQ_FS  ((mali_jd_core_req)1 << 0)

/**
 * Requires compute shaders
 * This covers any of the following Midgard Job types:
 * - Vertex Shader Job
 * - Geometry Shader Job
 * - An actual Compute Shader Job
 *
 * Compare this with @ref MALI_JD_REQ_ONLY_COMPUTE, which specifies that the
 * job is specifically just the "Compute Shader" job type, and not the "Vertex
 * Shader" nor the "Geometry Shader" job type.
 */
#define MALI_JD_REQ_CS  ((mali_jd_core_req)1 << 1)
#define MALI_JD_REQ_T   ((mali_jd_core_req)1 << 2)   /**< Requires tiling */
#define MALI_JD_REQ_CF  ((mali_jd_core_req)1 << 3)   /**< Requires cache flushes */
#define MALI_JD_REQ_V   ((mali_jd_core_req)1 << 4)   /**< Requires value writeback */

/* SW-only requirements - the HW does not expose these as part of the job slot
 * capabilities */

/* Requires fragment job with AFBC encoding */
#define MALI_JD_REQ_FS_AFBC  ((mali_jd_core_req)1 << 13)

/**
 * SW-only requirement: coalesce completion events.
 * If this bit is set then completion of this atom will not cause an event to
 * be sent to userspace, whether successful or not; completion events will be
 * deferred until an atom completes which does not have this bit set.
 *
 * This bit may not be used in combination with MALI_JD_REQ_EXTERNAL_RESOURCES.
 */
#define MALI_JD_REQ_EVENT_COALESCE ((mali_jd_core_req)1 << 5)

/**
 * SW Only requirement: the job chain requires a coherent core group. We don't
 * mind which coherent core group is used.
 */
#define MALI_JD_REQ_COHERENT_GROUP  ((mali_jd_core_req)1 << 6)

/**
 * SW Only requirement: The performance counters should be enabled only when
 * they are needed, to reduce power consumption.
 */

#define MALI_JD_REQ_PERMON               ((mali_jd_core_req)1 << 7)

/**
 * SW Only requirement: External resources are referenced by this atom.  When
 * external resources are referenced no syncsets can be bundled with the atom
 * but should instead be part of a NULL jobs inserted into the dependency
 * tree.  The first pre_dep object must be configured for the external
 * resouces to use, the second pre_dep object can be used to create other
 * dependencies.
 *
 * This bit may not be used in combination with MALI_JD_REQ_EVENT_COALESCE.
 */
#define MALI_JD_REQ_EXTERNAL_RESOURCES   ((mali_jd_core_req)1 << 8)

/**
 * SW Only requirement: Software defined job. Jobs with this bit set will not
 * be submitted to the hardware but will cause some action to happen within
 * the driver
 */
#define MALI_JD_REQ_SOFT_JOB        ((mali_jd_core_req)1 << 9)

#define MALI_JD_REQ_SOFT_DUMP_CPU_GPU_TIME      (MALI_JD_REQ_SOFT_JOB | 0x1)
#define MALI_JD_REQ_SOFT_FENCE_TRIGGER          (MALI_JD_REQ_SOFT_JOB | 0x2)
#define MALI_JD_REQ_SOFT_FENCE_WAIT             (MALI_JD_REQ_SOFT_JOB | 0x3)

/**
 * SW Only requirement : Replay job.
 *
 * If the preceding job fails, the replay job will cause the jobs specified in
 * the list of mali_jd_replay_payload pointed to by the jc pointer to be
 * replayed.
 *
 * A replay job will only cause jobs to be replayed up to MALIP_JD_REPLAY_LIMIT
 * times. If a job fails more than MALIP_JD_REPLAY_LIMIT times then the replay
 * job is failed, as well as any following dependencies.
 *
 * The replayed jobs will require a number of atom IDs. If there are not enough
 * free atom IDs then the replay job will fail.
 *
 * If the preceding job does not fail, then the replay job is returned as
 * completed.
 *
 * The replayed jobs will never be returned to userspace. The preceding failed
 * job will be returned to userspace as failed; the status of this job should
 * be ignored. Completion should be determined by the status of the replay soft
 * job.
 *
 * In order for the jobs to be replayed, the job headers will have to be
 * modified. The Status field will be reset to NOT_STARTED. If the Job Type
 * field indicates a Vertex Shader Job then it will be changed to Null Job.
 *
 * The replayed jobs have the following assumptions :
 *
 * - No external resources. Any required external resources will be held by the
 *   replay atom.
 * - Pre-dependencies are created based on job order.
 * - Atom numbers are automatically assigned.
 * - device_nr is set to 0. This is not relevant as
 *   MALI_JD_REQ_SPECIFIC_COHERENT_GROUP should not be set.
 * - Priority is inherited from the replay job.
 */
#define MALI_JD_REQ_SOFT_REPLAY                 (MALI_JD_REQ_SOFT_JOB | 0x4)
/**
 * SW only requirement: event wait/trigger job.
 *
 * - MALI_JD_REQ_SOFT_EVENT_WAIT: this job will block until the event is set.
 * - MALI_JD_REQ_SOFT_EVENT_SET: this job sets the event, thus unblocks the
 *   other waiting jobs. It completes immediately.
 * - MALI_JD_REQ_SOFT_EVENT_RESET: this job resets the event, making it
 *   possible for other jobs to wait upon. It completes immediately.
 */
#define MALI_JD_REQ_SOFT_EVENT_WAIT             (MALI_JD_REQ_SOFT_JOB | 0x5)
#define MALI_JD_REQ_SOFT_EVENT_SET              (MALI_JD_REQ_SOFT_JOB | 0x6)
#define MALI_JD_REQ_SOFT_EVENT_RESET            (MALI_JD_REQ_SOFT_JOB | 0x7)

#define MALI_JD_REQ_SOFT_DEBUG_COPY             (MALI_JD_REQ_SOFT_JOB | 0x8)

/**
 * SW only requirement: Just In Time allocation
 *
 * This job requests a JIT allocation based on the request in the
 * @base_jit_alloc_info structure which is passed via the jc element of
 * the atom.
 *
 * It should be noted that the id entry in @base_jit_alloc_info must not
 * be reused until it has been released via @MALI_JD_REQ_SOFT_JIT_FREE.
 *
 * Should this soft job fail it is expected that a @MALI_JD_REQ_SOFT_JIT_FREE
 * soft job to free the JIT allocation is still made.
 *
 * The job will complete immediately.
 */
#define MALI_JD_REQ_SOFT_JIT_ALLOC              (MALI_JD_REQ_SOFT_JOB | 0x9)
/**
 * SW only requirement: Just In Time free
 *
 * This job requests a JIT allocation created by @MALI_JD_REQ_SOFT_JIT_ALLOC
 * to be freed. The ID of the JIT allocation is passed via the jc element of
 * the atom.
 *
 * The job will complete immediately.
 */
#define MALI_JD_REQ_SOFT_JIT_FREE               (MALI_JD_REQ_SOFT_JOB | 0xa)

/**
 * SW only requirement: Map external resource
 *
 * This job requests external resource(s) are mapped once the dependencies
 * of the job have been satisfied. The list of external resources are
 * passed via the jc element of the atom which is a pointer to a
 * @base_external_resource_list.
 */
#define MALI_JD_REQ_SOFT_EXT_RES_MAP            (MALI_JD_REQ_SOFT_JOB | 0xb)
/**
 * SW only requirement: Unmap external resource
 *
 * This job requests external resource(s) are unmapped once the dependencies
 * of the job has been satisfied. The list of external resources are
 * passed via the jc element of the atom which is a pointer to a
 * @base_external_resource_list.
 */
#define MALI_JD_REQ_SOFT_EXT_RES_UNMAP          (MALI_JD_REQ_SOFT_JOB | 0xc)

/**
 * HW Requirement: Requires Compute shaders (but not Vertex or Geometry Shaders)
 *
 * This indicates that the Job Chain contains Midgard Jobs of the 'Compute
 * Shaders' type.
 *
 * In contrast to @ref MALI_JD_REQ_CS, this does \b not indicate that the Job
 * Chain contains 'Geometry Shader' or 'Vertex Shader' jobs.
 */
#define MALI_JD_REQ_ONLY_COMPUTE    ((mali_jd_core_req)1 << 10)

/**
 * HW Requirement: Use the mali_jd_atom::device_nr field to specify a
 * particular core group
 *
 * If both @ref MALI_JD_REQ_COHERENT_GROUP and this flag are set, this flag
 * takes priority
 *
 * This is only guaranteed to work for @ref MALI_JD_REQ_ONLY_COMPUTE atoms.
 *
 * If the core availability policy is keeping the required core group turned
 * off, then the job will fail with a @ref MALI_JD_EVENT_PM_EVENT error code.
 */
#define MALI_JD_REQ_SPECIFIC_COHERENT_GROUP ((mali_jd_core_req)1 << 11)

/**
 * SW Flag: If this bit is set then the successful completion of this atom
 * will not cause an event to be sent to userspace
 */
#define MALI_JD_REQ_EVENT_ONLY_ON_FAILURE   ((mali_jd_core_req)1 << 12)

/**
 * SW Flag: If this bit is set then completion of this atom will not cause an
 * event to be sent to userspace, whether successful or not.
 */
#define MALI_JD_REQ_EVENT_NEVER ((mali_jd_core_req)1 << 14)

/**
 * SW Flag: Skip GPU cache clean and invalidation before starting a GPU job.
 *
 * If this bit is set then the GPU's cache will not be cleaned and invalidated
 * until a GPU job starts which does not have this bit set or a job completes
 * which does not have the @ref MALI_JD_REQ_SKIP_CACHE_END bit set. Do not use if
 * the CPU may have written to memory addressed by the job since the last job
 * without this bit set was submitted.
 */
#define MALI_JD_REQ_SKIP_CACHE_START ((mali_jd_core_req)1 << 15)

/**
 * SW Flag: Skip GPU cache clean and invalidation after a GPU job completes.
 *
 * If this bit is set then the GPU's cache will not be cleaned and invalidated
 * until a GPU job completes which does not have this bit set or a job starts
 * which does not have the @ref MALI_JD_REQ_SKIP_CACHE_START bti set. Do not
 * use if the CPU may read from or partially overwrite memory addressed by the
 * job before the next job without this bit set completes.
 */
#define MALI_JD_REQ_SKIP_CACHE_END ((mali_jd_core_req)1 << 16)

/**
 * These requirement bits are currently unused in mali_jd_core_req
 */
#define MALIP_JD_REQ_RESERVED \
	(~(MALI_JD_REQ_ATOM_TYPE | MALI_JD_REQ_EXTERNAL_RESOURCES | \
	MALI_JD_REQ_EVENT_ONLY_ON_FAILURE | MALIP_JD_REQ_EVENT_NEVER | \
	MALI_JD_REQ_EVENT_COALESCE | \
	MALI_JD_REQ_COHERENT_GROUP | MALI_JD_REQ_SPECIFIC_COHERENT_GROUP | \
	MALI_JD_REQ_FS_AFBC | MALI_JD_REQ_PERMON | \
	MALI_JD_REQ_SKIP_CACHE_START | MALI_JD_REQ_SKIP_CACHE_END))

/**
 * Mask of all bits in mali_jd_core_req that control the type of the atom.
 *
 * This allows dependency only atoms to have flags set
 */
#define MALI_JD_REQ_ATOM_TYPE \
	(MALI_JD_REQ_FS | MALI_JD_REQ_CS | MALI_JD_REQ_T | MALI_JD_REQ_CF | \
	MALI_JD_REQ_V | MALI_JD_REQ_SOFT_JOB | MALI_JD_REQ_ONLY_COMPUTE)

/**
 * Mask of all bits in mali_jd_core_req that control the type of a soft job.
 */
#define MALI_JD_REQ_SOFT_JOB_TYPE (MALI_JD_REQ_SOFT_JOB | 0x1f)

/*
 * Returns non-zero value if core requirements passed define a soft job or
 * a dependency only job.
 */
#define MALI_JD_REQ_SOFT_JOB_OR_DEP(core_req) \
	((core_req & MALI_JD_REQ_SOFT_JOB) || \
	(core_req & MALI_JD_REQ_ATOM_TYPE) == MALI_JD_REQ_DEP)

struct mali_gpu_core_props {
	/**
	 * Product specific value.
	 */
	u32 product_id;

	/**
	 * Status of the GPU release.
	 * No defined values, but starts at 0 and increases by one for each
	 * release status (alpha, beta, EAC, etc.).
	 * 4 bit values (0-15).
	 */
	u16 version_status;

	/**
	 * Minor release number of the GPU. "P" part of an "RnPn" release
	 * number.
	 * 8 bit values (0-255).
	 */
	u16 minor_revision;

	/**
	 * Major release number of the GPU. "R" part of an "RnPn" release
	 * number.
	 * 4 bit values (0-15).
	 */
	u16 major_revision;

	u16 :16;

	/**
	 * @usecase GPU clock speed is not specified in the Midgard
	 * Architecture, but is <b>necessary for OpenCL's clGetDeviceInfo()
	 * function</b>.
	 */
	u32 gpu_speed_mhz;

	/**
	 * @usecase GPU clock max/min speed is required for computing
	 * best/worst case in tasks as job scheduling ant irq_throttling. (It
	 * is not specified in the Midgard Architecture).
	 */
	u32 gpu_freq_khz_max;
	u32 gpu_freq_khz_min;

	/**
	 * Size of the shader program counter, in bits.
	 */
	u32 log2_program_counter_size;

	/**
	 * TEXTURE_FEATURES_x registers, as exposed by the GPU. This is a
	 * bitpattern where a set bit indicates that the format is supported.
	 *
	 * Before using a texture format, it is recommended that the
	 * corresponding bit be checked.
	 */
	u32 texture_features[MALI_GPU_NUM_TEXTURE_FEATURES_REGISTERS];

	/**
	 * Theoretical maximum memory available to the GPU. It is unlikely
	 * that a client will be able to allocate all of this memory for their
	 * own purposes, but this at least provides an upper bound on the
	 * memory available to the GPU.
	 *
	 * This is required for OpenCL's clGetDeviceInfo() call when
	 * CL_DEVICE_GLOBAL_MEM_SIZE is requested, for OpenCL GPU devices. The
	 * client will not be expecting to allocate anywhere near this value.
	 */
	u64 gpu_available_memory_size;
};

struct mali_gpu_l2_cache_props {
	u8 log2_line_size;
	u8 log2_cache_size;
	u8 num_l2_slices; /* Number of L2C slices. 1 or higher */
	u64 :40;
};

struct mali_gpu_tiler_props {
	u32 bin_size_bytes;	/* Max is 4*2^15 */
	u32 max_active_levels;	/* Max is 2^15 */
};

struct mali_gpu_thread_props {
	u32 max_threads;            /* Max. number of threads per core */
	u32 max_workgroup_size;     /* Max. number of threads per workgroup */
	u32 max_barrier_size;       /* Max. number of threads that can
				       synchronize on a simple barrier */
	u16 max_registers;          /* Total size [1..65535] of the register
				       file available per core. */
	u8  max_task_queue;         /* Max. tasks [1..255] which may be sent
				       to a core before it becomes blocked. */
	u8  max_thread_group_split; /* Max. allowed value [1..15] of the
				       Thread Group Split field. */
	enum {
		MALI_GPU_IMPLEMENTATION_UNKNOWN = 0,
		MALI_GPU_IMPLEMENTATION_SILICON = 1,
		MALI_GPU_IMPLEMENTATION_FPGA    = 2,
		MALI_GPU_IMPLEMENTATION_SW      = 3,
	} impl_tech :8;
	u64 :56;
};

/**
 * @brief descriptor for a coherent group
 *
 * \c core_mask exposes all cores in that coherent group, and \c num_cores
 * provides a cached population-count for that mask.
 *
 * @note Whilst all cores are exposed in the mask, not all may be available to
 * the application, depending on the Kernel Power policy.
 *
 * @note if u64s must be 8-byte aligned, then this structure has 32-bits of
 * wastage.
 */
struct mali_ioctl_gpu_coherent_group {
	u64 core_mask;	       /**< Core restriction mask required for the
				 group */
	u16 num_cores;	       /**< Number of cores in the group */
	u64 :48;
};

/**
 * @brief Coherency group information
 *
 * Note that the sizes of the members could be reduced. However, the \c group
 * member might be 8-byte aligned to ensure the u64 core_mask is 8-byte
 * aligned, thus leading to wastage if the other members sizes were reduced.
 *
 * The groups are sorted by core mask. The core masks are non-repeating and do
 * not intersect.
 */
struct mali_gpu_coherent_group_info {
	u32 num_groups;

	/**
	 * Number of core groups (coherent or not) in the GPU. Equivalent to
	 * the number of L2 Caches.
	 *
	 * The GPU Counter dumping writes 2048 bytes per core group,
	 * regardless of whether the core groups are coherent or not. Hence
	 * this member is needed to calculate how much memory is required for
	 * dumping.
	 *
	 * @note Do not use it to work out how many valid elements are in the
	 * group[] member. Use num_groups instead.
	 */
	u32 num_core_groups;

	/**
	 * Coherency features of the memory, accessed by @ref gpu_mem_features
	 * methods
	 */
	u32 coherency;

	u32 :32;

	/**
	 * Descriptors of coherent groups
	 */
	struct mali_ioctl_gpu_coherent_group group[MALI_MAX_COHERENT_GROUPS];
};

/**
 * A complete description of the GPU's Hardware Configuration Discovery
 * registers.
 *
 * The information is presented inefficiently for access. For frequent access,
 * the values should be better expressed in an unpacked form in the
 * base_gpu_props structure.
 *
 * @usecase The raw properties in @ref gpu_raw_gpu_props are necessary to
 * allow a user of the Mali Tools (e.g. PAT) to determine "Why is this device
 * behaving differently?". In this case, all information about the
 * configuration is potentially useful, but it <b>does not need to be processed
 * by the driver</b>. Instead, the raw registers can be processed by the Mali
 * Tools software on the host PC.
 *
 */
struct mali_gpu_raw_props {
	u64 shader_present;
	u64 tiler_present;
	u64 l2_present;
	u64 stack_present;

	u32 l2_features;
	u32 suspend_size; /* API 8.2+ */
	u32 mem_features;
	u32 mmu_features;

	u32 as_present;

	u32 js_present;
	u32 js_features[MALI_GPU_MAX_JOB_SLOTS];
	u32 tiler_features;
	u32 texture_features[3];

	u32 gpu_id;

	u32 thread_max_threads;
	u32 thread_max_workgroup_size;
	u32 thread_max_barrier_size;
	u32 thread_features;

	/*
	 * Note: This is the _selected_ coherency mode rather than the
	 * available modes as exposed in the coherency_features register.
	 */
	u32 coherency_mode;
};

/*
 * The original mali driver from ARM has 64 bit memory pointers in most of the
 * ioctls argument structures, regardless of whether or not the system is 32
 * bit or 64 bit. For 32 bit systems, the upper 32 bits are ignored
 *
 * FIXME: confirm this actually works
 */
#ifdef  __LP64__
#define PAD_PTR(p) p
#else
#define PAD_PTR(p) p; u32 :32
#endif

/* FIXME: Again, they don't specify any of these as packed structs. However,
 * looking at these structs I'm worried that there is already spots where the
 * compiler is potentially sticking in padding...
 * Going to try something a little crazy, and just hope that our compiler
 * happens to add the same kind of offsets since we can't really compare sizes
 */

/*
 * Blob provided by the driver to store callback driver, not actually modified
 * by the driver itself
 */
struct mali_jd_udata {
	u64 blob[2];
};

struct mali_jd_dependency {
	mali_atom_id  atom_id;               /**< An atom number */
	mali_jd_dep_type dependency_type;    /**< Dependency type */
};

#define MALI_EXT_RES_MAX 10

/* The original header never explicitly defines any values for these. In C,
 * this -should- expand to SHARED == 0 and EXCLUSIVE == 1, so the only flag we
 * actually need to decode here is EXCLUSIVE
 */
enum mali_external_resource_access {
	MALI_EXT_RES_ACCESS_SHARED,
	MALI_EXT_RES_ACCESS_EXCLUSIVE,
};

struct mali_external_resource {
	/*
	 * Lyude:
	 * At first I thought this might be the count for what could be in
	 * ext_resource, but it looks like there's not really a reason for
	 * ext_resource being an array with a length of one?
	 * Maybe this is some weird compatibility crap...
	 *
	 * (source for this from original kdriver: mali_kbase_jd.c)
	 */
	u64 count;
	u64 ext_resource[1];
};

struct mali_jd_atom_v2 {
	u64 jc;                             /**< job-chain GPU address */
	struct mali_jd_udata udata;	    /**< user data */
	PAD_PTR(struct mali_external_resource *ext_res_list); /**< list of external resources */
	u16 nr_ext_res;			    /**< nr of external resources */
	u16 compat_core_req;	            /**< core requirements which
					      correspond to the legacy support
					      for UK 10.2 */
	struct mali_jd_dependency pre_dep[2];  /**< pre-dependencies, one need to
					      use SETTER function to assign
					      this field, this is done in
					      order to reduce possibility of
					      improper assigment of a
					      dependency field */
	mali_atom_id atom_number;	    /**< unique number to identify the
					      atom */
	mali_jd_prio prio;                  /**< Atom priority. Refer to @ref
					      mali_jd_prio for more details */
	u8 device_nr;			    /**< coregroup when
					      BASE_JD_REQ_SPECIFIC_COHERENT_GROUP
					      specified */
	u8 :8;
	mali_jd_core_req core_req;          /**< core requirements */
} __attribute__((packed));
ASSERT_SIZEOF_TYPE(struct mali_jd_atom_v2, 48);

/**
 * Header used by all ioctls
 */
union mali_ioctl_header {
	/* [in] The ID of the UK function being called */
	u32 id :32;
	/* [out] The return value of the UK function that was called */
	u32 rc :32;

	u64 :64;
} __attribute__((packed));
ASSERT_SIZEOF_TYPE(union mali_ioctl_header, 8);

struct mali_ioctl_get_version {
	union mali_ioctl_header header;
	u16 major; /* [out] */
	u16 minor; /* [out] */
	u32 :32;
} __attribute__((packed));
ASSERT_SIZEOF_TYPE(struct mali_ioctl_get_version, 16);

struct mali_ioctl_mem_alloc {
	union mali_ioctl_header header;
	/* [in] */
	u64 va_pages;
	u64 commit_pages;
	u64 extent;
	/* [in/out] */
	u64 flags;
	/* [out] */
	u64 gpu_va;
	u16 va_alignment;

	u32 :32;
	u16 :16;
} __attribute__((packed));
ASSERT_SIZEOF_TYPE(struct mali_ioctl_mem_alloc, 56);

struct mali_ioctl_mem_import {
	union mali_ioctl_header header;
	/* [in] */
	u64 phandle;
	enum {
		MALI_MEM_IMPORT_TYPE_INVALID = 0,
		MALI_MEM_IMPORT_TYPE_UMP = 1,
		MALI_MEM_IMPORT_TYPE_UMM = 2,
		MALI_MEM_IMPORT_TYPE_USER_BUFFER = 3,
	} type :32;
	u32 :32;
	/* [in/out] */
	u64 flags;
	/* [out] */
	u64 gpu_va;
	u64 va_pages;
} __attribute__((packed));
/* FIXME: Size unconfirmed (haven't seen in a trace yet) */

struct mali_ioctl_mem_commit {
	union mali_ioctl_header header;
	/* [in] */
	u64 gpu_addr;
	u64 pages;
	/* [out] */
	u32 result_subcode;
	u32 :32;
} __attribute__((packed));
ASSERT_SIZEOF_TYPE(struct mali_ioctl_mem_commit, 32);

struct mali_ioctl_mem_query {
	union mali_ioctl_header header;
	/* [in] */
	u64 gpu_addr;
	enum {
		MALI_MEM_QUERY_COMMIT_SIZE = 1,
		MALI_MEM_QUERY_VA_SIZE     = 2,
		MALI_MEM_QUERY_FLAGS       = 3
	} query :32;
	u32 :32;
	/* [out] */
	u64 value;
} __attribute__((packed));
ASSERT_SIZEOF_TYPE(struct mali_ioctl_mem_query, 32);

struct mali_ioctl_mem_free {
	union mali_ioctl_header header;
	u64 gpu_addr; /* [in] */
} __attribute__((packed));
/* FIXME: Size unconfirmed (haven't seen in a trace yet) */

struct mali_ioctl_mem_flags_change {
	union mali_ioctl_header header;
	/* [in] */
	u64 gpu_va;
	u64 flags;
	u64 mask;
} __attribute__((packed));
/* FIXME: Size unconfirmed (haven't seen in a trace yet) */

struct mali_ioctl_mem_alias {
	union mali_ioctl_header header;
	/* [in/out] */
	u64 flags;
	/* [in] */
	u64 stride;
	u64 nents;
	u64 ai;
	/* [out] */
	u64 gpu_va;
	u64 va_pages;
} __attribute__((packed));

struct mali_ioctl_sync {
	union mali_ioctl_header header;
	u64 handle;
	PAD_PTR(void* user_addr);
	u64 size;
	enum {
		MALI_SYNC_TO_DEVICE = (1 << 0),
		MALI_SYNC_TO_CPU = (1 << 1),
	} type :8;
	u64 :56;
} __attribute__((packed));
ASSERT_SIZEOF_TYPE(struct mali_ioctl_sync, 40);

struct mali_ioctl_gpu_props_reg_dump {
	union mali_ioctl_header header;
	struct mali_gpu_core_props core;
	struct mali_gpu_l2_cache_props l2;
	u64 :64;
	struct mali_gpu_tiler_props tiler;
	struct mali_gpu_thread_props thread;

	struct mali_gpu_raw_props raw;

	/** This must be last member of the structure */
	struct mali_gpu_coherent_group_info coherency_info;
} __attribute__((packed));
ASSERT_SIZEOF_TYPE(struct mali_ioctl_gpu_props_reg_dump, 536);

struct mali_ioctl_set_flags {
	union mali_ioctl_header header;
#define MALI_CONTEXT_CREATE_FLAG_NONE 0
	u32 create_flags; /* [in] */
	u32 :32;
} __attribute__((packed));
ASSERT_SIZEOF_TYPE(struct mali_ioctl_set_flags, 16);

struct mali_ioctl_stream_create {
	union mali_ioctl_header header;
	/* [in] */
	char name[32];
	/* [out] */
	s32 fd;
	u32 :32;
} __attribute__((packed));
ASSERT_SIZEOF_TYPE(struct mali_ioctl_stream_create, 48);

struct mali_ioctl_job_submit {
	union mali_ioctl_header header;
	/* [in] */
	PAD_PTR(struct mali_jd_atom_v2 *addr);
	u32 nr_atoms;
	u32 stride;
} __attribute__((packed));
ASSERT_SIZEOF_TYPE(struct mali_ioctl_job_submit, 24);

struct mali_ioctl_get_context_id {
	union mali_ioctl_header header;
	/* [out] */
	s64 id;
} __attribute__((packed));
ASSERT_SIZEOF_TYPE(struct mali_ioctl_get_context_id, 16);

#undef PAD_PTR

/* For ioctl's we haven't written decoding stuff for yet */
typedef struct {
	union mali_ioctl_header header;
} __ioctl_placeholder;

#define MALI_IOCTL_TYPE_BASE  0x80
#define MALI_IOCTL_TYPE_MAX   0x82
#define MALI_IOCTL_TYPE_COUNT (MALI_IOCTL_TYPE_MAX - MALI_IOCTL_TYPE_BASE + 1)

#define MALI_IOCTL_GET_VERSION             (_IOWR(0x80,  0, struct mali_ioctl_get_version))
#define MALI_IOCTL_MEM_ALLOC               (_IOWR(0x82,  0, struct mali_ioctl_mem_alloc))
#define MALI_IOCTL_MEM_IMPORT              (_IOWR(0x82,  1, struct mali_ioctl_mem_import))
#define MALI_IOCTL_MEM_COMMIT              (_IOWR(0x82,  2, struct mali_ioctl_mem_commit))
#define MALI_IOCTL_MEM_QUERY               (_IOWR(0x82,  3, struct mali_ioctl_mem_query))
#define MALI_IOCTL_MEM_FREE                (_IOWR(0x82,  4, struct mali_ioctl_mem_free))
#define MALI_IOCTL_MEM_FLAGS_CHANGE        (_IOWR(0x82,  5, struct mali_ioctl_mem_flags_change))
#define MALI_IOCTL_MEM_ALIAS               (_IOWR(0x82,  6, struct mali_ioctl_mem_alias))
#define MALI_IOCTL_SYNC                    (_IOWR(0x82,  8, struct mali_ioctl_sync))
#define MALI_IOCTL_POST_TERM               (_IOWR(0x82,  9, __ioctl_placeholder))
#define MALI_IOCTL_HWCNT_SETUP             (_IOWR(0x82, 10, __ioctl_placeholder))
#define MALI_IOCTL_HWCNT_DUMP              (_IOWR(0x82, 11, __ioctl_placeholder))
#define MALI_IOCTL_HWCNT_CLEAR             (_IOWR(0x82, 12, __ioctl_placeholder))
#define MALI_IOCTL_GPU_PROPS_REG_DUMP      (_IOWR(0x82, 14, struct mali_ioctl_gpu_props_reg_dump))
#define MALI_IOCTL_FIND_CPU_OFFSET         (_IOWR(0x82, 15, __ioctl_placeholder))
#define MALI_IOCTL_GET_VERSION_NEW         (_IOWR(0x82, 16, struct mali_ioctl_get_version))
#define MALI_IOCTL_SET_FLAGS               (_IOWR(0x82, 18, struct mali_ioctl_set_flags))
#define MALI_IOCTL_SET_TEST_DATA           (_IOWR(0x82, 19, __ioctl_placeholder))
#define MALI_IOCTL_INJECT_ERROR            (_IOWR(0x82, 20, __ioctl_placeholder))
#define MALI_IOCTL_MODEL_CONTROL           (_IOWR(0x82, 21, __ioctl_placeholder))
#define MALI_IOCTL_KEEP_GPU_POWERED        (_IOWR(0x82, 22, __ioctl_placeholder))
#define MALI_IOCTL_FENCE_VALIDATE          (_IOWR(0x82, 23, __ioctl_placeholder))
#define MALI_IOCTL_STREAM_CREATE           (_IOWR(0x82, 24, struct mali_ioctl_stream_create))
#define MALI_IOCTL_GET_PROFILING_CONTROLS  (_IOWR(0x82, 25, __ioctl_placeholder))
#define MALI_IOCTL_SET_PROFILING_CONTROLS  (_IOWR(0x82, 26, __ioctl_placeholder))
#define MALI_IOCTL_DEBUGFS_MEM_PROFILE_ADD (_IOWR(0x82, 27, __ioctl_placeholder))
#define MALI_IOCTL_JOB_SUBMIT              (_IOWR(0x82, 28, struct mali_ioctl_job_submit))
#define MALI_IOCTL_DISJOINT_QUERY          (_IOWR(0x82, 29, __ioctl_placeholder))
#define MALI_IOCTL_GET_CONTEXT_ID          (_IOWR(0x82, 31, struct mali_ioctl_get_context_id))
#define MALI_IOCTL_TLSTREAM_ACQUIRE_V10_4  (_IOWR(0x82, 32, __ioctl_placeholder))
#define MALI_IOCTL_TLSTREAM_TEST           (_IOWR(0x82, 33, __ioctl_placeholder))
#define MALI_IOCTL_TLSTREAM_STATS          (_IOWR(0x82, 34, __ioctl_placeholder))
#define MALI_IOCTL_TLSTREAM_FLUSH          (_IOWR(0x82, 35, __ioctl_placeholder))
#define MALI_IOCTL_HWCNT_READER_SETUP      (_IOWR(0x82, 36, __ioctl_placeholder))
#define MALI_IOCTL_SET_PRFCNT_VALUES       (_IOWR(0x82, 37, __ioctl_placeholder))
#define MALI_IOCTL_SOFT_EVENT_UPDATE       (_IOWR(0x82, 38, __ioctl_placeholder))
#define MALI_IOCTL_MEM_JIT_INIT            (_IOWR(0x82, 39, __ioctl_placeholder))
#define MALI_IOCTL_TLSTREAM_ACQUIRE        (_IOWR(0x82, 40, __ioctl_placeholder))

#endif /* __MALI_IOCTL_H__ */