|
bf9f69fc
|
2025-02-14T11:57:14
|
|
CL/Vulkan: Add support for LiteralSampler
The support for LiteralSampler in clspv is enabled through the SPIR-V
non-sementic clspv reflection instructions - LiteralSampler [1].
Setting the utility functions to process
these instructions.
Bug: angleproject:394767574
Change-Id: I47aa03dac050f723d2e3f2a7f9c6c1756e7ac9cf
Reviewed-on: https://chromium-review.googlesource.com/c/angle/angle/+/6239519
Commit-Queue: Austin Annestrand <a.annestrand@samsung.com>
Reviewed-by: Shahbaz Youssefi <syoussefi@chromium.org>
Reviewed-by: Austin Annestrand <a.annestrand@samsung.com>
|
|
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>
|
|
2e36e947
|
2025-01-16T15:12:22
|
|
CL/VK: kernel's arg.handle update properly when setArg
Fix arg.handle to be an underlying CL object type.
Bug: angleproject:390488741
Change-Id: I8a9d3981a46f17e9af906f0e345c17585a332b30
Reviewed-on: https://chromium-review.googlesource.com/c/angle/angle/+/6180553
Reviewed-by: Shahbaz Youssefi <syoussefi@chromium.org>
Reviewed-by: Austin Annestrand <a.annestrand@samsung.com>
Commit-Queue: Austin Annestrand <a.annestrand@samsung.com>
Reviewed-by: Geoff Lang <geofflang@chromium.org>
|
|
|
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>
|
|
6c1021ec
|
2024-11-22T16:48:45
|
|
Vulkan: Switch DescriptorSetLayout to use AtomicSharedPtr
SharedPtr has better semantics and safer to use. This CL removes direct
exposure of RefCounted object and also allows me to delete
BindingPointer class in later CL.
Bug: angleproject:372268711
Change-Id: I08a0dff3efcf794be843a4a548b9f2609bb9a5e1
Reviewed-on: https://chromium-review.googlesource.com/c/angle/angle/+/6044328
Commit-Queue: Shahbaz Youssefi <syoussefi@chromium.org>
Reviewed-by: Yuxin Hu <yuxinhu@google.com>
Reviewed-by: Shahbaz Youssefi <syoussefi@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>
|
|
|
a58b35bc
|
2024-08-07T15:01:56
|
|
CL/Vulkan: Implement image creation from buffer object
Add support for creation of image from buffer object for types
image1d_buffer. At the kernel side setup texel buffer descriptor sets
for these kernel arguments.
Bug: angleproject:378103913
Change-Id: I600692cd003b75396afd45dcc93c568bcf390b96
Signed-off-by: Gowtham Tammana <g.tammana@samsung.com>
Reviewed-on: https://chromium-review.googlesource.com/c/angle/angle/+/6005389
Reviewed-by: Shahbaz Youssefi <syoussefi@chromium.org>
Reviewed-by: Geoff Lang <geofflang@chromium.org>
|
|
c0a28403
|
2024-05-06T12:42:28
|
|
CL/Vulkan: Enable clEnqueueNDRangeKernel for Images and Samplers
* Performs clspv reflection for ArgumentSampler, ArgumentStorageImage,
and ArgumentSampledImage
* Creates descriptor sets for performing enqueue
Tests-Passing: OCLCTS.test_basic readimage, readimage_fp32,
readimage_int16, writeimage, writeimage_fp32, writeimage_int16,
mri_one, mri_multiple, imagenpot
Bug: angleproject:42266936
Change-Id: I2b6c631e76556870c4342d2046c267ff5cf5105d
Reviewed-on: https://chromium-review.googlesource.com/c/angle/angle/+/5955597
Reviewed-by: Shahbaz Youssefi <syoussefi@chromium.org>
Reviewed-by: Geoff Lang <geofflang@chromium.org>
Commit-Queue: Shahbaz Youssefi <syoussefi@chromium.org>
|
|
|
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>
|
|
|
fc65058c
|
2024-05-03T14:47:29
|
|
CL/Vulkan: Rework spec constant handling
Specialization constants can be combined into
a single hash map that's later iterated on
when we create the compute pipeline.
Bug: angleproject:364396920
Change-Id: I161356808ff0bd6a589f01854264210011bec512
Reviewed-on: https://chromium-review.googlesource.com/c/angle/angle/+/5834664
Reviewed-by: Geoff Lang <geofflang@chromium.org>
Commit-Queue: Austin Annestrand <a.annestrand@samsung.com>
Reviewed-by: Shahbaz Youssefi <syoussefi@chromium.org>
|
|
d193d51b
|
2024-06-17T22:46:08
|
|
Replace issue ids post migration to new issue tracker
This change replaces anglebug.com/NNNN links.
Bug: None
Change-Id: I8ac3aec8d2a8a844b3d7b99fc0a6b2be8da31761
Reviewed-on: https://chromium-review.googlesource.com/c/angle/angle/+/5637912
Reviewed-by: Geoff Lang <geofflang@chromium.org>
Commit-Queue: Shahbaz Youssefi <syoussefi@chromium.org>
|
|
5703bd61
|
2024-06-14T14:12:41
|
|
Vulkan: Further optimize ProgramExecutableVk::resetLayout
1. Handle compute pipelines similar to how we handle graphics pipelines
2. Track valid compute pipeline permutations
Bug: angleproject:8297
Change-Id: I58200517e5a44a2b3092777ea24d1529ceee00f5
Reviewed-on: https://chromium-review.googlesource.com/c/angle/angle/+/5634574
Reviewed-by: Charlie Lao <cclao@google.com>
Commit-Queue: mohan maiya <m.maiya@samsung.com>
Reviewed-by: Shahbaz Youssefi <syoussefi@chromium.org>
|
|
|
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>
|
|
|
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>
|