Skip to content

Commit fa58fdd

Browse files
SS-JIASS-JIA
committed
[ET-VK][Ez] Fix Validation Layer warnings about wrong image layout
Pull Request resolved: #2854 ## Context Currently, when executing a `ComputeGraph` with prepacked tensors with [Vulkan Validation Layers](https://github.com/KhronosGroup/Vulkan-ValidationLayers) turned on, the following Validation Errors can be observed. Note that Validation Layers can be turned on by running Vulkan binaries on Mac with the `vkconfig` app opened. ``` UNASSIGNED-CoreValidation-DrawState-InvalidImageLayout(ERROR / SPEC): msgNum: 1303270965 - Validation Error: [ UNASSIGNED-CoreValidation-DrawState-InvalidImageLayout ] Object 0: handle = 0x7fb76dbbf988, type = VK_OBJECT_TYPE_COMMAND_BUFFER; | MessageID = 0x4dae5635 | vkQueueSubmit(): pSubmits[0].pCommandBuffers[0] command buffer VkCommandBuffer 0x7fb76dbbf988[] expects VkImage 0xd79c8a0000000f09[] (subresource: aspectMask 0x1 array layer 0, mip level 0) to be in layout VK_IMAGE_LAYOUT_GENERAL--instead, current layout is VK_IMAGE_LAYOUT_SHADER_READ_ONLY_OPTIMAL. Objects: 1 [0] 0x7fb76dbbf988, type: 6, name: NULL ``` The reason for this is that prepacked textures are written to with `WRITE` memory access during packing, which means they will be in the `VK_IMAGE_LAYOUT_GENERAL` layout. However, they will subsequently be read from during `graph.execute()`, meaning the texture will have transitioned to `VK_IMAGE_LAYOUT_SHADER_READ_ONLY_OPTIMAL`, but will be bound using the `VK_IMAGE_LAYOUT_GENERAL` layout. Subsequent calls to `execute()` will therefore see that the prepacked texture has been bound with the wrong layout, since after the first graph execution the texture will have the `VK_IMAGE_LAYOUT_SHADER_READ_ONLY_OPTIMAL` layout. The solution is to submit a no-op shader dispatch during prepacking to trigger a transition to the `READ_ONLY_OPTIMAL` layout. ghstack-source-id: 221868425 Differential Revision: [D55772003](https://our.internmc.facebook.com/intern/diff/D55772003/)
1 parent 5d299fe commit fa58fdd

File tree

6 files changed

+115
-16
lines changed

6 files changed

+115
-16
lines changed

‎backends/vulkan/runtime/graph/ops/PrepackNode.cpp

Lines changed: 49 additions & 16 deletions
Original file line numberDiff line numberDiff line change
@@ -11,10 +11,19 @@
1111
#include <executorch/backends/vulkan/runtime/graph/ComputeGraph.h>
1212

1313
#include <executorch/backends/vulkan/runtime/graph/ops/utils/BindingUtils.h>
14+
#include <executorch/backends/vulkan/runtime/graph/ops/utils/ShaderNameUtils.h>
1415
#include <executorch/backends/vulkan/runtime/graph/ops/utils/StagingUtils.h>
1516

1617
namespace vkcompute {
1718

19+
api::ShaderInfo get_noop_shader(ComputeGraph& graph, const ValueRef packed) {
20+
std::stringstream noop_shader_name;
21+
noop_shader_name << "no_op";
22+
apply_ndim_suffix(noop_shader_name, graph.get_val(packed).toTensor());
23+
apply_dtype_suffix(noop_shader_name, graph.get_val(packed).toTensor());
24+
return VK_KERNEL_FROM_STR(noop_shader_name.str());
25+
}
26+
1827
PrepackNode::PrepackNode(
1928
ComputeGraph& graph,
2029
const api::ShaderInfo& shader,
@@ -24,17 +33,18 @@ PrepackNode::PrepackNode(
2433
const ValueRef packed,
2534
const std::vector<std::shared_ptr<api::UniformParamsBuffer>>& params)
2635
: shader_(shader),
36+
noop_shader_(get_noop_shader(graph, packed)),
2737
global_workgroup_size_(global_workgroup_size),
2838
local_workgroup_size_(local_workgroup_size),
2939
tref_(tref),
3040
packed_(packed),
3141
params_(params) {
3242
graph.update_descriptor_counts(shader, /*execute = */ false);
43+
graph.update_descriptor_counts(noop_shader_, /*execute = */ false);
3344
}
3445

3546
void PrepackNode::encode(ComputeGraph* graph) {
3647
api::Context* const context = graph->context();
37-
api::PipelineBarrier pipeline_barrier{};
3848

3949
TensorRef& tref = graph->get_val(tref_).toTensorRef();
4050
vTensor& packed = graph->get_val(packed_).toTensor();
@@ -46,21 +56,44 @@ void PrepackNode::encode(ComputeGraph* graph) {
4656

4757
std::unique_lock<std::mutex> cmd_lock = context->dispatch_lock();
4858

49-
api::DescriptorSet descriptor_set =
50-
context->get_descriptor_set(shader_, local_workgroup_size_);
51-
52-
uint32_t idx = 0;
53-
bind_tensor_to_descriptor_set(
54-
packed,
55-
pipeline_barrier,
56-
api::MemoryAccessType::WRITE,
57-
descriptor_set,
58-
idx++);
59-
bind_staging_to_descriptor_set(staging, descriptor_set, idx++);
60-
bind_params_to_descriptor_set(params_, descriptor_set, idx);
61-
62-
context->register_shader_dispatch(
63-
descriptor_set, pipeline_barrier, shader_, global_workgroup_size_);
59+
{
60+
api::PipelineBarrier pipeline_barrier{};
61+
api::DescriptorSet descriptor_set =
62+
context->get_descriptor_set(shader_, local_workgroup_size_);
63+
64+
uint32_t idx = 0;
65+
bind_tensor_to_descriptor_set(
66+
packed,
67+
pipeline_barrier,
68+
api::MemoryAccessType::WRITE,
69+
descriptor_set,
70+
idx++);
71+
bind_staging_to_descriptor_set(staging, descriptor_set, idx++);
72+
bind_params_to_descriptor_set(params_, descriptor_set, idx);
73+
74+
context->register_shader_dispatch(
75+
descriptor_set, pipeline_barrier, shader_, global_workgroup_size_);
76+
}
77+
78+
// Submit a compute shader that performs a no-op with the packed tensor in
79+
// order to trigger a image layout transition from GENERAL to
80+
// READ_ONLY_OPTIMAL. This ensures that future uses of the tensor will be
81+
// bound with the correct image layout.
82+
{
83+
api::PipelineBarrier pipeline_barrier{};
84+
api::DescriptorSet descriptor_set =
85+
context->get_descriptor_set(noop_shader_, {1, 1, 1});
86+
87+
bind_tensor_to_descriptor_set(
88+
packed,
89+
pipeline_barrier,
90+
api::MemoryAccessType::READ,
91+
descriptor_set,
92+
0);
93+
94+
context->register_shader_dispatch(
95+
descriptor_set, pipeline_barrier, noop_shader_, {1, 1, 1});
96+
}
6497
}
6598

6699
} // namespace vkcompute

‎backends/vulkan/runtime/graph/ops/PrepackNode.h

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -41,6 +41,7 @@ class PrepackNode final {
4141

4242
protected:
4343
const api::ShaderInfo shader_;
44+
api::ShaderInfo noop_shader_;
4445
const api::utils::uvec3 global_workgroup_size_;
4546
const api::utils::uvec3 local_workgroup_size_;
4647
const ValueRef tref_;

‎backends/vulkan/runtime/graph/ops/glsl/no_op.glsl

Lines changed: 24 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,24 @@
1+
/*
2+
* Copyright (c) Meta Platforms, Inc. and affiliates.
3+
* All rights reserved.
4+
*
5+
* This source code is licensed under the BSD-style license found in the
6+
* LICENSE file in the root directory of this source tree.
7+
*/
8+
9+
#version 450 core
10+
11+
#include "broadcasting_utils.h"
12+
#include "indexing_utils.h"
13+
14+
#define PRECISION ${PRECISION}
15+
16+
#define OP(X, Y, A) ${OPERATOR}
17+
18+
layout(std430) buffer;
19+
20+
layout(set = 0, binding = 0) uniform PRECISION ${SAMPLER_T[NDIM][DTYPE]} image_in;
21+
22+
layout(local_size_x_id = 0, local_size_y_id = 1, local_size_z_id = 2) in;
23+
24+
void main() {}

‎backends/vulkan/runtime/graph/ops/glsl/no_op.yaml

Lines changed: 26 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,26 @@
1+
# Copyright (c) Meta Platforms, Inc. and affiliates.
2+
# All rights reserved.
3+
#
4+
# This source code is licensed under the BSD-style license found in the
5+
# LICENSE file in the root directory of this source tree.
6+
7+
no_op:
8+
parameter_names_with_default_values:
9+
OPERATOR: X + A * Y
10+
NDIM: 3
11+
DTYPE: float
12+
generate_variant_forall:
13+
NDIM:
14+
- VALUE: 3
15+
SUFFIX: 3d
16+
- VALUE: 2
17+
SUFFIX: 2d
18+
DTYPE:
19+
- VALUE: half
20+
SUFFIX: half
21+
- VALUE: float
22+
SUFFIX: float
23+
- VALUE: int
24+
SUFFIX: int
25+
shader_variants:
26+
- NAME: no_op

‎backends/vulkan/runtime/graph/ops/utils/ShaderNameUtils.cpp

Lines changed: 13 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -26,6 +26,19 @@ void apply_dtype_suffix(std::stringstream& kernel_name, const vTensor& tensor) {
2626
}
2727
}
2828

29+
void apply_ndim_suffix(std::stringstream& kernel_name, const vTensor& tensor) {
30+
switch (tensor.storage_type()) {
31+
case api::StorageType::TEXTURE_3D:
32+
kernel_name << "_3d";
33+
break;
34+
case api::StorageType::TEXTURE_2D:
35+
kernel_name << "_2d";
36+
break;
37+
default:
38+
break;
39+
}
40+
}
41+
2942
void apply_memory_layout_suffix(
3043
std::stringstream& kernel_name,
3144
const vTensor& tensor) {

‎backends/vulkan/runtime/graph/ops/utils/ShaderNameUtils.h

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -16,6 +16,8 @@ namespace vkcompute {
1616

1717
void apply_dtype_suffix(std::stringstream& kernel_name, const vTensor& tensor);
1818

19+
void apply_ndim_suffix(std::stringstream& kernel_name, const vTensor& tensor);
20+
1921
void apply_memory_layout_suffix(
2022
std::stringstream& kernel_name,
2123
const vTensor& tensor);

0 commit comments

Comments
 (0)