57552f52b2
This patch adds support for compilation and execution of GLSL compute shaders. This, along with a few systematic changes and fixes, enable realtime compositor functionality with the Metal backend on macOS. A number of GLSL source modifications have been made to add the required level of type explicitness, allowing all compilations to succeed. GLSL Compute shader compilation follows a similar path to Vertex/Fragment translation, with added support for shader atomics, shared memory blocks and barriers. Texture flags have also been updated to ensure correct read/write specification for textures used within the compositor pipeline. GPU command submission changes have also been made in the high level path, when Metal is used, to address command buffer time-outs caused by certain expensive compute shaders. Authored by Apple: Michael Parkin-White Ref T96261 Ref T99210 Reviewed By: fclem Maniphest Tasks: T99210, T96261 Differential Revision: https://developer.blender.org/D16990
2160 lines
78 KiB
C++
2160 lines
78 KiB
C++
/* SPDX-License-Identifier: GPL-2.0-or-later */
|
|
|
|
/** \file
|
|
* \ingroup gpu
|
|
*/
|
|
|
|
#include "BKE_global.h"
|
|
|
|
#include "DNA_userdef_types.h"
|
|
|
|
#include "GPU_batch.h"
|
|
#include "GPU_batch_presets.h"
|
|
#include "GPU_capabilities.h"
|
|
#include "GPU_framebuffer.h"
|
|
#include "GPU_immediate.h"
|
|
#include "GPU_platform.h"
|
|
#include "GPU_state.h"
|
|
|
|
#include "mtl_backend.hh"
|
|
#include "mtl_common.hh"
|
|
#include "mtl_context.hh"
|
|
#include "mtl_debug.hh"
|
|
#include "mtl_texture.hh"
|
|
#include "mtl_vertex_buffer.hh"
|
|
|
|
#include "GHOST_C-api.h"
|
|
|
|
namespace blender::gpu {
|
|
|
|
/* -------------------------------------------------------------------- */
|
|
/** \name Creation & Deletion
|
|
* \{ */
|
|
|
|
void gpu::MTLTexture::mtl_texture_init()
|
|
{
|
|
BLI_assert(MTLContext::get() != nullptr);
|
|
|
|
/* Status. */
|
|
is_baked_ = false;
|
|
is_dirty_ = false;
|
|
resource_mode_ = MTL_TEXTURE_MODE_DEFAULT;
|
|
mtl_max_mips_ = 1;
|
|
|
|
/* Metal properties. */
|
|
texture_ = nil;
|
|
texture_buffer_ = nil;
|
|
mip_swizzle_view_ = nil;
|
|
|
|
/* Binding information. */
|
|
is_bound_ = false;
|
|
|
|
/* VBO. */
|
|
vert_buffer_ = nullptr;
|
|
vert_buffer_mtl_ = nil;
|
|
|
|
/* Default Swizzle. */
|
|
tex_swizzle_mask_[0] = 'r';
|
|
tex_swizzle_mask_[1] = 'g';
|
|
tex_swizzle_mask_[2] = 'b';
|
|
tex_swizzle_mask_[3] = 'a';
|
|
mtl_swizzle_mask_ = MTLTextureSwizzleChannelsMake(
|
|
MTLTextureSwizzleRed, MTLTextureSwizzleGreen, MTLTextureSwizzleBlue, MTLTextureSwizzleAlpha);
|
|
}
|
|
|
|
gpu::MTLTexture::MTLTexture(const char *name) : Texture(name)
|
|
{
|
|
/* Common Initialization. */
|
|
mtl_texture_init();
|
|
}
|
|
|
|
gpu::MTLTexture::MTLTexture(const char *name,
|
|
eGPUTextureFormat format,
|
|
eGPUTextureType type,
|
|
id<MTLTexture> metal_texture)
|
|
: Texture(name)
|
|
{
|
|
/* Common Initialization. */
|
|
mtl_texture_init();
|
|
|
|
/* Prep texture from METAL handle. */
|
|
BLI_assert(metal_texture != nil);
|
|
BLI_assert(type == GPU_TEXTURE_2D);
|
|
type_ = type;
|
|
init_2D(metal_texture.width, metal_texture.height, 0, 1, format);
|
|
|
|
/* Assign MTLTexture. */
|
|
texture_ = metal_texture;
|
|
[texture_ retain];
|
|
gpu_image_usage_flags_ = gpu_usage_from_mtl(metal_texture.usage);
|
|
|
|
/* Flag as Baked. */
|
|
is_baked_ = true;
|
|
is_dirty_ = false;
|
|
resource_mode_ = MTL_TEXTURE_MODE_EXTERNAL;
|
|
}
|
|
|
|
gpu::MTLTexture::~MTLTexture()
|
|
{
|
|
/* Unbind if bound. */
|
|
if (is_bound_) {
|
|
MTLContext *ctx = static_cast<MTLContext *>(unwrap(GPU_context_active_get()));
|
|
if (ctx != nullptr) {
|
|
ctx->state_manager->texture_unbind(this);
|
|
}
|
|
}
|
|
|
|
/* Free memory. */
|
|
this->reset();
|
|
}
|
|
|
|
/** \} */
|
|
|
|
/* -------------------------------------------------------------------- */
|
|
void gpu::MTLTexture::bake_mip_swizzle_view()
|
|
{
|
|
if (texture_view_dirty_flags_) {
|
|
|
|
/* Optimization: only generate texture view for mipmapped textures if base level > 0
|
|
* and max level does not match the existing number of mips.
|
|
* Only apply this if mipmap is the only change, and we have not previously generated
|
|
* a texture view. For textures which are created as views, this should also be skipped. */
|
|
if (resource_mode_ != MTL_TEXTURE_MODE_TEXTURE_VIEW &&
|
|
texture_view_dirty_flags_ == TEXTURE_VIEW_MIP_DIRTY && mip_swizzle_view_ == nil) {
|
|
|
|
if (mip_texture_base_level_ == 0 && mip_texture_max_level_ == mtl_max_mips_) {
|
|
texture_view_dirty_flags_ = TEXTURE_VIEW_NOT_DIRTY;
|
|
return;
|
|
}
|
|
}
|
|
|
|
/* Ensure we have texture view usage flagged. */
|
|
BLI_assert(gpu_image_usage_flags_ & GPU_TEXTURE_USAGE_MIP_SWIZZLE_VIEW);
|
|
|
|
/* if a texture view was previously created we release it. */
|
|
if (mip_swizzle_view_ != nil) {
|
|
[mip_swizzle_view_ release];
|
|
mip_swizzle_view_ = nil;
|
|
}
|
|
|
|
/* Determine num slices */
|
|
int num_slices = 1;
|
|
switch (type_) {
|
|
case GPU_TEXTURE_1D_ARRAY:
|
|
num_slices = h_;
|
|
break;
|
|
case GPU_TEXTURE_2D_ARRAY:
|
|
num_slices = d_;
|
|
break;
|
|
case GPU_TEXTURE_CUBE:
|
|
num_slices = 6;
|
|
break;
|
|
case GPU_TEXTURE_CUBE_ARRAY:
|
|
/* d_ is equal to array levels * 6, including face count. */
|
|
num_slices = d_;
|
|
break;
|
|
default:
|
|
num_slices = 1;
|
|
break;
|
|
}
|
|
|
|
int range_len = min_ii((mip_texture_max_level_ - mip_texture_base_level_) + 1,
|
|
texture_.mipmapLevelCount);
|
|
BLI_assert(range_len > 0);
|
|
BLI_assert(mip_texture_base_level_ < texture_.mipmapLevelCount);
|
|
BLI_assert(mip_texture_base_layer_ < num_slices);
|
|
mip_swizzle_view_ = [texture_
|
|
newTextureViewWithPixelFormat:texture_.pixelFormat
|
|
textureType:texture_.textureType
|
|
levels:NSMakeRange(mip_texture_base_level_, range_len)
|
|
slices:NSMakeRange(mip_texture_base_layer_, num_slices)
|
|
swizzle:mtl_swizzle_mask_];
|
|
MTL_LOG_INFO(
|
|
"Updating texture view - MIP TEXTURE BASE LEVEL: %d, MAX LEVEL: %d (Range len: %d)\n",
|
|
mip_texture_base_level_,
|
|
min_ii(mip_texture_max_level_, texture_.mipmapLevelCount),
|
|
range_len);
|
|
mip_swizzle_view_.label = [texture_ label];
|
|
texture_view_dirty_flags_ = TEXTURE_VIEW_NOT_DIRTY;
|
|
}
|
|
}
|
|
|
|
/** \name Operations
|
|
* \{ */
|
|
|
|
id<MTLTexture> gpu::MTLTexture::get_metal_handle()
|
|
{
|
|
|
|
/* Verify VBO texture shares same buffer. */
|
|
if (resource_mode_ == MTL_TEXTURE_MODE_VBO) {
|
|
id<MTLBuffer> buf = vert_buffer_->get_metal_buffer();
|
|
|
|
/* Source vertex buffer has been re-generated, require re-initialization. */
|
|
if (buf != vert_buffer_mtl_) {
|
|
MTL_LOG_INFO(
|
|
"MTLTexture '%p' using MTL_TEXTURE_MODE_VBO requires re-generation due to updated "
|
|
"Vertex-Buffer.\n",
|
|
this);
|
|
/* Clear state. */
|
|
this->reset();
|
|
|
|
/* Re-initialize. */
|
|
this->init_internal(wrap(vert_buffer_));
|
|
|
|
/* Update for assertion check below. */
|
|
buf = vert_buffer_->get_metal_buffer();
|
|
}
|
|
|
|
/* Ensure buffer is valid.
|
|
* Fetch-vert buffer handle directly in-case it changed above. */
|
|
BLI_assert(vert_buffer_mtl_ != nil);
|
|
BLI_assert(vert_buffer_->get_metal_buffer() == vert_buffer_mtl_);
|
|
}
|
|
|
|
/* ensure up to date and baked. */
|
|
this->ensure_baked();
|
|
|
|
if (is_baked_) {
|
|
/* For explicit texture views, ensure we always return the texture view. */
|
|
if (resource_mode_ == MTL_TEXTURE_MODE_TEXTURE_VIEW) {
|
|
BLI_assert_msg(mip_swizzle_view_, "Texture view should always have a valid handle.");
|
|
}
|
|
|
|
if (mip_swizzle_view_ != nil || texture_view_dirty_flags_) {
|
|
bake_mip_swizzle_view();
|
|
|
|
/* Optimization: If texture view does not change mip parameters, no texture view will be
|
|
* baked. This is because texture views remove the ability to perform lossless compression.
|
|
*/
|
|
if (mip_swizzle_view_ != nil) {
|
|
return mip_swizzle_view_;
|
|
}
|
|
}
|
|
return texture_;
|
|
}
|
|
return nil;
|
|
}
|
|
|
|
id<MTLTexture> gpu::MTLTexture::get_metal_handle_base()
|
|
{
|
|
|
|
/* ensure up to date and baked. */
|
|
this->ensure_baked();
|
|
|
|
/* For explicit texture views, always return the texture view. */
|
|
if (resource_mode_ == MTL_TEXTURE_MODE_TEXTURE_VIEW) {
|
|
BLI_assert_msg(mip_swizzle_view_, "Texture view should always have a valid handle.");
|
|
if (mip_swizzle_view_ != nil || texture_view_dirty_flags_) {
|
|
bake_mip_swizzle_view();
|
|
}
|
|
BLI_assert(mip_swizzle_view_ != nil);
|
|
return mip_swizzle_view_;
|
|
}
|
|
|
|
/* Return base handle. */
|
|
if (is_baked_) {
|
|
return texture_;
|
|
}
|
|
return nil;
|
|
}
|
|
|
|
void gpu::MTLTexture::blit(id<MTLBlitCommandEncoder> blit_encoder,
|
|
uint src_x_offset,
|
|
uint src_y_offset,
|
|
uint src_z_offset,
|
|
uint src_slice,
|
|
uint src_mip,
|
|
gpu::MTLTexture *dest,
|
|
uint dst_x_offset,
|
|
uint dst_y_offset,
|
|
uint dst_z_offset,
|
|
uint dst_slice,
|
|
uint dst_mip,
|
|
uint width,
|
|
uint height,
|
|
uint depth)
|
|
{
|
|
|
|
BLI_assert(this && dest);
|
|
BLI_assert(width > 0 && height > 0 && depth > 0);
|
|
MTLSize src_size = MTLSizeMake(width, height, depth);
|
|
MTLOrigin src_origin = MTLOriginMake(src_x_offset, src_y_offset, src_z_offset);
|
|
MTLOrigin dst_origin = MTLOriginMake(dst_x_offset, dst_y_offset, dst_z_offset);
|
|
|
|
if (this->format_get() != dest->format_get()) {
|
|
MTL_LOG_WARNING(
|
|
"[Warning] gpu::MTLTexture: Cannot copy between two textures of different types using a "
|
|
"blit encoder. TODO: Support this operation\n");
|
|
return;
|
|
}
|
|
|
|
/* TODO(Metal): Verify if we want to use the one with modified base-level/texture view
|
|
* or not. */
|
|
[blit_encoder copyFromTexture:this->get_metal_handle_base()
|
|
sourceSlice:src_slice
|
|
sourceLevel:src_mip
|
|
sourceOrigin:src_origin
|
|
sourceSize:src_size
|
|
toTexture:dest->get_metal_handle_base()
|
|
destinationSlice:dst_slice
|
|
destinationLevel:dst_mip
|
|
destinationOrigin:dst_origin];
|
|
}
|
|
|
|
void gpu::MTLTexture::blit(gpu::MTLTexture *dst,
|
|
uint src_x_offset,
|
|
uint src_y_offset,
|
|
uint dst_x_offset,
|
|
uint dst_y_offset,
|
|
uint src_mip,
|
|
uint dst_mip,
|
|
uint dst_slice,
|
|
int width,
|
|
int height)
|
|
{
|
|
BLI_assert(this->type_get() == dst->type_get());
|
|
|
|
GPUShader *shader = fullscreen_blit_sh_get();
|
|
BLI_assert(shader != nullptr);
|
|
BLI_assert(GPU_context_active_get());
|
|
|
|
/* Fetch restore framebuffer and blit target framebuffer from destination texture. */
|
|
GPUFrameBuffer *restore_fb = GPU_framebuffer_active_get();
|
|
GPUFrameBuffer *blit_target_fb = dst->get_blit_framebuffer(dst_slice, dst_mip);
|
|
BLI_assert(blit_target_fb);
|
|
GPU_framebuffer_bind(blit_target_fb);
|
|
|
|
/* Execute graphics draw call to perform the blit. */
|
|
GPUBatch *quad = GPU_batch_preset_quad();
|
|
GPU_batch_set_shader(quad, shader);
|
|
|
|
float w = dst->width_get();
|
|
float h = dst->height_get();
|
|
|
|
GPU_shader_uniform_2f(shader, "fullscreen", w, h);
|
|
GPU_shader_uniform_2f(shader, "src_offset", src_x_offset, src_y_offset);
|
|
GPU_shader_uniform_2f(shader, "dst_offset", dst_x_offset, dst_y_offset);
|
|
GPU_shader_uniform_2f(shader, "size", width, height);
|
|
|
|
GPU_shader_uniform_1i(shader, "mip", src_mip);
|
|
GPU_batch_texture_bind(quad, "imageTexture", wrap(this));
|
|
|
|
/* Caching previous pipeline state. */
|
|
bool depth_write_prev = GPU_depth_mask_get();
|
|
uint stencil_mask_prev = GPU_stencil_mask_get();
|
|
eGPUStencilTest stencil_test_prev = GPU_stencil_test_get();
|
|
eGPUFaceCullTest culling_test_prev = GPU_face_culling_get();
|
|
eGPUBlend blend_prev = GPU_blend_get();
|
|
eGPUDepthTest depth_test_prev = GPU_depth_test_get();
|
|
GPU_scissor_test(false);
|
|
|
|
/* Apply state for blit draw call. */
|
|
GPU_stencil_write_mask_set(0xFF);
|
|
GPU_stencil_reference_set(0);
|
|
GPU_face_culling(GPU_CULL_NONE);
|
|
GPU_stencil_test(GPU_STENCIL_ALWAYS);
|
|
GPU_depth_mask(false);
|
|
GPU_blend(GPU_BLEND_NONE);
|
|
GPU_depth_test(GPU_DEPTH_ALWAYS);
|
|
|
|
GPU_batch_draw(quad);
|
|
|
|
/* restoring old pipeline state. */
|
|
GPU_depth_mask(depth_write_prev);
|
|
GPU_stencil_write_mask_set(stencil_mask_prev);
|
|
GPU_stencil_test(stencil_test_prev);
|
|
GPU_face_culling(culling_test_prev);
|
|
GPU_depth_mask(depth_write_prev);
|
|
GPU_blend(blend_prev);
|
|
GPU_depth_test(depth_test_prev);
|
|
|
|
if (restore_fb != nullptr) {
|
|
GPU_framebuffer_bind(restore_fb);
|
|
}
|
|
else {
|
|
GPU_framebuffer_restore();
|
|
}
|
|
}
|
|
|
|
GPUFrameBuffer *gpu::MTLTexture::get_blit_framebuffer(uint dst_slice, uint dst_mip)
|
|
{
|
|
|
|
/* Check if layer has changed. */
|
|
bool update_attachments = false;
|
|
if (!blit_fb_) {
|
|
blit_fb_ = GPU_framebuffer_create("gpu_blit");
|
|
update_attachments = true;
|
|
}
|
|
|
|
/* Check if current blit FB has the correct attachment properties. */
|
|
if (blit_fb_) {
|
|
if (blit_fb_slice_ != dst_slice || blit_fb_mip_ != dst_mip) {
|
|
update_attachments = true;
|
|
}
|
|
}
|
|
|
|
if (update_attachments) {
|
|
if (format_flag_ & GPU_FORMAT_DEPTH || format_flag_ & GPU_FORMAT_STENCIL) {
|
|
/* DEPTH TEX */
|
|
GPU_framebuffer_ensure_config(
|
|
&blit_fb_,
|
|
{GPU_ATTACHMENT_TEXTURE_LAYER_MIP(wrap(static_cast<Texture *>(this)),
|
|
static_cast<int>(dst_slice),
|
|
static_cast<int>(dst_mip)),
|
|
GPU_ATTACHMENT_NONE});
|
|
}
|
|
else {
|
|
/* COLOR TEX */
|
|
GPU_framebuffer_ensure_config(
|
|
&blit_fb_,
|
|
{GPU_ATTACHMENT_NONE,
|
|
GPU_ATTACHMENT_TEXTURE_LAYER_MIP(wrap(static_cast<Texture *>(this)),
|
|
static_cast<int>(dst_slice),
|
|
static_cast<int>(dst_mip))});
|
|
}
|
|
blit_fb_slice_ = dst_slice;
|
|
blit_fb_mip_ = dst_mip;
|
|
}
|
|
|
|
BLI_assert(blit_fb_);
|
|
return blit_fb_;
|
|
}
|
|
|
|
MTLSamplerState gpu::MTLTexture::get_sampler_state()
|
|
{
|
|
MTLSamplerState sampler_state;
|
|
sampler_state.state = this->sampler_state;
|
|
/* Add more parameters as needed */
|
|
return sampler_state;
|
|
}
|
|
|
|
void gpu::MTLTexture::update_sub(
|
|
int mip, int offset[3], int extent[3], eGPUDataFormat type, const void *data)
|
|
{
|
|
/* Fetch active context. */
|
|
MTLContext *ctx = static_cast<MTLContext *>(unwrap(GPU_context_active_get()));
|
|
BLI_assert(ctx);
|
|
|
|
/* Do not update texture view. */
|
|
BLI_assert(resource_mode_ != MTL_TEXTURE_MODE_TEXTURE_VIEW);
|
|
|
|
/* Ensure mipmaps. */
|
|
this->ensure_mipmaps(mip);
|
|
|
|
/* Ensure texture is baked. */
|
|
this->ensure_baked();
|
|
|
|
/* Safety checks. */
|
|
#if TRUST_NO_ONE
|
|
BLI_assert(mip >= mip_min_ && mip <= mip_max_);
|
|
BLI_assert(mip < texture_.mipmapLevelCount);
|
|
BLI_assert(texture_.mipmapLevelCount >= mip_max_);
|
|
#endif
|
|
|
|
/* DEPTH FLAG - Depth formats cannot use direct BLIT - pass off to their own routine which will
|
|
* do a depth-only render. */
|
|
bool is_depth_format = (format_flag_ & GPU_FORMAT_DEPTH);
|
|
if (is_depth_format) {
|
|
switch (type_) {
|
|
|
|
case GPU_TEXTURE_2D:
|
|
update_sub_depth_2d(mip, offset, extent, type, data);
|
|
return;
|
|
default:
|
|
MTL_LOG_ERROR(
|
|
"[Error] gpu::MTLTexture::update_sub not yet supported for other depth "
|
|
"configurations\n");
|
|
return;
|
|
}
|
|
}
|
|
|
|
@autoreleasepool {
|
|
/* Determine totalsize of INPUT Data. */
|
|
int num_channels = to_component_len(format_);
|
|
int input_bytes_per_pixel = to_bytesize(format_, type);
|
|
int totalsize = 0;
|
|
|
|
/* If unpack row length is used, size of input data uses the unpack row length, rather than the
|
|
* image length. */
|
|
int expected_update_w = ((ctx->pipeline_state.unpack_row_length == 0) ?
|
|
extent[0] :
|
|
ctx->pipeline_state.unpack_row_length);
|
|
|
|
/* Ensure calculated total size isn't larger than remaining image data size */
|
|
switch (this->dimensions_count()) {
|
|
case 1:
|
|
totalsize = input_bytes_per_pixel * max_ii(expected_update_w, 1);
|
|
break;
|
|
case 2:
|
|
totalsize = input_bytes_per_pixel * max_ii(expected_update_w, 1) * extent[1];
|
|
break;
|
|
case 3:
|
|
totalsize = input_bytes_per_pixel * max_ii(expected_update_w, 1) * extent[1] * extent[2];
|
|
break;
|
|
default:
|
|
BLI_assert(false);
|
|
break;
|
|
}
|
|
|
|
/* Early exit if update size is zero. update_sub sometimes has a zero-sized
|
|
* extent when called from texture painting. */
|
|
if (totalsize <= 0 || extent[0] <= 0) {
|
|
MTL_LOG_WARNING(
|
|
"MTLTexture::update_sub called with extent size of zero for one or more dimensions. "
|
|
"(%d, %d, %d) - DimCount: %u\n",
|
|
extent[0],
|
|
extent[1],
|
|
extent[2],
|
|
this->dimensions_count());
|
|
return;
|
|
}
|
|
|
|
/* When unpack row length is used, provided data does not necessarily contain padding for last
|
|
* row, so we only include up to the end of updated data. */
|
|
if (ctx->pipeline_state.unpack_row_length > 0) {
|
|
BLI_assert(ctx->pipeline_state.unpack_row_length >= extent[0]);
|
|
totalsize -= (ctx->pipeline_state.unpack_row_length - extent[0]) * input_bytes_per_pixel;
|
|
}
|
|
|
|
/* Check */
|
|
BLI_assert(totalsize > 0);
|
|
|
|
/* Determine expected destination data size. */
|
|
MTLPixelFormat destination_format = gpu_texture_format_to_metal(format_);
|
|
int expected_dst_bytes_per_pixel = get_mtl_format_bytesize(destination_format);
|
|
int destination_num_channels = get_mtl_format_num_components(destination_format);
|
|
|
|
/* 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,
|
|
destination_num_channels};
|
|
|
|
/* Determine whether we can do direct BLIT or not. */
|
|
bool can_use_direct_blit = true;
|
|
if (expected_dst_bytes_per_pixel != input_bytes_per_pixel ||
|
|
num_channels != destination_num_channels) {
|
|
can_use_direct_blit = false;
|
|
}
|
|
|
|
if (is_depth_format) {
|
|
if (type_ == GPU_TEXTURE_2D || type_ == GPU_TEXTURE_2D_ARRAY) {
|
|
/* Workaround for crash in validation layer when blitting to depth2D target with
|
|
* dimensions (1, 1, 1); */
|
|
if (extent[0] == 1 && extent[1] == 1 && extent[2] == 1 && totalsize == 4) {
|
|
can_use_direct_blit = false;
|
|
}
|
|
}
|
|
}
|
|
|
|
if (format_ == GPU_SRGB8_A8 && !can_use_direct_blit) {
|
|
MTL_LOG_WARNING(
|
|
"SRGB data upload does not work correctly using compute upload. "
|
|
"texname '%s'\n",
|
|
name_);
|
|
}
|
|
|
|
/* Safety Checks. */
|
|
if (type == GPU_DATA_UINT_24_8 || type == GPU_DATA_10_11_11_REV) {
|
|
BLI_assert(can_use_direct_blit &&
|
|
"Special input data type must be a 1-1 mapping with destination texture as it "
|
|
"cannot easily be split");
|
|
}
|
|
|
|
/* Debug and verification. */
|
|
if (!can_use_direct_blit) {
|
|
MTL_LOG_WARNING(
|
|
"gpu::MTLTexture::update_sub supplied bpp is %d bytes (%d components per "
|
|
"pixel), but backing texture bpp is %d bytes (%d components per pixel) "
|
|
"(TODO(Metal): Channel Conversion needed) (w: %d, h: %d, d: %d)\n",
|
|
input_bytes_per_pixel,
|
|
num_channels,
|
|
expected_dst_bytes_per_pixel,
|
|
destination_num_channels,
|
|
w_,
|
|
h_,
|
|
d_);
|
|
|
|
/* Check mip compatibility. */
|
|
if (mip != 0) {
|
|
MTL_LOG_ERROR(
|
|
"[Error]: Updating texture layers other than mip=0 when data is mismatched is not "
|
|
"possible in METAL on macOS using texture->write\n");
|
|
return;
|
|
}
|
|
|
|
/* Check Format write-ability. */
|
|
if (mtl_format_get_writeable_view_format(destination_format) == MTLPixelFormatInvalid) {
|
|
MTL_LOG_ERROR(
|
|
"[Error]: Updating texture -- destination MTLPixelFormat '%d' does not support write "
|
|
"operations, and no suitable TextureView format exists.\n",
|
|
*(int *)(&destination_format));
|
|
return;
|
|
}
|
|
}
|
|
|
|
/* Prepare staging buffer for data. */
|
|
id<MTLBuffer> staging_buffer = nil;
|
|
uint64_t staging_buffer_offset = 0;
|
|
|
|
/* Fetch allocation from scratch buffer. */
|
|
MTLTemporaryBuffer allocation =
|
|
ctx->get_scratchbuffer_manager().scratch_buffer_allocate_range_aligned(totalsize, 256);
|
|
memcpy(allocation.data, data, totalsize);
|
|
staging_buffer = allocation.metal_buffer;
|
|
staging_buffer_offset = allocation.buffer_offset;
|
|
|
|
/* Common Properties. */
|
|
MTLPixelFormat compatible_write_format = mtl_format_get_writeable_view_format(
|
|
destination_format);
|
|
|
|
/* Some texture formats are not writeable so we need to use a texture view. */
|
|
if (compatible_write_format == MTLPixelFormatInvalid) {
|
|
MTL_LOG_ERROR("Cannot use compute update blit with texture-view format: %d\n",
|
|
*((int *)&compatible_write_format));
|
|
return;
|
|
}
|
|
|
|
/* Prepare command encoders. */
|
|
id<MTLBlitCommandEncoder> blit_encoder = nil;
|
|
id<MTLComputeCommandEncoder> compute_encoder = nil;
|
|
id<MTLTexture> staging_texture = nil;
|
|
id<MTLTexture> texture_handle = nil;
|
|
|
|
/* Use staging texture. */
|
|
bool use_staging_texture = false;
|
|
|
|
if (can_use_direct_blit) {
|
|
blit_encoder = ctx->main_command_buffer.ensure_begin_blit_encoder();
|
|
BLI_assert(blit_encoder != nil);
|
|
|
|
/* If we need to use a texture view to write texture data as the source
|
|
* format is unwritable, if our texture has not been initialized with
|
|
* texture view support, use a staging texture. */
|
|
if ((compatible_write_format != destination_format) &&
|
|
!(gpu_image_usage_flags_ & GPU_TEXTURE_USAGE_MIP_SWIZZLE_VIEW)) {
|
|
use_staging_texture = true;
|
|
}
|
|
}
|
|
else {
|
|
compute_encoder = ctx->main_command_buffer.ensure_begin_compute_encoder();
|
|
BLI_assert(compute_encoder != nil);
|
|
|
|
/* For compute, we should use a stating texture to avoid texture write usage,
|
|
* if it has not been specified for the texture. Using shader-write disables
|
|
* lossless texture compression, so this is best to avoid where possible. */
|
|
if (!(gpu_image_usage_flags_ & GPU_TEXTURE_USAGE_SHADER_WRITE)) {
|
|
use_staging_texture = true;
|
|
}
|
|
if (compatible_write_format != destination_format) {
|
|
if (!(gpu_image_usage_flags_ & GPU_TEXTURE_USAGE_MIP_SWIZZLE_VIEW)) {
|
|
use_staging_texture = true;
|
|
}
|
|
}
|
|
}
|
|
|
|
/* Allocate stating texture if needed. */
|
|
if (use_staging_texture) {
|
|
/* Create staging texture to avoid shader-write limiting optimization. */
|
|
BLI_assert(texture_descriptor_ != nullptr);
|
|
MTLTextureUsage original_usage = texture_descriptor_.usage;
|
|
texture_descriptor_.usage = original_usage | MTLTextureUsageShaderWrite |
|
|
MTLTextureUsagePixelFormatView;
|
|
staging_texture = [ctx->device newTextureWithDescriptor:texture_descriptor_];
|
|
staging_texture.label = @"Staging texture";
|
|
texture_descriptor_.usage = original_usage;
|
|
|
|
/* Create texture view if needed. */
|
|
texture_handle = ((compatible_write_format == destination_format)) ?
|
|
[staging_texture retain] :
|
|
[staging_texture newTextureViewWithPixelFormat:compatible_write_format];
|
|
}
|
|
else {
|
|
/* Use texture view. */
|
|
if (compatible_write_format != destination_format) {
|
|
BLI_assert(gpu_image_usage_flags_ & GPU_TEXTURE_USAGE_MIP_SWIZZLE_VIEW);
|
|
texture_handle = [texture_ newTextureViewWithPixelFormat:compatible_write_format];
|
|
}
|
|
else {
|
|
texture_handle = texture_;
|
|
[texture_handle retain];
|
|
}
|
|
}
|
|
|
|
switch (type_) {
|
|
|
|
/* 1D */
|
|
case GPU_TEXTURE_1D:
|
|
case GPU_TEXTURE_1D_ARRAY: {
|
|
if (can_use_direct_blit) {
|
|
/* Use Blit based update. */
|
|
int bytes_per_row = expected_dst_bytes_per_pixel *
|
|
((ctx->pipeline_state.unpack_row_length == 0) ?
|
|
extent[0] :
|
|
ctx->pipeline_state.unpack_row_length);
|
|
int bytes_per_image = bytes_per_row;
|
|
int max_array_index = ((type_ == GPU_TEXTURE_1D_ARRAY) ? extent[1] : 1);
|
|
for (int array_index = 0; array_index < max_array_index; array_index++) {
|
|
|
|
int buffer_array_offset = staging_buffer_offset + (bytes_per_image * array_index);
|
|
[blit_encoder
|
|
copyFromBuffer:staging_buffer
|
|
sourceOffset:buffer_array_offset
|
|
sourceBytesPerRow:bytes_per_row
|
|
sourceBytesPerImage:bytes_per_image
|
|
sourceSize:MTLSizeMake(extent[0], 1, 1)
|
|
toTexture:texture_handle
|
|
destinationSlice:((type_ == GPU_TEXTURE_1D_ARRAY) ? (array_index + offset[1]) :
|
|
0)
|
|
destinationLevel:mip
|
|
destinationOrigin:MTLOriginMake(offset[0], 0, 0)];
|
|
}
|
|
}
|
|
else {
|
|
/* Use Compute Based update. */
|
|
if (type_ == GPU_TEXTURE_1D) {
|
|
id<MTLComputePipelineState> pso = texture_update_1d_get_kernel(
|
|
compute_specialization_kernel);
|
|
TextureUpdateParams params = {mip,
|
|
{extent[0], 1, 1},
|
|
{offset[0], 0, 0},
|
|
((ctx->pipeline_state.unpack_row_length == 0) ?
|
|
extent[0] :
|
|
ctx->pipeline_state.unpack_row_length)};
|
|
|
|
/* Bind resources via compute state for optimal state caching performance. */
|
|
MTLComputeState &cs = ctx->main_command_buffer.get_compute_state();
|
|
cs.bind_pso(pso);
|
|
cs.bind_compute_bytes(¶ms, sizeof(params), 0);
|
|
cs.bind_compute_buffer(staging_buffer, staging_buffer_offset, 1);
|
|
cs.bind_compute_texture(texture_handle, 0);
|
|
[compute_encoder
|
|
dispatchThreads:MTLSizeMake(extent[0], 1, 1) /* Width, Height, Layer */
|
|
threadsPerThreadgroup:MTLSizeMake(64, 1, 1)];
|
|
}
|
|
else if (type_ == GPU_TEXTURE_1D_ARRAY) {
|
|
id<MTLComputePipelineState> pso = texture_update_1d_array_get_kernel(
|
|
compute_specialization_kernel);
|
|
TextureUpdateParams params = {mip,
|
|
{extent[0], extent[1], 1},
|
|
{offset[0], offset[1], 0},
|
|
((ctx->pipeline_state.unpack_row_length == 0) ?
|
|
extent[0] :
|
|
ctx->pipeline_state.unpack_row_length)};
|
|
|
|
/* Bind resources via compute state for optimal state caching performance. */
|
|
MTLComputeState &cs = ctx->main_command_buffer.get_compute_state();
|
|
cs.bind_pso(pso);
|
|
cs.bind_compute_bytes(¶ms, sizeof(params), 0);
|
|
cs.bind_compute_buffer(staging_buffer, staging_buffer_offset, 1);
|
|
cs.bind_compute_texture(texture_handle, 0);
|
|
[compute_encoder
|
|
dispatchThreads:MTLSizeMake(extent[0], extent[1], 1) /* Width, layers, nil */
|
|
threadsPerThreadgroup:MTLSizeMake(8, 8, 1)];
|
|
}
|
|
}
|
|
} break;
|
|
|
|
/* 2D */
|
|
case GPU_TEXTURE_2D:
|
|
case GPU_TEXTURE_2D_ARRAY: {
|
|
if (can_use_direct_blit) {
|
|
/* Use Blit encoder update. */
|
|
int bytes_per_row = expected_dst_bytes_per_pixel *
|
|
((ctx->pipeline_state.unpack_row_length == 0) ?
|
|
extent[0] :
|
|
ctx->pipeline_state.unpack_row_length);
|
|
int bytes_per_image = bytes_per_row * extent[1];
|
|
|
|
int texture_array_relative_offset = 0;
|
|
int base_slice = (type_ == GPU_TEXTURE_2D_ARRAY) ? offset[2] : 0;
|
|
int final_slice = base_slice + ((type_ == GPU_TEXTURE_2D_ARRAY) ? extent[2] : 1);
|
|
|
|
for (int array_slice = base_slice; array_slice < final_slice; array_slice++) {
|
|
|
|
if (array_slice > 0) {
|
|
BLI_assert(type_ == GPU_TEXTURE_2D_ARRAY);
|
|
BLI_assert(array_slice < d_);
|
|
}
|
|
|
|
[blit_encoder copyFromBuffer:staging_buffer
|
|
sourceOffset:staging_buffer_offset + texture_array_relative_offset
|
|
sourceBytesPerRow:bytes_per_row
|
|
sourceBytesPerImage:bytes_per_image
|
|
sourceSize:MTLSizeMake(extent[0], extent[1], 1)
|
|
toTexture:texture_handle
|
|
destinationSlice:array_slice
|
|
destinationLevel:mip
|
|
destinationOrigin:MTLOriginMake(offset[0], offset[1], 0)];
|
|
|
|
texture_array_relative_offset += bytes_per_image;
|
|
}
|
|
}
|
|
else {
|
|
/* Use Compute texture update. */
|
|
if (type_ == GPU_TEXTURE_2D) {
|
|
id<MTLComputePipelineState> pso = texture_update_2d_get_kernel(
|
|
compute_specialization_kernel);
|
|
TextureUpdateParams params = {mip,
|
|
{extent[0], extent[1], 1},
|
|
{offset[0], offset[1], 0},
|
|
((ctx->pipeline_state.unpack_row_length == 0) ?
|
|
extent[0] :
|
|
ctx->pipeline_state.unpack_row_length)};
|
|
|
|
/* Bind resources via compute state for optimal state caching performance. */
|
|
MTLComputeState &cs = ctx->main_command_buffer.get_compute_state();
|
|
cs.bind_pso(pso);
|
|
cs.bind_compute_bytes(¶ms, sizeof(params), 0);
|
|
cs.bind_compute_buffer(staging_buffer, staging_buffer_offset, 1);
|
|
cs.bind_compute_texture(texture_handle, 0);
|
|
[compute_encoder
|
|
dispatchThreads:MTLSizeMake(
|
|
extent[0], extent[1], 1) /* Width, Height, Layer */
|
|
threadsPerThreadgroup:MTLSizeMake(8, 8, 1)];
|
|
}
|
|
else if (type_ == GPU_TEXTURE_2D_ARRAY) {
|
|
id<MTLComputePipelineState> pso = texture_update_2d_array_get_kernel(
|
|
compute_specialization_kernel);
|
|
TextureUpdateParams params = {mip,
|
|
{extent[0], extent[1], extent[2]},
|
|
{offset[0], offset[1], offset[2]},
|
|
((ctx->pipeline_state.unpack_row_length == 0) ?
|
|
extent[0] :
|
|
ctx->pipeline_state.unpack_row_length)};
|
|
|
|
/* Bind resources via compute state for optimal state caching performance. */
|
|
MTLComputeState &cs = ctx->main_command_buffer.get_compute_state();
|
|
cs.bind_pso(pso);
|
|
cs.bind_compute_bytes(¶ms, sizeof(params), 0);
|
|
cs.bind_compute_buffer(staging_buffer, staging_buffer_offset, 1);
|
|
cs.bind_compute_texture(texture_handle, 0);
|
|
[compute_encoder dispatchThreads:MTLSizeMake(extent[0],
|
|
extent[1],
|
|
extent[2]) /* Width, Height, Layer */
|
|
threadsPerThreadgroup:MTLSizeMake(4, 4, 4)];
|
|
}
|
|
}
|
|
|
|
} break;
|
|
|
|
/* 3D */
|
|
case GPU_TEXTURE_3D: {
|
|
if (can_use_direct_blit) {
|
|
int bytes_per_row = expected_dst_bytes_per_pixel *
|
|
((ctx->pipeline_state.unpack_row_length == 0) ?
|
|
extent[0] :
|
|
ctx->pipeline_state.unpack_row_length);
|
|
int bytes_per_image = bytes_per_row * extent[1];
|
|
[blit_encoder copyFromBuffer:staging_buffer
|
|
sourceOffset:staging_buffer_offset
|
|
sourceBytesPerRow:bytes_per_row
|
|
sourceBytesPerImage:bytes_per_image
|
|
sourceSize:MTLSizeMake(extent[0], extent[1], extent[2])
|
|
toTexture:texture_handle
|
|
destinationSlice:0
|
|
destinationLevel:mip
|
|
destinationOrigin:MTLOriginMake(offset[0], offset[1], offset[2])];
|
|
}
|
|
else {
|
|
id<MTLComputePipelineState> pso = texture_update_3d_get_kernel(
|
|
compute_specialization_kernel);
|
|
TextureUpdateParams params = {mip,
|
|
{extent[0], extent[1], extent[2]},
|
|
{offset[0], offset[1], offset[2]},
|
|
((ctx->pipeline_state.unpack_row_length == 0) ?
|
|
extent[0] :
|
|
ctx->pipeline_state.unpack_row_length)};
|
|
|
|
/* Bind resources via compute state for optimal state caching performance. */
|
|
MTLComputeState &cs = ctx->main_command_buffer.get_compute_state();
|
|
cs.bind_pso(pso);
|
|
cs.bind_compute_bytes(¶ms, sizeof(params), 0);
|
|
cs.bind_compute_buffer(staging_buffer, staging_buffer_offset, 1);
|
|
cs.bind_compute_texture(texture_handle, 0);
|
|
[compute_encoder
|
|
dispatchThreads:MTLSizeMake(
|
|
extent[0], extent[1], extent[2]) /* Width, Height, Depth */
|
|
threadsPerThreadgroup:MTLSizeMake(4, 4, 4)];
|
|
}
|
|
} break;
|
|
|
|
/* CUBE */
|
|
case GPU_TEXTURE_CUBE: {
|
|
if (can_use_direct_blit) {
|
|
int bytes_per_row = expected_dst_bytes_per_pixel *
|
|
((ctx->pipeline_state.unpack_row_length == 0) ?
|
|
extent[0] :
|
|
ctx->pipeline_state.unpack_row_length);
|
|
int bytes_per_image = bytes_per_row * extent[1];
|
|
|
|
int texture_array_relative_offset = 0;
|
|
|
|
/* Iterate over all cube faces in range (offset[2], offset[2] + extent[2]). */
|
|
for (int i = 0; i < extent[2]; i++) {
|
|
int face_index = offset[2] + i;
|
|
|
|
[blit_encoder copyFromBuffer:staging_buffer
|
|
sourceOffset:staging_buffer_offset + texture_array_relative_offset
|
|
sourceBytesPerRow:bytes_per_row
|
|
sourceBytesPerImage:bytes_per_image
|
|
sourceSize:MTLSizeMake(extent[0], extent[1], 1)
|
|
toTexture:texture_handle
|
|
destinationSlice:face_index /* = cubeFace+arrayIndex*6 */
|
|
destinationLevel:mip
|
|
destinationOrigin:MTLOriginMake(offset[0], offset[1], 0)];
|
|
texture_array_relative_offset += bytes_per_image;
|
|
}
|
|
}
|
|
else {
|
|
MTL_LOG_ERROR(
|
|
"TODO(Metal): Support compute texture update for GPU_TEXTURE_CUBE %d, %d, %d\n",
|
|
w_,
|
|
h_,
|
|
d_);
|
|
}
|
|
} break;
|
|
|
|
case GPU_TEXTURE_CUBE_ARRAY: {
|
|
if (can_use_direct_blit) {
|
|
|
|
int bytes_per_row = expected_dst_bytes_per_pixel *
|
|
((ctx->pipeline_state.unpack_row_length == 0) ?
|
|
extent[0] :
|
|
ctx->pipeline_state.unpack_row_length);
|
|
int bytes_per_image = bytes_per_row * extent[1];
|
|
|
|
/* Upload to all faces between offset[2] (which is zero in most cases) AND extent[2]. */
|
|
int texture_array_relative_offset = 0;
|
|
for (int i = 0; i < extent[2]; i++) {
|
|
int face_index = offset[2] + i;
|
|
[blit_encoder copyFromBuffer:staging_buffer
|
|
sourceOffset:staging_buffer_offset + texture_array_relative_offset
|
|
sourceBytesPerRow:bytes_per_row
|
|
sourceBytesPerImage:bytes_per_image
|
|
sourceSize:MTLSizeMake(extent[0], extent[1], 1)
|
|
toTexture:texture_handle
|
|
destinationSlice:face_index /* = cubeFace+arrayIndex*6. */
|
|
destinationLevel:mip
|
|
destinationOrigin:MTLOriginMake(offset[0], offset[1], 0)];
|
|
texture_array_relative_offset += bytes_per_image;
|
|
}
|
|
}
|
|
else {
|
|
MTL_LOG_ERROR(
|
|
"TODO(Metal): Support compute texture update for GPU_TEXTURE_CUBE_ARRAY %d, %d, "
|
|
"%d\n",
|
|
w_,
|
|
h_,
|
|
d_);
|
|
}
|
|
} break;
|
|
|
|
case GPU_TEXTURE_BUFFER: {
|
|
/* TODO(Metal): Support Data upload to TEXTURE BUFFER
|
|
* Data uploads generally happen via GPUVertBuf instead. */
|
|
BLI_assert(false);
|
|
} break;
|
|
|
|
case GPU_TEXTURE_ARRAY:
|
|
/* Not an actual format - modifier flag for others. */
|
|
return;
|
|
}
|
|
|
|
/* If staging texture was used, copy contents to original texture. */
|
|
if (use_staging_texture) {
|
|
/* When using staging texture, copy results into existing texture. */
|
|
BLI_assert(staging_texture != nil);
|
|
blit_encoder = ctx->main_command_buffer.ensure_begin_blit_encoder();
|
|
|
|
/* Copy modified staging texture region back to original texture.
|
|
* Differing blit dimensions based on type. */
|
|
switch (type_) {
|
|
case GPU_TEXTURE_1D:
|
|
case GPU_TEXTURE_1D_ARRAY: {
|
|
int base_slice = (type_ == GPU_TEXTURE_1D_ARRAY) ? offset[1] : 0;
|
|
int final_slice = base_slice + ((type_ == GPU_TEXTURE_1D_ARRAY) ? extent[1] : 1);
|
|
for (int array_index = base_slice; array_index < final_slice; array_index++) {
|
|
[blit_encoder copyFromTexture:staging_texture
|
|
sourceSlice:array_index
|
|
sourceLevel:mip
|
|
sourceOrigin:MTLOriginMake(offset[0], 0, 0)
|
|
sourceSize:MTLSizeMake(extent[0], 1, 1)
|
|
toTexture:texture_
|
|
destinationSlice:array_index
|
|
destinationLevel:mip
|
|
destinationOrigin:MTLOriginMake(offset[0], 0, 0)];
|
|
}
|
|
} break;
|
|
case GPU_TEXTURE_2D:
|
|
case GPU_TEXTURE_2D_ARRAY: {
|
|
int base_slice = (type_ == GPU_TEXTURE_2D_ARRAY) ? offset[2] : 0;
|
|
int final_slice = base_slice + ((type_ == GPU_TEXTURE_2D_ARRAY) ? extent[2] : 1);
|
|
for (int array_index = base_slice; array_index < final_slice; array_index++) {
|
|
[blit_encoder copyFromTexture:staging_texture
|
|
sourceSlice:array_index
|
|
sourceLevel:mip
|
|
sourceOrigin:MTLOriginMake(offset[0], offset[1], 0)
|
|
sourceSize:MTLSizeMake(extent[0], extent[1], 1)
|
|
toTexture:texture_
|
|
destinationSlice:array_index
|
|
destinationLevel:mip
|
|
destinationOrigin:MTLOriginMake(offset[0], offset[1], 0)];
|
|
}
|
|
} break;
|
|
case GPU_TEXTURE_3D: {
|
|
[blit_encoder copyFromTexture:staging_texture
|
|
sourceSlice:0
|
|
sourceLevel:mip
|
|
sourceOrigin:MTLOriginMake(offset[0], offset[1], offset[2])
|
|
sourceSize:MTLSizeMake(extent[0], extent[1], extent[2])
|
|
toTexture:texture_
|
|
destinationSlice:0
|
|
destinationLevel:mip
|
|
destinationOrigin:MTLOriginMake(offset[0], offset[1], offset[2])];
|
|
} break;
|
|
case GPU_TEXTURE_CUBE:
|
|
case GPU_TEXTURE_CUBE_ARRAY: {
|
|
/* Iterate over all cube faces in range (offset[2], offset[2] + extent[2]). */
|
|
for (int i = 0; i < extent[2]; i++) {
|
|
int face_index = offset[2] + i;
|
|
[blit_encoder copyFromTexture:staging_texture
|
|
sourceSlice:face_index
|
|
sourceLevel:mip
|
|
sourceOrigin:MTLOriginMake(offset[0], offset[1], 0)
|
|
sourceSize:MTLSizeMake(extent[0], extent[1], 1)
|
|
toTexture:texture_
|
|
destinationSlice:face_index
|
|
destinationLevel:mip
|
|
destinationOrigin:MTLOriginMake(offset[0], offset[1], 0)];
|
|
}
|
|
} break;
|
|
case GPU_TEXTURE_ARRAY:
|
|
case GPU_TEXTURE_BUFFER:
|
|
BLI_assert_unreachable();
|
|
break;
|
|
}
|
|
|
|
[staging_texture release];
|
|
}
|
|
|
|
/* Finalize Blit Encoder. */
|
|
if (can_use_direct_blit) {
|
|
/* Textures which use MTLStorageModeManaged need to have updated contents
|
|
* synced back to CPU to avoid an automatic flush overwriting contents. */
|
|
if (texture_.storageMode == MTLStorageModeManaged) {
|
|
[blit_encoder synchronizeResource:texture_];
|
|
}
|
|
}
|
|
else {
|
|
/* Textures which use MTLStorageModeManaged need to have updated contents
|
|
* synced back to CPU to avoid an automatic flush overwriting contents. */
|
|
if (texture_.storageMode == MTLStorageModeManaged) {
|
|
blit_encoder = ctx->main_command_buffer.ensure_begin_blit_encoder();
|
|
[blit_encoder synchronizeResource:texture_];
|
|
}
|
|
}
|
|
|
|
/* Decrement texture reference counts. This ensures temporary texture views are released. */
|
|
[texture_handle release];
|
|
}
|
|
}
|
|
|
|
void MTLTexture::update_sub(int offset[3],
|
|
int extent[3],
|
|
eGPUDataFormat format,
|
|
GPUPixelBuffer *pixbuf)
|
|
{
|
|
/* Update texture from pixel buffer. */
|
|
BLI_assert(validate_data_format(format_, format));
|
|
BLI_assert(pixbuf != nullptr);
|
|
|
|
/* Fetch pixel buffer metal buffer. */
|
|
MTLPixelBuffer *mtl_pix_buf = static_cast<MTLPixelBuffer *>(unwrap(pixbuf));
|
|
id<MTLBuffer> buffer = mtl_pix_buf->get_metal_buffer();
|
|
BLI_assert(buffer != nil);
|
|
if (buffer == nil) {
|
|
return;
|
|
}
|
|
|
|
/* Ensure texture is ready. */
|
|
this->ensure_baked();
|
|
BLI_assert(texture_ != nil);
|
|
|
|
/* Calculate dimensions. */
|
|
int num_image_channels = to_component_len(format_);
|
|
|
|
uint bits_per_pixel = num_image_channels * to_bytesize(format);
|
|
uint bytes_per_row = bits_per_pixel * extent[0];
|
|
uint bytes_per_image = bytes_per_row * extent[1];
|
|
|
|
/* Currently only required for 2D textures. */
|
|
if (type_ == GPU_TEXTURE_2D) {
|
|
|
|
/* Create blit command encoder to copy data. */
|
|
MTLContext *ctx = MTLContext::get();
|
|
BLI_assert(ctx);
|
|
|
|
id<MTLBlitCommandEncoder> blit_encoder = ctx->main_command_buffer.ensure_begin_blit_encoder();
|
|
[blit_encoder copyFromBuffer:buffer
|
|
sourceOffset:0
|
|
sourceBytesPerRow:bytes_per_row
|
|
sourceBytesPerImage:bytes_per_image
|
|
sourceSize:MTLSizeMake(extent[0], extent[1], 1)
|
|
toTexture:texture_
|
|
destinationSlice:0
|
|
destinationLevel:0
|
|
destinationOrigin:MTLOriginMake(offset[0], offset[1], 0)];
|
|
|
|
if (texture_.storageMode == MTLStorageModeManaged) {
|
|
[blit_encoder synchronizeResource:texture_];
|
|
}
|
|
}
|
|
else {
|
|
BLI_assert(false);
|
|
}
|
|
}
|
|
|
|
void gpu::MTLTexture::ensure_mipmaps(int miplvl)
|
|
{
|
|
|
|
/* Do not update texture view. */
|
|
BLI_assert(resource_mode_ != MTL_TEXTURE_MODE_TEXTURE_VIEW);
|
|
|
|
/* Clamp level to maximum. */
|
|
int effective_h = (type_ == GPU_TEXTURE_1D_ARRAY) ? 0 : h_;
|
|
int effective_d = (type_ != GPU_TEXTURE_3D) ? 0 : d_;
|
|
int max_dimension = max_iii(w_, effective_h, effective_d);
|
|
int max_miplvl = floor(log2(max_dimension));
|
|
miplvl = min_ii(max_miplvl, miplvl);
|
|
|
|
/* Increase mipmap level. */
|
|
if (mipmaps_ < miplvl) {
|
|
mipmaps_ = miplvl;
|
|
|
|
/* Check if baked. */
|
|
if (is_baked_ && mipmaps_ > mtl_max_mips_) {
|
|
BLI_assert_msg(false,
|
|
"Texture requires a higher mipmap level count. Please specify the required "
|
|
"amount upfront.");
|
|
is_dirty_ = true;
|
|
MTL_LOG_WARNING("Texture requires regenerating due to increase in mip-count\n");
|
|
}
|
|
}
|
|
this->mip_range_set(0, mipmaps_);
|
|
}
|
|
|
|
void gpu::MTLTexture::generate_mipmap()
|
|
{
|
|
/* Fetch Active Context. */
|
|
MTLContext *ctx = reinterpret_cast<MTLContext *>(GPU_context_active_get());
|
|
BLI_assert(ctx);
|
|
|
|
if (!ctx->device) {
|
|
MTL_LOG_ERROR("Cannot Generate mip-maps -- metal device invalid\n");
|
|
BLI_assert(false);
|
|
return;
|
|
}
|
|
|
|
/* Ensure mipmaps. */
|
|
this->ensure_mipmaps(9999);
|
|
|
|
/* Ensure texture is baked. */
|
|
this->ensure_baked();
|
|
BLI_assert_msg(is_baked_ && texture_, "MTLTexture is not valid");
|
|
|
|
if (mipmaps_ == 1 || mtl_max_mips_ == 1) {
|
|
MTL_LOG_WARNING("Call to generate mipmaps on texture with 'mipmaps_=1\n'");
|
|
return;
|
|
}
|
|
|
|
/* Verify if we can perform mipmap generation. */
|
|
if (format_ == GPU_DEPTH_COMPONENT32F || format_ == GPU_DEPTH_COMPONENT24 ||
|
|
format_ == GPU_DEPTH_COMPONENT16 || format_ == GPU_DEPTH32F_STENCIL8 ||
|
|
format_ == GPU_DEPTH24_STENCIL8) {
|
|
MTL_LOG_WARNING("Cannot generate mipmaps for textures using DEPTH formats\n");
|
|
return;
|
|
}
|
|
|
|
@autoreleasepool {
|
|
|
|
/* Fetch active BlitCommandEncoder. */
|
|
id<MTLBlitCommandEncoder> enc = ctx->main_command_buffer.ensure_begin_blit_encoder();
|
|
if (G.debug & G_DEBUG_GPU) {
|
|
[enc insertDebugSignpost:@"Generate MipMaps"];
|
|
}
|
|
[enc generateMipmapsForTexture:texture_];
|
|
has_generated_mips_ = true;
|
|
}
|
|
return;
|
|
}
|
|
|
|
void gpu::MTLTexture::copy_to(Texture *dst)
|
|
{
|
|
/* Safety Checks. */
|
|
gpu::MTLTexture *mt_src = this;
|
|
gpu::MTLTexture *mt_dst = static_cast<gpu::MTLTexture *>(dst);
|
|
BLI_assert((mt_dst->w_ == mt_src->w_) && (mt_dst->h_ == mt_src->h_) &&
|
|
(mt_dst->d_ == mt_src->d_));
|
|
BLI_assert(mt_dst->format_ == mt_src->format_);
|
|
BLI_assert(mt_dst->type_ == mt_src->type_);
|
|
|
|
UNUSED_VARS_NDEBUG(mt_src);
|
|
|
|
/* Fetch active context. */
|
|
MTLContext *ctx = static_cast<MTLContext *>(unwrap(GPU_context_active_get()));
|
|
BLI_assert(ctx);
|
|
|
|
/* Ensure texture is baked. */
|
|
this->ensure_baked();
|
|
|
|
@autoreleasepool {
|
|
/* Setup blit encoder. */
|
|
id<MTLBlitCommandEncoder> blit_encoder = ctx->main_command_buffer.ensure_begin_blit_encoder();
|
|
BLI_assert(blit_encoder != nil);
|
|
|
|
/* TODO(Metal): Consider supporting multiple mip levels IF the GL implementation
|
|
* follows, currently it does not. */
|
|
int mip = 0;
|
|
|
|
/* NOTE: mip_size_get() won't override any dimension that is equal to 0. */
|
|
int extent[3] = {1, 1, 1};
|
|
this->mip_size_get(mip, extent);
|
|
|
|
switch (mt_dst->type_) {
|
|
case GPU_TEXTURE_2D_ARRAY:
|
|
case GPU_TEXTURE_CUBE_ARRAY:
|
|
case GPU_TEXTURE_3D: {
|
|
/* Do full texture copy for 3D textures */
|
|
BLI_assert(mt_dst->d_ == d_);
|
|
[blit_encoder copyFromTexture:this->get_metal_handle_base()
|
|
toTexture:mt_dst->get_metal_handle_base()];
|
|
} break;
|
|
default: {
|
|
int slice = 0;
|
|
this->blit(blit_encoder,
|
|
0,
|
|
0,
|
|
0,
|
|
slice,
|
|
mip,
|
|
mt_dst,
|
|
0,
|
|
0,
|
|
0,
|
|
slice,
|
|
mip,
|
|
extent[0],
|
|
extent[1],
|
|
extent[2]);
|
|
} break;
|
|
}
|
|
}
|
|
}
|
|
|
|
void gpu::MTLTexture::clear(eGPUDataFormat data_format, const void *data)
|
|
{
|
|
/* Ensure texture is baked. */
|
|
this->ensure_baked();
|
|
|
|
/* Create clear framebuffer. */
|
|
GPUFrameBuffer *prev_fb = GPU_framebuffer_active_get();
|
|
FrameBuffer *fb = reinterpret_cast<FrameBuffer *>(this->get_blit_framebuffer(0, 0));
|
|
fb->bind(true);
|
|
fb->clear_attachment(this->attachment_type(0), data_format, data);
|
|
GPU_framebuffer_bind(prev_fb);
|
|
}
|
|
|
|
static MTLTextureSwizzle swizzle_to_mtl(const char swizzle)
|
|
{
|
|
switch (swizzle) {
|
|
default:
|
|
case 'x':
|
|
case 'r':
|
|
return MTLTextureSwizzleRed;
|
|
case 'y':
|
|
case 'g':
|
|
return MTLTextureSwizzleGreen;
|
|
case 'z':
|
|
case 'b':
|
|
return MTLTextureSwizzleBlue;
|
|
case 'w':
|
|
case 'a':
|
|
return MTLTextureSwizzleAlpha;
|
|
case '0':
|
|
return MTLTextureSwizzleZero;
|
|
case '1':
|
|
return MTLTextureSwizzleOne;
|
|
}
|
|
}
|
|
|
|
void gpu::MTLTexture::swizzle_set(const char swizzle_mask[4])
|
|
{
|
|
if (memcmp(tex_swizzle_mask_, swizzle_mask, 4) != 0) {
|
|
memcpy(tex_swizzle_mask_, swizzle_mask, 4);
|
|
|
|
/* Creating the swizzle mask and flagging as dirty if changed. */
|
|
MTLTextureSwizzleChannels new_swizzle_mask = MTLTextureSwizzleChannelsMake(
|
|
swizzle_to_mtl(swizzle_mask[0]),
|
|
swizzle_to_mtl(swizzle_mask[1]),
|
|
swizzle_to_mtl(swizzle_mask[2]),
|
|
swizzle_to_mtl(swizzle_mask[3]));
|
|
|
|
BLI_assert_msg(gpu_image_usage_flags_ & GPU_TEXTURE_USAGE_MIP_SWIZZLE_VIEW,
|
|
"Texture view support is required to change swizzle parameters.");
|
|
mtl_swizzle_mask_ = new_swizzle_mask;
|
|
texture_view_dirty_flags_ |= TEXTURE_VIEW_SWIZZLE_DIRTY;
|
|
}
|
|
}
|
|
|
|
void gpu::MTLTexture::mip_range_set(int min, int max)
|
|
{
|
|
BLI_assert(min <= max && min >= 0 && max <= mipmaps_);
|
|
|
|
/* NOTE:
|
|
* - mip_min_ and mip_max_ are used to Clamp LODs during sampling.
|
|
* - Given functions like Framebuffer::recursive_downsample modifies the mip range
|
|
* between each layer, we do not want to be re-baking the texture.
|
|
* - For the time being, we are going to just need to generate a FULL mipmap chain
|
|
* as we do not know ahead of time whether mipmaps will be used.
|
|
*
|
|
* TODO(Metal): Add texture initialization flag to determine whether mipmaps are used
|
|
* or not. Will be important for saving memory for big textures. */
|
|
mip_min_ = min;
|
|
mip_max_ = max;
|
|
|
|
if ((type_ == GPU_TEXTURE_1D || type_ == GPU_TEXTURE_1D_ARRAY || type_ == GPU_TEXTURE_BUFFER) &&
|
|
max > 1) {
|
|
|
|
MTL_LOG_ERROR(
|
|
" MTLTexture of type TEXTURE_1D_ARRAY or TEXTURE_BUFFER cannot have a mipcount "
|
|
"greater than 1\n");
|
|
mip_min_ = 0;
|
|
mip_max_ = 0;
|
|
mipmaps_ = 0;
|
|
BLI_assert(false);
|
|
}
|
|
|
|
/* Mip range for texture view. */
|
|
mip_texture_base_level_ = mip_min_;
|
|
mip_texture_max_level_ = mip_max_;
|
|
texture_view_dirty_flags_ |= TEXTURE_VIEW_MIP_DIRTY;
|
|
}
|
|
|
|
void *gpu::MTLTexture::read(int mip, eGPUDataFormat type)
|
|
{
|
|
/* Prepare Array for return data. */
|
|
BLI_assert(!(format_flag_ & GPU_FORMAT_COMPRESSED));
|
|
BLI_assert(mip <= mipmaps_);
|
|
BLI_assert(validate_data_format_mtl(format_, type));
|
|
|
|
/* NOTE: mip_size_get() won't override any dimension that is equal to 0. */
|
|
int extent[3] = {1, 1, 1};
|
|
this->mip_size_get(mip, extent);
|
|
|
|
size_t sample_len = extent[0] * extent[1] * extent[2];
|
|
size_t sample_size = to_bytesize(format_, type);
|
|
size_t texture_size = sample_len * sample_size;
|
|
int num_channels = to_component_len(format_);
|
|
|
|
void *data = MEM_mallocN(texture_size + 8, "GPU_texture_read");
|
|
|
|
/* Ensure texture is baked. */
|
|
if (is_baked_) {
|
|
this->read_internal(
|
|
mip, 0, 0, 0, extent[0], extent[1], extent[2], type, num_channels, texture_size + 8, data);
|
|
}
|
|
else {
|
|
/* Clear return values? */
|
|
MTL_LOG_WARNING("MTLTexture::read - reading from texture with no image data\n");
|
|
}
|
|
|
|
return data;
|
|
}
|
|
|
|
/* Fetch the raw buffer data from a texture and copy to CPU host ptr. */
|
|
void gpu::MTLTexture::read_internal(int mip,
|
|
int x_off,
|
|
int y_off,
|
|
int z_off,
|
|
int width,
|
|
int height,
|
|
int depth,
|
|
eGPUDataFormat desired_output_format,
|
|
int num_output_components,
|
|
int debug_data_size,
|
|
void *r_data)
|
|
{
|
|
/* Verify textures are baked. */
|
|
if (!is_baked_) {
|
|
MTL_LOG_WARNING("gpu::MTLTexture::read_internal - Trying to read from a non-baked texture!\n");
|
|
return;
|
|
}
|
|
/* Fetch active context. */
|
|
MTLContext *ctx = static_cast<MTLContext *>(unwrap(GPU_context_active_get()));
|
|
BLI_assert(ctx);
|
|
|
|
/* Calculate Desired output size. */
|
|
int num_channels = to_component_len(format_);
|
|
BLI_assert(num_output_components <= num_channels);
|
|
uint desired_output_bpp = num_output_components * to_bytesize(desired_output_format);
|
|
|
|
/* Calculate Metal data output for trivial copy. */
|
|
uint image_bpp = get_mtl_format_bytesize(texture_.pixelFormat);
|
|
uint image_components = get_mtl_format_num_components(texture_.pixelFormat);
|
|
bool is_depth_format = (format_flag_ & GPU_FORMAT_DEPTH);
|
|
|
|
/* Verify if we need to use compute read. */
|
|
eGPUDataFormat data_format = to_mtl_internal_data_format(this->format_get());
|
|
bool format_conversion_needed = (data_format != desired_output_format);
|
|
bool can_use_simple_read = (desired_output_bpp == image_bpp) && (!format_conversion_needed) &&
|
|
(num_output_components == image_components);
|
|
|
|
/* Depth must be read using the compute shader -- Some safety checks to verify that params are
|
|
* correct. */
|
|
if (is_depth_format) {
|
|
can_use_simple_read = false;
|
|
/* TODO(Metal): Stencil data write not yet supported, so force components to one. */
|
|
image_components = 1;
|
|
BLI_assert(num_output_components == 1);
|
|
BLI_assert(image_components == 1);
|
|
BLI_assert(data_format == GPU_DATA_FLOAT || data_format == GPU_DATA_UINT_24_8);
|
|
BLI_assert(validate_data_format_mtl(format_, data_format));
|
|
}
|
|
|
|
/* SPECIAL Workaround for R11G11B10 textures requesting a read using: GPU_DATA_10_11_11_REV. */
|
|
if (desired_output_format == GPU_DATA_10_11_11_REV) {
|
|
BLI_assert(format_ == GPU_R11F_G11F_B10F);
|
|
|
|
/* override parameters - we'll be able to use simple copy, as bpp will match at 4 bytes. */
|
|
image_bpp = sizeof(int);
|
|
image_components = 1;
|
|
desired_output_bpp = sizeof(int);
|
|
num_output_components = 1;
|
|
|
|
data_format = GPU_DATA_INT;
|
|
format_conversion_needed = false;
|
|
can_use_simple_read = true;
|
|
}
|
|
|
|
/* Determine size of output data. */
|
|
uint bytes_per_row = desired_output_bpp * width;
|
|
uint bytes_per_image = bytes_per_row * height;
|
|
uint total_bytes = bytes_per_image * depth;
|
|
|
|
if (can_use_simple_read) {
|
|
/* DEBUG check that if direct copy is being used, then both the expected output size matches
|
|
* the METAL texture size. */
|
|
BLI_assert(
|
|
((num_output_components * to_bytesize(desired_output_format)) == desired_output_bpp) &&
|
|
(desired_output_bpp == image_bpp));
|
|
}
|
|
/* DEBUG check that the allocated data size matches the bytes we expect. */
|
|
BLI_assert(total_bytes <= debug_data_size);
|
|
|
|
/* Fetch allocation from scratch buffer. */
|
|
gpu::MTLBuffer *dest_buf = MTLContext::get_global_memory_manager()->allocate_aligned(
|
|
total_bytes, 256, true);
|
|
BLI_assert(dest_buf != nullptr);
|
|
|
|
id<MTLBuffer> destination_buffer = dest_buf->get_metal_buffer();
|
|
BLI_assert(destination_buffer != nil);
|
|
void *destination_buffer_host_ptr = dest_buf->get_host_ptr();
|
|
BLI_assert(destination_buffer_host_ptr != nullptr);
|
|
|
|
/* Prepare specialization struct (For non-trivial texture read routine). */
|
|
int depth_format_mode = 0;
|
|
if (is_depth_format) {
|
|
depth_format_mode = 1;
|
|
switch (desired_output_format) {
|
|
case GPU_DATA_FLOAT:
|
|
depth_format_mode = 1;
|
|
break;
|
|
case GPU_DATA_UINT_24_8:
|
|
depth_format_mode = 2;
|
|
break;
|
|
case GPU_DATA_UINT:
|
|
depth_format_mode = 4;
|
|
break;
|
|
default:
|
|
BLI_assert_msg(false, "Unhandled depth read format case");
|
|
break;
|
|
}
|
|
}
|
|
|
|
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 */
|
|
num_output_components, /* OUTPUT DATA COMPONENT COUNT */
|
|
depth_format_mode};
|
|
|
|
bool copy_successful = false;
|
|
@autoreleasepool {
|
|
|
|
/* TODO(Metal): Verify whether we need some form of barrier here to ensure reads
|
|
* happen after work with associated texture is finished. */
|
|
GPU_finish();
|
|
|
|
/* Texture View for SRGB special case. */
|
|
id<MTLTexture> read_texture = texture_;
|
|
if (format_ == GPU_SRGB8_A8) {
|
|
BLI_assert(gpu_image_usage_flags_ & GPU_TEXTURE_USAGE_MIP_SWIZZLE_VIEW);
|
|
read_texture = [texture_ newTextureViewWithPixelFormat:MTLPixelFormatRGBA8Unorm];
|
|
}
|
|
|
|
/* Perform per-texture type read. */
|
|
switch (type_) {
|
|
case GPU_TEXTURE_2D: {
|
|
if (can_use_simple_read) {
|
|
/* Use Blit Encoder READ. */
|
|
id<MTLBlitCommandEncoder> enc = ctx->main_command_buffer.ensure_begin_blit_encoder();
|
|
if (G.debug & G_DEBUG_GPU) {
|
|
[enc insertDebugSignpost:@"GPUTextureRead"];
|
|
}
|
|
[enc copyFromTexture:read_texture
|
|
sourceSlice:0
|
|
sourceLevel:mip
|
|
sourceOrigin:MTLOriginMake(x_off, y_off, 0)
|
|
sourceSize:MTLSizeMake(width, height, 1)
|
|
toBuffer:destination_buffer
|
|
destinationOffset:0
|
|
destinationBytesPerRow:bytes_per_row
|
|
destinationBytesPerImage:bytes_per_image];
|
|
copy_successful = true;
|
|
}
|
|
else {
|
|
|
|
/* Use Compute READ. */
|
|
id<MTLComputeCommandEncoder> compute_encoder =
|
|
ctx->main_command_buffer.ensure_begin_compute_encoder();
|
|
id<MTLComputePipelineState> pso = texture_read_2d_get_kernel(
|
|
compute_specialization_kernel);
|
|
TextureReadParams params = {
|
|
mip,
|
|
{width, height, 1},
|
|
{x_off, y_off, 0},
|
|
};
|
|
|
|
/* Bind resources via compute state for optimal state caching performance. */
|
|
MTLComputeState &cs = ctx->main_command_buffer.get_compute_state();
|
|
cs.bind_pso(pso);
|
|
cs.bind_compute_bytes(¶ms, sizeof(params), 0);
|
|
cs.bind_compute_buffer(destination_buffer, 0, 1);
|
|
cs.bind_compute_texture(read_texture, 0);
|
|
[compute_encoder dispatchThreads:MTLSizeMake(width, height, 1) /* Width, Height, Layer */
|
|
threadsPerThreadgroup:MTLSizeMake(8, 8, 1)];
|
|
copy_successful = true;
|
|
}
|
|
} break;
|
|
|
|
case GPU_TEXTURE_2D_ARRAY: {
|
|
if (can_use_simple_read) {
|
|
/* Use Blit Encoder READ. */
|
|
id<MTLBlitCommandEncoder> enc = ctx->main_command_buffer.ensure_begin_blit_encoder();
|
|
if (G.debug & G_DEBUG_GPU) {
|
|
[enc insertDebugSignpost:@"GPUTextureRead"];
|
|
}
|
|
int base_slice = z_off;
|
|
int final_slice = base_slice + depth;
|
|
int texture_array_relative_offset = 0;
|
|
|
|
for (int array_slice = base_slice; array_slice < final_slice; array_slice++) {
|
|
[enc copyFromTexture:read_texture
|
|
sourceSlice:0
|
|
sourceLevel:mip
|
|
sourceOrigin:MTLOriginMake(x_off, y_off, 0)
|
|
sourceSize:MTLSizeMake(width, height, 1)
|
|
toBuffer:destination_buffer
|
|
destinationOffset:texture_array_relative_offset
|
|
destinationBytesPerRow:bytes_per_row
|
|
destinationBytesPerImage:bytes_per_image];
|
|
texture_array_relative_offset += bytes_per_image;
|
|
}
|
|
copy_successful = true;
|
|
}
|
|
else {
|
|
|
|
/* Use Compute READ */
|
|
id<MTLComputeCommandEncoder> compute_encoder =
|
|
ctx->main_command_buffer.ensure_begin_compute_encoder();
|
|
id<MTLComputePipelineState> pso = texture_read_2d_array_get_kernel(
|
|
compute_specialization_kernel);
|
|
TextureReadParams params = {
|
|
mip,
|
|
{width, height, depth},
|
|
{x_off, y_off, z_off},
|
|
};
|
|
|
|
/* Bind resources via compute state for optimal state caching performance. */
|
|
MTLComputeState &cs = ctx->main_command_buffer.get_compute_state();
|
|
cs.bind_pso(pso);
|
|
cs.bind_compute_bytes(¶ms, sizeof(params), 0);
|
|
cs.bind_compute_buffer(destination_buffer, 0, 1);
|
|
cs.bind_compute_texture(read_texture, 0);
|
|
[compute_encoder
|
|
dispatchThreads:MTLSizeMake(width, height, depth) /* Width, Height, Layer */
|
|
threadsPerThreadgroup:MTLSizeMake(8, 8, 1)];
|
|
copy_successful = true;
|
|
}
|
|
} break;
|
|
|
|
case GPU_TEXTURE_CUBE_ARRAY: {
|
|
if (can_use_simple_read) {
|
|
id<MTLBlitCommandEncoder> enc = ctx->main_command_buffer.ensure_begin_blit_encoder();
|
|
if (G.debug & G_DEBUG_GPU) {
|
|
[enc insertDebugSignpost:@"GPUTextureRead"];
|
|
}
|
|
int base_slice = z_off;
|
|
int final_slice = base_slice + depth;
|
|
int texture_array_relative_offset = 0;
|
|
|
|
for (int array_slice = base_slice; array_slice < final_slice; array_slice++) {
|
|
[enc copyFromTexture:read_texture
|
|
sourceSlice:array_slice
|
|
sourceLevel:mip
|
|
sourceOrigin:MTLOriginMake(x_off, y_off, 0)
|
|
sourceSize:MTLSizeMake(width, height, 1)
|
|
toBuffer:destination_buffer
|
|
destinationOffset:texture_array_relative_offset
|
|
destinationBytesPerRow:bytes_per_row
|
|
destinationBytesPerImage:bytes_per_image];
|
|
|
|
texture_array_relative_offset += bytes_per_image;
|
|
}
|
|
MTL_LOG_INFO("Copying texture data to buffer GPU_TEXTURE_CUBE_ARRAY\n");
|
|
copy_successful = true;
|
|
}
|
|
else {
|
|
MTL_LOG_ERROR("TODO(Metal): unsupported compute copy of texture cube array");
|
|
}
|
|
} break;
|
|
|
|
default:
|
|
MTL_LOG_WARNING(
|
|
"[Warning] gpu::MTLTexture::read_internal simple-copy not yet supported for texture "
|
|
"type: %d\n",
|
|
(int)type_);
|
|
break;
|
|
}
|
|
|
|
if (copy_successful) {
|
|
|
|
/* Use Blit encoder to synchronize results back to CPU. */
|
|
if (dest_buf->get_resource_options() == MTLResourceStorageModeManaged) {
|
|
id<MTLBlitCommandEncoder> enc = ctx->main_command_buffer.ensure_begin_blit_encoder();
|
|
if (G.debug & G_DEBUG_GPU) {
|
|
[enc insertDebugSignpost:@"GPUTextureRead-syncResource"];
|
|
}
|
|
[enc synchronizeResource:destination_buffer];
|
|
}
|
|
|
|
/* Ensure GPU copy commands have completed. */
|
|
GPU_finish();
|
|
|
|
/* Copy data from Shared Memory into ptr. */
|
|
memcpy(r_data, destination_buffer_host_ptr, total_bytes);
|
|
MTL_LOG_INFO("gpu::MTLTexture::read_internal success! %d bytes read\n", total_bytes);
|
|
}
|
|
else {
|
|
MTL_LOG_WARNING(
|
|
"[Warning] gpu::MTLTexture::read_internal not yet supported for this config -- data "
|
|
"format different (src %d bytes, dst %d bytes) (src format: %d, dst format: %d), or "
|
|
"varying component counts (src %d, dst %d)\n",
|
|
image_bpp,
|
|
desired_output_bpp,
|
|
(int)data_format,
|
|
(int)desired_output_format,
|
|
image_components,
|
|
num_output_components);
|
|
}
|
|
|
|
/* Release destination buffer. */
|
|
dest_buf->free();
|
|
}
|
|
}
|
|
|
|
/* Remove once no longer required -- will just return 0 for now in MTL path. */
|
|
uint gpu::MTLTexture::gl_bindcode_get() const
|
|
{
|
|
return 0;
|
|
}
|
|
|
|
bool gpu::MTLTexture::init_internal()
|
|
{
|
|
if (format_ == GPU_DEPTH24_STENCIL8) {
|
|
/* Apple Silicon requires GPU_DEPTH32F_STENCIL8 instead of GPU_DEPTH24_STENCIL8. */
|
|
format_ = GPU_DEPTH32F_STENCIL8;
|
|
}
|
|
|
|
this->prepare_internal();
|
|
return true;
|
|
}
|
|
|
|
bool gpu::MTLTexture::init_internal(GPUVertBuf *vbo)
|
|
{
|
|
if (this->format_ == GPU_DEPTH24_STENCIL8) {
|
|
/* Apple Silicon requires GPU_DEPTH32F_STENCIL8 instead of GPU_DEPTH24_STENCIL8. */
|
|
this->format_ = GPU_DEPTH32F_STENCIL8;
|
|
}
|
|
|
|
MTLPixelFormat mtl_format = gpu_texture_format_to_metal(this->format_);
|
|
mtl_max_mips_ = 1;
|
|
mipmaps_ = 0;
|
|
this->mip_range_set(0, 0);
|
|
|
|
/* Create texture from GPUVertBuf's buffer. */
|
|
MTLVertBuf *mtl_vbo = static_cast<MTLVertBuf *>(unwrap(vbo));
|
|
mtl_vbo->bind();
|
|
mtl_vbo->flag_used();
|
|
|
|
/* Get Metal Buffer. */
|
|
id<MTLBuffer> source_buffer = mtl_vbo->get_metal_buffer();
|
|
BLI_assert(source_buffer);
|
|
|
|
/* Verify size. */
|
|
if (w_ <= 0) {
|
|
MTL_LOG_WARNING("Allocating texture buffer of width 0!\n");
|
|
w_ = 1;
|
|
}
|
|
|
|
/* Verify Texture and vertex buffer alignment. */
|
|
const GPUVertFormat *format = GPU_vertbuf_get_format(vbo);
|
|
int bytes_per_pixel = get_mtl_format_bytesize(mtl_format);
|
|
int bytes_per_row = bytes_per_pixel * w_;
|
|
|
|
MTLContext *mtl_ctx = MTLContext::get();
|
|
uint32_t align_requirement = static_cast<uint32_t>(
|
|
[mtl_ctx->device minimumLinearTextureAlignmentForPixelFormat:mtl_format]);
|
|
|
|
/* If stride is larger than bytes per pixel, but format has multiple attributes,
|
|
* split attributes across several pixels. */
|
|
if (format->stride > bytes_per_pixel && format->attr_len > 1) {
|
|
|
|
/* We need to increase the number of pixels available to store additional attributes.
|
|
* First ensure that the total stride of the vertex format fits uniformly into
|
|
* multiple pixels. If these sizes are different, then attributes are of differing
|
|
* sizes and this operation is unsupported. */
|
|
if (bytes_per_pixel * format->attr_len != format->stride) {
|
|
BLI_assert_msg(false,
|
|
"Cannot split attributes across multiple pixels as attribute format sizes do "
|
|
"not match.");
|
|
return false;
|
|
}
|
|
|
|
/* Provide a single pixel per attribute. */
|
|
/* Increase bytes per row to ensure there are enough bytes for all vertex attribute data. */
|
|
bytes_per_row *= format->attr_len;
|
|
BLI_assert(bytes_per_row == format->stride * w_);
|
|
|
|
/* Multiply width of image to provide one attribute per pixel. */
|
|
w_ *= format->attr_len;
|
|
BLI_assert(bytes_per_row == bytes_per_pixel * w_);
|
|
BLI_assert_msg(w_ == mtl_vbo->vertex_len * format->attr_len,
|
|
"Image should contain one pixel for each attribute in every vertex.");
|
|
}
|
|
else {
|
|
/* Verify per-vertex size aligns with texture size. */
|
|
BLI_assert(bytes_per_pixel == format->stride &&
|
|
"Pixel format stride MUST match the texture format stride -- These being different "
|
|
"is likely caused by Metal's VBO padding to a minimum of 4-bytes per-vertex."
|
|
" If multiple attributes are used. Each attribute is to be packed into its own "
|
|
"individual pixel when stride length is exceeded. ");
|
|
}
|
|
|
|
/* Create texture descriptor. */
|
|
BLI_assert(type_ == GPU_TEXTURE_BUFFER);
|
|
texture_descriptor_ = [[MTLTextureDescriptor alloc] init];
|
|
texture_descriptor_.pixelFormat = mtl_format;
|
|
texture_descriptor_.textureType = MTLTextureTypeTextureBuffer;
|
|
texture_descriptor_.width = w_;
|
|
texture_descriptor_.height = 1;
|
|
texture_descriptor_.depth = 1;
|
|
texture_descriptor_.arrayLength = 1;
|
|
texture_descriptor_.mipmapLevelCount = mtl_max_mips_;
|
|
texture_descriptor_.usage =
|
|
MTLTextureUsageShaderRead | MTLTextureUsageShaderWrite |
|
|
MTLTextureUsagePixelFormatView; /* TODO(Metal): Optimize usage flags. */
|
|
texture_descriptor_.storageMode = [source_buffer storageMode];
|
|
texture_descriptor_.sampleCount = 1;
|
|
texture_descriptor_.cpuCacheMode = [source_buffer cpuCacheMode];
|
|
texture_descriptor_.hazardTrackingMode = [source_buffer hazardTrackingMode];
|
|
|
|
texture_ = [source_buffer
|
|
newTextureWithDescriptor:texture_descriptor_
|
|
offset:0
|
|
bytesPerRow:ceil_to_multiple_u(bytes_per_row, align_requirement)];
|
|
aligned_w_ = bytes_per_row / bytes_per_pixel;
|
|
|
|
BLI_assert(texture_);
|
|
texture_.label = [NSString stringWithUTF8String:this->get_name()];
|
|
is_baked_ = true;
|
|
is_dirty_ = false;
|
|
resource_mode_ = MTL_TEXTURE_MODE_VBO;
|
|
|
|
/* Track Status. */
|
|
vert_buffer_ = mtl_vbo;
|
|
vert_buffer_mtl_ = source_buffer;
|
|
|
|
return true;
|
|
}
|
|
|
|
bool gpu::MTLTexture::init_internal(const GPUTexture *src, int mip_offset, int layer_offset)
|
|
{
|
|
BLI_assert(src);
|
|
|
|
/* Zero initialize. */
|
|
this->prepare_internal();
|
|
|
|
/* Flag as using texture view. */
|
|
resource_mode_ = MTL_TEXTURE_MODE_TEXTURE_VIEW;
|
|
source_texture_ = src;
|
|
mip_texture_base_level_ = mip_offset;
|
|
mip_texture_base_layer_ = layer_offset;
|
|
|
|
/* Assign usage. */
|
|
gpu_image_usage_flags_ = GPU_texture_usage(src);
|
|
|
|
/* Assign texture as view. */
|
|
const gpu::MTLTexture *mtltex = static_cast<const gpu::MTLTexture *>(unwrap(src));
|
|
texture_ = mtltex->texture_;
|
|
BLI_assert(texture_);
|
|
[texture_ retain];
|
|
|
|
/* Flag texture as baked -- we do not need explicit initialization. */
|
|
is_baked_ = true;
|
|
is_dirty_ = false;
|
|
|
|
/* Bake mip swizzle view. */
|
|
bake_mip_swizzle_view();
|
|
return true;
|
|
}
|
|
|
|
/** \} */
|
|
|
|
/* -------------------------------------------------------------------- */
|
|
/** \name METAL Resource creation and management
|
|
* \{ */
|
|
|
|
bool gpu::MTLTexture::texture_is_baked()
|
|
{
|
|
return is_baked_;
|
|
}
|
|
|
|
/* Prepare texture parameters after initialization, but before baking. */
|
|
void gpu::MTLTexture::prepare_internal()
|
|
{
|
|
/* Derive implicit usage flags for Depth/Stencil attachments. */
|
|
if (format_flag_ & GPU_FORMAT_DEPTH || format_flag_ & GPU_FORMAT_STENCIL) {
|
|
gpu_image_usage_flags_ |= GPU_TEXTURE_USAGE_ATTACHMENT;
|
|
}
|
|
|
|
/* Derive maximum number of mip levels by default.
|
|
* TODO(Metal): This can be removed if max mip counts are specified upfront. */
|
|
if (type_ == GPU_TEXTURE_1D || type_ == GPU_TEXTURE_1D_ARRAY || type_ == GPU_TEXTURE_BUFFER) {
|
|
mtl_max_mips_ = 1;
|
|
}
|
|
else {
|
|
/* Require correct explicit mipmap level counts. */
|
|
mtl_max_mips_ = mipmaps_;
|
|
}
|
|
}
|
|
|
|
void gpu::MTLTexture::ensure_baked()
|
|
{
|
|
|
|
/* If properties have changed, re-bake. */
|
|
id<MTLTexture> previous_texture = nil;
|
|
bool copy_previous_contents = false;
|
|
|
|
if (is_baked_ && is_dirty_) {
|
|
copy_previous_contents = true;
|
|
previous_texture = texture_;
|
|
[previous_texture retain];
|
|
this->reset();
|
|
}
|
|
|
|
if (!is_baked_) {
|
|
MTLContext *ctx = static_cast<MTLContext *>(unwrap(GPU_context_active_get()));
|
|
BLI_assert(ctx);
|
|
|
|
/* Ensure texture mode is valid. */
|
|
BLI_assert(resource_mode_ != MTL_TEXTURE_MODE_EXTERNAL);
|
|
BLI_assert(resource_mode_ != MTL_TEXTURE_MODE_TEXTURE_VIEW);
|
|
BLI_assert(resource_mode_ != MTL_TEXTURE_MODE_VBO);
|
|
|
|
/* Format and mip levels (TODO(Metal): Optimize mipmaps counts, specify up-front). */
|
|
MTLPixelFormat mtl_format = gpu_texture_format_to_metal(format_);
|
|
|
|
/* SRGB textures require a texture view for reading data and when rendering with SRGB
|
|
* disabled. Enabling the texture_view or texture_read usage flags disables lossless
|
|
* compression, so the situations in which it is used should be limited. */
|
|
if (format_ == GPU_SRGB8_A8) {
|
|
gpu_image_usage_flags_ = gpu_image_usage_flags_ | GPU_TEXTURE_USAGE_MIP_SWIZZLE_VIEW;
|
|
}
|
|
|
|
/* Create texture descriptor. */
|
|
switch (type_) {
|
|
|
|
/* 1D */
|
|
case GPU_TEXTURE_1D:
|
|
case GPU_TEXTURE_1D_ARRAY: {
|
|
BLI_assert(w_ > 0);
|
|
texture_descriptor_ = [[MTLTextureDescriptor alloc] init];
|
|
texture_descriptor_.pixelFormat = mtl_format;
|
|
texture_descriptor_.textureType = (type_ == GPU_TEXTURE_1D_ARRAY) ? MTLTextureType1DArray :
|
|
MTLTextureType1D;
|
|
texture_descriptor_.width = w_;
|
|
texture_descriptor_.height = 1;
|
|
texture_descriptor_.depth = 1;
|
|
texture_descriptor_.arrayLength = (type_ == GPU_TEXTURE_1D_ARRAY) ? h_ : 1;
|
|
texture_descriptor_.mipmapLevelCount = (mtl_max_mips_ > 0) ? mtl_max_mips_ : 1;
|
|
texture_descriptor_.usage = mtl_usage_from_gpu(gpu_image_usage_flags_);
|
|
texture_descriptor_.storageMode = MTLStorageModePrivate;
|
|
texture_descriptor_.sampleCount = 1;
|
|
texture_descriptor_.cpuCacheMode = MTLCPUCacheModeDefaultCache;
|
|
texture_descriptor_.hazardTrackingMode = MTLHazardTrackingModeDefault;
|
|
} break;
|
|
|
|
/* 2D */
|
|
case GPU_TEXTURE_2D:
|
|
case GPU_TEXTURE_2D_ARRAY: {
|
|
BLI_assert(w_ > 0 && h_ > 0);
|
|
texture_descriptor_ = [[MTLTextureDescriptor alloc] init];
|
|
texture_descriptor_.pixelFormat = mtl_format;
|
|
texture_descriptor_.textureType = (type_ == GPU_TEXTURE_2D_ARRAY) ? MTLTextureType2DArray :
|
|
MTLTextureType2D;
|
|
texture_descriptor_.width = w_;
|
|
texture_descriptor_.height = h_;
|
|
texture_descriptor_.depth = 1;
|
|
texture_descriptor_.arrayLength = (type_ == GPU_TEXTURE_2D_ARRAY) ? d_ : 1;
|
|
texture_descriptor_.mipmapLevelCount = (mtl_max_mips_ > 0) ? mtl_max_mips_ : 1;
|
|
texture_descriptor_.usage = mtl_usage_from_gpu(gpu_image_usage_flags_);
|
|
texture_descriptor_.storageMode = MTLStorageModePrivate;
|
|
texture_descriptor_.sampleCount = 1;
|
|
texture_descriptor_.cpuCacheMode = MTLCPUCacheModeDefaultCache;
|
|
texture_descriptor_.hazardTrackingMode = MTLHazardTrackingModeDefault;
|
|
} break;
|
|
|
|
/* 3D */
|
|
case GPU_TEXTURE_3D: {
|
|
BLI_assert(w_ > 0 && h_ > 0 && d_ > 0);
|
|
texture_descriptor_ = [[MTLTextureDescriptor alloc] init];
|
|
texture_descriptor_.pixelFormat = mtl_format;
|
|
texture_descriptor_.textureType = MTLTextureType3D;
|
|
texture_descriptor_.width = w_;
|
|
texture_descriptor_.height = h_;
|
|
texture_descriptor_.depth = d_;
|
|
texture_descriptor_.arrayLength = 1;
|
|
texture_descriptor_.mipmapLevelCount = (mtl_max_mips_ > 0) ? mtl_max_mips_ : 1;
|
|
texture_descriptor_.usage = mtl_usage_from_gpu(gpu_image_usage_flags_);
|
|
texture_descriptor_.storageMode = MTLStorageModePrivate;
|
|
texture_descriptor_.sampleCount = 1;
|
|
texture_descriptor_.cpuCacheMode = MTLCPUCacheModeDefaultCache;
|
|
texture_descriptor_.hazardTrackingMode = MTLHazardTrackingModeDefault;
|
|
} break;
|
|
|
|
/* CUBE TEXTURES */
|
|
case GPU_TEXTURE_CUBE:
|
|
case GPU_TEXTURE_CUBE_ARRAY: {
|
|
/* NOTE: For a cube-map 'Texture::d_' refers to total number of faces,
|
|
* not just array slices. */
|
|
BLI_assert(w_ > 0 && h_ > 0);
|
|
texture_descriptor_ = [[MTLTextureDescriptor alloc] init];
|
|
texture_descriptor_.pixelFormat = mtl_format;
|
|
texture_descriptor_.textureType = (type_ == GPU_TEXTURE_CUBE_ARRAY) ?
|
|
MTLTextureTypeCubeArray :
|
|
MTLTextureTypeCube;
|
|
texture_descriptor_.width = w_;
|
|
texture_descriptor_.height = h_;
|
|
texture_descriptor_.depth = 1;
|
|
texture_descriptor_.arrayLength = (type_ == GPU_TEXTURE_CUBE_ARRAY) ? d_ / 6 : 1;
|
|
texture_descriptor_.mipmapLevelCount = (mtl_max_mips_ > 0) ? mtl_max_mips_ : 1;
|
|
texture_descriptor_.usage = mtl_usage_from_gpu(gpu_image_usage_flags_);
|
|
texture_descriptor_.storageMode = MTLStorageModePrivate;
|
|
texture_descriptor_.sampleCount = 1;
|
|
texture_descriptor_.cpuCacheMode = MTLCPUCacheModeDefaultCache;
|
|
texture_descriptor_.hazardTrackingMode = MTLHazardTrackingModeDefault;
|
|
} break;
|
|
|
|
/* GPU_TEXTURE_BUFFER */
|
|
case GPU_TEXTURE_BUFFER: {
|
|
texture_descriptor_ = [[MTLTextureDescriptor alloc] init];
|
|
texture_descriptor_.pixelFormat = mtl_format;
|
|
texture_descriptor_.textureType = MTLTextureTypeTextureBuffer;
|
|
texture_descriptor_.width = w_;
|
|
texture_descriptor_.height = 1;
|
|
texture_descriptor_.depth = 1;
|
|
texture_descriptor_.arrayLength = 1;
|
|
texture_descriptor_.mipmapLevelCount = (mtl_max_mips_ > 0) ? mtl_max_mips_ : 1;
|
|
texture_descriptor_.usage = mtl_usage_from_gpu(gpu_image_usage_flags_);
|
|
texture_descriptor_.storageMode = MTLStorageModePrivate;
|
|
texture_descriptor_.sampleCount = 1;
|
|
texture_descriptor_.cpuCacheMode = MTLCPUCacheModeDefaultCache;
|
|
texture_descriptor_.hazardTrackingMode = MTLHazardTrackingModeDefault;
|
|
} break;
|
|
|
|
default: {
|
|
MTL_LOG_ERROR("[METAL] Error: Cannot create texture with unknown type: %d\n", type_);
|
|
return;
|
|
} break;
|
|
}
|
|
|
|
/* Determine Resource Mode. */
|
|
resource_mode_ = MTL_TEXTURE_MODE_DEFAULT;
|
|
|
|
/* Standard texture allocation. */
|
|
texture_ = [ctx->device newTextureWithDescriptor:texture_descriptor_];
|
|
|
|
texture_.label = [NSString stringWithUTF8String:this->get_name()];
|
|
BLI_assert(texture_);
|
|
is_baked_ = true;
|
|
is_dirty_ = false;
|
|
}
|
|
|
|
/* Re-apply previous contents. */
|
|
if (copy_previous_contents) {
|
|
/* TODO(Metal): May need to copy previous contents of texture into new texture. */
|
|
[previous_texture release];
|
|
}
|
|
}
|
|
|
|
void gpu::MTLTexture::reset()
|
|
{
|
|
|
|
MTL_LOG_INFO("Texture %s reset. Size %d, %d, %d\n", this->get_name(), w_, h_, d_);
|
|
/* Delete associated METAL resources. */
|
|
if (texture_ != nil) {
|
|
[texture_ release];
|
|
texture_ = nil;
|
|
is_baked_ = false;
|
|
is_dirty_ = true;
|
|
}
|
|
|
|
if (texture_no_srgb_ != nil) {
|
|
[texture_no_srgb_ release];
|
|
texture_no_srgb_ = nil;
|
|
}
|
|
|
|
if (mip_swizzle_view_ != nil) {
|
|
[mip_swizzle_view_ release];
|
|
mip_swizzle_view_ = nil;
|
|
}
|
|
|
|
if (texture_buffer_ != nil) {
|
|
[texture_buffer_ release];
|
|
}
|
|
|
|
/* Blit framebuffer. */
|
|
if (blit_fb_) {
|
|
GPU_framebuffer_free(blit_fb_);
|
|
blit_fb_ = nullptr;
|
|
}
|
|
|
|
/* Descriptor. */
|
|
if (texture_descriptor_ != nullptr) {
|
|
[texture_descriptor_ release];
|
|
texture_descriptor_ = nullptr;
|
|
}
|
|
|
|
/* Reset mipmap state. */
|
|
has_generated_mips_ = false;
|
|
|
|
BLI_assert(texture_ == nil);
|
|
BLI_assert(mip_swizzle_view_ == nil);
|
|
}
|
|
|
|
/** \} */
|
|
|
|
/* -------------------------------------------------------------------- */
|
|
/** \name SRGB Handling.
|
|
* \{ */
|
|
bool MTLTexture::is_format_srgb()
|
|
{
|
|
return (format_ == GPU_SRGB8_A8);
|
|
}
|
|
|
|
id<MTLTexture> MTLTexture::get_non_srgb_handle()
|
|
{
|
|
id<MTLTexture> base_tex = get_metal_handle_base();
|
|
BLI_assert(base_tex != nil);
|
|
if (texture_no_srgb_ == nil) {
|
|
texture_no_srgb_ = [base_tex newTextureViewWithPixelFormat:MTLPixelFormatRGBA8Unorm];
|
|
}
|
|
return texture_no_srgb_;
|
|
}
|
|
|
|
/** \} */
|
|
/* -------------------------------------------------------------------- */
|
|
/** \name Pixel Buffer
|
|
* \{ */
|
|
|
|
MTLPixelBuffer::MTLPixelBuffer(uint size) : PixelBuffer(size)
|
|
{
|
|
MTLContext *ctx = MTLContext::get();
|
|
BLI_assert(ctx);
|
|
/* Ensure buffer satisfies the alignment of 256 bytes for copying
|
|
* data between buffers and textures. As specified in:
|
|
* https://developer.apple.com/metal/Metal-Feature-Set-Tables.pdf */
|
|
BLI_assert(size >= 256);
|
|
|
|
MTLResourceOptions resource_options = ([ctx->device hasUnifiedMemory]) ?
|
|
MTLResourceStorageModeShared :
|
|
MTLResourceStorageModeManaged;
|
|
buffer_ = [ctx->device newBufferWithLength:size options:resource_options];
|
|
BLI_assert(buffer_ != nil);
|
|
}
|
|
|
|
MTLPixelBuffer::~MTLPixelBuffer()
|
|
{
|
|
if (buffer_) {
|
|
[buffer_ release];
|
|
buffer_ = nil;
|
|
}
|
|
}
|
|
|
|
void *MTLPixelBuffer::map()
|
|
{
|
|
if (buffer_ == nil) {
|
|
return nullptr;
|
|
}
|
|
|
|
return [buffer_ contents];
|
|
}
|
|
|
|
void MTLPixelBuffer::unmap()
|
|
{
|
|
if (buffer_ == nil) {
|
|
return;
|
|
}
|
|
|
|
/* Ensure changes are synchronized. */
|
|
if (buffer_.resourceOptions & MTLResourceStorageModeManaged) {
|
|
[buffer_ didModifyRange:NSMakeRange(0, size_)];
|
|
}
|
|
}
|
|
|
|
int64_t MTLPixelBuffer::get_native_handle()
|
|
{
|
|
if (buffer_ == nil) {
|
|
return 0;
|
|
}
|
|
|
|
return reinterpret_cast<int64_t>(buffer_);
|
|
}
|
|
|
|
uint MTLPixelBuffer::get_size()
|
|
{
|
|
return size_;
|
|
}
|
|
|
|
id<MTLBuffer> MTLPixelBuffer::get_metal_buffer()
|
|
{
|
|
return buffer_;
|
|
}
|
|
|
|
/** \} */
|
|
|
|
} // namespace blender::gpu
|