← Back to branch summary

~mmach/netext73/mesa-haswell

« back to all changes in this revision

Viewing changes to src/amd/vulkan/radv_meta_resolve_cs.c

  • Committer: mmach
  • Date: 2022-09-22 20:02:48 UTC
  • Revision ID: netbit73@gmail.com-20220922200248-7y4wybmdgipuwdiw
2022-09-22 21:17:09

Show diffs side-by-side

added added

removed removed

Lines of Context:
src/amd/vulkan/radv_meta_resolve_cs.c
61
61
static nir_shader *
62
62
build_resolve_compute_shader(struct radv_device *dev, bool is_integer, bool is_srgb, int samples)
63
63
{
 
64
   enum glsl_base_type img_base_type = is_integer ? GLSL_TYPE_UINT : GLSL_TYPE_FLOAT;
64
65
   const struct glsl_type *sampler_type =
65
 
      glsl_sampler_type(GLSL_SAMPLER_DIM_MS, false, false, GLSL_TYPE_FLOAT);
66
 
   const struct glsl_type *img_type = glsl_image_type(GLSL_SAMPLER_DIM_2D, false, GLSL_TYPE_FLOAT);
67
 
   nir_builder b = radv_meta_init_shader(MESA_SHADER_COMPUTE, "meta_resolve_cs-%d-%s", samples,
 
66
      glsl_sampler_type(GLSL_SAMPLER_DIM_MS, false, false, img_base_type);
 
67
   const struct glsl_type *img_type = glsl_image_type(GLSL_SAMPLER_DIM_2D, false, img_base_type);
 
68
   nir_builder b = radv_meta_init_shader(dev, MESA_SHADER_COMPUTE, "meta_resolve_cs-%d-%s", samples,
68
69
                                         is_integer ? "int" : (is_srgb ? "srgb" : "float"));
69
70
   b.shader->info.workgroup_size[0] = 8;
70
71
   b.shader->info.workgroup_size[1] = 8;
79
80
 
80
81
   nir_ssa_def *global_id = get_global_ids(&b, 2);
81
82
 
82
 
   nir_ssa_def *src_offset = nir_load_push_constant(&b, 2, 32, nir_imm_int(&b, 0), .range = 16);
 
83
   nir_ssa_def *src_offset = nir_load_push_constant(&b, 2, 32, nir_imm_int(&b, 0), .range = 8);
83
84
   nir_ssa_def *dst_offset = nir_load_push_constant(&b, 2, 32, nir_imm_int(&b, 8), .range = 16);
84
85
 
85
86
   nir_ssa_def *src_coord = nir_iadd(&b, global_id, src_offset);
130
131
build_depth_stencil_resolve_compute_shader(struct radv_device *dev, int samples, int index,
131
132
                                           VkResolveModeFlagBits resolve_mode)
132
133
{
 
134
   enum glsl_base_type img_base_type = index == DEPTH_RESOLVE ? GLSL_TYPE_FLOAT : GLSL_TYPE_UINT;
133
135
   const struct glsl_type *sampler_type =
134
 
      glsl_sampler_type(GLSL_SAMPLER_DIM_MS, false, true, GLSL_TYPE_FLOAT);
135
 
   const struct glsl_type *img_type = glsl_image_type(GLSL_SAMPLER_DIM_2D, true, GLSL_TYPE_FLOAT);
 
136
      glsl_sampler_type(GLSL_SAMPLER_DIM_MS, false, true, img_base_type);
 
137
   const struct glsl_type *img_type = glsl_image_type(GLSL_SAMPLER_DIM_2D, true, img_base_type);
136
138
 
137
 
   nir_builder b = radv_meta_init_shader(MESA_SHADER_COMPUTE, "meta_resolve_cs_%s-%s-%d",
 
139
   nir_builder b = radv_meta_init_shader(dev, MESA_SHADER_COMPUTE, "meta_resolve_cs_%s-%s-%d",
138
140
                                         index == DEPTH_RESOLVE ? "depth" : "stencil",
139
141
                                         get_resolve_mode_str(resolve_mode), samples);
140
142
   b.shader->info.workgroup_size[0] = 8;
373
375
 
374
376
   res = create_layout(device);
375
377
   if (res != VK_SUCCESS)
376
 
      goto fail;
 
378
      return res;
377
379
 
378
380
   if (on_demand)
379
381
      return VK_SUCCESS;
384
386
      res = create_resolve_pipeline(device, samples, false, false,
385
387
                                    &state->resolve_compute.rc[i].pipeline);
386
388
      if (res != VK_SUCCESS)
387
 
         goto fail;
 
389
         return res;
388
390
 
389
391
      res = create_resolve_pipeline(device, samples, true, false,
390
392
                                    &state->resolve_compute.rc[i].i_pipeline);
391
393
      if (res != VK_SUCCESS)
392
 
         goto fail;
 
394
         return res;
393
395
 
394
396
      res = create_resolve_pipeline(device, samples, false, true,
395
397
                                    &state->resolve_compute.rc[i].srgb_pipeline);
396
398
      if (res != VK_SUCCESS)
397
 
         goto fail;
 
399
         return res;
398
400
 
399
401
      res = create_depth_stencil_resolve_pipeline(
400
402
         device, samples, DEPTH_RESOLVE, VK_RESOLVE_MODE_AVERAGE_BIT,
401
403
         &state->resolve_compute.depth[i].average_pipeline);
402
404
      if (res != VK_SUCCESS)
403
 
         goto fail;
 
405
         return res;
404
406
 
405
407
      res = create_depth_stencil_resolve_pipeline(device, samples, DEPTH_RESOLVE,
406
408
                                                  VK_RESOLVE_MODE_MAX_BIT,
407
409
                                                  &state->resolve_compute.depth[i].max_pipeline);
408
410
      if (res != VK_SUCCESS)
409
 
         goto fail;
 
411
         return res;
410
412
 
411
413
      res = create_depth_stencil_resolve_pipeline(device, samples, DEPTH_RESOLVE,
412
414
                                                  VK_RESOLVE_MODE_MIN_BIT,
413
415
                                                  &state->resolve_compute.depth[i].min_pipeline);
414
416
      if (res != VK_SUCCESS)
415
 
         goto fail;
 
417
         return res;
416
418
 
417
419
      res = create_depth_stencil_resolve_pipeline(device, samples, STENCIL_RESOLVE,
418
420
                                                  VK_RESOLVE_MODE_MAX_BIT,
419
421
                                                  &state->resolve_compute.stencil[i].max_pipeline);
420
422
      if (res != VK_SUCCESS)
421
 
         goto fail;
 
423
         return res;
422
424
 
423
425
      res = create_depth_stencil_resolve_pipeline(device, samples, STENCIL_RESOLVE,
424
426
                                                  VK_RESOLVE_MODE_MIN_BIT,
425
427
                                                  &state->resolve_compute.stencil[i].min_pipeline);
426
428
      if (res != VK_SUCCESS)
427
 
         goto fail;
 
429
         return res;
428
430
   }
429
431
 
430
432
   res = create_depth_stencil_resolve_pipeline(device, 0, DEPTH_RESOLVE,
431
433
                                               VK_RESOLVE_MODE_SAMPLE_ZERO_BIT,
432
434
                                               &state->resolve_compute.depth_zero_pipeline);
433
435
   if (res != VK_SUCCESS)
434
 
      goto fail;
435
 
 
436
 
   res = create_depth_stencil_resolve_pipeline(device, 0, STENCIL_RESOLVE,
437
 
                                               VK_RESOLVE_MODE_SAMPLE_ZERO_BIT,
438
 
                                               &state->resolve_compute.stencil_zero_pipeline);
439
 
   if (res != VK_SUCCESS)
440
 
      goto fail;
441
 
 
442
 
   return VK_SUCCESS;
443
 
fail:
444
 
   radv_device_finish_meta_resolve_compute_state(device);
445
 
   return res;
 
436
      return res;
 
437
 
 
438
   return create_depth_stencil_resolve_pipeline(device, 0, STENCIL_RESOLVE,
 
439
                                                VK_RESOLVE_MODE_SAMPLE_ZERO_BIT,
 
440
                                                &state->resolve_compute.stencil_zero_pipeline);
446
441
}
447
442
 
448
443
void
481
476
   radv_DestroyPipeline(radv_device_to_handle(device), state->resolve_compute.stencil_zero_pipeline,
482
477
                        &state->alloc);
483
478
 
484
 
   radv_DestroyDescriptorSetLayout(radv_device_to_handle(device), state->resolve_compute.ds_layout,
485
 
                                   &state->alloc);
 
479
   device->vk.dispatch_table.DestroyDescriptorSetLayout(
 
480
      radv_device_to_handle(device), state->resolve_compute.ds_layout, &state->alloc);
486
481
   radv_DestroyPipelineLayout(radv_device_to_handle(device), state->resolve_compute.p_layout,
487
482
                              &state->alloc);
488
483
}
496
491
   uint32_t samples_log2 = ffs(samples) - 1;
497
492
   VkPipeline *pipeline;
498
493
 
499
 
   if (vk_format_is_int(src_iview->vk_format))
 
494
   if (vk_format_is_int(src_iview->vk.format))
500
495
      pipeline = &state->resolve_compute.rc[samples_log2].i_pipeline;
501
 
   else if (vk_format_is_srgb(src_iview->vk_format))
 
496
   else if (vk_format_is_srgb(src_iview->vk.format))
502
497
      pipeline = &state->resolve_compute.rc[samples_log2].srgb_pipeline;
503
498
   else
504
499
      pipeline = &state->resolve_compute.rc[samples_log2].pipeline;
506
501
   if (!*pipeline) {
507
502
      VkResult ret;
508
503
 
509
 
      ret = create_resolve_pipeline(device, samples, vk_format_is_int(src_iview->vk_format),
510
 
                                    vk_format_is_srgb(src_iview->vk_format), pipeline);
 
504
      ret = create_resolve_pipeline(device, samples, vk_format_is_int(src_iview->vk.format),
 
505
                                    vk_format_is_srgb(src_iview->vk.format), pipeline);
511
506
      if (ret != VK_SUCCESS) {
512
507
         cmd_buffer->record_result = ret;
513
508
         return NULL;
700
695
   const uint32_t dest_base_layer =
701
696
      radv_meta_get_iview_layer(dest_image, &region->dstSubresource, &region->dstOffset);
702
697
 
703
 
   const struct VkExtent3D extent = radv_sanitize_image_extent(src_image->type, region->extent);
704
 
   const struct VkOffset3D srcOffset =
705
 
      radv_sanitize_image_offset(src_image->type, region->srcOffset);
706
 
   const struct VkOffset3D dstOffset =
707
 
      radv_sanitize_image_offset(dest_image->type, region->dstOffset);
 
698
   const struct VkExtent3D extent = vk_image_sanitize_extent(&src_image->vk, region->extent);
 
699
   const struct VkOffset3D srcOffset = vk_image_sanitize_offset(&src_image->vk, region->srcOffset);
 
700
   const struct VkOffset3D dstOffset = vk_image_sanitize_offset(&dest_image->vk, region->dstOffset);
708
701
 
709
702
   for (uint32_t layer = 0; layer < region->srcSubresource.layerCount; ++layer) {
710
703
 
724
717
                                    .layerCount = 1,
725
718
                                 },
726
719
                           },
727
 
                           NULL);
 
720
                           0, NULL);
728
721
 
729
722
      struct radv_image_view dest_iview;
730
723
      radv_image_view_init(&dest_iview, cmd_buffer->device,
742
735
                                    .layerCount = 1,
743
736
                                 },
744
737
                           },
745
 
                           NULL);
 
738
                           0, NULL);
746
739
 
747
740
      emit_resolve(cmd_buffer, &src_iview, &dest_iview, &(VkOffset2D){srcOffset.x, srcOffset.y},
748
741
                   &(VkOffset2D){dstOffset.x, dstOffset.y},
810
803
         .srcSubresource =
811
804
            (VkImageSubresourceLayers){
812
805
               .aspectMask = VK_IMAGE_ASPECT_COLOR_BIT,
813
 
               .mipLevel = src_iview->base_mip,
814
 
               .baseArrayLayer = src_iview->base_layer,
 
806
               .mipLevel = src_iview->vk.base_mip_level,
 
807
               .baseArrayLayer = src_iview->vk.base_array_layer,
815
808
               .layerCount = layer_count,
816
809
            },
817
810
         .dstSubresource =
818
811
            (VkImageSubresourceLayers){
819
812
               .aspectMask = VK_IMAGE_ASPECT_COLOR_BIT,
820
 
               .mipLevel = dst_iview->base_mip,
821
 
               .baseArrayLayer = dst_iview->base_layer,
 
813
               .mipLevel = dst_iview->vk.base_mip_level,
 
814
               .baseArrayLayer = dst_iview->vk.base_array_layer,
822
815
               .layerCount = layer_count,
823
816
            },
824
817
         .srcOffset = (VkOffset3D){0, 0, 0},
825
818
         .dstOffset = (VkOffset3D){0, 0, 0},
826
819
      };
827
820
 
828
 
      radv_meta_resolve_compute_image(cmd_buffer, src_iview->image, src_iview->vk_format,
829
 
                                      src_att.layout, dst_iview->image, dst_iview->vk_format,
 
821
      radv_meta_resolve_compute_image(cmd_buffer, src_iview->image, src_iview->vk.format,
 
822
                                      src_att.layout, dst_iview->image, dst_iview->vk.format,
830
823
                                      dst_att.layout, &region);
831
824
   }
832
825
 
864
857
   region.sType = VK_STRUCTURE_TYPE_IMAGE_RESOLVE_2;
865
858
   region.srcSubresource.aspectMask = aspects;
866
859
   region.srcSubresource.mipLevel = 0;
867
 
   region.srcSubresource.baseArrayLayer = src_iview->base_layer;
 
860
   region.srcSubresource.baseArrayLayer = src_iview->vk.base_array_layer;
868
861
   region.srcSubresource.layerCount = layer_count;
869
862
 
870
863
   radv_decompress_resolve_src(cmd_buffer, src_image, src_att.layout, &region);
882
875
                           .sType = VK_STRUCTURE_TYPE_IMAGE_VIEW_CREATE_INFO,
883
876
                           .image = radv_image_to_handle(src_image),
884
877
                           .viewType = radv_meta_get_view_type(src_image),
885
 
                           .format = src_iview->vk_format,
 
878
                           .format = src_iview->vk.format,
886
879
                           .subresourceRange =
887
880
                              {
888
881
                                 .aspectMask = aspects,
889
 
                                 .baseMipLevel = src_iview->base_mip,
 
882
                                 .baseMipLevel = src_iview->vk.base_mip_level,
890
883
                                 .levelCount = 1,
891
 
                                 .baseArrayLayer = src_iview->base_layer,
 
884
                                 .baseArrayLayer = src_iview->vk.base_array_layer,
892
885
                                 .layerCount = layer_count,
893
886
                              },
894
887
                        },
895
 
                        NULL);
 
888
                        0, NULL);
896
889
 
897
890
   struct radv_image_view tdst_iview;
898
891
   radv_image_view_init(&tdst_iview, cmd_buffer->device,
900
893
                           .sType = VK_STRUCTURE_TYPE_IMAGE_VIEW_CREATE_INFO,
901
894
                           .image = radv_image_to_handle(dst_image),
902
895
                           .viewType = radv_meta_get_view_type(dst_image),
903
 
                           .format = dst_iview->vk_format,
 
896
                           .format = dst_iview->vk.format,
904
897
                           .subresourceRange =
905
898
                              {
906
899
                                 .aspectMask = aspects,
907
 
                                 .baseMipLevel = dst_iview->base_mip,
 
900
                                 .baseMipLevel = dst_iview->vk.base_mip_level,
908
901
                                 .levelCount = 1,
909
 
                                 .baseArrayLayer = dst_iview->base_layer,
 
902
                                 .baseArrayLayer = dst_iview->vk.base_array_layer,
910
903
                                 .layerCount = layer_count,
911
904
                              },
912
905
                        },
913
 
                        NULL);
 
906
                        0, NULL);
914
907
 
915
908
   emit_depth_stencil_resolve(cmd_buffer, &tsrc_iview, &tdst_iview,
916
909
                              &(VkExtent3D){fb->width, fb->height, layer_count}, aspects,
927
920
   if (radv_layout_is_htile_compressed(cmd_buffer->device, dst_image, layout, false, queue_mask)) {
928
921
      VkImageSubresourceRange range = {0};
929
922
      range.aspectMask = aspects;
930
 
      range.baseMipLevel = dst_iview->base_mip;
 
923
      range.baseMipLevel = dst_iview->vk.base_mip_level;
931
924
      range.levelCount = 1;
932
 
      range.baseArrayLayer = dst_iview->base_layer;
 
925
      range.baseArrayLayer = dst_iview->vk.base_array_layer;
933
926
      range.layerCount = layer_count;
934
927
 
935
928
      uint32_t htile_value = radv_get_htile_initial_value(cmd_buffer->device, dst_image);