Page MenuHome
Differential D15563 Diff 55217 source/blender/gpu/metal/mtl_texture.mm

Changeset View

Standalone View

View Options

source/blender/gpu/metal/mtl_texture.mm

Show First 20 Lines • Show All 473 Lines • ▼ Show 20 Lines @autoreleasepool {
/* Check */ /* Check */
BLI_assert(totalsize > 0); BLI_assert(totalsize > 0);
/* Determine expected destination data size. */ /* Determine expected destination data size. */
MTLPixelFormat destination_format = gpu_texture_format_to_metal(format_); MTLPixelFormat destination_format = gpu_texture_format_to_metal(format_);
int expected_dst_bytes_per_pixel = get_mtl_format_bytesize(destination_format); int expected_dst_bytes_per_pixel = get_mtl_format_bytesize(destination_format);
int destination_num_channels = get_mtl_format_num_components(destination_format); int destination_num_channels = get_mtl_format_num_components(destination_format);
/* Prepare specialisation struct (For texture update routine). */ /* Prepare specialization struct (For texture update routine). */
TextureUpdateRoutineSpecialisation compute_specialisation_kernel = { TextureUpdateRoutineSpecialisation compute_specialization_kernel = {
tex_data_format_to_msl_type_str(type), /* INPUT DATA FORMAT */ tex_data_format_to_msl_type_str(type), /* INPUT DATA FORMAT */
tex_data_format_to_msl_texture_template_type(type), /* TEXTURE DATA FORMAT */ tex_data_format_to_msl_texture_template_type(type), /* TEXTURE DATA FORMAT */
num_channels, num_channels,
destination_num_channels}; destination_num_channels};
/* Determine whether we can do direct BLIT or not. */ /* Determine whether we can do direct BLIT or not. */
bool can_use_direct_blit = true; bool can_use_direct_blit = true;
if (expected_dst_bytes_per_pixel != input_bytes_per_pixel || if (expected_dst_bytes_per_pixel != input_bytes_per_pixel ||
▲ Show 20 Lines • Show All 123 Lines • ▼ Show 20 Lines switch (type_) {
destinationLevel:mip destinationLevel:mip
destinationOrigin:MTLOriginMake(offset[0], 0, 0)]; destinationOrigin:MTLOriginMake(offset[0], 0, 0)];
} }
} }
else { else {
/* Use Compute Based update. */ /* Use Compute Based update. */
if (type_ == GPU_TEXTURE_1D) { if (type_ == GPU_TEXTURE_1D) {
id<MTLComputePipelineState> pso = texture_update_1d_get_kernel( id<MTLComputePipelineState> pso = texture_update_1d_get_kernel(
compute_specialisation_kernel); compute_specialization_kernel);
TextureUpdateParams params = {mip, TextureUpdateParams params = {mip,
{extent[0], 1, 1}, {extent[0], 1, 1},
{offset[0], 0, 0}, {offset[0], 0, 0},
((ctx->pipeline_state.unpack_row_length == 0) ? ((ctx->pipeline_state.unpack_row_length == 0) ?
extent[0] : extent[0] :
ctx->pipeline_state.unpack_row_length)}; ctx->pipeline_state.unpack_row_length)};
[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:staging_buffer offset:staging_buffer_offset atIndex:1]; [compute_encoder setBuffer:staging_buffer offset:staging_buffer_offset atIndex:1];
[compute_encoder setTexture:texture_handle atIndex:0]; [compute_encoder setTexture:texture_handle atIndex:0];
[compute_encoder [compute_encoder
dispatchThreads:MTLSizeMake(extent[0], 1, 1) /* Width, Height, Layer */ dispatchThreads:MTLSizeMake(extent[0], 1, 1) /* Width, Height, Layer */
threadsPerThreadgroup:MTLSizeMake(64, 1, 1)]; threadsPerThreadgroup:MTLSizeMake(64, 1, 1)];
} }
else if (type_ == GPU_TEXTURE_1D_ARRAY) { else if (type_ == GPU_TEXTURE_1D_ARRAY) {
id<MTLComputePipelineState> pso = texture_update_1d_array_get_kernel( id<MTLComputePipelineState> pso = texture_update_1d_array_get_kernel(
compute_specialisation_kernel); compute_specialization_kernel);
TextureUpdateParams params = {mip, TextureUpdateParams params = {mip,
{extent[0], extent[1], 1}, {extent[0], extent[1], 1},
{offset[0], offset[1], 0}, {offset[0], offset[1], 0},
((ctx->pipeline_state.unpack_row_length == 0) ? ((ctx->pipeline_state.unpack_row_length == 0) ?
extent[0] : extent[0] :
ctx->pipeline_state.unpack_row_length)}; ctx->pipeline_state.unpack_row_length)};
[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];
Show All 40 Lines switch (type_) {
texture_array_relative_offset += bytes_per_image; texture_array_relative_offset += bytes_per_image;
} }
} }
else { else {
/* Use Compute texture update. */ /* Use Compute texture update. */
if (type_ == GPU_TEXTURE_2D) { if (type_ == GPU_TEXTURE_2D) {
id<MTLComputePipelineState> pso = texture_update_2d_get_kernel( id<MTLComputePipelineState> pso = texture_update_2d_get_kernel(
compute_specialisation_kernel); compute_specialization_kernel);
TextureUpdateParams params = {mip, TextureUpdateParams params = {mip,
{extent[0], extent[1], 1}, {extent[0], extent[1], 1},
{offset[0], offset[1], 0}, {offset[0], offset[1], 0},
((ctx->pipeline_state.unpack_row_length == 0) ? ((ctx->pipeline_state.unpack_row_length == 0) ?
extent[0] : extent[0] :
ctx->pipeline_state.unpack_row_length)}; ctx->pipeline_state.unpack_row_length)};
[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:staging_buffer offset:staging_buffer_offset atIndex:1]; [compute_encoder setBuffer:staging_buffer offset:staging_buffer_offset atIndex:1];
[compute_encoder setTexture:texture_handle atIndex:0]; [compute_encoder setTexture:texture_handle atIndex:0];
[compute_encoder [compute_encoder
dispatchThreads:MTLSizeMake( dispatchThreads:MTLSizeMake(
extent[0], extent[1], 1) /* Width, Height, Layer */ extent[0], extent[1], 1) /* Width, Height, Layer */
threadsPerThreadgroup:MTLSizeMake(8, 8, 1)]; threadsPerThreadgroup:MTLSizeMake(8, 8, 1)];
} }
else if (type_ == GPU_TEXTURE_2D_ARRAY) { else if (type_ == GPU_TEXTURE_2D_ARRAY) {
id<MTLComputePipelineState> pso = texture_update_2d_array_get_kernel( id<MTLComputePipelineState> pso = texture_update_2d_array_get_kernel(
compute_specialisation_kernel); compute_specialization_kernel);
TextureUpdateParams params = {mip, TextureUpdateParams params = {mip,
{extent[0], extent[1], extent[2]}, {extent[0], extent[1], extent[2]},
{offset[0], offset[1], offset[2]}, {offset[0], offset[1], offset[2]},
((ctx->pipeline_state.unpack_row_length == 0) ? ((ctx->pipeline_state.unpack_row_length == 0) ?
extent[0] : extent[0] :
ctx->pipeline_state.unpack_row_length)}; ctx->pipeline_state.unpack_row_length)};
[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];
Show All 23 Lines switch (type_) {
sourceSize:MTLSizeMake(extent[0], extent[1], extent[2]) sourceSize:MTLSizeMake(extent[0], extent[1], extent[2])
toTexture:texture_handle toTexture:texture_handle
destinationSlice:0 destinationSlice:0
destinationLevel:mip destinationLevel:mip
destinationOrigin:MTLOriginMake(offset[0], offset[1], offset[2])]; destinationOrigin:MTLOriginMake(offset[0], offset[1], offset[2])];
} }
else { else {
id<MTLComputePipelineState> pso = texture_update_3d_get_kernel( id<MTLComputePipelineState> pso = texture_update_3d_get_kernel(
compute_specialisation_kernel); compute_specialization_kernel);
TextureUpdateParams params = {mip, TextureUpdateParams params = {mip,
{extent[0], extent[1], extent[2]}, {extent[0], extent[1], extent[2]},
{offset[0], offset[1], offset[2]}, {offset[0], offset[1], offset[2]},
((ctx->pipeline_state.unpack_row_length == 0) ? ((ctx->pipeline_state.unpack_row_length == 0) ?
extent[0] : extent[0] :
ctx->pipeline_state.unpack_row_length)}; ctx->pipeline_state.unpack_row_length)};
[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];
▲ Show 20 Lines • Show All 447 Lines • ▼ Show 20 Linesvoid gpu::MTLTexture::read_internal(int mip,
/* TODO(Metal): Optimize buffer allocation. */ /* TODO(Metal): Optimize buffer allocation. */
MTLResourceOptions bufferOptions = MTLResourceStorageModeManaged; MTLResourceOptions bufferOptions = MTLResourceStorageModeManaged;
destination_buffer = [ctx->device newBufferWithLength:max_ii(total_bytes, 256) destination_buffer = [ctx->device newBufferWithLength:max_ii(total_bytes, 256)
options:bufferOptions]; options:bufferOptions];
destination_offset = 0; destination_offset = 0;
destination_buffer_host_ptr = (void *)((uint8_t *)([destination_buffer contents]) + destination_buffer_host_ptr = (void *)((uint8_t *)([destination_buffer contents]) +
destination_offset); 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; 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;
break; break;
case GPU_DATA_UINT_24_8: case GPU_DATA_UINT_24_8:
depth_format_mode = 2; depth_format_mode = 2;
break; break;
case GPU_DATA_UINT: case GPU_DATA_UINT:
depth_format_mode = 4; depth_format_mode = 4;
break; break;
default: default:
BLI_assert(false && "Unhandled depth read format case"); BLI_assert(false && "Unhandled depth read format case");
break; break;
} }
} }
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_texture_template_type(data_format), /* TEXTURE DATA TYPE */
tex_data_format_to_msl_type_str(desired_output_format), /* OUTPUT DATA TYPE */ tex_data_format_to_msl_type_str(desired_output_format), /* OUTPUT DATA TYPE */
num_channels, /* TEXTURE COMPONENT COUNT */ num_channels, /* TEXTURE COMPONENT COUNT */
num_output_components, /* OUTPUT DATA COMPONENT COUNT */ num_output_components, /* OUTPUT DATA COMPONENT COUNT */
depth_format_mode}; depth_format_mode};
bool copy_successful = false; bool copy_successful = false;
@autoreleasepool { @autoreleasepool {
Show All 30 Lines switch (type_) {
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_specialisation_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:destination_offset atIndex:1];
Show All 39 Lines switch (type_) {
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_specialisation_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:destination_offset atIndex:1];
▲ Show 20 Lines • Show All 365 Lines • Show Last 20 Lines