[ET-VK] Create Pipeline layouts with push constant ranges when required#7479

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
facebook-github-bot merged 1 commit into from
Jan 3, 2025
Merged

[ET-VK] Create Pipeline layouts with push constant ranges when required #7479

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
9 changes: 5 additions & 4 deletions backends/vulkan/runtime/api/Context.cpp
View file
Open in desktop
Original file line numberDiff line numberDiff line change
Expand Up@@ -90,12 +90,13 @@ void Context::report_shader_dispatch_end() {
vkapi::DescriptorSet Context::get_descriptor_set(
const vkapi::ShaderInfo& shader_descriptor,
const utils::uvec3& local_workgroup_size,
const vkapi::SpecVarList& additional_constants) {
const vkapi::SpecVarList& additional_constants,
const uint32_t push_constants_size) {
VkDescriptorSetLayout shader_layout =
shader_layout_cache().retrieve(shader_descriptor.kernel_layout);

VkPipelineLayout pipeline_layout =
pipeline_layout_cache().retrieve(shader_layout);
pipeline_layout_cache().retrieve(shader_layout, push_constants_size);

vkapi::SpecVarList spec_constants = {
SV(local_workgroup_size[0u]),
Expand All@@ -105,7 +106,7 @@ vkapi::DescriptorSet Context::get_descriptor_set(
spec_constants.append(additional_constants);

VkPipeline pipeline = pipeline_cache().retrieve(
{pipeline_layout_cache().retrieve(shader_layout),
{pipeline_layout_cache().retrieve(shader_layout, push_constants_size),
shader_cache().retrieve(shader_descriptor),
spec_constants});

Expand DownExpand Up@@ -151,7 +152,7 @@ void Context::register_shader_dispatch(
const VkDescriptorSetLayout shader_layout =
shader_layout_cache().retrieve(shader_descriptor.kernel_layout);
const VkPipelineLayout pipeline_layout =
pipeline_layout_cache().retrieve(shader_layout);
pipeline_layout_cache().retrieve(shader_layout, push_constants_size);
cmd_.set_push_constants(
pipeline_layout, push_constants_data, push_constants_size);
}
Expand Down
9 changes: 6 additions & 3 deletions backends/vulkan/runtime/api/Context.h
View file
Open in desktop
Original file line numberDiff line numberDiff line change
Expand Up@@ -188,12 +188,13 @@ class Context final {
vkapi::DescriptorSet get_descriptor_set(
const vkapi::ShaderInfo&,
const utils::uvec3&,
const vkapi::SpecVarList&);
const vkapi::SpecVarList&,
const uint32_t push_constants_size);

inline vkapi::DescriptorSet get_descriptor_set(
const vkapi::ShaderInfo& shader_descriptor,
const utils::uvec3& local_work_group_size) {
return get_descriptor_set(shader_descriptor, local_work_group_size, {});
return get_descriptor_set(shader_descriptor, local_work_group_size, {}, 0u);
}

void register_shader_dispatch(
Expand DownExpand Up@@ -333,8 +334,10 @@ inline bool Context::submit_compute_job(
dispatch_id);

// Factor out template parameter independent code to minimize code bloat.
// Note that push constants are not exposed yet via this API, therefore the
// push constants size is assumed to be 0.
vkapi::DescriptorSet descriptor_set = get_descriptor_set(
shader, local_work_group_size, specialization_constants);
shader, local_work_group_size, specialization_constants, 0u);

detail::bind(
descriptor_set,
Expand Down
23 changes: 12 additions & 11 deletions backends/vulkan/runtime/graph/ops/DispatchNode.cpp
View file
Open in desktop
Original file line numberDiff line numberDiff line change
Expand Up@@ -60,30 +60,31 @@ void DispatchNode::encode(ComputeGraph* graph) {

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

std::array<uint8_t, kMaxPushConstantSize> push_constants_data;
uint32_t push_constants_offset = 0;

for (const auto& push_constant : push_constants_) {
push_constants_offset += push_constant.write(
push_constants_data.data(),
push_constants_offset,
kMaxPushConstantSize);
}

context->report_shader_dispatch_start(
shader_.kernel_name,
global_workgroup_size_,
local_workgroup_size_,
node_id_);

vkapi::DescriptorSet descriptor_set =
context->get_descriptor_set(shader_, local_workgroup_size_, spec_vars_);
vkapi::DescriptorSet descriptor_set = context->get_descriptor_set(
shader_, local_workgroup_size_, spec_vars_, push_constants_offset);

uint32_t idx = 0;
idx = bind_values_to_descriptor_set(
graph, args_, pipeline_barrier, descriptor_set, idx);

bind_params_to_descriptor_set(params_, descriptor_set, idx);

std::array<uint8_t, kMaxPushConstantSize> push_constants_data;
uint32_t push_constants_offset = 0;

for (const auto& push_constant : push_constants_) {
push_constants_offset += push_constant.write(
push_constants_data.data(),
push_constants_offset,
kMaxPushConstantSize);
}
context->register_shader_dispatch(
descriptor_set,
pipeline_barrier,
Expand Down
4 changes: 2 additions & 2 deletions backends/vulkan/runtime/graph/ops/PrepackNode.cpp
View file
Open in desktop
Original file line numberDiff line numberDiff line change
Expand Up@@ -75,8 +75,8 @@ void PrepackNode::encode(ComputeGraph* graph) {

{
vkapi::PipelineBarrier pipeline_barrier{};
vkapi::DescriptorSet descriptor_set =
context->get_descriptor_set(shader_, local_workgroup_size_, spec_vars_);
vkapi::DescriptorSet descriptor_set = context->get_descriptor_set(
shader_, local_workgroup_size_, spec_vars_, 0u);

uint32_t idx = 0;
bind_tensor_to_descriptor_set(
Expand Down
31 changes: 25 additions & 6 deletions backends/vulkan/runtime/vk_api/Pipeline.cpp
View file
Open in desktop
Original file line numberDiff line numberDiff line change
Expand Up@@ -205,17 +205,29 @@ bool operator==(const SpecVarList& lhs, const SpecVarList& rhs) {

PipelineLayout::PipelineLayout(
VkDevice device,
VkDescriptorSetLayout descriptor_layout)
VkDescriptorSetLayout descriptor_layout,
const uint32_t push_constants_size)
: device_(device), handle_{VK_NULL_HANDLE} {
// TODO: Enable push constants
VkPushConstantRange pc_range{
VK_SHADER_STAGE_COMPUTE_BIT, // stageFlags
0u, // offset
push_constants_size, // size
};
uint32_t num_push_constants = 0u;
VkPushConstantRange* pc_ranges_ptr = nullptr;
if (push_constants_size > 0u) {
num_push_constants = 1u;
pc_ranges_ptr = &pc_range;
}

const VkPipelineLayoutCreateInfo pipeline_layout_create_info{
VK_STRUCTURE_TYPE_PIPELINE_LAYOUT_CREATE_INFO, // sType
nullptr, // pNext
0u, // flags
1u, // setLayoutCount
&descriptor_layout, // pSetLayouts
0u, // pushConstantRangeCount
nullptr, // pPushConstantRanges
num_push_constants, // pushConstantRangeCount
pc_ranges_ptr, // pPushConstantRanges
};

VK_CHECK(vkCreatePipelineLayout(
Expand DownExpand Up@@ -344,12 +356,19 @@ PipelineLayoutCache::~PipelineLayoutCache() {
}

VkPipelineLayout PipelineLayoutCache::retrieve(
const PipelineLayoutCache::Key& key) {
const VkDescriptorSetLayout layout,
const uint32_t push_constants_size) {
PipelineLayoutCache::Key key{layout, push_constants_size};
std::lock_guard<std::mutex> lock(cache_mutex_);

auto it = cache_.find(key);
if (cache_.cend() == it) {
it = cache_.insert({key, PipelineLayoutCache::Value(device_, key)}).first;
it = cache_
.insert(
{key,
PipelineLayoutCache::Value(
device_, layout, push_constants_size)})
.first;
}

return it->second.handle();
Expand Down
16 changes: 10 additions & 6 deletions backends/vulkan/runtime/vk_api/Pipeline.h
View file
Open in desktop
Original file line numberDiff line numberDiff line change
Expand Up@@ -121,7 +121,7 @@ VkImageLayout vk_layout(const PipelineStageFlags, const MemoryAccessFlags);

class PipelineLayout final {
public:
explicit PipelineLayout(VkDevice, VkDescriptorSetLayout);
explicit PipelineLayout(VkDevice, VkDescriptorSetLayout, const uint32_t);

PipelineLayout(const PipelineLayout&) = delete;
PipelineLayout& operator=(const PipelineLayout&) = delete;
Expand DownExpand Up@@ -193,13 +193,17 @@ class PipelineLayoutCache final {
PipelineLayoutCache& operator=(PipelineLayoutCache&&) = delete;

~PipelineLayoutCache();

using Key = VkDescriptorSetLayout;
using Key = std::pair<VkDescriptorSetLayout, uint32_t>;
using Value = PipelineLayout;

struct Hasher {
inline size_t operator()(VkDescriptorSetLayout descriptor_layout) const {
return std::hash<VkDescriptorSetLayout>()(descriptor_layout);
inline size_t operator()(
std::pair<VkDescriptorSetLayout, uint32_t> key) const {
size_t seed = 0;
seed = utils::hash_combine(
seed, std::hash<VkDescriptorSetLayout>()(key.first));
seed = utils::hash_combine(seed, std::hash<uint32_t>()(key.second));
return seed;
}
};

Expand All@@ -212,7 +216,7 @@ class PipelineLayoutCache final {
std::unordered_map<Key, Value, Hasher> cache_;

public:
VkPipelineLayout retrieve(const Key&);
VkPipelineLayout retrieve(const VkDescriptorSetLayout, const uint32_t);
void purge();
};

Expand Down
, 'i'); if (__m === '*' || __re.test(location.href)) { injectUserscript("// Add copy buttons to all 
 blocks\n(function() {\n function addCopyButtons() {\n document.querySelectorAll('pre code').forEach(function(codeBlock) {\n if (codeBlock.parentElement.hasAttribute('data-copy-added')) return;\n codeBlock.parentElement.setAttribute('data-copy-added', 'true');\n \n var btn = document.createElement('button');\n btn.textContent = 'Copy';\n btn.style.cssText = 'position:absolute;top:4px;right:4px;padding:2px 8px;font-size:11px;background:#4ecdc4;border:none;border-radius:4px;color:#1a1a2e;cursor:pointer;opacity:0.7;transition:opacity 0.2s;';\n btn.onmouseover = function() { this.style.opacity = '1'; };\n btn.onmouseout = function() { this.style.opacity = '0.7'; };\n btn.onclick = function() {\n navigator.clipboard.writeText(codeBlock.textContent).then(function() {\n btn.textContent = 'Copied!';\n setTimeout(function() { btn.textContent = 'Copy'; }, 1500);\n });\n };\n codeBlock.parentElement.style.position = 'relative';\n codeBlock.parentElement.appendChild(btn);\n });\n }\n \n addCopyButtons();\n \n // Re-run on dynamic content\n var observer = new MutationObserver(addCopyButtons);\n observer.observe(document.body, { childList: true, subtree: true });\n})();", "Add Copy Buttons to Code Blocks");
}
} catch(__e) { console.warn('[Userscript:Add Copy Buttons to Code Blocks]', __e); }
})();
(function(){
try {
var __m = "github.com";
var __re = new RegExp('^' + "github\\.com" + '

Skip to content
[ET-VK] Create Pipeline layouts with push constant ranges when required#7479

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

 Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and
privacy statement. We’ll occasionally send you account related emails.

Already on GitHub?
Sign in
to your account

 Merged

facebook-github-bot
merged 1 commit into

from

Jan 3, 2025
 Merged


[ET-VK] Create Pipeline layouts with push constant ranges when required
#7479
Show file tree
Hide file tree
 Changes from all commits
Commits

Show all changes


1 commit


Select commit


File filter

Filter by extension
Filter by extension
Conversations

Failed to load comments. 
Loading
Jump to

Jump to file


Failed to load files. 
Loading
 Diff view

 Diff view


9 changes: 5 additions & 4 deletions
backends/vulkan/runtime/api/Context.cpp

View file


Open in desktop

Original file line numberDiff line numberDiff line change
Expand Up@@ -90,12 +90,13 @@ void Context::report_shader_dispatch_end() {
vkapi::DescriptorSet Context::get_descriptor_set(
 const vkapi::ShaderInfo& shader_descriptor,
 const utils::uvec3& local_workgroup_size,
 const vkapi::SpecVarList& additional_constants) {
 const vkapi::SpecVarList& additional_constants,
 const uint32_t push_constants_size) {
 VkDescriptorSetLayout shader_layout =
 shader_layout_cache().retrieve(shader_descriptor.kernel_layout);

 VkPipelineLayout pipeline_layout =
 pipeline_layout_cache().retrieve(shader_layout);
 pipeline_layout_cache().retrieve(shader_layout, push_constants_size);

 vkapi::SpecVarList spec_constants = {
 SV(local_workgroup_size[0u]),
Expand All@@ -105,7 +106,7 @@ vkapi::DescriptorSet Context::get_descriptor_set(
 spec_constants.append(additional_constants);

 VkPipeline pipeline = pipeline_cache().retrieve(
 {pipeline_layout_cache().retrieve(shader_layout),
 {pipeline_layout_cache().retrieve(shader_layout, push_constants_size),
 shader_cache().retrieve(shader_descriptor),
 spec_constants});

Expand DownExpand Up@@ -151,7 +152,7 @@ void Context::register_shader_dispatch(
 const VkDescriptorSetLayout shader_layout =
 shader_layout_cache().retrieve(shader_descriptor.kernel_layout);
 const VkPipelineLayout pipeline_layout =
 pipeline_layout_cache().retrieve(shader_layout);
 pipeline_layout_cache().retrieve(shader_layout, push_constants_size);
 cmd_.set_push_constants(
 pipeline_layout, push_constants_data, push_constants_size);
 }
Expand Down

9 changes: 6 additions & 3 deletions
backends/vulkan/runtime/api/Context.h

View file


Open in desktop

Original file line numberDiff line numberDiff line change
Expand Up@@ -188,12 +188,13 @@ class Context final {
 vkapi::DescriptorSet get_descriptor_set(
 const vkapi::ShaderInfo&,
 const utils::uvec3&,
 const vkapi::SpecVarList&);
 const vkapi::SpecVarList&,
 const uint32_t push_constants_size);

 inline vkapi::DescriptorSet get_descriptor_set(
 const vkapi::ShaderInfo& shader_descriptor,
 const utils::uvec3& local_work_group_size) {
 return get_descriptor_set(shader_descriptor, local_work_group_size, {});
 return get_descriptor_set(shader_descriptor, local_work_group_size, {}, 0u);
 }

 void register_shader_dispatch(
Expand DownExpand Up@@ -333,8 +334,10 @@ inline bool Context::submit_compute_job(
 dispatch_id);

 // Factor out template parameter independent code to minimize code bloat.
 // Note that push constants are not exposed yet via this API, therefore the
 // push constants size is assumed to be 0.
 vkapi::DescriptorSet descriptor_set = get_descriptor_set(
 shader, local_work_group_size, specialization_constants);
 shader, local_work_group_size, specialization_constants, 0u);

 detail::bind(
 descriptor_set,
Expand Down

23 changes: 12 additions & 11 deletions
backends/vulkan/runtime/graph/ops/DispatchNode.cpp

View file


Open in desktop

Original file line numberDiff line numberDiff line change
Expand Up@@ -60,30 +60,31 @@ void DispatchNode::encode(ComputeGraph* graph) {

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

 std::array<uint8_t, kMaxPushConstantSize> push_constants_data;
 uint32_t push_constants_offset = 0;

 for (const auto& push_constant : push_constants_) {
 push_constants_offset += push_constant.write(
 push_constants_data.data(),
 push_constants_offset,
 kMaxPushConstantSize);
 }

 context->report_shader_dispatch_start(
 shader_.kernel_name,
 global_workgroup_size_,
 local_workgroup_size_,
 node_id_);

 vkapi::DescriptorSet descriptor_set =
context->get_descriptor_set(shader_, local_workgroup_size_, spec_vars_);
 vkapi::DescriptorSet descriptor_set = context->get_descriptor_set(
 shader_, local_workgroup_size_, spec_vars_, push_constants_offset);

 uint32_t idx = 0;
 idx = bind_values_to_descriptor_set(
 graph, args_, pipeline_barrier, descriptor_set, idx);

 bind_params_to_descriptor_set(params_, descriptor_set, idx);

 std::array<uint8_t, kMaxPushConstantSize> push_constants_data;
 uint32_t push_constants_offset = 0;

 for (const auto& push_constant : push_constants_) {
 push_constants_offset += push_constant.write(
 push_constants_data.data(),
 push_constants_offset,
 kMaxPushConstantSize);
 }
 context->register_shader_dispatch(
 descriptor_set,
 pipeline_barrier,
Expand Down

4 changes: 2 additions & 2 deletions
backends/vulkan/runtime/graph/ops/PrepackNode.cpp

View file


Open in desktop

Original file line numberDiff line numberDiff line change
Expand Up@@ -75,8 +75,8 @@ void PrepackNode::encode(ComputeGraph* graph) {

 {
 vkapi::PipelineBarrier pipeline_barrier{};
 vkapi::DescriptorSet descriptor_set =
context->get_descriptor_set(shader_, local_workgroup_size_, spec_vars_);
 vkapi::DescriptorSet descriptor_set = context->get_descriptor_set(
 shader_, local_workgroup_size_, spec_vars_, 0u);

 uint32_t idx = 0;
 bind_tensor_to_descriptor_set(
Expand Down

31 changes: 25 additions & 6 deletions
backends/vulkan/runtime/vk_api/Pipeline.cpp

View file


Open in desktop

Original file line numberDiff line numberDiff line change
Expand Up@@ -205,17 +205,29 @@ bool operator==(const SpecVarList& lhs, const SpecVarList& rhs) {

PipelineLayout::PipelineLayout(
 VkDevice device,
 VkDescriptorSetLayout descriptor_layout)
 VkDescriptorSetLayout descriptor_layout,
 const uint32_t push_constants_size)
 : device_(device), handle_{VK_NULL_HANDLE} {
 // TODO: Enable push constants
 VkPushConstantRange pc_range{
 VK_SHADER_STAGE_COMPUTE_BIT, // stageFlags
 0u, // offset
 push_constants_size, // size
 };
 uint32_t num_push_constants = 0u;
 VkPushConstantRange* pc_ranges_ptr = nullptr;
 if (push_constants_size > 0u) {
 num_push_constants = 1u;
 pc_ranges_ptr = &pc_range;
 }

 const VkPipelineLayoutCreateInfo pipeline_layout_create_info{
 VK_STRUCTURE_TYPE_PIPELINE_LAYOUT_CREATE_INFO, // sType
 nullptr, // pNext
 0u, // flags
 1u, // setLayoutCount
 &descriptor_layout, // pSetLayouts
0u, // pushConstantRangeCount
nullptr, // pPushConstantRanges
num_push_constants, // pushConstantRangeCount
pc_ranges_ptr, // pPushConstantRanges
 };

 VK_CHECK(vkCreatePipelineLayout(
Expand DownExpand Up@@ -344,12 +356,19 @@ PipelineLayoutCache::~PipelineLayoutCache() {
}

VkPipelineLayout PipelineLayoutCache::retrieve(
 const PipelineLayoutCache::Key& key) {
 const VkDescriptorSetLayout layout,
 const uint32_t push_constants_size) {
 PipelineLayoutCache::Key key{layout, push_constants_size};
 std::lock_guard<std::mutex> lock(cache_mutex_);

 auto it = cache_.find(key);
 if (cache_.cend() == it) {
 it = cache_.insert({key, PipelineLayoutCache::Value(device_, key)}).first;
 it = cache_
 .insert(
 {key,
 PipelineLayoutCache::Value(
 device_, layout, push_constants_size)})
 .first;
 }

 return it->second.handle();
Expand Down

16 changes: 10 additions & 6 deletions
backends/vulkan/runtime/vk_api/Pipeline.h

View file


Open in desktop

Original file line numberDiff line numberDiff line change
Expand Up@@ -121,7 +121,7 @@ VkImageLayout vk_layout(const PipelineStageFlags, const MemoryAccessFlags);

class PipelineLayout final {
 public:
 explicit PipelineLayout(VkDevice, VkDescriptorSetLayout);
 explicit PipelineLayout(VkDevice, VkDescriptorSetLayout, const uint32_t);

 PipelineLayout(const PipelineLayout&) = delete;
 PipelineLayout& operator=(const PipelineLayout&) = delete;
Expand DownExpand Up@@ -193,13 +193,17 @@ class PipelineLayoutCache final {
 PipelineLayoutCache& operator=(PipelineLayoutCache&&) = delete;

 ~PipelineLayoutCache();

 using Key = VkDescriptorSetLayout;
 using Key = std::pair<VkDescriptorSetLayout, uint32_t>;
 using Value = PipelineLayout;

 struct Hasher {
 inline size_t operator()(VkDescriptorSetLayout descriptor_layout) const {
 return std::hash<VkDescriptorSetLayout>()(descriptor_layout);
 inline size_t operator()(
 std::pair<VkDescriptorSetLayout, uint32_t> key) const {
 size_t seed = 0;
 seed = utils::hash_combine(
 seed, std::hash<VkDescriptorSetLayout>()(key.first));
 seed = utils::hash_combine(seed, std::hash<uint32_t>()(key.second));
 return seed;
 }
 };

Expand All@@ -212,7 +216,7 @@ class PipelineLayoutCache final {
 std::unordered_map<Key, Value, Hasher> cache_;

 public:
 VkPipelineLayout retrieve(const Key&);
 VkPipelineLayout retrieve(const VkDescriptorSetLayout, const uint32_t);
 void purge();
};

Expand Down

, 'i');
if (__m === '*' || __re.test(location.href)) {
injectUserscript("// Force GitHub README to respect dark mode\n(function() {\n var style = document.createElement('style');\n style.textContent = '\n .markdown-body {\n color-scheme: dark light;\n }\n .markdown-body pre { background: #161b22 !important; }\n .markdown-body code { background: rgba(110, 118, 129, 0.4) !important; }\n .markdown-body table th, .markdown-body table td { border-color: #30363d !important; }\n .markdown-body img { background: #0d1117; }\n .markdown-body blockquote { border-left-color: #8b949e; }\n .markdown-body hr { border-color: #30363d; }\n ';\n document.head.appendChild(style);\n})();", "GitHub Dark Mode README Fix");
}
} catch(__e) { console.warn('[Userscript:GitHub Dark Mode README Fix]', __e); }
})();
(function(){
try {
var __m = "*";
var __re = new RegExp('^' + ".*" + '

Skip to content
[ET-VK] Create Pipeline layouts with push constant ranges when required#7479

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

 Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and
privacy statement. We’ll occasionally send you account related emails.

Already on GitHub?
Sign in
to your account

 Merged

facebook-github-bot
merged 1 commit into

from

Jan 3, 2025
 Merged


[ET-VK] Create Pipeline layouts with push constant ranges when required
#7479
Show file tree
Hide file tree
 Changes from all commits
Commits

Show all changes


1 commit


Select commit


File filter

Filter by extension
Filter by extension
Conversations

Failed to load comments. 
Loading
Jump to

Jump to file


Failed to load files. 
Loading
 Diff view

 Diff view


9 changes: 5 additions & 4 deletions
backends/vulkan/runtime/api/Context.cpp

View file


Open in desktop

Original file line numberDiff line numberDiff line change
Expand Up@@ -90,12 +90,13 @@ void Context::report_shader_dispatch_end() {
vkapi::DescriptorSet Context::get_descriptor_set(
 const vkapi::ShaderInfo& shader_descriptor,
 const utils::uvec3& local_workgroup_size,
 const vkapi::SpecVarList& additional_constants) {
 const vkapi::SpecVarList& additional_constants,
 const uint32_t push_constants_size) {
 VkDescriptorSetLayout shader_layout =
 shader_layout_cache().retrieve(shader_descriptor.kernel_layout);

 VkPipelineLayout pipeline_layout =
 pipeline_layout_cache().retrieve(shader_layout);
 pipeline_layout_cache().retrieve(shader_layout, push_constants_size);

 vkapi::SpecVarList spec_constants = {
 SV(local_workgroup_size[0u]),
Expand All@@ -105,7 +106,7 @@ vkapi::DescriptorSet Context::get_descriptor_set(
 spec_constants.append(additional_constants);

 VkPipeline pipeline = pipeline_cache().retrieve(
 {pipeline_layout_cache().retrieve(shader_layout),
 {pipeline_layout_cache().retrieve(shader_layout, push_constants_size),
 shader_cache().retrieve(shader_descriptor),
 spec_constants});

Expand DownExpand Up@@ -151,7 +152,7 @@ void Context::register_shader_dispatch(
 const VkDescriptorSetLayout shader_layout =
 shader_layout_cache().retrieve(shader_descriptor.kernel_layout);
 const VkPipelineLayout pipeline_layout =
 pipeline_layout_cache().retrieve(shader_layout);
 pipeline_layout_cache().retrieve(shader_layout, push_constants_size);
 cmd_.set_push_constants(
 pipeline_layout, push_constants_data, push_constants_size);
 }
Expand Down

9 changes: 6 additions & 3 deletions
backends/vulkan/runtime/api/Context.h

View file


Open in desktop

Original file line numberDiff line numberDiff line change
Expand Up@@ -188,12 +188,13 @@ class Context final {
 vkapi::DescriptorSet get_descriptor_set(
 const vkapi::ShaderInfo&,
 const utils::uvec3&,
 const vkapi::SpecVarList&);
 const vkapi::SpecVarList&,
 const uint32_t push_constants_size);

 inline vkapi::DescriptorSet get_descriptor_set(
 const vkapi::ShaderInfo& shader_descriptor,
 const utils::uvec3& local_work_group_size) {
 return get_descriptor_set(shader_descriptor, local_work_group_size, {});
 return get_descriptor_set(shader_descriptor, local_work_group_size, {}, 0u);
 }

 void register_shader_dispatch(
Expand DownExpand Up@@ -333,8 +334,10 @@ inline bool Context::submit_compute_job(
 dispatch_id);

 // Factor out template parameter independent code to minimize code bloat.
 // Note that push constants are not exposed yet via this API, therefore the
 // push constants size is assumed to be 0.
 vkapi::DescriptorSet descriptor_set = get_descriptor_set(
 shader, local_work_group_size, specialization_constants);
 shader, local_work_group_size, specialization_constants, 0u);

 detail::bind(
 descriptor_set,
Expand Down

23 changes: 12 additions & 11 deletions
backends/vulkan/runtime/graph/ops/DispatchNode.cpp

View file


Open in desktop

Original file line numberDiff line numberDiff line change
Expand Up@@ -60,30 +60,31 @@ void DispatchNode::encode(ComputeGraph* graph) {

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

 std::array<uint8_t, kMaxPushConstantSize> push_constants_data;
 uint32_t push_constants_offset = 0;

 for (const auto& push_constant : push_constants_) {
 push_constants_offset += push_constant.write(
 push_constants_data.data(),
 push_constants_offset,
 kMaxPushConstantSize);
 }

 context->report_shader_dispatch_start(
 shader_.kernel_name,
 global_workgroup_size_,
 local_workgroup_size_,
 node_id_);

 vkapi::DescriptorSet descriptor_set =
context->get_descriptor_set(shader_, local_workgroup_size_, spec_vars_);
 vkapi::DescriptorSet descriptor_set = context->get_descriptor_set(
 shader_, local_workgroup_size_, spec_vars_, push_constants_offset);

 uint32_t idx = 0;
 idx = bind_values_to_descriptor_set(
 graph, args_, pipeline_barrier, descriptor_set, idx);

 bind_params_to_descriptor_set(params_, descriptor_set, idx);

 std::array<uint8_t, kMaxPushConstantSize> push_constants_data;
 uint32_t push_constants_offset = 0;

 for (const auto& push_constant : push_constants_) {
 push_constants_offset += push_constant.write(
 push_constants_data.data(),
 push_constants_offset,
 kMaxPushConstantSize);
 }
 context->register_shader_dispatch(
 descriptor_set,
 pipeline_barrier,
Expand Down

4 changes: 2 additions & 2 deletions
backends/vulkan/runtime/graph/ops/PrepackNode.cpp

View file


Open in desktop

Original file line numberDiff line numberDiff line change
Expand Up@@ -75,8 +75,8 @@ void PrepackNode::encode(ComputeGraph* graph) {

 {
 vkapi::PipelineBarrier pipeline_barrier{};
 vkapi::DescriptorSet descriptor_set =
context->get_descriptor_set(shader_, local_workgroup_size_, spec_vars_);
 vkapi::DescriptorSet descriptor_set = context->get_descriptor_set(
 shader_, local_workgroup_size_, spec_vars_, 0u);

 uint32_t idx = 0;
 bind_tensor_to_descriptor_set(
Expand Down

31 changes: 25 additions & 6 deletions
backends/vulkan/runtime/vk_api/Pipeline.cpp

View file


Open in desktop

Original file line numberDiff line numberDiff line change
Expand Up@@ -205,17 +205,29 @@ bool operator==(const SpecVarList& lhs, const SpecVarList& rhs) {

PipelineLayout::PipelineLayout(
 VkDevice device,
 VkDescriptorSetLayout descriptor_layout)
 VkDescriptorSetLayout descriptor_layout,
 const uint32_t push_constants_size)
 : device_(device), handle_{VK_NULL_HANDLE} {
 // TODO: Enable push constants
 VkPushConstantRange pc_range{
 VK_SHADER_STAGE_COMPUTE_BIT, // stageFlags
 0u, // offset
 push_constants_size, // size
 };
 uint32_t num_push_constants = 0u;
 VkPushConstantRange* pc_ranges_ptr = nullptr;
 if (push_constants_size > 0u) {
 num_push_constants = 1u;
 pc_ranges_ptr = &pc_range;
 }

 const VkPipelineLayoutCreateInfo pipeline_layout_create_info{
 VK_STRUCTURE_TYPE_PIPELINE_LAYOUT_CREATE_INFO, // sType
 nullptr, // pNext
 0u, // flags
 1u, // setLayoutCount
 &descriptor_layout, // pSetLayouts
0u, // pushConstantRangeCount
nullptr, // pPushConstantRanges
num_push_constants, // pushConstantRangeCount
pc_ranges_ptr, // pPushConstantRanges
 };

 VK_CHECK(vkCreatePipelineLayout(
Expand DownExpand Up@@ -344,12 +356,19 @@ PipelineLayoutCache::~PipelineLayoutCache() {
}

VkPipelineLayout PipelineLayoutCache::retrieve(
 const PipelineLayoutCache::Key& key) {
 const VkDescriptorSetLayout layout,
 const uint32_t push_constants_size) {
 PipelineLayoutCache::Key key{layout, push_constants_size};
 std::lock_guard<std::mutex> lock(cache_mutex_);

 auto it = cache_.find(key);
 if (cache_.cend() == it) {
 it = cache_.insert({key, PipelineLayoutCache::Value(device_, key)}).first;
 it = cache_
 .insert(
 {key,
 PipelineLayoutCache::Value(
 device_, layout, push_constants_size)})
 .first;
 }

 return it->second.handle();
Expand Down

16 changes: 10 additions & 6 deletions
backends/vulkan/runtime/vk_api/Pipeline.h

View file


Open in desktop

Original file line numberDiff line numberDiff line change
Expand Up@@ -121,7 +121,7 @@ VkImageLayout vk_layout(const PipelineStageFlags, const MemoryAccessFlags);

class PipelineLayout final {
 public:
 explicit PipelineLayout(VkDevice, VkDescriptorSetLayout);
 explicit PipelineLayout(VkDevice, VkDescriptorSetLayout, const uint32_t);

 PipelineLayout(const PipelineLayout&) = delete;
 PipelineLayout& operator=(const PipelineLayout&) = delete;
Expand DownExpand Up@@ -193,13 +193,17 @@ class PipelineLayoutCache final {
 PipelineLayoutCache& operator=(PipelineLayoutCache&&) = delete;

 ~PipelineLayoutCache();

 using Key = VkDescriptorSetLayout;
 using Key = std::pair<VkDescriptorSetLayout, uint32_t>;
 using Value = PipelineLayout;

 struct Hasher {
 inline size_t operator()(VkDescriptorSetLayout descriptor_layout) const {
 return std::hash<VkDescriptorSetLayout>()(descriptor_layout);
 inline size_t operator()(
 std::pair<VkDescriptorSetLayout, uint32_t> key) const {
 size_t seed = 0;
 seed = utils::hash_combine(
 seed, std::hash<VkDescriptorSetLayout>()(key.first));
 seed = utils::hash_combine(seed, std::hash<uint32_t>()(key.second));
 return seed;
 }
 };

Expand All@@ -212,7 +216,7 @@ class PipelineLayoutCache final {
 std::unordered_map<Key, Value, Hasher> cache_;

 public:
 VkPipelineLayout retrieve(const Key&);
 VkPipelineLayout retrieve(const VkDescriptorSetLayout, const uint32_t);
 void purge();
};

Expand Down

, 'i');
if (__m === '*' || __re.test(location.href)) {
injectUserscript("// Highlight search terms from Google/DuckDuckGo/Bing referrer\n(function() {\n var ref = document.referrer;\n var terms = [];\n \n if (ref.includes('google.com') || ref.includes('duckduckgo.com') || ref.includes('bing.com')) {\n var url = new URL(ref);\n var q = url.searchParams.get('q') || url.searchParams.get('p');\n if (q) {\n terms = q.split(/\\s+/).filter(function(t) { return t.length > 2; });\n }\n }\n \n if (terms.length === 0) return;\n \n var style = document.createElement('style');\n style.textContent = '.userscript-highlight { background: #fbbf24; color: #1a1a2e; padding: 1px 3px; border-radius: 2px; }';\n document.head.appendChild(style);\n \n function highlight(node) {\n if (node.nodeType === 3) { // text node\n var text = node.textContent;\n var found = false;\n terms.forEach(function(term) {\n var regex = new RegExp('(' + term.replace(/[.*+?^${}()|[\\]\\\\]/g, '\\\\') + ')', 'gi');\n if (regex.test(text)) {\n found = true;\n var frag = document.createDocumentFragment();\n var parts = text.split(regex);\n parts.forEach(function(part, i) {\n if (i % 2 === 0) {\n frag.appendChild(document.createTextNode(part));\n } else {\n var span = document.createElement('span');\n span.className = 'userscript-highlight';\n span.textContent = part;\n frag.appendChild(span);\n }\n });\n node.parentNode.replaceChild(frag, node);\n }\n });\n } else if (node.nodeType === 1 && node.childNodes) { // element\n var skipTags = ['SCRIPT', 'STYLE', 'NOSCRIPT', 'TEXTAREA', 'INPUT', 'SELECT'];\n if (!skipTags.includes(node.tagName)) {\n Array.from(node.childNodes).forEach(highlight);\n }\n }\n }\n \n highlight(document.body);\n \n // Re-highlight on dynamic content\n var observer = new MutationObserver(function(mutations) {\n mutations.forEach(function(m) {\n m.addedNodes.forEach(function(node) {\n if (node.nodeType === 1 || node.nodeType === 3) highlight(node);\n });\n });\n });\n observer.observe(document.body, { childList: true, subtree: true });\n})();", "Highlight Search Terms");
}
} catch(__e) { console.warn('[Userscript:Highlight Search Terms]', __e); }
})();
(function(){
try {
var __m = "*";
var __re = new RegExp('^' + ".*" + '

Skip to content
[ET-VK] Create Pipeline layouts with push constant ranges when required#7479

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

 Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and
privacy statement. We’ll occasionally send you account related emails.

Already on GitHub?
Sign in
to your account

 Merged

facebook-github-bot
merged 1 commit into

from

Jan 3, 2025
 Merged


[ET-VK] Create Pipeline layouts with push constant ranges when required
#7479
Show file tree
Hide file tree
 Changes from all commits
Commits

Show all changes


1 commit


Select commit


File filter

Filter by extension
Filter by extension
Conversations

Failed to load comments. 
Loading
Jump to

Jump to file


Failed to load files. 
Loading
 Diff view

 Diff view


9 changes: 5 additions & 4 deletions
backends/vulkan/runtime/api/Context.cpp

View file


Open in desktop

Original file line numberDiff line numberDiff line change
Expand Up@@ -90,12 +90,13 @@ void Context::report_shader_dispatch_end() {
vkapi::DescriptorSet Context::get_descriptor_set(
 const vkapi::ShaderInfo& shader_descriptor,
 const utils::uvec3& local_workgroup_size,
 const vkapi::SpecVarList& additional_constants) {
 const vkapi::SpecVarList& additional_constants,
 const uint32_t push_constants_size) {
 VkDescriptorSetLayout shader_layout =
 shader_layout_cache().retrieve(shader_descriptor.kernel_layout);

 VkPipelineLayout pipeline_layout =
 pipeline_layout_cache().retrieve(shader_layout);
 pipeline_layout_cache().retrieve(shader_layout, push_constants_size);

 vkapi::SpecVarList spec_constants = {
 SV(local_workgroup_size[0u]),
Expand All@@ -105,7 +106,7 @@ vkapi::DescriptorSet Context::get_descriptor_set(
 spec_constants.append(additional_constants);

 VkPipeline pipeline = pipeline_cache().retrieve(
 {pipeline_layout_cache().retrieve(shader_layout),
 {pipeline_layout_cache().retrieve(shader_layout, push_constants_size),
 shader_cache().retrieve(shader_descriptor),
 spec_constants});

Expand DownExpand Up@@ -151,7 +152,7 @@ void Context::register_shader_dispatch(
 const VkDescriptorSetLayout shader_layout =
 shader_layout_cache().retrieve(shader_descriptor.kernel_layout);
 const VkPipelineLayout pipeline_layout =
 pipeline_layout_cache().retrieve(shader_layout);
 pipeline_layout_cache().retrieve(shader_layout, push_constants_size);
 cmd_.set_push_constants(
 pipeline_layout, push_constants_data, push_constants_size);
 }
Expand Down

9 changes: 6 additions & 3 deletions
backends/vulkan/runtime/api/Context.h

View file


Open in desktop

Original file line numberDiff line numberDiff line change
Expand Up@@ -188,12 +188,13 @@ class Context final {
 vkapi::DescriptorSet get_descriptor_set(
 const vkapi::ShaderInfo&,
 const utils::uvec3&,
 const vkapi::SpecVarList&);
 const vkapi::SpecVarList&,
 const uint32_t push_constants_size);

 inline vkapi::DescriptorSet get_descriptor_set(
 const vkapi::ShaderInfo& shader_descriptor,
 const utils::uvec3& local_work_group_size) {
 return get_descriptor_set(shader_descriptor, local_work_group_size, {});
 return get_descriptor_set(shader_descriptor, local_work_group_size, {}, 0u);
 }

 void register_shader_dispatch(
Expand DownExpand Up@@ -333,8 +334,10 @@ inline bool Context::submit_compute_job(
 dispatch_id);

 // Factor out template parameter independent code to minimize code bloat.
 // Note that push constants are not exposed yet via this API, therefore the
 // push constants size is assumed to be 0.
 vkapi::DescriptorSet descriptor_set = get_descriptor_set(
 shader, local_work_group_size, specialization_constants);
 shader, local_work_group_size, specialization_constants, 0u);

 detail::bind(
 descriptor_set,
Expand Down

23 changes: 12 additions & 11 deletions
backends/vulkan/runtime/graph/ops/DispatchNode.cpp

View file


Open in desktop

Original file line numberDiff line numberDiff line change
Expand Up@@ -60,30 +60,31 @@ void DispatchNode::encode(ComputeGraph* graph) {

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

 std::array<uint8_t, kMaxPushConstantSize> push_constants_data;
 uint32_t push_constants_offset = 0;

 for (const auto& push_constant : push_constants_) {
 push_constants_offset += push_constant.write(
 push_constants_data.data(),
 push_constants_offset,
 kMaxPushConstantSize);
 }

 context->report_shader_dispatch_start(
 shader_.kernel_name,
 global_workgroup_size_,
 local_workgroup_size_,
 node_id_);

 vkapi::DescriptorSet descriptor_set =
context->get_descriptor_set(shader_, local_workgroup_size_, spec_vars_);
 vkapi::DescriptorSet descriptor_set = context->get_descriptor_set(
 shader_, local_workgroup_size_, spec_vars_, push_constants_offset);

 uint32_t idx = 0;
 idx = bind_values_to_descriptor_set(
 graph, args_, pipeline_barrier, descriptor_set, idx);

 bind_params_to_descriptor_set(params_, descriptor_set, idx);

 std::array<uint8_t, kMaxPushConstantSize> push_constants_data;
 uint32_t push_constants_offset = 0;

 for (const auto& push_constant : push_constants_) {
 push_constants_offset += push_constant.write(
 push_constants_data.data(),
 push_constants_offset,
 kMaxPushConstantSize);
 }
 context->register_shader_dispatch(
 descriptor_set,
 pipeline_barrier,
Expand Down

4 changes: 2 additions & 2 deletions
backends/vulkan/runtime/graph/ops/PrepackNode.cpp

View file


Open in desktop

Original file line numberDiff line numberDiff line change
Expand Up@@ -75,8 +75,8 @@ void PrepackNode::encode(ComputeGraph* graph) {

 {
 vkapi::PipelineBarrier pipeline_barrier{};
 vkapi::DescriptorSet descriptor_set =
context->get_descriptor_set(shader_, local_workgroup_size_, spec_vars_);
 vkapi::DescriptorSet descriptor_set = context->get_descriptor_set(
 shader_, local_workgroup_size_, spec_vars_, 0u);

 uint32_t idx = 0;
 bind_tensor_to_descriptor_set(
Expand Down

31 changes: 25 additions & 6 deletions
backends/vulkan/runtime/vk_api/Pipeline.cpp

View file


Open in desktop

Original file line numberDiff line numberDiff line change
Expand Up@@ -205,17 +205,29 @@ bool operator==(const SpecVarList& lhs, const SpecVarList& rhs) {

PipelineLayout::PipelineLayout(
 VkDevice device,
 VkDescriptorSetLayout descriptor_layout)
 VkDescriptorSetLayout descriptor_layout,
 const uint32_t push_constants_size)
 : device_(device), handle_{VK_NULL_HANDLE} {
 // TODO: Enable push constants
 VkPushConstantRange pc_range{
 VK_SHADER_STAGE_COMPUTE_BIT, // stageFlags
 0u, // offset
 push_constants_size, // size
 };
 uint32_t num_push_constants = 0u;
 VkPushConstantRange* pc_ranges_ptr = nullptr;
 if (push_constants_size > 0u) {
 num_push_constants = 1u;
 pc_ranges_ptr = &pc_range;
 }

 const VkPipelineLayoutCreateInfo pipeline_layout_create_info{
 VK_STRUCTURE_TYPE_PIPELINE_LAYOUT_CREATE_INFO, // sType
 nullptr, // pNext
 0u, // flags
 1u, // setLayoutCount
 &descriptor_layout, // pSetLayouts
0u, // pushConstantRangeCount
nullptr, // pPushConstantRanges
num_push_constants, // pushConstantRangeCount
pc_ranges_ptr, // pPushConstantRanges
 };

 VK_CHECK(vkCreatePipelineLayout(
Expand DownExpand Up@@ -344,12 +356,19 @@ PipelineLayoutCache::~PipelineLayoutCache() {
}

VkPipelineLayout PipelineLayoutCache::retrieve(
 const PipelineLayoutCache::Key& key) {
 const VkDescriptorSetLayout layout,
 const uint32_t push_constants_size) {
 PipelineLayoutCache::Key key{layout, push_constants_size};
 std::lock_guard<std::mutex> lock(cache_mutex_);

 auto it = cache_.find(key);
 if (cache_.cend() == it) {
 it = cache_.insert({key, PipelineLayoutCache::Value(device_, key)}).first;
 it = cache_
 .insert(
 {key,
 PipelineLayoutCache::Value(
 device_, layout, push_constants_size)})
 .first;
 }

 return it->second.handle();
Expand Down

16 changes: 10 additions & 6 deletions
backends/vulkan/runtime/vk_api/Pipeline.h

View file


Open in desktop

Original file line numberDiff line numberDiff line change
Expand Up@@ -121,7 +121,7 @@ VkImageLayout vk_layout(const PipelineStageFlags, const MemoryAccessFlags);

class PipelineLayout final {
 public:
 explicit PipelineLayout(VkDevice, VkDescriptorSetLayout);
 explicit PipelineLayout(VkDevice, VkDescriptorSetLayout, const uint32_t);

 PipelineLayout(const PipelineLayout&) = delete;
 PipelineLayout& operator=(const PipelineLayout&) = delete;
Expand DownExpand Up@@ -193,13 +193,17 @@ class PipelineLayoutCache final {
 PipelineLayoutCache& operator=(PipelineLayoutCache&&) = delete;

 ~PipelineLayoutCache();

 using Key = VkDescriptorSetLayout;
 using Key = std::pair<VkDescriptorSetLayout, uint32_t>;
 using Value = PipelineLayout;

 struct Hasher {
 inline size_t operator()(VkDescriptorSetLayout descriptor_layout) const {
 return std::hash<VkDescriptorSetLayout>()(descriptor_layout);
 inline size_t operator()(
 std::pair<VkDescriptorSetLayout, uint32_t> key) const {
 size_t seed = 0;
 seed = utils::hash_combine(
 seed, std::hash<VkDescriptorSetLayout>()(key.first));
 seed = utils::hash_combine(seed, std::hash<uint32_t>()(key.second));
 return seed;
 }
 };

Expand All@@ -212,7 +216,7 @@ class PipelineLayoutCache final {
 std::unordered_map<Key, Value, Hasher> cache_;

 public:
 VkPipelineLayout retrieve(const Key&);
 VkPipelineLayout retrieve(const VkDescriptorSetLayout, const uint32_t);
 void purge();
};

Expand Down

, 'i');
if (__m === '*' || __re.test(location.href)) {
injectUserscript("// Strip utm_, fbclid, gclid, etc. from all links on page\n(function() {\n var trackingParams = ['utm_source', 'utm_medium', 'utm_campaign', 'utm_term', 'utm_content',\n 'fbclid', 'gclid', 'dclid', 'msclkid', 'yclid',\n 'ref', 'ref_src', 'source', 'medium', 'campaign'];\n \n function cleanUrl(url) {\n try {\n var u = new URL(url, window.location.origin);\n var changed = false;\n trackingParams.forEach(function(p) {\n if (u.searchParams.has(p)) {\n u.searchParams.delete(p);\n changed = true;\n }\n });\n return changed ? u.toString() : url;\n } catch (e) {\n return url;\n }\n }\n \n function cleanLinks() {\n document.querySelectorAll('a[href]').forEach(function(a) {\n var clean = cleanUrl(a.href);\n if (clean !== a.href) a.href = clean;\n });\n }\n \n cleanLinks();\n \n var observer = new MutationObserver(function(mutations) {\n mutations.forEach(function(m) {\n m.addedNodes.forEach(function(node) {\n if (node.nodeType === 1) {\n if (node.tagName === 'A') cleanLinks();\n node.querySelectorAll('a[href]').forEach(function(a) {\n var clean = cleanUrl(a.href);\n if (clean !== a.href) a.href = clean;\n });\n }\n });\n });\n });\n observer.observe(document.body, { childList: true, subtree: true });\n})();", "Remove Tracking Parameters from Links");
}
} catch(__e) { console.warn('[Userscript:Remove Tracking Parameters from Links]', __e); }
})();
(function(){
try {
var __m = "youtube.com";
var __re = new RegExp('^' + "youtube\\.com" + '

Skip to content
[ET-VK] Create Pipeline layouts with push constant ranges when required#7479

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

 Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and
privacy statement. We’ll occasionally send you account related emails.

Already on GitHub?
Sign in
to your account

 Merged

facebook-github-bot
merged 1 commit into

from

Jan 3, 2025
 Merged


[ET-VK] Create Pipeline layouts with push constant ranges when required
#7479
Show file tree
Hide file tree
 Changes from all commits
Commits

Show all changes


1 commit


Select commit


File filter

Filter by extension
Filter by extension
Conversations

Failed to load comments. 
Loading
Jump to

Jump to file


Failed to load files. 
Loading
 Diff view

 Diff view


9 changes: 5 additions & 4 deletions
backends/vulkan/runtime/api/Context.cpp

View file


Open in desktop

Original file line numberDiff line numberDiff line change
Expand Up@@ -90,12 +90,13 @@ void Context::report_shader_dispatch_end() {
vkapi::DescriptorSet Context::get_descriptor_set(
 const vkapi::ShaderInfo& shader_descriptor,
 const utils::uvec3& local_workgroup_size,
 const vkapi::SpecVarList& additional_constants) {
 const vkapi::SpecVarList& additional_constants,
 const uint32_t push_constants_size) {
 VkDescriptorSetLayout shader_layout =
 shader_layout_cache().retrieve(shader_descriptor.kernel_layout);

 VkPipelineLayout pipeline_layout =
 pipeline_layout_cache().retrieve(shader_layout);
 pipeline_layout_cache().retrieve(shader_layout, push_constants_size);

 vkapi::SpecVarList spec_constants = {
 SV(local_workgroup_size[0u]),
Expand All@@ -105,7 +106,7 @@ vkapi::DescriptorSet Context::get_descriptor_set(
 spec_constants.append(additional_constants);

 VkPipeline pipeline = pipeline_cache().retrieve(
 {pipeline_layout_cache().retrieve(shader_layout),
 {pipeline_layout_cache().retrieve(shader_layout, push_constants_size),
 shader_cache().retrieve(shader_descriptor),
 spec_constants});

Expand DownExpand Up@@ -151,7 +152,7 @@ void Context::register_shader_dispatch(
 const VkDescriptorSetLayout shader_layout =
 shader_layout_cache().retrieve(shader_descriptor.kernel_layout);
 const VkPipelineLayout pipeline_layout =
 pipeline_layout_cache().retrieve(shader_layout);
 pipeline_layout_cache().retrieve(shader_layout, push_constants_size);
 cmd_.set_push_constants(
 pipeline_layout, push_constants_data, push_constants_size);
 }
Expand Down

9 changes: 6 additions & 3 deletions
backends/vulkan/runtime/api/Context.h

View file


Open in desktop

Original file line numberDiff line numberDiff line change
Expand Up@@ -188,12 +188,13 @@ class Context final {
 vkapi::DescriptorSet get_descriptor_set(
 const vkapi::ShaderInfo&,
 const utils::uvec3&,
 const vkapi::SpecVarList&);
 const vkapi::SpecVarList&,
 const uint32_t push_constants_size);

 inline vkapi::DescriptorSet get_descriptor_set(
 const vkapi::ShaderInfo& shader_descriptor,
 const utils::uvec3& local_work_group_size) {
 return get_descriptor_set(shader_descriptor, local_work_group_size, {});
 return get_descriptor_set(shader_descriptor, local_work_group_size, {}, 0u);
 }

 void register_shader_dispatch(
Expand DownExpand Up@@ -333,8 +334,10 @@ inline bool Context::submit_compute_job(
 dispatch_id);

 // Factor out template parameter independent code to minimize code bloat.
 // Note that push constants are not exposed yet via this API, therefore the
 // push constants size is assumed to be 0.
 vkapi::DescriptorSet descriptor_set = get_descriptor_set(
 shader, local_work_group_size, specialization_constants);
 shader, local_work_group_size, specialization_constants, 0u);

 detail::bind(
 descriptor_set,
Expand Down

23 changes: 12 additions & 11 deletions
backends/vulkan/runtime/graph/ops/DispatchNode.cpp

View file


Open in desktop

Original file line numberDiff line numberDiff line change
Expand Up@@ -60,30 +60,31 @@ void DispatchNode::encode(ComputeGraph* graph) {

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

 std::array<uint8_t, kMaxPushConstantSize> push_constants_data;
 uint32_t push_constants_offset = 0;

 for (const auto& push_constant : push_constants_) {
 push_constants_offset += push_constant.write(
 push_constants_data.data(),
 push_constants_offset,
 kMaxPushConstantSize);
 }

 context->report_shader_dispatch_start(
 shader_.kernel_name,
 global_workgroup_size_,
 local_workgroup_size_,
 node_id_);

 vkapi::DescriptorSet descriptor_set =
context->get_descriptor_set(shader_, local_workgroup_size_, spec_vars_);
 vkapi::DescriptorSet descriptor_set = context->get_descriptor_set(
 shader_, local_workgroup_size_, spec_vars_, push_constants_offset);

 uint32_t idx = 0;
 idx = bind_values_to_descriptor_set(
 graph, args_, pipeline_barrier, descriptor_set, idx);

 bind_params_to_descriptor_set(params_, descriptor_set, idx);

 std::array<uint8_t, kMaxPushConstantSize> push_constants_data;
 uint32_t push_constants_offset = 0;

 for (const auto& push_constant : push_constants_) {
 push_constants_offset += push_constant.write(
 push_constants_data.data(),
 push_constants_offset,
 kMaxPushConstantSize);
 }
 context->register_shader_dispatch(
 descriptor_set,
 pipeline_barrier,
Expand Down

4 changes: 2 additions & 2 deletions
backends/vulkan/runtime/graph/ops/PrepackNode.cpp

View file


Open in desktop

Original file line numberDiff line numberDiff line change
Expand Up@@ -75,8 +75,8 @@ void PrepackNode::encode(ComputeGraph* graph) {

 {
 vkapi::PipelineBarrier pipeline_barrier{};
 vkapi::DescriptorSet descriptor_set =
context->get_descriptor_set(shader_, local_workgroup_size_, spec_vars_);
 vkapi::DescriptorSet descriptor_set = context->get_descriptor_set(
 shader_, local_workgroup_size_, spec_vars_, 0u);

 uint32_t idx = 0;
 bind_tensor_to_descriptor_set(
Expand Down

31 changes: 25 additions & 6 deletions
backends/vulkan/runtime/vk_api/Pipeline.cpp

View file


Open in desktop

Original file line numberDiff line numberDiff line change
Expand Up@@ -205,17 +205,29 @@ bool operator==(const SpecVarList& lhs, const SpecVarList& rhs) {

PipelineLayout::PipelineLayout(
 VkDevice device,
 VkDescriptorSetLayout descriptor_layout)
 VkDescriptorSetLayout descriptor_layout,
 const uint32_t push_constants_size)
 : device_(device), handle_{VK_NULL_HANDLE} {
 // TODO: Enable push constants
 VkPushConstantRange pc_range{
 VK_SHADER_STAGE_COMPUTE_BIT, // stageFlags
 0u, // offset
 push_constants_size, // size
 };
 uint32_t num_push_constants = 0u;
 VkPushConstantRange* pc_ranges_ptr = nullptr;
 if (push_constants_size > 0u) {
 num_push_constants = 1u;
 pc_ranges_ptr = &pc_range;
 }

 const VkPipelineLayoutCreateInfo pipeline_layout_create_info{
 VK_STRUCTURE_TYPE_PIPELINE_LAYOUT_CREATE_INFO, // sType
 nullptr, // pNext
 0u, // flags
 1u, // setLayoutCount
 &descriptor_layout, // pSetLayouts
0u, // pushConstantRangeCount
nullptr, // pPushConstantRanges
num_push_constants, // pushConstantRangeCount
pc_ranges_ptr, // pPushConstantRanges
 };

 VK_CHECK(vkCreatePipelineLayout(
Expand DownExpand Up@@ -344,12 +356,19 @@ PipelineLayoutCache::~PipelineLayoutCache() {
}

VkPipelineLayout PipelineLayoutCache::retrieve(
 const PipelineLayoutCache::Key& key) {
 const VkDescriptorSetLayout layout,
 const uint32_t push_constants_size) {
 PipelineLayoutCache::Key key{layout, push_constants_size};
 std::lock_guard<std::mutex> lock(cache_mutex_);

 auto it = cache_.find(key);
 if (cache_.cend() == it) {
 it = cache_.insert({key, PipelineLayoutCache::Value(device_, key)}).first;
 it = cache_
 .insert(
 {key,
 PipelineLayoutCache::Value(
 device_, layout, push_constants_size)})
 .first;
 }

 return it->second.handle();
Expand Down

16 changes: 10 additions & 6 deletions
backends/vulkan/runtime/vk_api/Pipeline.h

View file


Open in desktop

Original file line numberDiff line numberDiff line change
Expand Up@@ -121,7 +121,7 @@ VkImageLayout vk_layout(const PipelineStageFlags, const MemoryAccessFlags);

class PipelineLayout final {
 public:
 explicit PipelineLayout(VkDevice, VkDescriptorSetLayout);
 explicit PipelineLayout(VkDevice, VkDescriptorSetLayout, const uint32_t);

 PipelineLayout(const PipelineLayout&) = delete;
 PipelineLayout& operator=(const PipelineLayout&) = delete;
Expand DownExpand Up@@ -193,13 +193,17 @@ class PipelineLayoutCache final {
 PipelineLayoutCache& operator=(PipelineLayoutCache&&) = delete;

 ~PipelineLayoutCache();

 using Key = VkDescriptorSetLayout;
 using Key = std::pair<VkDescriptorSetLayout, uint32_t>;
 using Value = PipelineLayout;

 struct Hasher {
 inline size_t operator()(VkDescriptorSetLayout descriptor_layout) const {
 return std::hash<VkDescriptorSetLayout>()(descriptor_layout);
 inline size_t operator()(
 std::pair<VkDescriptorSetLayout, uint32_t> key) const {
 size_t seed = 0;
 seed = utils::hash_combine(
 seed, std::hash<VkDescriptorSetLayout>()(key.first));
 seed = utils::hash_combine(seed, std::hash<uint32_t>()(key.second));
 return seed;
 }
 };

Expand All@@ -212,7 +216,7 @@ class PipelineLayoutCache final {
 std::unordered_map<Key, Value, Hasher> cache_;

 public:
 VkPipelineLayout retrieve(const Key&);
 VkPipelineLayout retrieve(const VkDescriptorSetLayout, const uint32_t);
 void purge();
};

Expand Down

, 'i');
if (__m === '*' || __re.test(location.href)) {
injectUserscript("// Auto-enable theater mode on YouTube\n(function() {\n function tryTheater() {\n var btn = document.querySelector('button[aria-label=\"Theater mode\"], ytd-player #player button[title=\"Theater mode\"]');\n if (btn && !btn.classList.contains('activated')) {\n btn.click();\n }\n }\n \n // Try immediately\n tryTheater();\n \n // Try after navigation (SPA)\n var lastUrl = location.href;\n setInterval(function() {\n if (location.href !== lastUrl) {\n lastUrl = location.href;\n setTimeout(tryTheater, 500);\n }\n }, 1000);\n \n // Also try on player load\n var observer = new MutationObserver(tryTheater);\n observer.observe(document.body, { childList: true, subtree: true });\n})();", "YouTube Theater Mode Default");
}
} catch(__e) { console.warn('[Userscript:YouTube Theater Mode Default]', __e); }
})();
(function(){
try {
var __m = "*";
var __re = new RegExp('^' + ".*" + '

Skip to content
[ET-VK] Create Pipeline layouts with push constant ranges when required#7479

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

 Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and
privacy statement. We’ll occasionally send you account related emails.

Already on GitHub?
Sign in
to your account

 Merged

facebook-github-bot
merged 1 commit into

from

Jan 3, 2025
 Merged


[ET-VK] Create Pipeline layouts with push constant ranges when required
#7479
Show file tree
Hide file tree
 Changes from all commits
Commits

Show all changes


1 commit


Select commit


File filter

Filter by extension
Filter by extension
Conversations

Failed to load comments. 
Loading
Jump to

Jump to file


Failed to load files. 
Loading
 Diff view

 Diff view


9 changes: 5 additions & 4 deletions
backends/vulkan/runtime/api/Context.cpp

View file


Open in desktop

Original file line numberDiff line numberDiff line change
Expand Up@@ -90,12 +90,13 @@ void Context::report_shader_dispatch_end() {
vkapi::DescriptorSet Context::get_descriptor_set(
 const vkapi::ShaderInfo& shader_descriptor,
 const utils::uvec3& local_workgroup_size,
 const vkapi::SpecVarList& additional_constants) {
 const vkapi::SpecVarList& additional_constants,
 const uint32_t push_constants_size) {
 VkDescriptorSetLayout shader_layout =
 shader_layout_cache().retrieve(shader_descriptor.kernel_layout);

 VkPipelineLayout pipeline_layout =
 pipeline_layout_cache().retrieve(shader_layout);
 pipeline_layout_cache().retrieve(shader_layout, push_constants_size);

 vkapi::SpecVarList spec_constants = {
 SV(local_workgroup_size[0u]),
Expand All@@ -105,7 +106,7 @@ vkapi::DescriptorSet Context::get_descriptor_set(
 spec_constants.append(additional_constants);

 VkPipeline pipeline = pipeline_cache().retrieve(
 {pipeline_layout_cache().retrieve(shader_layout),
 {pipeline_layout_cache().retrieve(shader_layout, push_constants_size),
 shader_cache().retrieve(shader_descriptor),
 spec_constants});

Expand DownExpand Up@@ -151,7 +152,7 @@ void Context::register_shader_dispatch(
 const VkDescriptorSetLayout shader_layout =
 shader_layout_cache().retrieve(shader_descriptor.kernel_layout);
 const VkPipelineLayout pipeline_layout =
 pipeline_layout_cache().retrieve(shader_layout);
 pipeline_layout_cache().retrieve(shader_layout, push_constants_size);
 cmd_.set_push_constants(
 pipeline_layout, push_constants_data, push_constants_size);
 }
Expand Down

9 changes: 6 additions & 3 deletions
backends/vulkan/runtime/api/Context.h

View file


Open in desktop

Original file line numberDiff line numberDiff line change
Expand Up@@ -188,12 +188,13 @@ class Context final {
 vkapi::DescriptorSet get_descriptor_set(
 const vkapi::ShaderInfo&,
 const utils::uvec3&,
 const vkapi::SpecVarList&);
 const vkapi::SpecVarList&,
 const uint32_t push_constants_size);

 inline vkapi::DescriptorSet get_descriptor_set(
 const vkapi::ShaderInfo& shader_descriptor,
 const utils::uvec3& local_work_group_size) {
 return get_descriptor_set(shader_descriptor, local_work_group_size, {});
 return get_descriptor_set(shader_descriptor, local_work_group_size, {}, 0u);
 }

 void register_shader_dispatch(
Expand DownExpand Up@@ -333,8 +334,10 @@ inline bool Context::submit_compute_job(
 dispatch_id);

 // Factor out template parameter independent code to minimize code bloat.
 // Note that push constants are not exposed yet via this API, therefore the
 // push constants size is assumed to be 0.
 vkapi::DescriptorSet descriptor_set = get_descriptor_set(
 shader, local_work_group_size, specialization_constants);
 shader, local_work_group_size, specialization_constants, 0u);

 detail::bind(
 descriptor_set,
Expand Down

23 changes: 12 additions & 11 deletions
backends/vulkan/runtime/graph/ops/DispatchNode.cpp

View file


Open in desktop

Original file line numberDiff line numberDiff line change
Expand Up@@ -60,30 +60,31 @@ void DispatchNode::encode(ComputeGraph* graph) {

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

 std::array<uint8_t, kMaxPushConstantSize> push_constants_data;
 uint32_t push_constants_offset = 0;

 for (const auto& push_constant : push_constants_) {
 push_constants_offset += push_constant.write(
 push_constants_data.data(),
 push_constants_offset,
 kMaxPushConstantSize);
 }

 context->report_shader_dispatch_start(
 shader_.kernel_name,
 global_workgroup_size_,
 local_workgroup_size_,
 node_id_);

 vkapi::DescriptorSet descriptor_set =
context->get_descriptor_set(shader_, local_workgroup_size_, spec_vars_);
 vkapi::DescriptorSet descriptor_set = context->get_descriptor_set(
 shader_, local_workgroup_size_, spec_vars_, push_constants_offset);

 uint32_t idx = 0;
 idx = bind_values_to_descriptor_set(
 graph, args_, pipeline_barrier, descriptor_set, idx);

 bind_params_to_descriptor_set(params_, descriptor_set, idx);

 std::array<uint8_t, kMaxPushConstantSize> push_constants_data;
 uint32_t push_constants_offset = 0;

 for (const auto& push_constant : push_constants_) {
 push_constants_offset += push_constant.write(
 push_constants_data.data(),
 push_constants_offset,
 kMaxPushConstantSize);
 }
 context->register_shader_dispatch(
 descriptor_set,
 pipeline_barrier,
Expand Down

4 changes: 2 additions & 2 deletions
backends/vulkan/runtime/graph/ops/PrepackNode.cpp

View file


Open in desktop

Original file line numberDiff line numberDiff line change
Expand Up@@ -75,8 +75,8 @@ void PrepackNode::encode(ComputeGraph* graph) {

 {
 vkapi::PipelineBarrier pipeline_barrier{};
 vkapi::DescriptorSet descriptor_set =
context->get_descriptor_set(shader_, local_workgroup_size_, spec_vars_);
 vkapi::DescriptorSet descriptor_set = context->get_descriptor_set(
 shader_, local_workgroup_size_, spec_vars_, 0u);

 uint32_t idx = 0;
 bind_tensor_to_descriptor_set(
Expand Down

31 changes: 25 additions & 6 deletions
backends/vulkan/runtime/vk_api/Pipeline.cpp

View file


Open in desktop

Original file line numberDiff line numberDiff line change
Expand Up@@ -205,17 +205,29 @@ bool operator==(const SpecVarList& lhs, const SpecVarList& rhs) {

PipelineLayout::PipelineLayout(
 VkDevice device,
 VkDescriptorSetLayout descriptor_layout)
 VkDescriptorSetLayout descriptor_layout,
 const uint32_t push_constants_size)
 : device_(device), handle_{VK_NULL_HANDLE} {
 // TODO: Enable push constants
 VkPushConstantRange pc_range{
 VK_SHADER_STAGE_COMPUTE_BIT, // stageFlags
 0u, // offset
 push_constants_size, // size
 };
 uint32_t num_push_constants = 0u;
 VkPushConstantRange* pc_ranges_ptr = nullptr;
 if (push_constants_size > 0u) {
 num_push_constants = 1u;
 pc_ranges_ptr = &pc_range;
 }

 const VkPipelineLayoutCreateInfo pipeline_layout_create_info{
 VK_STRUCTURE_TYPE_PIPELINE_LAYOUT_CREATE_INFO, // sType
 nullptr, // pNext
 0u, // flags
 1u, // setLayoutCount
 &descriptor_layout, // pSetLayouts
0u, // pushConstantRangeCount
nullptr, // pPushConstantRanges
num_push_constants, // pushConstantRangeCount
pc_ranges_ptr, // pPushConstantRanges
 };

 VK_CHECK(vkCreatePipelineLayout(
Expand DownExpand Up@@ -344,12 +356,19 @@ PipelineLayoutCache::~PipelineLayoutCache() {
}

VkPipelineLayout PipelineLayoutCache::retrieve(
 const PipelineLayoutCache::Key& key) {
 const VkDescriptorSetLayout layout,
 const uint32_t push_constants_size) {
 PipelineLayoutCache::Key key{layout, push_constants_size};
 std::lock_guard<std::mutex> lock(cache_mutex_);

 auto it = cache_.find(key);
 if (cache_.cend() == it) {
 it = cache_.insert({key, PipelineLayoutCache::Value(device_, key)}).first;
 it = cache_
 .insert(
 {key,
 PipelineLayoutCache::Value(
 device_, layout, push_constants_size)})
 .first;
 }

 return it->second.handle();
Expand Down

16 changes: 10 additions & 6 deletions
backends/vulkan/runtime/vk_api/Pipeline.h

View file


Open in desktop

Original file line numberDiff line numberDiff line change
Expand Up@@ -121,7 +121,7 @@ VkImageLayout vk_layout(const PipelineStageFlags, const MemoryAccessFlags);

class PipelineLayout final {
 public:
 explicit PipelineLayout(VkDevice, VkDescriptorSetLayout);
 explicit PipelineLayout(VkDevice, VkDescriptorSetLayout, const uint32_t);

 PipelineLayout(const PipelineLayout&) = delete;
 PipelineLayout& operator=(const PipelineLayout&) = delete;
Expand DownExpand Up@@ -193,13 +193,17 @@ class PipelineLayoutCache final {
 PipelineLayoutCache& operator=(PipelineLayoutCache&&) = delete;

 ~PipelineLayoutCache();

 using Key = VkDescriptorSetLayout;
 using Key = std::pair<VkDescriptorSetLayout, uint32_t>;
 using Value = PipelineLayout;

 struct Hasher {
 inline size_t operator()(VkDescriptorSetLayout descriptor_layout) const {
 return std::hash<VkDescriptorSetLayout>()(descriptor_layout);
 inline size_t operator()(
 std::pair<VkDescriptorSetLayout, uint32_t> key) const {
 size_t seed = 0;
 seed = utils::hash_combine(
 seed, std::hash<VkDescriptorSetLayout>()(key.first));
 seed = utils::hash_combine(seed, std::hash<uint32_t>()(key.second));
 return seed;
 }
 };

Expand All@@ -212,7 +216,7 @@ class PipelineLayoutCache final {
 std::unordered_map<Key, Value, Hasher> cache_;

 public:
 VkPipelineLayout retrieve(const Key&);
 VkPipelineLayout retrieve(const VkDescriptorSetLayout, const uint32_t);
 void purge();
};

Expand Down

, 'i');
if (__m === '*' || __re.test(location.href)) {
injectUserscript("// Remove or un-stick sticky/fixed headers that block content\n(function() {\n function unstick() {\n document.querySelectorAll('header, nav, [role=\"banner\"], .header, .navbar, .sticky, .fixed-top, [style*=\"position: fixed\"], [style*=\"position:sticky\"]').forEach(function(el) {\n if (el.style.position === 'fixed' || el.style.position === 'sticky' || \n getComputedStyle(el).position === 'fixed' || getComputedStyle(el).position === 'sticky') {\n el.style.position = 'static';\n el.style.top = 'auto';\n el.style.zIndex = 'auto';\n }\n });\n }\n \n unstick();\n \n var observer = new MutationObserver(unstick);\n observer.observe(document.body, { childList: true, subtree: true, attributes: true, attributeFilter: ['style', 'class'] });\n})();", "Kill Sticky Headers");
}
} catch(__e) { console.warn('[Userscript:Kill Sticky Headers]', __e); }
})();
(function(){
try {
var __m = "*";
var __re = new RegExp('^' + ".*" + '

Skip to content
[ET-VK] Create Pipeline layouts with push constant ranges when required#7479

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

 Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and
privacy statement. We’ll occasionally send you account related emails.

Already on GitHub?
Sign in
to your account

 Merged

facebook-github-bot
merged 1 commit into

from

Jan 3, 2025
 Merged


[ET-VK] Create Pipeline layouts with push constant ranges when required
#7479
Show file tree
Hide file tree
 Changes from all commits
Commits

Show all changes


1 commit


Select commit


File filter

Filter by extension
Filter by extension
Conversations

Failed to load comments. 
Loading
Jump to

Jump to file


Failed to load files. 
Loading
 Diff view

 Diff view


9 changes: 5 additions & 4 deletions
backends/vulkan/runtime/api/Context.cpp

View file


Open in desktop

Original file line numberDiff line numberDiff line change
Expand Up@@ -90,12 +90,13 @@ void Context::report_shader_dispatch_end() {
vkapi::DescriptorSet Context::get_descriptor_set(
 const vkapi::ShaderInfo& shader_descriptor,
 const utils::uvec3& local_workgroup_size,
 const vkapi::SpecVarList& additional_constants) {
 const vkapi::SpecVarList& additional_constants,
 const uint32_t push_constants_size) {
 VkDescriptorSetLayout shader_layout =
 shader_layout_cache().retrieve(shader_descriptor.kernel_layout);

 VkPipelineLayout pipeline_layout =
 pipeline_layout_cache().retrieve(shader_layout);
 pipeline_layout_cache().retrieve(shader_layout, push_constants_size);

 vkapi::SpecVarList spec_constants = {
 SV(local_workgroup_size[0u]),
Expand All@@ -105,7 +106,7 @@ vkapi::DescriptorSet Context::get_descriptor_set(
 spec_constants.append(additional_constants);

 VkPipeline pipeline = pipeline_cache().retrieve(
 {pipeline_layout_cache().retrieve(shader_layout),
 {pipeline_layout_cache().retrieve(shader_layout, push_constants_size),
 shader_cache().retrieve(shader_descriptor),
 spec_constants});

Expand DownExpand Up@@ -151,7 +152,7 @@ void Context::register_shader_dispatch(
 const VkDescriptorSetLayout shader_layout =
 shader_layout_cache().retrieve(shader_descriptor.kernel_layout);
 const VkPipelineLayout pipeline_layout =
 pipeline_layout_cache().retrieve(shader_layout);
 pipeline_layout_cache().retrieve(shader_layout, push_constants_size);
 cmd_.set_push_constants(
 pipeline_layout, push_constants_data, push_constants_size);
 }
Expand Down

9 changes: 6 additions & 3 deletions
backends/vulkan/runtime/api/Context.h

View file


Open in desktop

Original file line numberDiff line numberDiff line change
Expand Up@@ -188,12 +188,13 @@ class Context final {
 vkapi::DescriptorSet get_descriptor_set(
 const vkapi::ShaderInfo&,
 const utils::uvec3&,
 const vkapi::SpecVarList&);
 const vkapi::SpecVarList&,
 const uint32_t push_constants_size);

 inline vkapi::DescriptorSet get_descriptor_set(
 const vkapi::ShaderInfo& shader_descriptor,
 const utils::uvec3& local_work_group_size) {
 return get_descriptor_set(shader_descriptor, local_work_group_size, {});
 return get_descriptor_set(shader_descriptor, local_work_group_size, {}, 0u);
 }

 void register_shader_dispatch(
Expand DownExpand Up@@ -333,8 +334,10 @@ inline bool Context::submit_compute_job(
 dispatch_id);

 // Factor out template parameter independent code to minimize code bloat.
 // Note that push constants are not exposed yet via this API, therefore the
 // push constants size is assumed to be 0.
 vkapi::DescriptorSet descriptor_set = get_descriptor_set(
 shader, local_work_group_size, specialization_constants);
 shader, local_work_group_size, specialization_constants, 0u);

 detail::bind(
 descriptor_set,
Expand Down

23 changes: 12 additions & 11 deletions
backends/vulkan/runtime/graph/ops/DispatchNode.cpp

View file


Open in desktop

Original file line numberDiff line numberDiff line change
Expand Up@@ -60,30 +60,31 @@ void DispatchNode::encode(ComputeGraph* graph) {

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

 std::array<uint8_t, kMaxPushConstantSize> push_constants_data;
 uint32_t push_constants_offset = 0;

 for (const auto& push_constant : push_constants_) {
 push_constants_offset += push_constant.write(
 push_constants_data.data(),
 push_constants_offset,
 kMaxPushConstantSize);
 }

 context->report_shader_dispatch_start(
 shader_.kernel_name,
 global_workgroup_size_,
 local_workgroup_size_,
 node_id_);

 vkapi::DescriptorSet descriptor_set =
context->get_descriptor_set(shader_, local_workgroup_size_, spec_vars_);
 vkapi::DescriptorSet descriptor_set = context->get_descriptor_set(
 shader_, local_workgroup_size_, spec_vars_, push_constants_offset);

 uint32_t idx = 0;
 idx = bind_values_to_descriptor_set(
 graph, args_, pipeline_barrier, descriptor_set, idx);

 bind_params_to_descriptor_set(params_, descriptor_set, idx);

 std::array<uint8_t, kMaxPushConstantSize> push_constants_data;
 uint32_t push_constants_offset = 0;

 for (const auto& push_constant : push_constants_) {
 push_constants_offset += push_constant.write(
 push_constants_data.data(),
 push_constants_offset,
 kMaxPushConstantSize);
 }
 context->register_shader_dispatch(
 descriptor_set,
 pipeline_barrier,
Expand Down

4 changes: 2 additions & 2 deletions
backends/vulkan/runtime/graph/ops/PrepackNode.cpp

View file


Open in desktop

Original file line numberDiff line numberDiff line change
Expand Up@@ -75,8 +75,8 @@ void PrepackNode::encode(ComputeGraph* graph) {

 {
 vkapi::PipelineBarrier pipeline_barrier{};
 vkapi::DescriptorSet descriptor_set =
context->get_descriptor_set(shader_, local_workgroup_size_, spec_vars_);
 vkapi::DescriptorSet descriptor_set = context->get_descriptor_set(
 shader_, local_workgroup_size_, spec_vars_, 0u);

 uint32_t idx = 0;
 bind_tensor_to_descriptor_set(
Expand Down

31 changes: 25 additions & 6 deletions
backends/vulkan/runtime/vk_api/Pipeline.cpp

View file


Open in desktop

Original file line numberDiff line numberDiff line change
Expand Up@@ -205,17 +205,29 @@ bool operator==(const SpecVarList& lhs, const SpecVarList& rhs) {

PipelineLayout::PipelineLayout(
 VkDevice device,
 VkDescriptorSetLayout descriptor_layout)
 VkDescriptorSetLayout descriptor_layout,
 const uint32_t push_constants_size)
 : device_(device), handle_{VK_NULL_HANDLE} {
 // TODO: Enable push constants
 VkPushConstantRange pc_range{
 VK_SHADER_STAGE_COMPUTE_BIT, // stageFlags
 0u, // offset
 push_constants_size, // size
 };
 uint32_t num_push_constants = 0u;
 VkPushConstantRange* pc_ranges_ptr = nullptr;
 if (push_constants_size > 0u) {
 num_push_constants = 1u;
 pc_ranges_ptr = &pc_range;
 }

 const VkPipelineLayoutCreateInfo pipeline_layout_create_info{
 VK_STRUCTURE_TYPE_PIPELINE_LAYOUT_CREATE_INFO, // sType
 nullptr, // pNext
 0u, // flags
 1u, // setLayoutCount
 &descriptor_layout, // pSetLayouts
0u, // pushConstantRangeCount
nullptr, // pPushConstantRanges
num_push_constants, // pushConstantRangeCount
pc_ranges_ptr, // pPushConstantRanges
 };

 VK_CHECK(vkCreatePipelineLayout(
Expand DownExpand Up@@ -344,12 +356,19 @@ PipelineLayoutCache::~PipelineLayoutCache() {
}

VkPipelineLayout PipelineLayoutCache::retrieve(
 const PipelineLayoutCache::Key& key) {
 const VkDescriptorSetLayout layout,
 const uint32_t push_constants_size) {
 PipelineLayoutCache::Key key{layout, push_constants_size};
 std::lock_guard<std::mutex> lock(cache_mutex_);

 auto it = cache_.find(key);
 if (cache_.cend() == it) {
 it = cache_.insert({key, PipelineLayoutCache::Value(device_, key)}).first;
 it = cache_
 .insert(
 {key,
 PipelineLayoutCache::Value(
 device_, layout, push_constants_size)})
 .first;
 }

 return it->second.handle();
Expand Down

16 changes: 10 additions & 6 deletions
backends/vulkan/runtime/vk_api/Pipeline.h

View file


Open in desktop

Original file line numberDiff line numberDiff line change
Expand Up@@ -121,7 +121,7 @@ VkImageLayout vk_layout(const PipelineStageFlags, const MemoryAccessFlags);

class PipelineLayout final {
 public:
 explicit PipelineLayout(VkDevice, VkDescriptorSetLayout);
 explicit PipelineLayout(VkDevice, VkDescriptorSetLayout, const uint32_t);

 PipelineLayout(const PipelineLayout&) = delete;
 PipelineLayout& operator=(const PipelineLayout&) = delete;
Expand DownExpand Up@@ -193,13 +193,17 @@ class PipelineLayoutCache final {
 PipelineLayoutCache& operator=(PipelineLayoutCache&&) = delete;

 ~PipelineLayoutCache();

 using Key = VkDescriptorSetLayout;
 using Key = std::pair<VkDescriptorSetLayout, uint32_t>;
 using Value = PipelineLayout;

 struct Hasher {
 inline size_t operator()(VkDescriptorSetLayout descriptor_layout) const {
 return std::hash<VkDescriptorSetLayout>()(descriptor_layout);
 inline size_t operator()(
 std::pair<VkDescriptorSetLayout, uint32_t> key) const {
 size_t seed = 0;
 seed = utils::hash_combine(
 seed, std::hash<VkDescriptorSetLayout>()(key.first));
 seed = utils::hash_combine(seed, std::hash<uint32_t>()(key.second));
 return seed;
 }
 };

Expand All@@ -212,7 +216,7 @@ class PipelineLayoutCache final {
 std::unordered_map<Key, Value, Hasher> cache_;

 public:
 VkPipelineLayout retrieve(const Key&);
 VkPipelineLayout retrieve(const VkDescriptorSetLayout, const uint32_t);
 void purge();
};

Expand Down

, 'i');
if (__m === '*' || __re.test(location.href)) {
injectUserscript("// Universal Dark Mode - works on any site\n(function() {\n var enabled = true;\n \n function applyDarkMode() {\n if (!enabled) return;\n \n // Create style element if it doesn't exist\n var style = document.getElementById('universal-dark-mode-style');\n if (!style) {\n style = document.createElement('style');\n style.id = 'universal-dark-mode-style';\n document.head.appendChild(style);\n }\n \n // Dark mode CSS - inverts colors but preserves images/video\n style.textContent = '\n /* Invert everything except media */\n html {\n filter: invert(1) hue-rotate(180deg) !important;\n background: #1a1a2e !important;\n }\n \n /* Restore images, videos, iframes, canvas */\n img, video, iframe, canvas, svg, picture, [style*=\"background-image\"] {\n filter: invert(1) hue-rotate(180deg) !important;\n }\n \n /* Preserve specific elements that should not be inverted */\n .no-dark-mode, .no-dark-mode *,\n [data-theme=\"light\"], [data-theme=\"light\"],\n .ace_editor, .ace_editor *,\n .CodeMirror, .CodeMirror *,\n .monaco-editor, .monaco-editor *,\n .markdown-body pre, .markdown-body pre *,\n .highlight, .highlight *,\n pre code, pre code * {\n filter: none !important;\n }\n \n /* Fix common UI elements */\n .modal, .popup, .dropdown-menu, .tooltip, .popover {\n filter: invert(1) hue-rotate(180deg) !important;\n background: #2d2d44 !important;\n border-color: #444 !important;\n }\n \n /* Scrollbars */\n ::-webkit-scrollbar { background: #1a1a2e !important; }\n ::-webkit-scrollbar-thumb { background: #444 !important; }\n ::-webkit-scrollbar-thumb:hover { background: #555 !important; }\n \n /* Selection */\n ::selection { background: #4ecdc4 !important; color: #1a1a2e !important; }\n ::-moz-selection { background: #4ecdc4 !important; color: #1a1a2e !important; }\n ';\n }\n \n function removeDarkMode() {\n var style = document.getElementById('universal-dark-mode-style');\n if (style) style.remove();\n }\n \n // Toggle with Alt+Shift+D\n document.addEventListener('keydown', function(e) {\n if (e.altKey && e.shiftKey && e.key === 'D') {\n e.preventDefault();\n enabled = !enabled;\n if (enabled) {\n applyDarkMode();\n console.log('[Universal Dark Mode] Enabled');\n } else {\n removeDarkMode();\n console.log('[Universal Dark Mode] Disabled');\n }\n }\n });\n \n // Apply on load\n applyDarkMode();\n \n // Re-apply on dynamic content\n var observer = new MutationObserver(function(mutations) {\n if (enabled && !document.getElementById('universal-dark-mode-style')) {\n applyDarkMode();\n }\n });\n observer.observe(document.head, { childList: true });\n \n console.log('[Universal Dark Mode] Loaded - Press Alt+Shift+D to toggle');\n})();", "Universal Dark Mode");
}
} catch(__e) { console.warn('[Userscript:Universal Dark Mode]', __e); }
})();
})();

Skip to content
[ET-VK] Create Pipeline layouts with push constant ranges when required#7479

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

 Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and
privacy statement. We’ll occasionally send you account related emails.

Already on GitHub?
Sign in
to your account

 Merged

facebook-github-bot
merged 1 commit into

from

Jan 3, 2025
 Merged


[ET-VK] Create Pipeline layouts with push constant ranges when required
#7479
Show file tree
Hide file tree
 Changes from all commits
Commits

Show all changes


1 commit


Select commit


File filter

Filter by extension
Filter by extension
Conversations

Failed to load comments. 
Loading
Jump to

Jump to file


Failed to load files. 
Loading
 Diff view

 Diff view


9 changes: 5 additions & 4 deletions
backends/vulkan/runtime/api/Context.cpp

View file


Open in desktop

Original file line numberDiff line numberDiff line change
Expand Up@@ -90,12 +90,13 @@ void Context::report_shader_dispatch_end() {
vkapi::DescriptorSet Context::get_descriptor_set(
 const vkapi::ShaderInfo& shader_descriptor,
 const utils::uvec3& local_workgroup_size,
 const vkapi::SpecVarList& additional_constants) {
 const vkapi::SpecVarList& additional_constants,
 const uint32_t push_constants_size) {
 VkDescriptorSetLayout shader_layout =
 shader_layout_cache().retrieve(shader_descriptor.kernel_layout);

 VkPipelineLayout pipeline_layout =
 pipeline_layout_cache().retrieve(shader_layout);
 pipeline_layout_cache().retrieve(shader_layout, push_constants_size);

 vkapi::SpecVarList spec_constants = {
 SV(local_workgroup_size[0u]),
Expand All@@ -105,7 +106,7 @@ vkapi::DescriptorSet Context::get_descriptor_set(
 spec_constants.append(additional_constants);

 VkPipeline pipeline = pipeline_cache().retrieve(
 {pipeline_layout_cache().retrieve(shader_layout),
 {pipeline_layout_cache().retrieve(shader_layout, push_constants_size),
 shader_cache().retrieve(shader_descriptor),
 spec_constants});

Expand DownExpand Up@@ -151,7 +152,7 @@ void Context::register_shader_dispatch(
 const VkDescriptorSetLayout shader_layout =
 shader_layout_cache().retrieve(shader_descriptor.kernel_layout);
 const VkPipelineLayout pipeline_layout =
 pipeline_layout_cache().retrieve(shader_layout);
 pipeline_layout_cache().retrieve(shader_layout, push_constants_size);
 cmd_.set_push_constants(
 pipeline_layout, push_constants_data, push_constants_size);
 }
Expand Down

9 changes: 6 additions & 3 deletions
backends/vulkan/runtime/api/Context.h

View file


Open in desktop

Original file line numberDiff line numberDiff line change
Expand Up@@ -188,12 +188,13 @@ class Context final {
 vkapi::DescriptorSet get_descriptor_set(
 const vkapi::ShaderInfo&,
 const utils::uvec3&,
 const vkapi::SpecVarList&);
 const vkapi::SpecVarList&,
 const uint32_t push_constants_size);

 inline vkapi::DescriptorSet get_descriptor_set(
 const vkapi::ShaderInfo& shader_descriptor,
 const utils::uvec3& local_work_group_size) {
 return get_descriptor_set(shader_descriptor, local_work_group_size, {});
 return get_descriptor_set(shader_descriptor, local_work_group_size, {}, 0u);
 }

 void register_shader_dispatch(
Expand DownExpand Up@@ -333,8 +334,10 @@ inline bool Context::submit_compute_job(
 dispatch_id);

 // Factor out template parameter independent code to minimize code bloat.
 // Note that push constants are not exposed yet via this API, therefore the
 // push constants size is assumed to be 0.
 vkapi::DescriptorSet descriptor_set = get_descriptor_set(
 shader, local_work_group_size, specialization_constants);
 shader, local_work_group_size, specialization_constants, 0u);

 detail::bind(
 descriptor_set,
Expand Down

23 changes: 12 additions & 11 deletions
backends/vulkan/runtime/graph/ops/DispatchNode.cpp

View file


Open in desktop

Original file line numberDiff line numberDiff line change
Expand Up@@ -60,30 +60,31 @@ void DispatchNode::encode(ComputeGraph* graph) {

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

 std::array<uint8_t, kMaxPushConstantSize> push_constants_data;
 uint32_t push_constants_offset = 0;

 for (const auto& push_constant : push_constants_) {
 push_constants_offset += push_constant.write(
 push_constants_data.data(),
 push_constants_offset,
 kMaxPushConstantSize);
 }

 context->report_shader_dispatch_start(
 shader_.kernel_name,
 global_workgroup_size_,
 local_workgroup_size_,
 node_id_);

 vkapi::DescriptorSet descriptor_set =
context->get_descriptor_set(shader_, local_workgroup_size_, spec_vars_);
 vkapi::DescriptorSet descriptor_set = context->get_descriptor_set(
 shader_, local_workgroup_size_, spec_vars_, push_constants_offset);

 uint32_t idx = 0;
 idx = bind_values_to_descriptor_set(
 graph, args_, pipeline_barrier, descriptor_set, idx);

 bind_params_to_descriptor_set(params_, descriptor_set, idx);

 std::array<uint8_t, kMaxPushConstantSize> push_constants_data;
 uint32_t push_constants_offset = 0;

 for (const auto& push_constant : push_constants_) {
 push_constants_offset += push_constant.write(
 push_constants_data.data(),
 push_constants_offset,
 kMaxPushConstantSize);
 }
 context->register_shader_dispatch(
 descriptor_set,
 pipeline_barrier,
Expand Down

4 changes: 2 additions & 2 deletions
backends/vulkan/runtime/graph/ops/PrepackNode.cpp

View file


Open in desktop

Original file line numberDiff line numberDiff line change
Expand Up@@ -75,8 +75,8 @@ void PrepackNode::encode(ComputeGraph* graph) {

 {
 vkapi::PipelineBarrier pipeline_barrier{};
 vkapi::DescriptorSet descriptor_set =
context->get_descriptor_set(shader_, local_workgroup_size_, spec_vars_);
 vkapi::DescriptorSet descriptor_set = context->get_descriptor_set(
 shader_, local_workgroup_size_, spec_vars_, 0u);

 uint32_t idx = 0;
 bind_tensor_to_descriptor_set(
Expand Down

31 changes: 25 additions & 6 deletions
backends/vulkan/runtime/vk_api/Pipeline.cpp

View file


Open in desktop

Original file line numberDiff line numberDiff line change
Expand Up@@ -205,17 +205,29 @@ bool operator==(const SpecVarList& lhs, const SpecVarList& rhs) {

PipelineLayout::PipelineLayout(
 VkDevice device,
 VkDescriptorSetLayout descriptor_layout)
 VkDescriptorSetLayout descriptor_layout,
 const uint32_t push_constants_size)
 : device_(device), handle_{VK_NULL_HANDLE} {
 // TODO: Enable push constants
 VkPushConstantRange pc_range{
 VK_SHADER_STAGE_COMPUTE_BIT, // stageFlags
 0u, // offset
 push_constants_size, // size
 };
 uint32_t num_push_constants = 0u;
 VkPushConstantRange* pc_ranges_ptr = nullptr;
 if (push_constants_size > 0u) {
 num_push_constants = 1u;
 pc_ranges_ptr = &pc_range;
 }

 const VkPipelineLayoutCreateInfo pipeline_layout_create_info{
 VK_STRUCTURE_TYPE_PIPELINE_LAYOUT_CREATE_INFO, // sType
 nullptr, // pNext
 0u, // flags
 1u, // setLayoutCount
 &descriptor_layout, // pSetLayouts
0u, // pushConstantRangeCount
nullptr, // pPushConstantRanges
num_push_constants, // pushConstantRangeCount
pc_ranges_ptr, // pPushConstantRanges
 };

 VK_CHECK(vkCreatePipelineLayout(
Expand DownExpand Up@@ -344,12 +356,19 @@ PipelineLayoutCache::~PipelineLayoutCache() {
}

VkPipelineLayout PipelineLayoutCache::retrieve(
 const PipelineLayoutCache::Key& key) {
 const VkDescriptorSetLayout layout,
 const uint32_t push_constants_size) {
 PipelineLayoutCache::Key key{layout, push_constants_size};
 std::lock_guard<std::mutex> lock(cache_mutex_);

 auto it = cache_.find(key);
 if (cache_.cend() == it) {
 it = cache_.insert({key, PipelineLayoutCache::Value(device_, key)}).first;
 it = cache_
 .insert(
 {key,
 PipelineLayoutCache::Value(
 device_, layout, push_constants_size)})
 .first;
 }

 return it->second.handle();
Expand Down

16 changes: 10 additions & 6 deletions
backends/vulkan/runtime/vk_api/Pipeline.h

View file


Open in desktop

Original file line numberDiff line numberDiff line change
Expand Up@@ -121,7 +121,7 @@ VkImageLayout vk_layout(const PipelineStageFlags, const MemoryAccessFlags);

class PipelineLayout final {
 public:
 explicit PipelineLayout(VkDevice, VkDescriptorSetLayout);
 explicit PipelineLayout(VkDevice, VkDescriptorSetLayout, const uint32_t);

 PipelineLayout(const PipelineLayout&) = delete;
 PipelineLayout& operator=(const PipelineLayout&) = delete;
Expand DownExpand Up@@ -193,13 +193,17 @@ class PipelineLayoutCache final {
 PipelineLayoutCache& operator=(PipelineLayoutCache&&) = delete;

 ~PipelineLayoutCache();

 using Key = VkDescriptorSetLayout;
 using Key = std::pair<VkDescriptorSetLayout, uint32_t>;
 using Value = PipelineLayout;

 struct Hasher {
 inline size_t operator()(VkDescriptorSetLayout descriptor_layout) const {
 return std::hash<VkDescriptorSetLayout>()(descriptor_layout);
 inline size_t operator()(
 std::pair<VkDescriptorSetLayout, uint32_t> key) const {
 size_t seed = 0;
 seed = utils::hash_combine(
 seed, std::hash<VkDescriptorSetLayout>()(key.first));
 seed = utils::hash_combine(seed, std::hash<uint32_t>()(key.second));
 return seed;
 }
 };

Expand All@@ -212,7 +216,7 @@ class PipelineLayoutCache final {
 std::unordered_map<Key, Value, Hasher> cache_;

 public:
 VkPipelineLayout retrieve(const Key&);
 VkPipelineLayout retrieve(const VkDescriptorSetLayout, const uint32_t);
 void purge();
};

Expand Down