AuroraMiddleware/SPIRV-Cross
1
0
Fork 0
Code Issues Pull Requests Packages Projects Releases Wiki Activity

Initial commit.

Browse Source
This commit is contained in:
Hans-Kristian Arntzen 2016-03-02 18:09:16 +01:00
commit 75471fbb98
98 changed files with 14084 additions and 0 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

13
.gitignore vendored Normal file
View File

@ -0,0 +1,13 @@
*.o
*.d
*.txt
/test
/spir2cross
*.spv
/obj
/msvc/x64
/msvc/Debug
/msvc/Release
*.suo
*.sdf
*.opensdf

131
GLSL.std.450.h Normal file
View File

@ -0,0 +1,131 @@
/*
** Copyright (c) 2014-2016 The Khronos Group Inc.
**
** Permission is hereby granted, free of charge, to any person obtaining a copy
** of this software and/or associated documentation files (the "Materials"),
** to deal in the Materials without restriction, including without limitation
** the rights to use, copy, modify, merge, publish, distribute, sublicense,
** and/or sell copies of the Materials, and to permit persons to whom the
** Materials are furnished to do so, subject to the following conditions:
**
** The above copyright notice and this permission notice shall be included in
** all copies or substantial portions of the Materials.
**
** MODIFICATIONS TO THIS FILE MAY MEAN IT NO LONGER ACCURATELY REFLECTS KHRONOS
** STANDARDS. THE UNMODIFIED, NORMATIVE VERSIONS OF KHRONOS SPECIFICATIONS AND
** HEADER INFORMATION ARE LOCATED AT https://www.khronos.org/registry/
**
** THE MATERIALS ARE 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 MATERIALS OR THE USE OR OTHER DEALINGS
** IN THE MATERIALS.
*/
#ifndef GLSLstd450_H
#define GLSLstd450_H
static const int GLSLstd450Version = 100;
static const int GLSLstd450Revision = 1;
enum GLSLstd450 {
GLSLstd450Bad = 0, // Don't use
GLSLstd450Round = 1,
GLSLstd450RoundEven = 2,
GLSLstd450Trunc = 3,
GLSLstd450FAbs = 4,
GLSLstd450SAbs = 5,
GLSLstd450FSign = 6,
GLSLstd450SSign = 7,
GLSLstd450Floor = 8,
GLSLstd450Ceil = 9,
GLSLstd450Fract = 10,
GLSLstd450Radians = 11,
GLSLstd450Degrees = 12,
GLSLstd450Sin = 13,
GLSLstd450Cos = 14,
GLSLstd450Tan = 15,
GLSLstd450Asin = 16,
GLSLstd450Acos = 17,
GLSLstd450Atan = 18,
GLSLstd450Sinh = 19,
GLSLstd450Cosh = 20,
GLSLstd450Tanh = 21,
GLSLstd450Asinh = 22,
GLSLstd450Acosh = 23,
GLSLstd450Atanh = 24,
GLSLstd450Atan2 = 25,
GLSLstd450Pow = 26,
GLSLstd450Exp = 27,
GLSLstd450Log = 28,
GLSLstd450Exp2 = 29,
GLSLstd450Log2 = 30,
GLSLstd450Sqrt = 31,
GLSLstd450InverseSqrt = 32,
GLSLstd450Determinant = 33,
GLSLstd450MatrixInverse = 34,
GLSLstd450Modf = 35, // second operand needs an OpVariable to write to
GLSLstd450ModfStruct = 36, // no OpVariable operand
GLSLstd450FMin = 37,
GLSLstd450UMin = 38,
GLSLstd450SMin = 39,
GLSLstd450FMax = 40,
GLSLstd450UMax = 41,
GLSLstd450SMax = 42,
GLSLstd450FClamp = 43,
GLSLstd450UClamp = 44,
GLSLstd450SClamp = 45,
GLSLstd450FMix = 46,
GLSLstd450IMix = 47, // Reserved
GLSLstd450Step = 48,
GLSLstd450SmoothStep = 49,
GLSLstd450Fma = 50,
GLSLstd450Frexp = 51, // second operand needs an OpVariable to write to
GLSLstd450FrexpStruct = 52, // no OpVariable operand
GLSLstd450Ldexp = 53,
GLSLstd450PackSnorm4x8 = 54,
GLSLstd450PackUnorm4x8 = 55,
GLSLstd450PackSnorm2x16 = 56,
GLSLstd450PackUnorm2x16 = 57,
GLSLstd450PackHalf2x16 = 58,
GLSLstd450PackDouble2x32 = 59,
GLSLstd450UnpackSnorm2x16 = 60,
GLSLstd450UnpackUnorm2x16 = 61,
GLSLstd450UnpackHalf2x16 = 62,
GLSLstd450UnpackSnorm4x8 = 63,
GLSLstd450UnpackUnorm4x8 = 64,
GLSLstd450UnpackDouble2x32 = 65,
GLSLstd450Length = 66,
GLSLstd450Distance = 67,
GLSLstd450Cross = 68,
GLSLstd450Normalize = 69,
GLSLstd450FaceForward = 70,
GLSLstd450Reflect = 71,
GLSLstd450Refract = 72,
GLSLstd450FindILsb = 73,
GLSLstd450FindSMsb = 74,
GLSLstd450FindUMsb = 75,
GLSLstd450InterpolateAtCentroid = 76,
GLSLstd450InterpolateAtSample = 77,
GLSLstd450InterpolateAtOffset = 78,
GLSLstd450NMin = 79,
GLSLstd450NMax = 80,
GLSLstd450NClamp = 81,
GLSLstd450Count
};
#endif // #ifndef GLSLstd450_H

202
LICENSE Normal file
View File

@ -0,0 +1,202 @@
Apache License
Version 2.0, January 2004
http://www.apache.org/licenses/
TERMS AND CONDITIONS FOR USE, REPRODUCTION, AND DISTRIBUTION
1. Definitions.
"License" shall mean the terms and conditions for use, reproduction,
and distribution as defined by Sections 1 through 9 of this document.
"Licensor" shall mean the copyright owner or entity authorized by
the copyright owner that is granting the License.
"Legal Entity" shall mean the union of the acting entity and all
other entities that control, are controlled by, or are under common
control with that entity. For the purposes of this definition,
"control" means (i) the power, direct or indirect, to cause the
direction or management of such entity, whether by contract or
otherwise, or (ii) ownership of fifty percent (50%) or more of the
outstanding shares, or (iii) beneficial ownership of such entity.
"You" (or "Your") shall mean an individual or Legal Entity
exercising permissions granted by this License.
"Source" form shall mean the preferred form for making modifications,
including but not limited to software source code, documentation
source, and configuration files.
"Object" form shall mean any form resulting from mechanical
transformation or translation of a Source form, including but
not limited to compiled object code, generated documentation,
and conversions to other media types.
"Work" shall mean the work of authorship, whether in Source or
Object form, made available under the License, as indicated by a
copyright notice that is included in or attached to the work
(an example is provided in the Appendix below).
"Derivative Works" shall mean any work, whether in Source or Object
form, that is based on (or derived from) the Work and for which the
editorial revisions, annotations, elaborations, or other modifications
represent, as a whole, an original work of authorship. For the purposes
of this License, Derivative Works shall not include works that remain
separable from, or merely link (or bind by name) to the interfaces of,
the Work and Derivative Works thereof.
"Contribution" shall mean any work of authorship, including
the original version of the Work and any modifications or additions
to that Work or Derivative Works thereof, that is intentionally
submitted to Licensor for inclusion in the Work by the copyright owner
or by an individual or Legal Entity authorized to submit on behalf of
the copyright owner. For the purposes of this definition, "submitted"
means any form of electronic, verbal, or written communication sent
to the Licensor or its representatives, including but not limited to
communication on electronic mailing lists, source code control systems,
and issue tracking systems that are managed by, or on behalf of, the
Licensor for the purpose of discussing and improving the Work, but
excluding communication that is conspicuously marked or otherwise
designated in writing by the copyright owner as "Not a Contribution."
"Contributor" shall mean Licensor and any individual or Legal Entity
on behalf of whom a Contribution has been received by Licensor and
subsequently incorporated within the Work.
2. Grant of Copyright License. Subject to the terms and conditions of
this License, each Contributor hereby grants to You a perpetual,
worldwide, non-exclusive, no-charge, royalty-free, irrevocable
copyright license to reproduce, prepare Derivative Works of,
publicly display, publicly perform, sublicense, and distribute the
Work and such Derivative Works in Source or Object form.
3. Grant of Patent License. Subject to the terms and conditions of
this License, each Contributor hereby grants to You a perpetual,
worldwide, non-exclusive, no-charge, royalty-free, irrevocable
(except as stated in this section) patent license to make, have made,
use, offer to sell, sell, import, and otherwise transfer the Work,
where such license applies only to those patent claims licensable
by such Contributor that are necessarily infringed by their
Contribution(s) alone or by combination of their Contribution(s)
with the Work to which such Contribution(s) was submitted. If You
institute patent litigation against any entity (including a
cross-claim or counterclaim in a lawsuit) alleging that the Work
or a Contribution incorporated within the Work constitutes direct
or contributory patent infringement, then any patent licenses
granted to You under this License for that Work shall terminate
as of the date such litigation is filed.
4. Redistribution. You may reproduce and distribute copies of the
Work or Derivative Works thereof in any medium, with or without
modifications, and in Source or Object form, provided that You
meet the following conditions:
(a) You must give any other recipients of the Work or
Derivative Works a copy of this License; and
(b) You must cause any modified files to carry prominent notices
stating that You changed the files; and
(c) You must retain, in the Source form of any Derivative Works
that You distribute, all copyright, patent, trademark, and
attribution notices from the Source form of the Work,
excluding those notices that do not pertain to any part of
the Derivative Works; and
(d) If the Work includes a "NOTICE" text file as part of its
distribution, then any Derivative Works that You distribute must
include a readable copy of the attribution notices contained
within such NOTICE file, excluding those notices that do not
pertain to any part of the Derivative Works, in at least one
of the following places: within a NOTICE text file distributed
as part of the Derivative Works; within the Source form or
documentation, if provided along with the Derivative Works; or,
within a display generated by the Derivative Works, if and
wherever such third-party notices normally appear. The contents
of the NOTICE file are for informational purposes only and
do not modify the License. You may add Your own attribution
notices within Derivative Works that You distribute, alongside
or as an addendum to the NOTICE text from the Work, provided
that such additional attribution notices cannot be construed
as modifying the License.
You may add Your own copyright statement to Your modifications and
may provide additional or different license terms and conditions
for use, reproduction, or distribution of Your modifications, or
for any such Derivative Works as a whole, provided Your use,
reproduction, and distribution of the Work otherwise complies with
the conditions stated in this License.
5. Submission of Contributions. Unless You explicitly state otherwise,
any Contribution intentionally submitted for inclusion in the Work
by You to the Licensor shall be under the terms and conditions of
this License, without any additional terms or conditions.
Notwithstanding the above, nothing herein shall supersede or modify
the terms of any separate license agreement you may have executed
with Licensor regarding such Contributions.
6. Trademarks. This License does not grant permission to use the trade
names, trademarks, service marks, or product names of the Licensor,
except as required for reasonable and customary use in describing the
origin of the Work and reproducing the content of the NOTICE file.
7. Disclaimer of Warranty. Unless required by applicable law or
agreed to in writing, Licensor provides the Work (and each
Contributor provides its Contributions) on an "AS IS" BASIS,
WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or
implied, including, without limitation, any warranties or conditions
of TITLE, NON-INFRINGEMENT, MERCHANTABILITY, or FITNESS FOR A
PARTICULAR PURPOSE. You are solely responsible for determining the
appropriateness of using or redistributing the Work and assume any
risks associated with Your exercise of permissions under this License.
8. Limitation of Liability. In no event and under no legal theory,
whether in tort (including negligence), contract, or otherwise,
unless required by applicable law (such as deliberate and grossly
negligent acts) or agreed to in writing, shall any Contributor be
liable to You for damages, including any direct, indirect, special,
incidental, or consequential damages of any character arising as a
result of this License or out of the use or inability to use the
Work (including but not limited to damages for loss of goodwill,
work stoppage, computer failure or malfunction, or any and all
other commercial damages or losses), even if such Contributor
has been advised of the possibility of such damages.
9. Accepting Warranty or Additional Liability. While redistributing
the Work or Derivative Works thereof, You may choose to offer,
and charge a fee for, acceptance of support, warranty, indemnity,
or other liability obligations and/or rights consistent with this
License. However, in accepting such obligations, You may act only
on Your own behalf and on Your sole responsibility, not on behalf
of any other Contributor, and only if You agree to indemnify,
defend, and hold each Contributor harmless for any liability
incurred by, or claims asserted against, such Contributor by reason
of your accepting any such warranty or additional liability.
END OF TERMS AND CONDITIONS
APPENDIX: How to apply the Apache License to your work.
To apply the Apache License to your work, attach the following
boilerplate notice, with the fields enclosed by brackets "[]"
replaced with your own identifying information. (Don't include
the brackets!) The text should be enclosed in the appropriate
comment syntax for the file format. We also recommend that a
file or class name and description of purpose be included on the
same "printed page" as the copyright notice for easier
identification within third-party archives.
Copyright [yyyy] [name of copyright owner]
Licensed under the Apache License, Version 2.0 (the "License");
you may not use this file except in compliance with the License.
You may obtain a copy of the License at
http://www.apache.org/licenses/LICENSE-2.0
Unless required by applicable law or agreed to in writing, software
distributed under the License is distributed on an "AS IS" BASIS,
WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
See the License for the specific language governing permissions and
limitations under the License.

27
Makefile Normal file
View File

@ -0,0 +1,27 @@
TARGET := spir2cross
SOURCES := $(wildcard *.cpp)
OBJECTS := $(SOURCES:.cpp=.o)
DEPS := $(OBJECTS:.o=.d)
CXXFLAGS += -std=c++11 -Wall -Wextra
ifeq ($(DEBUG), 1)
CXXFLAGS += -O0 -gdwarf-2
else
CXXFLAGS += -O2 -gdwarf-2
endif
all: $(TARGET)
-include $(DEPS)
$(TARGET): $(OBJECTS)
$(CXX) -o $@ $(OBJECTS) $(LDFLAGS)
%.o: %.cpp
$(CXX) -c -o $@ $< $(CXXFLAGS) -MMD
clean:
rm -f $(TARGET) $(OBJECTS)
.PHONY: clean

75
README.md Normal file
View File

@ -0,0 +1,75 @@
# SPIR2CROSS
SPIR2CROSS is a tool designed for parsing and converting SPIR-V to other shader languages.
## Features
- Convert SPIR-V to readable, usable and efficient GLSL
- Convert SPIR-V to debuggable C++ [EXPERIMENTAL]
- Reflection API to simplify the creation of Vulkan pipeline layouts
- Reflection API to modify and tweak OpDecorations
- Supports "all" of vertex, fragment, tessellation, geometry and compute shaders.
SPIR2CROSS tries hard to emit readable and clean output from the SPIR-V.
The goal is to emit GLSL that looks like it was written by a human and not awkward IR/assembly-like code.
NOTE: Individual features are expected to be mostly complete, but it is possible that certain obscure GLSL features are not yet supported.
However, most missing features are expected to be "trivial" improvements at this stage.
Occasionally, missing features is due to glslangValidator's lack of proper support for that feature making testing hard.
## Building
SPIR2CROSS has been tested on Linux, OSX and Windows.
### Linux and OSX
Just run `make` on the command line. A recent GCC (4.8+) or Clang (3.x+) compiler is required as SPIR2CROSS uses C++11 extensively.
### Windows
MinGW-w64 based compilation works with `make`, and an MSVC 2013 solution is also included.
## Usage
### Creating a SPIR-V file from GLSL with glslang
```
glslangValidator -H -V -o test.spv test.frag
```
### Converting a SPIR-V file to GLSL ES
```
glslangValidator -H -V -o test.spv shaders/comp/basic.comp
./spir2cross --version 310 --es test.spv
```
#### Converting to desktop GLSL
```
glslangValidator -H -V -o test.spv shaders/comp/basic.comp
./spir2cross --version 330 test.spv --output test.comp
```
#### Disable prettifying optimizations
```
glslangValidator -H -V -o test.spv shaders/comp/basic.comp
./spir2cross --version 310 --es test.spv --output test.comp --force-temporary
```
## ABI concerns
### SPIR-V headers
The current repository uses the latest SPIR-V and GLSL.std.450 headers.
SPIR-V files created from older headers could have ABI issues.
## Regression testing
In shaders/ a collection of shaders are maintained for purposes of regression testing.
The current reference output is contained in reference/.
`./test_shaders.py shaders` can be run to perform regression testing.
Currently, the Mali Offline Compiler `malisc` is used to verify the outputs from SPIR2CROSS.

79
include/spir2cross/barrier.hpp Normal file
View File

@ -0,0 +1,79 @@
/*
* Copyright 2015-2016 ARM Limited
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
* You may obtain a copy of the License at
*
* http://www.apache.org/licenses/LICENSE-2.0
*
* Unless required by applicable law or agreed to in writing, software
* distributed under the License is distributed on an "AS IS" BASIS,
* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
* See the License for the specific language governing permissions and
* limitations under the License.
*/
#ifndef SPIR2CROSS_BARRIER_HPP
#define SPIR2CROSS_BARRIER_HPP
#include <thread>
#include <atomic>
namespace spir2cross
{
class Barrier
{
public:
Barrier()
{
count.store(0);
iteration.store(0);
}
void set_release_divisor(unsigned divisor)
{
this->divisor = divisor;
}
static inline void memoryBarrier()
{
std::atomic_thread_fence(std::memory_order_seq_cst);
}
void reset_counter()
{
count.store(0);
iteration.store(0);
}
void wait()
{
unsigned target_iteration = iteration.load(std::memory_order_relaxed) + 1;
// Overflows cleanly.
unsigned target_count = divisor * target_iteration;
// Barriers don't enforce memory ordering.
// Be as relaxed about the barrier as we possibly can!
unsigned c = count.fetch_add(1u, std::memory_order_relaxed);
if (c + 1 == target_count)
{
iteration.store(target_iteration, std::memory_order_relaxed);
}
else
{
// If we have more threads than the CPU, don't hog the CPU for very long periods of time.
while (iteration.load(std::memory_order_relaxed) != target_iteration)
std::this_thread::yield();
}
}
private:
unsigned divisor = 1;
std::atomic<unsigned> count;
std::atomic<unsigned> iteration;
};
}
#endif

137
include/spir2cross/external_interface.h Normal file
View File

@ -0,0 +1,137 @@
/*
* Copyright 2015-2016 ARM Limited
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
* You may obtain a copy of the License at
*
* http://www.apache.org/licenses/LICENSE-2.0
*
* Unless required by applicable law or agreed to in writing, software
* distributed under the License is distributed on an "AS IS" BASIS,
* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
* See the License for the specific language governing permissions and
* limitations under the License.
*/
#ifndef SPIR2CROSS_EXTERNAL_INTERFACE_H
#define SPIR2CROSS_EXTERNAL_INTERFACE_H
#ifdef __cplusplus
extern "C" {
#endif
#include <stddef.h>
typedef struct spir2cross_shader spir2cross_shader_t;
struct spir2cross_interface
{
spir2cross_shader_t* (*construct)(void);
void (*destruct)(spir2cross_shader_t *thiz);
void (*invoke)(spir2cross_shader_t *thiz);
};
void spir2cross_set_stage_input(spir2cross_shader_t *thiz,
unsigned location, void *data, size_t size);
void spir2cross_set_stage_output(spir2cross_shader_t *thiz,
unsigned location, void *data, size_t size);
void spir2cross_set_push_constant(spir2cross_shader_t *thiz,
void *data, size_t size);
void spir2cross_set_uniform_constant(spir2cross_shader_t *thiz,
unsigned location,
void *data, size_t size);
void spir2cross_set_resource(spir2cross_shader_t *thiz,
unsigned set, unsigned binding,
void **data, size_t size);
const struct spir2cross_interface* spir2cross_get_interface(void);
typedef enum spir2cross_builtin
{
SPIR2CROSS_BUILTIN_POSITION = 0,
SPIR2CROSS_BUILTIN_FRAG_COORD = 1,
SPIR2CROSS_BUILTIN_WORK_GROUP_ID = 2,
SPIR2CROSS_BUILTIN_NUM_WORK_GROUPS = 3,
SPIR2CROSS_NUM_BUILTINS
} spir2cross_builtin;
void spir2cross_set_builtin(spir2cross_shader_t *thiz,
spir2cross_builtin builtin,
void *data, size_t size);
#define SPIR2CROSS_NUM_DESCRIPTOR_SETS 4
#define SPIR2CROSS_NUM_DESCRIPTOR_BINDINGS 16
#define SPIR2CROSS_NUM_STAGE_INPUTS 16
#define SPIR2CROSS_NUM_STAGE_OUTPUTS 16
#define SPIR2CROSS_NUM_UNIFORM_CONSTANTS 32
enum spir2cross_format
{
SPIR2CROSS_FORMAT_R8_UNORM = 0,
SPIR2CROSS_FORMAT_R8G8_UNORM = 1,
SPIR2CROSS_FORMAT_R8G8B8_UNORM = 2,
SPIR2CROSS_FORMAT_R8G8B8A8_UNORM = 3,
SPIR2CROSS_NUM_FORMATS
};
enum spir2cross_wrap
{
SPIR2CROSS_WRAP_CLAMP_TO_EDGE = 0,
SPIR2CROSS_WRAP_REPEAT = 1,
SPIR2CROSS_NUM_WRAP
};
enum spir2cross_filter
{
SPIR2CROSS_FILTER_NEAREST = 0,
SPIR2CROSS_FILTER_LINEAR = 1,
SPIR2CROSS_NUM_FILTER
};
enum spir2cross_mipfilter
{
SPIR2CROSS_MIPFILTER_BASE = 0,
SPIR2CROSS_MIPFILTER_NEAREST = 1,
SPIR2CROSS_MIPFILTER_LINEAR = 2,
SPIR2CROSS_NUM_MIPFILTER
};
struct spir2cross_miplevel
{
const void *data;
unsigned width, height;
size_t stride;
};
struct spir2cross_sampler_info
{
const struct spir2cross_miplevel *mipmaps;
unsigned num_mipmaps;
enum spir2cross_format format;
enum spir2cross_wrap wrap_s;
enum spir2cross_wrap wrap_t;
enum spir2cross_filter min_filter;
enum spir2cross_filter mag_filter;
enum spir2cross_mipfilter mip_filter;
};
typedef struct spir2cross_sampler_2d spir2cross_sampler_2d_t;
spir2cross_sampler_2d_t *spir2cross_create_sampler_2d(const struct spir2cross_sampler_info *info);
void spir2cross_destroy_sampler_2d(spir2cross_sampler_2d_t *samp);
#ifdef __cplusplus
}
#endif
#endif

45
include/spir2cross/image.hpp Normal file
View File

@ -0,0 +1,45 @@
/*
* Copyright 2015-2016 ARM Limited
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
* You may obtain a copy of the License at
*
* http://www.apache.org/licenses/LICENSE-2.0
*
* Unless required by applicable law or agreed to in writing, software
* distributed under the License is distributed on an "AS IS" BASIS,
* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
* See the License for the specific language governing permissions and
* limitations under the License.
*/
#ifndef SPIR2CROSS_IMAGE_HPP
#define SPIR2CROSS_IMAGE_HPP
#ifndef GLM_SWIZZLE
#define GLM_SWIZZLE
#endif
#ifndef GLM_FORCE_RADIANS
#define GLM_FORCE_RADIANS
#endif
#include <glm/glm.hpp>
namespace spir2cross
{
template<typename T>
struct image2DBase
{
virtual ~image2DBase() = default;
inline virtual T load(glm::ivec2 coord) { return T(0, 0, 0, 1); }
inline virtual void store(glm::ivec2 coord, const T &v) {}
};
typedef image2DBase<glm::vec4> image2D;
typedef image2DBase<glm::ivec4> iimage2D;
typedef image2DBase<glm::uvec4> uimage2D;
}
#endif

531
include/spir2cross/internal_interface.hpp Normal file
View File

@ -0,0 +1,531 @@
/*
* Copyright 2015-2016 ARM Limited
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
* You may obtain a copy of the License at
*
* http://www.apache.org/licenses/LICENSE-2.0
*
* Unless required by applicable law or agreed to in writing, software
* distributed under the License is distributed on an "AS IS" BASIS,
* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
* See the License for the specific language governing permissions and
* limitations under the License.
*/
#ifndef SPIR2CROSS_INTERNAL_INTERFACE_HPP
#define SPIR2CROSS_INTERNAL_INTERFACE_HPP
// This file must only be included by the shader generated by spir2cross!
#ifndef GLM_SWIZZLE
#define GLM_SWIZZLE
#endif
#ifndef GLM_FORCE_RADIANS
#define GLM_FORCE_RADIANS
#endif
#include <glm/glm.hpp>
#include <assert.h>
#include <stdint.h>
#include "external_interface.h"
#include "barrier.hpp"
#include "thread_group.hpp"
#include "sampler.hpp"
#include "image.hpp"
namespace internal
{
// Adaptor helpers to adapt GLSL access chain syntax to C++.
// Don't bother with arrays of arrays on uniforms ...
// Would likely need horribly complex variadic template munging.
template<typename T>
struct Interface
{
enum { ArraySize = 1, Size = sizeof(T) };
Interface() : ptr(0) {}
T& get()
{
assert(ptr);
return *ptr;
}
T *ptr;
};
// For array types, return a pointer instead.
template<typename T, unsigned U>
struct Interface<T[U]>
{
enum { ArraySize = U, Size = U * sizeof(T) };
Interface() : ptr(0) {}
T* get()
{
assert(ptr);
return ptr;
}
T *ptr;
};
// For case when array size is 1, avoid double dereference.
template<typename T>
struct PointerInterface
{
enum { ArraySize = 1, Size = sizeof(T*) };
enum { PreDereference = true };
PointerInterface() : ptr(0) {}
T& get()
{
assert(ptr);
return *ptr;
}
T *ptr;
};
// Automatically converts a pointer down to reference to match GLSL syntax.
template<typename T>
struct DereferenceAdaptor
{
DereferenceAdaptor(T **ptr) : ptr(ptr) {}
T& operator[](unsigned index) const { return *(ptr[index]); }
T **ptr;
};
// We can't have a linear array of T* since T* can be an abstract type in case of samplers.
// We also need a list of pointers since we can have run-time length SSBOs.
template<typename T, unsigned U>
struct PointerInterface<T[U]>
{
enum { ArraySize = U, Size = sizeof(T*) * U };
enum { PreDereference = false };
PointerInterface() : ptr(0) {}
DereferenceAdaptor<T> get()
{
assert(ptr);
return DereferenceAdaptor<T>(ptr);
}
T **ptr;
};
// Resources can be more abstract and be unsized,
// so we need to have an array of pointers for those cases.
template<typename T> struct Resource : PointerInterface<T> {};
// POD with no unknown sizes, so we can express these as flat arrays.
template<typename T> struct UniformConstant : Interface<T> {};
template<typename T> struct StageInput : Interface<T> {};
template<typename T> struct StageOutput : Interface<T> {};
template<typename T> struct PushConstant : Interface<T> {};
}
struct spir2cross_shader
{
struct PPSize
{
PPSize() : ptr(0), size(0) {}
void **ptr;
size_t size;
};
struct PPSizeResource
{
PPSizeResource() : ptr(0), size(0), pre_dereference(false) {}
void **ptr;
size_t size;
bool pre_dereference;
};
PPSizeResource resources[SPIR2CROSS_NUM_DESCRIPTOR_SETS][SPIR2CROSS_NUM_DESCRIPTOR_BINDINGS];
PPSize stage_inputs[SPIR2CROSS_NUM_STAGE_INPUTS];
PPSize stage_outputs[SPIR2CROSS_NUM_STAGE_OUTPUTS];
PPSize uniform_constants[SPIR2CROSS_NUM_UNIFORM_CONSTANTS];
PPSize push_constant;
PPSize builtins[SPIR2CROSS_NUM_BUILTINS];
template<typename U>
void register_builtin(spir2cross_builtin builtin, const U& value)
{
assert(!builtins[builtin].ptr);
builtins[builtin].ptr = (void**)&value.ptr;
builtins[builtin].size = sizeof(*value.ptr) * U::ArraySize;
}
void set_builtin(spir2cross_builtin builtin, void *data, size_t size)
{
assert(builtins[builtin].ptr);
assert(size >= builtins[builtin].size);
*builtins[builtin].ptr = data;
}
template<typename U>
void register_resource(const internal::Resource<U> &value, unsigned set, unsigned binding)
{
assert(set < SPIR2CROSS_NUM_DESCRIPTOR_SETS);
assert(binding < SPIR2CROSS_NUM_DESCRIPTOR_BINDINGS);
assert(!resources[set][binding].ptr);
resources[set][binding].ptr = (void**)&value.ptr;
resources[set][binding].size = internal::Resource<U>::Size;
resources[set][binding].pre_dereference = internal::Resource<U>::PreDereference;
}
template<typename U>
void register_stage_input(const internal::StageInput<U> &value, unsigned location)
{
assert(location < SPIR2CROSS_NUM_STAGE_INPUTS);
assert(!stage_inputs[location].ptr);
stage_inputs[location].ptr = (void**)&value.ptr;
stage_inputs[location].size = internal::StageInput<U>::Size;
}
template<typename U>
void register_stage_output(const internal::StageOutput<U> &value, unsigned location)
{
assert(location < SPIR2CROSS_NUM_STAGE_OUTPUTS);
assert(!stage_outputs[location].ptr);
stage_outputs[location].ptr = (void**)&value.ptr;
stage_outputs[location].size = internal::StageOutput<U>::Size;
}
template<typename U>
void register_uniform_constant(const internal::UniformConstant<U> &value, unsigned location)
{
assert(location < SPIR2CROSS_NUM_UNIFORM_CONSTANTS);
assert(!uniform_constants[location].ptr);
uniform_constants[location].ptr = (void**)&value.ptr;
uniform_constants[location].size = internal::UniformConstant<U>::Size;
}
template<typename U>
void register_push_constant(const internal::PushConstant<U> &value)
{
assert(!push_constant.ptr);
push_constant.ptr = (void**)&value.ptr;
push_constant.size = internal::PushConstant<U>::Size;
}
void set_stage_input(unsigned location, void *data, size_t size)
{
assert(location < SPIR2CROSS_NUM_STAGE_INPUTS);
assert(stage_inputs[location].ptr);
assert(size >= stage_inputs[location].size);
*stage_inputs[location].ptr = data;
}
void set_stage_output(unsigned location, void *data, size_t size)
{
assert(location < SPIR2CROSS_NUM_STAGE_OUTPUTS);
assert(stage_outputs[location].ptr);
assert(size >= stage_outputs[location].size);
*stage_outputs[location].ptr = data;
}
void set_uniform_constant(unsigned location, void *data, size_t size)
{
assert(location < SPIR2CROSS_NUM_UNIFORM_CONSTANTS);
assert(uniform_constants[location].ptr);
assert(size >= uniform_constants[location].size);
*uniform_constants[location].ptr = data;
}
void set_push_constant(void *data, size_t size)
{
assert(push_constant.ptr);
assert(size >= push_constant.size);
*push_constant.ptr = data;
}
void set_resource(unsigned set, unsigned binding,
void **data, size_t size)
{
assert(set < SPIR2CROSS_NUM_DESCRIPTOR_SETS);
assert(binding < SPIR2CROSS_NUM_DESCRIPTOR_BINDINGS);
assert(resources[set][binding].ptr);
assert(size >= resources[set][binding].size);
// We're using the regular PointerInterface, dereference ahead of time.
if (resources[set][binding].pre_dereference)
*resources[set][binding].ptr = *data;
else
*resources[set][binding].ptr = data;
}
};
namespace spir2cross
{
template<typename T>
struct BaseShader : spir2cross_shader
{
void invoke()
{
static_cast<T*>(this)->main();
}
};
struct FragmentResources
{
internal::StageOutput<glm::vec4> gl_FragCoord;
void init(spir2cross_shader &s)
{
s.register_builtin(SPIR2CROSS_BUILTIN_FRAG_COORD, gl_FragCoord);
}
#define gl_FragCoord __res->gl_FragCoord.get()
};
template<typename T, typename Res>
struct FragmentShader : BaseShader<FragmentShader<T, Res> >
{
inline void main()
{
impl.main();
}
FragmentShader()
{
resources.init(*this);
impl.__res = &resources;
}
T impl;
Res resources;
};
struct VertexResources
{
internal::StageOutput<glm::vec4> gl_Position;
void init(spir2cross_shader &s)
{
s.register_builtin(SPIR2CROSS_BUILTIN_POSITION, gl_Position);
}
#define gl_Position __res->gl_Position.get()
};
template<typename T, typename Res>
struct VertexShader : BaseShader<VertexShader<T, Res> >
{
inline void main()
{
impl.main();
}
VertexShader()
{
resources.init(*this);
impl.__res = &resources;
}
T impl;
Res resources;
};
struct TessEvaluationResources
{
inline void init(spir2cross_shader&) {}
};
template<typename T, typename Res>
struct TessEvaluationShader : BaseShader<TessEvaluationShader<T, Res> >
{
inline void main()
{
impl.main();
}
TessEvaluationShader()
{
resources.init(*this);
impl.__res = &resources;
}
T impl;
Res resources;
};
struct TessControlResources
{
inline void init(spir2cross_shader&) {}
};
template<typename T, typename Res>
struct TessControlShader : BaseShader<TessControlShader<T, Res> >
{
inline void main()
{
impl.main();
}
TessControlShader()
{
resources.init(*this);
impl.__res = &resources;
}
T impl;
Res resources;
};
struct GeometryResources
{
inline void init(spir2cross_shader&) {}
};
template<typename T, typename Res>
struct GeometryShader : BaseShader<GeometryShader<T, Res> >
{
inline void main()
{
impl.main();
}
GeometryShader()
{
resources.init(*this);
impl.__res = &resources;
}
T impl;
Res resources;
};
struct ComputeResources
{
internal::StageInput<glm::uvec3> gl_WorkGroupID__;
internal::StageInput<glm::uvec3> gl_NumWorkGroups__;
void init(spir2cross_shader &s)
{
s.register_builtin(SPIR2CROSS_BUILTIN_WORK_GROUP_ID, gl_WorkGroupID__);
s.register_builtin(SPIR2CROSS_BUILTIN_NUM_WORK_GROUPS, gl_NumWorkGroups__);
}
#define gl_WorkGroupID __res->gl_WorkGroupID__.get()
#define gl_NumWorkGroups __res->gl_NumWorkGroups.get()
Barrier barrier__;
#define barrier() __res->barrier__.wait()
};
struct ComputePrivateResources
{
uint32_t gl_LocalInvocationIndex__;
#define gl_LocalInvocationIndex __priv_res.gl_LocalInvocationIndex__
glm::uvec3 gl_LocalInvocationID__;
#define gl_LocalInvocationID __priv_res.gl_LocalInvocationID__
glm::uvec3 gl_GlobalInvocationID__;
#define gl_GlobalInvocationID __priv_res.gl_GlobalInvocationID__
};
template<typename T, typename Res, unsigned WorkGroupX, unsigned WorkGroupY, unsigned WorkGroupZ>
struct ComputeShader : BaseShader<ComputeShader<T, Res, WorkGroupX, WorkGroupY, WorkGroupZ> >
{
inline void main()
{
resources.barrier__.reset_counter();
for (unsigned z = 0; z < WorkGroupZ; z++)
for (unsigned y = 0; y < WorkGroupY; y++)
for (unsigned x = 0; x < WorkGroupX; x++)
impl[z][y][x].__priv_res.gl_GlobalInvocationID__ =
glm::uvec3(WorkGroupX, WorkGroupY, WorkGroupZ) * resources.gl_WorkGroupID__.get() +
glm::uvec3(x, y, z);
group.run();
group.wait();
}
ComputeShader()
: group(&impl[0][0][0])
{
resources.init(*this);
resources.barrier__.set_release_divisor(WorkGroupX * WorkGroupY * WorkGroupZ);
unsigned i = 0;
for (unsigned z = 0; z < WorkGroupZ; z++)
{
for (unsigned y = 0; y < WorkGroupY; y++)
{
for (unsigned x = 0; x < WorkGroupX; x++)
{
impl[z][y][x].__priv_res.gl_LocalInvocationID__ =
glm::uvec3(x, y, z);
impl[z][y][x].__priv_res.gl_LocalInvocationIndex__ = i++;
impl[z][y][x].__res = &resources;
}
}
}
}
T impl[WorkGroupZ][WorkGroupY][WorkGroupX];
ThreadGroup<T, WorkGroupX * WorkGroupY * WorkGroupZ> group;
Res resources;
};
inline void memoryBarrierShared() { Barrier::memoryBarrier(); }
inline void memoryBarrier() { Barrier::memoryBarrier(); }
// TODO: Rest of the barriers.
// Atomics
template<typename T>
inline T atomicAdd(T &v, T a)
{
static_assert(sizeof(std::atomic<T>) == sizeof(T), "Cannot cast properly to std::atomic<T>.");
// We need explicit memory barriers in GLSL to enfore any ordering.
// FIXME: Can we really cast this? There is no other way I think ...
return std::atomic_fetch_add_explicit(reinterpret_cast<std::atomic<T>*>(&v),
a, std::memory_order_relaxed);
}
}
void spir2cross_set_stage_input(spir2cross_shader_t *shader, unsigned location, void *data, size_t size)
{
shader->set_stage_input(location, data, size);
}
void spir2cross_set_stage_output(spir2cross_shader_t *shader, unsigned location, void *data, size_t size)
{
shader->set_stage_output(location, data, size);
}
void spir2cross_set_uniform_constant(spir2cross_shader_t *shader, unsigned location, void *data, size_t size)
{
shader->set_uniform_constant(location, data, size);
}
void spir2cross_set_resource(spir2cross_shader_t *shader, unsigned set, unsigned binding, void **data, size_t size)
{
shader->set_resource(set, binding, data, size);
}
void spir2cross_set_push_constant(spir2cross_shader_t *shader, void *data, size_t size)
{
shader->set_push_constant(data, size);
}
void spir2cross_set_builtin(spir2cross_shader_t *shader, spir2cross_builtin builtin,
void *data, size_t size)
{
shader->set_builtin(builtin, data, size);
}
#endif

101
include/spir2cross/sampler.hpp Normal file
View File

@ -0,0 +1,101 @@
/*
* Copyright 2015-2016 ARM Limited
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
* You may obtain a copy of the License at
*
* http://www.apache.org/licenses/LICENSE-2.0
*
* Unless required by applicable law or agreed to in writing, software
* distributed under the License is distributed on an "AS IS" BASIS,
* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
* See the License for the specific language governing permissions and
* limitations under the License.
*/
#ifndef SPIR2CROSS_SAMPLER_HPP
#define SPIR2CROSS_SAMPLER_HPP
#include <vector>
namespace spir2cross
{
struct spir2cross_sampler_2d
{
inline virtual ~spir2cross_sampler_2d() {}
};
template<typename T>
struct sampler2DBase : spir2cross_sampler_2d
{
sampler2DBase(const spir2cross_sampler_info *info)
{
mips.insert(mips.end(), info->mipmaps, info->mipmaps + info->num_mipmaps);
format = info->format;
wrap_s = info->wrap_s;
wrap_t = info->wrap_t;
min_filter = info->min_filter;
mag_filter = info->mag_filter;
mip_filter = info->mip_filter;
}
inline virtual T sample(glm::vec2 uv, float bias)
{
return sampleLod(uv, bias);
}
inline virtual T sampleLod(glm::vec2 uv, float lod)
{
if (mag_filter == SPIR2CROSS_FILTER_NEAREST)
{
uv.x = wrap(uv.x, wrap_s, mips[0].width);
uv.y = wrap(uv.y, wrap_t, mips[0].height);
glm::vec2 uv_full = uv * glm::vec2(mips[0].width, mips[0].height);
int x = int(uv_full.x);
int y = int(uv_full.y);
return sample(x, y, 0);
}
else
{
return T(0, 0, 0, 1);
}
}
inline float wrap(float v, spir2cross_wrap wrap, unsigned size)
{
switch (wrap)
{
case SPIR2CROSS_WRAP_REPEAT:
return v - glm::floor(v);
case SPIR2CROSS_WRAP_CLAMP_TO_EDGE:
{
float half = 0.5f / size;
return glm::clamp(v, half, 1.0f - half);
}
default:
return 0.0f;
}
}
std::vector<spir2cross_miplevel> mips;
spir2cross_format format;
spir2cross_wrap wrap_s;
spir2cross_format wrap_t;
spir2cross_filter min_filter;
spir2cross_filter mag_filter;
spir2cross_mipfilter mip_filter;
};
typedef sampler2DBase<glm::vec4> sampler2D;
typedef sampler2DBase<glm::ivec4> isampler2D;
typedef sampler2DBase<glm::uvec4> usampler2D;
template<typename T>
inline T texture(const sampler2DBase<T> &samp, const glm::vec2 &uv, float bias = 0.0f) { return samp.sample(uv, bias); }
}
#endif

113
include/spir2cross/thread_group.hpp Normal file
View File

@ -0,0 +1,113 @@
/*
* Copyright 2015-2016 ARM Limited
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
* You may obtain a copy of the License at
*
* http://www.apache.org/licenses/LICENSE-2.0
*
* Unless required by applicable law or agreed to in writing, software
* distributed under the License is distributed on an "AS IS" BASIS,
* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
* See the License for the specific language governing permissions and
* limitations under the License.
*/
#ifndef SPIR2CROSS_THREAD_GROUP_HPP
#define SPIR2CROSS_THREAD_GROUP_HPP
#include <thread>
#include <condition_variable>
#include <mutex>
namespace spir2cross
{
template<typename T, unsigned Size>
class ThreadGroup
{
public:
ThreadGroup(T *impl)
{
for (unsigned i = 0; i < Size; i++)
workers[i].start(&impl[i]);
}
void run()
{
for (auto &worker : workers)
worker.run();
}
void wait()
{
for (auto &worker : workers)
worker.wait();
}
private:
struct Thread
{
enum State
{
Idle,
Running,
Dying
};
State state = Idle;
void start(T *impl)
{
worker = std::thread([impl, this] {
for (;;)
{
{
std::unique_lock<std::mutex> l{lock};
cond.wait(l, [this] { return state != Idle; });
if (state == Dying)
break;
}
impl->main();
std::lock_guard<std::mutex> l{lock};
state = Idle;
cond.notify_one();
}
});
}
void wait()
{
std::unique_lock<std::mutex> l{lock};
cond.wait(l, [this] { return state == Idle; });
}
void run()
{
std::lock_guard<std::mutex> l{lock};
state = Running;
cond.notify_one();
}
~Thread()
{
if (worker.joinable())
{
{
std::lock_guard<std::mutex> l{lock};
state = Dying;
cond.notify_one();
}
worker.join();
}
}
std::thread worker;
std::condition_variable cond;
std::mutex lock;
};
Thread workers[Size];
};
}
#endif

11
jni/Android.mk Normal file
View File

@ -0,0 +1,11 @@
LOCAL_PATH := $(call my-dir)
include $(CLEAR_VARS)
LOCAL_CFLAGS += -std=c++11 -Wall -Wextra
LOCAL_MODULE := spir2cross
LOCAL_SRC_FILES := ../spir2cross.cpp ../spir2glsl.cpp ../spir2cpp.cpp
LOCAL_CPP_FEATURES := exceptions
LOCAL_ARM_MODE := arm
include $(BUILD_STATIC_LIBRARY)

2
jni/Application.mk Normal file
View File

@ -0,0 +1,2 @@
APP_STL := c++_static
APP_ABI := armeabi-v7a

425
main.cpp Normal file
View File

@ -0,0 +1,425 @@
/*
* Copyright 2015-2016 ARM Limited
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
* You may obtain a copy of the License at
*
* http://www.apache.org/licenses/LICENSE-2.0
*
* Unless required by applicable law or agreed to in writing, software
* distributed under the License is distributed on an "AS IS" BASIS,
* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
* See the License for the specific language governing permissions and
* limitations under the License.
*/
#include "spir2cpp.hpp"
#include <cstdio>
#include <stdexcept>
#include <functional>
#include <limits>
#include <memory>
#include <unordered_map>
#include <cstring>
#include <unordered_set>
using namespace spv;
using namespace spir2cross;
using namespace std;
struct CLIParser;
struct CLICallbacks
{
void add(const char *cli, const function<void (CLIParser&)> &func)
{
callbacks[cli] = func;
}
unordered_map<string, function<void (CLIParser&)>> callbacks;
function<void ()> error_handler;
function<void (const char*)> default_handler;
};
struct CLIParser
{
CLIParser(CLICallbacks cbs, int argc, char *argv[])
: cbs(move(cbs)), argc(argc), argv(argv)
{}
bool parse()
{
try
{
while (argc && !ended_state)
{
const char *next = *argv++;
argc--;
if (*next != '-' && cbs.default_handler)
{
cbs.default_handler(next);
}
else
{
auto itr = cbs.callbacks.find(next);
if (itr == ::end(cbs.callbacks))
{
throw logic_error("Invalid argument.\n");
}
itr->second(*this);
}
}
return true;
}
catch (...)
{
if (cbs.error_handler)
{
cbs.error_handler();
}
return false;
}
}
void end()
{
ended_state = true;
}
uint32_t next_uint()
{
if (!argc)
{
throw logic_error("Tried to parse uint, but nothing left in arguments.\n");
}
uint32_t val = stoul(*argv);
if (val > numeric_limits<uint32_t>::max())
{
throw out_of_range("next_uint() out of range.\n");
}
argc--;
argv++;
return val;
}
double next_double()
{
if (!argc)
{
throw logic_error("Tried to parse double, but nothing left in arguments.\n");
}
double val = stod(*argv);
argc--;
argv++;
return val;
}
const char *next_string()
{
if (!argc)
{
throw logic_error("Tried to parse string, but nothing left in arguments.\n");
}
const char *ret = *argv;
argc--;
argv++;
return ret;
}
CLICallbacks cbs;
int argc;
char **argv;
bool ended_state = false;
};
static vector<uint32_t> read_spirv_file(const char *path)
{
FILE *file = fopen(path, "rb");
if (!file)
{
fprintf(stderr, "Failed to open SPIRV file: %s\n", path);
return {};
}
fseek(file, 0, SEEK_END);
long len = ftell(file) / sizeof(uint32_t);
rewind(file);
vector<uint32_t> spirv(len);
if (fread(spirv.data(), sizeof(uint32_t), len, file) != size_t(len))
spirv.clear();
fclose(file);
return spirv;
}
static bool write_string_to_file(const char *path, const char *string)
{
FILE *file = fopen(path, "w");
if (!file)
{
fprintf(file, "Failed to write file: %s\n", path);
return false;
}
fprintf(file, "%s", string);
fclose(file);
return true;
}
static void print_resources(const Compiler &compiler, const char *tag, const vector<Resource> &resources)
{
fprintf(stderr, "%s\n", tag);
fprintf(stderr, "=============\n\n");
for (auto &res : resources)
{
auto &type = compiler.get_type(res.type_id);
auto mask = compiler.get_decoration_mask(res.id);
// If we don't have a name, use the fallback for the type instead of the variable
// for SSBOs and UBOs since those are the only meaningful names to use externally.
// Push constant blocks are still accessed by name and not block name, even though they are technically Blocks.
bool is_push_constant = compiler.get_storage_class(res.id) == StorageClassPushConstant;
bool is_block = (compiler.get_decoration_mask(type.self) &
((1ull << DecorationBlock) | (1ull << DecorationBufferBlock))) != 0;
uint32_t fallback_id = !is_push_constant && is_block ? res.type_id : res.id;
fprintf(stderr, " ID %03u : %s", res.id,
!res.name.empty() ? res.name.c_str() : compiler.get_fallback_name(fallback_id).c_str());
if (mask & (1ull << DecorationLocation))
fprintf(stderr, " (Location : %u)", compiler.get_decoration(res.id, DecorationLocation));
if (mask & (1ull << DecorationDescriptorSet))
fprintf(stderr, " (Set : %u)", compiler.get_decoration(res.id, DecorationDescriptorSet));
if (mask & (1ull << DecorationBinding))
fprintf(stderr, " (Binding : %u)", compiler.get_decoration(res.id, DecorationBinding));
fprintf(stderr, "\n");
}
fprintf(stderr, "=============\n\n");
}
static void print_resources(const Compiler &compiler, const ShaderResources &res)
{
print_resources(compiler, "subpass inputs", res.subpass_inputs);
print_resources(compiler, "inputs", res.stage_inputs);
print_resources(compiler, "outputs", res.stage_outputs);
print_resources(compiler, "textures", res.sampled_images);
print_resources(compiler, "images", res.storage_images);
print_resources(compiler, "ssbos", res.storage_buffers);
print_resources(compiler, "ubos", res.uniform_buffers);
print_resources(compiler, "push", res.push_constant_buffers);
print_resources(compiler, "counters", res.atomic_counters);
}
static void print_push_constant_resources(const Compiler &compiler, const vector<Resource> &res)
{
for (auto &block : res)
{
auto ranges = compiler.get_active_buffer_ranges(block.id);
fprintf(stderr, "Active members in buffer: %s\n",
!block.name.empty() ? block.name.c_str() : compiler.get_fallback_name(block.id).c_str());
fprintf(stderr, "==================\n\n");
for (auto &range : ranges)
{
const auto &name = compiler.get_member_name(block.type_id, range.index);
fprintf(stderr, "Member #%3u (%s): Offset: %4u, Range: %4u\n",
range.index, !name.empty() ? name.c_str() : compiler.get_fallback_member_name(range.index).c_str(),
unsigned(range.offset), unsigned(range.range));
}
fprintf(stderr, "==================\n\n");
}
}
struct PLSArg
{
PlsFormat format;
string name;
};
struct CLIArguments
{
const char *input = nullptr;
const char *output = nullptr;
uint32_t version = 0;
bool es = false;
bool set_version = false;
bool set_es = false;
bool dump_resources = false;
bool force_temporary = false;
bool flatten_ubo = false;
bool fixup = false;
vector<PLSArg> pls_in;
vector<PLSArg> pls_out;
uint32_t iterations = 1;
bool cpp = false;
};
static void print_help()
{
fprintf(stderr, "Usage: spir2cross [--output <output path>] [SPIR-V file] [--es] [--no-es] [--version <GLSL version>] [--dump-resources] [--help] [--force-temporary] [-cpp] [--flatten-ubo] [--fixup-clipspace] [--iterations iter] [--pls-in format input-name] [--pls-out format output-name]\n");
}
static vector<PlsRemap> remap_pls(const vector<PLSArg> &pls_variables, const vector<Resource> &resources, const vector<Resource> *secondary_resources)
{
vector<PlsRemap> ret;
for (auto &pls : pls_variables)
{
bool found = false;
for (auto &res : resources)
{
if (res.name == pls.name)
{
ret.push_back({ res.id, pls.format });
found = true;
break;
}
}
if (!found && secondary_resources)
{
for (auto &res : *secondary_resources)
{
if (res.name == pls.name)
{
ret.push_back({ res.id, pls.format });
found = true;
break;
}
}
}
if (!found)
fprintf(stderr, "Did not find stage input/output/target with name \"%s\".\n",
pls.name.c_str());
}
return ret;
}
static PlsFormat pls_format(const char *str)
{
if (!strcmp(str, "r11f_g11f_b10f")) return PlsR11FG11FB10F;
else if (!strcmp(str, "r32f")) return PlsR32F;
else if (!strcmp(str, "rg16f")) return PlsRG16F;
else if (!strcmp(str, "rg16")) return PlsRG16;
else if (!strcmp(str, "rgb10_a2")) return PlsRGB10A2;
else if (!strcmp(str, "rgba8")) return PlsRGBA8;
else if (!strcmp(str, "rgba8i")) return PlsRGBA8I;
else if (!strcmp(str, "rgba8ui")) return PlsRGBA8UI;
else if (!strcmp(str, "rg16i")) return PlsRG16I;
else if (!strcmp(str, "rgb10_a2ui")) return PlsRGB10A2UI;
else if (!strcmp(str, "rg16ui")) return PlsRG16UI;
else if (!strcmp(str, "r32ui")) return PlsR32UI;
else return PlsNone;
}
int main(int argc, char *argv[])
{
CLIArguments args;
CLICallbacks cbs;
cbs.add("--help", [](CLIParser &parser) { print_help(); parser.end(); });
cbs.add("--output", [&args](CLIParser &parser) { args.output = parser.next_string(); });
cbs.add("--es", [&args](CLIParser &) { args.es = true; args.set_es = true; });
cbs.add("--no-es", [&args](CLIParser &) { args.es = false; args.set_es = true; });
cbs.add("--version", [&args](CLIParser &parser) { args.version = parser.next_uint(); args.set_version = true; });
cbs.add("--dump-resources", [&args](CLIParser &) { args.dump_resources = true; });
cbs.add("--force-temporary", [&args](CLIParser &) { args.force_temporary = true; });
cbs.add("--flatten-ubo", [&args](CLIParser &) { args.flatten_ubo = true; });
cbs.add("--fixup-clipspace", [&args](CLIParser &) { args.fixup = true; });
cbs.add("--iterations", [&args](CLIParser &parser) { args.iterations = parser.next_uint(); });
cbs.add("--cpp", [&args](CLIParser &) { args.cpp = true; });
cbs.add("--pls-in", [&args](CLIParser &parser) {
auto fmt = pls_format(parser.next_string());
auto name = parser.next_string();
args.pls_in.push_back({ move(fmt), move(name) });
});
cbs.add("--pls-out", [&args](CLIParser &parser) {
auto fmt = pls_format(parser.next_string());
auto name = parser.next_string();
args.pls_out.push_back({ move(fmt), move(name) });
});
cbs.default_handler = [&args](const char *value) { args.input = value; };
cbs.error_handler = []{ print_help(); };
CLIParser parser{move(cbs), argc - 1, argv + 1};
if (!parser.parse())
{
return EXIT_FAILURE;
}
else if (parser.ended_state)
{
return EXIT_SUCCESS;
}
if (!args.input)
{
fprintf(stderr, "Didn't specify input file.\n");
print_help();
return EXIT_FAILURE;
}
unique_ptr<CompilerGLSL> compiler;
if (args.cpp)
compiler = unique_ptr<CompilerGLSL>(new CompilerCPP(read_spirv_file(args.input)));
else
compiler = unique_ptr<CompilerGLSL>(new CompilerGLSL(read_spirv_file(args.input)));
if (!args.set_version && !compiler->get_options().version)
{
fprintf(stderr, "Didn't specify GLSL version and SPIR-V did not specify language.\n");
print_help();
return EXIT_FAILURE;
}
CompilerGLSL::Options opts = compiler->get_options();
if (args.set_version)
opts.version = args.version;
if (args.set_es)
opts.es = args.es;
opts.force_temporary = args.force_temporary;
opts.vertex.fixup_clipspace = args.fixup;
compiler->set_options(opts);
auto res = compiler->get_shader_resources();
if (args.flatten_ubo)
for (auto &ubo : res.uniform_buffers)
compiler->flatten_interface_block(ubo.id);
auto pls_inputs = remap_pls(args.pls_in, res.stage_inputs, &res.subpass_inputs);
auto pls_outputs = remap_pls(args.pls_out, res.stage_outputs, nullptr);
compiler->remap_pixel_local_storage(move(pls_inputs), move(pls_outputs));
if (args.dump_resources)
{
print_resources(*compiler, res);
print_push_constant_resources(*compiler, res.push_constant_buffers);
}
string glsl;
for (uint32_t i = 0; i < args.iterations; i++)
glsl = compiler->compile();
if (args.output)
write_string_to_file(args.output, glsl.c_str());
else
printf("%s", glsl.c_str());
}

28
msvc/spir2cross.sln Normal file
View File

@ -0,0 +1,28 @@

Microsoft Visual Studio Solution File, Format Version 12.00
# Visual Studio Express 2013 for Windows Desktop
VisualStudioVersion = 12.0.31101.0
MinimumVisualStudioVersion = 10.0.40219.1
Project("{8BC9CEB8-8B4A-11D0-8D11-00A0C91BC942}") = "spir2cross", "spir2cross.vcxproj", "{977E3701-1A21-4425-B7E5-6BDF5EA062CD}"
EndProject
Global
GlobalSection(SolutionConfigurationPlatforms) = preSolution
Debug|Win32 = Debug|Win32
Debug|x64 = Debug|x64
Release|Win32 = Release|Win32
Release|x64 = Release|x64
EndGlobalSection
GlobalSection(ProjectConfigurationPlatforms) = postSolution
{977E3701-1A21-4425-B7E5-6BDF5EA062CD}.Debug|Win32.ActiveCfg = Debug|Win32
{977E3701-1A21-4425-B7E5-6BDF5EA062CD}.Debug|Win32.Build.0 = Debug|Win32
{977E3701-1A21-4425-B7E5-6BDF5EA062CD}.Debug|x64.ActiveCfg = Debug|x64
{977E3701-1A21-4425-B7E5-6BDF5EA062CD}.Debug|x64.Build.0 = Debug|x64
{977E3701-1A21-4425-B7E5-6BDF5EA062CD}.Release|Win32.ActiveCfg = Release|Win32
{977E3701-1A21-4425-B7E5-6BDF5EA062CD}.Release|Win32.Build.0 = Release|Win32
{977E3701-1A21-4425-B7E5-6BDF5EA062CD}.Release|x64.ActiveCfg = Release|x64
{977E3701-1A21-4425-B7E5-6BDF5EA062CD}.Release|x64.Build.0 = Release|x64
EndGlobalSection
GlobalSection(SolutionProperties) = preSolution
HideSolutionNode = FALSE
EndGlobalSection
EndGlobal

160
msvc/spir2cross.vcxproj Normal file
View File

@ -0,0 +1,160 @@
<?xml version="1.0" encoding="utf-8"?>
<Project DefaultTargets="Build" ToolsVersion="12.0" xmlns="http://schemas.microsoft.com/developer/msbuild/2003">
<ItemGroup Label="ProjectConfigurations">
<ProjectConfiguration Include="Debug|Win32">
<Configuration>Debug</Configuration>
<Platform>Win32</Platform>
</ProjectConfiguration>
<ProjectConfiguration Include="Debug|x64">
<Configuration>Debug</Configuration>
<Platform>x64</Platform>
</ProjectConfiguration>
<ProjectConfiguration Include="Release|Win32">
<Configuration>Release</Configuration>
<Platform>Win32</Platform>
</ProjectConfiguration>
<ProjectConfiguration Include="Release|x64">
<Configuration>Release</Configuration>
<Platform>x64</Platform>
</ProjectConfiguration>
</ItemGroup>
<PropertyGroup Label="Globals">
<ProjectGuid>{977E3701-1A21-4425-B7E5-6BDF5EA062CD}</ProjectGuid>
<RootNamespace>spir2cross</RootNamespace>
</PropertyGroup>
<Import Project="$(VCTargetsPath)\Microsoft.Cpp.Default.props" />
<PropertyGroup Condition="'$(Configuration)|$(Platform)'=='Debug|Win32'" Label="Configuration">
<ConfigurationType>Application</ConfigurationType>
<UseDebugLibraries>true</UseDebugLibraries>
<PlatformToolset>v120</PlatformToolset>
<CharacterSet>MultiByte</CharacterSet>
</PropertyGroup>
<PropertyGroup Condition="'$(Configuration)|$(Platform)'=='Debug|x64'" Label="Configuration">
<ConfigurationType>Application</ConfigurationType>
<UseDebugLibraries>true</UseDebugLibraries>
<PlatformToolset>v120</PlatformToolset>
<CharacterSet>MultiByte</CharacterSet>
</PropertyGroup>
<PropertyGroup Condition="'$(Configuration)|$(Platform)'=='Release|Win32'" Label="Configuration">
<ConfigurationType>Application</ConfigurationType>
<UseDebugLibraries>false</UseDebugLibraries>
<PlatformToolset>v120</PlatformToolset>
<WholeProgramOptimization>true</WholeProgramOptimization>
<CharacterSet>MultiByte</CharacterSet>
</PropertyGroup>
<PropertyGroup Condition="'$(Configuration)|$(Platform)'=='Release|x64'" Label="Configuration">
<ConfigurationType>Application</ConfigurationType>
<UseDebugLibraries>false</UseDebugLibraries>
<PlatformToolset>v120</PlatformToolset>
<WholeProgramOptimization>true</WholeProgramOptimization>
<CharacterSet>MultiByte</CharacterSet>
</PropertyGroup>
<Import Project="$(VCTargetsPath)\Microsoft.Cpp.props" />
<ImportGroup Label="ExtensionSettings">
</ImportGroup>
<ImportGroup Label="PropertySheets" Condition="'$(Configuration)|$(Platform)'=='Debug|Win32'">
<Import Project="$(UserRootDir)\Microsoft.Cpp.$(Platform).user.props" Condition="exists('$(UserRootDir)\Microsoft.Cpp.$(Platform).user.props')" Label="LocalAppDataPlatform" />
</ImportGroup>
<ImportGroup Condition="'$(Configuration)|$(Platform)'=='Debug|x64'" Label="PropertySheets">
<Import Project="$(UserRootDir)\Microsoft.Cpp.$(Platform).user.props" Condition="exists('$(UserRootDir)\Microsoft.Cpp.$(Platform).user.props')" Label="LocalAppDataPlatform" />
</ImportGroup>
<ImportGroup Label="PropertySheets" Condition="'$(Configuration)|$(Platform)'=='Release|Win32'">
<Import Project="$(UserRootDir)\Microsoft.Cpp.$(Platform).user.props" Condition="exists('$(UserRootDir)\Microsoft.Cpp.$(Platform).user.props')" Label="LocalAppDataPlatform" />
</ImportGroup>
<ImportGroup Condition="'$(Configuration)|$(Platform)'=='Release|x64'" Label="PropertySheets">
<Import Project="$(UserRootDir)\Microsoft.Cpp.$(Platform).user.props" Condition="exists('$(UserRootDir)\Microsoft.Cpp.$(Platform).user.props')" Label="LocalAppDataPlatform" />
</ImportGroup>
<PropertyGroup Label="UserMacros" />
<PropertyGroup />
<ItemDefinitionGroup Condition="'$(Configuration)|$(Platform)'=='Debug|Win32'">
<ClCompile>
<WarningLevel>Level3</WarningLevel>
<Optimization>Disabled</Optimization>
<SDLCheck>true</SDLCheck>
<PreprocessorDefinitions>_CRT_SECURE_NO_WARNINGS;%(PreprocessorDefinitions)</PreprocessorDefinitions>
<RuntimeLibrary>MultiThreadedDebugDLL</RuntimeLibrary>
</ClCompile>
<Link>
<GenerateDebugInformation>true</GenerateDebugInformation>
</Link>
</ItemDefinitionGroup>
<ItemDefinitionGroup Condition="'$(Configuration)|$(Platform)'=='Debug|x64'">
<ClCompile>
<WarningLevel>Level3</WarningLevel>
<Optimization>Disabled</Optimization>
<SDLCheck>true</SDLCheck>
<PreprocessorDefinitions>_CRT_SECURE_NO_WARNINGS;%(PreprocessorDefinitions)</PreprocessorDefinitions>
<RuntimeLibrary>MultiThreadedDebugDLL</RuntimeLibrary>
</ClCompile>
<Link>
<GenerateDebugInformation>true</GenerateDebugInformation>
</Link>
</ItemDefinitionGroup>
<ItemDefinitionGroup Condition="'$(Configuration)|$(Platform)'=='Release|Win32'">
<ClCompile>
<WarningLevel>Level3</WarningLevel>
<Optimization>MaxSpeed</Optimization>
<FunctionLevelLinking>true</FunctionLevelLinking>
<IntrinsicFunctions>true</IntrinsicFunctions>
<SDLCheck>true</SDLCheck>
<PreprocessorDefinitions>_CRT_SECURE_NO_WARNINGS;%(PreprocessorDefinitions)</PreprocessorDefinitions>
<RuntimeLibrary>MultiThreadedDLL</RuntimeLibrary>
</ClCompile>
<Link>
<GenerateDebugInformation>true</GenerateDebugInformation>
<EnableCOMDATFolding>true</EnableCOMDATFolding>
<OptimizeReferences>true</OptimizeReferences>
</Link>
</ItemDefinitionGroup>
<ItemDefinitionGroup Condition="'$(Configuration)|$(Platform)'=='Release|x64'">
<ClCompile>
<WarningLevel>Level3</WarningLevel>
<Optimization>MaxSpeed</Optimization>
<FunctionLevelLinking>true</FunctionLevelLinking>
<IntrinsicFunctions>true</IntrinsicFunctions>
<SDLCheck>true</SDLCheck>
<PreprocessorDefinitions>_CRT_SECURE_NO_WARNINGS;%(PreprocessorDefinitions)</PreprocessorDefinitions>
<RuntimeLibrary>MultiThreadedDLL</RuntimeLibrary>
</ClCompile>
<Link>
<GenerateDebugInformation>true</GenerateDebugInformation>
<EnableCOMDATFolding>true</EnableCOMDATFolding>
<OptimizeReferences>true</OptimizeReferences>
</Link>
</ItemDefinitionGroup>
<ItemGroup>
<ClCompile Include="..\cpp\compute.cpp">
<ExcludedFromBuild Condition="'$(Configuration)|$(Platform)'=='Release|x64'">true</ExcludedFromBuild>
<ExcludedFromBuild Condition="'$(Configuration)|$(Platform)'=='Debug|Win32'">true</ExcludedFromBuild>
<ExcludedFromBuild Condition="'$(Configuration)|$(Platform)'=='Release|Win32'">true</ExcludedFromBuild>
<ExcludedFromBuild Condition="'$(Configuration)|$(Platform)'=='Debug|x64'">true</ExcludedFromBuild>
</ClCompile>
<ClCompile Include="..\cpp\culling.cpp">
<ExcludedFromBuild Condition="'$(Configuration)|$(Platform)'=='Release|x64'">true</ExcludedFromBuild>
<ExcludedFromBuild Condition="'$(Configuration)|$(Platform)'=='Debug|Win32'">true</ExcludedFromBuild>
<ExcludedFromBuild Condition="'$(Configuration)|$(Platform)'=='Release|Win32'">true</ExcludedFromBuild>
<ExcludedFromBuild Condition="'$(Configuration)|$(Platform)'=='Debug|x64'">true</ExcludedFromBuild>
</ClCompile>
<ClCompile Include="..\cpp\shadertoy.cpp">
<ExcludedFromBuild Condition="'$(Configuration)|$(Platform)'=='Release|x64'">true</ExcludedFromBuild>
<ExcludedFromBuild Condition="'$(Configuration)|$(Platform)'=='Debug|Win32'">true</ExcludedFromBuild>
<ExcludedFromBuild Condition="'$(Configuration)|$(Platform)'=='Release|Win32'">true</ExcludedFromBuild>
<ExcludedFromBuild Condition="'$(Configuration)|$(Platform)'=='Debug|x64'">true</ExcludedFromBuild>
</ClCompile>
<ClCompile Include="..\main.cpp" />
<ClCompile Include="..\spir2cpp.cpp" />
<ClCompile Include="..\spir2cross.cpp" />
<ClCompile Include="..\spir2glsl.cpp" />
</ItemGroup>
<ItemGroup>
<ClInclude Include="..\GLSL.std.450.h" />
<ClInclude Include="..\spir2common.hpp" />
<ClInclude Include="..\spir2cpp.hpp" />
<ClInclude Include="..\spir2cross.hpp" />
<ClInclude Include="..\spir2glsl.hpp" />
<ClInclude Include="..\spirv.hpp" />
</ItemGroup>
<Import Project="$(VCTargetsPath)\Microsoft.Cpp.targets" />
<ImportGroup Label="ExtensionTargets">
</ImportGroup>
</Project>

63
msvc/spir2cross.vcxproj.filters Normal file
View File

@ -0,0 +1,63 @@
<?xml version="1.0" encoding="utf-8"?>
<Project ToolsVersion="4.0" xmlns="http://schemas.microsoft.com/developer/msbuild/2003">
<ItemGroup>
<Filter Include="Source Files">
<UniqueIdentifier>{4FC737F1-C7A5-4376-A066-2A32D752A2FF}</UniqueIdentifier>
<Extensions>cpp;c;cc;cxx;def;odl;idl;hpj;bat;asm;asmx</Extensions>
</Filter>
<Filter Include="Header Files">
<UniqueIdentifier>{93995380-89BD-4b04-88EB-625FBE52EBFB}</UniqueIdentifier>
<Extensions>h;hh;hpp;hxx;hm;inl;inc;xsd</Extensions>
</Filter>
<Filter Include="Resource Files">
<UniqueIdentifier>{67DA6AB6-F800-4c08-8B7A-83BB121AAD01}</UniqueIdentifier>
<Extensions>rc;ico;cur;bmp;dlg;rc2;rct;bin;rgs;gif;jpg;jpeg;jpe;resx;tiff;tif;png;wav;mfcribbon-ms</Extensions>
</Filter>
<Filter Include="Source Files\cpp">
<UniqueIdentifier>{61390b44-2b95-4b9a-8910-9ea1bc3a4920}</UniqueIdentifier>
</Filter>
</ItemGroup>
<ItemGroup>
<ClCompile Include="..\main.cpp">
<Filter>Source Files</Filter>
</ClCompile>
<ClCompile Include="..\spir2cross.cpp">
<Filter>Source Files</Filter>
</ClCompile>
<ClCompile Include="..\spir2glsl.cpp">
<Filter>Source Files</Filter>
</ClCompile>
<ClCompile Include="..\spir2cpp.cpp">
<Filter>Source Files</Filter>
</ClCompile>
<ClCompile Include="..\cpp\compute.cpp">
<Filter>Source Files\cpp</Filter>
</ClCompile>
<ClCompile Include="..\cpp\culling.cpp">
<Filter>Source Files\cpp</Filter>
</ClCompile>
<ClCompile Include="..\cpp\shadertoy.cpp">
<Filter>Source Files\cpp</Filter>
</ClCompile>
</ItemGroup>
<ItemGroup>
<ClInclude Include="..\GLSL.std.450.h">
<Filter>Header Files</Filter>
</ClInclude>
<ClInclude Include="..\spir2cross.hpp">
<Filter>Header Files</Filter>
</ClInclude>
<ClInclude Include="..\spir2glsl.hpp">
<Filter>Header Files</Filter>
</ClInclude>
<ClInclude Include="..\spirv.hpp">
<Filter>Header Files</Filter>
</ClInclude>
<ClInclude Include="..\spir2common.hpp">
<Filter>Header Files</Filter>
</ClInclude>
<ClInclude Include="..\spir2cpp.hpp">
<Filter>Header Files</Filter>
</ClInclude>
</ItemGroup>
</Project>

47
reference/shaders/comp/atomic.comp Normal file
View File

@ -0,0 +1,47 @@
#version 310 es
#extension GL_OES_shader_image_atomic : require
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
layout(binding = 2, std430) buffer SSBO
{
uint u32;
int i32;
} ssbo;
layout(binding = 0, r32ui) uniform highp uimage2D uImage;
layout(binding = 1, r32i) uniform highp iimage2D iImage;
void main()
{
uint _19 = imageAtomicAdd(uImage, ivec2(1, 5), 1u);
uint _21 = imageAtomicOr(uImage, ivec2(1, 5), 1u);
uint _23 = imageAtomicXor(uImage, ivec2(1, 5), 1u);
uint _25 = imageAtomicAnd(uImage, ivec2(1, 5), 1u);
uint _27 = imageAtomicMin(uImage, ivec2(1, 5), 1u);
uint _29 = imageAtomicMax(uImage, ivec2(1, 5), 1u);
uint _33 = imageAtomicCompSwap(uImage, ivec2(1, 5), 10u, 2u);
int _41 = imageAtomicAdd(iImage, ivec2(1, 6), 1);
int _43 = imageAtomicOr(iImage, ivec2(1, 6), 1);
int _45 = imageAtomicXor(iImage, ivec2(1, 6), 1);
int _47 = imageAtomicAnd(iImage, ivec2(1, 6), 1);
int _49 = imageAtomicMin(iImage, ivec2(1, 6), 1);
int _51 = imageAtomicMax(iImage, ivec2(1, 6), 1);
int _55 = imageAtomicCompSwap(iImage, ivec2(1, 5), 10, 2);
uint _62 = atomicAdd(ssbo.u32, 1u);
uint _64 = atomicOr(ssbo.u32, 1u);
uint _66 = atomicXor(ssbo.u32, 1u);
uint _68 = atomicAnd(ssbo.u32, 1u);
uint _70 = atomicMin(ssbo.u32, 1u);
uint _72 = atomicMax(ssbo.u32, 1u);
uint _74 = atomicExchange(ssbo.u32, 1u);
uint _76 = atomicCompSwap(ssbo.u32, 10u, 2u);
int _79 = atomicAdd(ssbo.i32, 1);
int _81 = atomicOr(ssbo.i32, 1);
int _83 = atomicXor(ssbo.i32, 1);
int _85 = atomicAnd(ssbo.i32, 1);
int _87 = atomicMin(ssbo.i32, 1);
int _89 = atomicMax(ssbo.i32, 1);
int _91 = atomicExchange(ssbo.i32, 1);
int _93 = atomicCompSwap(ssbo.i32, 10, 2);
}

39
reference/shaders/comp/bake_gradient.comp Normal file
View File

@ -0,0 +1,39 @@
#version 310 es
layout(local_size_x = 8, local_size_y = 8, local_size_z = 1) in;
layout(binding = 4, std140) uniform UBO
{
vec4 uInvSize;
vec4 uScale;
} _46;
layout(binding = 0) uniform mediump sampler2D uHeight;
layout(binding = 1) uniform mediump sampler2D uDisplacement;
layout(binding = 2, rgba16f) uniform mediump writeonly image2D iHeightDisplacement;
layout(binding = 3, rgba16f) uniform mediump writeonly image2D iGradJacobian;
mediump float jacobian(mediump vec2 dDdx, mediump vec2 dDdy)
{
return (((1.000000 + dDdx.x) * (1.000000 + dDdy.y)) - (dDdx.y * dDdy.x));
}
void main()
{
vec4 uv = ((vec2(gl_GlobalInvocationID.xy) * _46.uInvSize.xy).xyxy + (_46.uInvSize * 0.500000));
float h = textureLod(uHeight, uv.xy, 0.000000).x;
float x0 = textureLodOffset(uHeight, uv.xy, 0.000000, ivec2(-1, 0)).x;
float x1 = textureLodOffset(uHeight, uv.xy, 0.000000, ivec2(1, 0)).x;
float y0 = textureLodOffset(uHeight, uv.xy, 0.000000, ivec2(0, -1)).x;
float y1 = textureLodOffset(uHeight, uv.xy, 0.000000, ivec2(0, 1)).x;
vec2 grad = ((_46.uScale.xy * 0.500000) * vec2((x1 - x0), (y1 - y0)));
vec2 displacement = (textureLod(uDisplacement, uv.zw, 0.000000).xy * 1.200000);
vec2 dDdx = ((textureLodOffset(uDisplacement, uv.zw, 0.000000, ivec2(1, 0)).xy - textureLodOffset(uDisplacement, uv.zw, 0.000000, ivec2(-1, 0)).xy) * 0.600000);
vec2 dDdy = ((textureLodOffset(uDisplacement, uv.zw, 0.000000, ivec2(0, 1)).xy - textureLodOffset(uDisplacement, uv.zw, 0.000000, ivec2(0, -1)).xy) * 0.600000);
vec2 param = (dDdx * _46.uScale.z);
vec2 param_1 = (dDdy * _46.uScale.z);
float j = jacobian(param, param_1);
displacement = vec2(0.000000);
imageStore(iHeightDisplacement, ivec2(gl_GlobalInvocationID.xy), vec4(h, displacement, 0.000000));
imageStore(iGradJacobian, ivec2(gl_GlobalInvocationID.xy), vec4(grad, j, 0.000000));
}

29
reference/shaders/comp/basic.comp Normal file
View File

@ -0,0 +1,29 @@
#version 310 es
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
layout(binding = 0, std430) buffer SSBO
{
vec4 in_data[];
} _23;
layout(binding = 1, std430) buffer SSBO2
{
vec4 out_data[];
} _45;
layout(binding = 2, std430) buffer SSBO3
{
uint counter;
} _48;
void main()
{
uint ident = gl_GlobalInvocationID.x;
vec4 idata = _23.in_data[ident];
if ((dot(idata, vec4(1.000000, 5.000000, 6.000000, 2.000000)) > 8.200000))
{
uint _52 = atomicAdd(_48.counter, 1u);
_45.out_data[_52] = idata;
}
}

29
reference/shaders/comp/culling.comp Normal file
View File

@ -0,0 +1,29 @@
#version 310 es
layout(local_size_x = 4, local_size_y = 1, local_size_z = 1) in;
layout(binding = 0, std430) buffer SSBO
{
float in_data[];
} _22;
layout(binding = 1, std430) buffer SSBO2
{
float out_data[];
} _38;
layout(binding = 2, std430) buffer SSBO3
{
uint count;
} _41;
void main()
{
uint ident = gl_GlobalInvocationID.x;
float idata = _22.in_data[ident];
if ((idata > 12.000000))
{
uint _45 = atomicAdd(_41.count, 1u);
_38.out_data[_45] = idata;
}
}

29
reference/shaders/comp/dowhile.comp Normal file
View File

@ -0,0 +1,29 @@
#version 310 es
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
layout(binding = 0, std430) buffer SSBO
{
mat4 mvp;
vec4 in_data[];
} _28;
layout(binding = 1, std430) buffer SSBO2
{
vec4 out_data[];
} _52;
int i;
void main()
{
uint ident = gl_GlobalInvocationID.x;
i = 0;
vec4 idat = _28.in_data[ident];
do
{
idat = (_28.mvp * idat);
i = (i + 1);
} while ((i < 16));
_52.out_data[ident] = idat;
}

96
reference/shaders/comp/generate_height.comp Normal file
View File

@ -0,0 +1,96 @@
#version 310 es
layout(local_size_x = 64, local_size_y = 1, local_size_z = 1) in;
layout(binding = 0, std430) buffer Distribution
{
vec2 distribution[];
} _190;
layout(binding = 2, std140) uniform UBO
{
vec4 uModTime;
} _218;
layout(binding = 1, std430) buffer HeightmapFFT
{
uint heights[];
} _276;
uvec2 workaround_mix(uvec2 a, uvec2 b, bvec2 sel)
{
uint _137;
uint _148;
if (sel.x)
{
_137 = b.x;
}
else
{
_137 = a.x;
}
uint _147 = _137;
if (sel.y)
{
_148 = b.y;
}
else
{
_148 = a.y;
}
return uvec2(_147, _148);
}
vec2 alias(vec2 i, vec2 N)
{
return mix(i, (i - N), greaterThan(i, (N * 0.500000)));
}
vec2 cmul(vec2 a, vec2 b)
{
vec2 r3 = a.yx;
vec2 r1 = b.xx;
vec2 R0 = (a * r1);
vec2 r2 = b.yy;
vec2 R1 = (r2 * r3);
return (R0 + vec2((-R1.x), R1.y));
}
uint pack2(vec2 v)
{
return packHalf2x16(v);
}
void generate_heightmap()
{
uvec2 N = (uvec2(64u, 1u) * gl_NumWorkGroups.xy);
uvec2 i = gl_GlobalInvocationID.xy;
uvec2 param = (N - i);
uvec2 param_1 = uvec2(0u);
bvec2 param_2 = equal(i, uvec2(0u));
uvec2 wi = workaround_mix(param, param_1, param_2);
vec2 a = _190.distribution[((i.y * N.x) + i.x)];
vec2 b = _190.distribution[((wi.y * N.x) + wi.x)];
vec2 param_3 = vec2(i);
vec2 param_4 = vec2(N);
vec2 k = (_218.uModTime.xy * alias(param_3, param_4));
float k_len = length(k);
float w = (sqrt((9.810000 * k_len)) * _218.uModTime.z);
float cw = cos(w);
float sw = sin(w);
vec2 param_5 = a;
vec2 param_6 = vec2(cw, sw);
a = cmul(param_5, param_6);
vec2 param_7 = b;
vec2 param_8 = vec2(cw, sw);
b = cmul(param_7, param_8);
b = vec2(b.x, (-b.y));
vec2 res = (a + b);
vec2 param_9 = res;
_276.heights[((i.y * N.x) + i.x)] = pack2(param_9);
}
void main()
{
generate_heightmap();
}

12
reference/shaders/comp/image.comp Normal file
View File

@ -0,0 +1,12 @@
#version 310 es
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
layout(binding = 0, rgba8) uniform mediump readonly image2D uImageIn;
layout(binding = 1, rgba8) uniform mediump writeonly image2D uImageOut;
void main()
{
vec4 v = imageLoad(uImageIn, (ivec2(gl_GlobalInvocationID.xy) + imageSize(uImageIn)));
imageStore(uImageOut, ivec2(gl_GlobalInvocationID.xy), v);
}

61
reference/shaders/comp/inout-struct.comp Normal file
View File

@ -0,0 +1,61 @@
#version 310 es
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
struct Foo
{
vec4 a;
vec4 b;
vec4 c;
vec4 d;
};
layout(binding = 1, std430) buffer SSBO2
{
vec4 data[];
} indata;
layout(binding = 0, std430) buffer SSBO
{
vec4 data[];
} outdata;
layout(binding = 2, std430) buffer SSBO3
{
Foo foos[];
} foobar;
void baz(out Foo foo)
{
uint ident = gl_GlobalInvocationID.x;
foo.a = indata.data[((4u * ident) + 0u)];
foo.b = indata.data[((4u * ident) + 1u)];
foo.c = indata.data[((4u * ident) + 2u)];
foo.d = indata.data[((4u * ident) + 3u)];
}
void meow(inout Foo foo)
{
foo.a = (foo.a + vec4(10.000000));
foo.b = (foo.b + vec4(20.000000));
foo.c = (foo.c + vec4(30.000000));
foo.d = (foo.d + vec4(40.000000));
}
vec4 bar(Foo foo)
{
return (((foo.a + foo.b) + foo.c) + foo.d);
}
void main()
{
Foo param;
baz(param);
Foo foo = param;
Foo param_1 = foo;
meow(param_1);
foo = param_1;
Foo param_2 = foo;
Foo param_3 = foobar.foos[gl_GlobalInvocationID.x];
outdata.data[gl_GlobalInvocationID.x] = (bar(param_2) + bar(param_3));
}

19
reference/shaders/comp/insert.comp Normal file
View File

@ -0,0 +1,19 @@
#version 310 es
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
layout(binding = 0, std430) buffer SSBO
{
vec4 out_data[];
} _27;
void main()
{
vec4 v;
v.x = 10.000000;
v.y = 30.000000;
v.z = 70.000000;
v.w = 90.000000;
_27.out_data[gl_GlobalInvocationID.x] = v;
_27.out_data[gl_GlobalInvocationID.x].y = 20.000000;
}

72
reference/shaders/comp/loop.comp Normal file
View File

@ -0,0 +1,72 @@
#version 310 es
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
layout(binding = 0, std430) buffer SSBO
{
mat4 mvp;
vec4 in_data[];
} _24;
layout(binding = 1, std430) buffer SSBO2
{
vec4 out_data[];
} _125;
void main()
{
uint ident = gl_GlobalInvocationID.x;
vec4 idat = _24.in_data[ident];
int k = 0;
uint i;
uint j;
int l;
while ((k < 10))
{
idat = (idat * 2.000000);
k = (k + 1);
}
i = 0u;
for (; (i < 16u); i = (i + uint(1)), k = (k + 1))
{
j = 0u;
for (; (j < 30u); j = (j + uint(1)))
{
idat = (_24.mvp * idat);
}
}
k = 0;
for (;;)
{
k = (k + 1);
if ((k > 10))
{
k = (k + 2);
}
else
{
k = (k + 3);
continue;
}
k = (k + 10);
continue;
}
k = 0;
do
{
k = (k + 1);
} while ((k > 10));
l = 0;
for (;;)
{
if ((l == 5))
{
l = (l + 1);
continue;
}
idat = (idat + vec4(1.000000));
l = (l + 1);
continue;
}
_125.out_data[ident] = idat;
}

14
reference/shaders/comp/mat3.comp Normal file
View File

@ -0,0 +1,14 @@
#version 310 es
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
layout(binding = 1, std430) buffer SSBO2
{
mat3 out_data[];
} _22;
void main()
{
uint ident = gl_GlobalInvocationID.x;
_22.out_data[ident] = mat3(vec3(10.000000), vec3(20.000000), vec3(40.000000));
}

22
reference/shaders/comp/modf.comp Normal file
View File

@ -0,0 +1,22 @@
#version 310 es
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
layout(binding = 0, std430) buffer SSBO
{
vec4 in_data[];
} _23;
layout(binding = 1, std430) buffer SSBO2
{
vec4 out_data[];
} _35;
void main()
{
uint ident = gl_GlobalInvocationID.x;
vec4 i;
vec4 _31 = modf(_23.in_data[ident], i);
vec4 v = _31;
_35.out_data[ident] = v;
}

36
reference/shaders/comp/return.comp Normal file
View File

@ -0,0 +1,36 @@
#version 310 es
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
layout(binding = 1, std430) buffer SSBO2
{
vec4 out_data[];
} _27;
void main()
{
uint ident = gl_GlobalInvocationID.x;
int i;
if ((ident == 2u))
{
_27.out_data[ident] = vec4(20.000000);
}
else
{
if ((ident == 4u))
{
_27.out_data[ident] = vec4(10.000000);
return;
}
}
i = 0;
for (; (i < 20); i = (i + 1))
{
if ((i == 10))
{
break;
}
return;
}
_27.out_data[ident] = vec4(10.000000);
}

25
reference/shaders/comp/shared.comp Normal file
View File

@ -0,0 +1,25 @@
#version 310 es
layout(local_size_x = 4, local_size_y = 1, local_size_z = 1) in;
layout(binding = 0, std430) buffer SSBO
{
float in_data[];
} _22;
layout(binding = 1, std430) buffer SSBO2
{
float out_data[];
} _44;
shared float sShared[4];
void main()
{
uint ident = gl_GlobalInvocationID.x;
float idata = _22.in_data[ident];
sShared[gl_LocalInvocationIndex] = idata;
memoryBarrierShared();
barrier();
_44.out_data[ident] = sShared[((4u - gl_LocalInvocationIndex) - 1u)];
}

24
reference/shaders/comp/struct-layout.comp Normal file
View File

@ -0,0 +1,24 @@
#version 310 es
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
struct Foo
{
mat4 m;
};
layout(binding = 1, std430) buffer SSBO2
{
Foo out_data[];
} _23;
layout(binding = 0, std430) buffer SSBO
{
Foo in_data[];
} _30;
void main()
{
uint ident = gl_GlobalInvocationID.x;
_23.out_data[ident].m = (_30.in_data[ident].m * _30.in_data[ident].m);
}

68
reference/shaders/comp/struct-packing.comp Normal file
View File

@ -0,0 +1,68 @@
#version 310 es
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
struct S0
{
vec2 a[1];
float b;
};
struct S1
{
vec3 a;
float b;
};
struct S2
{
vec3 a[1];
float b;
};
struct S3
{
vec2 a;
float b;
};
struct Content
{
S0 m0s[1];
S1 m1s[1];
S2 m2s[1];
S0 m0;
S1 m1;
S2 m2;
S3 m3;
float m4;
};
layout(binding = 1, std140) buffer SSBO1
{
Content content;
Content content1[2];
Content content2;
mat2 m0;
mat2 m1;
mat2x3 m2[4];
mat3x2 m3;
layout(row_major) mat2 m4;
layout(row_major) mat2 m5[9];
layout(row_major) mat2x3 m6[2][4];
layout(row_major) mat3x2 m7;
float array[];
} ssbo_430;
layout(binding = 0, std140) buffer SSBO0
{
Content content;
Content content1[2];
Content content2;
float array[];
} ssbo_140;
void main()
{
ssbo_430.content = ssbo_140.content;
}

66
reference/shaders/comp/torture-loop.comp Normal file
View File

@ -0,0 +1,66 @@
#version 310 es
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
layout(binding = 0, std430) buffer SSBO
{
mat4 mvp;
vec4 in_data[];
} _24;
layout(binding = 1, std430) buffer SSBO2
{
vec4 out_data[];
} _89;
void main()
{
uint ident = gl_GlobalInvocationID.x;
vec4 idat = _24.in_data[ident];
int k = 0;
uint i;
uint j;
for (;;)
{
int _39 = k;
int _40 = _39 + 1;
k = _40;
if ((_40 < 10))
{
idat = (idat * 2.000000);
int _47 = k;
k = (_47 + 1);
continue;
}
else
{
break;
}
}
i = 0u;
int _76;
for (; (i < 16u); i = (i + uint(1)), _76 = k, k = (_76 + 1))
{
j = 0u;
for (; (j < 30u); j = (j + uint(1)))
{
idat = (_24.mvp * idat);
}
}
int _84;
for (;;)
{
int _82 = k;
k = (_82 + 1);
int _84 = k;
if ((_84 > 10))
{
continue;
}
else
{
break;
}
}
_89.out_data[ident] = idat;
}

18
reference/shaders/comp/udiv.comp Normal file
View File

@ -0,0 +1,18 @@
#version 310 es
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
layout(binding = 0, std430) buffer SSBO2
{
uint outputs[];
} _10;
layout(binding = 0, std430) buffer SSBO
{
uint inputs[];
} _23;
void main()
{
_10.outputs[gl_GlobalInvocationID.x] = (_23.inputs[gl_GlobalInvocationID.x] / 29u);
}

15
reference/shaders/frag/basic.frag Normal file
View File

@ -0,0 +1,15 @@
#version 310 es
precision mediump float;
precision highp int;
layout(binding = 0) uniform mediump sampler2D uTex;
layout(location = 0) out vec4 FragColor;
in vec4 vColor;
in vec2 vTex;
void main()
{
FragColor = (vColor * texture(uTex, vTex));
}

28
reference/shaders/frag/constant-array.frag Normal file
View File

@ -0,0 +1,28 @@
#version 310 es
precision mediump float;
precision highp int;
struct Foobar
{
float a;
float b;
};
layout(location = 0) out vec4 FragColor;
layout(location = 0) in mediump flat int index;
vec4 resolve(Foobar f)
{
return vec4((f.a + f.b));
}
void main()
{
highp vec4 indexable[3] = vec4[](vec4(1.000000), vec4(2.000000), vec4(3.000000));
highp vec4 indexable_1[2][2] = vec4[][](vec4[](vec4(1.000000), vec4(2.000000)), vec4[](vec4(8.000000), vec4(10.000000)));
Foobar param = Foobar(10.000000, 20.000000);
Foobar indexable_2[2] = Foobar[](Foobar(10.000000, 40.000000), Foobar(90.000000, 70.000000));
Foobar param_1 = indexable_2[index];
FragColor = (((indexable[index] + indexable_1[index][(index + 1)]) + resolve(param)) + resolve(param_1));
}

30
reference/shaders/frag/flush_params.frag Normal file
View File

@ -0,0 +1,30 @@
#version 310 es
precision mediump float;
precision highp int;
struct Structy
{
vec4 c;
};
layout(location = 0) out vec4 FragColor;
void foo2(out Structy f)
{
f.c = vec4(10.000000);
}
Structy foo()
{
Structy param;
foo2(param);
Structy f = param;
return f;
}
void main()
{
Structy s = foo();
FragColor = s.c;
}

62
reference/shaders/frag/ground.frag Normal file
View File

@ -0,0 +1,62 @@
#version 310 es
precision mediump float;
precision highp int;
layout(binding = 4, std140) uniform GlobalPSData
{
vec4 g_CamPos;
vec4 g_SunDir;
vec4 g_SunColor;
vec4 g_ResolutionParams;
vec4 g_TimeParams;
vec4 g_FogColor_Distance;
} _56;
layout(binding = 2) uniform mediump sampler2D TexNormalmap;
layout(location = 3) out vec4 LightingOut;
layout(location = 2) out vec4 NormalOut;
layout(location = 1) out vec4 SpecularOut;
layout(location = 0) out vec4 AlbedoOut;
layout(location = 0) in vec2 TexCoord;
layout(location = 1) in vec3 EyeVec;
float saturate(float x)
{
return clamp(x, 0.000000, 1.000000);
}
void Resolve(vec3 Albedo, vec3 Normal, float Roughness, float Metallic)
{
LightingOut = vec4(0.000000);
NormalOut = vec4(((Normal * 0.500000) + vec3(0.500000)), 0.000000);
SpecularOut = vec4(Roughness, Metallic, 0.000000, 0.000000);
AlbedoOut = vec4(Albedo, 1.000000);
}
void main()
{
vec3 Normal = ((texture(TexNormalmap, TexCoord).xyz * 2.000000) - vec3(1.000000));
Normal = normalize(Normal);
highp float param = (length(EyeVec) / 1000.000000);
vec2 scatter_uv;
scatter_uv.x = saturate(param);
vec3 nEye = normalize(EyeVec);
scatter_uv.y = 0.000000;
vec3 Color = vec3(0.100000, 0.300000, 0.100000);
vec3 grass = vec3(0.100000, 0.300000, 0.100000);
vec3 dirt = vec3(0.100000);
vec3 snow = vec3(0.800000);
float grass_snow = smoothstep(0.000000, 0.150000, ((_56.g_CamPos.y + EyeVec.y) / 200.000000));
vec3 base = mix(grass, snow, vec3(grass_snow));
float edge = smoothstep(0.700000, 0.750000, Normal.y);
Color = mix(dirt, base, vec3(edge));
Color = (Color * Color);
float Roughness = (1.000000 - (edge * grass_snow));
highp vec3 param_1 = Color;
highp vec3 param_2 = Normal;
highp float param_3 = Roughness;
highp float param_4 = 0.000000;
Resolve(param_1, param_2, param_3, param_4);
}

38
reference/shaders/frag/mix.frag Normal file
View File

@ -0,0 +1,38 @@
#version 310 es
precision mediump float;
precision highp int;
layout(location = 0) out vec4 FragColor;
layout(location = 0) in vec4 vIn0;
layout(location = 1) in vec4 vIn1;
layout(location = 2) in float vIn2;
layout(location = 3) in float vIn3;
void main()
{
bvec4 l = bvec4(false, true, false, false);
FragColor = mix(vIn0, vIn1, l);
bool f = true;
FragColor = vec4(mix(vIn2, vIn3, f));
highp vec4 _35;
highp float _44;
if (f)
{
_35 = vIn0;
}
else
{
_35 = vIn1;
}
FragColor = _35;
if (f)
{
_44 = vIn2;
}
else
{
_44 = vIn3;
}
FragColor = vec4(_44);
}

21
reference/shaders/frag/pls.frag Normal file
View File

@ -0,0 +1,21 @@
#version 310 es
precision mediump float;
precision highp int;
layout(location = 0) out vec4 PLSOut0;
layout(location = 0) in vec4 PLSIn0;
layout(location = 1) out vec4 PLSOut1;
layout(location = 1) in vec4 PLSIn1;
layout(location = 2) out vec4 PLSOut2;
in vec4 PLSIn2;
layout(location = 3) out vec4 PLSOut3;
in vec4 PLSIn3;
void main()
{
PLSOut0 = (PLSIn0 * 2.000000);
PLSOut1 = (PLSIn1 * 6.000000);
PLSOut2 = (PLSIn2 * 7.000000);
PLSOut3 = (PLSIn3 * 4.000000);
}

21
reference/shaders/frag/sampler.frag Normal file
View File

@ -0,0 +1,21 @@
#version 310 es
precision mediump float;
precision highp int;
layout(binding = 0) uniform mediump sampler2D uTex;
layout(location = 0) out vec4 FragColor;
in vec4 vColor;
in vec2 vTex;
vec4 sample_texture(mediump sampler2D tex, vec2 uv)
{
return texture(tex, uv);
}
void main()
{
highp vec2 param = vTex;
FragColor = (vColor * sample_texture(uTex, param));
}

20
reference/shaders/frag/swizzle.frag Normal file
View File

@ -0,0 +1,20 @@
#version 310 es
precision mediump float;
precision highp int;
layout(location = 0) uniform mediump sampler2D samp;
layout(location = 0) out vec4 FragColor;
layout(location = 2) in vec2 vUV;
layout(location = 1) in vec3 vNormal;
void main()
{
FragColor = vec4(texture(samp, vUV).xyz, 1.000000);
FragColor = vec4(texture(samp, vUV).xz, 1.000000, 4.000000);
FragColor = vec4(texture(samp, vUV).xx, texture(samp, (vUV + vec2(0.100000))).yy);
FragColor = vec4(vNormal, 1.000000);
FragColor = vec4((vNormal + vec3(1.800000)), 1.000000);
FragColor = vec4(vUV, (vUV + vec2(1.800000)));
}

26
reference/shaders/frag/ubo_layout.frag Normal file
View File

@ -0,0 +1,26 @@
#version 310 es
precision mediump float;
precision highp int;
struct Str
{
mat4 foo;
};
layout(binding = 0, std140) uniform UBO1
{
layout(row_major) Str foo;
} ubo1;
layout(binding = 1, std140) uniform UBO2
{
Str foo;
} ubo0;
layout(location = 0) out vec4 FragColor;
void main()
{
FragColor = (ubo1.foo.foo[0] + ubo0.foo.foo[0]);
}

26
reference/shaders/geom/basic.geom Normal file
View File

@ -0,0 +1,26 @@
#version 310 es
#extension GL_EXT_geometry_shader : require
layout(invocations = 4, triangles) in;
layout(max_vertices = 3, triangle_strip) out;
out vec3 vNormal;
in VertexData
{
vec3 normal;
} vin[3];
void main()
{
gl_Position = gl_in[0].gl_Position;
vNormal = (vin[0].normal + vec3(float(gl_InvocationID)));
EmitVertex();
gl_Position = gl_in[1].gl_Position;
vNormal = (vin[1].normal + vec3((4.000000 * float(gl_InvocationID))));
EmitVertex();
gl_Position = gl_in[2].gl_Position;
vNormal = (vin[2].normal + vec3((2.000000 * float(gl_InvocationID))));
EmitVertex();
EndPrimitive();
}

17
reference/shaders/tesc/basic.tesc Normal file
View File

@ -0,0 +1,17 @@
#version 310 es
#extension GL_EXT_tessellation_shader : require
layout(vertices = 1) out;
out patch vec3 vFoo;
void main()
{
gl_TessLevelInner[0] = 8.900000;
gl_TessLevelInner[1] = 6.900000;
gl_TessLevelOuter[0] = 8.900000;
gl_TessLevelOuter[1] = 6.900000;
gl_TessLevelOuter[2] = 3.900000;
gl_TessLevelOuter[3] = 4.900000;
vFoo = vec3(1.000000);
}

117
reference/shaders/tesc/water_tess.tesc Normal file
View File

@ -0,0 +1,117 @@
#version 310 es
#extension GL_EXT_tessellation_shader : require
layout(vertices = 1) out;
layout(std140) uniform UBO
{
vec4 uScale;
vec3 uCamPos;
vec2 uPatchSize;
vec2 uMaxTessLevel;
float uDistanceMod;
vec4 uFrustum[6];
} _41;
out patch vec2 vOutPatchPosBase;
out patch vec4 vPatchLods;
in vec2 vPatchPosBase[32];
bool frustum_cull(vec2 p0)
{
vec2 min_xz = ((p0 - vec2(10.000000)) * _41.uScale.xy);
vec2 max_xz = (((p0 + _41.uPatchSize) + vec2(10.000000)) * _41.uScale.xy);
vec3 bb_min = vec3(min_xz.x, -10.000000, min_xz.y);
vec3 bb_max = vec3(max_xz.x, 10.000000, max_xz.y);
vec3 center = ((bb_min + bb_max) * 0.500000);
float radius = (0.500000 * length((bb_max - bb_min)));
vec3 f0 = vec3(dot(_41.uFrustum[0], vec4(center, 1.000000)), dot(_41.uFrustum[1], vec4(center, 1.000000)), dot(_41.uFrustum[2], vec4(center, 1.000000)));
vec3 f1 = vec3(dot(_41.uFrustum[3], vec4(center, 1.000000)), dot(_41.uFrustum[4], vec4(center, 1.000000)), dot(_41.uFrustum[5], vec4(center, 1.000000)));
bool _205 = any(lessThanEqual(f0, vec3((-radius))));
bool _215;
if ((!_205))
{
_215 = any(lessThanEqual(f1, vec3((-radius))));
}
else
{
_215 = _205;
}
return (!_215);
}
float lod_factor(vec2 pos_)
{
vec2 pos = (pos_ * _41.uScale.xy);
vec3 dist_to_cam = (_41.uCamPos - vec3(pos.x, 0.000000, pos.y));
float level = log2(((length(dist_to_cam) + 0.000100) * _41.uDistanceMod));
return clamp(level, 0.000000, _41.uMaxTessLevel.x);
}
vec4 tess_level(vec4 lod)
{
return (exp2((-lod)) * _41.uMaxTessLevel.y);
}
float tess_level(float lod)
{
return (_41.uMaxTessLevel.y * exp2((-lod)));
}
void compute_tess_levels(vec2 p0)
{
vOutPatchPosBase = p0;
vec2 param = (p0 + (vec2(-0.500000) * _41.uPatchSize));
float l00 = lod_factor(param);
vec2 param_1 = (p0 + (vec2(0.500000, -0.500000) * _41.uPatchSize));
float l10 = lod_factor(param_1);
vec2 param_2 = (p0 + (vec2(1.500000, -0.500000) * _41.uPatchSize));
float l20 = lod_factor(param_2);
vec2 param_3 = (p0 + (vec2(-0.500000, 0.500000) * _41.uPatchSize));
float l01 = lod_factor(param_3);
vec2 param_4 = (p0 + (vec2(0.500000) * _41.uPatchSize));
float l11 = lod_factor(param_4);
vec2 param_5 = (p0 + (vec2(1.500000, 0.500000) * _41.uPatchSize));
float l21 = lod_factor(param_5);
vec2 param_6 = (p0 + (vec2(-0.500000, 1.500000) * _41.uPatchSize));
float l02 = lod_factor(param_6);
vec2 param_7 = (p0 + (vec2(0.500000, 1.500000) * _41.uPatchSize));
float l12 = lod_factor(param_7);
vec2 param_8 = (p0 + (vec2(1.500000) * _41.uPatchSize));
float l22 = lod_factor(param_8);
vec4 lods = vec4(dot(vec4(l01, l11, l02, l12), vec4(0.250000)), dot(vec4(l00, l10, l01, l11), vec4(0.250000)), dot(vec4(l10, l20, l11, l21), vec4(0.250000)), dot(vec4(l11, l21, l12, l22), vec4(0.250000)));
vPatchLods = lods;
vec4 outer_lods = min(lods, lods.yzwx);
vec4 param_9 = outer_lods;
vec4 levels = tess_level(param_9);
gl_TessLevelOuter[0] = levels.x;
gl_TessLevelOuter[1] = levels.y;
gl_TessLevelOuter[2] = levels.z;
gl_TessLevelOuter[3] = levels.w;
float min_lod = min(min(lods.x, lods.y), min(lods.z, lods.w));
float param_10 = min(min_lod, l11);
float inner = tess_level(param_10);
gl_TessLevelInner[0] = inner;
gl_TessLevelInner[1] = inner;
}
void main()
{
vec2 p0 = vPatchPosBase[0];
vec2 param = p0;
vec2 param_1;
if ((!frustum_cull(param)))
{
gl_TessLevelOuter[0] = -1.000000;
gl_TessLevelOuter[1] = -1.000000;
gl_TessLevelOuter[2] = -1.000000;
gl_TessLevelOuter[3] = -1.000000;
gl_TessLevelInner[0] = -1.000000;
gl_TessLevelInner[1] = -1.000000;
}
else
{
param_1 = p0;
compute_tess_levels(param_1);
}
}

61
reference/shaders/tese/water_tess.tese Normal file
View File

@ -0,0 +1,61 @@
#version 310 es
#extension GL_EXT_tessellation_shader : require
layout(quads, cw, fractional_even_spacing) in;
layout(binding = 1, std140) uniform UBO
{
mat4 uMVP;
vec4 uScale;
vec2 uInvScale;
vec3 uCamPos;
vec2 uPatchSize;
vec2 uInvHeightmapSize;
} _31;
layout(binding = 0) uniform mediump sampler2D uHeightmapDisplacement;
in patch vec2 vOutPatchPosBase;
in patch vec4 vPatchLods;
out vec4 vGradNormalTex;
out vec3 vWorld;
vec2 lerp_vertex(vec2 tess_coord)
{
return (vOutPatchPosBase + (tess_coord * _31.uPatchSize));
}
mediump vec2 lod_factor(vec2 tess_coord)
{
mediump vec2 x = mix(vPatchLods.yx, vPatchLods.zw, vec2(tess_coord.x));
mediump float level = mix(x.x, x.y, tess_coord.y);
mediump float floor_level = floor(level);
mediump float fract_level = (level - floor_level);
return vec2(floor_level, fract_level);
}
mediump vec3 sample_height_displacement(vec2 uv, vec2 off, mediump vec2 lod)
{
return mix(textureLod(uHeightmapDisplacement, (uv + (off * 0.500000)), lod.x).xyz, textureLod(uHeightmapDisplacement, (uv + (off * 1.000000)), (lod.x + 1.000000)).xyz, vec3(lod.y));
}
void main()
{
vec2 tess_coord = gl_TessCoord.xy;
vec2 param = tess_coord;
vec2 pos = lerp_vertex(param);
vec2 param_1 = tess_coord;
mediump vec2 lod = lod_factor(param_1);
vec2 tex = (pos * _31.uInvHeightmapSize);
pos = (pos * _31.uScale.xy);
mediump float delta_mod = exp2(lod.x);
vec2 off = (_31.uInvHeightmapSize * delta_mod);
vGradNormalTex = vec4((tex + (_31.uInvHeightmapSize * 0.500000)), (tex * _31.uScale.zw));
vec2 param_2 = tex;
vec2 param_3 = off;
vec2 param_4 = lod;
vec3 height_displacement = sample_height_displacement(param_2, param_3, param_4);
pos = (pos + height_displacement.yz);
vWorld = vec3(pos.x, height_displacement.x, pos.y);
gl_Position = (_31.uMVP * vec4(vWorld, 1.000000));
}

17
reference/shaders/vert/basic.vert Normal file
View File

@ -0,0 +1,17 @@
#version 310 es
layout(std140) uniform UBO
{
mat4 uMVP;
} _16;
in vec4 aVertex;
out vec3 vNormal;
in vec3 aNormal;
void main()
{
gl_Position = (_16.uMVP * aVertex);
vNormal = aNormal;
}

111
reference/shaders/vert/ground.vert Normal file
View File

@ -0,0 +1,111 @@
#version 310 es
struct PatchData
{
vec4 Position;
vec4 LODs;
};
layout(binding = 0, std140) uniform GlobalVSData
{
vec4 g_ViewProj_Row0;
vec4 g_ViewProj_Row1;
vec4 g_ViewProj_Row2;
vec4 g_ViewProj_Row3;
vec4 g_CamPos;
vec4 g_CamRight;
vec4 g_CamUp;
vec4 g_CamFront;
vec4 g_SunDir;
vec4 g_SunColor;
vec4 g_TimeParams;
vec4 g_ResolutionParams;
vec4 g_CamAxisRight;
vec4 g_FogColor_Distance;
vec4 g_ShadowVP_Row0;
vec4 g_ShadowVP_Row1;
vec4 g_ShadowVP_Row2;
vec4 g_ShadowVP_Row3;
} _58;
layout(binding = 0, std140) uniform PerPatch
{
PatchData Patches[256];
} _284;
layout(binding = 2, std140) uniform GlobalGround
{
vec4 GroundScale;
vec4 GroundPosition;
vec4 InvGroundSize_PatchScale;
} _381;
layout(binding = 1) uniform mediump sampler2D TexLOD;
layout(binding = 0) uniform mediump sampler2D TexHeightmap;
layout(location = 1) in vec4 LODWeights;
layout(location = 0) in vec2 Position;
layout(location = 1) out vec3 EyeVec;
layout(location = 0) out vec2 TexCoord;
vec2 warp_position()
{
float vlod = dot(LODWeights, _284.Patches[gl_InstanceID].LODs);
vlod = mix(vlod, _284.Patches[gl_InstanceID].Position.w, all(equal(LODWeights, vec4(0.000000))));
float floor_lod = floor(vlod);
float fract_lod = (vlod - floor_lod);
uint ufloor_lod = uint(floor_lod);
uvec2 uPosition = uvec2(Position);
uvec2 mask = ((uvec2(1u) << uvec2(ufloor_lod, (ufloor_lod + 1u))) - uvec2(1u));
uvec2 rounding;
uint _332;
uint _343;
vec4 lower_upper_snapped;
if ((uPosition.x < 32u))
{
_332 = mask.x;
}
else
{
_332 = 0u;
}
uint _342 = _332;
if ((uPosition.y < 32u))
{
_343 = mask.y;
}
else
{
_343 = 0u;
}
rounding = uvec2(_342, _343);
lower_upper_snapped = vec4(((uPosition + rounding).xyxy & (~mask).xxyy));
return mix(lower_upper_snapped.xy, lower_upper_snapped.zw, vec2(fract_lod));
}
vec2 lod_factor(vec2 uv)
{
float level = (textureLod(TexLOD, uv, 0.000000).x * 7.968750);
float floor_level = floor(level);
float fract_level = (level - floor_level);
return vec2(floor_level, fract_level);
}
void main()
{
vec2 PatchPos = (_284.Patches[gl_InstanceID].Position.xz * _381.InvGroundSize_PatchScale.zw);
vec2 WarpedPos = warp_position();
vec2 VertexPos = (PatchPos + WarpedPos);
vec2 NormalizedPos = (VertexPos * _381.InvGroundSize_PatchScale.xy);
vec2 param = NormalizedPos;
vec2 lod = lod_factor(param);
vec2 Offset = (_381.InvGroundSize_PatchScale.xy * exp2(lod.x));
float Elevation = mix(textureLod(TexHeightmap, (NormalizedPos + (Offset * 0.500000)), lod.x).x, textureLod(TexHeightmap, (NormalizedPos + (Offset * 1.000000)), (lod.x + 1.000000)).x, lod.y);
vec3 WorldPos = vec3(NormalizedPos.x, Elevation, NormalizedPos.y);
WorldPos = (WorldPos * _381.GroundScale.xyz);
WorldPos = (WorldPos + _381.GroundPosition.xyz);
EyeVec = (WorldPos - _58.g_CamPos.xyz);
TexCoord = (NormalizedPos + (_381.InvGroundSize_PatchScale.xy * 0.500000));
gl_Position = ((((_58.g_ViewProj_Row0 * WorldPos.x) + (_58.g_ViewProj_Row1 * WorldPos.y)) + (_58.g_ViewProj_Row2 * WorldPos.z)) + _58.g_ViewProj_Row3);
}

133
reference/shaders/vert/ocean.vert Normal file
View File

@ -0,0 +1,133 @@
#version 310 es
struct PatchData
{
vec4 Position;
vec4 LODs;
};
layout(binding = 0, std140) uniform GlobalVSData
{
vec4 g_ViewProj_Row0;
vec4 g_ViewProj_Row1;
vec4 g_ViewProj_Row2;
vec4 g_ViewProj_Row3;
vec4 g_CamPos;
vec4 g_CamRight;
vec4 g_CamUp;
vec4 g_CamFront;
vec4 g_SunDir;
vec4 g_SunColor;
vec4 g_TimeParams;
vec4 g_ResolutionParams;
vec4 g_CamAxisRight;
vec4 g_FogColor_Distance;
vec4 g_ShadowVP_Row0;
vec4 g_ShadowVP_Row1;
vec4 g_ShadowVP_Row2;
vec4 g_ShadowVP_Row3;
} _58;
layout(binding = 0, std140) uniform Offsets
{
PatchData Patches[256];
} _284;
layout(binding = 4, std140) uniform GlobalOcean
{
vec4 OceanScale;
vec4 OceanPosition;
vec4 InvOceanSize_PatchScale;
vec4 NormalTexCoordScale;
} _405;
layout(binding = 1) uniform mediump sampler2D TexLOD;
layout(binding = 0) uniform mediump sampler2D TexDisplacement;
layout(location = 1) in vec4 LODWeights;
layout(location = 0) in vec4 Position;
layout(location = 0) out vec3 EyeVec;
layout(location = 1) out vec4 TexCoord;
vec2 warp_position()
{
float vlod = dot(LODWeights, _284.Patches[gl_InstanceID].LODs);
vlod = mix(vlod, _284.Patches[gl_InstanceID].Position.w, all(equal(LODWeights, vec4(0.000000))));
float floor_lod = floor(vlod);
float fract_lod = (vlod - floor_lod);
uint ufloor_lod = uint(floor_lod);
uvec4 uPosition = uvec4(Position);
uvec2 mask = ((uvec2(1u) << uvec2(ufloor_lod, (ufloor_lod + 1u))) - uvec2(1u));
uvec4 rounding;
uint _333;
uint _345;
uint _356;
uint _368;
vec4 lower_upper_snapped;
if ((uPosition.x < 32u))
{
_333 = mask.x;
}
else
{
_333 = 0u;
}
rounding.x = _333;
if ((uPosition.y < 32u))
{
_345 = mask.x;
}
else
{
_345 = 0u;
}
rounding.y = _345;
if ((uPosition.x < 32u))
{
_356 = mask.y;
}
else
{
_356 = 0u;
}
rounding.z = _356;
if ((uPosition.y < 32u))
{
_368 = mask.y;
}
else
{
_368 = 0u;
}
rounding.w = _368;
lower_upper_snapped = vec4(((uPosition.xyxy + rounding) & (~mask).xxyy));
return mix(lower_upper_snapped.xy, lower_upper_snapped.zw, vec2(fract_lod));
}
vec2 lod_factor(vec2 uv)
{
float level = (textureLod(TexLOD, uv, 0.000000).x * 7.968750);
float floor_level = floor(level);
float fract_level = (level - floor_level);
return vec2(floor_level, fract_level);
}
void main()
{
vec2 PatchPos = (_284.Patches[gl_InstanceID].Position.xz * _405.InvOceanSize_PatchScale.zw);
vec2 WarpedPos = warp_position();
vec2 VertexPos = (PatchPos + WarpedPos);
vec2 NormalizedPos = (VertexPos * _405.InvOceanSize_PatchScale.xy);
vec2 NormalizedTex = (NormalizedPos * _405.NormalTexCoordScale.zw);
vec2 param = NormalizedPos;
vec2 lod = lod_factor(param);
vec2 Offset = ((_405.InvOceanSize_PatchScale.xy * exp2(lod.x)) * _405.NormalTexCoordScale.zw);
vec3 Displacement = mix(textureLod(TexDisplacement, (NormalizedTex + (Offset * 0.500000)), lod.x).yxz, textureLod(TexDisplacement, (NormalizedTex + (Offset * 1.000000)), (lod.x + 1.000000)).yxz, vec3(lod.y));
vec3 WorldPos = (vec3(NormalizedPos.x, 0.000000, NormalizedPos.y) + Displacement);
WorldPos = (WorldPos * _405.OceanScale.xyz);
WorldPos = (WorldPos + _405.OceanPosition.xyz);
EyeVec = (WorldPos - _58.g_CamPos.xyz);
TexCoord = (vec4(NormalizedTex, (NormalizedTex * _405.NormalTexCoordScale.xy)) + ((_405.InvOceanSize_PatchScale.xyxy * 0.500000) * _405.NormalTexCoordScale.zwzw));
gl_Position = ((((_58.g_ViewProj_Row0 * WorldPos.x) + (_58.g_ViewProj_Row1 * WorldPos.y)) + (_58.g_ViewProj_Row2 * WorldPos.z)) + _58.g_ViewProj_Row3);
}

11
reference/shaders/vert/texture_buffer.vert Normal file
View File

@ -0,0 +1,11 @@
#version 310 es
#extension GL_OES_texture_buffer : require
layout(binding = 4) uniform highp samplerBuffer uSamp;
layout(binding = 5, rgba32f) uniform highp readonly imageBuffer uSampo;
void main()
{
gl_Position = (texelFetch(uSamp, 10) + imageLoad(uSampo, 100));
}

17
reference/shaders/vert/ubo.vert Normal file
View File

@ -0,0 +1,17 @@
#version 310 es
layout(binding = 0, std140) uniform UBO
{
mat4 mvp;
} _16;
in vec4 aVertex;
out vec3 vNormal;
in vec3 aNormal;
void main()
{
gl_Position = (_16.mvp * aVertex);
vNormal = aNormal;
}

51
shaders/comp/atomic.comp Normal file
View File

@ -0,0 +1,51 @@
#version 310 es
#extension GL_OES_shader_image_atomic : require
layout(local_size_x = 1) in;
layout(r32ui, binding = 0) uniform highp uimage2D uImage;
layout(r32i, binding = 1) uniform highp iimage2D iImage;
layout(binding = 2, std430) buffer SSBO
{
uint u32;
int i32;
} ssbo;
void main()
{
imageAtomicAdd(uImage, ivec2(1, 5), 1u);
imageAtomicOr(uImage, ivec2(1, 5), 1u);
imageAtomicXor(uImage, ivec2(1, 5), 1u);
imageAtomicAnd(uImage, ivec2(1, 5), 1u);
imageAtomicMin(uImage, ivec2(1, 5), 1u);
imageAtomicMax(uImage, ivec2(1, 5), 1u);
//imageAtomicExchange(uImage, ivec2(1, 5), 1u);
imageAtomicCompSwap(uImage, ivec2(1, 5), 10u, 2u);
imageAtomicAdd(iImage, ivec2(1, 6), 1);
imageAtomicOr(iImage, ivec2(1, 6), 1);
imageAtomicXor(iImage, ivec2(1, 6), 1);
imageAtomicAnd(iImage, ivec2(1, 6), 1);
imageAtomicMin(iImage, ivec2(1, 6), 1);
imageAtomicMax(iImage, ivec2(1, 6), 1);
//imageAtomicExchange(iImage, ivec2(1, 5), 1u);
imageAtomicCompSwap(iImage, ivec2(1, 5), 10, 2);
atomicAdd(ssbo.u32, 1u);
atomicOr(ssbo.u32, 1u);
atomicXor(ssbo.u32, 1u);
atomicAnd(ssbo.u32, 1u);
atomicMin(ssbo.u32, 1u);
atomicMax(ssbo.u32, 1u);
atomicExchange(ssbo.u32, 1u);
atomicCompSwap(ssbo.u32, 10u, 2u);
atomicAdd(ssbo.i32, 1);
atomicOr(ssbo.i32, 1);
atomicXor(ssbo.i32, 1);
atomicAnd(ssbo.i32, 1);
atomicMin(ssbo.i32, 1);
atomicMax(ssbo.i32, 1);
atomicExchange(ssbo.i32, 1);
atomicCompSwap(ssbo.i32, 10, 2);
}

55
shaders/comp/bake_gradient.comp Normal file
View File

@ -0,0 +1,55 @@
#version 310 es
layout(local_size_x = 8, local_size_y = 8) in;
layout(binding = 0) uniform sampler2D uHeight;
layout(binding = 1) uniform sampler2D uDisplacement;
layout(rgba16f, binding = 2) uniform writeonly mediump image2D iHeightDisplacement;
layout(rgba16f, binding = 3) uniform writeonly mediump image2D iGradJacobian;
layout(binding = 4) uniform UBO
{
vec4 uInvSize;
vec4 uScale;
};
mediump float jacobian(mediump vec2 dDdx, mediump vec2 dDdy)
{
return (1.0 + dDdx.x) * (1.0 + dDdy.y) - dDdx.y * dDdy.x;
}
#define LAMBDA 1.2
void main()
{
vec4 uv = (vec2(gl_GlobalInvocationID.xy) * uInvSize.xy).xyxy + 0.5 * uInvSize;
float h = textureLod(uHeight, uv.xy, 0.0).x;
// Compute the heightmap gradient by simple differentiation.
float x0 = textureLodOffset(uHeight, uv.xy, 0.0, ivec2(-1, 0)).x;
float x1 = textureLodOffset(uHeight, uv.xy, 0.0, ivec2(+1, 0)).x;
float y0 = textureLodOffset(uHeight, uv.xy, 0.0, ivec2(0, -1)).x;
float y1 = textureLodOffset(uHeight, uv.xy, 0.0, ivec2(0, +1)).x;
vec2 grad = uScale.xy * 0.5 * vec2(x1 - x0, y1 - y0);
// Displacement map must be sampled with a different offset since it's a smaller texture.
vec2 displacement = LAMBDA * textureLod(uDisplacement, uv.zw, 0.0).xy;
// Compute jacobian.
vec2 dDdx = 0.5 * LAMBDA * (
textureLodOffset(uDisplacement, uv.zw, 0.0, ivec2(+1, 0)).xy -
textureLodOffset(uDisplacement, uv.zw, 0.0, ivec2(-1, 0)).xy);
vec2 dDdy = 0.5 * LAMBDA * (
textureLodOffset(uDisplacement, uv.zw, 0.0, ivec2(0, +1)).xy -
textureLodOffset(uDisplacement, uv.zw, 0.0, ivec2(0, -1)).xy);
float j = jacobian(dDdx * uScale.z, dDdy * uScale.z);
displacement = vec2(0.0);
// Read by vertex shader/tess shader.
imageStore(iHeightDisplacement, ivec2(gl_GlobalInvocationID.xy), vec4(h, displacement, 0.0));
// Read by fragment shader.
imageStore(iGradJacobian, ivec2(gl_GlobalInvocationID.xy), vec4(grad, j, 0.0));
}

28
shaders/comp/basic.comp Normal file
View File

@ -0,0 +1,28 @@
#version 310 es
layout(local_size_x = 1) in;
layout(std430, binding = 0) readonly buffer SSBO
{
vec4 in_data[];
};
layout(std430, binding = 1) writeonly buffer SSBO2
{
vec4 out_data[];
};
layout(std430, binding = 2) buffer SSBO3
{
uint counter;
};
void main()
{
uint ident = gl_GlobalInvocationID.x;
vec4 idata = in_data[ident];
if (dot(idata, vec4(1.0, 5.0, 6.0, 2.0)) > 8.2)
{
out_data[atomicAdd(counter, 1u)] = idata;
}
}

26
shaders/comp/culling.comp Normal file
View File

@ -0,0 +1,26 @@
#version 310 es
layout(local_size_x = 4) in;
layout(std430, binding = 0) readonly buffer SSBO
{
float in_data[];
};
layout(std430, binding = 1) writeonly buffer SSBO2
{
float out_data[];
};
layout(std430, binding = 2) buffer SSBO3
{
uint count;
};
void main()
{
uint ident = gl_GlobalInvocationID.x;
float idata = in_data[ident];
if (idata > 12.0)
out_data[atomicAdd(count, 1u)] = idata;
}

31
shaders/comp/dowhile.comp Normal file
View File

@ -0,0 +1,31 @@
#version 310 es
layout(local_size_x = 1) in;
layout(std430, binding = 0) readonly buffer SSBO
{
mat4 mvp;
vec4 in_data[];
};
layout(std430, binding = 1) writeonly buffer SSBO2
{
vec4 out_data[];
};
int i;
void main()
{
uint ident = gl_GlobalInvocationID.x;
i = 0;
vec4 idat = in_data[ident];
do
{
idat = mvp * idat;
i++;
} while(i < 16);
out_data[ident] = idat;
}

97
shaders/comp/generate_height.comp Normal file
View File

@ -0,0 +1,97 @@
#version 310 es
layout(local_size_x = 64) in;
layout(std430, binding = 0) readonly buffer Distribution
{
vec2 distribution[];
};
layout(std430, binding = 1) writeonly buffer HeightmapFFT
{
uint heights[];
};
layout(binding = 2, std140) uniform UBO
{
vec4 uModTime;
};
vec2 alias(vec2 i, vec2 N)
{
return mix(i, i - N, greaterThan(i, 0.5 * N));
}
vec4 cmul(vec4 a, vec4 b)
{
vec4 r3 = a.yxwz;
vec4 r1 = b.xxzz;
vec4 R0 = a * r1;
vec4 r2 = b.yyww;
vec4 R1 = r2 * r3;
return R0 + vec4(-R1.x, R1.y, -R1.z, R1.w);
}
vec2 cmul(vec2 a, vec2 b)
{
vec2 r3 = a.yx;
vec2 r1 = b.xx;
vec2 R0 = a * r1;
vec2 r2 = b.yy;
vec2 R1 = r2 * r3;
return R0 + vec2(-R1.x, R1.y);
}
uint pack2(vec2 v)
{
return packHalf2x16(v);
}
uvec2 pack4(vec4 v)
{
return uvec2(packHalf2x16(v.xy), packHalf2x16(v.zw));
}
uvec2 workaround_mix(uvec2 a, uvec2 b, bvec2 sel)
{
return uvec2(sel.x ? b.x : a.x, sel.y ? b.y : a.y);
}
void generate_heightmap()
{
uvec2 N = gl_WorkGroupSize.xy * gl_NumWorkGroups.xy;
uvec2 i = gl_GlobalInvocationID.xy;
// Pick out the negative frequency variant.
uvec2 wi = workaround_mix(N - i, uvec2(0u), equal(i, uvec2(0u)));
// Pick out positive and negative travelling waves.
vec2 a = distribution[i.y * N.x + i.x];
vec2 b = distribution[wi.y * N.x + wi.x];
vec2 k = uModTime.xy * alias(vec2(i), vec2(N));
float k_len = length(k);
const float G = 9.81;
// If this sample runs for hours on end, the cosines of very large numbers will eventually become unstable.
// It is fairly easy to fix this by wrapping uTime,
// and quantizing w such that wrapping uTime does not change the result.
// See Tessendorf's paper for how to do it.
// The sqrt(G * k_len) factor represents how fast ocean waves at different frequencies propagate.
float w = sqrt(G * k_len) * uModTime.z;
float cw = cos(w);
float sw = sin(w);
// Complex multiply to rotate our frequency samples.
a = cmul(a, vec2(cw, sw));
b = cmul(b, vec2(cw, sw));
b = vec2(b.x, -b.y); // Complex conjugate since we picked a frequency with the opposite direction.
vec2 res = a + b; // Sum up forward and backwards travelling waves.
heights[i.y * N.x + i.x] = pack2(res);
}
void main()
{
generate_heightmap();
}

12
shaders/comp/image.comp Normal file
View File

@ -0,0 +1,12 @@
#version 310 es
layout(local_size_x = 1) in;
layout(rgba8, binding = 0) uniform readonly mediump image2D uImageIn;
layout(rgba8, binding = 1) uniform writeonly mediump image2D uImageOut;
void main()
{
vec4 v = imageLoad(uImageIn, ivec2(gl_GlobalInvocationID.xy) + imageSize(uImageIn));
imageStore(uImageOut, ivec2(gl_GlobalInvocationID.xy), v);
}

55
shaders/comp/inout-struct.comp Normal file
View File

@ -0,0 +1,55 @@
#version 310 es
layout(local_size_x = 1) in;
layout(std430, binding = 0) writeonly buffer SSBO
{
vec4 data[];
} outdata;
layout(std430, binding = 1) readonly buffer SSBO2
{
vec4 data[];
} indata;
struct Foo
{
vec4 a;
vec4 b;
vec4 c;
vec4 d;
};
layout(std430, binding = 2) readonly buffer SSBO3
{
Foo foos[];
} foobar;
vec4 bar(Foo foo)
{
return foo.a + foo.b + foo.c + foo.d;
}
void baz(out Foo foo)
{
uint ident = gl_GlobalInvocationID.x;
foo.a = indata.data[4u * ident + 0u];
foo.b = indata.data[4u * ident + 1u];
foo.c = indata.data[4u * ident + 2u];
foo.d = indata.data[4u * ident + 3u];
}
void meow(inout Foo foo)
{
foo.a += 10.0;
foo.b += 20.0;
foo.c += 30.0;
foo.d += 40.0;
}
void main()
{
Foo foo;
baz(foo);
meow(foo);
outdata.data[gl_GlobalInvocationID.x] = bar(foo) + bar(foobar.foos[gl_GlobalInvocationID.x]);
}

18
shaders/comp/insert.comp Normal file
View File

@ -0,0 +1,18 @@
#version 310 es
layout(local_size_x = 1) in;
layout(std430, binding = 0) writeonly buffer SSBO
{
vec4 out_data[];
};
void main()
{
vec4 v;
v.x = 10.0;
v.y = 30.0;
v.z = 70.0;
v.w = 90.0;
out_data[gl_GlobalInvocationID.x] = v;
out_data[gl_GlobalInvocationID.x].y = 20.0;
}

66
shaders/comp/loop.comp Normal file
View File

@ -0,0 +1,66 @@
#version 310 es
layout(local_size_x = 1) in;
layout(std430, binding = 0) readonly buffer SSBO
{
mat4 mvp;
vec4 in_data[];
};
layout(std430, binding = 1) writeonly buffer SSBO2
{
vec4 out_data[];
};
void main()
{
uint ident = gl_GlobalInvocationID.x;
vec4 idat = in_data[ident];
int k = 0;
while (k < 10)
{
idat *= 2.0;
k++;
}
for (uint i = 0u; i < 16u; i++, k++)
for (uint j = 0u; j < 30u; j++)
idat = mvp * idat;
k = 0;
for (;;)
{
k++;
if (k > 10)
{
k += 2;
}
else
{
k += 3;
continue;
}
k += 10;
}
k = 0;
do
{
k++;
} while (k > 10);
int l = 0;
for (;; l++)
{
if (l == 5)
{
continue;
}
idat += 1.0;
}
out_data[ident] = idat;
}

14
shaders/comp/mat3.comp Normal file
View File

@ -0,0 +1,14 @@
#version 310 es
layout(local_size_x = 1) in;
layout(std430, binding = 1) writeonly buffer SSBO2
{
mat3 out_data[];
};
void main()
{
uint ident = gl_GlobalInvocationID.x;
out_data[ident] = mat3(vec3(10.0), vec3(20.0), vec3(40.0));
}

23
shaders/comp/modf.comp Normal file
View File

@ -0,0 +1,23 @@
#version 310 es
layout(local_size_x = 1) in;
layout(std430, binding = 0) readonly buffer SSBO
{
vec4 in_data[];
};
layout(std430, binding = 1) writeonly buffer SSBO2
{
vec4 out_data[];
};
void main()
{
uint ident = gl_GlobalInvocationID.x;
vec4 i;
//vec4 v = frexp(in_data[ident], i);
//out_data[ident] = ldexp(v, i);
vec4 v = modf(in_data[ident], i);
out_data[ident] = v;
}

33
shaders/comp/return.comp Normal file
View File

@ -0,0 +1,33 @@
#version 310 es
layout(local_size_x = 1) in;
layout(std430, binding = 1) writeonly buffer SSBO2
{
vec4 out_data[];
};
void main()
{
uint ident = gl_GlobalInvocationID.x;
if (ident == 2u)
{
out_data[ident] = vec4(20.0);
}
else if (ident == 4u)
{
out_data[ident] = vec4(10.0);
return;
}
for (int i = 0; i < 20; i++)
{
if (i == 10)
break;
return;
}
out_data[ident] = vec4(10.0);
}

27
shaders/comp/shared.comp Normal file
View File

@ -0,0 +1,27 @@
#version 310 es
layout(local_size_x = 4) in;
shared float sShared[gl_WorkGroupSize.x];
layout(std430, binding = 0) readonly buffer SSBO
{
float in_data[];
};
layout(std430, binding = 1) writeonly buffer SSBO2
{
float out_data[];
};
void main()
{
uint ident = gl_GlobalInvocationID.x;
float idata = in_data[ident];
sShared[gl_LocalInvocationIndex] = idata;
memoryBarrierShared();
barrier();
out_data[ident] = sShared[gl_WorkGroupSize.x - gl_LocalInvocationIndex - 1u];
}

24
shaders/comp/struct-layout.comp Normal file
View File

@ -0,0 +1,24 @@
#version 310 es
layout(local_size_x = 1) in;
struct Foo
{
mat4 m;
};
layout(std430, binding = 0) readonly buffer SSBO
{
Foo in_data[];
};
layout(std430, binding = 1) writeonly buffer SSBO2
{
Foo out_data[];
};
void main()
{
uint ident = gl_GlobalInvocationID.x;
out_data[ident].m = in_data[ident].m * in_data[ident].m;
}

69
shaders/comp/struct-packing.comp Normal file
View File

@ -0,0 +1,69 @@
#version 310 es
struct S0
{
vec2 a[1];
float b;
};
struct S1
{
vec3 a;
float b;
};
struct S2
{
vec3 a[1];
float b;
};
struct S3
{
vec2 a;
float b;
};
struct Content
{
S0 m0s[1];
S1 m1s[1];
S2 m2s[1];
S0 m0;
S1 m1;
S2 m2;
S3 m3;
float m4;
};
layout(binding = 0, std140) buffer SSBO0
{
Content content;
Content content1[2];
Content content2;
float array[];
} ssbo_140;
layout(binding = 1, std430) buffer SSBO1
{
Content content;
Content content1[2];
Content content2;
layout(column_major) mat2 m0;
layout(column_major) mat2 m1;
layout(column_major) mat2x3 m2[4];
layout(column_major) mat3x2 m3;
layout(row_major) mat2 m4;
layout(row_major) mat2 m5[9];
layout(row_major) mat2x3 m6[4][2];
layout(row_major) mat3x2 m7;
float array[];
} ssbo_430;
void main()
{
ssbo_430.content = ssbo_140.content;
}

40
shaders/comp/torture-loop.comp Normal file
View File

@ -0,0 +1,40 @@
#version 310 es
layout(local_size_x = 1) in;
layout(std430, binding = 0) readonly buffer SSBO
{
mat4 mvp;
vec4 in_data[];
};
layout(std430, binding = 1) writeonly buffer SSBO2
{
vec4 out_data[];
};
void main()
{
uint ident = gl_GlobalInvocationID.x;
vec4 idat = in_data[ident];
int k = 0;
// Continue with side effects.
while (++k < 10)
{
idat *= 2.0;
k++;
}
// Again used here ...
for (uint i = 0u; i < 16u; i++, k++)
for (uint j = 0u; j < 30u; j++)
idat = mvp * idat;
do
{
k++;
} while (k > 10);
out_data[ident] = idat;
}

17
shaders/comp/udiv.comp Normal file
View File

@ -0,0 +1,17 @@
#version 310 es
layout(local_size_x = 1) in;
layout(std430, binding = 0) buffer SSBO
{
uint inputs[];
};
layout(std430, binding = 0) buffer SSBO2
{
uint outputs[];
};
void main()
{
outputs[gl_GlobalInvocationID.x] = inputs[gl_GlobalInvocationID.x] / 29u;
}

13
shaders/frag/basic.frag Normal file
View File

@ -0,0 +1,13 @@
#version 310 es
precision mediump float;
in vec4 vColor;
in vec2 vTex;
layout(binding = 0) uniform sampler2D uTex;
layout(location = 0) out vec4 FragColor;
void main()
{
FragColor = vColor * texture(uTex, vTex);
}

21
shaders/frag/constant-array.frag Normal file
View File

@ -0,0 +1,21 @@
#version 310 es
precision mediump float;
layout(location = 0) out vec4 FragColor;
layout(location = 0) flat in int index;
struct Foobar { float a; float b; };
vec4 resolve(Foobar f)
{
return vec4(f.a + f.b);
}
void main()
{
const vec4 foo[3] = vec4[](vec4(1.0), vec4(2.0), vec4(3.0));
const vec4 foobars[2][2] = vec4[][](vec4[](vec4(1.0), vec4(2.0)), vec4[](vec4(8.0), vec4(10.0)));
const Foobar foos[2] = Foobar[](Foobar(10.0, 40.0), Foobar(90.0, 70.0));
FragColor = foo[index] + foobars[index][index + 1] + resolve(Foobar(10.0, 20.0)) + resolve(foos[index]);
}

27
shaders/frag/flush_params.frag Normal file
View File

@ -0,0 +1,27 @@
#version 310 es
precision mediump float;
layout(location = 0) out vec4 FragColor;
struct Structy
{
vec4 c;
};
void foo2(out Structy f)
{
f.c = vec4(10.0);
}
Structy foo()
{
Structy f;
foo2(f);
return f;
}
void main()
{
Structy s = foo();
FragColor = s.c;
}

162
shaders/frag/ground.frag Executable file
View File

@ -0,0 +1,162 @@
#version 310 es
precision mediump float;
#define DEBUG_NONE 0
#define DEBUG_DIFFUSE 1
#define DEBUG_SPECULAR 2
#define DEBUG_LIGHTING 3
#define DEBUG_FOG 4
#define DEBUG DEBUG_NONE
#define FORWARD 0
#define DEFERRED 1
#define DEFERRED_VTEX 2
float saturate(float x) { return clamp(x, 0.0, 1.0); }
layout(std140, binding = 4) uniform GlobalPSData
{
vec4 g_CamPos;
vec4 g_SunDir;
vec4 g_SunColor;
vec4 g_ResolutionParams;
vec4 g_TimeParams;
vec4 g_FogColor_Distance;
};
vec4 ComputeFogFactor(vec3 WorldPos)
{
vec4 FogData;
vec3 vEye = WorldPos - g_CamPos.xyz;
vec3 nEye = normalize(vEye);
FogData.w = exp(-dot(vEye, vEye) * g_FogColor_Distance.w * 0.75);
float fog_sun_factor = pow(saturate(dot(nEye, g_SunDir.xyz)), 8.0);
FogData.xyz = mix(vec3(1.0, 1.0, 1.0), vec3(0.6, 0.6, 0.9), nEye.y * 0.5 + 0.5);
FogData.xyz = mix(FogData.xyz, vec3(0.95, 0.87, 0.78), fog_sun_factor);
return FogData;
}
void ApplyFog(inout vec3 Color, vec4 FogData)
{
Color = mix(FogData.xyz, Color, FogData.w);
}
void ApplyLighting(inout mediump vec3 Color, mediump float DiffuseFactor)
{
mediump vec3 DiffuseLight = g_SunColor.xyz * DiffuseFactor;
mediump vec3 AmbientLight = vec3(0.2, 0.35, 0.55) * 0.5;
mediump vec3 Lighting = DiffuseLight + AmbientLight;
#if DEBUG == DEBUG_LIGHTING
Color = Lighting;
#else
Color *= Lighting;
#endif
}
#define SPECULAR 0
#define GLOSSMAP 0
void ApplySpecular(inout mediump vec3 Color, mediump vec3 EyeVec, mediump vec3 Normal, mediump vec3 SpecularColor, mediump float Shininess, mediump float FresnelAmount)
{
mediump vec3 HalfAngle = normalize(-EyeVec + g_SunDir.xyz);
mediump float v_dot_h = saturate(dot(HalfAngle, -EyeVec));
mediump float n_dot_l = saturate(dot(Normal, g_SunDir.xyz));
mediump float n_dot_h = saturate(dot(Normal, HalfAngle));
mediump float n_dot_v = saturate(dot(-EyeVec, Normal));
mediump float h_dot_l = saturate(dot(g_SunDir.xyz, HalfAngle));
const mediump float roughness_value = 0.25;
mediump float r_sq = roughness_value * roughness_value;
mediump float n_dot_h_sq = n_dot_h * n_dot_h;
mediump float roughness_a = 1.0 / (4.0 * r_sq * n_dot_h_sq * n_dot_h_sq);
mediump float roughness_b = n_dot_h_sq - 1.0;
mediump float roughness_c = r_sq * n_dot_h_sq;
mediump float roughness = saturate(roughness_a * exp(roughness_b / roughness_c));
FresnelAmount = 0.5;
mediump float fresnel_term = pow(1.0 - n_dot_v, 5.0) * (1.0 - FresnelAmount) + FresnelAmount;
mediump float geo_numerator = 2.0 * n_dot_h;
mediump float geo_denominator = 1.0 / v_dot_h;
mediump float geo_term = min(1.0, min(n_dot_v, n_dot_l) * geo_numerator * geo_denominator);
#if SPECULAR || GLOSSMAP
Color += SpecularColor * g_SunColor.xyz * fresnel_term * roughness * n_dot_l * geo_term / (n_dot_v * n_dot_l + 0.0001);
#endif
//Color = vec3(0.025 * 1.0 / (n_dot_v * n_dot_l));
}
layout(location = 0) in vec2 TexCoord;
layout(location = 1) in vec3 EyeVec;
layout(binding = 2) uniform sampler2D TexNormalmap;
//layout(binding = 3) uniform sampler2D TexScatteringLUT;
#define DIFFUSE_ONLY 0
#define GLOBAL_RENDERER DEFERRED
#define OUTPUT_FEEDBACK_TEXTURE 0
#if DIFFUSE_ONLY
layout(location = 0) out vec4 ColorOut;
layout(location = 1) out vec4 NormalOut;
#else
layout(location = 0) out vec4 AlbedoOut;
layout(location = 1) out vec4 SpecularOut;
layout(location = 2) out vec4 NormalOut;
layout(location = 3) out vec4 LightingOut;
#endif
void Resolve(vec3 Albedo, vec3 Normal, float Roughness, float Metallic)
{
#if (GLOBAL_RENDERER == FORWARD) || OUTPUT_FEEDBACK_TEXTURE
float Lighting = saturate(dot(Normal, normalize(vec3(1.0, 0.5, 1.0))));
ColorOut.xyz = Albedo * Lighting;
ColorOut.w = 1.0;
#elif DIFFUSE_ONLY
ColorOut = vec4(Albedo, 0.0);
NormalOut.xyz = Normal * 0.5 + 0.5;
NormalOut.w = 1.0;
// linearize and map to 0..255 range
ColorOut.w = -0.003921569 / (gl_FragCoord.z - 1.003921569);
ColorOut.w = log2(1.0 + saturate(length(EyeVec.xyz) / 200.0));
ColorOut.w -= 1.0 / 255.0;
#else
LightingOut = vec4(0.0);
NormalOut = vec4(Normal * 0.5 + 0.5, 0.0);
SpecularOut = vec4(Roughness, Metallic, 0.0, 0.0);
AlbedoOut = vec4(Albedo, 1.0);
#endif
}
void main()
{
vec3 Normal = texture(TexNormalmap, TexCoord).xyz * 2.0 - 1.0;
Normal = normalize(Normal);
vec2 scatter_uv;
scatter_uv.x = saturate(length(EyeVec) / 1000.0);
vec3 nEye = normalize(EyeVec);
scatter_uv.y = 0.0; //nEye.x * 0.5 + 0.5;
vec3 Color = vec3(0.1, 0.3, 0.1);
vec3 grass = vec3(0.1, 0.3, 0.1);
vec3 dirt = vec3(0.1, 0.1, 0.1);
vec3 snow = vec3(0.8, 0.8, 0.8);
float grass_snow = smoothstep(0.0, 0.15, (g_CamPos.y + EyeVec.y) / 200.0);
vec3 base = mix(grass, snow, grass_snow);
float edge = smoothstep(0.7, 0.75, Normal.y);
Color = mix(dirt, base, edge);
Color *= Color;
float Roughness = 1.0 - edge * grass_snow;
Resolve(Color, Normal, Roughness, 0.0);
}

20
shaders/frag/mix.frag Normal file
View File

@ -0,0 +1,20 @@
#version 310 es
precision mediump float;
layout(location = 0) in vec4 vIn0;
layout(location = 1) in vec4 vIn1;
layout(location = 2) in float vIn2;
layout(location = 3) in float vIn3;
layout(location = 0) out vec4 FragColor;
void main()
{
bvec4 l = bvec4(false, true, false, false);
FragColor = mix(vIn0, vIn1, l);
bool f = true;
FragColor = vec4(mix(vIn2, vIn3, f));
FragColor = f ? vIn0 : vIn1;
FragColor = vec4(f ? vIn2 : vIn3);
}

20
shaders/frag/pls.frag Normal file
View File

@ -0,0 +1,20 @@
#version 310 es
precision mediump float;
layout(location = 0) in vec4 PLSIn0;
layout(location = 1) in vec4 PLSIn1;
in vec4 PLSIn2;
in vec4 PLSIn3;
layout(location = 0) out vec4 PLSOut0;
layout(location = 1) out vec4 PLSOut1;
layout(location = 2) out vec4 PLSOut2;
layout(location = 3) out vec4 PLSOut3;
void main()
{
PLSOut0 = 2.0 * PLSIn0;
PLSOut1 = 6.0 * PLSIn1;
PLSOut2 = 7.0 * PLSIn2;
PLSOut3 = 4.0 * PLSIn3;
}

18
shaders/frag/sampler.frag Normal file
View File

@ -0,0 +1,18 @@
#version 310 es
precision mediump float;
in vec4 vColor;
in vec2 vTex;
layout(binding = 0) uniform sampler2D uTex;
layout(location = 0) out vec4 FragColor;
vec4 sample_texture(sampler2D tex, vec2 uv)
{
return texture(tex, uv);
}
void main()
{
FragColor = vColor * sample_texture(uTex, vTex);
}

17
shaders/frag/swizzle.frag Normal file
View File

@ -0,0 +1,17 @@
#version 310 es
precision mediump float;
layout(location = 0) uniform sampler2D samp;
layout(location = 0) out vec4 FragColor;
layout(location = 1) in vec3 vNormal;
layout(location = 2) in vec2 vUV;
void main()
{
FragColor = vec4(texture(samp, vUV).xyz, 1.0);
FragColor = vec4(texture(samp, vUV).xz, 1.0, 4.0);
FragColor = vec4(texture(samp, vUV).xx, texture(samp, vUV + vec2(0.1)).yy);
FragColor = vec4(vNormal, 1.0);
FragColor = vec4(vNormal + 1.8, 1.0);
FragColor = vec4(vUV, vUV + 1.8);
}

24
shaders/frag/ubo_layout.frag Normal file
View File

@ -0,0 +1,24 @@
#version 310 es
precision mediump float;
layout(location = 0) out vec4 FragColor;
struct Str
{
mat4 foo;
};
layout(binding = 0, std140) uniform UBO1
{
layout(row_major) Str foo;
} ubo1;
layout(binding = 1, std140) uniform UBO2
{
layout(column_major) Str foo;
} ubo0;
void main()
{
FragColor = ubo1.foo.foo[0] + ubo0.foo.foo[0];
}

28
shaders/geom/basic.geom Normal file
View File

@ -0,0 +1,28 @@
#version 310 es
#extension GL_EXT_geometry_shader : require
layout(triangles, invocations = 4) in;
layout(triangle_strip, max_vertices = 3) out;
in VertexData {
vec3 normal;
} vin[];
out vec3 vNormal;
void main()
{
gl_Position = gl_in[0].gl_Position;
vNormal = vin[0].normal + float(gl_InvocationID);
EmitVertex();
gl_Position = gl_in[1].gl_Position;
vNormal = vin[1].normal + 4.0 * float(gl_InvocationID);
EmitVertex();
gl_Position = gl_in[2].gl_Position;
vNormal = vin[2].normal + 2.0 * float(gl_InvocationID);
EmitVertex();
EndPrimitive();
}

17
shaders/tesc/basic.tesc Normal file
View File

@ -0,0 +1,17 @@
#version 310 es
#extension GL_EXT_tessellation_shader : require
patch out vec3 vFoo;
layout(vertices = 1) out;
void main()
{
gl_TessLevelInner[0] = 8.9;
gl_TessLevelInner[1] = 6.9;
gl_TessLevelOuter[0] = 8.9;
gl_TessLevelOuter[1] = 6.9;
gl_TessLevelOuter[2] = 3.9;
gl_TessLevelOuter[3] = 4.9;
vFoo = vec3(1.0);
}

115
shaders/tesc/water_tess.tesc Normal file
View File

@ -0,0 +1,115 @@
#version 310 es
#extension GL_EXT_tessellation_shader : require
layout(vertices = 1) out;
in vec2 vPatchPosBase[];
layout(std140) uniform UBO
{
vec4 uScale;
highp vec3 uCamPos;
vec2 uPatchSize;
vec2 uMaxTessLevel;
float uDistanceMod;
vec4 uFrustum[6];
};
patch out vec2 vOutPatchPosBase;
patch out vec4 vPatchLods;
float lod_factor(vec2 pos_)
{
vec2 pos = pos_ * uScale.xy;
vec3 dist_to_cam = uCamPos - vec3(pos.x, 0.0, pos.y);
float level = log2((length(dist_to_cam) + 0.0001) * uDistanceMod);
return clamp(level, 0.0, uMaxTessLevel.x);
}
float tess_level(float lod)
{
return uMaxTessLevel.y * exp2(-lod);
}
vec4 tess_level(vec4 lod)
{
return uMaxTessLevel.y * exp2(-lod);
}
// Guard band for vertex displacement.
#define GUARD_BAND 10.0
bool frustum_cull(vec2 p0)
{
vec2 min_xz = (p0 - GUARD_BAND) * uScale.xy;
vec2 max_xz = (p0 + uPatchSize + GUARD_BAND) * uScale.xy;
vec3 bb_min = vec3(min_xz.x, -GUARD_BAND, min_xz.y);
vec3 bb_max = vec3(max_xz.x, +GUARD_BAND, max_xz.y);
vec3 center = 0.5 * (bb_min + bb_max);
float radius = 0.5 * length(bb_max - bb_min);
vec3 f0 = vec3(
dot(uFrustum[0], vec4(center, 1.0)),
dot(uFrustum[1], vec4(center, 1.0)),
dot(uFrustum[2], vec4(center, 1.0)));
vec3 f1 = vec3(
dot(uFrustum[3], vec4(center, 1.0)),
dot(uFrustum[4], vec4(center, 1.0)),
dot(uFrustum[5], vec4(center, 1.0)));
return !(any(lessThanEqual(f0, vec3(-radius))) || any(lessThanEqual(f1, vec3(-radius))));
}
void compute_tess_levels(vec2 p0)
{
vOutPatchPosBase = p0;
float l00 = lod_factor(p0 + vec2(-0.5, -0.5) * uPatchSize);
float l10 = lod_factor(p0 + vec2(+0.5, -0.5) * uPatchSize);
float l20 = lod_factor(p0 + vec2(+1.5, -0.5) * uPatchSize);
float l01 = lod_factor(p0 + vec2(-0.5, +0.5) * uPatchSize);
float l11 = lod_factor(p0 + vec2(+0.5, +0.5) * uPatchSize);
float l21 = lod_factor(p0 + vec2(+1.5, +0.5) * uPatchSize);
float l02 = lod_factor(p0 + vec2(-0.5, +1.5) * uPatchSize);
float l12 = lod_factor(p0 + vec2(+0.5, +1.5) * uPatchSize);
float l22 = lod_factor(p0 + vec2(+1.5, +1.5) * uPatchSize);
vec4 lods = vec4(
dot(vec4(l01, l11, l02, l12), vec4(0.25)),
dot(vec4(l00, l10, l01, l11), vec4(0.25)),
dot(vec4(l10, l20, l11, l21), vec4(0.25)),
dot(vec4(l11, l21, l12, l22), vec4(0.25)));
vPatchLods = lods;
vec4 outer_lods = min(lods.xyzw, lods.yzwx);
vec4 levels = tess_level(outer_lods);
gl_TessLevelOuter[0] = levels.x;
gl_TessLevelOuter[1] = levels.y;
gl_TessLevelOuter[2] = levels.z;
gl_TessLevelOuter[3] = levels.w;
float min_lod = min(min(lods.x, lods.y), min(lods.z, lods.w));
float inner = tess_level(min(min_lod, l11));
gl_TessLevelInner[0] = inner;
gl_TessLevelInner[1] = inner;
}
void main()
{
vec2 p0 = vPatchPosBase[0];
if (!frustum_cull(p0))
{
gl_TessLevelOuter[0] = -1.0;
gl_TessLevelOuter[1] = -1.0;
gl_TessLevelOuter[2] = -1.0;
gl_TessLevelOuter[3] = -1.0;
gl_TessLevelInner[0] = -1.0;
gl_TessLevelInner[1] = -1.0;
}
else
{
compute_tess_levels(p0);
}
}

65
shaders/tese/water_tess.tese Normal file
View File

@ -0,0 +1,65 @@
#version 310 es
#extension GL_EXT_tessellation_shader : require
precision highp int;
layout(cw, quads, fractional_even_spacing) in;
patch in vec2 vOutPatchPosBase;
patch in vec4 vPatchLods;
layout(binding = 1, std140) uniform UBO
{
mat4 uMVP;
vec4 uScale;
vec2 uInvScale;
vec3 uCamPos;
vec2 uPatchSize;
vec2 uInvHeightmapSize;
};
layout(binding = 0) uniform mediump sampler2D uHeightmapDisplacement;
highp out vec3 vWorld;
highp out vec4 vGradNormalTex;
vec2 lerp_vertex(vec2 tess_coord)
{
return vOutPatchPosBase + tess_coord * uPatchSize;
}
mediump vec2 lod_factor(vec2 tess_coord)
{
mediump vec2 x = mix(vPatchLods.yx, vPatchLods.zw, tess_coord.x);
mediump float level = mix(x.x, x.y, tess_coord.y);
mediump float floor_level = floor(level);
mediump float fract_level = level - floor_level;
return vec2(floor_level, fract_level);
}
mediump vec3 sample_height_displacement(vec2 uv, vec2 off, mediump vec2 lod)
{
return mix(
textureLod(uHeightmapDisplacement, uv + 0.5 * off, lod.x).xyz,
textureLod(uHeightmapDisplacement, uv + 1.0 * off, lod.x + 1.0).xyz,
lod.y);
}
void main()
{
vec2 tess_coord = gl_TessCoord.xy;
vec2 pos = lerp_vertex(tess_coord);
mediump vec2 lod = lod_factor(tess_coord);
vec2 tex = pos * uInvHeightmapSize.xy;
pos *= uScale.xy;
mediump float delta_mod = exp2(lod.x);
vec2 off = uInvHeightmapSize.xy * delta_mod;
vGradNormalTex = vec4(tex + 0.5 * uInvHeightmapSize.xy, tex * uScale.zw);
vec3 height_displacement = sample_height_displacement(tex, off, lod);
pos += height_displacement.yz;
vWorld = vec3(pos.x, height_displacement.x, pos.y);
gl_Position = uMVP * vec4(vWorld, 1.0);
}

15
shaders/vert/basic.vert Normal file
View File

@ -0,0 +1,15 @@
#version 310 es
layout(std140) uniform UBO
{
uniform mat4 uMVP;
};
in vec4 aVertex;
in vec3 aNormal;
out vec3 vNormal;
void main()
{
gl_Position = uMVP * aVertex;
vNormal = aNormal;
}

196
shaders/vert/ground.vert Executable file
View File

@ -0,0 +1,196 @@
#version 310 es
#define YFLIP 0
#define SPECULAR 0
#define GLOSSMAP 0
#define DEBUG_NONE 0
#define DEBUG_DIFFUSE 1
#define DEBUG_SPECULAR 2
#define DEBUG_LIGHTING 3
#define DEBUG_FOG 4
#define DEBUG DEBUG_NONE
#define FORWARD 0
#define DEFERRED 1
#define DEFERRED_VTEX 2
float saturate(float x) { return clamp(x, 0.0, 1.0); }
layout(std140, binding = 0) uniform GlobalVSData
{
vec4 g_ViewProj_Row0;
vec4 g_ViewProj_Row1;
vec4 g_ViewProj_Row2;
vec4 g_ViewProj_Row3;
vec4 g_CamPos;
vec4 g_CamRight;
vec4 g_CamUp;
vec4 g_CamFront;
vec4 g_SunDir;
vec4 g_SunColor;
vec4 g_TimeParams;
vec4 g_ResolutionParams;
vec4 g_CamAxisRight;
vec4 g_FogColor_Distance;
vec4 g_ShadowVP_Row0;
vec4 g_ShadowVP_Row1;
vec4 g_ShadowVP_Row2;
vec4 g_ShadowVP_Row3;
};
vec4 ComputeFogFactor(vec3 WorldPos)
{
vec4 FogData;
vec3 vEye = WorldPos - g_CamPos.xyz;
vec3 nEye = normalize(vEye);
FogData.w = exp(-dot(vEye, vEye) * g_FogColor_Distance.w * 0.75);
float fog_sun_factor = pow(saturate(dot(nEye, g_SunDir.xyz)), 8.0);
FogData.xyz = mix(vec3(1.0, 1.0, 1.0), vec3(0.6, 0.6, 0.9), nEye.y * 0.5 + 0.5);
FogData.xyz = mix(FogData.xyz, vec3(0.95, 0.87, 0.78), fog_sun_factor);
return FogData;
}
void ApplyFog(inout vec3 Color, vec4 FogData)
{
Color = mix(FogData.xyz, Color, FogData.w);
}
void ApplyLighting(inout mediump vec3 Color, mediump float DiffuseFactor)
{
mediump vec3 DiffuseLight = g_SunColor.xyz * DiffuseFactor;
mediump vec3 AmbientLight = vec3(0.2, 0.35, 0.55) * 0.5;
mediump vec3 Lighting = DiffuseLight + AmbientLight;
#if DEBUG == DEBUG_LIGHTING
Color = Lighting;
#else
Color *= Lighting;
#endif
}
#pragma VARIANT SPECULAR
#pragma VARIANT GLOSSMAP
void ApplySpecular(inout mediump vec3 Color, mediump vec3 EyeVec, mediump vec3 Normal, mediump vec3 SpecularColor, mediump float Shininess, mediump float FresnelAmount)
{
mediump vec3 HalfAngle = normalize(-EyeVec + g_SunDir.xyz);
mediump float v_dot_h = saturate(dot(HalfAngle, -EyeVec));
mediump float n_dot_l = saturate(dot(Normal, g_SunDir.xyz));
mediump float n_dot_h = saturate(dot(Normal, HalfAngle));
mediump float n_dot_v = saturate(dot(-EyeVec, Normal));
mediump float h_dot_l = saturate(dot(g_SunDir.xyz, HalfAngle));
const mediump float roughness_value = 0.25;
mediump float r_sq = roughness_value * roughness_value;
mediump float n_dot_h_sq = n_dot_h * n_dot_h;
mediump float roughness_a = 1.0 / (4.0 * r_sq * n_dot_h_sq * n_dot_h_sq);
mediump float roughness_b = n_dot_h_sq - 1.0;
mediump float roughness_c = r_sq * n_dot_h_sq;
mediump float roughness = saturate(roughness_a * exp(roughness_b / roughness_c));
FresnelAmount = 0.5;
mediump float fresnel_term = pow(1.0 - n_dot_v, 5.0) * (1.0 - FresnelAmount) + FresnelAmount;
mediump float geo_numerator = 2.0 * n_dot_h;
mediump float geo_denominator = 1.0 / v_dot_h;
mediump float geo_term = min(1.0, min(n_dot_v, n_dot_l) * geo_numerator * geo_denominator);
#if SPECULAR || GLOSSMAP
Color += SpecularColor * g_SunColor.xyz * fresnel_term * roughness * n_dot_l * geo_term / (n_dot_v * n_dot_l + 0.0001);
#endif
//Color = vec3(0.025 * 1.0 / (n_dot_v * n_dot_l));
}
layout(location = 0) in vec2 Position;
layout(location = 1) in vec4 LODWeights;
layout(location = 0) out vec2 TexCoord;
layout(location = 1) out vec3 EyeVec;
layout(std140, binding = 2) uniform GlobalGround
{
vec4 GroundScale;
vec4 GroundPosition;
vec4 InvGroundSize_PatchScale;
};
struct PatchData
{
vec4 Position;
vec4 LODs;
};
layout(std140, binding = 0) uniform PerPatch
{
PatchData Patches[256];
};
layout(binding = 0) uniform sampler2D TexHeightmap;
layout(binding = 1) uniform sampler2D TexLOD;
vec2 lod_factor(vec2 uv)
{
float level = textureLod(TexLOD, uv, 0.0).x * (255.0 / 32.0);
float floor_level = floor(level);
float fract_level = level - floor_level;
return vec2(floor_level, fract_level);
}
vec2 warp_position()
{
float vlod = dot(LODWeights, Patches[gl_InstanceID].LODs);
vlod = mix(vlod, Patches[gl_InstanceID].Position.w, all(equal(LODWeights, vec4(0.0))));
#ifdef DEBUG_LOD_HEIGHT
LODFactor = vec4(vlod);
#endif
float floor_lod = floor(vlod);
float fract_lod = vlod - floor_lod;
uint ufloor_lod = uint(floor_lod);
#ifdef DEBUG_LOD_HEIGHT
LODFactor = vec4(fract_lod);
#endif
uvec2 uPosition = uvec2(Position);
uvec2 mask = (uvec2(1u) << uvec2(ufloor_lod, ufloor_lod + 1u)) - 1u;
//uvec2 rounding = mix(uvec2(0u), mask, lessThan(uPosition, uvec2(32u)));
uvec2 rounding = uvec2(
uPosition.x < 32u ? mask.x : 0u,
uPosition.y < 32u ? mask.y : 0u);
vec4 lower_upper_snapped = vec4((uPosition + rounding).xyxy & (~mask).xxyy);
return mix(lower_upper_snapped.xy, lower_upper_snapped.zw, fract_lod);
}
void main()
{
vec2 PatchPos = Patches[gl_InstanceID].Position.xz * InvGroundSize_PatchScale.zw;
vec2 WarpedPos = warp_position();
vec2 VertexPos = PatchPos + WarpedPos;
vec2 NormalizedPos = VertexPos * InvGroundSize_PatchScale.xy;
vec2 lod = lod_factor(NormalizedPos);
vec2 Offset = exp2(lod.x) * InvGroundSize_PatchScale.xy;
float Elevation =
mix(textureLod(TexHeightmap, NormalizedPos + 0.5 * Offset, lod.x).x,
textureLod(TexHeightmap, NormalizedPos + 1.0 * Offset, lod.x + 1.0).x,
lod.y);
vec3 WorldPos = vec3(NormalizedPos.x, Elevation, NormalizedPos.y);
WorldPos *= GroundScale.xyz;
WorldPos += GroundPosition.xyz;
EyeVec = WorldPos - g_CamPos.xyz;
TexCoord = NormalizedPos + 0.5 * InvGroundSize_PatchScale.xy;
gl_Position = WorldPos.x * g_ViewProj_Row0 + WorldPos.y * g_ViewProj_Row1 + WorldPos.z * g_ViewProj_Row2 + g_ViewProj_Row3;
}

194
shaders/vert/ocean.vert Normal file
View File

@ -0,0 +1,194 @@
#version 310 es
#define YFLIP 0
#define SPECULAR 0
#define GLOSSMAP 0
#define DEBUG_NONE 0
#define DEBUG_DIFFUSE 1
#define DEBUG_SPECULAR 2
#define DEBUG_LIGHTING 3
#define DEBUG_FOG 4
#define DEBUG DEBUG_NONE
#define FORWARD 0
#define DEFERRED 1
#define DEFERRED_VTEX 2
float saturate(float x) { return clamp(x, 0.0, 1.0); }
layout(std140, binding = 0) uniform GlobalVSData
{
vec4 g_ViewProj_Row0;
vec4 g_ViewProj_Row1;
vec4 g_ViewProj_Row2;
vec4 g_ViewProj_Row3;
vec4 g_CamPos;
vec4 g_CamRight;
vec4 g_CamUp;
vec4 g_CamFront;
vec4 g_SunDir;
vec4 g_SunColor;
vec4 g_TimeParams;
vec4 g_ResolutionParams;
vec4 g_CamAxisRight;
vec4 g_FogColor_Distance;
vec4 g_ShadowVP_Row0;
vec4 g_ShadowVP_Row1;
vec4 g_ShadowVP_Row2;
vec4 g_ShadowVP_Row3;
};
vec4 ComputeFogFactor(vec3 WorldPos)
{
vec4 FogData;
vec3 vEye = WorldPos - g_CamPos.xyz;
vec3 nEye = normalize(vEye);
FogData.w = exp(-dot(vEye, vEye) * g_FogColor_Distance.w * 0.75);
float fog_sun_factor = pow(saturate(dot(nEye, g_SunDir.xyz)), 8.0);
FogData.xyz = mix(vec3(1.0, 1.0, 1.0), vec3(0.6, 0.6, 0.9), nEye.y * 0.5 + 0.5);
FogData.xyz = mix(FogData.xyz, vec3(0.95, 0.87, 0.78), fog_sun_factor);
return FogData;
}
void ApplyFog(inout vec3 Color, vec4 FogData)
{
Color = mix(FogData.xyz, Color, FogData.w);
}
void ApplyLighting(inout mediump vec3 Color, mediump float DiffuseFactor)
{
mediump vec3 DiffuseLight = g_SunColor.xyz * DiffuseFactor;
mediump vec3 AmbientLight = vec3(0.2, 0.35, 0.55) * 0.5;
mediump vec3 Lighting = DiffuseLight + AmbientLight;
#if DEBUG == DEBUG_LIGHTING
Color = Lighting;
#else
Color *= Lighting;
#endif
}
void ApplySpecular(inout mediump vec3 Color, mediump vec3 EyeVec, mediump vec3 Normal, mediump vec3 SpecularColor, mediump float Shininess, mediump float FresnelAmount)
{
mediump vec3 HalfAngle = normalize(-EyeVec + g_SunDir.xyz);
mediump float v_dot_h = saturate(dot(HalfAngle, -EyeVec));
mediump float n_dot_l = saturate(dot(Normal, g_SunDir.xyz));
mediump float n_dot_h = saturate(dot(Normal, HalfAngle));
mediump float n_dot_v = saturate(dot(-EyeVec, Normal));
mediump float h_dot_l = saturate(dot(g_SunDir.xyz, HalfAngle));
const mediump float roughness_value = 0.25;
mediump float r_sq = roughness_value * roughness_value;
mediump float n_dot_h_sq = n_dot_h * n_dot_h;
mediump float roughness_a = 1.0 / (4.0 * r_sq * n_dot_h_sq * n_dot_h_sq);
mediump float roughness_b = n_dot_h_sq - 1.0;
mediump float roughness_c = r_sq * n_dot_h_sq;
mediump float roughness = saturate(roughness_a * exp(roughness_b / roughness_c));
FresnelAmount = 0.5;
mediump float fresnel_term = pow(1.0 - n_dot_v, 5.0) * (1.0 - FresnelAmount) + FresnelAmount;
mediump float geo_numerator = 2.0 * n_dot_h;
mediump float geo_denominator = 1.0 / v_dot_h;
mediump float geo_term = min(1.0, min(n_dot_v, n_dot_l) * geo_numerator * geo_denominator);
#if SPECULAR || GLOSSMAP
Color += SpecularColor * g_SunColor.xyz * fresnel_term * roughness * n_dot_l * geo_term / (n_dot_v * n_dot_l + 0.0001);
#endif
//Color = vec3(0.025 * 1.0 / (n_dot_v * n_dot_l));
}
precision highp int;
layout(binding = 0) uniform mediump sampler2D TexDisplacement;
layout(binding = 1) uniform mediump sampler2D TexLOD;
layout(location = 0) in vec4 Position;
layout(location = 1) in vec4 LODWeights;
layout(location = 0) out highp vec3 EyeVec;
layout(location = 1) out highp vec4 TexCoord;
layout(std140, binding = 4) uniform GlobalOcean
{
vec4 OceanScale;
vec4 OceanPosition;
vec4 InvOceanSize_PatchScale;
vec4 NormalTexCoordScale;
};
struct PatchData
{
vec4 Position;
vec4 LODs;
};
layout(std140, binding = 0) uniform Offsets
{
PatchData Patches[256];
};
vec2 lod_factor(vec2 uv)
{
float level = textureLod(TexLOD, uv, 0.0).x * (255.0 / 32.0);
float floor_level = floor(level);
float fract_level = level - floor_level;
return vec2(floor_level, fract_level);
}
vec2 warp_position()
{
float vlod = dot(LODWeights, Patches[gl_InstanceID].LODs);
vlod = mix(vlod, Patches[gl_InstanceID].Position.w, all(equal(LODWeights, vec4(0.0))));
float floor_lod = floor(vlod);
float fract_lod = vlod - floor_lod;
uint ufloor_lod = uint(floor_lod);
uvec4 uPosition = uvec4(Position);
uvec2 mask = (uvec2(1u) << uvec2(ufloor_lod, ufloor_lod + 1u)) - 1u;
uvec4 rounding;
rounding.x = uPosition.x < 32u ? mask.x : 0u;
rounding.y = uPosition.y < 32u ? mask.x : 0u;
rounding.z = uPosition.x < 32u ? mask.y : 0u;
rounding.w = uPosition.y < 32u ? mask.y : 0u;
//rounding = uPosition.xyxy * mask.xxyy;
vec4 lower_upper_snapped = vec4((uPosition.xyxy + rounding) & (~mask).xxyy);
return mix(lower_upper_snapped.xy, lower_upper_snapped.zw, fract_lod);
}
void main()
{
vec2 PatchPos = Patches[gl_InstanceID].Position.xz * InvOceanSize_PatchScale.zw;
vec2 WarpedPos = warp_position();
vec2 VertexPos = PatchPos + WarpedPos;
vec2 NormalizedPos = VertexPos * InvOceanSize_PatchScale.xy;
vec2 NormalizedTex = NormalizedPos * NormalTexCoordScale.zw;
vec2 lod = lod_factor(NormalizedPos);
vec2 Offset = exp2(lod.x) * InvOceanSize_PatchScale.xy * NormalTexCoordScale.zw;
vec3 Displacement =
mix(textureLod(TexDisplacement, NormalizedTex + 0.5 * Offset, lod.x).yxz,
textureLod(TexDisplacement, NormalizedTex + 1.0 * Offset, lod.x + 1.0).yxz,
lod.y);
vec3 WorldPos = vec3(NormalizedPos.x, 0.0, NormalizedPos.y) + Displacement;
WorldPos *= OceanScale.xyz;
WorldPos += OceanPosition.xyz;
EyeVec = WorldPos - g_CamPos.xyz;
TexCoord = vec4(NormalizedTex, NormalizedTex * NormalTexCoordScale.xy) + 0.5 * InvOceanSize_PatchScale.xyxy * NormalTexCoordScale.zwzw;
gl_Position = WorldPos.x * g_ViewProj_Row0 + WorldPos.y * g_ViewProj_Row1 + WorldPos.z * g_ViewProj_Row2 + g_ViewProj_Row3;
#if YFLIP
gl_Position *= vec4(1.0, -1.0, 1.0, 1.0);
#endif
}

10
shaders/vert/texture_buffer.vert Normal file
View File

@ -0,0 +1,10 @@
#version 310 es
#extension GL_OES_texture_buffer : require
layout(binding = 4) uniform highp samplerBuffer uSamp;
layout(rgba32f, binding = 5) uniform readonly highp imageBuffer uSampo;
void main()
{
gl_Position = texelFetch(uSamp, 10) + imageLoad(uSampo, 100);
}

16
shaders/vert/ubo.vert Normal file
View File

@ -0,0 +1,16 @@
#version 310 es
layout(binding = 0, std140) uniform UBO
{
mat4 mvp;
};
in vec4 aVertex;
in vec3 aNormal;
out vec3 vNormal;
void main()
{
gl_Position = mvp * aVertex;
vNormal = aNormal;
}

635
spir2common.hpp Normal file
View File

@ -0,0 +1,635 @@
/*
* Copyright 2015-2016 ARM Limited
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
* You may obtain a copy of the License at
*
* http://www.apache.org/licenses/LICENSE-2.0
*
* Unless required by applicable law or agreed to in writing, software
* distributed under the License is distributed on an "AS IS" BASIS,
* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
* See the License for the specific language governing permissions and
* limitations under the License.
*/
#ifndef SPIR2COMMON_HPP
#define SPIR2COMMON_HPP
#include <sstream>
#include <stdio.h>
namespace spir2cross
{
class CompilerError : public std::runtime_error
{
public:
CompilerError(const std::string &str) : std::runtime_error(str) {}
};
namespace inner
{
template<typename T>
void join_helper(std::ostringstream &stream, T&& t)
{
stream << std::forward<T>(t);
}
template<typename T, typename... Ts>
void join_helper(std::ostringstream &stream, T&& t, Ts&&... ts)
{
stream << std::forward<T>(t);
join_helper(stream, std::forward<Ts>(ts)...);
}
}
// Helper template to avoid lots of nasty string temporary munging.
template<typename... Ts>
std::string join(Ts&&... ts)
{
std::ostringstream stream;
inner::join_helper(stream, std::forward<Ts>(ts)...);
return stream.str();
}
inline std::string merge(const std::vector<std::string> &list)
{
std::string s;
for (auto &elem : list)
{
s += elem;
if (&elem != &list.back())
s += ", ";
}
return s;
}
template<typename T>
inline std::string convert_to_string(T&& t)
{
return std::to_string(std::forward<T>(t));
}
inline std::string convert_to_string(float t)
{
// std::to_string for floating point values is broken.
// Fallback to something more sane.
char buf[64];
sprintf(buf, "%.32g", t);
return buf;
}
inline std::string convert_to_string(double t)
{
// std::to_string for floating point values is broken.
// Fallback to something more sane.
char buf[64];
sprintf(buf, "%.32g", t);
return buf;
}
struct Instruction
{
Instruction(const std::vector<uint32_t> &spirv, uint32_t &index);
uint16_t op;
uint16_t count;
uint32_t offset;
uint32_t length;
};
// Helper for Variant interface.
struct IVariant
{
virtual ~IVariant() = default;
uint32_t self = 0;
};
enum Types
{
TypeNone,
TypeType,
TypeVariable,
TypeConstant,
TypeFunction,
TypeFunctionPrototype,
TypePointer,
TypeBlock,
TypeExtension,
TypeExpression,
TypeUndef
};
struct SPIRUndef : IVariant
{
enum { type = TypeUndef };
SPIRUndef(uint32_t basetype) : basetype(basetype) {}
uint32_t basetype;
};
struct SPIRType : IVariant
{
enum { type = TypeType };
enum BaseType
{
Unknown,
Void,
Bool,
Int,
UInt,
AtomicCounter,
Float,
Struct,
Image,
SampledImage,
Sampler
};
// Scalar/vector/matrix support.
BaseType basetype = Unknown;
uint32_t width = 0;
uint32_t vecsize = 1;
uint32_t columns = 1;
// Arrays, suport array of arrays by having a vector of array sizes.
std::vector<uint32_t> array;
// Pointers
bool pointer = false;
spv::StorageClass storage = spv::StorageClassGeneric;
std::vector<uint32_t> member_types;
struct Image
{
uint32_t type;
spv::Dim dim;
bool depth;
bool arrayed;
bool ms;
uint32_t sampled;
spv::ImageFormat format;
} image;
};
struct SPIRExtension : IVariant
{
enum { type = TypeExtension };
enum Extension
{
GLSL
};
SPIRExtension(Extension ext)
: ext(ext) {}
Extension ext;
};
struct SPIRExpression : IVariant
{
enum { type = TypeExpression };
// Only created by the backend target to avoid creating tons of temporaries.
SPIRExpression(std::string expr, uint32_t expression_type, bool immutable)
: expression(move(expr)), expression_type(expression_type), immutable(immutable) {}
// If non-zero, prepend expression with to_expression(base_expression).
// Used in amortizing multiple calls to to_expression()
// where in certain cases that would quickly force a temporary when not needed.
uint32_t base_expression = 0;
std::string expression;
uint32_t expression_type = 0;
// If this expression is a forwarded load,
// allow us to reference the original variable.
uint32_t loaded_from = 0;
// If this expression will never change, we can avoid lots of temporaries
// in high level source.
bool immutable = false;
// If this expression has been used while invalidated.
bool used_while_invalidated = false;
// A list of a variables for which this expression was invalidated by.
std::vector<uint32_t> invalidated_by;
};
struct SPIRFunctionPrototype : IVariant
{
enum { type = TypeFunctionPrototype };
SPIRFunctionPrototype(uint32_t return_type)
: return_type(return_type) {}
uint32_t return_type;
std::vector<uint32_t> parameter_types;
};
struct SPIRBlock : IVariant
{
enum { type = TypeBlock };
enum Terminator
{
Unknown,
Direct, // Emit next block directly without a particular condition.
Select, // Block ends with an if/else block.
MultiSelect, // Block ends with switch statement.
Loop, // Block ends with a loop.
Return, // Block ends with return.
Unreachable, // Noop
Kill // Discard
};
enum Merge
{
MergeNone,
MergeLoop,
MergeSelection
};
enum Method
{
MergeToSelectForLoop,
MergeToDirectForLoop
};
enum ContinueBlockType
{
ContinueNone,
// Continue block is branchless and has at least one instruction.
ForLoop,
// Noop continue block.
WhileLoop,
// Continue block is conditional.
DoWhileLoop,
// Highly unlikely that anything will use this,
// since it is really awkward/impossible to express in GLSL.
ComplexLoop
};
enum { NoDominator = 0xffffffffu };
Terminator terminator = Unknown;
Merge merge = MergeNone;
uint32_t next_block = 0;
uint32_t merge_block = 0;
uint32_t continue_block = 0;
uint32_t return_value = 0; // If 0, return nothing (void).
uint32_t condition = 0;
uint32_t true_block = 0;
uint32_t false_block = 0;
uint32_t default_block = 0;
std::vector<Instruction> ops;
struct Phi
{
uint32_t local_variable; // flush local variable ...
uint32_t parent; // If we're in from_block and want to branch into this block ...
uint32_t function_variable; // to this function-global "phi" variable first.
};
// Before entering this block flush out local variables to magical "phi" variables.
std::vector<Phi> phi_variables;
// Declare these temporaries before beginning the block.
// Used for handling complex continue blocks which have side effects.
std::vector<std::pair<uint32_t, uint32_t>> declare_temporary;
struct Case
{
uint32_t value;
uint32_t block;
};
std::vector<Case> cases;
// If we have tried to optimize code for this block but failed,
// keep track of this.
bool disable_block_optimization = false;
// If the continue block is complex, fallback to "dumb" for loops.
bool complex_continue = false;
// The dominating block which this block might be within.
// Used in continue; blocks to determine if we really need to write continue.
uint32_t loop_dominator = 0;
};
struct SPIRFunction : IVariant
{
enum { type = TypeFunction };
SPIRFunction(uint32_t return_type, uint32_t function_type)
: return_type(return_type), function_type(function_type)
{}
struct Parameter
{
uint32_t type;
uint32_t id;
uint32_t read_count;
uint32_t write_count;
};
uint32_t return_type;
uint32_t function_type;
std::vector<Parameter> arguments;
std::vector<uint32_t> local_variables;
uint32_t entry_block = 0;
std::vector<uint32_t> blocks;
void add_local_variable(uint32_t id)
{
local_variables.push_back(id);
}
void add_parameter(uint32_t type, uint32_t id)
{
// Arguments are read-only until proven otherwise.
arguments.push_back({ type, id, 0u, 0u });
}
bool active = false;
bool flush_undeclared = true;
};
struct SPIRVariable : IVariant
{
enum { type = TypeVariable };
SPIRVariable() = default;
SPIRVariable(uint32_t basetype, spv::StorageClass storage, uint32_t initializer = 0)
: basetype(basetype), storage(storage), initializer(initializer)
{}
uint32_t basetype = 0;
spv::StorageClass storage = spv::StorageClassGeneric;
uint32_t decoration = 0;
uint32_t initializer = 0;
std::vector<uint32_t> dereference_chain;
bool compat_builtin = false;
// If a variable is shadowed, we only statically assign to it
// and never actually emit a statement for it.
// When we read the variable as an expression, just forward
// shadowed_id as the expression.
bool statically_assigned = false;
uint32_t static_expression = 0;
// Temporaries which can remain forwarded as long as this variable is not modified.
std::vector<uint32_t> dependees;
bool forwardable = true;
bool deferred_declaration = false;
bool phi_variable = false;
bool remapped_variable = false;
SPIRFunction::Parameter *parameter = nullptr;
};
struct SPIRConstant : IVariant
{
enum { type = TypeConstant };
union Constant
{
uint32_t u32;
int32_t i32;
float f32;
};
struct ConstantVector
{
Constant r[4];
uint32_t vecsize;
};
struct ConstantMatrix
{
ConstantVector c[4];
uint32_t columns;
};
inline uint32_t scalar(uint32_t col = 0, uint32_t row = 0) const
{
return m.c[col].r[row].u32;
}
inline float scalar_f32(uint32_t col = 0, uint32_t row = 0) const
{
return m.c[col].r[row].f32;
}
inline int scalar_i32(uint32_t col = 0, uint32_t row = 0) const
{
return m.c[col].r[row].i32;
}
inline const ConstantVector& vector() const { return m.c[0]; }
inline uint32_t vector_size() const { return m.c[0].vecsize; }
inline uint32_t columns() const { return m.columns; }
SPIRConstant(uint32_t constant_type, const uint32_t *elements, uint32_t num_elements) :
constant_type(constant_type)
{
subconstants.insert(end(subconstants), elements, elements + num_elements);
}
SPIRConstant(uint32_t constant_type, uint32_t v0) :
constant_type(constant_type)
{
m.c[0].r[0].u32 = v0;
m.c[0].vecsize = 1;
m.columns = 1;
}
SPIRConstant(uint32_t constant_type, uint32_t v0, uint32_t v1) :
constant_type(constant_type)
{
m.c[0].r[0].u32 = v0;
m.c[0].r[1].u32 = v1;
m.c[0].vecsize = 2;
m.columns = 1;
}
SPIRConstant(uint32_t constant_type, uint32_t v0, uint32_t v1, uint32_t v2) :
constant_type(constant_type)
{
m.c[0].r[0].u32 = v0;
m.c[0].r[1].u32 = v1;
m.c[0].r[2].u32 = v2;
m.c[0].vecsize = 3;
m.columns = 1;
}
SPIRConstant(uint32_t constant_type, uint32_t v0, uint32_t v1, uint32_t v2, uint32_t v3) :
constant_type(constant_type)
{
m.c[0].r[0].u32 = v0;
m.c[0].r[1].u32 = v1;
m.c[0].r[2].u32 = v2;
m.c[0].r[3].u32 = v3;
m.c[0].vecsize = 4;
m.columns = 1;
}
SPIRConstant(uint32_t constant_type,
const ConstantVector &vec0) :
constant_type(constant_type)
{
m.columns = 1;
m.c[0] = vec0;
}
SPIRConstant(uint32_t constant_type,
const ConstantVector &vec0, const ConstantVector &vec1) :
constant_type(constant_type)
{
m.columns = 2;
m.c[0] = vec0;
m.c[1] = vec1;
}
SPIRConstant(uint32_t constant_type,
const ConstantVector &vec0, const ConstantVector &vec1,
const ConstantVector &vec2) :
constant_type(constant_type)
{
m.columns = 3;
m.c[0] = vec0;
m.c[1] = vec1;
m.c[2] = vec2;
}
SPIRConstant(uint32_t constant_type,
const ConstantVector &vec0, const ConstantVector &vec1,
const ConstantVector &vec2, const ConstantVector &vec3) :
constant_type(constant_type)
{
m.columns = 4;
m.c[0] = vec0;
m.c[1] = vec1;
m.c[2] = vec2;
m.c[3] = vec3;
}
uint32_t constant_type;
ConstantMatrix m;
bool specialization = false; // If the constant is a specialization constant.
// For composites which are constant arrays, etc.
std::vector<uint32_t> subconstants;
};
class Variant
{
public:
// MSVC 2013 workaround, we shouldn't need these constructors.
Variant() = default;
Variant(Variant&& other) { *this = std::move(other); }
Variant& operator=(Variant&& other)
{
if (this != &other)
{
holder = move(other.holder);
type = other.type;
other.type = TypeNone;
}
return *this;
}
void set(std::unique_ptr<IVariant> val, uint32_t type)
{
holder = std::move(val);
if (this->type != TypeNone && this->type != type)
throw CompilerError("Overwriting a variant with new type.");
this->type = type;
}
template<typename T>
T& get()
{
if (!holder)
throw CompilerError("nullptr");
if (T::type != type)
throw CompilerError("Bad cast");
return *static_cast<T*>(holder.get());
}
template<typename T>
const T& get() const
{
if (!holder)
throw CompilerError("nullptr");
if (T::type != type)
throw CompilerError("Bad cast");
return *static_cast<const T*>(holder.get());
}
uint32_t get_type() const { return type; }
bool empty() const { return !holder; }
void reset() { holder.reset(); type = TypeNone; }
private:
std::unique_ptr<IVariant> holder;
uint32_t type = TypeNone;
};
template<typename T>
T& variant_get(Variant &var)
{
return var.get<T>();
}
template<typename T>
const T& variant_get(const Variant &var)
{
return var.get<T>();
}
template<typename T, typename... P>
T& variant_set(Variant &var, P&&... args)
{
auto uptr = std::unique_ptr<T>(new T(std::forward<P>(args)...));
auto ptr = uptr.get();
var.set(std::move(uptr), T::type);
return *ptr;
}
struct Meta
{
struct Decoration
{
std::string alias;
uint64_t decoration_flags = 0;
spv::BuiltIn builtin_type;
uint32_t location = 0;
uint32_t set = 0;
uint32_t binding = 0;
uint32_t offset = 0;
uint32_t array_stride = 0;
bool builtin = false;
};
Decoration decoration;
std::vector<Decoration> members;
};
}
#endif

436
spir2cpp.cpp Normal file
View File

@ -0,0 +1,436 @@
/*
* Copyright 2015-2016 ARM Limited
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
* You may obtain a copy of the License at
*
* http://www.apache.org/licenses/LICENSE-2.0
*
* Unless required by applicable law or agreed to in writing, software
* distributed under the License is distributed on an "AS IS" BASIS,
* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
* See the License for the specific language governing permissions and
* limitations under the License.
*/
#include "spir2cpp.hpp"
using namespace spv;
using namespace spir2cross;
using namespace std;
void CompilerCPP::emit_buffer_block(const SPIRVariable &var)
{
auto &type = get<SPIRType>(var.basetype);
auto instance_name = to_name(var.self);
uint32_t set = meta[var.self].decoration.set;
uint32_t binding = meta[var.self].decoration.binding;
emit_struct(type);
statement("internal::Resource<", type_to_glsl(type), type_to_array_glsl(type), "> ", instance_name, "__;");
statement_no_indent("#define ", instance_name, " __res->", instance_name, "__.get()");
resource_registrations.push_back(join("s.register_resource(", instance_name, "__", ", ", set, ", ", binding, ");"));
statement("");
}
void CompilerCPP::emit_interface_block(const SPIRVariable &var)
{
auto &type = get<SPIRType>(var.basetype);
const char *qual = var.storage == StorageClassInput ? "StageInput" : "StageOutput";
const char *lowerqual = var.storage == StorageClassInput ? "stage_input" : "stage_output";
auto instance_name = to_name(var.self);
uint32_t location = meta[var.self].decoration.location;
auto flags = meta[type.self].decoration.decoration_flags;
if (flags & (1ull << DecorationBlock))
emit_struct(type);
statement("internal::", qual, "<", type_to_glsl(type), type_to_array_glsl(type), "> ", instance_name, "__;");
statement_no_indent("#define ", instance_name, " __res->", instance_name, "__.get()");
resource_registrations.push_back(join("s.register_", lowerqual, "(", instance_name, "__", ", ", location, ");"));
statement("");
}
void CompilerCPP::emit_shared(const SPIRVariable &var)
{
auto instance_name = to_name(var.self);
statement(variable_decl(var), ";");
statement_no_indent("#define ", instance_name, " __res->", instance_name);
}
void CompilerCPP::emit_uniform(const SPIRVariable &var)
{
auto &type = get<SPIRType>(var.basetype);
auto instance_name = to_name(var.self);
uint32_t set = meta[var.self].decoration.set;
uint32_t binding = meta[var.self].decoration.binding;
uint32_t location = meta[var.self].decoration.location;
if (type.basetype == SPIRType::Image || type.basetype == SPIRType::SampledImage || type.basetype == SPIRType::AtomicCounter)
{
statement("internal::Resource<", type_to_glsl(type), type_to_array_glsl(type), "> ", instance_name, "__;");
statement_no_indent("#define ", instance_name, " __res->", instance_name, "__.get()");
resource_registrations.push_back(join("s.register_resource(", instance_name, "__", ", ", set, ", ", binding, ");"));
}
else
{
statement("internal::UniformConstant<", type_to_glsl(type), type_to_array_glsl(type), "> ", instance_name, "__;");
statement_no_indent("#define ", instance_name, " __res->", instance_name, "__.get()");
resource_registrations.push_back(join("s.register_uniform_constant(", instance_name, "__", ", ", location, ");"));
}
statement("");
}
void CompilerCPP::emit_push_constant_block(const SPIRVariable &var)
{
auto &type = get<SPIRType>(var.basetype);
auto &flags = meta[var.self].decoration.decoration_flags;
if ((flags & (1ull << DecorationBinding)) || (flags & (1ull << DecorationDescriptorSet)))
throw CompilerError("Push constant blocks cannot be compiled to GLSL with Binding or Set syntax. "
"Remap to location with reflection API first or disable these decorations.");
emit_struct(type);
auto instance_name = to_name(var.self);
statement("internal::PushConstant<", type_to_glsl(type), type_to_array_glsl(type), "> ", instance_name, ";");
statement_no_indent("#define ", instance_name, " __res->", instance_name, ".get()");
resource_registrations.push_back(join("s.register_push_constant(", instance_name, "__", ");"));
statement("");
}
void CompilerCPP::emit_resources()
{
// Output all basic struct types which are not Block or BufferBlock as these are declared inplace
// when such variables are instantiated.
for (auto &id : ids)
{
if (id.get_type() == TypeType)
{
auto &type = id.get<SPIRType>();
if (type.basetype == SPIRType::Struct &&
type.array.empty() &&
!type.pointer &&
(meta[type.self].decoration.decoration_flags & ((1ull << DecorationBlock) | (1ull << DecorationBufferBlock))) == 0)
{
emit_struct(type);
}
}
}
statement("struct Resources : ", resource_type);
begin_scope();
// Output UBOs and SSBOs
for (auto &id : ids)
{
if (id.get_type() == TypeVariable)
{
auto &var = id.get<SPIRVariable>();
auto &type = get<SPIRType>(var.basetype);
if (type.pointer && type.storage == StorageClassUniform &&
!is_builtin_variable(var) &&
(meta[type.self].decoration.decoration_flags & ((1ull << DecorationBlock) | (1ull << DecorationBufferBlock))))
{
emit_buffer_block(var);
}
}
}
// Output push constant blocks
for (auto &id : ids)
{
if (id.get_type() == TypeVariable)
{
auto &var = id.get<SPIRVariable>();
auto &type = get<SPIRType>(var.basetype);
if (type.pointer && type.storage == StorageClassPushConstant)
emit_push_constant_block(var);
}
}
// Output in/out interfaces.
for (auto &id : ids)
{
if (id.get_type() == TypeVariable)
{
auto &var = id.get<SPIRVariable>();
auto &type = get<SPIRType>(var.basetype);
if (!is_builtin_variable(var) &&
!var.remapped_variable &&
type.pointer &&
(var.storage == StorageClassInput || var.storage == StorageClassOutput))
{
emit_interface_block(var);
}
}
}
// Output Uniform Constants (values, samplers, images, etc).
for (auto &id : ids)
{
if (id.get_type() == TypeVariable)
{
auto &var = id.get<SPIRVariable>();
auto &type = get<SPIRType>(var.basetype);
if (!is_builtin_variable(var) && !var.remapped_variable && type.pointer &&
(type.storage == StorageClassUniformConstant || type.storage == StorageClassAtomicCounter))
{
emit_uniform(var);
}
}
}
// Global variables.
bool emitted = false;
for (auto global : global_variables)
{
auto &var = get<SPIRVariable>(global);
if (var.storage == StorageClassWorkgroup)
{
emit_shared(var);
emitted = true;
}
}
if (emitted)
statement("");
statement("inline void init(spir2cross_shader& s)");
begin_scope();
statement(resource_type, "::init(s);");
for (auto &reg : resource_registrations)
statement(reg);
end_scope();
resource_registrations.clear();
end_scope_decl();
statement("");
statement("Resources* __res;");
if (execution.model == ExecutionModelGLCompute)
statement("ComputePrivateResources __priv_res;");
statement("");
// Emit regular globals which are allocated per invocation.
emitted = false;
for (auto global : global_variables)
{
auto &var = get<SPIRVariable>(global);
if (var.storage == StorageClassPrivate)
{
if (var.storage == StorageClassWorkgroup)
emit_shared(var);
else
statement(variable_decl(var), ";");
emitted = true;
}
}
if (emitted)
statement("");
}
string CompilerCPP::compile()
{
// Do not deal with ES-isms like precision, older extensions and such.
options.es = false;
options.version = 450;
backend.float_literal_suffix = true;
backend.uint32_t_literal_suffix = true;
backend.basic_int_type = "int32_t";
backend.basic_uint_type = "uint32_t";
backend.swizzle_is_function = true;
backend.shared_is_implied = true;
uint32_t pass_count = 0;
do
{
if (pass_count >= 2)
throw CompilerError("Over 2 compilation loops detected. Must be a bug!");
resource_registrations.clear();
reset();
// Move constructor for this type is broken on GCC 4.9 ...
buffer = unique_ptr<ostringstream>(new ostringstream());
emit_header();
emit_resources();
emit_function(get<SPIRFunction>(execution.entry_point), 0);
pass_count++;
} while (force_recompile);
// Match opening scope of emit_header().
end_scope_decl();
// namespace
end_scope();
// Emit C entry points
emit_c_linkage();
return buffer->str();
}
void CompilerCPP::emit_c_linkage()
{
statement("");
statement("spir2cross_shader_t* spir2cross_construct(void)");
begin_scope();
statement("return new ", impl_type, "();");
end_scope();
statement("");
statement("void spir2cross_destruct(spir2cross_shader_t *shader)");
begin_scope();
statement("delete static_cast<", impl_type, "*>(shader);");
end_scope();
statement("");
statement("void spir2cross_invoke(spir2cross_shader_t *shader)");
begin_scope();
statement("static_cast<", impl_type, "*>(shader)->invoke();");
end_scope();
statement("");
statement("static const struct spir2cross_interface vtable =");
begin_scope();
statement("spir2cross_construct,");
statement("spir2cross_destruct,");
statement("spir2cross_invoke,");
end_scope_decl();
statement("");
statement("const struct spir2cross_interface* spir2cross_get_interface(void)");
begin_scope();
statement("return &vtable;");
end_scope();
}
void CompilerCPP::emit_function_prototype(SPIRFunction &func, uint64_t)
{
local_variables.clear();
string decl;
auto &type = get<SPIRType>(func.return_type);
decl += "inline ";
decl += type_to_glsl(type);
decl += " ";
if (func.self == execution.entry_point)
{
decl += "main";
processing_entry_point = true;
}
else
decl += to_name(func.self);
decl += "(";
for (auto &arg : func.arguments)
{
add_local_variable(arg.id);
decl += argument_decl(arg);
if (&arg != &func.arguments.back())
decl += ", ";
// Hold a pointer to the parameter so we can invalidate the readonly field if needed.
auto *var = maybe_get<SPIRVariable>(arg.id);
if (var)
var->parameter = &arg;
}
decl += ")";
statement(decl);
}
string CompilerCPP::argument_decl(const SPIRFunction::Parameter &arg)
{
auto &type = expression_type(arg.id);
bool constref = !type.pointer || arg.write_count == 0;
auto &var = get<SPIRVariable>(arg.id);
return join(constref ? "const " : "",
type_to_glsl(type), "& ", to_name(var.self), type_to_array_glsl(type));
}
void CompilerCPP::emit_header()
{
statement("// This C++ shader is autogenerated by spir2cross.");
statement("#include \"spir2cross/internal_interface.hpp\"");
statement("#include \"spir2cross/external_interface.h\"");
statement("#include <stdint.h>");
statement("");
statement("using namespace spir2cross;");
statement("using namespace glm;");
statement("");
statement("namespace Impl");
begin_scope();
switch (execution.model)
{
case ExecutionModelGeometry:
case ExecutionModelTessellationControl:
case ExecutionModelTessellationEvaluation:
case ExecutionModelGLCompute:
case ExecutionModelFragment:
case ExecutionModelVertex:
statement("struct Shader");
begin_scope();
break;
default:
throw CompilerError("Unsupported execution model.");
}
switch (execution.model)
{
case ExecutionModelGeometry:
impl_type = "GeometryShader<Impl::Shader, Impl::Shader::Resources>";
resource_type = "GeometryResources";
break;
case ExecutionModelVertex:
impl_type = "VertexShader<Impl::Shader, Impl::Shader::Resources>";
resource_type = "VertexResources";
break;
case ExecutionModelFragment:
impl_type = "FragmentShader<Impl::Shader, Impl::Shader::Resources>";
resource_type = "FragmentResources";
break;
case ExecutionModelGLCompute:
impl_type = join("ComputeShader<Impl::Shader, Impl::Shader::Resources, ",
execution.workgroup_size.x, ", ",
execution.workgroup_size.y, ", ",
execution.workgroup_size.z, ">");
resource_type = "ComputeResources";
break;
case ExecutionModelTessellationControl:
impl_type = "TessControlShader<Impl::Shader, Impl::Shader::Resources>";
resource_type = "TessControlResources";
break;
case ExecutionModelTessellationEvaluation:
impl_type = "TessEvaluationShader<Impl::Shader, Impl::Shader::Resources>";
resource_type = "TessEvaluationResources";
break;
default:
throw CompilerError("Unsupported execution model.");
}
}

53
spir2cpp.hpp Normal file
View File

@ -0,0 +1,53 @@
/*
* Copyright 2015-2016 ARM Limited
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
* You may obtain a copy of the License at
*
* http://www.apache.org/licenses/LICENSE-2.0
*
* Unless required by applicable law or agreed to in writing, software
* distributed under the License is distributed on an "AS IS" BASIS,
* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
* See the License for the specific language governing permissions and
* limitations under the License.
*/
#ifndef SPIR2CPP
#define SPIR2CPP
#include "spir2glsl.hpp"
#include <vector>
namespace spir2cross
{
class CompilerCPP : public CompilerGLSL
{
public:
CompilerCPP(std::vector<uint32_t> spirv) : CompilerGLSL(move(spirv)) {}
std::string compile() override;
private:
void emit_header() override;
void emit_c_linkage();
void emit_function_prototype(SPIRFunction &func, uint64_t return_flags) override;
void emit_resources();
void emit_buffer_block(const SPIRVariable &type);
void emit_push_constant_block(const SPIRVariable &var);
void emit_interface_block(const SPIRVariable &type);
void emit_block_chain(SPIRBlock &block);
void emit_uniform(const SPIRVariable &var);
void emit_shared(const SPIRVariable &var);
std::string argument_decl(const SPIRFunction::Parameter &arg);
std::vector<std::string> resource_registrations;
std::string impl_type;
std::string resource_type;
uint32_t shared_counter = 0;
};
}
#endif

1802
spir2cross.cpp Normal file
View File

File diff suppressed because it is too large Load Diff

345
spir2cross.hpp Normal file
View File

@ -0,0 +1,345 @@
/*
* Copyright 2015-2016 ARM Limited
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
* You may obtain a copy of the License at
*
* http://www.apache.org/licenses/LICENSE-2.0
*
* Unless required by applicable law or agreed to in writing, software
* distributed under the License is distributed on an "AS IS" BASIS,
* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
* See the License for the specific language governing permissions and
* limitations under the License.
*/
#ifndef SPIR2CROSS_HPP
#define SPIR2CROSS_HPP
#include "spirv.hpp"
#include <vector>
#include <string>
#include <stdexcept>
#include <memory>
#include <utility>
#include <unordered_set>
#include <unordered_map>
#include "spir2common.hpp"
namespace spir2cross
{
struct Resource
{
// Resources are identified with their SPIR-V ID.
// This is the ID of the OpVariable.
uint32_t id;
// The type of the declared resource.
uint32_t type_id;
// The declared name (OpName) of the resource.
// For Buffer blocks, the name actually reflects the externally
// visible Block name.
//
// This name can be retrieved again by using either
// get_name(id) or get_name(type_id) depending if it's a buffer block or not.
//
// This name can be an empty string in which case get_fallback_name(id) can be
// used which obtains a suitable fallback identifier for an ID.
std::string name;
};
struct ShaderResources
{
std::vector<Resource> uniform_buffers;
std::vector<Resource> storage_buffers;
std::vector<Resource> stage_inputs;
std::vector<Resource> stage_outputs;
std::vector<Resource> subpass_inputs;
std::vector<Resource> storage_images;
std::vector<Resource> sampled_images;
std::vector<Resource> atomic_counters;
// There can only be one push constant block,
// but keep the vector in case this restriction is lifted in the future.
std::vector<Resource> push_constant_buffers;
};
struct BufferRange
{
unsigned index;
size_t offset;
size_t range;
};
class Compiler
{
public:
// The constructor takes a buffer of SPIR-V words and parses it.
Compiler(std::vector<uint32_t> ir);
// After parsing, API users can modify the SPIR-V via reflection and call this
// to disassemble the SPIR-V into the desired langauage.
// Sub-classes actually implement this.
virtual std::string compile();
// Gets the identifier (OpName) of an ID. If not defined, an empty string will be returned.
const std::string& get_name(uint32_t id) const;
// Applies a decoration to an ID. Effectively injects OpDecorate.
void set_decoration(uint32_t id, spv::Decoration decoration, uint32_t argument = 0);
// Overrides the identifier OpName of an ID.
// Identifiers beginning with underscores or identifiers which contain double underscores
// are reserved by the implementation.
void set_name(uint32_t id, const std::string& name);
// Gets a bitmask for the decorations which are applied to ID.
// I.e. (1ull << spv::DecorationFoo) | (1ull << spv::DecorationBar)
uint64_t get_decoration_mask(uint32_t id) const;
// Gets the value for decorations which take arguments.
// If decoration doesn't exist or decoration is not recognized,
// 0 will be returned.
uint32_t get_decoration(uint32_t id, spv::Decoration decoration) const;
// Removes the decoration for a an ID.
void unset_decoration(uint32_t id, spv::Decoration decoration);
// Gets the SPIR-V associated with ID.
// Mostly used with Resource::type_id to parse the underlying type of a resource.
const SPIRType& get_type(uint32_t id) const;
// Gets the underlying storage class for an OpVariable.
spv::StorageClass get_storage_class(uint32_t id) const;
// If get_name() is an empty string, get the fallback name which will be used
// instead in the disassembled source.
virtual const std::string get_fallback_name(uint32_t id) const
{
return join("_", id);
}
// Given an OpTypeStruct in ID, obtain the identifier for member number "index".
// This may be an empty string.
const std::string& get_member_name(uint32_t id, uint32_t index) const;
// Given an OpTypeStruct in ID, obtain the OpMemberDecoration for member number "index".
uint32_t get_member_decoration(uint32_t id, uint32_t index, spv::Decoration decoration) const;
// Sets the member identifier for OpTypeStruct ID, member number "index".
void set_member_name(uint32_t id, uint32_t index, const std::string& name);
// Gets the decoration mask for a member of a struct, similar to get_decoration_mask.
uint64_t get_member_decoration_mask(uint32_t id, uint32_t index) const;
// Similar to set_decoration, but for struct members.
void set_member_decoration(uint32_t id, uint32_t index,
spv::Decoration decoration, uint32_t argument = 0);
// Unsets a member decoration, similar to unset_decoration.
void unset_member_decoration(uint32_t id, uint32_t index, spv::Decoration decoration);
// Gets the fallback name for a member, similar to get_fallback_name.
virtual const std::string get_fallback_member_name(uint32_t index) const
{
return join("_", index);
}
// Returns a vector of which members of a struct are potentially in use by a
// SPIR-V shader. The granularity of this analysis is per-member of a struct.
// This can be used for Buffer (UBO), BufferBlock (SSBO) and PushConstant blocks.
// ID is the Resource::id obtained from get_shader_resources().
std::vector<BufferRange> get_active_buffer_ranges(unsigned id) const;
// Returns the effective size of a buffer block.
size_t get_declared_struct_size(const SPIRType &struct_type) const;
// Legacy GLSL compatibility method.
// Takes a variable with a block interface and flattens it into a T array[N]; array instead.
// For this to work, all types in the block must not themselves be composites
// (except vectors and matrices), and all types must be the same.
// The name of the uniform will be the same as the interface block name.
void flatten_interface_block(uint32_t id);
// Query shader resources, use ids with reflection interface to modify or query binding points, etc.
ShaderResources get_shader_resources() const;
protected:
const uint32_t* stream(uint32_t offset) const
{
if (offset > spirv.size())
throw CompilerError("Compiler::stream() out of range.");
return &spirv[offset];
}
std::vector<uint32_t> spirv;
std::vector<Instruction> inst;
std::vector<Variant> ids;
std::vector<Meta> meta;
SPIRFunction *function = nullptr;
SPIRBlock *block = nullptr;
std::vector<uint32_t> global_variables;
std::vector<uint32_t> aliased_variables;
// If our IDs are out of range here as part of opcodes, throw instead of
// undefined behavior.
template<typename T, typename... P>
T& set(uint32_t id, P&&... args)
{
auto &var = variant_set<T>(ids.at(id), std::forward<P>(args)...);
var.self = id;
return var;
}
template<typename T>
T& get(uint32_t id)
{
return variant_get<T>(ids.at(id));
}
template<typename T>
T* maybe_get(uint32_t id)
{
if (ids.at(id).get_type() == T::type)
return &get<T>(id);
else
return nullptr;
}
template<typename T>
const T& get(uint32_t id) const
{
return variant_get<T>(ids.at(id));
}
template<typename T>
const T* maybe_get(uint32_t id) const
{
if (ids.at(id).get_type() == T::type)
return &get<T>(id);
else
return nullptr;
}
struct
{
uint64_t flags = 0;
spv::ExecutionModel model;
uint32_t entry_point = 0;
struct
{
uint32_t x = 0, y = 0, z = 0;
} workgroup_size;
uint32_t invocations = 0;
uint32_t output_vertices = 0;
} execution;
struct
{
uint32_t version = 0;
bool es = false;
bool known = false;
} source;
std::unordered_set<uint32_t> loop_block;
std::unordered_set<uint32_t> continue_block;
std::unordered_set<uint32_t> loop_merge_target;
std::unordered_set<uint32_t> selection_merge_target;
std::unordered_set<uint32_t> multiselect_merge_target;
std::string to_name(uint32_t id);
bool is_builtin_variable(const SPIRVariable &var) const;
bool is_immutable(uint32_t id) const;
bool is_member_builtin(const SPIRType &type, uint32_t index, spv::BuiltIn *builtin) const;
const SPIRType& expression_type(uint32_t id) const;
bool expression_is_lvalue(uint32_t id) const;
bool variable_storage_is_aliased(const SPIRVariable &var);
SPIRVariable* maybe_get_backing_variable(uint32_t chain);
void register_read(uint32_t expr, uint32_t chain, bool forwarded);
void register_write(uint32_t chain);
inline bool is_continue(uint32_t next) const
{
return continue_block.find(next) != end(continue_block);
}
inline bool is_break(uint32_t next) const
{
return loop_merge_target.find(next) != end(loop_merge_target) ||
multiselect_merge_target.find(next) != end(multiselect_merge_target);
}
inline bool is_conditional(uint32_t next) const
{
return selection_merge_target.find(next) != end(selection_merge_target) &&
multiselect_merge_target.find(next) == end(multiselect_merge_target);
}
// Dependency tracking for temporaries read from variables.
void flush_dependees(SPIRVariable &var);
void flush_all_active_variables();
void flush_all_atomic_capable_variables();
void flush_all_aliased_variables();
void register_global_read_dependencies(const SPIRBlock &func, uint32_t id);
void register_global_read_dependencies(const SPIRFunction &func, uint32_t id);
std::unordered_set<uint32_t> invalid_expressions;
void update_name_cache(std::unordered_set<std::string> &cache, std::string &name);
std::unordered_set<std::string> global_struct_cache;
bool function_is_pure(const SPIRFunction &func);
bool block_is_pure(const SPIRBlock &block);
bool block_is_outside_flow_control_from_block(const SPIRBlock &from, const SPIRBlock &to);
bool execution_is_branchless(const SPIRBlock &from, const SPIRBlock &to) const;
bool execution_is_noop(const SPIRBlock &from, const SPIRBlock &to) const;
SPIRBlock::ContinueBlockType continue_block_type(const SPIRBlock &continue_block) const;
bool force_recompile = false;
uint32_t type_struct_member_offset(const SPIRType &type, uint32_t index) const;
uint32_t type_struct_member_array_stride(const SPIRType &type, uint32_t index) const;
bool block_is_loop_candidate(const SPIRBlock &block, SPIRBlock::Method method) const;
private:
void parse();
void parse(const Instruction &i);
// Used internally to implement various traversals for queries.
struct OpcodeHandler
{
virtual ~OpcodeHandler() = default;
// Return true if traversal should continue.
// If false, traversal will end immediately.
virtual bool handle(spv::Op opcode, const uint32_t *args, uint32_t length) = 0;
};
struct BufferAccessHandler : OpcodeHandler
{
BufferAccessHandler(const Compiler &compiler, std::vector<BufferRange> &ranges, unsigned id)
: compiler(compiler), ranges(ranges), id(id) {}
bool handle(spv::Op opcode, const uint32_t *args, uint32_t length) override;
const Compiler &compiler;
std::vector<BufferRange> &ranges;
uint32_t id;
std::unordered_set<uint32_t> seen;
};
bool traverse_all_reachable_opcodes(const SPIRBlock &block, OpcodeHandler &handler) const;
bool traverse_all_reachable_opcodes(const SPIRFunction &block, OpcodeHandler &handler) const;
size_t get_declared_struct_member_size(const SPIRType &struct_type, uint32_t index) const;
};
}
#endif

4225
spir2glsl.cpp Normal file
View File

File diff suppressed because it is too large Load Diff

298
spir2glsl.hpp Normal file
View File

@ -0,0 +1,298 @@
/*
* Copyright 2015-2016 ARM Limited
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
* You may obtain a copy of the License at
*
* http://www.apache.org/licenses/LICENSE-2.0
*
* Unless required by applicable law or agreed to in writing, software
* distributed under the License is distributed on an "AS IS" BASIS,
* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
* See the License for the specific language governing permissions and
* limitations under the License.
*/
#ifndef SPIR2GLSL
#define SPIR2GLSL
#include "spir2cross.hpp"
#include <sstream>
#include <unordered_set>
#include <unordered_map>
#include <utility>
namespace spir2cross
{
enum PlsFormat
{
PlsNone = 0,
PlsR11FG11FB10F,
PlsR32F,
PlsRG16F,
PlsRGB10A2,
PlsRGBA8,
PlsRG16,
PlsRGBA8I,
PlsRG16I,
PlsRGB10A2UI,
PlsRGBA8UI,
PlsRG16UI,
PlsR32UI
};
struct PlsRemap
{
uint32_t id;
PlsFormat format;
};
class CompilerGLSL : public Compiler
{
public:
struct Options
{
uint32_t version = 450;
bool es = false;
bool force_temporary = false;
enum Precision
{
DontCare,
Lowp,
Mediump,
Highp
};
struct
{
// In vertex shaders, rewrite [0, w] depth (Vulkan/D3D style) to [-w, w] depth (GL style).
bool fixup_clipspace = true;
} vertex;
struct
{
// Add precision mediump float in ES targets when emitting GLES source.
// Add precision highp int in ES targets when emitting GLES source.
Precision default_float_precision = Mediump;
Precision default_int_precision = Highp;
} fragment;
};
void remap_pixel_local_storage(std::vector<PlsRemap> inputs, std::vector<PlsRemap> outputs)
{
pls_inputs = std::move(inputs);
pls_outputs = std::move(outputs);
remap_pls_variables();
}
CompilerGLSL(std::vector<uint32_t> spirv) : Compiler(move(spirv))
{
if (source.known)
{
options.es = source.es;
options.version = source.version;
}
}
const Options& get_options() const { return options; }
void set_options(Options &opts) { options = opts; }
std::string compile() override;
protected:
void reset();
void emit_function(SPIRFunction &func, uint64_t return_flags);
// Virtualize methods which need to be overridden by subclass targets like C++ and such.
virtual void emit_function_prototype(SPIRFunction &func, uint64_t return_flags);
virtual void emit_header();
std::unique_ptr<std::ostringstream> buffer;
template<typename T>
inline void statement_inner(T&& t)
{
(*buffer) << std::forward<T>(t);
statement_count++;
}
template<typename T, typename... Ts>
inline void statement_inner(T&& t, Ts&&... ts)
{
(*buffer) << std::forward<T>(t);
statement_count++;
statement_inner(std::forward<Ts>(ts)...);
}
template<typename... Ts>
inline void statement(Ts&&... ts)
{
if (redirect_statement)
redirect_statement->push_back(join(std::forward<Ts>(ts)...));
else
{
for (uint32_t i = 0; i < indent; i++)
(*buffer) << " ";
statement_inner(std::forward<Ts>(ts)...);
(*buffer) << '\n';
}
}
template<typename... Ts>
inline void statement_no_indent(Ts&&... ts)
{
auto old_indent = indent;
indent = 0;
statement(std::forward<Ts>(ts)...);
indent = old_indent;
}
// Used for implementing continue blocks where
// we want to obtain a list of statements we can merge
// on a single line separated by comma.
std::vector<std::string> *redirect_statement = nullptr;
const SPIRBlock *current_continue_block = nullptr;
void begin_scope();
void end_scope();
void end_scope_decl();
void end_scope_decl(const std::string &decl);
Options options;
std::string type_to_glsl(const SPIRType &type);
std::string type_to_array_glsl(const SPIRType &type);
std::string variable_decl(const SPIRVariable &variable);
void add_local_variable(uint32_t id);
std::unordered_set<std::string> local_variables;
bool processing_entry_point = false;
// Can be overriden by subclass backends for trivial things which
// shouldn't need polymorphism.
struct BackendVariations
{
bool float_literal_suffix = false;
bool uint32_t_literal_suffix = true;
const char *basic_int_type = "int";
const char *basic_uint_type = "uint";
bool swizzle_is_function = false;
bool shared_is_implied = false;
} backend;
void emit_struct(const SPIRType &type);
void emit_instruction(const Instruction &instr);
private:
void emit_resources();
void emit_buffer_block(const SPIRVariable &type);
void emit_push_constant_block(const SPIRVariable &var);
void emit_interface_block(const SPIRVariable &type);
void emit_block_chain(SPIRBlock &block);
std::string emit_continue_block(uint32_t continue_block);
bool attempt_emit_loop_header(SPIRBlock &block, SPIRBlock::Method method);
void emit_uniform(const SPIRVariable &var);
void propagate_loop_dominators(const SPIRBlock &block);
void branch(uint32_t from, uint32_t to);
void branch(uint32_t from, uint32_t cond, uint32_t true_block, uint32_t false_block);
void flush_phi(uint32_t from, uint32_t to);
bool flush_phi_required(uint32_t from, uint32_t to);
void flush_variable_declaration(uint32_t id);
void flush_undeclared_variables();
bool should_forward(uint32_t id);
void emit_texture_op(const Instruction &i);
void emit_mix_op(uint32_t result_type, uint32_t id, uint32_t left, uint32_t right, uint32_t lerp);
void emit_glsl_op(uint32_t result_type, uint32_t result_id, uint32_t op, const uint32_t *args, uint32_t count);
void emit_quaternary_func_op(uint32_t result_type, uint32_t result_id, uint32_t op0, uint32_t op1, uint32_t op2, uint32_t op3, const char *op);
void emit_trinary_func_op(uint32_t result_type, uint32_t result_id, uint32_t op0, uint32_t op1, uint32_t op2, const char *op);
void emit_binary_func_op(uint32_t result_type, uint32_t result_id, uint32_t op0, uint32_t op1, const char *op);
void emit_unary_func_op(uint32_t result_type, uint32_t result_id, uint32_t op0, const char *op);
void emit_binary_op(uint32_t result_type, uint32_t result_id, uint32_t op0, uint32_t op1, const char *op);
void emit_unary_op(uint32_t result_type, uint32_t result_id, uint32_t op0, const char *op);
bool expression_is_forwarded(uint32_t id);
SPIRExpression& emit_op(uint32_t result_type, uint32_t result_id, const std::string &rhs, bool forward_rhs, bool extra_parens, bool suppress_usage_tracking = false);
std::string access_chain(uint32_t base, const uint32_t *indices, uint32_t count, bool index_is_literal, bool chain_only = false);
const char* index_to_swizzle(uint32_t index);
std::string remap_swizzle(uint32_t result_type, uint32_t input_components, uint32_t expr);
std::string declare_temporary(uint32_t type, uint32_t id);
std::string to_expression(uint32_t id);
std::string to_member_name(const SPIRType &type, uint32_t index);
std::string type_to_glsl_constructor(const SPIRType &type);
std::string argument_decl(const SPIRFunction::Parameter &arg);
std::string member_decl(const SPIRType &type, const SPIRType &member_type, uint32_t member);
std::string image_type_glsl(const SPIRType &type);
std::string to_qualifiers_glsl(uint32_t id);
const char* to_precision_qualifiers_glsl(uint32_t id);
const char* flags_to_precision_qualifiers_glsl(const SPIRType &type, uint64_t flags);
std::string constant_expression(const SPIRConstant &c);
std::string constant_expression_vector(const SPIRConstant &c, uint32_t vector);
const char* format_to_glsl(spv::ImageFormat format);
std::string layout_for_member(const SPIRType &type, uint32_t index);
uint64_t combined_decoration_for_member(const SPIRType &type, uint32_t index);
std::string layout_for_variable(const SPIRVariable &variable);
bool ssbo_is_std430_packing(const SPIRType &type);
uint32_t type_to_std430_alignment(const SPIRType &type, uint64_t flags);
uint32_t type_to_std430_array_stride(const SPIRType &type, uint64_t flags);
uint32_t type_to_std430_size(const SPIRType &type, uint64_t flags);
std::string bitcast_glsl(uint32_t result_type, uint32_t arg);
std::string bitcast_glsl_op(uint32_t result_type, uint32_t arg);
const char* builtin_to_glsl(spv::BuiltIn builtin);
std::string build_composite_combiner(const uint32_t *elems, uint32_t length);
bool remove_duplicate_swizzle(std::string &op);
bool remove_unity_swizzle(uint32_t base, std::string &op);
// Can modify flags to remote readonly/writeonly if image type
// and force recompile.
bool check_atomic_image(uint32_t id);
void require_extension(const std::string &ext);
void replace_fragment_output(SPIRVariable &var);
void replace_fragment_outputs();
std::string legacy_tex_op(const std::string &op, const SPIRType &imgtype);
uint32_t indent = 0;
void emit_fixup();
std::unordered_set<uint32_t> emitted_functions;
// Usage tracking. If a temporary is used more than once, use the temporary instead to
// avoid AST explosion when SPIRV is generated with pure SSA and doesn't write stuff to variables.
std::unordered_map<uint32_t, uint32_t> expression_usage_counts;
std::unordered_set<uint32_t> forced_temporaries;
std::unordered_set<uint32_t> forwarded_temporaries;
void track_expression_read(uint32_t id);
std::unordered_set<std::string> forced_extensions;
uint32_t statement_count;
inline bool is_legacy() const { return (options.es && options.version < 300) || (!options.es && options.version < 130); }
bool args_will_forward(uint32_t id, const uint32_t *args, uint32_t num_args, bool pure);
void register_call_out_argument(uint32_t id);
void register_impure_function_call();
// GL_EXT_shader_pixel_local_storage support.
std::vector<PlsRemap> pls_inputs;
std::vector<PlsRemap> pls_outputs;
std::string pls_decl(const PlsRemap &variable);
const char* to_pls_qualifiers_glsl(const SPIRVariable &variable);
void emit_pls();
void remap_pls_variables();
};
}
#endif

879
spirv.hpp Normal file
View File

@ -0,0 +1,879 @@
// Copyright (c) 2014-2016 The Khronos Group Inc.
//
// Permission is hereby granted, free of charge, to any person obtaining a copy
// of this software and/or associated documentation files (the "Materials"),
// to deal in the Materials without restriction, including without limitation
// the rights to use, copy, modify, merge, publish, distribute, sublicense,
// and/or sell copies of the Materials, and to permit persons to whom the
// Materials are furnished to do so, subject to the following conditions:
//
// The above copyright notice and this permission notice shall be included in
// all copies or substantial portions of the Materials.
//
// MODIFICATIONS TO THIS FILE MAY MEAN IT NO LONGER ACCURATELY REFLECTS KHRONOS
// STANDARDS. THE UNMODIFIED, NORMATIVE VERSIONS OF KHRONOS SPECIFICATIONS AND
// HEADER INFORMATION ARE LOCATED AT https://www.khronos.org/registry/
//
// THE MATERIALS ARE 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 MATERIALS OR THE USE OR OTHER DEALINGS
// IN THE MATERIALS.
// This header is automatically generated by the same tool that creates
// the Binary Section of the SPIR-V specification.
// Enumeration tokens for SPIR-V, in various styles:
// C, C++, C++11, JSON, Lua, Python
//
// - C will have tokens with a "Spv" prefix, e.g.: SpvSourceLanguageGLSL
// - C++ will have tokens in the "spv" name space, e.g.: spv::SourceLanguageGLSL
// - C++11 will use enum classes in the spv namespace, e.g.: spv::SourceLanguage::GLSL
// - Lua will use tables, e.g.: spv.SourceLanguage.GLSL
// - Python will use dictionaries, e.g.: spv['SourceLanguage']['GLSL']
//
// Some tokens act like mask values, which can be OR'd together,
// while others are mutually exclusive. The mask-like ones have
// "Mask" in their name, and a parallel enum that has the shift
// amount (1 << x) for each corresponding enumerant.
#ifndef spirv_HPP
#define spirv_HPP
namespace spv {
typedef unsigned int Id;
#define SPV_VERSION 0x10000
#define SPV_REVISION 3
static const unsigned int MagicNumber = 0x07230203;
static const unsigned int Version = 0x00010000;
static const unsigned int Revision = 3;
static const unsigned int OpCodeMask = 0xffff;
static const unsigned int WordCountShift = 16;
enum SourceLanguage {
SourceLanguageUnknown = 0,
SourceLanguageESSL = 1,
SourceLanguageGLSL = 2,
SourceLanguageOpenCL_C = 3,
SourceLanguageOpenCL_CPP = 4,
};
enum ExecutionModel {
ExecutionModelVertex = 0,
ExecutionModelTessellationControl = 1,
ExecutionModelTessellationEvaluation = 2,
ExecutionModelGeometry = 3,
ExecutionModelFragment = 4,
ExecutionModelGLCompute = 5,
ExecutionModelKernel = 6,
};
enum AddressingModel {
AddressingModelLogical = 0,
AddressingModelPhysical32 = 1,
AddressingModelPhysical64 = 2,
};
enum MemoryModel {
MemoryModelSimple = 0,
MemoryModelGLSL450 = 1,
MemoryModelOpenCL = 2,
};
enum ExecutionMode {
ExecutionModeInvocations = 0,
ExecutionModeSpacingEqual = 1,
ExecutionModeSpacingFractionalEven = 2,
ExecutionModeSpacingFractionalOdd = 3,
ExecutionModeVertexOrderCw = 4,
ExecutionModeVertexOrderCcw = 5,
ExecutionModePixelCenterInteger = 6,
ExecutionModeOriginUpperLeft = 7,
ExecutionModeOriginLowerLeft = 8,
ExecutionModeEarlyFragmentTests = 9,
ExecutionModePointMode = 10,
ExecutionModeXfb = 11,
ExecutionModeDepthReplacing = 12,
ExecutionModeDepthGreater = 14,
ExecutionModeDepthLess = 15,
ExecutionModeDepthUnchanged = 16,
ExecutionModeLocalSize = 17,
ExecutionModeLocalSizeHint = 18,
ExecutionModeInputPoints = 19,
ExecutionModeInputLines = 20,
ExecutionModeInputLinesAdjacency = 21,
ExecutionModeTriangles = 22,
ExecutionModeInputTrianglesAdjacency = 23,
ExecutionModeQuads = 24,
ExecutionModeIsolines = 25,
ExecutionModeOutputVertices = 26,
ExecutionModeOutputPoints = 27,
ExecutionModeOutputLineStrip = 28,
ExecutionModeOutputTriangleStrip = 29,
ExecutionModeVecTypeHint = 30,
ExecutionModeContractionOff = 31,
};
enum StorageClass {
StorageClassUniformConstant = 0,
StorageClassInput = 1,
StorageClassUniform = 2,
StorageClassOutput = 3,
StorageClassWorkgroup = 4,
StorageClassCrossWorkgroup = 5,
StorageClassPrivate = 6,
StorageClassFunction = 7,
StorageClassGeneric = 8,
StorageClassPushConstant = 9,
StorageClassAtomicCounter = 10,
StorageClassImage = 11,
};
enum Dim {
Dim1D = 0,
Dim2D = 1,
Dim3D = 2,
DimCube = 3,
DimRect = 4,
DimBuffer = 5,
DimSubpassData = 6,
};
enum SamplerAddressingMode {
SamplerAddressingModeNone = 0,
SamplerAddressingModeClampToEdge = 1,
SamplerAddressingModeClamp = 2,
SamplerAddressingModeRepeat = 3,
SamplerAddressingModeRepeatMirrored = 4,
};
enum SamplerFilterMode {
SamplerFilterModeNearest = 0,
SamplerFilterModeLinear = 1,
};
enum ImageFormat {
ImageFormatUnknown = 0,
ImageFormatRgba32f = 1,
ImageFormatRgba16f = 2,
ImageFormatR32f = 3,
ImageFormatRgba8 = 4,
ImageFormatRgba8Snorm = 5,
ImageFormatRg32f = 6,
ImageFormatRg16f = 7,
ImageFormatR11fG11fB10f = 8,
ImageFormatR16f = 9,
ImageFormatRgba16 = 10,
ImageFormatRgb10A2 = 11,
ImageFormatRg16 = 12,
ImageFormatRg8 = 13,
ImageFormatR16 = 14,
ImageFormatR8 = 15,
ImageFormatRgba16Snorm = 16,
ImageFormatRg16Snorm = 17,
ImageFormatRg8Snorm = 18,
ImageFormatR16Snorm = 19,
ImageFormatR8Snorm = 20,
ImageFormatRgba32i = 21,
ImageFormatRgba16i = 22,
ImageFormatRgba8i = 23,
ImageFormatR32i = 24,
ImageFormatRg32i = 25,
ImageFormatRg16i = 26,
ImageFormatRg8i = 27,
ImageFormatR16i = 28,
ImageFormatR8i = 29,
ImageFormatRgba32ui = 30,
ImageFormatRgba16ui = 31,
ImageFormatRgba8ui = 32,
ImageFormatR32ui = 33,
ImageFormatRgb10a2ui = 34,
ImageFormatRg32ui = 35,
ImageFormatRg16ui = 36,
ImageFormatRg8ui = 37,
ImageFormatR16ui = 38,
ImageFormatR8ui = 39,
};
enum ImageChannelOrder {
ImageChannelOrderR = 0,
ImageChannelOrderA = 1,
ImageChannelOrderRG = 2,
ImageChannelOrderRA = 3,
ImageChannelOrderRGB = 4,
ImageChannelOrderRGBA = 5,
ImageChannelOrderBGRA = 6,
ImageChannelOrderARGB = 7,
ImageChannelOrderIntensity = 8,
ImageChannelOrderLuminance = 9,
ImageChannelOrderRx = 10,
ImageChannelOrderRGx = 11,
ImageChannelOrderRGBx = 12,
ImageChannelOrderDepth = 13,
ImageChannelOrderDepthStencil = 14,
ImageChannelOrdersRGB = 15,
ImageChannelOrdersRGBx = 16,
ImageChannelOrdersRGBA = 17,
ImageChannelOrdersBGRA = 18,
};
enum ImageChannelDataType {
ImageChannelDataTypeSnormInt8 = 0,
ImageChannelDataTypeSnormInt16 = 1,
ImageChannelDataTypeUnormInt8 = 2,
ImageChannelDataTypeUnormInt16 = 3,
ImageChannelDataTypeUnormShort565 = 4,
ImageChannelDataTypeUnormShort555 = 5,
ImageChannelDataTypeUnormInt101010 = 6,
ImageChannelDataTypeSignedInt8 = 7,
ImageChannelDataTypeSignedInt16 = 8,
ImageChannelDataTypeSignedInt32 = 9,
ImageChannelDataTypeUnsignedInt8 = 10,
ImageChannelDataTypeUnsignedInt16 = 11,
ImageChannelDataTypeUnsignedInt32 = 12,
ImageChannelDataTypeHalfFloat = 13,
ImageChannelDataTypeFloat = 14,
ImageChannelDataTypeUnormInt24 = 15,
ImageChannelDataTypeUnormInt101010_2 = 16,
};
enum ImageOperandsShift {
ImageOperandsBiasShift = 0,
ImageOperandsLodShift = 1,
ImageOperandsGradShift = 2,
ImageOperandsConstOffsetShift = 3,
ImageOperandsOffsetShift = 4,
ImageOperandsConstOffsetsShift = 5,
ImageOperandsSampleShift = 6,
ImageOperandsMinLodShift = 7,
};
enum ImageOperandsMask {
ImageOperandsMaskNone = 0,
ImageOperandsBiasMask = 0x00000001,
ImageOperandsLodMask = 0x00000002,
ImageOperandsGradMask = 0x00000004,
ImageOperandsConstOffsetMask = 0x00000008,
ImageOperandsOffsetMask = 0x00000010,
ImageOperandsConstOffsetsMask = 0x00000020,
ImageOperandsSampleMask = 0x00000040,
ImageOperandsMinLodMask = 0x00000080,
};
enum FPFastMathModeShift {
FPFastMathModeNotNaNShift = 0,
FPFastMathModeNotInfShift = 1,
FPFastMathModeNSZShift = 2,
FPFastMathModeAllowRecipShift = 3,
FPFastMathModeFastShift = 4,
};
enum FPFastMathModeMask {
FPFastMathModeMaskNone = 0,
FPFastMathModeNotNaNMask = 0x00000001,
FPFastMathModeNotInfMask = 0x00000002,
FPFastMathModeNSZMask = 0x00000004,
FPFastMathModeAllowRecipMask = 0x00000008,
FPFastMathModeFastMask = 0x00000010,
};
enum FPRoundingMode {
FPRoundingModeRTE = 0,
FPRoundingModeRTZ = 1,
FPRoundingModeRTP = 2,
FPRoundingModeRTN = 3,
};
enum LinkageType {
LinkageTypeExport = 0,
LinkageTypeImport = 1,
};
enum AccessQualifier {
AccessQualifierReadOnly = 0,
AccessQualifierWriteOnly = 1,
AccessQualifierReadWrite = 2,
};
enum FunctionParameterAttribute {
FunctionParameterAttributeZext = 0,
FunctionParameterAttributeSext = 1,
FunctionParameterAttributeByVal = 2,
FunctionParameterAttributeSret = 3,
FunctionParameterAttributeNoAlias = 4,
FunctionParameterAttributeNoCapture = 5,
FunctionParameterAttributeNoWrite = 6,
FunctionParameterAttributeNoReadWrite = 7,
};
enum Decoration {
DecorationRelaxedPrecision = 0,
DecorationSpecId = 1,
DecorationBlock = 2,
DecorationBufferBlock = 3,
DecorationRowMajor = 4,
DecorationColMajor = 5,
DecorationArrayStride = 6,
DecorationMatrixStride = 7,
DecorationGLSLShared = 8,
DecorationGLSLPacked = 9,
DecorationCPacked = 10,
DecorationBuiltIn = 11,
DecorationNoPerspective = 13,
DecorationFlat = 14,
DecorationPatch = 15,
DecorationCentroid = 16,
DecorationSample = 17,
DecorationInvariant = 18,
DecorationRestrict = 19,
DecorationAliased = 20,
DecorationVolatile = 21,
DecorationConstant = 22,
DecorationCoherent = 23,
DecorationNonWritable = 24,
DecorationNonReadable = 25,
DecorationUniform = 26,
DecorationSaturatedConversion = 28,
DecorationStream = 29,
DecorationLocation = 30,
DecorationComponent = 31,
DecorationIndex = 32,
DecorationBinding = 33,
DecorationDescriptorSet = 34,
DecorationOffset = 35,
DecorationXfbBuffer = 36,
DecorationXfbStride = 37,
DecorationFuncParamAttr = 38,
DecorationFPRoundingMode = 39,
DecorationFPFastMathMode = 40,
DecorationLinkageAttributes = 41,
DecorationNoContraction = 42,
DecorationInputAttachmentIndex = 43,
DecorationAlignment = 44,
};
enum BuiltIn {
BuiltInPosition = 0,
BuiltInPointSize = 1,
BuiltInClipDistance = 3,
BuiltInCullDistance = 4,
BuiltInVertexId = 5,
BuiltInInstanceId = 6,
BuiltInPrimitiveId = 7,
BuiltInInvocationId = 8,
BuiltInLayer = 9,
BuiltInViewportIndex = 10,
BuiltInTessLevelOuter = 11,
BuiltInTessLevelInner = 12,
BuiltInTessCoord = 13,
BuiltInPatchVertices = 14,
BuiltInFragCoord = 15,
BuiltInPointCoord = 16,
BuiltInFrontFacing = 17,
BuiltInSampleId = 18,
BuiltInSamplePosition = 19,
BuiltInSampleMask = 20,
BuiltInFragDepth = 22,
BuiltInHelperInvocation = 23,
BuiltInNumWorkgroups = 24,
BuiltInWorkgroupSize = 25,
BuiltInWorkgroupId = 26,
BuiltInLocalInvocationId = 27,
BuiltInGlobalInvocationId = 28,
BuiltInLocalInvocationIndex = 29,
BuiltInWorkDim = 30,
BuiltInGlobalSize = 31,
BuiltInEnqueuedWorkgroupSize = 32,
BuiltInGlobalOffset = 33,
BuiltInGlobalLinearId = 34,
BuiltInSubgroupSize = 36,
BuiltInSubgroupMaxSize = 37,
BuiltInNumSubgroups = 38,
BuiltInNumEnqueuedSubgroups = 39,
BuiltInSubgroupId = 40,
BuiltInSubgroupLocalInvocationId = 41,
BuiltInVertexIndex = 42,
BuiltInInstanceIndex = 43,
};
enum SelectionControlShift {
SelectionControlFlattenShift = 0,
SelectionControlDontFlattenShift = 1,
};
enum SelectionControlMask {
SelectionControlMaskNone = 0,
SelectionControlFlattenMask = 0x00000001,
SelectionControlDontFlattenMask = 0x00000002,
};
enum LoopControlShift {
LoopControlUnrollShift = 0,
LoopControlDontUnrollShift = 1,
};
enum LoopControlMask {
LoopControlMaskNone = 0,
LoopControlUnrollMask = 0x00000001,
LoopControlDontUnrollMask = 0x00000002,
};
enum FunctionControlShift {
FunctionControlInlineShift = 0,
FunctionControlDontInlineShift = 1,
FunctionControlPureShift = 2,
FunctionControlConstShift = 3,
};
enum FunctionControlMask {
FunctionControlMaskNone = 0,
FunctionControlInlineMask = 0x00000001,
FunctionControlDontInlineMask = 0x00000002,
FunctionControlPureMask = 0x00000004,
FunctionControlConstMask = 0x00000008,
};
enum MemorySemanticsShift {
MemorySemanticsAcquireShift = 1,
MemorySemanticsReleaseShift = 2,
MemorySemanticsAcquireReleaseShift = 3,
MemorySemanticsSequentiallyConsistentShift = 4,
MemorySemanticsUniformMemoryShift = 6,
MemorySemanticsSubgroupMemoryShift = 7,
MemorySemanticsWorkgroupMemoryShift = 8,
MemorySemanticsCrossWorkgroupMemoryShift = 9,
MemorySemanticsAtomicCounterMemoryShift = 10,
MemorySemanticsImageMemoryShift = 11,
};
enum MemorySemanticsMask {
MemorySemanticsMaskNone = 0,
MemorySemanticsAcquireMask = 0x00000002,
MemorySemanticsReleaseMask = 0x00000004,
MemorySemanticsAcquireReleaseMask = 0x00000008,
MemorySemanticsSequentiallyConsistentMask = 0x00000010,
MemorySemanticsUniformMemoryMask = 0x00000040,
MemorySemanticsSubgroupMemoryMask = 0x00000080,
MemorySemanticsWorkgroupMemoryMask = 0x00000100,
MemorySemanticsCrossWorkgroupMemoryMask = 0x00000200,
MemorySemanticsAtomicCounterMemoryMask = 0x00000400,
MemorySemanticsImageMemoryMask = 0x00000800,
};
enum MemoryAccessShift {
MemoryAccessVolatileShift = 0,
MemoryAccessAlignedShift = 1,
MemoryAccessNontemporalShift = 2,
};
enum MemoryAccessMask {
MemoryAccessMaskNone = 0,
MemoryAccessVolatileMask = 0x00000001,
MemoryAccessAlignedMask = 0x00000002,
MemoryAccessNontemporalMask = 0x00000004,
};
enum Scope {
ScopeCrossDevice = 0,
ScopeDevice = 1,
ScopeWorkgroup = 2,
ScopeSubgroup = 3,
ScopeInvocation = 4,
};
enum GroupOperation {
GroupOperationReduce = 0,
GroupOperationInclusiveScan = 1,
GroupOperationExclusiveScan = 2,
};
enum KernelEnqueueFlags {
KernelEnqueueFlagsNoWait = 0,
KernelEnqueueFlagsWaitKernel = 1,
KernelEnqueueFlagsWaitWorkGroup = 2,
};
enum KernelProfilingInfoShift {
KernelProfilingInfoCmdExecTimeShift = 0,
};
enum KernelProfilingInfoMask {
KernelProfilingInfoMaskNone = 0,
KernelProfilingInfoCmdExecTimeMask = 0x00000001,
};
enum Capability {
CapabilityMatrix = 0,
CapabilityShader = 1,
CapabilityGeometry = 2,
CapabilityTessellation = 3,
CapabilityAddresses = 4,
CapabilityLinkage = 5,
CapabilityKernel = 6,
CapabilityVector16 = 7,
CapabilityFloat16Buffer = 8,
CapabilityFloat16 = 9,
CapabilityFloat64 = 10,
CapabilityInt64 = 11,
CapabilityInt64Atomics = 12,
CapabilityImageBasic = 13,
CapabilityImageReadWrite = 14,
CapabilityImageMipmap = 15,
CapabilityPipes = 17,
CapabilityGroups = 18,
CapabilityDeviceEnqueue = 19,
CapabilityLiteralSampler = 20,
CapabilityAtomicStorage = 21,
CapabilityInt16 = 22,
CapabilityTessellationPointSize = 23,
CapabilityGeometryPointSize = 24,
CapabilityImageGatherExtended = 25,
CapabilityStorageImageMultisample = 27,
CapabilityUniformBufferArrayDynamicIndexing = 28,
CapabilitySampledImageArrayDynamicIndexing = 29,
CapabilityStorageBufferArrayDynamicIndexing = 30,
CapabilityStorageImageArrayDynamicIndexing = 31,
CapabilityClipDistance = 32,
CapabilityCullDistance = 33,
CapabilityImageCubeArray = 34,
CapabilitySampleRateShading = 35,
CapabilityImageRect = 36,
CapabilitySampledRect = 37,
CapabilityGenericPointer = 38,
CapabilityInt8 = 39,
CapabilityInputAttachment = 40,
CapabilitySparseResidency = 41,
CapabilityMinLod = 42,
CapabilitySampled1D = 43,
CapabilityImage1D = 44,
CapabilitySampledCubeArray = 45,
CapabilitySampledBuffer = 46,
CapabilityImageBuffer = 47,
CapabilityImageMSArray = 48,
CapabilityStorageImageExtendedFormats = 49,
CapabilityImageQuery = 50,
CapabilityDerivativeControl = 51,
CapabilityInterpolationFunction = 52,
CapabilityTransformFeedback = 53,
CapabilityGeometryStreams = 54,
CapabilityStorageImageReadWithoutFormat = 55,
CapabilityStorageImageWriteWithoutFormat = 56,
CapabilityMultiViewport = 57,
};
enum Op {
OpNop = 0,
OpUndef = 1,
OpSourceContinued = 2,
OpSource = 3,
OpSourceExtension = 4,
OpName = 5,
OpMemberName = 6,
OpString = 7,
OpLine = 8,
OpExtension = 10,
OpExtInstImport = 11,
OpExtInst = 12,
OpMemoryModel = 14,
OpEntryPoint = 15,
OpExecutionMode = 16,
OpCapability = 17,
OpTypeVoid = 19,
OpTypeBool = 20,
OpTypeInt = 21,
OpTypeFloat = 22,
OpTypeVector = 23,
OpTypeMatrix = 24,
OpTypeImage = 25,
OpTypeSampler = 26,
OpTypeSampledImage = 27,
OpTypeArray = 28,
OpTypeRuntimeArray = 29,
OpTypeStruct = 30,
OpTypeOpaque = 31,
OpTypePointer = 32,
OpTypeFunction = 33,
OpTypeEvent = 34,
OpTypeDeviceEvent = 35,
OpTypeReserveId = 36,
OpTypeQueue = 37,
OpTypePipe = 38,
OpTypeForwardPointer = 39,
OpConstantTrue = 41,
OpConstantFalse = 42,
OpConstant = 43,
OpConstantComposite = 44,
OpConstantSampler = 45,
OpConstantNull = 46,
OpSpecConstantTrue = 48,
OpSpecConstantFalse = 49,
OpSpecConstant = 50,
OpSpecConstantComposite = 51,
OpSpecConstantOp = 52,
OpFunction = 54,
OpFunctionParameter = 55,
OpFunctionEnd = 56,
OpFunctionCall = 57,
OpVariable = 59,
OpImageTexelPointer = 60,
OpLoad = 61,
OpStore = 62,
OpCopyMemory = 63,
OpCopyMemorySized = 64,
OpAccessChain = 65,
OpInBoundsAccessChain = 66,
OpPtrAccessChain = 67,
OpArrayLength = 68,
OpGenericPtrMemSemantics = 69,
OpInBoundsPtrAccessChain = 70,
OpDecorate = 71,
OpMemberDecorate = 72,
OpDecorationGroup = 73,
OpGroupDecorate = 74,
OpGroupMemberDecorate = 75,
OpVectorExtractDynamic = 77,
OpVectorInsertDynamic = 78,
OpVectorShuffle = 79,
OpCompositeConstruct = 80,
OpCompositeExtract = 81,
OpCompositeInsert = 82,
OpCopyObject = 83,
OpTranspose = 84,
OpSampledImage = 86,
OpImageSampleImplicitLod = 87,
OpImageSampleExplicitLod = 88,
OpImageSampleDrefImplicitLod = 89,
OpImageSampleDrefExplicitLod = 90,
OpImageSampleProjImplicitLod = 91,
OpImageSampleProjExplicitLod = 92,
OpImageSampleProjDrefImplicitLod = 93,
OpImageSampleProjDrefExplicitLod = 94,
OpImageFetch = 95,
OpImageGather = 96,
OpImageDrefGather = 97,
OpImageRead = 98,
OpImageWrite = 99,
OpImage = 100,
OpImageQueryFormat = 101,
OpImageQueryOrder = 102,
OpImageQuerySizeLod = 103,
OpImageQuerySize = 104,
OpImageQueryLod = 105,
OpImageQueryLevels = 106,
OpImageQuerySamples = 107,
OpConvertFToU = 109,
OpConvertFToS = 110,
OpConvertSToF = 111,
OpConvertUToF = 112,
OpUConvert = 113,
OpSConvert = 114,
OpFConvert = 115,
OpQuantizeToF16 = 116,
OpConvertPtrToU = 117,
OpSatConvertSToU = 118,
OpSatConvertUToS = 119,
OpConvertUToPtr = 120,
OpPtrCastToGeneric = 121,
OpGenericCastToPtr = 122,
OpGenericCastToPtrExplicit = 123,
OpBitcast = 124,
OpSNegate = 126,
OpFNegate = 127,
OpIAdd = 128,
OpFAdd = 129,
OpISub = 130,
OpFSub = 131,
OpIMul = 132,
OpFMul = 133,
OpUDiv = 134,
OpSDiv = 135,
OpFDiv = 136,
OpUMod = 137,
OpSRem = 138,
OpSMod = 139,
OpFRem = 140,
OpFMod = 141,
OpVectorTimesScalar = 142,
OpMatrixTimesScalar = 143,
OpVectorTimesMatrix = 144,
OpMatrixTimesVector = 145,
OpMatrixTimesMatrix = 146,
OpOuterProduct = 147,
OpDot = 148,
OpIAddCarry = 149,
OpISubBorrow = 150,
OpUMulExtended = 151,
OpSMulExtended = 152,
OpAny = 154,
OpAll = 155,
OpIsNan = 156,
OpIsInf = 157,
OpIsFinite = 158,
OpIsNormal = 159,
OpSignBitSet = 160,
OpLessOrGreater = 161,
OpOrdered = 162,
OpUnordered = 163,
OpLogicalEqual = 164,
OpLogicalNotEqual = 165,
OpLogicalOr = 166,
OpLogicalAnd = 167,
OpLogicalNot = 168,
OpSelect = 169,
OpIEqual = 170,
OpINotEqual = 171,
OpUGreaterThan = 172,
OpSGreaterThan = 173,
OpUGreaterThanEqual = 174,
OpSGreaterThanEqual = 175,
OpULessThan = 176,
OpSLessThan = 177,
OpULessThanEqual = 178,
OpSLessThanEqual = 179,
OpFOrdEqual = 180,
OpFUnordEqual = 181,
OpFOrdNotEqual = 182,
OpFUnordNotEqual = 183,
OpFOrdLessThan = 184,
OpFUnordLessThan = 185,
OpFOrdGreaterThan = 186,
OpFUnordGreaterThan = 187,
OpFOrdLessThanEqual = 188,
OpFUnordLessThanEqual = 189,
OpFOrdGreaterThanEqual = 190,
OpFUnordGreaterThanEqual = 191,
OpShiftRightLogical = 194,
OpShiftRightArithmetic = 195,
OpShiftLeftLogical = 196,
OpBitwiseOr = 197,
OpBitwiseXor = 198,
OpBitwiseAnd = 199,
OpNot = 200,
OpBitFieldInsert = 201,
OpBitFieldSExtract = 202,
OpBitFieldUExtract = 203,
OpBitReverse = 204,
OpBitCount = 205,
OpDPdx = 207,
OpDPdy = 208,
OpFwidth = 209,
OpDPdxFine = 210,
OpDPdyFine = 211,
OpFwidthFine = 212,
OpDPdxCoarse = 213,
OpDPdyCoarse = 214,
OpFwidthCoarse = 215,
OpEmitVertex = 218,
OpEndPrimitive = 219,
OpEmitStreamVertex = 220,
OpEndStreamPrimitive = 221,
OpControlBarrier = 224,
OpMemoryBarrier = 225,
OpAtomicLoad = 227,
OpAtomicStore = 228,
OpAtomicExchange = 229,
OpAtomicCompareExchange = 230,
OpAtomicCompareExchangeWeak = 231,
OpAtomicIIncrement = 232,
OpAtomicIDecrement = 233,
OpAtomicIAdd = 234,
OpAtomicISub = 235,
OpAtomicSMin = 236,
OpAtomicUMin = 237,
OpAtomicSMax = 238,
OpAtomicUMax = 239,
OpAtomicAnd = 240,
OpAtomicOr = 241,
OpAtomicXor = 242,
OpPhi = 245,
OpLoopMerge = 246,
OpSelectionMerge = 247,
OpLabel = 248,
OpBranch = 249,
OpBranchConditional = 250,
OpSwitch = 251,
OpKill = 252,
OpReturn = 253,
OpReturnValue = 254,
OpUnreachable = 255,
OpLifetimeStart = 256,
OpLifetimeStop = 257,
OpGroupAsyncCopy = 259,
OpGroupWaitEvents = 260,
OpGroupAll = 261,
OpGroupAny = 262,
OpGroupBroadcast = 263,
OpGroupIAdd = 264,
OpGroupFAdd = 265,
OpGroupFMin = 266,
OpGroupUMin = 267,
OpGroupSMin = 268,
OpGroupFMax = 269,
OpGroupUMax = 270,
OpGroupSMax = 271,
OpReadPipe = 274,
OpWritePipe = 275,
OpReservedReadPipe = 276,
OpReservedWritePipe = 277,
OpReserveReadPipePackets = 278,
OpReserveWritePipePackets = 279,
OpCommitReadPipe = 280,
OpCommitWritePipe = 281,
OpIsValidReserveId = 282,
OpGetNumPipePackets = 283,
OpGetMaxPipePackets = 284,
OpGroupReserveReadPipePackets = 285,
OpGroupReserveWritePipePackets = 286,
OpGroupCommitReadPipe = 287,
OpGroupCommitWritePipe = 288,
OpEnqueueMarker = 291,
OpEnqueueKernel = 292,
OpGetKernelNDrangeSubGroupCount = 293,
OpGetKernelNDrangeMaxSubGroupSize = 294,
OpGetKernelWorkGroupSize = 295,
OpGetKernelPreferredWorkGroupSizeMultiple = 296,
OpRetainEvent = 297,
OpReleaseEvent = 298,
OpCreateUserEvent = 299,
OpIsValidEvent = 300,
OpSetUserEventStatus = 301,
OpCaptureEventProfilingInfo = 302,
OpGetDefaultQueue = 303,
OpBuildNDRange = 304,
OpImageSparseSampleImplicitLod = 305,
OpImageSparseSampleExplicitLod = 306,
OpImageSparseSampleDrefImplicitLod = 307,
OpImageSparseSampleDrefExplicitLod = 308,
OpImageSparseSampleProjImplicitLod = 309,
OpImageSparseSampleProjExplicitLod = 310,
OpImageSparseSampleProjDrefImplicitLod = 311,
OpImageSparseSampleProjDrefExplicitLod = 312,
OpImageSparseFetch = 313,
OpImageSparseGather = 314,
OpImageSparseDrefGather = 315,
OpImageSparseTexelsResident = 316,
OpNoLine = 317,
OpAtomicFlagTestAndSet = 318,
OpAtomicFlagClear = 319,
OpImageSparseRead = 320,
};
// Overload operator| for mask bit combining
inline ImageOperandsMask operator|(ImageOperandsMask a, ImageOperandsMask b) { return ImageOperandsMask(unsigned(a) | unsigned(b)); }
inline FPFastMathModeMask operator|(FPFastMathModeMask a, FPFastMathModeMask b) { return FPFastMathModeMask(unsigned(a) | unsigned(b)); }
inline SelectionControlMask operator|(SelectionControlMask a, SelectionControlMask b) { return SelectionControlMask(unsigned(a) | unsigned(b)); }
inline LoopControlMask operator|(LoopControlMask a, LoopControlMask b) { return LoopControlMask(unsigned(a) | unsigned(b)); }
inline FunctionControlMask operator|(FunctionControlMask a, FunctionControlMask b) { return FunctionControlMask(unsigned(a) | unsigned(b)); }
inline MemorySemanticsMask operator|(MemorySemanticsMask a, MemorySemanticsMask b) { return MemorySemanticsMask(unsigned(a) | unsigned(b)); }
inline MemoryAccessMask operator|(MemoryAccessMask a, MemoryAccessMask b) { return MemoryAccessMask(unsigned(a) | unsigned(b)); }
inline KernelProfilingInfoMask operator|(KernelProfilingInfoMask a, KernelProfilingInfoMask b) { return KernelProfilingInfoMask(unsigned(a) | unsigned(b)); }
} // end namespace spv
#endif // #ifndef spirv_HPP

127
test_shaders.py Executable file
View File

@ -0,0 +1,127 @@
#!/usr/bin/env python3
import sys
import os
import subprocess
import tempfile
import re
import itertools
import hashlib
import shutil
def parse_stats(stats):
m = re.search('([0-9]+) work registers', stats)
registers = int(m.group(1)) if m else 0
m = re.search('([0-9]+) uniform registers', stats)
uniform_regs = int(m.group(1)) if m else 0
m_list = re.findall('(-?[0-9]+)\s+(-?[0-9]+)\s+(-?[0-9]+)', stats)
alu_short = float(m_list[1][0]) if m_list else 0
ls_short = float(m_list[1][1]) if m_list else 0
tex_short = float(m_list[1][2]) if m_list else 0
alu_long = float(m_list[2][0]) if m_list else 0
ls_long = float(m_list[2][1]) if m_list else 0
tex_long = float(m_list[2][2]) if m_list else 0
return (registers, uniform_regs, alu_short, ls_short, tex_short, alu_long, ls_long, tex_long)
def get_shader_type(shader):
_, ext = os.path.splitext(shader)
if ext == '.vert':
return '--vertex'
elif ext == '.frag':
return '--fragment'
elif ext == '.comp':
return '--compute'
elif ext == '.tesc':
return '--tessellation_control'
elif ext == '.tese':
return '--tessellation_evaluation'
elif ext == '.geom':
return '--geometry'
else:
return ''
def get_shader_stats(shader):
f, path = tempfile.mkstemp()
os.close(f)
p = subprocess.Popen(['malisc', get_shader_type(shader), '--core', 'Mali-T760', '-V', shader], stdout = subprocess.PIPE, stderr = subprocess.PIPE)
stdout, stderr = p.communicate()
os.remove(path)
if p.returncode != 0:
print(stderr.decode('utf-8'))
raise OSError('malisc failed')
p.wait()
returned = stdout.decode('utf-8')
return parse_stats(returned)
def cross_compile(shader):
spirv_f, spirv_path = tempfile.mkstemp()
glsl_f, glsl_path = tempfile.mkstemp(suffix = os.path.basename(shader))
os.close(spirv_f)
os.close(glsl_f)
subprocess.check_call(['glslangValidator', '-G', '-o', spirv_path, shader])
subprocess.check_call(['./spir2cross', '--output', glsl_path, spirv_path])
return (spirv_path, glsl_path)
def md5_for_file(path):
md5 = hashlib.md5()
with open(path, 'rb') as f:
for chunk in iter(lambda: f.read(8192), b''):
md5.update(chunk)
return md5.digest()
def make_reference_dir(path):
base = os.path.dirname(path)
if not os.path.exists(base):
os.makedirs(base)
def regression_check(shader, glsl):
reference = os.path.join('./reference', shader)
if os.path.exists(reference):
if md5_for_file(glsl) != md5_for_file(reference):
print('Generated GLSL has changed for {}!'.format(reference))
if os.path.exists(reference):
os.remove(reference)
make_reference_dir(reference)
shutil.move(glsl, reference)
else:
os.remove(glsl)
else:
print('Found new shader {}. Placing GLSL in {}'.format(shader, reference))
make_reference_dir(reference)
shutil.move(glsl, reference)
def test_shader(stats, shader):
print('Testing shader:', shader)
pristine_stats = get_shader_stats(shader)
spirv, glsl = cross_compile(shader)
cross_stats = get_shader_stats(glsl)
regression_check(shader, glsl)
os.remove(spirv)
a = []
a.append(shader)
for i in pristine_stats:
a.append(str(i))
for i in cross_stats:
a.append(str(i))
print(','.join(a), file = stats)
def test_shaders(shader_dir):
with open('stats.csv', 'w') as stats:
print('Shader,OrigRegs,OrigUniRegs,OrigALUShort,OrigLSShort,OrigTEXShort,OrigALULong,OrigLSLong,OrigTEXLong,CrossRegs,CrossUniRegs,CrossALUShort,CrossLSShort,CrossTEXShort,CrossALULong,CrossLSLong,CrossTEXLong', file = stats)
for f in os.walk(os.path.join(shader_dir)):
for i in f[2]:
shader = os.path.join(f[0], i)
test_shader(stats, shader)
if __name__ == '__main__':
test_shaders(sys.argv[1])
print('Stats in stats.csv!')