Differential D16415 Diff 58437 source/blender/gpu/metal/mtl_texture.mm

Changeset View

Standalone View

source/blender/gpu/metal/mtl_texture.mm

Show First 20 Lines • Show All 1,266 Lines • ▼ Show 20 Lines	if (can_use_simple_read) {
BLI_assert(		BLI_assert(
((num_output_components * to_bytesize(desired_output_format)) == desired_output_bpp) &&		((num_output_components * to_bytesize(desired_output_format)) == desired_output_bpp) &&
(desired_output_bpp == image_bpp));		(desired_output_bpp == image_bpp));
}		}
/* DEBUG check that the allocated data size matches the bytes we expect. */		/* DEBUG check that the allocated data size matches the bytes we expect. */
BLI_assert(total_bytes <= debug_data_size);		BLI_assert(total_bytes <= debug_data_size);

/* Fetch allocation from scratch buffer. */		/* Fetch allocation from scratch buffer. */
id<MTLBuffer> destination_buffer = nil;		gpu::MTLBuffer *dest_buf = MTLContext::get_global_memory_manager()->allocate_aligned(
uint destination_offset = 0;		total_bytes, 256, true);
void *destination_buffer_host_ptr = nullptr;		BLI_assert(dest_buf != nullptr);

/* TODO(Metal): Optimize buffer allocation. */		id<MTLBuffer> destination_buffer = dest_buf->get_metal_buffer();
MTLResourceOptions bufferOptions = MTLResourceStorageModeManaged;		BLI_assert(destination_buffer != nil);
destination_buffer = [ctx->device newBufferWithLength:max_ii(total_bytes, 256)		void *destination_buffer_host_ptr = dest_buf->get_host_ptr();
options:bufferOptions];		BLI_assert(destination_buffer_host_ptr != nullptr);
destination_offset = 0;
destination_buffer_host_ptr = (void )((uint8_t )([destination_buffer contents]) +
destination_offset);

/* Prepare specialization struct (For non-trivial texture read routine). */		/* Prepare specialization struct (For non-trivial texture read routine). */
int depth_format_mode = 0;		int depth_format_mode = 0;
if (is_depth_format) {		if (is_depth_format) {
depth_format_mode = 1;		depth_format_mode = 1;
switch (desired_output_format) {		switch (desired_output_format) {
case GPU_DATA_FLOAT:		case GPU_DATA_FLOAT:
depth_format_mode = 1;		depth_format_mode = 1;
Show All 40 Lines	switch (type_) {
[enc insertDebugSignpost:@"GPUTextureRead"];		[enc insertDebugSignpost:@"GPUTextureRead"];
}		}
[enc copyFromTexture:read_texture		[enc copyFromTexture:read_texture
sourceSlice:0		sourceSlice:0
sourceLevel:mip		sourceLevel:mip
sourceOrigin:MTLOriginMake(x_off, y_off, 0)		sourceOrigin:MTLOriginMake(x_off, y_off, 0)
sourceSize:MTLSizeMake(width, height, 1)		sourceSize:MTLSizeMake(width, height, 1)
toBuffer:destination_buffer		toBuffer:destination_buffer
destinationOffset:destination_offset		destinationOffset:0
destinationBytesPerRow:bytes_per_row		destinationBytesPerRow:bytes_per_row
destinationBytesPerImage:bytes_per_image];		destinationBytesPerImage:bytes_per_image];
[enc synchronizeResource:destination_buffer];
copy_successful = true;		copy_successful = true;
}		}
else {		else {

/* Use Compute READ. */		/* Use Compute READ. */
id<MTLComputeCommandEncoder> compute_encoder =		id<MTLComputeCommandEncoder> compute_encoder =
ctx->main_command_buffer.ensure_begin_compute_encoder();		ctx->main_command_buffer.ensure_begin_compute_encoder();
id<MTLComputePipelineState> pso = texture_read_2d_get_kernel(		id<MTLComputePipelineState> pso = texture_read_2d_get_kernel(
compute_specialization_kernel);		compute_specialization_kernel);
TextureReadParams params = {		TextureReadParams params = {
mip,		mip,
{width, height, 1},		{width, height, 1},
{x_off, y_off, 0},		{x_off, y_off, 0},
};		};
[compute_encoder setComputePipelineState:pso];		[compute_encoder setComputePipelineState:pso];
[compute_encoder setBytes:&params length:sizeof(params) atIndex:0];		[compute_encoder setBytes:&params length:sizeof(params) atIndex:0];
[compute_encoder setBuffer:destination_buffer offset:destination_offset atIndex:1];		[compute_encoder setBuffer:destination_buffer offset:0 atIndex:1];
[compute_encoder setTexture:read_texture atIndex:0];		[compute_encoder setTexture:read_texture atIndex:0];
[compute_encoder dispatchThreads:MTLSizeMake(width, height, 1) /* Width, Height, Layer */		[compute_encoder dispatchThreads:MTLSizeMake(width, height, 1) /* Width, Height, Layer */
threadsPerThreadgroup:MTLSizeMake(8, 8, 1)];		threadsPerThreadgroup:MTLSizeMake(8, 8, 1)];

/* Use Blit encoder to synchronize results back to CPU. */
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];
copy_successful = true;		copy_successful = true;
}		}
} break;		} break;

case GPU_TEXTURE_2D_ARRAY: {		case GPU_TEXTURE_2D_ARRAY: {
if (can_use_simple_read) {		if (can_use_simple_read) {
/* Use Blit Encoder READ. */		/* Use Blit Encoder READ. */
id<MTLBlitCommandEncoder> enc = ctx->main_command_buffer.ensure_begin_blit_encoder();		id<MTLBlitCommandEncoder> enc = ctx->main_command_buffer.ensure_begin_blit_encoder();
if (G.debug & G_DEBUG_GPU) {		if (G.debug & G_DEBUG_GPU) {
[enc insertDebugSignpost:@"GPUTextureRead"];		[enc insertDebugSignpost:@"GPUTextureRead"];
}		}
int base_slice = z_off;		int base_slice = z_off;
int final_slice = base_slice + depth;		int final_slice = base_slice + depth;
int texture_array_relative_offset = 0;		int texture_array_relative_offset = 0;

for (int array_slice = base_slice; array_slice < final_slice; array_slice++) {		for (int array_slice = base_slice; array_slice < final_slice; array_slice++) {
[enc copyFromTexture:read_texture		[enc copyFromTexture:read_texture
sourceSlice:0		sourceSlice:0
sourceLevel:mip		sourceLevel:mip
sourceOrigin:MTLOriginMake(x_off, y_off, 0)		sourceOrigin:MTLOriginMake(x_off, y_off, 0)
sourceSize:MTLSizeMake(width, height, 1)		sourceSize:MTLSizeMake(width, height, 1)
toBuffer:destination_buffer		toBuffer:destination_buffer
destinationOffset:destination_offset + texture_array_relative_offset		destinationOffset:texture_array_relative_offset
destinationBytesPerRow:bytes_per_row		destinationBytesPerRow:bytes_per_row
destinationBytesPerImage:bytes_per_image];		destinationBytesPerImage:bytes_per_image];
[enc synchronizeResource:destination_buffer];

texture_array_relative_offset += bytes_per_image;		texture_array_relative_offset += bytes_per_image;
}		}
copy_successful = true;		copy_successful = true;
}		}
else {		else {

/* Use Compute READ */		/* Use Compute READ */
id<MTLComputeCommandEncoder> compute_encoder =		id<MTLComputeCommandEncoder> compute_encoder =
ctx->main_command_buffer.ensure_begin_compute_encoder();		ctx->main_command_buffer.ensure_begin_compute_encoder();
id<MTLComputePipelineState> pso = texture_read_2d_array_get_kernel(		id<MTLComputePipelineState> pso = texture_read_2d_array_get_kernel(
compute_specialization_kernel);		compute_specialization_kernel);
TextureReadParams params = {		TextureReadParams params = {
mip,		mip,
{width, height, depth},		{width, height, depth},
{x_off, y_off, z_off},		{x_off, y_off, z_off},
};		};
[compute_encoder setComputePipelineState:pso];		[compute_encoder setComputePipelineState:pso];
[compute_encoder setBytes:&params length:sizeof(params) atIndex:0];		[compute_encoder setBytes:&params length:sizeof(params) atIndex:0];
[compute_encoder setBuffer:destination_buffer offset:destination_offset atIndex:1];		[compute_encoder setBuffer:destination_buffer offset:0 atIndex:1];
[compute_encoder setTexture:read_texture atIndex:0];		[compute_encoder setTexture:read_texture atIndex:0];
[compute_encoder		[compute_encoder
dispatchThreads:MTLSizeMake(width, height, depth) /* Width, Height, Layer */		dispatchThreads:MTLSizeMake(width, height, depth) /* Width, Height, Layer */
threadsPerThreadgroup:MTLSizeMake(8, 8, 1)];		threadsPerThreadgroup:MTLSizeMake(8, 8, 1)];

/* Use Blit encoder to synchronize results back to CPU. */
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];
copy_successful = true;		copy_successful = true;
}		}
} break;		} break;

case GPU_TEXTURE_CUBE_ARRAY: {		case GPU_TEXTURE_CUBE_ARRAY: {
if (can_use_simple_read) {		if (can_use_simple_read) {
id<MTLBlitCommandEncoder> enc = ctx->main_command_buffer.ensure_begin_blit_encoder();		id<MTLBlitCommandEncoder> enc = ctx->main_command_buffer.ensure_begin_blit_encoder();
if (G.debug & G_DEBUG_GPU) {		if (G.debug & G_DEBUG_GPU) {
[enc insertDebugSignpost:@"GPUTextureRead"];		[enc insertDebugSignpost:@"GPUTextureRead"];
}		}
int base_slice = z_off;		int base_slice = z_off;
int final_slice = base_slice + depth;		int final_slice = base_slice + depth;
int texture_array_relative_offset = 0;		int texture_array_relative_offset = 0;

for (int array_slice = base_slice; array_slice < final_slice; array_slice++) {		for (int array_slice = base_slice; array_slice < final_slice; array_slice++) {
[enc copyFromTexture:read_texture		[enc copyFromTexture:read_texture
sourceSlice:array_slice		sourceSlice:array_slice
sourceLevel:mip		sourceLevel:mip
sourceOrigin:MTLOriginMake(x_off, y_off, 0)		sourceOrigin:MTLOriginMake(x_off, y_off, 0)
sourceSize:MTLSizeMake(width, height, 1)		sourceSize:MTLSizeMake(width, height, 1)
toBuffer:destination_buffer		toBuffer:destination_buffer
destinationOffset:destination_offset + texture_array_relative_offset		destinationOffset:texture_array_relative_offset
destinationBytesPerRow:bytes_per_row		destinationBytesPerRow:bytes_per_row
destinationBytesPerImage:bytes_per_image];		destinationBytesPerImage:bytes_per_image];
[enc synchronizeResource:destination_buffer];

texture_array_relative_offset += bytes_per_image;		texture_array_relative_offset += bytes_per_image;
}		}
MTL_LOG_INFO("Copying texture data to buffer GPU_TEXTURE_CUBE_ARRAY\n");		MTL_LOG_INFO("Copying texture data to buffer GPU_TEXTURE_CUBE_ARRAY\n");
copy_successful = true;		copy_successful = true;
}		}
else {		else {
MTL_LOG_ERROR("TODO(Metal): unsupported compute copy of texture cube array");		MTL_LOG_ERROR("TODO(Metal): unsupported compute copy of texture cube array");
}		}
} break;		} break;

default:		default:
MTL_LOG_WARNING(		MTL_LOG_WARNING(
"[Warning] gpu::MTLTexture::read_internal simple-copy not yet supported for texture "		"[Warning] gpu::MTLTexture::read_internal simple-copy not yet supported for texture "
"type: %d\n",		"type: %d\n",
(int)type_);		(int)type_);
break;		break;
}		}

if (copy_successful) {		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. */		/* Ensure GPU copy commands have completed. */
GPU_finish();		GPU_finish();

/* Copy data from Shared Memory into ptr. */		/* Copy data from Shared Memory into ptr. */
memcpy(r_data, destination_buffer_host_ptr, total_bytes);		memcpy(r_data, destination_buffer_host_ptr, total_bytes);
MTL_LOG_INFO("gpu::MTLTexture::read_internal success! %d bytes read\n", total_bytes);		MTL_LOG_INFO("gpu::MTLTexture::read_internal success! %d bytes read\n", total_bytes);
}		}
else {		else {
MTL_LOG_WARNING(		MTL_LOG_WARNING(
"[Warning] gpu::MTLTexture::read_internal not yet supported for this config -- data "		"[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 "		"format different (src %d bytes, dst %d bytes) (src format: %d, dst format: %d), or "
"varying component counts (src %d, dst %d)\n",		"varying component counts (src %d, dst %d)\n",
image_bpp,		image_bpp,
desired_output_bpp,		desired_output_bpp,
(int)data_format,		(int)data_format,
(int)desired_output_format,		(int)desired_output_format,
image_components,		image_components,
num_output_components);		num_output_components);
}		}

		/* Release destination buffer. */
		dest_buf->free();
}		}
}		}

/* Remove once no longer required -- will just return 0 for now in MTL path. */		/* Remove once no longer required -- will just return 0 for now in MTL path. */
uint gpu::MTLTexture::gl_bindcode_get() const		uint gpu::MTLTexture::gl_bindcode_get() const
{		{
return 0;		return 0;
}		}
▲ Show 20 Lines • Show All 151 Lines • ▼ Show 20 Lines	else {
mtl_max_mips_ = max_miplvl;		mtl_max_mips_ = max_miplvl;
}		}
}		}

void gpu::MTLTexture::ensure_baked()		void gpu::MTLTexture::ensure_baked()
{		{

/* If properties have changed, re-bake. */		/* If properties have changed, re-bake. */
		id<MTLTexture> previous_texture = nil;
bool copy_previous_contents = false;		bool copy_previous_contents = false;

if (is_baked_ && is_dirty_) {		if (is_baked_ && is_dirty_) {
copy_previous_contents = true;		copy_previous_contents = true;
id<MTLTexture> previous_texture = texture_;		previous_texture = texture_;
[previous_texture retain];		[previous_texture retain];

this->reset();		this->reset();
}		}

if (!is_baked_) {		if (!is_baked_) {
MTLContext ctx = static_cast<MTLContext >(unwrap(GPU_context_active_get()));		MTLContext ctx = static_cast<MTLContext >(unwrap(GPU_context_active_get()));
BLI_assert(ctx);		BLI_assert(ctx);

/* Ensure texture mode is valid. */		/* Ensure texture mode is valid. */
▲ Show 20 Lines • Show All 132 Lines • ▼ Show 20 Lines	if (!is_baked_) {
texture_.label = [NSString stringWithUTF8String:this->get_name()];		texture_.label = [NSString stringWithUTF8String:this->get_name()];
BLI_assert(texture_);		BLI_assert(texture_);
is_baked_ = true;		is_baked_ = true;
is_dirty_ = false;		is_dirty_ = false;
}		}

/* Re-apply previous contents. */		/* Re-apply previous contents. */
if (copy_previous_contents) {		if (copy_previous_contents) {
id<MTLTexture> previous_texture;
/* TODO(Metal): May need to copy previous contents of texture into new texture. */		/* TODO(Metal): May need to copy previous contents of texture into new texture. */
/[previous_texture release]; /		[previous_texture release];
UNUSED_VARS(previous_texture);
}		}
}		}

void gpu::MTLTexture::reset()		void gpu::MTLTexture::reset()
{		{

MTL_LOG_INFO("Texture %s reset. Size %d, %d, %d\n", this->get_name(), w_, h_, d_);		MTL_LOG_INFO("Texture %s reset. Size %d, %d, %d\n", this->get_name(), w_, h_, d_);
/* Delete associated METAL resources. */		/* Delete associated METAL resources. */
▲ Show 20 Lines • Show All 99 Lines • Show Last 20 Lines