Skip to content

[ET-VK] Move files using the Vulkan API to vk_api/ #4125

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

Closed
junpi3 wants to merge 3 commits into from
Closed

[ET-VK] Move files using the Vulkan API to vk_api/ #4125

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
3 commits
Select commit Hold shift + click to select a range
3edd7fd
[ET-VK] Move files using the Vulkan API to `vk_api/`
Jul 2, 2024
1f0fe59
Update on "[ET-VK] Move files using the Vulkan API to `vk_api/`"
Jul 2, 2024
7828080
Update on "[ET-VK] Move files using the Vulkan API to `vk_api/`"
Jul 2, 2024
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
44 changes: 22 additions & 22 deletions backends/vulkan/runtime/VulkanBackend.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -59,47 +59,47 @@ const uint8_t* get_constant_data_ptr(
return constant_data + constant_bytes->offset();
}

api::ScalarType get_scalar_type(const vkgraph::VkDataType& vk_datatype) {
vkapi::ScalarType get_scalar_type(const vkgraph::VkDataType& vk_datatype) {
switch (vk_datatype) {
case vkgraph::VkDataType::BOOL:
return api::kBool;
return vkapi::kBool;
case vkgraph::VkDataType::UINT8:
return api::kByte;
return vkapi::kByte;
case vkgraph::VkDataType::INT8:
return api::kChar;
return vkapi::kChar;
case vkgraph::VkDataType::INT32:
return api::kInt;
return vkapi::kInt;
case vkgraph::VkDataType::FLOAT16:
return api::kHalf;
return vkapi::kHalf;
case vkgraph::VkDataType::FLOAT32:
return api::kFloat;
return vkapi::kFloat;
}
}

api::StorageType get_storage_type(
vkapi::StorageType get_storage_type(
const vkgraph::VkStorageType& vk_storage_type) {
switch (vk_storage_type) {
case vkgraph::VkStorageType::BUFFER:
return api::kBuffer;
return vkapi::kBuffer;
case vkgraph::VkStorageType::TEXTURE_3D:
return api::kTexture3D;
return vkapi::kTexture3D;
case vkgraph::VkStorageType::TEXTURE_2D:
return api::kTexture2D;
return vkapi::kTexture2D;
default:
break;
}
VK_THROW("Invalid storage type encountered!");
}

api::GPUMemoryLayout get_memory_layout(
vkapi::GPUMemoryLayout get_memory_layout(
const vkgraph::VkMemoryLayout& vk_memory_layout) {
switch (vk_memory_layout) {
case vkgraph::VkMemoryLayout::TENSOR_WIDTH_PACKED:
return api::kWidthPacked;
return vkapi::kWidthPacked;
case vkgraph::VkMemoryLayout::TENSOR_HEIGHT_PACKED:
return api::kHeightPacked;
return vkapi::kHeightPacked;
case vkgraph::VkMemoryLayout::TENSOR_CHANNELS_PACKED:
return api::kChannelsPacked;
return vkapi::kChannelsPacked;
default:
break;
}
Expand All @@ -115,16 +115,16 @@ GraphConfig get_graph_config(ArrayRef<CompileSpec>& compile_specs) {
if (strcmp(spec.key, "storage_type_override") == 0) {
ET_CHECK_MSG(value_size == sizeof(int32_t), "Unexpected value size!");
int value_as_int = static_cast<int>(getUInt32LE(value_data));
api::StorageType storage_type =
static_cast<api::StorageType>(value_as_int);
vkapi::StorageType storage_type =
static_cast<vkapi::StorageType>(value_as_int);

config.set_storage_type_override(storage_type);
}
if (strcmp(spec.key, "memory_layout_override") == 0) {
ET_CHECK_MSG(value_size == sizeof(uint32_t), "Unexpected value size!");
uint32_t value_as_int = getUInt32LE(value_data);
api::GPUMemoryLayout memory_layout =
static_cast<api::GPUMemoryLayout>(value_as_int);
vkapi::GPUMemoryLayout memory_layout =
static_cast<vkapi::GPUMemoryLayout>(value_as_int);

config.set_memory_layout_override(memory_layout);
}
Expand Down Expand Up @@ -171,16 +171,16 @@ class GraphBuilder {
}

void add_tensor_to_graph(const uint32_t fb_id, VkTensorPtr tensor_fb) {
const api::ScalarType& dtype = get_scalar_type(tensor_fb->datatype());
api::StorageType storage_type =
const vkapi::ScalarType& dtype = get_scalar_type(tensor_fb->datatype());
vkapi::StorageType storage_type =
tensor_fb->storage_type() == vkgraph::VkStorageType::DEFAULT_STORAGE
? compute_graph_->suggested_storage_type()
: get_storage_type(tensor_fb->storage_type());

UIntVector dims_fb = tensor_fb->dims();
const std::vector<int64_t> dims_vector(dims_fb->cbegin(), dims_fb->cend());

api::GPUMemoryLayout memory_layout =
vkapi::GPUMemoryLayout memory_layout =
tensor_fb->memory_layout() == vkgraph::VkMemoryLayout::DEFAULT_LAYOUT
? compute_graph_->suggested_memory_layout(dims_vector)
: get_memory_layout(tensor_fb->memory_layout());
Expand Down
31 changes: 16 additions & 15 deletions backends/vulkan/runtime/api/Context.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -7,7 +7,8 @@
*/

#include <executorch/backends/vulkan/runtime/api/Context.h>
#include <executorch/backends/vulkan/runtime/api/VkUtils.h>

#include <executorch/backends/vulkan/runtime/api/vk_api/VkUtils.h>

#ifndef VULKAN_DESCRIPTOR_POOL_SIZE
#define VULKAN_DESCRIPTOR_POOL_SIZE 1024u
Expand All @@ -23,7 +24,7 @@ namespace api {
Context::Context(size_t adapter_i, const ContextConfig& config)
: config_(config),
// Important handles
adapter_p_(runtime()->get_adapter_p(adapter_i)),
adapter_p_(vkapi::runtime()->get_adapter_p(adapter_i)),
device_(adapter_p_->device_handle()),
queue_(adapter_p_->request_queue()),
// Resource pools
Expand Down Expand Up @@ -72,8 +73,8 @@ void Context::report_shader_dispatch_start(
cmd_,
dispatch_id,
shader_name,
create_extent3d(global_wg_size),
create_extent3d(local_wg_size));
vkapi::create_extent3d(global_wg_size),
vkapi::create_extent3d(local_wg_size));
}
}

Expand All @@ -83,17 +84,17 @@ void Context::report_shader_dispatch_end() {
}
}

DescriptorSet Context::get_descriptor_set(
const ShaderInfo& shader_descriptor,
vkapi::DescriptorSet Context::get_descriptor_set(
const vkapi::ShaderInfo& shader_descriptor,
const utils::uvec3& local_workgroup_size,
const SpecVarList& additional_constants) {
const vkapi::SpecVarList& additional_constants) {
VkDescriptorSetLayout shader_layout =
shader_layout_cache().retrieve(shader_descriptor.kernel_layout);

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

SpecVarList spec_constants = {
vkapi::SpecVarList spec_constants = {
SV(local_workgroup_size.data[0u]),
SV(local_workgroup_size.data[1u]),
SV(local_workgroup_size.data[2u])};
Expand All @@ -112,9 +113,9 @@ DescriptorSet Context::get_descriptor_set(
}

void Context::register_shader_dispatch(
const DescriptorSet& descriptors,
PipelineBarrier& pipeline_barrier,
const ShaderInfo& shader_descriptor,
const vkapi::DescriptorSet& descriptors,
vkapi::PipelineBarrier& pipeline_barrier,
const vkapi::ShaderInfo& shader_descriptor,
const utils::uvec3& global_workgroup_size) {
// Adjust the global workgroup size based on the output tile size
uint32_t global_wg_w = utils::div_up(
Expand Down Expand Up @@ -180,12 +181,12 @@ Context* context() {
try {
const uint32_t cmd_submit_frequency = 16u;

const CommandPoolConfig cmd_config{
const vkapi::CommandPoolConfig cmd_config{
32u, // cmdPoolInitialSize
8u, // cmdPoolBatchSize
};

const DescriptorPoolConfig descriptor_pool_config{
const vkapi::DescriptorPoolConfig descriptor_pool_config{
VULKAN_DESCRIPTOR_POOL_SIZE, // descriptorPoolMaxSets
VULKAN_DESCRIPTOR_POOL_SIZE, // descriptorUniformBufferCount
VULKAN_DESCRIPTOR_POOL_SIZE, // descriptorStorageBufferCount
Expand All @@ -194,7 +195,7 @@ Context* context() {
32u, // descriptorPileSizes
};

const QueryPoolConfig query_pool_config{
const vkapi::QueryPoolConfig query_pool_config{
VULKAN_QUERY_POOL_SIZE, // maxQueryCount
256u, // initialReserveSize
};
Expand All @@ -206,7 +207,7 @@ Context* context() {
query_pool_config,
};

return new Context(runtime()->default_adapter_i(), config);
return new Context(vkapi::runtime()->default_adapter_i(), config);
} catch (...) {
}

Expand Down
Loading