|
7c7fcb92
|
2025-07-14T07:56:22
|
|
CL/Vulkan: Allocate descriptor sets under lock
Multiple threads in the application are accessing the OpenCL
runtime simultaneously, resulting in multiple SharedPtr of
mDynamicDescriptorPools[] being created for a single context.
This can lead to incorrect reference counting and
potentially cause `DynamicDescriptorPool::destroy` to be
called unexpectedly. This fixes the concurrent access issue
through acquiring the descriptor set mutex
just before initializing mDynamicDescriptorPools to ensure that
only one mDynamicDescriptorPools[] is created per context.
Bug: angleproject:383999367
Tests-Passing: test_integer_ops int_logic
Change-Id: Iba0fad6813a08e1631b73d5efae4f4639892b36f
Signed-off-by: hoonee.cho <hoonee.cho@samsung.com>
Reviewed-on: https://chromium-review.googlesource.com/c/angle/angle/+/6981091
Reviewed-by: Shahbaz Youssefi <syoussefi@chromium.org>
Reviewed-by: Geoff Lang <geofflang@chromium.org>
Commit-Queue: Austin Annestrand <a.annestrand@samsung.com>
|
|
|
1df3b59f
|
2024-10-30T16:40:51
|
|
CL/VK: PrintfBufferPointerPushConstant support
Bug: angleproject:442950569
Change-Id: I64edba5a04c2f1f2d3eb7bb34e6629f12269a69c
Signed-off-by: hoonee.cho <hoonee.cho@samsung.com>
Reviewed-on: https://chromium-review.googlesource.com/c/angle/angle/+/6916342
Commit-Queue: Austin Annestrand <a.annestrand@samsung.com>
Reviewed-by: Geoff Lang <geofflang@chromium.org>
Reviewed-by: Shahbaz Youssefi <syoussefi@chromium.org>
|
|
ce158355
|
2025-05-06T17:11:33
|
|
CL/VK: Add isReadOnly() query for CLKernelArgument
A new helper function `isReadOnly()` is added to CLKernelArgument. This
makes the determination of memory dependencies more flexible.
As a result, updated the `addMemoryDependencies` function to take in the
write usage boolean.
Bug: angleproject:441240590
Change-Id: I8826f88b7ff84ee7c689a68df4c5121ee9f12619
Signed-off-by: Gowtham Tammana <g.tammana@samsung.com>
Reviewed-on: https://chromium-review.googlesource.com/c/angle/angle/+/6885850
Reviewed-by: Geoff Lang <geofflang@chromium.org>
Reviewed-by: Shahbaz Youssefi <syoussefi@chromium.org>
Commit-Queue: Austin Annestrand <a.annestrand@samsung.com>
|
|
239763d9
|
2025-04-04T17:52:10
|
|
CL/VK: Fix ArgumentWorkgroup setArg logic
Issue was that clSetKernelArg for ArgumentWorkgroup
type arguments would blindly push new values into
kernel's spec-constant FastVector on every clSetKernelArg
(even on same arg updates).
This would lead to over-pushing due to same arg updates,
which caused all kinds of issues, mainly erroneous misses
in compute pipeline cache since the key is based on
VkSpecializationInfo.
Since kernel object already keeps a vector of kernel args,
we don't need a separate spec-constant FastVector in CLKernelVk to
track this. Remove it and derive the spec-constant data
from the kernel args themselves.
Bug: angleproject:366415134
Tests-Passing: Geekbench-6.2.2 - Workloads: [ 401 & 601 ]
Change-Id: Iab7f27fdfdfede33881e1dd717ba3b771cffb985
Reviewed-on: https://chromium-review.googlesource.com/c/angle/angle/+/6773615
Reviewed-by: Shahbaz Youssefi <syoussefi@chromium.org>
Reviewed-by: Geoff Lang <geofflang@chromium.org>
Commit-Queue: Shahbaz Youssefi <syoussefi@chromium.org>
|
|
|
bab3f9ab
|
2025-01-30T14:30:16
|
|
CL/VK: Cosmetic kernel arg updates
- Move POD argument buffer init to kernel-init
- Remove unused CLKernelVK members/types
- Switch inside of setArg loop
- Rename kernel members that best aligns to their resources
Bug: angleproject:42267001
Change-Id: I837ceeceaf1ff903b67b7f100298d4a3159b97d7
Reviewed-on: https://chromium-review.googlesource.com/c/angle/angle/+/6220895
Reviewed-by: Shahbaz Youssefi <syoussefi@chromium.org>
Commit-Queue: Austin Annestrand <a.annestrand@samsung.com>
Reviewed-by: Geoff Lang <geofflang@chromium.org>
|
|
|
e0cbdbb5
|
2025-01-15T14:08:45
|
|
CL/VK: Enable dispatch region chunking
This change does two things:
- Allow non-uniform workgroup size
- Dispatches that go beyond VK workgroup count limit
Bug: angleproject:42267085
Change-Id: I1f2da93252e466b811273ee34d9d38e454f11686
Reviewed-on: https://chromium-review.googlesource.com/c/angle/angle/+/6180550
Reviewed-by: Shahbaz Youssefi <syoussefi@chromium.org>
Commit-Queue: Austin Annestrand <a.annestrand@samsung.com>
Reviewed-by: Geoff Lang <geofflang@chromium.org>
|
|
|
95635ef0
|
2025-01-23T16:30:41
|
|
CL/VK: Implementation of Compute Pipeline Cache.
Implemented ComputePipelineCache, hash map from OpenCL and
OpenGL compute state vectors to compiled pipelines. Implemented
ComputePipelineDesc, a tightly packed description of the current
compute state.
Compute Pipeline State includes the specialization constants,
Pipeline Options (Protected, Robust).
Updated-by: Austin Annestrand <a.annestrand@samsung.com>
Bug: angleproject:391672281
Change-Id: I88944dc169d194d1b2c75747769d7346b041fa75
Reviewed-on: https://chromium-review.googlesource.com/c/angle/angle/+/6191437
Reviewed-by: Shahbaz Youssefi <syoussefi@chromium.org>
Reviewed-by: Charlie Lao <cclao@google.com>
Commit-Queue: Austin Annestrand <a.annestrand@samsung.com>
|
|
|
b5c12605
|
2025-01-22T13:44:13
|
|
CL: Limit max work size/offset to UINT32_MAX
Set a hard limit of UINT32_MAX on work item/offset for OpenCL's
clEnqueueNDRangeKernel cmd for all backends (simplifies handling).
Bug: angleproject:42267067
Change-Id: Ied1685609b3818e4c3a366a03770dc361198639c
Reviewed-on: https://chromium-review.googlesource.com/c/angle/angle/+/6191436
Reviewed-by: Shahbaz Youssefi <syoussefi@chromium.org>
Commit-Queue: Austin Annestrand <a.annestrand@samsung.com>
|
|
810c19fa
|
2025-01-02T15:19:51
|
|
CL/VK: NonSemanticClspvReflectionArgumentPodUniform support
Missing NonSemanticClspvReflectionArgumentPodUniform support,
which gets used when the number of POD arguments exceed the
limit of the maximum amount of push constants that can be used.
Bug: angleproject:384549637
Change-Id: Ia0e2ea141cfb2c41c492ff4b125a1547c380c1cb
Reviewed-on: https://chromium-review.googlesource.com/c/angle/angle/+/6097413
Reviewed-by: Shahbaz Youssefi <syoussefi@chromium.org>
Reviewed-by: Geoff Lang <geofflang@chromium.org>
Commit-Queue: Austin Annestrand <a.annestrand@samsung.com>
|
|
|
3772d47e
|
2024-12-11T00:25:37
|
|
CL/Vulkan: Allocate descriptor sets under lock
The descriptor pool must be externally synchronized as per the vulkan
spec, as such acquire a lock when allocating descriptor sets from it.
Bug: angleproject:383999367
Change-Id: I98448770681ad39dae0dc7e413e28ec7dfa89f87
Signed-off-by: Gowtham Tammana <g.tammana@samsung.com>
Reviewed-on: https://chromium-review.googlesource.com/c/angle/angle/+/6099129
Reviewed-by: Shahbaz Youssefi <syoussefi@chromium.org>
Reviewed-by: Geoff Lang <geofflang@chromium.org>
|
|
2dc072ec
|
2024-11-22T16:14:52
|
|
Vulkan: Switch PipelineLayout from AtomicBind* to AtomicSharedPtr
AtomicSharedPtr/SharedPtr has better semantics and safer to use. This
will allow deleting BindingPointer in later CL.
Bug: angleproject:372268711
Change-Id: Ife20f68b2277a1913b06be0de153770214ac964a
Reviewed-on: https://chromium-review.googlesource.com/c/angle/angle/+/6044326
Reviewed-by: Shahbaz Youssefi <syoussefi@chromium.org>
Reviewed-by: Yuxin Hu <yuxinhu@google.com>
Commit-Queue: Charlie Lao <cclao@google.com>
|
|
|
5b4609de
|
2024-10-07T14:40:49
|
|
CL/Vulkan: Adjust the pushConstant size/offset to multple of 4
The spec[1] requires the push constants size/offset to be multiple of 4.
Adjust them as needed.
[1]: https://registry.khronos.org/vulkan/specs/1.3-extensions/man/html/vkCmdPushConstants.html#_description
Bug: angleproject:372157565
Change-Id: I8ea788dbd68e3aea262e12af56e40ac84087ceef
Signed-off-by: Gowtham Tammana <g.tammana@samsung.com>
Reviewed-on: https://chromium-review.googlesource.com/c/angle/angle/+/5916154
Reviewed-by: Shahbaz Youssefi <syoussefi@chromium.org>
Reviewed-by: Geoff Lang <geofflang@chromium.org>
|
|
|
dd54eeec
|
2024-10-11T13:26:46
|
|
Reland "Vulkan: Track GPU progress for individual DescriptorSet"
This is a reland of commit 292102944add2ab30f4aa12a971cac456cc7726b
with the fix of garbage being added back to garbage list.
Original change's description:
> Vulkan: Track GPU progress for individual DescriptorSet
>
> Right now ProgramExecutableVk keeps VkDescriptorSet object, and
> DescriptorSetHelper is created when a cache entry becomes invalid.
> Further, DescriptorSetCache keeps the cache of {VkDescriptorSet,
> RefCountedDescriptorPoolHelper} pair. So we are having three different
> type of objects at different stages of life: VkDescriptorSet,
> DescriptorSetHelper, and {VkDescriptorSet,
> RefCountedDescriptorPoolHelper. This CL makes DescriptorSetHelper at
> creation and at cache and at garbage. With this change, you have a
> reference counted DescriptorSetHelper object (i.e, DescriptorSetPointer)
> during entire life cycle and is passed around between cache and program
> as is. This CL is preparation for the future CL where we may disable
> cache for descriptorSet. The descriptorSet will be added to garbage list
> and reused constantly without go through the cache code. We need to
> track the individual descriptorSet with ResourceUse so that it won't
> reuse until GPU is finished. This CL is making DescriptorSetHelper a GPU
> tracking object so that it will still just work when cache is disabled.
>
> Bug: angleproject:372268711
> Change-Id: I1cfb77cc5069b202d870388fd8809e265cdca90b
> Reviewed-on: https://chromium-review.googlesource.com/c/angle/angle/+/5918586
> Reviewed-by: Shahbaz Youssefi <syoussefi@chromium.org>
> Commit-Queue: Charlie Lao <cclao@google.com>
> Reviewed-by: Yuxin Hu <yuxinhu@google.com>
Bug: angleproject:372268711
Change-Id: Ic920f99cc78cde1e94690bdbee3b885844fa155b
Reviewed-on: https://chromium-review.googlesource.com/c/angle/angle/+/5954701
Commit-Queue: Charlie Lao <cclao@google.com>
Reviewed-by: Yuxin Hu <yuxinhu@google.com>
|
|
|
45cc47af
|
2024-10-22T21:41:22
|
|
Revert "Vulkan: Track GPU progress for individual DescriptorSet"
This reverts commit 292102944add2ab30f4aa12a971cac456cc7726b.
Reason for revert: Causing bot failure in later CLs
Original change's description:
> Vulkan: Track GPU progress for individual DescriptorSet
>
> Right now ProgramExecutableVk keeps VkDescriptorSet object, and
> DescriptorSetHelper is created when a cache entry becomes invalid.
> Further, DescriptorSetCache keeps the cache of {VkDescriptorSet,
> RefCountedDescriptorPoolHelper} pair. So we are having three different
> type of objects at different stages of life: VkDescriptorSet,
> DescriptorSetHelper, and {VkDescriptorSet,
> RefCountedDescriptorPoolHelper. This CL makes DescriptorSetHelper at
> creation and at cache and at garbage. With this change, you have a
> reference counted DescriptorSetHelper object (i.e, DescriptorSetPointer)
> during entire life cycle and is passed around between cache and program
> as is. This CL is preparation for the future CL where we may disable
> cache for descriptorSet. The descriptorSet will be added to garbage list
> and reused constantly without go through the cache code. We need to
> track the individual descriptorSet with ResourceUse so that it won't
> reuse until GPU is finished. This CL is making DescriptorSetHelper a GPU
> tracking object so that it will still just work when cache is disabled.
>
> Bug: angleproject:372268711
> Change-Id: I1cfb77cc5069b202d870388fd8809e265cdca90b
> Reviewed-on: https://chromium-review.googlesource.com/c/angle/angle/+/5918586
> Reviewed-by: Shahbaz Youssefi <syoussefi@chromium.org>
> Commit-Queue: Charlie Lao <cclao@google.com>
> Reviewed-by: Yuxin Hu <yuxinhu@google.com>
Bug: angleproject:372268711
Change-Id: I4d3c34058d100112a098144276b52c0faf8d593a
No-Presubmit: true
No-Tree-Checks: true
No-Try: true
Reviewed-on: https://chromium-review.googlesource.com/c/angle/angle/+/5955529
Auto-Submit: Charlie Lao <cclao@google.com>
Commit-Queue: Rubber Stamper <rubber-stamper@appspot.gserviceaccount.com>
Bot-Commit: Rubber Stamper <rubber-stamper@appspot.gserviceaccount.com>
|
|
|
29210294
|
2024-10-11T13:26:46
|
|
Vulkan: Track GPU progress for individual DescriptorSet
Right now ProgramExecutableVk keeps VkDescriptorSet object, and
DescriptorSetHelper is created when a cache entry becomes invalid.
Further, DescriptorSetCache keeps the cache of {VkDescriptorSet,
RefCountedDescriptorPoolHelper} pair. So we are having three different
type of objects at different stages of life: VkDescriptorSet,
DescriptorSetHelper, and {VkDescriptorSet,
RefCountedDescriptorPoolHelper. This CL makes DescriptorSetHelper at
creation and at cache and at garbage. With this change, you have a
reference counted DescriptorSetHelper object (i.e, DescriptorSetPointer)
during entire life cycle and is passed around between cache and program
as is. This CL is preparation for the future CL where we may disable
cache for descriptorSet. The descriptorSet will be added to garbage list
and reused constantly without go through the cache code. We need to
track the individual descriptorSet with ResourceUse so that it won't
reuse until GPU is finished. This CL is making DescriptorSetHelper a GPU
tracking object so that it will still just work when cache is disabled.
Bug: angleproject:372268711
Change-Id: I1cfb77cc5069b202d870388fd8809e265cdca90b
Reviewed-on: https://chromium-review.googlesource.com/c/angle/angle/+/5918586
Reviewed-by: Shahbaz Youssefi <syoussefi@chromium.org>
Commit-Queue: Charlie Lao <cclao@google.com>
Reviewed-by: Yuxin Hu <yuxinhu@google.com>
|
|
|
f9709279
|
2024-09-20T16:02:56
|
|
CL/Vulkan: Add support for printf builtin processing
The support for printf builtin in clspv enabled through the SPIR-V
non-semantic clspv reflection instructions - PrintfInfo and
PrintfBufferStorageBuffer [1]. The printf buffer is setup with a
separate descriptor layout and the pipeline layout is updated
accordingly.
Also, printf is enabled as default option now for clspv.
[1]: https://github.com/KhronosGroup/SPIRV-Registry/blob/master/nonsemantic/NonSemantic.ClspvReflection.html
Bug: angleproject:369724757
Change-Id: I20b245eb0fea69941bd1aeb42534f8b729ec17e8
Signed-off-by: Gowtham Tammana <g.tammana@samsung.com>
Reviewed-on: https://chromium-review.googlesource.com/c/angle/angle/+/5893958
Reviewed-by: Charlie Lao <cclao@google.com>
Reviewed-by: Shahbaz Youssefi <syoussefi@chromium.org>
|
|
|
0a452697
|
2024-09-19T16:33:26
|
|
CL/Vulkan: Enable support for multiple descriptor set handling
The clspv transcompiler can generate multiple descriptor sets depending
on the kernel e.g. printf storage buffer is setup in a separate
descriptor set, and kernel arguments in a separate descriptor set. This
change enables setting up of multiple descriptor sets and appropriately
bind them.
Also, creation of descriptor set and pipeline layout are deferred to
kernel enqueue time as it is quite possible that kernels are built in a
batch before their use.
Bug: angleproject:369724757
Change-Id: I80eb93f4a3b8afc7461c299cc283526cc4b872fa
Signed-off-by: Gowtham Tammana <g.tammana@samsung.com>
Reviewed-on: https://chromium-review.googlesource.com/c/angle/angle/+/5893955
Reviewed-by: Charlie Lao <cclao@google.com>
Reviewed-by: Shahbaz Youssefi <syoussefi@chromium.org>
|
|
|
59f1e537
|
2024-09-03T17:56:47
|
|
CL/VK: Add ArgumentWorkgroup handling
Each OpenCL kernel can have
ArgumentWorkgroup args - which are
treated as spec constants.
Bug: angleproject:366415134
Change-Id: I2761010610e5b991e959006e7b78f1241a687960
Reviewed-on: https://chromium-review.googlesource.com/c/angle/angle/+/5860304
Commit-Queue: Austin Annestrand <a.annestrand@samsung.com>
Reviewed-by: Geoff Lang <geofflang@chromium.org>
Reviewed-by: Shahbaz Youssefi <syoussefi@chromium.org>
|
|
|
cc50c75b
|
2024-04-19T16:40:26
|
|
CL/Vulkan: Fix missing kernel enqueue retain/release
According to spec:
https://registry.khronos.org/OpenCL/specs/3.0-unified/html/OpenCL_API.html#clReleaseKernel
We also need to implicitly retain/release the Kernel object
on Kernel enqueue APIs along with other kernel resource
retain/release(s) (i.e. memobjs).
Bug: angleproject:356328473
Change-Id: Ie6ca228ebece2988d511d9ffd2a617abdc057d0a
Reviewed-on: https://chromium-review.googlesource.com/c/angle/angle/+/5753766
Reviewed-by: Shahbaz Youssefi <syoussefi@chromium.org>
Reviewed-by: Geoff Lang <geofflang@chromium.org>
Commit-Queue: Austin Annestrand <a.annestrand@samsung.com>
|
|
|
d4abe622
|
2024-04-03T17:46:38
|
|
CL/VK: Implement enqueue NDRangeKernel & Task
Adding support for:
clEnqueueNDRangeKernel
clEnqueueTask
Bug: angleproject:8631
Change-Id: If57002be3ea00a55215e89ca47ab8fe9a422c6e7
Reviewed-on: https://chromium-review.googlesource.com/c/angle/angle/+/5406614
Reviewed-by: Shahbaz Youssefi <syoussefi@chromium.org>
Commit-Queue: Austin Annestrand <a.annestrand@samsung.com>
Reviewed-by: Geoff Lang <geofflang@chromium.org>
|
|
|
76636ddb
|
2024-04-04T15:39:08
|
|
CL/VK: Update missing reflection ops & DS creation
Adding CLProgramVk routine to allocate DS.
Also adding reflection parsing support for
WGS IDs and PushConstantGlobalSize.
Additionally, we now check VK implementation
if non-semantic reflection data (in SPIR-V) is
supported. If not, we strip that data from binary
(via vk_feature check).
Bug: angleproject:8631
Change-Id: Ife02867c7c30b919abf663865adc92858e1bff8d
Reviewed-on: https://chromium-review.googlesource.com/c/angle/angle/+/5421574
Reviewed-by: Geoff Lang <geofflang@chromium.org>
Reviewed-by: Shahbaz Youssefi <syoussefi@chromium.org>
Commit-Queue: Austin Annestrand <a.annestrand@samsung.com>
|
|
|
9fee9156
|
2024-01-18T13:50:38
|
|
OpenCL/Vulkan: Add initial CL Kernel routines
Enables the following OpenCL APIs for Vulkan Backend:
clCreateKernel
clCreateKernels
clSetKernelArg
Bug: angleproject:8572
Change-Id: Iaaabd295d7951a48bf9bd6717ed3957960052dbd
Reviewed-on: https://chromium-review.googlesource.com/c/angle/angle/+/5341376
Reviewed-by: Geoff Lang <geofflang@chromium.org>
Reviewed-by: Shahbaz Youssefi <syoussefi@chromium.org>
Commit-Queue: Austin Annestrand <a.annestrand@samsung.com>
|
|
|
1ceddbf6
|
2024-01-11T16:15:08
|
|
OpenCL/Vulkan: Add createProgram routines
Introduce createProgramWithSource and
createProgramWithBinary.
Additionally introduce KernelArgument types,
SPIR-V Reflection info types, and DeviceProgram types.
Bug: angleproject:8549
Change-Id: I611627d747a5ba6718778fd27fd28477b77fbbe1
Reviewed-on: https://chromium-review.googlesource.com/c/angle/angle/+/5303563
Commit-Queue: Austin Annestrand <a.annestrand@samsung.com>
Reviewed-by: Geoff Lang <geofflang@chromium.org>
|
|
|
66d29149
|
2023-10-16T12:43:43
|
|
OpenCL: Add scaffolding to renderer/vulkan code
Start off ANCLE (OpenCL on ANGLE) with backend scaffolding code with
placeholder API error returns until they are later implemented.
Bug: angleproject:8377
Change-Id: I223d3482cce097ccb298e121fc03ec416e8958fd
Reviewed-on: https://chromium-review.googlesource.com/c/angle/angle/+/4950556
Reviewed-by: Shahbaz Youssefi <syoussefi@chromium.org>
Commit-Queue: Shahbaz Youssefi <syoussefi@chromium.org>
|
|
9459456b
|
2021-05-27T13:43:41
|
|
CL: kernel creation for front end & passthrough
Add kernel object to back end and implement creation and info query.
Bug: angleproject:6001
Change-Id: I6e3fdd2b35d7e73ed56144dc938bf21436d3559f
Reviewed-on: https://chromium-review.googlesource.com/c/angle/angle/+/2922150
Commit-Queue: John Plate <jplate@google.com>
Reviewed-by: Jamie Madill <jmadill@chromium.org>
Reviewed-by: Cody Northrop <cnorthrop@google.com>
|