Changeset View
Changeset View
Standalone View
Standalone View
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:¶ms length:sizeof(params) atIndex:0]; | [compute_encoder setBytes:¶ms 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:¶ms length:sizeof(params) atIndex:0]; | [compute_encoder setBytes:¶ms 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:¶ms length:sizeof(params) atIndex:0]; | [compute_encoder setBytes:¶ms 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:¶ms length:sizeof(params) atIndex:0]; | [compute_encoder setBytes:¶ms 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:¶ms length:sizeof(params) atIndex:0]; | [compute_encoder setBytes:¶ms length:sizeof(params) atIndex:0]; | ||||
| ▲ Show 20 Lines • Show All 447 Lines • ▼ Show 20 Lines | void 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:¶ms length:sizeof(params) atIndex:0]; | [compute_encoder setBytes:¶ms 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:¶ms length:sizeof(params) atIndex:0]; | [compute_encoder setBytes:¶ms 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 | |||||