archive/blender-archive
Archived
1
1
Fork 0
Code Pull Requests Packages Activity

Metal: MTLShader and MTLShaderGenerator implementation.

Browse Source 
Full support for translation and compilation of shaders in Metal, using
GPUShaderCreateInfo. Includes render pipeline state creation and management,
enabling all standard GPU viewport rendering features in Metal.

Authored by Apple: Michael Parkin-White, Marco Giordano

Ref T96261

Reviewed By: fclem

Maniphest Tasks: T96261

Differential Revision: https://developer.blender.org/D15563
This commit is contained in:
Thomas Dinges
2022-09-01 22:22:32 +02:00
committed by Clément Foucault
parent ac07fb38a1
commit cc8ea6ac67
42 changed files with 9252 additions and 213 deletions
Download Patch File Download Diff File Expand all files Collapse all files

61
source/blender/gpu/CMakeLists.txt
View File

@@ -194,6 +194,9 @@ set(METAL_SRC
metal/mtl_index_buffer.mm
metal/mtl_memory.mm
metal/mtl_query.mm
metal/mtl_shader.mm
metal/mtl_shader_generator.mm
metal/mtl_shader_interface.mm
metal/mtl_state.mm
metal/mtl_texture.mm
metal/mtl_texture_util.mm
@@ -207,7 +210,14 @@ set(METAL_SRC
metal/mtl_framebuffer.hh
metal/mtl_index_buffer.hh
metal/mtl_memory.hh
metal/mtl_pso_descriptor_state.hh
metal/mtl_primitive.hh
metal/mtl_query.hh
metal/mtl_shader.hh
metal/mtl_shader_generator.hh
metal/mtl_shader_interface_type.hh
metal/mtl_shader_interface.hh
metal/mtl_shader_shared.h
metal/mtl_state.hh
metal/mtl_texture.hh
metal/mtl_uniform_buffer.hh
@@ -227,6 +237,9 @@ set(LIB
)
set(MSL_SRC
shaders/metal/mtl_shader_defines.msl
shaders/metal/mtl_shader_common.msl
metal/kernels/compute_texture_update.msl
metal/kernels/compute_texture_read.msl
metal/kernels/depth_2d_update_float_frag.glsl
@@ -458,21 +471,44 @@ set(GLSL_SRC
GPU_shader_shared_utils.h
)
set(MTL_BACKEND_GLSL_SRC
metal/kernels/compute_texture_update.msl
metal/kernels/compute_texture_read.msl
metal/kernels/depth_2d_update_float_frag.glsl
metal/kernels/depth_2d_update_int24_frag.glsl
metal/kernels/depth_2d_update_int32_frag.glsl
metal/kernels/depth_2d_update_vert.glsl
metal/kernels/gpu_shader_fullscreen_blit_vert.glsl
metal/kernels/gpu_shader_fullscreen_blit_frag.glsl
)
set(MSL_SRC
shaders/metal/mtl_shader_defines.msl
shaders/metal/mtl_shader_common.msl
metal/mtl_shader_shared.h
)
if(WITH_METAL_BACKEND)
list(APPEND GLSL_SRC ${MTL_BACKEND_GLSL_SRC})
set(MSL_C)
foreach(MSL_FILE ${MSL_SRC})
data_to_c_simple(${MSL_FILE} MSL_C)
endforeach()
endif()
set(GLSL_C)
foreach(GLSL_FILE ${GLSL_SRC})
data_to_c_simple(${GLSL_FILE} GLSL_C)
endforeach()
set(SHADER_C)
list(APPEND SHADER_C ${GLSL_C})
if(WITH_METAL_BACKEND)
set(MSL_C)
foreach(MSL_FILE ${MSL_SRC})
data_to_c_simple(${MSL_FILE} MSL_C)
endforeach()
list(APPEND GLSL_C ${MSL_C})
list(APPEND SHADER_C ${MSL_C})
endif()
blender_add_lib(bf_gpu_shaders "${GLSL_C}" "" "" "")
blender_add_lib(bf_gpu_shaders "${SHADER_C}" "" "" "")
list(APPEND LIB
bf_gpu_shaders
@@ -587,6 +623,16 @@ set(SRC_SHADER_CREATE_INFOS
shaders/compositor/infos/compositor_split_viewer_info.hh
)
set(SRC_SHADER_CREATE_INFOS_MTL
metal/kernels/depth_2d_update_info.hh
metal/kernels/gpu_shader_fullscreen_blit_info.hh
)
if(WITH_METAL_BACKEND)
list(APPEND SRC_SHADER_CREATE_INFOS ${SRC_SHADER_CREATE_INFOS_MTL})
endif()
set(SHADER_CREATE_INFOS_CONTENT "")
foreach(DESCRIPTOR_FILE ${SRC_SHADER_CREATE_INFOS})
string(APPEND SHADER_CREATE_INFOS_CONTENT "#include \"${DESCRIPTOR_FILE}\"\n")
@@ -629,6 +675,7 @@ if(WITH_GPU_BUILDTIME_SHADER_BUILDER)
if(APPLE)
add_executable(shader_builder
intern/gpu_shader_builder.cc
intern/gpu_shader_builder_stubs.cc
${shader_create_info_list_file}
)

1
source/blender/gpu/GPU_capabilities.h
View File

@@ -30,6 +30,7 @@ int GPU_max_batch_indices(void);
int GPU_max_batch_vertices(void);
int GPU_max_vertex_attribs(void);
int GPU_max_varying_floats(void);
int GPU_max_samplers(void);
int GPU_max_shader_storage_buffer_bindings(void);
int GPU_max_compute_shader_storage_blocks(void);
int GPU_max_samplers(void);

5
source/blender/gpu/GPU_shader_shared_utils.h
View File

@@ -43,6 +43,9 @@
# define sqrtf sqrt
# define expf exp
# define bool1 bool
/* Type name collision with Metal shading language - These typenames are already defined. */
# ifndef GPU_METAL
# define float2 vec2
# define float3 vec3
# define float4 vec4
@@ -53,10 +56,10 @@
# define uint2 uvec2
# define uint3 uvec3
# define uint4 uvec4
# define bool1 bool
# define bool2 bvec2
# define bool3 bvec3
# define bool4 bvec4
# endif
#else /* C / C++ */
# pragma once

4
source/blender/gpu/intern/gpu_context.cc
View File

@@ -56,11 +56,15 @@ static void gpu_backend_discard();
namespace blender::gpu {
int Context::context_counter = 0;
Context::Context()
{
thread_ = pthread_self();
is_active_ = false;
matrix_state = GPU_matrix_state_create();
context_id = Context::context_counter;
Context::context_counter++;
}
Context::~Context()

8
source/blender/gpu/intern/gpu_context_private.hh
View File

@@ -48,6 +48,14 @@ class Context {
DebugStack debug_stack;
/* GPUContext counter used to assign a unique ID to each GPUContext.
* NOTE(Metal): This is required by the Metal Backend, as a bug exists in the global OS shader
* cache wherein compilation of identical source from two distinct threads can result in an
* invalid cache collision, result in a broken shader object. Appending the unique context ID
* onto compiled sources ensures the source hashes are different. */
static int context_counter;
int context_id = 0;
protected:
/** Thread on which this context is active. */
pthread_t thread_;

3
source/blender/gpu/intern/gpu_shader.cc
View File

@@ -95,6 +95,9 @@ static void standard_defines(Vector<const char *> &sources)
case GPU_BACKEND_OPENGL:
sources.append("#define GPU_OPENGL\n");
break;
case GPU_BACKEND_METAL:
sources.append("#define GPU_METAL\n");
break;
default:
BLI_assert(false && "Invalid GPU Backend Type");
break;

52
source/blender/gpu/intern/gpu_shader_create_info.hh
View File

@@ -32,6 +32,7 @@ namespace blender::gpu::shader {
#endif
enum class Type {
/* Types supported natively across all GPU backends. */
FLOAT = 0,
VEC2,
VEC3,
@@ -47,6 +48,21 @@ enum class Type {
IVEC3,
IVEC4,
BOOL,
/* Additionally supported types to enable data optimisation and native
* support in some GPUBackends.
* NOTE: These types must be representable in all APIs. E.g. VEC3_101010I2 is aliased as vec3 in
* the GL backend, as implicit type conversions from packed normal attribute data to vec3 is
* supported. UCHAR/CHAR types are natively supported in Metal and can be used to avoid
* additional data conversions for GPU_COMP_U8 vertex attributes. */
VEC3_101010I2,
UCHAR,
UCHAR2,
UCHAR3,
UCHAR4,
CHAR,
CHAR2,
CHAR3,
CHAR4
};
/* All of these functions is a bit out of place */
@@ -86,6 +102,40 @@ static inline std::ostream &operator<<(std::ostream &stream, const Type type)
return stream << "mat3";
case Type::MAT4:
return stream << "mat4";
case Type::VEC3_101010I2:
return stream << "vec3_1010102_Inorm";
case Type::UCHAR:
return stream << "uchar";
case Type::UCHAR2:
return stream << "uchar2";
case Type::UCHAR3:
return stream << "uchar3";
case Type::UCHAR4:
return stream << "uchar4";
case Type::CHAR:
return stream << "char";
case Type::CHAR2:
return stream << "char2";
case Type::CHAR3:
return stream << "char3";
case Type::CHAR4:
return stream << "char4";
case Type::INT:
return stream << "int";
case Type::IVEC2:
return stream << "ivec2";
case Type::IVEC3:
return stream << "ivec3";
case Type::IVEC4:
return stream << "ivec4";
case Type::UINT:
return stream << "uint";
case Type::UVEC2:
return stream << "uvec2";
case Type::UVEC3:
return stream << "uvec3";
case Type::UVEC4:
return stream << "uvec4";
default:
BLI_assert(0);
return stream;
@@ -228,6 +278,8 @@ enum class PrimitiveOut {
POINTS = 0,
LINE_STRIP,
TRIANGLE_STRIP,
LINES,
TRIANGLES,
};
struct StageInterfaceInfo {

2
source/blender/gpu/metal/kernels/compute_texture_read.msl
View File

@@ -74,7 +74,7 @@ template<> uchar convert_type<uchar>(float val)
template<> uint convert_type<uint>(float val)
{
return uint(val * double(0xFFFFFFFFu));
return uint(val * float(0xFFFFFFFFu));
}
struct TextureReadParams {

16
source/blender/gpu/metal/kernels/compute_texture_update.msl
View File

@@ -38,22 +38,6 @@ using namespace metal;
# define POSITION_TYPE uint3
#endif
float3 mtl_linear_to_srgb_attr(float3 c)
{
c = max(c, float3(0.0));
float3 c1 = c * 12.92;
float3 c2 = 1.055 * pow(c, float3(1.0 / 2.4)) - 0.055;
return mix(c1, c2, step(float3(0.0031308), c));
}
float3 mtl_srgb_to_linear_attr(float3 c)
{
c = max(c, float3(0.0));
float3 c1 = c * (1.0 / 12.92);
float3 c2 = pow((c + 0.055) * (1.0 / 1.055), float3(2.4));
return mix(c1, c2, step(float3(0.04045), c));
}
struct TextureUpdateParams {
int mip_index;
int extent[3];

5
source/blender/gpu/metal/kernels/depth_2d_update_float_frag.glsl
View File

@@ -1,9 +1,4 @@
uniform sampler2D source_data;
uniform int mip;
in vec2 texCoord_interp;
void main()
{
gl_FragDepth = textureLod(source_data, texCoord_interp, mip).r;

35
source/blender/gpu/metal/kernels/depth_2d_update_info.hh Normal file
View File

@@ -0,0 +1,35 @@
/* SPDX-License-Identifier: GPL-2.0-or-later */
/** \file
* \ingroup gpu
*/
#include "gpu_shader_create_info.hh"
GPU_SHADER_INTERFACE_INFO(depth_2d_update_iface, "").smooth(Type::VEC2, "texCoord_interp");
GPU_SHADER_CREATE_INFO(depth_2d_update_info_base)
.vertex_in(0, Type::VEC2, "pos")
.vertex_out(depth_2d_update_iface)
.fragment_out(0, Type::VEC4, "fragColor")
.push_constant(Type::VEC2, "extent")
.push_constant(Type::VEC2, "offset")
.push_constant(Type::VEC2, "size")
.push_constant(Type::INT, "mip")
.sampler(0, ImageType::FLOAT_2D, "source_data", Frequency::PASS)
.vertex_source("depth_2d_update_vert.glsl");
GPU_SHADER_CREATE_INFO(depth_2d_update_float)
.fragment_source("depth_2d_update_float_frag.glsl")
.additional_info("depth_2d_update_info_base")
.do_static_compilation(true);
GPU_SHADER_CREATE_INFO(depth_2d_update_int24)
.fragment_source("depth_2d_update_int24_frag.glsl")
.additional_info("depth_2d_update_info_base")
.do_static_compilation(true);
GPU_SHADER_CREATE_INFO(depth_2d_update_int32)
.fragment_source("depth_2d_update_int32_frag.glsl")
.additional_info("depth_2d_update_info_base")
.do_static_compilation(true);

4
source/blender/gpu/metal/kernels/depth_2d_update_int24_frag.glsl
View File

@@ -1,8 +1,4 @@
uniform isampler2D source_data;
uniform int mip;
in vec2 texCoord_interp;
void main()
{

5
source/blender/gpu/metal/kernels/depth_2d_update_int32_frag.glsl
View File

@@ -1,9 +1,4 @@
uniform isampler2D source_data;
uniform int mip;
in vec2 texCoord_interp;
void main()
{
uint val = textureLod(source_data, texCoord_interp, mip).r;

6
source/blender/gpu/metal/kernels/depth_2d_update_vert.glsl
View File

@@ -1,10 +1,4 @@
uniform vec2 extent;
uniform vec2 offset;
uniform vec2 size;
out vec2 texCoord_interp;
in vec2 pos;
void main()
{
vec4 rect = vec4(offset.x, offset.y, offset.x + extent.x, offset.y + extent.y);

5
source/blender/gpu/metal/kernels/gpu_shader_fullscreen_blit_frag.glsl
View File

@@ -1,10 +1,5 @@
in vec4 uvcoordsvar;
uniform sampler2D imageTexture;
uniform int mip;
out vec4 fragColor;
void main()
{
vec4 tex_color = textureLod(imageTexture, uvcoordsvar.xy, mip);

23
source/blender/gpu/metal/kernels/gpu_shader_fullscreen_blit_info.hh Normal file
View File

@@ -0,0 +1,23 @@
/* SPDX-License-Identifier: GPL-2.0-or-later */
/** \file
* \ingroup gpu
*/
#include "gpu_shader_create_info.hh"
GPU_SHADER_INTERFACE_INFO(fullscreen_blit_iface, "").smooth(Type::VEC4, "uvcoordsvar");
GPU_SHADER_CREATE_INFO(fullscreen_blit)
.vertex_in(0, Type::VEC2, "pos")
.vertex_out(fullscreen_blit_iface)
.fragment_out(0, Type::VEC4, "fragColor")
.push_constant(Type::VEC2, "fullscreen")
.push_constant(Type::VEC2, "size")
.push_constant(Type::VEC2, "dst_offset")
.push_constant(Type::VEC2, "src_offset")
.push_constant(Type::INT, "mip")
.sampler(0, ImageType::FLOAT_2D, "imageTexture", Frequency::PASS)
.vertex_source("gpu_shader_fullscreen_blit_vert.glsl")
.fragment_source("gpu_shader_fullscreen_blit_frag.glsl")
.do_static_compilation(true);

8
source/blender/gpu/metal/kernels/gpu_shader_fullscreen_blit_vert.glsl
View File

@@ -1,12 +1,4 @@
out vec4 uvcoordsvar;
in vec2 pos;
uniform vec2 fullscreen;
uniform vec2 size;
uniform vec2 dst_offset;
uniform vec2 src_offset;
void main()
{
/* The position represents a 0-1 square, we first scale it by the size we want to have it on

11
source/blender/gpu/metal/mtl_backend.mm
View File

@@ -12,6 +12,7 @@
#include "mtl_framebuffer.hh"
#include "mtl_index_buffer.hh"
#include "mtl_query.hh"
#include "mtl_shader.hh"
#include "mtl_uniform_buffer.hh"
#include "gpu_capabilities_private.hh"
@@ -71,8 +72,8 @@ QueryPool *MTLBackend::querypool_alloc()
Shader *MTLBackend::shader_alloc(const char *name)
{
/* TODO(Metal): Implement MTLShader. */
return nullptr;
MTLContext *mtl_context = MTLContext::get();
return new MTLShader(mtl_context, name);
};
Texture *MTLBackend::texture_alloc(const char *name)
@@ -168,7 +169,7 @@ void MTLBackend::platform_init(MTLContext *ctx)
eGPUSupportLevel support_level = GPU_SUPPORT_LEVEL_SUPPORTED;
BLI_assert(ctx);
id<MTLDevice> mtl_device = nil; /*ctx->device; TODO(Metal): Implement MTLContext. */
id<MTLDevice> mtl_device = ctx->device;
BLI_assert(device);
NSString *gpu_name = [mtl_device name];
@@ -187,7 +188,7 @@ void MTLBackend::platform_init(MTLContext *ctx)
os = GPU_OS_UNIX;
#endif
BLI_assert(os == GPU_OS_MAC && "Platform must be macOS");
BLI_assert_msg(os == GPU_OS_MAC, "Platform must be macOS");
/* Determine Vendor from name. */
if (strstr(vendor, "ATI") || strstr(vendor, "AMD")) {
@@ -334,7 +335,7 @@ bool MTLBackend::metal_is_supported()
void MTLBackend::capabilities_init(MTLContext *ctx)
{
BLI_assert(ctx);
id<MTLDevice> device = nil; /*ctx->device TODO(Metal): Implement MTLContext. */
id<MTLDevice> device = ctx->device;
BLI_assert(device);
/* Initialize Capabilities. */

2
source/blender/gpu/metal/mtl_capabilities.hh
View File

@@ -14,6 +14,8 @@ namespace gpu {
#define MTL_MAX_TEXTURE_SLOTS 128
#define MTL_MAX_SAMPLER_SLOTS MTL_MAX_TEXTURE_SLOTS
/* Max limit without using bindless for samplers. */
#define MTL_MAX_DEFAULT_SAMPLERS 16
#define MTL_MAX_UNIFORM_BUFFER_BINDINGS 31
#define MTL_MAX_VERTEX_INPUT_ATTRIBUTES 31
#define MTL_MAX_UNIFORMS_PER_BLOCK 64

2
source/blender/gpu/metal/mtl_common.hh
View File

@@ -13,4 +13,6 @@
* Set as number of GPU frames in flight, plus an additional value for extra possible CPU frame. */
#define MTL_NUM_SAFE_FRAMES (MTL_MAX_DRAWABLES + 1)
/* Display debug information about missing attributes and incorrect vertex formats. */
#define MTL_DEBUG_SHADER_ATTRIBUTES 0
#endif

16
source/blender/gpu/metal/mtl_context.hh
View File

@@ -17,6 +17,8 @@
#include "mtl_common.hh"
#include "mtl_framebuffer.hh"
#include "mtl_memory.hh"
#include "mtl_shader.hh"
#include "mtl_shader_interface.hh"
#include "mtl_texture.hh"
#include <Cocoa/Cocoa.h>
@@ -32,7 +34,6 @@ namespace blender::gpu {
/* Forward Declarations */
class MTLContext;
class MTLCommandBufferManager;
class MTLShader;
class MTLUniformBuf;
/* Structs containing information on current binding state for textures and samplers. */
@@ -40,7 +41,7 @@ struct MTLTextureBinding {
bool used;
/* Same value as index in bindings array. */
uint texture_slot_index;
uint slot_index;
gpu::MTLTexture *texture_resource;
};
@@ -56,9 +57,10 @@ struct MTLSamplerBinding {
/* Metal Context Render Pass State -- Used to track active RenderCommandEncoder state based on
* bound MTLFrameBuffer's.Owned by MTLContext. */
struct MTLRenderPassState {
class MTLRenderPassState {
friend class MTLContext;
public:
MTLRenderPassState(MTLContext &context, MTLCommandBufferManager &command_buffer_manager)
: ctx(context), cmd(command_buffer_manager){};
@@ -570,6 +572,11 @@ class MTLContext : public Context {
friend class MTLBackend;
private:
/* Null buffers for empty/unintialized bindings.
* Null attribute buffer follows default attribute format of OpenGL Backend. */
id<MTLBuffer> null_buffer_; /* All zero's. */
id<MTLBuffer> null_attribute_buffer_; /* Value float4(0.0,0.0,0.0,1.0). */
/* Compute and specialization caches. */
MTLContextTextureUtils texture_utils_;
@@ -713,6 +720,9 @@ class MTLContext : public Context {
{
return MTLContext::global_memory_manager;
}
/* Uniform Buffer Bindings to command encoders. */
id<MTLBuffer> get_null_buffer();
id<MTLBuffer> get_null_attribute_buffer();
};
} // namespace blender::gpu

80
source/blender/gpu/metal/mtl_context.mm
View File

@@ -5,6 +5,8 @@
*/
#include "mtl_context.hh"
#include "mtl_debug.hh"
#include "mtl_shader.hh"
#include "mtl_shader_interface.hh"
#include "mtl_state.hh"
#include "DNA_userdef_types.h"
@@ -29,19 +31,33 @@ MTLContext::MTLContext(void *ghost_window) : memory_manager(*this), main_command
/* Init debug. */
debug::mtl_debug_init();
/* Device creation.
* TODO(Metal): This is a temporary initialisation path to enable testing of features
* and shader compilation tests. Future functionality should fetch the existing device
* from GHOST_ContextCGL.mm. Plumbing to be updated in future. */
this->device = MTLCreateSystemDefaultDevice();
/* Initialize command buffer state. */
this->main_command_buffer.prepare();
/* Initialise imm and pipeline state */
this->pipeline_state.initialised = false;
/* Frame management. */
is_inside_frame_ = false;
current_frame_index_ = 0;
/* Prepare null data buffer */
null_buffer_ = nil;
null_attribute_buffer_ = nil;
/* Create FrameBuffer handles. */
MTLFrameBuffer *mtl_front_left = new MTLFrameBuffer(this, "front_left");
MTLFrameBuffer *mtl_back_left = new MTLFrameBuffer(this, "back_left");
this->front_left = mtl_front_left;
this->back_left = mtl_back_left;
this->active_fb = this->back_left;
/* Prepare platform and capabilities. (NOTE: With METAL, this needs to be done after CTX
* initialization). */
MTLBackend::platform_init(this);
@@ -93,6 +109,12 @@ MTLContext::~MTLContext()
sampler_state_cache_[i] = nil;
}
}
if (null_buffer_) {
[null_buffer_ release];
}
if (null_attribute_buffer_) {
[null_attribute_buffer_ release];
}
}
void MTLContext::begin_frame()
@@ -227,6 +249,50 @@ MTLFrameBuffer *MTLContext::get_default_framebuffer()
return static_cast<MTLFrameBuffer *>(this->back_left);
}
MTLShader *MTLContext::get_active_shader()
{
return this->pipeline_state.active_shader;
}
id<MTLBuffer> MTLContext::get_null_buffer()
{
if (null_buffer_ != nil) {
return null_buffer_;
}
static const int null_buffer_size = 4096;
null_buffer_ = [this->device newBufferWithLength:null_buffer_size
options:MTLResourceStorageModeManaged];
[null_buffer_ retain];
uint32_t *null_data = (uint32_t *)calloc(0, null_buffer_size);
memcpy([null_buffer_ contents], null_data, null_buffer_size);
[null_buffer_ didModifyRange:NSMakeRange(0, null_buffer_size)];
free(null_data);
BLI_assert(null_buffer_ != nil);
return null_buffer_;
}
id<MTLBuffer> MTLContext::get_null_attribute_buffer()
{
if (null_attribute_buffer_ != nil) {
return null_attribute_buffer_;
}
/* Allocate Null buffer if it has not yet been created.
* Min buffer size is 256 bytes -- though we only need 64 bytes of data. */
static const int null_buffer_size = 256;
null_attribute_buffer_ = [this->device newBufferWithLength:null_buffer_size
options:MTLResourceStorageModeManaged];
BLI_assert(null_attribute_buffer_ != nil);
[null_attribute_buffer_ retain];
float data[4] = {0.0f, 0.0f, 0.0f, 1.0f};
memcpy([null_attribute_buffer_ contents], data, sizeof(float) * 4);
[null_attribute_buffer_ didModifyRange:NSMakeRange(0, null_buffer_size)];
return null_attribute_buffer_;
}
/** \} */
/* -------------------------------------------------------------------- */
@@ -239,20 +305,20 @@ void MTLContext::pipeline_state_init()
/*** Initialize state only once. ***/
if (!this->pipeline_state.initialised) {
this->pipeline_state.initialised = true;
this->pipeline_state.active_shader = NULL;
this->pipeline_state.active_shader = nullptr;
/* Clear bindings state. */
for (int t = 0; t < GPU_max_textures(); t++) {
this->pipeline_state.texture_bindings[t].used = false;
this->pipeline_state.texture_bindings[t].texture_slot_index = t;
this->pipeline_state.texture_bindings[t].texture_resource = NULL;
this->pipeline_state.texture_bindings[t].slot_index = -1;
this->pipeline_state.texture_bindings[t].texture_resource = nullptr;
}
for (int s = 0; s < MTL_MAX_SAMPLER_SLOTS; s++) {
this->pipeline_state.sampler_bindings[s].used = false;
}
for (int u = 0; u < MTL_MAX_UNIFORM_BUFFER_BINDINGS; u++) {
this->pipeline_state.ubo_bindings[u].bound = false;
this->pipeline_state.ubo_bindings[u].ubo = NULL;
this->pipeline_state.ubo_bindings[u].ubo = nullptr;
}
}
@@ -487,11 +553,6 @@ id<MTLSamplerState> MTLContext::get_sampler_from_state(MTLSamplerState sampler_s
id<MTLSamplerState> MTLContext::generate_sampler_from_state(MTLSamplerState sampler_state)
{
/* Check if sampler already exists for given state. */
id<MTLSamplerState> st = sampler_state_cache_[(uint)sampler_state];
if (st != nil) {
return st;
}
else {
MTLSamplerDescriptor *descriptor = [[MTLSamplerDescriptor alloc] init];
descriptor.normalizedCoordinates = true;
@@ -533,7 +594,6 @@ id<MTLSamplerState> MTLContext::generate_sampler_from_state(MTLSamplerState samp
[descriptor autorelease];
return state;
}
}
id<MTLSamplerState> MTLContext::get_default_sampler_state()
{

13
source/blender/gpu/metal/mtl_memory.mm
View File

@@ -73,7 +73,9 @@ gpu::MTLBuffer *MTLBufferPool::allocate_with_data(uint64_t size,
return this->allocate_aligned_with_data(size, 256, cpu_visible, data);
}
gpu::MTLBuffer *MTLBufferPool::allocate_aligned(uint64_t size, uint alignment, bool cpu_visible)
gpu::MTLBuffer *MTLBufferPool::allocate_aligned(uint64_t size,
uint32_t alignment,
bool cpu_visible)
{
/* Check not required. Main GPU module usage considered thread-safe. */
// BLI_assert(BLI_thread_is_main());
@@ -167,7 +169,7 @@ gpu::MTLBuffer *MTLBufferPool::allocate_aligned(uint64_t size, uint alignment, b
}
gpu::MTLBuffer *MTLBufferPool::allocate_aligned_with_data(uint64_t size,
uint alignment,
uint32_t alignment,
bool cpu_visible,
const void *data)
{
@@ -548,7 +550,8 @@ void gpu::MTLBuffer::set_label(NSString *str)
void gpu::MTLBuffer::debug_ensure_used()
{
/* Debug: If buffer is not flagged as in-use, this is a problem. */
BLI_assert(in_use_ &&
BLI_assert_msg(
in_use_,
"Buffer should be marked as 'in-use' if being actively used by an instance. Buffer "
"has likely already been freed.");
}
@@ -665,9 +668,9 @@ MTLTemporaryBuffer MTLScratchBufferManager::scratch_buffer_allocate_range_aligne
/* Ensure scratch buffer allocation alignment adheres to offset alignment requirements. */
alignment = max_uu(alignment, 256);
BLI_assert(current_scratch_buffer_ >= 0 && "Scratch Buffer index not set");
BLI_assert_msg(current_scratch_buffer_ >= 0, "Scratch Buffer index not set");
MTLCircularBuffer *current_scratch_buff = this->scratch_buffers_[current_scratch_buffer_];
BLI_assert(current_scratch_buff != nullptr && "Scratch Buffer does not exist");
BLI_assert_msg(current_scratch_buff != nullptr, "Scratch Buffer does not exist");
MTLTemporaryBuffer allocated_range = current_scratch_buff->allocate_range_aligned(alloc_size,
alignment);
BLI_assert(allocated_range.size >= alloc_size && allocated_range.size <= alloc_size + alignment);

100
source/blender/gpu/metal/mtl_primitive.hh Normal file
View File

@@ -0,0 +1,100 @@
/* SPDX-License-Identifier: GPL-2.0-or-later */
/** \file
* \ingroup gpu
*
* Encapsulation of Frame-buffer states (attached textures, viewport, scissors).
*/
#pragma once
#include "BLI_assert.h"
#include "GPU_primitive.h"
#include <Metal/Metal.h>
namespace blender::gpu {
/** Utility functions **/
static inline MTLPrimitiveTopologyClass mtl_prim_type_to_topology_class(MTLPrimitiveType prim_type)
{
switch (prim_type) {
case MTLPrimitiveTypePoint:
return MTLPrimitiveTopologyClassPoint;
case MTLPrimitiveTypeLine:
case MTLPrimitiveTypeLineStrip:
return MTLPrimitiveTopologyClassLine;
case MTLPrimitiveTypeTriangle:
case MTLPrimitiveTypeTriangleStrip:
return MTLPrimitiveTopologyClassTriangle;
}
return MTLPrimitiveTopologyClassUnspecified;
}
static inline MTLPrimitiveType gpu_prim_type_to_metal(GPUPrimType prim_type)
{
switch (prim_type) {
case GPU_PRIM_POINTS:
return MTLPrimitiveTypePoint;
case GPU_PRIM_LINES:
case GPU_PRIM_LINES_ADJ:
case GPU_PRIM_LINE_LOOP:
return MTLPrimitiveTypeLine;
case GPU_PRIM_LINE_STRIP:
case GPU_PRIM_LINE_STRIP_ADJ:
return MTLPrimitiveTypeLineStrip;
case GPU_PRIM_TRIS:
case GPU_PRIM_TRI_FAN:
case GPU_PRIM_TRIS_ADJ:
return MTLPrimitiveTypeTriangle;
case GPU_PRIM_TRI_STRIP:
return MTLPrimitiveTypeTriangleStrip;
case GPU_PRIM_NONE:
return MTLPrimitiveTypePoint;
};
}
/* Certain primitive types are not supported in Metal, and require emulation.
* `GPU_PRIM_LINE_LOOP` and `GPU_PRIM_TRI_FAN` required index buffer patching.
* Adjacency types do not need emulation as the input structure is the same,
* and access is controlled from the vertex shader through SSBO vertex fetch.
* -- These Adj cases are only used in geometry shaders in OpenGL. */
static inline bool mtl_needs_topology_emulation(GPUPrimType prim_type)
{
BLI_assert(prim_type != GPU_PRIM_NONE);
switch (prim_type) {
case GPU_PRIM_LINE_LOOP:
case GPU_PRIM_TRI_FAN:
return true;
default:
return false;
}
return false;
}
static inline bool mtl_vertex_count_fits_primitive_type(uint32_t vertex_count,
MTLPrimitiveType prim_type)
{
if (vertex_count == 0) {
return false;
}
switch (prim_type) {
case MTLPrimitiveTypeLineStrip:
return (vertex_count > 1);
case MTLPrimitiveTypeLine:
return (vertex_count % 2 == 0);
case MTLPrimitiveTypePoint:
return (vertex_count > 0);
case MTLPrimitiveTypeTriangle:
return (vertex_count % 3 == 0);
case MTLPrimitiveTypeTriangleStrip:
return (vertex_count > 2);
}
BLI_assert(false);
return false;
}
} // namespace blender::gpu

250
source/blender/gpu/metal/mtl_pso_descriptor_state.hh Normal file
View File

@@ -0,0 +1,250 @@
/* SPDX-License-Identifier: GPL-2.0-or-later */
/** \file
* \ingroup gpu
*/
#pragma once
#include "GPU_vertex_format.h"
#include <Metal/Metal.h>
namespace blender::gpu {
/** Vertex attribute and buffer descriptor wrappers
* for use in PSO construction and caching. */
struct MTLVertexAttributeDescriptorPSO {
MTLVertexFormat format;
int offset;
int buffer_index;
GPUVertFetchMode format_conversion_mode;
bool operator==(const MTLVertexAttributeDescriptorPSO &other) const
{
return (format == other.format) && (offset == other.offset) &&
(buffer_index == other.buffer_index) &&
(format_conversion_mode == other.format_conversion_mode);
}
uint64_t hash() const
{
return (uint64_t)((uint64_t)this->format ^ (this->offset << 4) ^ (this->buffer_index << 8) ^
(this->format_conversion_mode << 12));
}
};
struct MTLVertexBufferLayoutDescriptorPSO {
MTLVertexStepFunction step_function;
int step_rate;
int stride;
bool operator==(const MTLVertexBufferLayoutDescriptorPSO &other) const
{
return (step_function == other.step_function) && (step_rate == other.step_rate) &&
(stride == other.stride);
}
uint64_t hash() const
{
return (uint64_t)((uint64_t)this->step_function ^ (this->step_rate << 4) ^
(this->stride << 8));
}
};
/* SSBO attribute state caching. */
struct MTLSSBOAttribute {
int mtl_attribute_index;
int vbo_id;
int attribute_offset;
int per_vertex_stride;
int attribute_format;
bool is_instance;
MTLSSBOAttribute(){};
MTLSSBOAttribute(
int attribute_ind, int vertexbuffer_ind, int offset, int stride, int format, bool instanced)
: mtl_attribute_index(attribute_ind),
vbo_id(vertexbuffer_ind),
attribute_offset(offset),
per_vertex_stride(stride),
attribute_format(format),
is_instance(instanced)
{
}
bool operator==(const MTLSSBOAttribute &other) const
{
return (memcmp(this, &other, sizeof(MTLSSBOAttribute)) == 0);
}
};
struct MTLVertexDescriptor {
/* Core Vertex Attributes. */
MTLVertexAttributeDescriptorPSO attributes[GPU_VERT_ATTR_MAX_LEN];
MTLVertexBufferLayoutDescriptorPSO
buffer_layouts[GPU_BATCH_VBO_MAX_LEN + GPU_BATCH_INST_VBO_MAX_LEN];
int num_attributes;
int num_vert_buffers;
MTLPrimitiveTopologyClass prim_topology_class;
/* WORKAROUND: SSBO Vertex-fetch attributes -- These follow the same structure
* but have slightly different binding rules, passed in via uniform
* push constant data block. */
bool uses_ssbo_vertex_fetch;
MTLSSBOAttribute ssbo_attributes[GPU_VERT_ATTR_MAX_LEN];
int num_ssbo_attributes;
bool operator==(const MTLVertexDescriptor &other) const
{
if ((this->num_attributes != other.num_attributes) ||
(this->num_vert_buffers != other.num_vert_buffers)) {
return false;
}
if (this->prim_topology_class != other.prim_topology_class) {
return false;
};
for (const int a : IndexRange(this->num_attributes)) {
if (!(this->attributes[a] == other.attributes[a])) {
return false;
}
}
for (const int b : IndexRange(this->num_vert_buffers)) {
if (!(this->buffer_layouts[b] == other.buffer_layouts[b])) {
return false;
}
}
/* NOTE: No need to compare SSBO attributes, as these will match attribute bindings for the
* given shader. These are simply extra pre-resolved properties we want to include in the
* cache. */
return true;
}
uint64_t hash() const
{
uint64_t hash = (uint64_t)(this->num_attributes ^ this->num_vert_buffers);
for (const int a : IndexRange(this->num_attributes)) {
hash ^= this->attributes[a].hash() << a;
}
for (const int b : IndexRange(this->num_vert_buffers)) {
hash ^= this->buffer_layouts[b].hash() << (b + 10);
}
/* NOTE: SSBO vertex fetch members not hashed as these will match attribute bindings. */
return hash;
}
};
/* Metal Render Pipeline State Descriptor -- All unique information which feeds PSO creation. */
struct MTLRenderPipelineStateDescriptor {
/* This state descriptor will contain ALL parameters which generate a unique PSO.
* We will then use this state-object to efficiently look-up or create a
* new PSO for the current shader.
*
* Unlike the 'MTLContextGlobalShaderPipelineState', this struct contains a subset of
* parameters used to distinguish between unique PSOs. This struct is hashable and only contains
* those parameters which are required by PSO generation. Non-unique state such as bound
* resources is not tracked here, as it does not require a unique PSO permutation if changed. */
/* Input Vertex Descriptor. */
MTLVertexDescriptor vertex_descriptor;
/* Render Target attachment state.
* Assign to MTLPixelFormatInvalid if not used. */
int num_color_attachments;
MTLPixelFormat color_attachment_format[GPU_FB_MAX_COLOR_ATTACHMENT];
MTLPixelFormat depth_attachment_format;
MTLPixelFormat stencil_attachment_format;
/* Render Pipeline State affecting PSO creation. */
bool blending_enabled;
MTLBlendOperation alpha_blend_op;
MTLBlendOperation rgb_blend_op;
MTLBlendFactor dest_alpha_blend_factor;
MTLBlendFactor dest_rgb_blend_factor;
MTLBlendFactor src_alpha_blend_factor;
MTLBlendFactor src_rgb_blend_factor;
/* Global colour write mask as this cannot be specified per attachment. */
MTLColorWriteMask color_write_mask;
/* Point size required by point primitives. */
float point_size = 0.0f;
/* Comparison Operator for caching. */
bool operator==(const MTLRenderPipelineStateDescriptor &other) const
{
if (!(vertex_descriptor == other.vertex_descriptor)) {
return false;
}
if ((num_color_attachments != other.num_color_attachments) ||
(depth_attachment_format != other.depth_attachment_format) ||
(stencil_attachment_format != other.stencil_attachment_format) ||
(color_write_mask != other.color_write_mask) ||
(blending_enabled != other.blending_enabled) || (alpha_blend_op != other.alpha_blend_op) ||
(rgb_blend_op != other.rgb_blend_op) ||
(dest_alpha_blend_factor != other.dest_alpha_blend_factor) ||
(dest_rgb_blend_factor != other.dest_rgb_blend_factor) ||
(src_alpha_blend_factor != other.src_alpha_blend_factor) ||
(src_rgb_blend_factor != other.src_rgb_blend_factor) ||
(vertex_descriptor.prim_topology_class != other.vertex_descriptor.prim_topology_class) ||
(point_size != other.point_size)) {
return false;
}
/* Attachments can be skipped, so num_color_attachments will not define the range. */
for (const int c : IndexRange(GPU_FB_MAX_COLOR_ATTACHMENT)) {
if (color_attachment_format[c] != other.color_attachment_format[c]) {
return false;
}
}
return true;
}
uint64_t hash() const
{
/* NOTE(Metal): Current setup aims to minimise overlap of parameters
* which are more likely to be different, to ensure earlier hash
* differences without having to fallback to comparisons.
* Though this could likely be further improved to remove
* has collisions. */
uint64_t hash = this->vertex_descriptor.hash();
hash ^= (uint64_t)this->num_color_attachments << 16; /* up to 6 (3 bits). */
hash ^= (uint64_t)this->depth_attachment_format << 18; /* up to 555 (9 bits). */
hash ^= (uint64_t)this->stencil_attachment_format << 20; /* up to 555 (9 bits). */
hash ^= (uint64_t)(*(
(uint64_t *)&this->vertex_descriptor.prim_topology_class)); /* Up to 3 (2 bits). */
/* Only include elements in Hash if they are needed - avoids variable null assignments
* influencing hash. */
if (this->num_color_attachments > 0) {
hash ^= (uint64_t)this->color_write_mask << 22; /* 4 bit bitmask. */
hash ^= (uint64_t)this->alpha_blend_op << 26; /* Up to 4 (3 bits). */
hash ^= (uint64_t)this->rgb_blend_op << 29; /* Up to 4 (3 bits). */
hash ^= (uint64_t)this->dest_alpha_blend_factor << 32; /* Up to 18 (5 bits). */
hash ^= (uint64_t)this->dest_rgb_blend_factor << 37; /* Up to 18 (5 bits). */
hash ^= (uint64_t)this->src_alpha_blend_factor << 42; /* Up to 18 (5 bits). */
hash ^= (uint64_t)this->src_rgb_blend_factor << 47; /* Up to 18 (5 bits). */
}
for (const uint c : IndexRange(GPU_FB_MAX_COLOR_ATTACHMENT)) {
hash ^= (uint64_t)this->color_attachment_format[c] << (c + 52); // up to 555 (9 bits)
}
hash |= (uint64_t)((this->blending_enabled && (this->num_color_attachments > 0)) ? 1 : 0)
<< 62;
hash ^= (uint64_t)this->point_size;
return hash;
}
};
} // namespace blender::gpu

1164
source/blender/gpu/metal/mtl_shader.hh Normal file
View File

File diff suppressed because it is too large Load Diff

1263
source/blender/gpu/metal/mtl_shader.mm Normal file
View File

File diff suppressed because it is too large Load Diff

724
source/blender/gpu/metal/mtl_shader_generator.hh Normal file
View File

@@ -0,0 +1,724 @@
/* SPDX-License-Identifier: GPL-2.0-or-later */
#pragma once
#include "gpu_shader_create_info.hh"
#include "gpu_shader_private.hh"
/** -- Metal Shader Generator for GLSL -> MSL conversion --
*
* The Metal shader generator class is used as a conversion utility for generating
* a compatible MSL shader from a source GLSL shader. There are several steps
* involved in creating a shader, and structural changes which enable the source
* to function in the same way.
*
* 1) Extraction and conversion of shaders input's and output's to their Metal-compatible
* version. This is a subtle data transformation from GPUShaderCreateInfo, allowing
* for Metal-specific parameters.
*
* 2) Determine usage of shader features such as GL global variable usage, depth write output,
* clip distances, multilayered rendering, barycentric coordinates etc;
*
* 3) Generate MSL shader.
*
* 4) Populate MTLShaderInterface, describing input/output structure, bindpoints, buffer size and
* alignment, shader feature usage etc; Everything required by the Metal backend to successfully
* enable use of shaders and GPU backend features.
*
*
*
* For each shading stage, we generate an MSL shader following these steps:
*
* 1) Output custom shader defines describing modes e.g. whether we are using
* sampler bindings or argument buffers; at the top of the shader.
*
* 2) Inject common Metal headers.
* - mtl_shader_defines.msl is used to map GLSL functions to MSL.
* - mtl_shader_common.msl is added to ALL MSL shaders to provide
* common functionality required by the backend. This primarily
* contains function-constant hooks, used in PSO generation.
*
* 3) Create a class Scope which wraps the GLSL shader. This is used to
* create a global per-thread scope around the shader source, to allow
* access to common shader members (GLSL globals, shader inputs/outptus etc)
*
* 4) Generate shader interface structs and populate local members where required for:
* - VertexInputs
* - VertexOutputs
* - Uniforms
* - Uniform Blocks
* - textures;
* etc;
*
* 5) Inject GLSL source.
*
* 6) Generate MSL shader entry point function. Every Metal shader must have a
* vertex/fragment/kernel entrypoint, which contains the function binding table.
* This is where bindings are specified and passed into the shader.
*
* For converted shaders, the MSL entry-point will also instantiate a shader
* class per thread, and pass over bound resource references into the class.
*
* Finally, the shaders "main()" method will be called, and outputs are copied.
*
* Note: For position outputs, the default output position will be converted to
* the Metal coordinate space, which involves flipping the Y coordinate and
* re-mapping the depth range between 0 and 1, as with Vulkan.
*
*
* The final shader structure looks as follows:
*
* -- Shader defines --
* #define USE_ARGUMENT_BUFFER_FOR_SAMPLERS 0
* ... etc ...;
*
* class MetalShaderVertexImp {
*
* -- Common shader interface structs --
* struct VertexIn {
* vec4 pos [[attribute(0)]]
* }
* struct VertexOut {...}
* struct PushConstantBlock {...}
* struct drw_Globals {...}
* ...
*
* -- GLSL source code --
* ...
* };
*
* vertex MetalShaderVertexImp::VertexOut vertex_function_entry(
* MetalShaderVertexImp::VertexIn v_in [[stage_in]],
* constant PushConstantBlock& globals [[buffer(MTL_uniform_buffer_base_index)]]) {
*
* MetalShaderVertexImp impl;
* -- Copy input members into impl instance --
* -- Execute GLSL main function --
* impl.main();
*
* -- Copy outputs and return --
* MetalShaderVertexImp::VertexOut out;
* out.pos = impl.pos;
* -- transform position to Metal coordinate system --
* return v_out;
* }
*
* -- SSBO-vertex-fetchmode --
*
* SSBO-vertex-fetchmode is a special option wherein vertex buffers are bound directly
* as buffers in the shader, rather than using the VertexDescriptor and [[stage_in]] vertex
* assembly.
*
* The purpose of this mode is to enable random-access reading of all vertex data. This is
* particularly useful for efficiently converting geometry shaders to Metal shading language,
* as these techniques are not supported natively in Metal.
*
* Geometry shaders can be re-created by firing off a vertex shader with the desired number of
* total output vertices. Each vertex can then read whichever input attributes it needs to
* achieve the output result.
* This manual reading is also used to provide support for GPU_provoking_vertex, wherein the
* output vertex for flat shading needs to change. In these cases, the manual vertex assembly
* can flip which vertices are read within the primitive.
*
* From an efficiency perspective, this is more GPU-friendly than geometry shading, due to improved
* parallelism throughout the whole pipe, and for Apple hardware specifically, there is no
* significant performance loss from manual vertex assembly vs under-the-hood assembly.
*
* This mode works by passing the required vertex descriptor information into the shader
* as uniform data, describing the type, stride, offset, stepmode and buffer index of each
* attribute, such that the shader ssbo-vertex-fetch utility functions know how to extract data.
*
* This also works with indexed rendering, by similarly binding the index buffer as a manul buffer.
*
* When this mode is used, the code generation and shader interface generation varies to accomodate
* the required features.
*
* This mode can be enabled in a shader with:
*
* `#pragma USE_SSBO_VERTEX_FETCH(TriangleList/LineList, output_vertex_count_per_input_primitive)`
*
* This mirrors the geometry shader interface `layout(triangle_strip, max_vertices = 3) out;`
*/
/* SSBO vertex fetch attribute uniform parameter names.
* These uniforms are used to pass the information
* required to perform manual vertex assembly within
* the vertex shader.
* Each vertex attribute requires a number of properties
* in order to correctly extract data from the bound vertex
* buffers. */
#ifndef NDEBUG
/* Global. */
# define UNIFORM_SSBO_USES_INDEXED_RENDERING_STR "uniform_ssbo_uses_indexed_rendering"
# define UNIFORM_SSBO_INDEX_MODE_U16_STR "uniform_ssbo_index_mode_u16"
# define UNIFORM_SSBO_INPUT_PRIM_TYPE_STR "uniform_ssbo_input_prim_type"
# define UNIFORM_SSBO_INPUT_VERT_COUNT_STR "uniform_ssbo_input_vert_count"
/* Per-attribute. */
# define UNIFORM_SSBO_OFFSET_STR "uniform_ssbo_offset_"
# define UNIFORM_SSBO_STRIDE_STR "uniform_ssbo_stride_"
# define UNIFORM_SSBO_FETCHMODE_STR "uniform_ssbo_fetchmode_"
# define UNIFORM_SSBO_VBO_ID_STR "uniform_ssbo_vbo_id_"
# define UNIFORM_SSBO_TYPE_STR "uniform_ssbo_type_"
#else
/* Global. */
# define UNIFORM_SSBO_USES_INDEXED_RENDERING_STR "_ir"
# define UNIFORM_SSBO_INDEX_MODE_U16_STR "_mu"
# define UNIFORM_SSBO_INPUT_PRIM_TYPE_STR "_pt"
# define UNIFORM_SSBO_INPUT_VERT_COUNT_STR "_vc"
/* Per-attribute. */
# define UNIFORM_SSBO_OFFSET_STR "_so"
# define UNIFORM_SSBO_STRIDE_STR "_ss"
# define UNIFORM_SSBO_FETCHMODE_STR "_sf"
# define UNIFORM_SSBO_VBO_ID_STR "_sv"
# define UNIFORM_SSBO_TYPE_STR "_st"
#endif
namespace blender::gpu {
struct MSLUniform {
shader::Type type;
std::string name;
bool is_array;
int array_elems;
ShaderStage stage;
MSLUniform(shader::Type uniform_type,
std::string uniform_name,
bool is_array_type,
uint32_t num_elems = 1)
: type(uniform_type), name(uniform_name), is_array(is_array_type), array_elems(num_elems)
{
}
bool operator==(const MSLUniform &right) const
{
return (type == right.type && name == right.name && is_array == right.is_array &&
array_elems == right.array_elems);
}
};
struct MSLUniformBlock {
std::string type_name;
std::string name;
ShaderStage stage;
bool is_array;
bool operator==(const MSLUniformBlock &right) const
{
return (type_name == right.type_name && name == right.name);
}
};
enum MSLTextureSamplerAccess {
TEXTURE_ACCESS_NONE = 0,
TEXTURE_ACCESS_SAMPLE,
TEXTURE_ACCESS_READ,
TEXTURE_ACCESS_WRITE,
TEXTURE_ACCESS_READWRITE,
};
struct MSLTextureSampler {
ShaderStage stage;
shader::ImageType type;
std::string name;
MSLTextureSamplerAccess access;
uint location;
eGPUTextureType get_texture_binding_type() const;
void resolve_binding_indices();
MSLTextureSampler(ShaderStage in_stage,
shader::ImageType in_sampler_type,
std::string in_sampler_name,
MSLTextureSamplerAccess in_access,
uint in_location)
: stage(in_stage),
type(in_sampler_type),
name(in_sampler_name),
access(in_access),
location(in_location)
{
}
bool operator==(const MSLTextureSampler &right) const
{
/* We do not compare stage as we want to avoid duplication of resources used across multiple
* stages. */
return (type == right.type && name == right.name && access == right.access);
}
std::string get_msl_access_str() const
{
switch (access) {
case TEXTURE_ACCESS_SAMPLE:
return "access::sample";
case TEXTURE_ACCESS_READ:
return "access::read";
case TEXTURE_ACCESS_WRITE:
return "access::write";
case TEXTURE_ACCESS_READWRITE:
return "access::read_write";
default:
BLI_assert(false);
return "";
}
return "";
}
/* Get typestring for wrapped texture class members.
* wrapper struct type contains combined texture and sampler, templated
* against the texture type.
* See `COMBINED_SAMPLER_TYPE` in `mtl_shader_defines.msl`. */
std::string get_msl_typestring_wrapper(bool is_addr) const
{
std::string str;
str = this->get_msl_wrapper_type_str() + "<" + this->get_msl_return_type_str() + "," +
this->get_msl_access_str() + ">" + ((is_addr) ? "* " : " ") + this->name;
return str;
}
/* Get raw texture typestring -- used in entry-point function argument table. */
std::string get_msl_typestring(bool is_addr) const
{
std::string str;
str = this->get_msl_texture_type_str() + "<" + this->get_msl_return_type_str() + "," +
this->get_msl_access_str() + ">" + ((is_addr) ? "* " : " ") + this->name;
return str;
}
std::string get_msl_return_type_str() const;
std::string get_msl_texture_type_str() const;
std::string get_msl_wrapper_type_str() const;
};
struct MSLVertexInputAttribute {
/* layout_location of -1 means unspecified and will
* be populated manually. */
int layout_location;
shader::Type type;
std::string name;
bool operator==(const MSLVertexInputAttribute &right) const
{
return (layout_location == right.layout_location && type == right.type && name == right.name);
}
};
struct MSLVertexOutputAttribute {
std::string type;
std::string name;
/* Instance name specified if attributes belong to a struct. */
std::string instance_name;
/* Interpolation qualifier can be any of smooth (default), flat, no_perspective. */
std::string interpolation_qualifier;
bool is_array;
int array_elems;
bool operator==(const MSLVertexOutputAttribute &right) const
{
return (type == right.type && name == right.name &&
interpolation_qualifier == right.interpolation_qualifier &&
is_array == right.is_array && array_elems == right.array_elems);
}
std::string get_mtl_interpolation_qualifier() const
{
if (interpolation_qualifier == "" || interpolation_qualifier == "smooth") {
return "";
}
else if (interpolation_qualifier == "flat") {
return " [[flat]]";
}
else if (interpolation_qualifier == "noperspective") {
return " [[center_no_perspective]]";
}
return "";
}
};
struct MSLFragmentOutputAttribute {
/* Explicit output binding location N for [[color(N)]] -1 = unspecified. */
int layout_location;
/* Output index for dual source blending. -1 = unspecified. */
int layout_index;
shader::Type type;
std::string name;
bool operator==(const MSLFragmentOutputAttribute &right) const
{
return (layout_location == right.layout_location && type == right.type && name == right.name &&
layout_index == right.layout_index);
}
};
class MSLGeneratorInterface {
static char *msl_patch_default;
public:
/** Shader stage input/output binding information.
* Derived from shader source reflection or GPUShaderCreateInfo. */
blender::Vector<MSLUniformBlock> uniform_blocks;
blender::Vector<MSLUniform> uniforms;
blender::Vector<MSLTextureSampler> texture_samplers;
blender::Vector<MSLVertexInputAttribute> vertex_input_attributes;
blender::Vector<MSLVertexOutputAttribute> vertex_output_varyings;
/* Should match vertex outputs, but defined separately as
* some shader permutations will not utilise all inputs/outputs.
* Final shader uses the intersection between the two sets. */
blender::Vector<MSLVertexOutputAttribute> fragment_input_varyings;
blender::Vector<MSLFragmentOutputAttribute> fragment_outputs;
/* Transform feedback interface. */
blender::Vector<MSLVertexOutputAttribute> vertex_output_varyings_tf;
/* Clip Distances. */
blender::Vector<std::string> clip_distances;
/** GL Global usage. */
/* Whether GL position is used, or an alternative vertex output should be the default. */
bool uses_gl_Position;
/* Whether gl_FragColor is used, or whether an alternative fragment output
* should be the default. */
bool uses_gl_FragColor;
/* Whether gl_PointCoord is used in the fragment shader. If so,
* we define float2 gl_PointCoord [[point_coord]]. */
bool uses_gl_PointCoord;
/* Writes out to gl_PointSize in the vertex shader output. */
bool uses_gl_PointSize;
bool uses_gl_VertexID;
bool uses_gl_InstanceID;
bool uses_gl_BaseInstanceARB;
bool uses_gl_FrontFacing;
/* Sets the output render target array index when using multilayered rendering. */
bool uses_gl_FragDepth;
bool uses_mtl_array_index_;
bool uses_transform_feedback;
bool uses_barycentrics;
/* Parameters. */
shader::DepthWrite depth_write;
/* Shader buffer bind indices for argument buffers. */
int sampler_argument_buffer_bind_index[2] = {-1, -1};
/*** SSBO Vertex fetch mode. ***/
/* Indicates whether to pass in Vertex Buffer's as a regular buffers instead of using vertex
* assembly in the PSO descriptor. Enabled with special pragma. */
bool uses_ssbo_vertex_fetch_mode;
private:
/* Parent shader instance. */
MTLShader &parent_shader_;
/* If prepared from Create info. */
const shader::ShaderCreateInfo *create_info_;
public:
MSLGeneratorInterface(MTLShader &shader) : parent_shader_(shader){};
/** Prepare MSLGeneratorInterface from create-info. **/
void prepare_from_createinfo(const shader::ShaderCreateInfo *info);
/* When SSBO Vertex Fetch mode is used, uniforms are used to pass on the required information
* about vertex attribute bindings, in order to perform manual vertex assembly and random-access
* vertex lookup throughout the bound VBOs.
*
* Some parameters are global for the shader, others change with the currently bound
* VertexBuffers, and their format, as they do with regular GPUBatch's.
*
* (Where ##attr is the attributes name)
* uniform_ssbo_stride_##attr -- Representing the stride between elements of attribute(attr)
* uniform_ssbo_offset_##attr -- Representing the base offset within the vertex
* uniform_ssbo_fetchmode_##attr -- Whether using per-vertex fetch or per-instance fetch
* (0=vert, 1=inst) uniform_ssbo_vbo_id_##attr -- index of the vertex buffer within which the
* data for this attribute is contained uniform_ssbo_type_##attr - The type of data in the
* currently bound buffer -- Could be a mismatch with the Officially reported type. */
void prepare_ssbo_vertex_fetch_uniforms();
/* Samplers. */
bool use_argument_buffer_for_samplers() const;
uint32_t num_samplers_for_stage(ShaderStage stage) const;
/* Returns the bind index, relative to MTL_uniform_buffer_base_index. */
uint32_t get_sampler_argument_buffer_bind_index(ShaderStage stage);
/* Code generation utility functions. */
std::string generate_msl_uniform_structs(ShaderStage shader_stage);
std::string generate_msl_vertex_in_struct();
std::string generate_msl_vertex_out_struct(ShaderStage shader_stage);
std::string generate_msl_vertex_transform_feedback_out_struct(ShaderStage shader_stage);
std::string generate_msl_fragment_out_struct();
std::string generate_msl_vertex_inputs_string();
std::string generate_msl_fragment_inputs_string();
std::string generate_msl_vertex_entry_stub();
std::string generate_msl_fragment_entry_stub();
std::string generate_msl_global_uniform_population(ShaderStage stage);
std::string generate_ubo_block_macro_chain(MSLUniformBlock block);
std::string generate_msl_uniform_block_population(ShaderStage stage);
std::string generate_msl_vertex_attribute_input_population();
std::string generate_msl_vertex_output_population();
std::string generate_msl_vertex_output_tf_population();
std::string generate_msl_fragment_input_population();
std::string generate_msl_fragment_output_population();
std::string generate_msl_uniform_undefs(ShaderStage stage);
std::string generate_ubo_block_undef_chain(ShaderStage stage);
std::string generate_msl_texture_vars(ShaderStage shader_stage);
void generate_msl_textures_input_string(std::stringstream &out, ShaderStage stage);
void generate_msl_uniforms_input_string(std::stringstream &out, ShaderStage stage);
/* Location is not always specified, so this will resolve outstanding locations. */
void resolve_input_attribute_locations();
void resolve_fragment_output_locations();
/* Create shader interface for converted GLSL shader. */
MTLShaderInterface *bake_shader_interface(const char *name);
/* Fetch combined shader source header. */
char *msl_patch_default_get();
MEM_CXX_CLASS_ALLOC_FUNCS("MSLGeneratorInterface");
};
inline std::string get_stage_class_name(ShaderStage stage)
{
switch (stage) {
case ShaderStage::VERTEX:
return "MTLShaderVertexImpl";
case ShaderStage::FRAGMENT:
return "MTLShaderFragmentImpl";
default:
BLI_assert_unreachable();
return "";
}
return "";
}
inline bool is_builtin_type(std::string type)
{
/* Add Types as needed. */
/* TODO(Metal): Consider replacing this with a switch and constexpr hash and switch.
* Though most efficient and maintainable approach to be determined. */
static std::map<std::string, eMTLDataType> glsl_builtin_types = {
{"float", MTL_DATATYPE_FLOAT},
{"vec2", MTL_DATATYPE_FLOAT2},
{"vec3", MTL_DATATYPE_FLOAT3},
{"vec4", MTL_DATATYPE_FLOAT4},
{"int", MTL_DATATYPE_INT},
{"ivec2", MTL_DATATYPE_INT2},
{"ivec3", MTL_DATATYPE_INT3},
{"ivec4", MTL_DATATYPE_INT4},
{"uint32_t", MTL_DATATYPE_UINT},
{"uvec2", MTL_DATATYPE_UINT2},
{"uvec3", MTL_DATATYPE_UINT3},
{"uvec4", MTL_DATATYPE_UINT4},
{"mat3", MTL_DATATYPE_FLOAT3x3},
{"mat4", MTL_DATATYPE_FLOAT4x4},
{"bool", MTL_DATATYPE_INT},
{"uchar", MTL_DATATYPE_UCHAR},
{"uchar2", MTL_DATATYPE_UCHAR2},
{"uchar2", MTL_DATATYPE_UCHAR3},
{"uchar4", MTL_DATATYPE_UCHAR4},
{"vec3_1010102_Unorm", MTL_DATATYPE_UINT1010102_NORM},
{"vec3_1010102_Inorm", MTL_DATATYPE_INT1010102_NORM},
};
return (glsl_builtin_types.find(type) != glsl_builtin_types.end());
}
inline bool is_matrix_type(const std::string &type)
{
/* Matrix type support. Add types as necessary. */
return (type == "mat4");
}
inline bool is_matrix_type(const shader::Type &type)
{
/* Matrix type support. Add types as necessary. */
return (type == shader::Type::MAT4 || type == shader::Type::MAT3);
}
inline int get_matrix_location_count(const std::string &type)
{
/* Matrix type support. Add types as necessary. */
if (type == "mat4") {
return 4;
}
if (type == "mat3") {
return 3;
}
return 1;
}
inline int get_matrix_location_count(const shader::Type &type)
{
/* Matrix type support. Add types as necessary. */
if (type == shader::Type::MAT4) {
return 4;
}
else if (type == shader::Type::MAT3) {
return 3;
}
return 1;
}
inline std::string get_matrix_subtype(const std::string &type)
{
if (type == "mat4") {
return "vec4";
}
return type;
}
inline shader::Type get_matrix_subtype(const shader::Type &type)
{
if (type == shader::Type::MAT4) {
return shader::Type::VEC4;
}
if (type == shader::Type::MAT3) {
return shader::Type::VEC3;
}
return type;
}
inline std::string get_attribute_conversion_function(bool *uses_conversion,
const shader::Type &type)
{
/* NOTE(Metal): Add more attribute types as required. */
if (type == shader::Type::FLOAT) {
*uses_conversion = true;
return "internal_vertex_attribute_convert_read_float";
}
else if (type == shader::Type::VEC2) {
*uses_conversion = true;
return "internal_vertex_attribute_convert_read_float2";
}
else if (type == shader::Type::VEC3) {
*uses_conversion = true;
return "internal_vertex_attribute_convert_read_float3";
}
else if (type == shader::Type::VEC4) {
*uses_conversion = true;
return "internal_vertex_attribute_convert_read_float4";
}
*uses_conversion = false;
return "";
}
inline const char *to_string(const shader::PrimitiveOut &layout)
{
switch (layout) {
case shader::PrimitiveOut::POINTS:
return "points";
case shader::PrimitiveOut::LINE_STRIP:
return "line_strip";
case shader::PrimitiveOut::TRIANGLE_STRIP:
return "triangle_strip";
default:
BLI_assert(false);
return "unknown";
}
}
inline const char *to_string(const shader::PrimitiveIn &layout)
{
switch (layout) {
case shader::PrimitiveIn::POINTS:
return "points";
case shader::PrimitiveIn::LINES:
return "lines";
case shader::PrimitiveIn::LINES_ADJACENCY:
return "lines_adjacency";
case shader::PrimitiveIn::TRIANGLES:
return "triangles";
case shader::PrimitiveIn::TRIANGLES_ADJACENCY:
return "triangles_adjacency";
default:
BLI_assert(false);
return "unknown";
}
}
inline const char *to_string(const shader::Interpolation &interp)
{
switch (interp) {
case shader::Interpolation::SMOOTH:
return "smooth";
case shader::Interpolation::FLAT:
return "flat";
case shader::Interpolation::NO_PERSPECTIVE:
return "noperspective";
default:
BLI_assert(false);
return "unkown";
}
}
inline const char *to_string_msl(const shader::Interpolation &interp)
{
switch (interp) {
case shader::Interpolation::SMOOTH:
return "[[smooth]]";
case shader::Interpolation::FLAT:
return "[[flat]]";
case shader::Interpolation::NO_PERSPECTIVE:
return "[[center_no_perspective]]";
default:
return "";
}
}
inline const char *to_string(const shader::Type &type)
{
switch (type) {
case shader::Type::FLOAT:
return "float";
case shader::Type::VEC2:
return "vec2";
case shader::Type::VEC3:
return "vec3";
case shader::Type::VEC3_101010I2:
return "vec3_1010102_Inorm";
case shader::Type::VEC4:
return "vec4";
case shader::Type::MAT3:
return "mat3";
case shader::Type::MAT4:
return "mat4";
case shader::Type::UINT:
return "uint32_t";
case shader::Type::UVEC2:
return "uvec2";
case shader::Type::UVEC3:
return "uvec3";
case shader::Type::UVEC4:
return "uvec4";
case shader::Type::INT:
return "int";
case shader::Type::IVEC2:
return "ivec2";
case shader::Type::IVEC3:
return "ivec3";
case shader::Type::IVEC4:
return "ivec4";
case shader::Type::BOOL:
return "bool";
case shader::Type::UCHAR:
return "uchar";
case shader::Type::UCHAR2:
return "uchar2";
case shader::Type::UCHAR3:
return "uchar3";
case shader::Type::UCHAR4:
return "uchar4";
case shader::Type::CHAR:
return "char";
case shader::Type::CHAR2:
return "char2";
case shader::Type::CHAR3:
return "char3";
case shader::Type::CHAR4:
return "char4";
default:
BLI_assert(false);
return "unkown";
}
}
} // namespace blender::gpu

2976
source/blender/gpu/metal/mtl_shader_generator.mm Normal file
View File

File diff suppressed because it is too large Load Diff

267
source/blender/gpu/metal/mtl_shader_interface.hh Normal file
View File

@@ -0,0 +1,267 @@
/* SPDX-License-Identifier: GPL-2.0-or-later */
/** \file
* \ingroup gpu
*/
#pragma once
#include "MEM_guardedalloc.h"
#include "BLI_vector.hh"
#include "gpu_shader_interface.hh"
#include "mtl_capabilities.hh"
#include "mtl_shader_interface_type.hh"
#include "GPU_common.h"
#include "GPU_common_types.h"
#include "GPU_texture.h"
#include "gpu_texture_private.hh"
#include <Metal/Metal.h>
#include <functional>
namespace blender::gpu {
/* MTLShaderInterface describes the layout and properties of a given shader,
* including input and output bindings, and any special properties or modes
* that the shader may require.
*
* -- Shader input/output bindings --
*
* We require custom datastructures for the binding information in Metal.
* This is because certain bindings contain and require more information to
* be stored than can be tracked solely within the `ShaderInput` struct.
* e.g. data sizes and offsets.
*
* Upon interface completion, `prepare_common_shader_inputs` is used to
* populate the global ShaderInput* array to enable correct functionality
* of shader binding location lookups. These returned locations act as indices
* into the arrays stored here in the MTLShaderInterace, such that extraction
* of required information can be performed within the backend.
*
* e.g. `int loc = GPU_shader_get_uniform(...)`
* `loc` will match the index into the MTLShaderUniform uniforms_[] array
* to fetch the required Metal specific information.
*
*
*
* -- Argument Buffers and Argument Encoders --
*
* We can use ArgumentBuffers (AB's) in Metal to extend the resource bind limitations
* by providing bindless support.
*
* Argument Buffers are used for sampler bindings when the builtin
* sampler limit of 16 is exceeded, as in all cases for Blender,
* each individual texture is associated with a given sampler, and this
* lower limit would otherwise reduce the total availability of textures
* used in shaders.
*
* In future, argument buffers may be extended to support other resource
* types, if overall bind limits are ever increased within Blender.
*
* The ArgumentEncoder cache used to store the generated ArgumentEncoders for a given
* shader permutation. The ArgumentEncoder is the resource used to write resource binding
* information to a specified buffer, and is unique to the shader's resource interface.
*/
enum class ShaderStage : uint32_t {
VERTEX = 1 << 0,
FRAGMENT = 1 << 1,
BOTH = (ShaderStage::VERTEX | ShaderStage::FRAGMENT),
};
ENUM_OPERATORS(ShaderStage, ShaderStage::BOTH);
inline uint get_shader_stage_index(ShaderStage stage)
{
switch (stage) {
case ShaderStage::VERTEX:
return 0;
case ShaderStage::FRAGMENT:
return 1;
default:
BLI_assert_unreachable();
return 0;
}
return 0;
}
/* Shader input/output binding information. */
struct MTLShaderInputAttribute {
uint32_t name_offset;
MTLVertexFormat format;
uint32_t index;
uint32_t location;
uint32_t size;
uint32_t buffer_index;
uint32_t offset;
/* For attributes of Matrix/array types, we need to insert "fake" attributes for
* each element, as matrix types are not natively supported.
*
* > 1 if matrix/arrays are used, specifying number of elements.
* = 1 for non-matrix types
* = 0 if used as a dummy slot for "fake" matrix attributes. */
uint32_t matrix_element_count;
};
struct MTLShaderUniformBlock {
uint32_t name_offset;
uint32_t size = 0;
/* Buffer resouce bind index in shader [[buffer(index)]]. */
uint32_t buffer_index;
/* Tracking for manual uniform addition. */
uint32_t current_offset;
ShaderStage stage_mask;
};
struct MTLShaderUniform {
uint32_t name_offset;
/* Index of `MTLShaderUniformBlock` this uniform belongs to. */
uint32_t size_in_bytes;
uint32_t byte_offset;
eMTLDataType type;
uint32_t array_len;
};
struct MTLShaderTexture {
bool used;
uint32_t name_offset;
/* Texture resource bind slot in shader [[texture(n)]]. */
int slot_index;
eGPUTextureType type;
ShaderStage stage_mask;
};
struct MTLShaderSampler {
uint32_t name_offset;
/* Sampler resource bind slot in shader [[sampler(n)]]. */
uint32_t slot_index = 0;
};
/* Utility Functions. */
MTLVertexFormat mtl_datatype_to_vertex_type(eMTLDataType type);
/**
* Implementation of Shader interface for Metal Backend.
**/
class MTLShaderInterface : public ShaderInterface {
private:
/* Argument encoders caching.
* Static size is based on common input permutation variations. */
static const int ARGUMENT_ENCODERS_CACHE_SIZE = 3;
struct ArgumentEncoderCacheEntry {
id<MTLArgumentEncoder> encoder;
int buffer_index;
};
ArgumentEncoderCacheEntry arg_encoders_[ARGUMENT_ENCODERS_CACHE_SIZE] = {};
/* Vertex input Attribues. */
uint32_t total_attributes_;
uint32_t total_vert_stride_;
MTLShaderInputAttribute attributes_[MTL_MAX_VERTEX_INPUT_ATTRIBUTES];
/* Uniforms. */
uint32_t total_uniforms_;
MTLShaderUniform uniforms_[MTL_MAX_UNIFORMS_PER_BLOCK];
/* Uniform Blocks. */
uint32_t total_uniform_blocks_;
MTLShaderUniformBlock ubos_[MTL_MAX_UNIFORM_BUFFER_BINDINGS];
MTLShaderUniformBlock push_constant_block_;
/* Textures. */
/* Textures support explicit binding indices, so some texture slots
* remain unused. */
uint32_t total_textures_;
int max_texture_index_;
MTLShaderTexture textures_[MTL_MAX_TEXTURE_SLOTS];
/* Whether argument buffers are used for sampler bindings. */
bool sampler_use_argument_buffer_;
int sampler_argument_buffer_bind_index_vert_;
int sampler_argument_buffer_bind_index_frag_;
/* Attribute Mask. */
uint32_t enabled_attribute_mask_;
/* Debug. */
char name[256];
public:
MTLShaderInterface(const char *name);
~MTLShaderInterface();
void init();
void add_input_attribute(uint32_t name_offset,
uint32_t attribute_location,
MTLVertexFormat format,
uint32_t buffer_index,
uint32_t size,
uint32_t offset,
int matrix_element_count = 1);
uint32_t add_uniform_block(uint32_t name_offset,
uint32_t buffer_index,
uint32_t size,
ShaderStage stage_mask = ShaderStage::BOTH);
void add_uniform(uint32_t name_offset, eMTLDataType type, int array_len = 1);
void add_texture(uint32_t name_offset,
uint32_t texture_slot,
eGPUTextureType tex_binding_type,
ShaderStage stage_mask = ShaderStage::FRAGMENT);
void add_push_constant_block(uint32_t name_offset);
/* Resolve and cache locations of builtin uniforms and uniform blocks. */
void map_builtins();
void set_sampler_properties(bool use_argument_buffer,
uint32_t argument_buffer_bind_index_vert,
uint32_t argument_buffer_bind_index_frag);
/* Prepare ShaderInput interface for binding resolution. */
void prepare_common_shader_inputs();
/* Fetch Uniforms. */
const MTLShaderUniform &get_uniform(uint index) const;
uint32_t get_total_uniforms() const;
/* Fetch Uniform Blocks. */
const MTLShaderUniformBlock &get_uniform_block(uint index) const;
uint32_t get_total_uniform_blocks() const;
bool has_uniform_block(uint32_t block_index) const;
uint32_t get_uniform_block_size(uint32_t block_index) const;
/* Push constant uniform data block should always be available. */
const MTLShaderUniformBlock &get_push_constant_block() const;
/* Fetch textures. */
const MTLShaderTexture &get_texture(uint index) const;
uint32_t get_total_textures() const;
uint32_t get_max_texture_index() const;
bool get_use_argument_buffer_for_samplers(int *vertex_arg_buffer_bind_index,
int *fragment_arg_buffer_bind_index) const;
/* Fetch Attributes. */
const MTLShaderInputAttribute &get_attribute(uint index) const;
uint32_t get_total_attributes() const;
uint32_t get_total_vertex_stride() const;
uint32_t get_enabled_attribute_mask() const;
/* Name buffer fetching. */
const char *get_name_at_offset(uint32_t offset) const;
/* Interface name. */
const char *get_name() const
{
return this->name;
}
/* Argument buffer encoder management. */
id<MTLArgumentEncoder> find_argument_encoder(int buffer_index) const;
void insert_argument_encoder(int buffer_index, id encoder);
MEM_CXX_CLASS_ALLOC_FUNCS("MTLShaderInterface");
};
} // namespace blender::gpu

604
source/blender/gpu/metal/mtl_shader_interface.mm Normal file
View File

@@ -0,0 +1,604 @@
/* SPDX-License-Identifier: GPL-2.0-or-later */
/** \file
* \ingroup gpu
*
* GPU shader interface (C --> GLSL)
*/
#include "BLI_bitmap.h"
#include "GPU_capabilities.h"
#include "mtl_common.hh"
#include "mtl_debug.hh"
#include "mtl_shader_interface.hh"
#include "mtl_shader_interface_type.hh"
#include "BLI_blenlib.h"
#include "BLI_math_base.h"
#include "BLI_utildefines.h"
#include "MEM_guardedalloc.h"
namespace blender::gpu {
MTLShaderInterface::MTLShaderInterface(const char *name)
{
/* Shared ShaderInputs array is populated later on in `prepare_common_shader_inputs`
* after Metal Shader Interface preparation. */
inputs_ = nullptr;
if (name != nullptr) {
strcpy(this->name, name);
}
/* Ensure ShaderInterface parameters are cleared. */
this->init();
}
MTLShaderInterface::~MTLShaderInterface()
{
for (const int i : IndexRange(ARGUMENT_ENCODERS_CACHE_SIZE)) {
if (arg_encoders_[i].encoder != nil) {
id<MTLArgumentEncoder> enc = arg_encoders_[i].encoder;
[enc release];
}
}
}
const char *MTLShaderInterface::get_name_at_offset(uint32_t offset) const
{
return name_buffer_ + offset;
}
void MTLShaderInterface::init()
{
total_attributes_ = 0;
total_uniform_blocks_ = 0;
total_uniforms_ = 0;
total_textures_ = 0;
max_texture_index_ = -1;
enabled_attribute_mask_ = 0;
total_vert_stride_ = 0;
sampler_use_argument_buffer_ = false;
sampler_argument_buffer_bind_index_vert_ = -1;
sampler_argument_buffer_bind_index_frag_ = -1;
/* NULL initialise uniform location markers for builtins. */
for (const int u : IndexRange(GPU_NUM_UNIFORMS)) {
builtins_[u] = -1;
}
for (const int ubo : IndexRange(GPU_NUM_UNIFORM_BLOCKS)) {
builtin_blocks_[ubo] = -1;
}
for (const int tex : IndexRange(MTL_MAX_TEXTURE_SLOTS)) {
textures_[tex].used = false;
textures_[tex].slot_index = -1;
}
/* Null initialisation for argument encoders. */
for (const int i : IndexRange(ARGUMENT_ENCODERS_CACHE_SIZE)) {
arg_encoders_[i].encoder = nil;
arg_encoders_[i].buffer_index = -1;
}
}
void MTLShaderInterface::add_input_attribute(uint32_t name_offset,
uint32_t attribute_location,
MTLVertexFormat format,
uint32_t buffer_index,
uint32_t size,
uint32_t offset,
int matrix_element_count)
{
MTLShaderInputAttribute &input_attr = attributes_[total_attributes_];
input_attr.name_offset = name_offset;
input_attr.format = format;
input_attr.location = attribute_location;
input_attr.size = size;
input_attr.buffer_index = buffer_index;
input_attr.offset = offset;
input_attr.matrix_element_count = matrix_element_count;
input_attr.index = total_attributes_;
total_attributes_++;
total_vert_stride_ = max_ii(total_vert_stride_, offset + size);
enabled_attribute_mask_ |= (1 << attribute_location);
}
uint32_t MTLShaderInterface::add_uniform_block(uint32_t name_offset,
uint32_t buffer_index,
uint32_t size,
ShaderStage stage_mask)
{
/* Ensure Size is 16 byte aligned to guarantees alignment rules are satisfied. */
if ((size % 16) != 0) {
size += 16 - (size % 16);
}
MTLShaderUniformBlock &uni_block = ubos_[total_uniform_blocks_];
uni_block.name_offset = name_offset;
/* We offset the buffer bidning index by one, as the first slot is reserved for push constant
* data. */
uni_block.buffer_index = buffer_index + 1;
uni_block.size = size;
uni_block.current_offset = 0;
uni_block.stage_mask = ShaderStage::BOTH;
return (total_uniform_blocks_++);
}
void MTLShaderInterface::add_push_constant_block(uint32_t name_offset)
{
push_constant_block_.name_offset = name_offset;
/* Push constant data block is always uniform buffer index 0. */
push_constant_block_.buffer_index = 0;
/* Size starts at zero and grows as uniforms are added. */
push_constant_block_.size = 0;
push_constant_block_.current_offset = 0;
push_constant_block_.stage_mask = ShaderStage::BOTH;
}
void MTLShaderInterface::add_uniform(uint32_t name_offset, eMTLDataType type, int array_len)
{
BLI_assert(array_len > 0);
BLI_assert(total_uniforms_ < MTL_MAX_UNIFORMS_PER_BLOCK);
if (total_uniforms_ >= MTL_MAX_UNIFORMS_PER_BLOCK) {
MTL_LOG_WARNING(
"[Warning] Cannot add uniform '%s' to shader interface '%s' as the uniform limit of %d "
"has been reached.\n",
name,
name,
MTL_MAX_UNIFORMS_PER_BLOCK);
return;
}
MTLShaderUniform &uniform = uniforms_[total_uniforms_];
uniform.name_offset = name_offset;
/* Determine size and offset alignment -- C++ struct alignment rules: Base address of value must
* match alignment of type. GLSL follows minimum type alignment of 4. */
int data_type_size = mtl_get_data_type_size(type) * array_len;
int data_type_alignment = max_ii(mtl_get_data_type_alignment(type), 4);
int current_offset = push_constant_block_.current_offset;
if ((current_offset % data_type_alignment) != 0) {
current_offset += data_type_alignment - (current_offset % data_type_alignment);
}
uniform.size_in_bytes = data_type_size;
uniform.byte_offset = current_offset;
uniform.type = type;
uniform.array_len = array_len;
total_uniforms_++;
/* Update Push constant block-- update offset, re-size and re-align total memory requirement to
* be 16-byte aligned. Following GLSL std140. */
push_constant_block_.current_offset = current_offset + data_type_size;
if (push_constant_block_.current_offset > push_constant_block_.size) {
push_constant_block_.size = push_constant_block_.current_offset;
if ((push_constant_block_.size % 16) != 0) {
push_constant_block_.size += 16 - (push_constant_block_.size % 16);
}
}
/* Validate properties. */
BLI_assert(uniform.size_in_bytes > 0);
BLI_assert_msg(
current_offset + data_type_size <= push_constant_block_.size,
"Uniform size and offset sits outside the specified size range for the uniform block");
}
void MTLShaderInterface::add_texture(uint32_t name_offset,
uint32_t texture_slot,
eGPUTextureType tex_binding_type,
ShaderStage stage_mask)
{
BLI_assert(texture_slot >= 0 && texture_slot < GPU_max_textures());
if (texture_slot >= 0 && texture_slot < GPU_max_textures()) {
MTLShaderTexture &tex = textures_[texture_slot];
BLI_assert_msg(tex.used == false, "Texture slot already in-use by another binding");
tex.name_offset = name_offset;
tex.slot_index = texture_slot;
tex.type = tex_binding_type;
tex.stage_mask = stage_mask;
tex.used = true;
total_textures_++;
max_texture_index_ = max_ii(max_texture_index_, texture_slot);
}
else {
BLI_assert_msg(false, "Exceeding maximum supported texture count.");
MTL_LOG_WARNING(
"Could not add additional texture with index %d to shader interface. Maximum "
"supported texture count is %d\n",
texture_slot,
GPU_max_textures());
}
}
void MTLShaderInterface::map_builtins()
{
/* Clear builtin arrays to NULL locations. */
for (const int u : IndexRange(GPU_NUM_UNIFORMS)) {
builtins_[u] = -1;
}
for (const int ubo : IndexRange(GPU_NUM_UNIFORM_BLOCKS)) {
builtin_blocks_[ubo] = -1;
}
/* Resolve and cache uniform locations for bultin uniforms. */
for (const int u : IndexRange(GPU_NUM_UNIFORMS)) {
const ShaderInput *uni = this->uniform_get(builtin_uniform_name((GPUUniformBuiltin)u));
if (uni != nullptr) {
BLI_assert(uni->location >= 0);
if (uni->location >= 0) {
builtins_[u] = uni->location;
MTL_LOG_INFO("Mapped builtin uniform '%s' NB: '%s' to location: %d\n",
builtin_uniform_name((GPUUniformBuiltin)u),
get_name_at_offset(uni->name_offset),
uni->location);
}
}
}
/* Resolve and cache uniform locations for bultin uniform blocks. */
for (const int u : IndexRange(GPU_NUM_UNIFORM_BLOCKS)) {
const ShaderInput *uni = this->ubo_get(builtin_uniform_block_name((GPUUniformBlockBuiltin)u));
if (uni != nullptr) {
BLI_assert(uni->location >= 0);
if (uni->location >= 0) {
builtin_blocks_[u] = uni->binding;
MTL_LOG_INFO("Mapped builtin uniform block '%s' to location %d\n",
builtin_uniform_block_name((GPUUniformBlockBuiltin)u),
uni->location);
}
}
}
}
/* Populate ShaderInput struct based on interface. */
void MTLShaderInterface::prepare_common_shader_inputs()
{
/* ShaderInput inputs_ maps a uniform name to an external
* uniform location, which is used as an array index to look-up
* information in the local MTLShaderInterface input structs.
*
* ShaderInput population follows the ordering rules in gpu_shader_interface. */
/* Populate ShaderInterface counts. */
attr_len_ = this->get_total_attributes();
ubo_len_ = this->get_total_uniform_blocks();
uniform_len_ = this->get_total_uniforms() + this->get_total_textures();
/* TODO(Metal): Support storage buffer bindings. Pending compute shader support. */
ssbo_len_ = 0;
/* Calculate total inputs and allocate ShaderInput array. */
/* NOTE: We use the existing name_buffer_ allocated for internal input structs. */
int input_tot_len = attr_len_ + ubo_len_ + uniform_len_ + ssbo_len_;
inputs_ = (ShaderInput *)MEM_callocN(sizeof(ShaderInput) * input_tot_len, __func__);
ShaderInput *current_input = inputs_;
/* Attributes. */
for (const int attr_index : IndexRange(total_attributes_)) {
MTLShaderInputAttribute &shd_attr = attributes_[attr_index];
current_input->name_offset = shd_attr.name_offset;
current_input->name_hash = BLI_hash_string(this->get_name_at_offset(shd_attr.name_offset));
current_input->location = attr_index;
current_input->binding = attr_index;
current_input++;
}
/* UBOs. */
BLI_assert(&inputs_[attr_len_] >= current_input);
current_input = &inputs_[attr_len_];
for (const int ubo_index : IndexRange(total_uniform_blocks_)) {
MTLShaderUniformBlock &shd_ubo = ubos_[ubo_index];
current_input->name_offset = shd_ubo.name_offset;
current_input->name_hash = BLI_hash_string(this->get_name_at_offset(shd_ubo.name_offset));
/* Location refers to the index in the ubos_ array. */
current_input->location = ubo_index;
/* Final binding location refers to the buffer binding index within the shader (Relative to
* MTL_uniform_buffer_base_index). */
current_input->binding = shd_ubo.buffer_index;
current_input++;
}
/* Uniforms. */
BLI_assert(&inputs_[attr_len_ + ubo_len_] >= current_input);
current_input = &inputs_[attr_len_ + ubo_len_];
for (const int uniform_index : IndexRange(total_uniforms_)) {
MTLShaderUniform &shd_uni = uniforms_[uniform_index];
current_input->name_offset = shd_uni.name_offset;
current_input->name_hash = BLI_hash_string(this->get_name_at_offset(shd_uni.name_offset));
current_input->location = uniform_index;
current_input->binding = uniform_index;
current_input++;
}
/* Textures.
* NOTE(Metal): Textures are externally treated as uniforms in gpu_shader_interface.
* Location for textures resolved as `binding` value. This
* is the index into the local MTLShaderTexture textures[] array.
*
* In MSL, we cannot trivially remap which texture slot a given texture
* handle points to, unlike in GLSL, where a uniform sampler/image can be updated
* and queried as both a texture and a uniform. */
for (int texture_index = 0; texture_index <= max_texture_index_; texture_index++) {
const MTLShaderTexture &shd_tex = textures_[texture_index];
/* Not all texture entries are used when explicit texture locations are specified. */
if (shd_tex.used) {
BLI_assert_msg(shd_tex.slot_index == texture_index,
"Texture binding slot should match array index for texture.");
current_input->name_offset = shd_tex.name_offset;
current_input->name_hash = BLI_hash_string(this->get_name_at_offset(shd_tex.name_offset));
/* Location represents look-up address.
* For Metal, this location is a unique value offset by
* total_uniforms such that it does not overlap.
*
* This range offset allows a check in the uniform look-up
* to ensure texture handles are not treated as standard uniforms in Metal. */
current_input->location = texture_index + total_uniforms_;
/* Binding represents texture slot [[texture(n)]]. */
current_input->binding = shd_tex.slot_index;
current_input++;
}
}
/* SSBO bindings.
* TODO(Metal): Support SSBOs. Pending compute support. */
BLI_assert(&inputs_[attr_len_ + ubo_len_ + uniform_len_] >= current_input);
current_input = &inputs_[attr_len_ + ubo_len_ + uniform_len_];
/* Map builtin uniform indices to uniform binding locations. */
this->map_builtins();
}
void MTLShaderInterface::set_sampler_properties(bool use_argument_buffer,
uint32_t argument_buffer_bind_index_vert,
uint32_t argument_buffer_bind_index_frag)
{
sampler_use_argument_buffer_ = use_argument_buffer;
sampler_argument_buffer_bind_index_vert_ = argument_buffer_bind_index_vert;
sampler_argument_buffer_bind_index_frag_ = argument_buffer_bind_index_frag;
}
/* Attributes. */
const MTLShaderInputAttribute &MTLShaderInterface::get_attribute(uint index) const
{
BLI_assert(index < MTL_MAX_VERTEX_INPUT_ATTRIBUTES);
BLI_assert(index < get_total_attributes());
return attributes_[index];
}
uint32_t MTLShaderInterface::get_total_attributes() const
{
return total_attributes_;
}
uint32_t MTLShaderInterface::get_total_vertex_stride() const
{
return total_vert_stride_;
}
uint32_t MTLShaderInterface::get_enabled_attribute_mask() const
{
return enabled_attribute_mask_;
}
/* Uniforms. */
const MTLShaderUniform &MTLShaderInterface::get_uniform(uint index) const
{
BLI_assert(index < MTL_MAX_UNIFORMS_PER_BLOCK);
BLI_assert(index < get_total_uniforms());
return uniforms_[index];
}
uint32_t MTLShaderInterface::get_total_uniforms() const
{
return total_uniforms_;
}
/* Uniform Blocks. */
const MTLShaderUniformBlock &MTLShaderInterface::get_uniform_block(uint index) const
{
BLI_assert(index < MTL_MAX_UNIFORM_BUFFER_BINDINGS);
BLI_assert(index < get_total_uniform_blocks());
return ubos_[index];
}
const MTLShaderUniformBlock &MTLShaderInterface::get_push_constant_block() const
{
return push_constant_block_;
}
uint32_t MTLShaderInterface::get_total_uniform_blocks() const
{
return total_uniform_blocks_;
}
bool MTLShaderInterface::has_uniform_block(uint32_t block_index) const
{
return (block_index < total_uniform_blocks_);
}
uint32_t MTLShaderInterface::get_uniform_block_size(uint32_t block_index) const
{
return (block_index < total_uniform_blocks_) ? ubos_[block_index].size : 0;
}
/* Textures. */
const MTLShaderTexture &MTLShaderInterface::get_texture(uint index) const
{
BLI_assert(index < MTL_MAX_TEXTURE_SLOTS);
BLI_assert(index <= get_max_texture_index());
return textures_[index];
}
uint32_t MTLShaderInterface::get_total_textures() const
{
return total_textures_;
}
uint32_t MTLShaderInterface::get_max_texture_index() const
{
return max_texture_index_;
}
bool MTLShaderInterface::get_use_argument_buffer_for_samplers(
int *vertex_arg_buffer_bind_index, int *fragment_arg_buffer_bind_index) const
{
/* Returns argument buffer binding slot for each shader stage.
* The exact bind slot may be different, as each stage has different buffer inputs. */
*vertex_arg_buffer_bind_index = sampler_argument_buffer_bind_index_vert_;
*fragment_arg_buffer_bind_index = sampler_argument_buffer_bind_index_frag_;
return sampler_use_argument_buffer_;
}
id<MTLArgumentEncoder> MTLShaderInterface::find_argument_encoder(int buffer_index) const
{
id encoder = nil;
for (const int i : IndexRange(ARGUMENT_ENCODERS_CACHE_SIZE)) {
encoder = arg_encoders_[i].buffer_index == buffer_index ? arg_encoders_[i].encoder : encoder;
}
return encoder;
}
void MTLShaderInterface::insert_argument_encoder(int buffer_index, id encoder)
{
for (const int i : IndexRange(ARGUMENT_ENCODERS_CACHE_SIZE)) {
if (arg_encoders_[i].encoder == nil) {
arg_encoders_[i].encoder = encoder;
arg_encoders_[i].buffer_index = buffer_index;
return;
}
}
MTL_LOG_WARNING("could not insert encoder into cache!");
}
MTLVertexFormat mtl_datatype_to_vertex_type(eMTLDataType type)
{
switch (type) {
case MTL_DATATYPE_CHAR:
return MTLVertexFormatChar;
case MTL_DATATYPE_UCHAR:
return MTLVertexFormatUChar;
case MTL_DATATYPE_BOOL:
return MTLVertexFormatUChar;
case MTL_DATATYPE_CHAR2:
return MTLVertexFormatChar2;
case MTL_DATATYPE_UCHAR2:
return MTLVertexFormatUChar2;
case MTL_DATATYPE_BOOL2:
return MTLVertexFormatUChar2;
case MTL_DATATYPE_SHORT:
return MTLVertexFormatShort;
case MTL_DATATYPE_USHORT:
return MTLVertexFormatUShort;
case MTL_DATATYPE_CHAR3:
return MTLVertexFormatChar3;
case MTL_DATATYPE_UCHAR3:
return MTLVertexFormatUChar3;
case MTL_DATATYPE_BOOL3:
return MTLVertexFormatUChar3;
case MTL_DATATYPE_CHAR4:
return MTLVertexFormatChar4;
case MTL_DATATYPE_UCHAR4:
return MTLVertexFormatUChar4;
case MTL_DATATYPE_INT:
return MTLVertexFormatInt;
case MTL_DATATYPE_UINT:
return MTLVertexFormatUInt;
case MTL_DATATYPE_BOOL4:
return MTLVertexFormatUChar4;
case MTL_DATATYPE_SHORT2:
return MTLVertexFormatShort2;
case MTL_DATATYPE_USHORT2:
return MTLVertexFormatUShort2;
case MTL_DATATYPE_FLOAT:
return MTLVertexFormatFloat;
case MTL_DATATYPE_HALF2x2:
case MTL_DATATYPE_HALF3x2:
case MTL_DATATYPE_HALF4x2:
BLI_assert_msg(false, "Unsupported raw vertex attribute types in Blender.");
return MTLVertexFormatInvalid;
case MTL_DATATYPE_SHORT3:
return MTLVertexFormatShort3;
case MTL_DATATYPE_USHORT3:
return MTLVertexFormatUShort3;
case MTL_DATATYPE_SHORT4:
return MTLVertexFormatShort4;
case MTL_DATATYPE_USHORT4:
return MTLVertexFormatUShort4;
case MTL_DATATYPE_INT2:
return MTLVertexFormatInt2;
case MTL_DATATYPE_UINT2:
return MTLVertexFormatUInt2;
case MTL_DATATYPE_FLOAT2:
return MTLVertexFormatFloat2;
case MTL_DATATYPE_LONG:
return MTLVertexFormatInt;
case MTL_DATATYPE_ULONG:
return MTLVertexFormatUInt;
case MTL_DATATYPE_HALF2x3:
case MTL_DATATYPE_HALF2x4:
case MTL_DATATYPE_HALF3x3:
case MTL_DATATYPE_HALF3x4:
case MTL_DATATYPE_HALF4x3:
case MTL_DATATYPE_HALF4x4:
case MTL_DATATYPE_FLOAT2x2:
case MTL_DATATYPE_FLOAT3x2:
case MTL_DATATYPE_FLOAT4x2:
BLI_assert_msg(false, "Unsupported raw vertex attribute types in Blender.");
return MTLVertexFormatInvalid;
case MTL_DATATYPE_INT3:
return MTLVertexFormatInt3;
case MTL_DATATYPE_INT4:
return MTLVertexFormatInt4;
case MTL_DATATYPE_UINT3:
return MTLVertexFormatUInt3;
case MTL_DATATYPE_UINT4:
return MTLVertexFormatUInt4;
case MTL_DATATYPE_FLOAT3:
return MTLVertexFormatFloat3;
case MTL_DATATYPE_FLOAT4:
return MTLVertexFormatFloat4;
case MTL_DATATYPE_LONG2:
return MTLVertexFormatInt2;
case MTL_DATATYPE_ULONG2:
return MTLVertexFormatUInt2;
case MTL_DATATYPE_FLOAT2x3:
case MTL_DATATYPE_FLOAT2x4:
case MTL_DATATYPE_FLOAT3x3:
case MTL_DATATYPE_FLOAT3x4:
case MTL_DATATYPE_FLOAT4x3:
case MTL_DATATYPE_FLOAT4x4:
BLI_assert_msg(false, "Unsupported raw vertex attribute types in Blender.");
return MTLVertexFormatInvalid;
case MTL_DATATYPE_LONG3:
return MTLVertexFormatInt3;
case MTL_DATATYPE_LONG4:
return MTLVertexFormatInt4;
case MTL_DATATYPE_ULONG3:
return MTLVertexFormatUInt3;
case MTL_DATATYPE_ULONG4:
return MTLVertexFormatUInt4;
/* Special Types */
case MTL_DATATYPE_UINT1010102_NORM:
return MTLVertexFormatUInt1010102Normalized;
case MTL_DATATYPE_INT1010102_NORM:
return MTLVertexFormatInt1010102Normalized;
default:
BLI_assert(false);
return MTLVertexFormatInvalid;
};
}
} // namespace blender::gpu

251
source/blender/gpu/metal/mtl_shader_interface_type.hh Normal file
View File

@@ -0,0 +1,251 @@
/* SPDX-License-Identifier: GPL-2.0-or-later */
/** \file
* \ingroup gpu
*/
#pragma once
#include "BLI_assert.h"
enum eMTLDataType {
MTL_DATATYPE_CHAR,
MTL_DATATYPE_CHAR2,
MTL_DATATYPE_CHAR3,
MTL_DATATYPE_CHAR4,
MTL_DATATYPE_UCHAR,
MTL_DATATYPE_UCHAR2,
MTL_DATATYPE_UCHAR3,
MTL_DATATYPE_UCHAR4,
MTL_DATATYPE_BOOL,
MTL_DATATYPE_BOOL2,
MTL_DATATYPE_BOOL3,
MTL_DATATYPE_BOOL4,
MTL_DATATYPE_SHORT,
MTL_DATATYPE_SHORT2,
MTL_DATATYPE_SHORT3,
MTL_DATATYPE_SHORT4,
MTL_DATATYPE_USHORT,
MTL_DATATYPE_USHORT2,
MTL_DATATYPE_USHORT3,
MTL_DATATYPE_USHORT4,
MTL_DATATYPE_INT,
MTL_DATATYPE_INT2,
MTL_DATATYPE_INT3,
MTL_DATATYPE_INT4,
MTL_DATATYPE_UINT,
MTL_DATATYPE_UINT2,
MTL_DATATYPE_UINT3,
MTL_DATATYPE_UINT4,
MTL_DATATYPE_FLOAT,
MTL_DATATYPE_FLOAT2,
MTL_DATATYPE_FLOAT3,
MTL_DATATYPE_FLOAT4,
MTL_DATATYPE_LONG,
MTL_DATATYPE_LONG2,
MTL_DATATYPE_LONG3,
MTL_DATATYPE_LONG4,
MTL_DATATYPE_ULONG,
MTL_DATATYPE_ULONG2,
MTL_DATATYPE_ULONG3,
MTL_DATATYPE_ULONG4,
MTL_DATATYPE_HALF2x2,
MTL_DATATYPE_HALF2x3,
MTL_DATATYPE_HALF2x4,
MTL_DATATYPE_HALF3x2,
MTL_DATATYPE_HALF3x3,
MTL_DATATYPE_HALF3x4,
MTL_DATATYPE_HALF4x2,
MTL_DATATYPE_HALF4x3,
MTL_DATATYPE_HALF4x4,
MTL_DATATYPE_FLOAT2x2,
MTL_DATATYPE_FLOAT2x3,
MTL_DATATYPE_FLOAT2x4,
MTL_DATATYPE_FLOAT3x2,
MTL_DATATYPE_FLOAT3x3,
MTL_DATATYPE_FLOAT3x4,
MTL_DATATYPE_FLOAT4x2,
MTL_DATATYPE_FLOAT4x3,
MTL_DATATYPE_FLOAT4x4,
MTL_DATATYPE_UINT1010102_NORM,
MTL_DATATYPE_INT1010102_NORM
};
inline uint mtl_get_data_type_size(eMTLDataType type)
{
switch (type) {
case MTL_DATATYPE_CHAR:
case MTL_DATATYPE_UCHAR:
case MTL_DATATYPE_BOOL:
return 1;
case MTL_DATATYPE_CHAR2:
case MTL_DATATYPE_UCHAR2:
case MTL_DATATYPE_BOOL2:
case MTL_DATATYPE_SHORT:
case MTL_DATATYPE_USHORT:
return 2;
case MTL_DATATYPE_CHAR3:
case MTL_DATATYPE_UCHAR3:
case MTL_DATATYPE_BOOL3:
return 3;
case MTL_DATATYPE_CHAR4:
case MTL_DATATYPE_UCHAR4:
case MTL_DATATYPE_INT:
case MTL_DATATYPE_UINT:
case MTL_DATATYPE_BOOL4:
case MTL_DATATYPE_SHORT2:
case MTL_DATATYPE_USHORT2:
case MTL_DATATYPE_FLOAT:
case MTL_DATATYPE_UINT1010102_NORM:
case MTL_DATATYPE_INT1010102_NORM:
return 4;
case MTL_DATATYPE_SHORT3:
case MTL_DATATYPE_USHORT3:
case MTL_DATATYPE_SHORT4:
case MTL_DATATYPE_USHORT4:
case MTL_DATATYPE_INT2:
case MTL_DATATYPE_UINT2:
case MTL_DATATYPE_FLOAT2:
case MTL_DATATYPE_LONG:
case MTL_DATATYPE_ULONG:
case MTL_DATATYPE_HALF2x2:
return 8;
case MTL_DATATYPE_HALF3x2:
return 12;
case MTL_DATATYPE_INT3:
case MTL_DATATYPE_INT4:
case MTL_DATATYPE_UINT3:
case MTL_DATATYPE_UINT4:
case MTL_DATATYPE_FLOAT3:
case MTL_DATATYPE_FLOAT4:
case MTL_DATATYPE_LONG2:
case MTL_DATATYPE_ULONG2:
case MTL_DATATYPE_HALF2x3:
case MTL_DATATYPE_HALF2x4:
case MTL_DATATYPE_HALF4x2:
return 16;
case MTL_DATATYPE_HALF3x3:
case MTL_DATATYPE_HALF3x4:
case MTL_DATATYPE_FLOAT3x2:
return 24;
case MTL_DATATYPE_LONG3:
case MTL_DATATYPE_LONG4:
case MTL_DATATYPE_ULONG3:
case MTL_DATATYPE_ULONG4:
case MTL_DATATYPE_HALF4x3:
case MTL_DATATYPE_HALF4x4:
case MTL_DATATYPE_FLOAT2x3:
case MTL_DATATYPE_FLOAT2x4:
case MTL_DATATYPE_FLOAT4x2:
return 32;
case MTL_DATATYPE_FLOAT3x3:
case MTL_DATATYPE_FLOAT3x4:
return 48;
case MTL_DATATYPE_FLOAT4x3:
case MTL_DATATYPE_FLOAT4x4:
return 64;
default:
BLI_assert(false);
return 0;
};
}
inline uint mtl_get_data_type_alignment(eMTLDataType type)
{
switch (type) {
case MTL_DATATYPE_CHAR:
case MTL_DATATYPE_UCHAR:
case MTL_DATATYPE_BOOL:
return 1;
case MTL_DATATYPE_CHAR2:
case MTL_DATATYPE_UCHAR2:
case MTL_DATATYPE_BOOL2:
case MTL_DATATYPE_SHORT:
case MTL_DATATYPE_USHORT:
return 2;
case MTL_DATATYPE_CHAR3:
case MTL_DATATYPE_UCHAR3:
case MTL_DATATYPE_BOOL3:
return 3;
case MTL_DATATYPE_CHAR4:
case MTL_DATATYPE_UCHAR4:
case MTL_DATATYPE_INT:
case MTL_DATATYPE_UINT:
case MTL_DATATYPE_BOOL4:
case MTL_DATATYPE_SHORT2:
case MTL_DATATYPE_USHORT2:
case MTL_DATATYPE_FLOAT:
case MTL_DATATYPE_HALF2x2:
case MTL_DATATYPE_HALF3x2:
case MTL_DATATYPE_HALF4x2:
case MTL_DATATYPE_UINT1010102_NORM:
case MTL_DATATYPE_INT1010102_NORM:
return 4;
case MTL_DATATYPE_SHORT3:
case MTL_DATATYPE_USHORT3:
case MTL_DATATYPE_SHORT4:
case MTL_DATATYPE_USHORT4:
case MTL_DATATYPE_INT2:
case MTL_DATATYPE_UINT2:
case MTL_DATATYPE_FLOAT2:
case MTL_DATATYPE_LONG:
case MTL_DATATYPE_ULONG:
case MTL_DATATYPE_HALF2x3:
case MTL_DATATYPE_HALF2x4:
case MTL_DATATYPE_HALF3x3:
case MTL_DATATYPE_HALF3x4:
case MTL_DATATYPE_HALF4x3:
case MTL_DATATYPE_HALF4x4:
case MTL_DATATYPE_FLOAT2x2:
case MTL_DATATYPE_FLOAT3x2:
case MTL_DATATYPE_FLOAT4x2:
return 8;
case MTL_DATATYPE_INT3:
case MTL_DATATYPE_INT4:
case MTL_DATATYPE_UINT3:
case MTL_DATATYPE_UINT4:
case MTL_DATATYPE_FLOAT3:
case MTL_DATATYPE_FLOAT4:
case MTL_DATATYPE_LONG2:
case MTL_DATATYPE_ULONG2:
case MTL_DATATYPE_FLOAT2x3:
case MTL_DATATYPE_FLOAT2x4:
case MTL_DATATYPE_FLOAT3x3:
case MTL_DATATYPE_FLOAT3x4:
case MTL_DATATYPE_FLOAT4x3:
case MTL_DATATYPE_FLOAT4x4:
return 16;
case MTL_DATATYPE_LONG3:
case MTL_DATATYPE_LONG4:
case MTL_DATATYPE_ULONG3:
case MTL_DATATYPE_ULONG4:
return 32;
default:
BLI_assert_msg(false, "Unrecognised MTL datatype.");
return 0;
};
}

32
source/blender/gpu/metal/mtl_shader_shared.h Normal file
View File

@@ -0,0 +1,32 @@
/* SPDX-License-Identifier: GPL-2.0-or-later */
/* Global parameters. */
#define MTL_SSBO_VERTEX_FETCH_MAX_VBOS 6 /* buffer bind 0..5 */
#define MTL_SSBO_VERTEX_FETCH_IBO_INDEX MTL_SSBO_VERTEX_FETCH_MAX_VBOS
/* Add Types as needed (Also need to be added to mtl_shader.h). */
#define GPU_SHADER_ATTR_TYPE_FLOAT 0
#define GPU_SHADER_ATTR_TYPE_INT 1
#define GPU_SHADER_ATTR_TYPE_SHORT 2
#define GPU_SHADER_ATTR_TYPE_CHAR 3
#define GPU_SHADER_ATTR_TYPE_VEC2 4
#define GPU_SHADER_ATTR_TYPE_VEC3 5
#define GPU_SHADER_ATTR_TYPE_VEC4 6
#define GPU_SHADER_ATTR_TYPE_UVEC2 7
#define GPU_SHADER_ATTR_TYPE_UVEC3 8
#define GPU_SHADER_ATTR_TYPE_UVEC4 9
#define GPU_SHADER_ATTR_TYPE_IVEC2 10
#define GPU_SHADER_ATTR_TYPE_IVEC3 11
#define GPU_SHADER_ATTR_TYPE_IVEC4 12
#define GPU_SHADER_ATTR_TYPE_MAT3 13
#define GPU_SHADER_ATTR_TYPE_MAT4 14
#define GPU_SHADER_ATTR_TYPE_UCHAR_NORM 15
#define GPU_SHADER_ATTR_TYPE_UCHAR2_NORM 16
#define GPU_SHADER_ATTR_TYPE_UCHAR3_NORM 17
#define GPU_SHADER_ATTR_TYPE_UCHAR4_NORM 18
#define GPU_SHADER_ATTR_TYPE_INT1010102_NORM 19
#define GPU_SHADER_ATTR_TYPE_SHORT3_NORM 20
#define GPU_SHADER_ATTR_TYPE_CHAR2 21
#define GPU_SHADER_ATTR_TYPE_CHAR3 22
#define GPU_SHADER_ATTR_TYPE_CHAR4 23
#define GPU_SHADER_ATTR_TYPE_UINT 24

14
source/blender/gpu/metal/mtl_state.hh
View File

@@ -3,6 +3,7 @@
/** \file
* \ingroup gpu
*/
#pragma once
#include "MEM_guardedalloc.h"
@@ -11,6 +12,8 @@
#include "GPU_state.h"
#include "gpu_state_private.hh"
#include "mtl_pso_descriptor_state.hh"
namespace blender::gpu {
/* Forward Declarations. */
@@ -21,7 +24,7 @@ class MTLContext;
* Metal Implementation.
**/
class MTLStateManager : public StateManager {
public:
private:
/* Current state of the associated MTLContext.
* Avoids resetting the whole state for every change. */
@@ -29,6 +32,9 @@ class MTLStateManager : public StateManager {
GPUStateMutable current_mutable_;
MTLContext *context_;
/* Global pipeline descriptors. */
MTLRenderPipelineStateDescriptor pipeline_descriptor_;
public:
MTLStateManager(MTLContext *ctx);
@@ -47,6 +53,12 @@ class MTLStateManager : public StateManager {
void texture_unpack_row_length_set(uint len) override;
/* Global pipeline descriptors. */
MTLRenderPipelineStateDescriptor &get_pipeline_descriptor()
{
return pipeline_descriptor_;
}
private:
void set_write_mask(const eGPUWriteMask value);
void set_depth_test(const eGPUDepthTest value);

1
source/blender/gpu/metal/mtl_state.mm
View File

@@ -11,6 +11,7 @@
#include "mtl_context.hh"
#include "mtl_framebuffer.hh"
#include "mtl_shader_interface_type.hh"
#include "mtl_state.hh"
namespace blender::gpu {

30
source/blender/gpu/metal/mtl_texture.hh
View File

@@ -363,20 +363,20 @@ class MTLTexture : public Texture {
};
id<MTLComputePipelineState> texture_update_1d_get_kernel(
TextureUpdateRoutineSpecialisation specialisation);
TextureUpdateRoutineSpecialisation specialization);
id<MTLComputePipelineState> texture_update_1d_array_get_kernel(
TextureUpdateRoutineSpecialisation specialisation);
TextureUpdateRoutineSpecialisation specialization);
id<MTLComputePipelineState> texture_update_2d_get_kernel(
TextureUpdateRoutineSpecialisation specialisation);
TextureUpdateRoutineSpecialisation specialization);
id<MTLComputePipelineState> texture_update_2d_array_get_kernel(
TextureUpdateRoutineSpecialisation specialisation);
TextureUpdateRoutineSpecialisation specialization);
id<MTLComputePipelineState> texture_update_3d_get_kernel(
TextureUpdateRoutineSpecialisation specialisation);
TextureUpdateRoutineSpecialisation specialization);
id<MTLComputePipelineState> mtl_texture_update_impl(
TextureUpdateRoutineSpecialisation specialisation_params,
TextureUpdateRoutineSpecialisation specialization_params,
blender::Map<TextureUpdateRoutineSpecialisation, id<MTLComputePipelineState>>
&specialisation_cache,
&specialization_cache,
eGPUTextureType texture_type);
/* Depth Update Utilities */
@@ -384,7 +384,7 @@ class MTLTexture : public Texture {
* use a compute shader to write to depth, so we must instead render to a depth target.
* These processes use vertex/fragment shaders to render texture data from an intermediate
* source, in order to prime the depth buffer*/
GPUShader *depth_2d_update_sh_get(DepthTextureUpdateRoutineSpecialisation specialisation);
GPUShader *depth_2d_update_sh_get(DepthTextureUpdateRoutineSpecialisation specialization);
void update_sub_depth_2d(
int mip, int offset[3], int extent[3], eGPUDataFormat type, const void *data);
@@ -397,20 +397,20 @@ class MTLTexture : public Texture {
};
id<MTLComputePipelineState> texture_read_1d_get_kernel(
TextureReadRoutineSpecialisation specialisation);
TextureReadRoutineSpecialisation specialization);
id<MTLComputePipelineState> texture_read_1d_array_get_kernel(
TextureReadRoutineSpecialisation specialisation);
TextureReadRoutineSpecialisation specialization);
id<MTLComputePipelineState> texture_read_2d_get_kernel(
TextureReadRoutineSpecialisation specialisation);
TextureReadRoutineSpecialisation specialization);
id<MTLComputePipelineState> texture_read_2d_array_get_kernel(
TextureReadRoutineSpecialisation specialisation);
TextureReadRoutineSpecialisation specialization);
id<MTLComputePipelineState> texture_read_3d_get_kernel(
TextureReadRoutineSpecialisation specialisation);
TextureReadRoutineSpecialisation specialization);
id<MTLComputePipelineState> mtl_texture_read_impl(
TextureReadRoutineSpecialisation specialisation_params,
TextureReadRoutineSpecialisation specialization_params,
blender::Map<TextureReadRoutineSpecialisation, id<MTLComputePipelineState>>
&specialisation_cache,
&specialization_cache,
eGPUTextureType texture_type);
/* fullscreen blit utilities. */

22
source/blender/gpu/metal/mtl_texture.mm
View File

@@ -479,8 +479,8 @@ void gpu::MTLTexture::update_sub(
int expected_dst_bytes_per_pixel = get_mtl_format_bytesize(destination_format);
int destination_num_channels = get_mtl_format_num_components(destination_format);
/* Prepare specialisation struct (For texture update routine). */
TextureUpdateRoutineSpecialisation compute_specialisation_kernel = {
/* Prepare specialization struct (For texture update routine). */
TextureUpdateRoutineSpecialisation compute_specialization_kernel = {
tex_data_format_to_msl_type_str(type), /* INPUT DATA FORMAT */
tex_data_format_to_msl_texture_template_type(type), /* TEXTURE DATA FORMAT */
num_channels,
@@ -620,7 +620,7 @@ void gpu::MTLTexture::update_sub(
/* Use Compute Based update. */
if (type_ == GPU_TEXTURE_1D) {
id<MTLComputePipelineState> pso = texture_update_1d_get_kernel(
compute_specialisation_kernel);
compute_specialization_kernel);
TextureUpdateParams params = {mip,
{extent[0], 1, 1},
{offset[0], 0, 0},
@@ -637,7 +637,7 @@ void gpu::MTLTexture::update_sub(
}
else if (type_ == GPU_TEXTURE_1D_ARRAY) {
id<MTLComputePipelineState> pso = texture_update_1d_array_get_kernel(
compute_specialisation_kernel);
compute_specialization_kernel);
TextureUpdateParams params = {mip,
{extent[0], extent[1], 1},
{offset[0], offset[1], 0},
@@ -694,7 +694,7 @@ void gpu::MTLTexture::update_sub(
/* Use Compute texture update. */
if (type_ == GPU_TEXTURE_2D) {
id<MTLComputePipelineState> pso = texture_update_2d_get_kernel(
compute_specialisation_kernel);
compute_specialization_kernel);
TextureUpdateParams params = {mip,
{extent[0], extent[1], 1},
{offset[0], offset[1], 0},
@@ -712,7 +712,7 @@ void gpu::MTLTexture::update_sub(
}
else if (type_ == GPU_TEXTURE_2D_ARRAY) {
id<MTLComputePipelineState> pso = texture_update_2d_array_get_kernel(
compute_specialisation_kernel);
compute_specialization_kernel);
TextureUpdateParams params = {mip,
{extent[0], extent[1], extent[2]},
{offset[0], offset[1], offset[2]},
@@ -752,7 +752,7 @@ void gpu::MTLTexture::update_sub(
}
else {
id<MTLComputePipelineState> pso = texture_update_3d_get_kernel(
compute_specialisation_kernel);
compute_specialization_kernel);
TextureUpdateParams params = {mip,
{extent[0], extent[1], extent[2]},
{offset[0], offset[1], offset[2]},
@@ -1216,7 +1216,7 @@ void gpu::MTLTexture::read_internal(int mip,
destination_buffer_host_ptr = (void *)((uint8_t *)([destination_buffer contents]) +
destination_offset);
/* Prepare specialisation struct (For non-trivial texture read routine). */
/* Prepare specialization struct (For non-trivial texture read routine). */
int depth_format_mode = 0;
if (is_depth_format) {
depth_format_mode = 1;
@@ -1236,7 +1236,7 @@ void gpu::MTLTexture::read_internal(int mip,
}
}
TextureReadRoutineSpecialisation compute_specialisation_kernel = {
TextureReadRoutineSpecialisation compute_specialization_kernel = {
tex_data_format_to_msl_texture_template_type(data_format), /* TEXTURE DATA TYPE */
tex_data_format_to_msl_type_str(desired_output_format), /* OUTPUT DATA TYPE */
num_channels, /* TEXTURE COMPONENT COUNT */
@@ -1283,7 +1283,7 @@ void gpu::MTLTexture::read_internal(int mip,
id<MTLComputeCommandEncoder> compute_encoder =
ctx->main_command_buffer.ensure_begin_compute_encoder();
id<MTLComputePipelineState> pso = texture_read_2d_get_kernel(
compute_specialisation_kernel);
compute_specialization_kernel);
TextureReadParams params = {
mip,
{width, height, 1},
@@ -1339,7 +1339,7 @@ void gpu::MTLTexture::read_internal(int mip,
id<MTLComputeCommandEncoder> compute_encoder =
ctx->main_command_buffer.ensure_begin_compute_encoder();
id<MTLComputePipelineState> pso = texture_read_2d_array_get_kernel(
compute_specialisation_kernel);
compute_specialization_kernel);
TextureReadParams params = {
mip,
{width, height, depth},

108
source/blender/gpu/metal/mtl_texture_util.mm
View File

@@ -305,13 +305,13 @@ bool mtl_format_supports_blending(MTLPixelFormat format)
* \{ */
id<MTLComputePipelineState> gpu::MTLTexture::mtl_texture_update_impl(
TextureUpdateRoutineSpecialisation specialisation_params,
TextureUpdateRoutineSpecialisation specialization_params,
blender::Map<TextureUpdateRoutineSpecialisation, id<MTLComputePipelineState>>
&specialisation_cache,
&specialization_cache,
eGPUTextureType texture_type)
{
/* Check whether the Kernel exists. */
id<MTLComputePipelineState> *result = specialisation_cache.lookup_ptr(specialisation_params);
id<MTLComputePipelineState> *result = specialization_cache.lookup_ptr(specialization_params);
if (result != nullptr) {
return *result;
}
@@ -332,18 +332,18 @@ id<MTLComputePipelineState> gpu::MTLTexture::mtl_texture_update_impl(
options.languageVersion = MTLLanguageVersion2_2;
options.preprocessorMacros = @{
@"INPUT_DATA_TYPE" :
[NSString stringWithUTF8String:specialisation_params.input_data_type.c_str()],
[NSString stringWithUTF8String:specialization_params.input_data_type.c_str()],
@"OUTPUT_DATA_TYPE" :
[NSString stringWithUTF8String:specialisation_params.output_data_type.c_str()],
[NSString stringWithUTF8String:specialization_params.output_data_type.c_str()],
@"COMPONENT_COUNT_INPUT" :
[NSNumber numberWithInt:specialisation_params.component_count_input],
[NSNumber numberWithInt:specialization_params.component_count_input],
@"COMPONENT_COUNT_OUTPUT" :
[NSNumber numberWithInt:specialisation_params.component_count_output],
[NSNumber numberWithInt:specialization_params.component_count_output],
@"TEX_TYPE" : [NSNumber numberWithInt:((int)(texture_type))]
};
/* Prepare shader library for conversion routine. */
NSError *error = NULL;
NSError *error = nullptr;
id<MTLLibrary> temp_lib = [[ctx->device newLibraryWithSource:tex_update_kernel_src
options:options
error:&error] autorelease];
@@ -370,7 +370,7 @@ id<MTLComputePipelineState> gpu::MTLTexture::mtl_texture_update_impl(
/* Store PSO. */
[compute_pso retain];
specialisation_cache.add_new(specialisation_params, compute_pso);
specialization_cache.add_new(specialization_params, compute_pso);
return_pso = compute_pso;
}
@@ -379,53 +379,53 @@ id<MTLComputePipelineState> gpu::MTLTexture::mtl_texture_update_impl(
}
id<MTLComputePipelineState> gpu::MTLTexture::texture_update_1d_get_kernel(
TextureUpdateRoutineSpecialisation specialisation)
TextureUpdateRoutineSpecialisation specialization)
{
MTLContext *mtl_context = static_cast<MTLContext *>(unwrap(GPU_context_active_get()));
BLI_assert(mtl_context != nullptr);
return mtl_texture_update_impl(specialisation,
return mtl_texture_update_impl(specialization,
mtl_context->get_texture_utils().texture_1d_update_compute_psos,
GPU_TEXTURE_1D);
}
id<MTLComputePipelineState> gpu::MTLTexture::texture_update_1d_array_get_kernel(
TextureUpdateRoutineSpecialisation specialisation)
TextureUpdateRoutineSpecialisation specialization)
{
MTLContext *mtl_context = static_cast<MTLContext *>(unwrap(GPU_context_active_get()));
BLI_assert(mtl_context != nullptr);
return mtl_texture_update_impl(
specialisation,
specialization,
mtl_context->get_texture_utils().texture_1d_array_update_compute_psos,
GPU_TEXTURE_1D_ARRAY);
}
id<MTLComputePipelineState> gpu::MTLTexture::texture_update_2d_get_kernel(
TextureUpdateRoutineSpecialisation specialisation)
TextureUpdateRoutineSpecialisation specialization)
{
MTLContext *mtl_context = static_cast<MTLContext *>(unwrap(GPU_context_active_get()));
BLI_assert(mtl_context != nullptr);
return mtl_texture_update_impl(specialisation,
return mtl_texture_update_impl(specialization,
mtl_context->get_texture_utils().texture_2d_update_compute_psos,
GPU_TEXTURE_2D);
}
id<MTLComputePipelineState> gpu::MTLTexture::texture_update_2d_array_get_kernel(
TextureUpdateRoutineSpecialisation specialisation)
TextureUpdateRoutineSpecialisation specialization)
{
MTLContext *mtl_context = static_cast<MTLContext *>(unwrap(GPU_context_active_get()));
BLI_assert(mtl_context != nullptr);
return mtl_texture_update_impl(
specialisation,
specialization,
mtl_context->get_texture_utils().texture_2d_array_update_compute_psos,
GPU_TEXTURE_2D_ARRAY);
}
id<MTLComputePipelineState> gpu::MTLTexture::texture_update_3d_get_kernel(
TextureUpdateRoutineSpecialisation specialisation)
TextureUpdateRoutineSpecialisation specialization)
{
MTLContext *mtl_context = static_cast<MTLContext *>(unwrap(GPU_context_active_get()));
BLI_assert(mtl_context != nullptr);
return mtl_texture_update_impl(specialisation,
return mtl_texture_update_impl(specialization,
mtl_context->get_texture_utils().texture_3d_update_compute_psos,
GPU_TEXTURE_3D);
}
@@ -434,7 +434,7 @@ id<MTLComputePipelineState> gpu::MTLTexture::texture_update_3d_get_kernel(
* Currently does not appear to be hit. */
GPUShader *gpu::MTLTexture::depth_2d_update_sh_get(
DepthTextureUpdateRoutineSpecialisation specialisation)
DepthTextureUpdateRoutineSpecialisation specialization)
{
/* Check whether the Kernel exists. */
@@ -442,13 +442,13 @@ GPUShader *gpu::MTLTexture::depth_2d_update_sh_get(
BLI_assert(mtl_context != nullptr);
GPUShader **result = mtl_context->get_texture_utils().depth_2d_update_shaders.lookup_ptr(
specialisation);
specialization);
if (result != nullptr) {
return *result;
}
const char *fragment_source = nullptr;
switch (specialisation.data_mode) {
switch (specialization.data_mode) {
case MTL_DEPTH_UPDATE_MODE_FLOAT:
fragment_source = datatoc_depth_2d_update_float_frag_glsl;
break;
@@ -469,7 +469,7 @@ GPUShader *gpu::MTLTexture::depth_2d_update_sh_get(
nullptr,
nullptr,
"depth_2d_update_sh_get");
mtl_context->get_texture_utils().depth_2d_update_shaders.add_new(specialisation, shader);
mtl_context->get_texture_utils().depth_2d_update_shaders.add_new(specialization, shader);
return shader;
}
@@ -507,18 +507,18 @@ void gpu::MTLTexture::update_sub_depth_2d(
eGPUTextureFormat format = (is_float) ? GPU_R32F : GPU_R32I;
/* Shader key - Add parameters here for different configurations. */
DepthTextureUpdateRoutineSpecialisation specialisation;
DepthTextureUpdateRoutineSpecialisation specialization;
switch (type) {
case GPU_DATA_FLOAT:
specialisation.data_mode = MTL_DEPTH_UPDATE_MODE_FLOAT;
specialization.data_mode = MTL_DEPTH_UPDATE_MODE_FLOAT;
break;
case GPU_DATA_UINT_24_8:
specialisation.data_mode = MTL_DEPTH_UPDATE_MODE_INT24;
specialization.data_mode = MTL_DEPTH_UPDATE_MODE_INT24;
break;
case GPU_DATA_UINT:
specialisation.data_mode = MTL_DEPTH_UPDATE_MODE_INT32;
specialization.data_mode = MTL_DEPTH_UPDATE_MODE_INT32;
break;
default:
@@ -544,7 +544,7 @@ void gpu::MTLTexture::update_sub_depth_2d(
GPU_framebuffer_clear_stencil(depth_fb_temp, 0);
}
GPUShader *depth_2d_update_sh = depth_2d_update_sh_get(specialisation);
GPUShader *depth_2d_update_sh = depth_2d_update_sh_get(specialization);
BLI_assert(depth_2d_update_sh != nullptr);
GPUBatch *quad = GPU_batch_preset_quad();
GPU_batch_set_shader(quad, depth_2d_update_sh);
@@ -591,13 +591,13 @@ void gpu::MTLTexture::update_sub_depth_2d(
* \{ */
id<MTLComputePipelineState> gpu::MTLTexture::mtl_texture_read_impl(
TextureReadRoutineSpecialisation specialisation_params,
TextureReadRoutineSpecialisation specialization_params,
blender::Map<TextureReadRoutineSpecialisation, id<MTLComputePipelineState>>
&specialisation_cache,
&specialization_cache,
eGPUTextureType texture_type)
{
/* Check whether the Kernel exists. */
id<MTLComputePipelineState> *result = specialisation_cache.lookup_ptr(specialisation_params);
id<MTLComputePipelineState> *result = specialization_cache.lookup_ptr(specialization_params);
if (result != nullptr) {
return *result;
}
@@ -615,10 +615,10 @@ id<MTLComputePipelineState> gpu::MTLTexture::mtl_texture_read_impl(
/* Defensive Debug Checks. */
long long int depth_scale_factor = 1;
if (specialisation_params.depth_format_mode > 0) {
BLI_assert(specialisation_params.component_count_input == 1);
BLI_assert(specialisation_params.component_count_output == 1);
switch (specialisation_params.depth_format_mode) {
if (specialization_params.depth_format_mode > 0) {
BLI_assert(specialization_params.component_count_input == 1);
BLI_assert(specialization_params.component_count_output == 1);
switch (specialization_params.depth_format_mode) {
case 1:
/* FLOAT */
depth_scale_factor = 1;
@@ -642,24 +642,24 @@ id<MTLComputePipelineState> gpu::MTLTexture::mtl_texture_read_impl(
options.languageVersion = MTLLanguageVersion2_2;
options.preprocessorMacros = @{
@"INPUT_DATA_TYPE" :
[NSString stringWithUTF8String:specialisation_params.input_data_type.c_str()],
[NSString stringWithUTF8String:specialization_params.input_data_type.c_str()],
@"OUTPUT_DATA_TYPE" :
[NSString stringWithUTF8String:specialisation_params.output_data_type.c_str()],
[NSString stringWithUTF8String:specialization_params.output_data_type.c_str()],
@"COMPONENT_COUNT_INPUT" :
[NSNumber numberWithInt:specialisation_params.component_count_input],
[NSNumber numberWithInt:specialization_params.component_count_input],
@"COMPONENT_COUNT_OUTPUT" :
[NSNumber numberWithInt:specialisation_params.component_count_output],
[NSNumber numberWithInt:specialization_params.component_count_output],
@"WRITE_COMPONENT_COUNT" :
[NSNumber numberWithInt:min_ii(specialisation_params.component_count_input,
specialisation_params.component_count_output)],
[NSNumber numberWithInt:min_ii(specialization_params.component_count_input,
specialization_params.component_count_output)],
@"IS_DEPTH_FORMAT" :
[NSNumber numberWithInt:((specialisation_params.depth_format_mode > 0) ? 1 : 0)],
[NSNumber numberWithInt:((specialization_params.depth_format_mode > 0) ? 1 : 0)],
@"DEPTH_SCALE_FACTOR" : [NSNumber numberWithLongLong:depth_scale_factor],
@"TEX_TYPE" : [NSNumber numberWithInt:((int)(texture_type))]
};
/* Prepare shader library for conversion routine. */
NSError *error = NULL;
NSError *error = nullptr;
id<MTLLibrary> temp_lib = [[ctx->device newLibraryWithSource:tex_update_kernel_src
options:options
error:&error] autorelease];
@@ -687,7 +687,7 @@ id<MTLComputePipelineState> gpu::MTLTexture::mtl_texture_read_impl(
/* Store PSO. */
[compute_pso retain];
specialisation_cache.add_new(specialisation_params, compute_pso);
specialization_cache.add_new(specialization_params, compute_pso);
return_pso = compute_pso;
}
@@ -696,51 +696,51 @@ id<MTLComputePipelineState> gpu::MTLTexture::mtl_texture_read_impl(
}
id<MTLComputePipelineState> gpu::MTLTexture::texture_read_2d_get_kernel(
TextureReadRoutineSpecialisation specialisation)
TextureReadRoutineSpecialisation specialization)
{
MTLContext *mtl_context = static_cast<MTLContext *>(unwrap(GPU_context_active_get()));
BLI_assert(mtl_context != nullptr);
return mtl_texture_read_impl(specialisation,
return mtl_texture_read_impl(specialization,
mtl_context->get_texture_utils().texture_2d_read_compute_psos,
GPU_TEXTURE_2D);
}
id<MTLComputePipelineState> gpu::MTLTexture::texture_read_2d_array_get_kernel(
TextureReadRoutineSpecialisation specialisation)
TextureReadRoutineSpecialisation specialization)
{
MTLContext *mtl_context = static_cast<MTLContext *>(unwrap(GPU_context_active_get()));
BLI_assert(mtl_context != nullptr);
return mtl_texture_read_impl(specialisation,
return mtl_texture_read_impl(specialization,
mtl_context->get_texture_utils().texture_2d_array_read_compute_psos,
GPU_TEXTURE_2D_ARRAY);
}
id<MTLComputePipelineState> gpu::MTLTexture::texture_read_1d_get_kernel(
TextureReadRoutineSpecialisation specialisation)
TextureReadRoutineSpecialisation specialization)
{
MTLContext *mtl_context = static_cast<MTLContext *>(unwrap(GPU_context_active_get()));
BLI_assert(mtl_context != nullptr);
return mtl_texture_read_impl(specialisation,
return mtl_texture_read_impl(specialization,
mtl_context->get_texture_utils().texture_1d_read_compute_psos,
GPU_TEXTURE_1D);
}
id<MTLComputePipelineState> gpu::MTLTexture::texture_read_1d_array_get_kernel(
TextureReadRoutineSpecialisation specialisation)
TextureReadRoutineSpecialisation specialization)
{
MTLContext *mtl_context = static_cast<MTLContext *>(unwrap(GPU_context_active_get()));
BLI_assert(mtl_context != nullptr);
return mtl_texture_read_impl(specialisation,
return mtl_texture_read_impl(specialization,
mtl_context->get_texture_utils().texture_1d_array_read_compute_psos,
GPU_TEXTURE_1D_ARRAY);
}
id<MTLComputePipelineState> gpu::MTLTexture::texture_read_3d_get_kernel(
TextureReadRoutineSpecialisation specialisation)
TextureReadRoutineSpecialisation specialization)
{
MTLContext *mtl_context = static_cast<MTLContext *>(unwrap(GPU_context_active_get()));
BLI_assert(mtl_context != nullptr);
return mtl_texture_read_impl(specialisation,
return mtl_texture_read_impl(specialization,
mtl_context->get_texture_utils().texture_3d_read_compute_psos,
GPU_TEXTURE_3D);
}

1
source/blender/gpu/opengl/gl_backend.cc
View File

@@ -497,6 +497,7 @@ void GLBackend::capabilities_init()
glGetIntegerv(GL_NUM_EXTENSIONS, &GCaps.extensions_len);
GCaps.extension_get = gl_extension_get;
GCaps.max_samplers = GCaps.max_textures;
GCaps.mem_stats_support = epoxy_has_gl_extension("GL_NVX_gpu_memory_info") ||
epoxy_has_gl_extension("GL_ATI_meminfo");
GCaps.shader_image_load_store_support = epoxy_has_gl_extension("GL_ARB_shader_image_load_store");

109
source/blender/gpu/shaders/metal/mtl_shader_common.msl Normal file
View File

@@ -0,0 +1,109 @@
/* SPDX-License-Identifier: GPL-2.0-or-later */
/* Common Metal header to be included in all compiled Metal shaders.
* Both native MSL shaders and GLSL shaders. */
using namespace metal;
/* Should match GPUVertFetchMode. */
typedef enum {
GPU_FETCH_FLOAT = 0,
GPU_FETCH_INT,
GPU_FETCH_INT_TO_FLOAT_UNIT,
GPU_FETCH_INT_TO_FLOAT,
} GPUVertFetchMode;
/* Consant to flag base binding index of uniform buffers. */
constant int MTL_uniform_buffer_base_index [[function_constant(0)]];
/* Default Point Size.
* Unused if function constant not set. */
constant float MTL_global_pointsize [[function_constant(1)]];
/* Attribute conversions flags (Up to 16 attributes supported in Blender). */
constant int MTL_AttributeConvert0 [[function_constant(2)]];
constant int MTL_AttributeConvert1 [[function_constant(3)]];
constant int MTL_AttributeConvert2 [[function_constant(4)]];
constant int MTL_AttributeConvert3 [[function_constant(5)]];
constant int MTL_AttributeConvert4 [[function_constant(6)]];
constant int MTL_AttributeConvert5 [[function_constant(7)]];
constant int MTL_AttributeConvert6 [[function_constant(8)]];
constant int MTL_AttributeConvert7 [[function_constant(9)]];
constant int MTL_AttributeConvert8 [[function_constant(10)]];
constant int MTL_AttributeConvert9 [[function_constant(11)]];
constant int MTL_AttributeConvert10 [[function_constant(12)]];
constant int MTL_AttributeConvert11 [[function_constant(13)]];
constant int MTL_AttributeConvert12 [[function_constant(14)]];
constant int MTL_AttributeConvert13 [[function_constant(15)]];
constant int MTL_AttributeConvert14 [[function_constant(16)]];
constant int MTL_AttributeConvert15 [[function_constant(17)]];
/* Consant to flag binding index of transform feedback buffer.
* Unused if function constant not set. */
constant int MTL_transform_feedback_buffer_index [[function_constant(18)]];
/** Internal attribute conversion functionality. */
/* Following descriptions in mtl_shader.hh, Metal only supports some implicit
* attribute type conversions. These conversions occur when there is a difference
* between the type specified in the vertex descriptor (In the input vertex buffers),
* and the attribute type in the shader's VertexIn struct (ShaderInterface).
*
* The supported implicit conversions are described here:
* https://developer.apple.com/documentation/metal/mtlvertexattributedescriptor/1516081-format?language=objc
*
* For unsupported conversions, the mtl_shader_generator will create an attribute reading function
* which performs this conversion manually upon read, depending on the requested fetchmode.
*
* These conversions use the function constants above, so any branching is optimized out during
* backend shader compilation (PSO creation).
*
* NOTE: Not all possibilities have been covered here, any additional conversion routines should
* be added as needed, and mtl_shader_generator should also be updated with any newly required
* read functions.
*
* These paths are only needed for cases where implicit conversion will not happen, in which
* case the value will be read as the type in the shader.
*/
#define internal_vertex_attribute_convert_read_float(ATTR, v_in, v_out) \
if (ATTR == GPU_FETCH_INT_TO_FLOAT) { \
v_out = float(as_type<int>(v_in)); \
} \
else if (ATTR == GPU_FETCH_INT_TO_FLOAT_UNIT) { \
v_out = float(as_type<int>(v_in)) / float(__INT_MAX__); \
} \
else { \
v_out = v_in; \
}
#define internal_vertex_attribute_convert_read_float2(ATTR, v_in, v_out) \
if (ATTR == GPU_FETCH_INT_TO_FLOAT) { \
v_out = float2(as_type<int2>(v_in)); \
} \
else if (ATTR == GPU_FETCH_INT_TO_FLOAT_UNIT) { \
v_out = float2(as_type<int2>(v_in)) / float2(__INT_MAX__); \
} \
else { \
v_out = v_in; \
}
#define internal_vertex_attribute_convert_read_float3(ATTR, v_in, v_out) \
if (ATTR == GPU_FETCH_INT_TO_FLOAT) { \
v_out = float3(as_type<int3>(v_in)); \
} \
else if (ATTR == GPU_FETCH_INT_TO_FLOAT_UNIT) { \
v_out = float3(as_type<int3>(v_in)) / float3(__INT_MAX__); \
} \
else { \
v_out = v_in; \
}
#define internal_vertex_attribute_convert_read_float4(ATTR, v_in, v_out) \
if (ATTR == GPU_FETCH_INT_TO_FLOAT) { \
v_out = float4(as_type<int4>(v_in)); \
} \
else if (ATTR == GPU_FETCH_INT_TO_FLOAT_UNIT) { \
v_out = float4(as_type<int4>(v_in)) / float4(__INT_MAX__); \
} \
else { \
v_out = v_in; \
}

1065
source/blender/gpu/shaders/metal/mtl_shader_defines.msl Normal file
View File

File diff suppressed because it is too large Load Diff

15
source/blender/python/gpu/gpu_py_shader_create_info.cc
View File

@@ -673,6 +673,9 @@ static int constant_type_size(Type type)
case Type::FLOAT:
case Type::INT:
case Type::UINT:
case Type::UCHAR4:
case Type::CHAR4:
case blender::gpu::shader::Type::VEC3_101010I2:
return 4;
break;
case Type::VEC2:
@@ -695,6 +698,18 @@ static int constant_type_size(Type type)
case Type::MAT4:
return 64;
break;
case blender::gpu::shader::Type::UCHAR:
case blender::gpu::shader::Type::CHAR:
return 1;
break;
case blender::gpu::shader::Type::UCHAR2:
case blender::gpu::shader::Type::CHAR2:
return 2;
break;
case blender::gpu::shader::Type::UCHAR3:
case blender::gpu::shader::Type::CHAR3:
return 3;
break;
}
BLI_assert(false);
return -1;