aa82808645
printf: extract clovers printf impl
...
Also make the code cleaner and simplier.
Signed-off-by: Karol Herbst <kherbst@redhat.com >
Acked-by: Jesse Natalie <jenatali@microsoft.com >
Reviewed-by: Jason Ekstrand <jason.ekstrand@collabora.com >
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/17334 >
2022-08-04 23:53:49 +00:00
0a4c0bc0dd
clover: Set images/samplers_used when lowering images
...
Also, stop using BITSET_SET_RANGE_INSIDE_WORD for textures so we can
handle more than 32 of them.
Reviewed-by: Karol Herbst <kherbst@redhat.com >
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/16435 >
2022-05-12 01:32:58 +00:00
9c5fd100cc
nir: add a nir_remove_non_entrypoints helper
...
This code just got duplicated a lot. There is still more, but the
remaining instances do a bit more than just removing other functions.
Signed-off-by: Karol Herbst <kherbst@redhat.com >
Reviewed-by: Alyssa Rosenzweig <alyssa@collabora.com >
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/16348 >
2022-05-10 03:37:44 +00:00
f452317849
clover/nir: respect lower to scalar options.
...
This just calls the lower alu to scalar pass like mesa/st
Reviewed-by: Karol Herbst <kherbst@redhat.com >
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/15433 >
2022-03-17 22:00:49 +00:00
99cda38c81
clover/nir: Don't remove texture variables
...
Reviewed-by: Jesse Natalie <jenatali@microsoft.com >
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/13389 >
2021-10-16 05:49:34 +00:00
d68bedbb45
clover: Use nir_foreach_image_variable for images
...
This splits image and sampler handling into two separate loops.
Reviewed-by: Caio Marcelo de Oliveira Filho <caio.oliveira@intel.com >
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/4743 >
2021-10-15 14:58:56 +00:00
aefa22ddb5
clover: Insert dummy uniform variables for images
...
Instead of making images have a well-defined size, insert a dummy
variable of the appropriate type which we can use for the parameter
block layout. This will work much better when we switch over to
nir_var_mem_image.
Acked-by: Caio Marcelo de Oliveira Filho <caio.oliveira@intel.com >
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/4743 >
2021-10-15 14:58:56 +00:00
1f880a2ea8
clover: Rename module -> binary, because C++20 makes module a keyword
...
Reviewed-by: Karol Herbst <kherbst@redhat.com >
Reviewed-by: Francisco Jerez <currojerez@riseup.net >
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/12273 >
2021-09-27 18:50:09 +00:00
3d65cea6ee
util/bitset: s/BITSET_SET_RANGE/BITSET_SET_RANGE_INSIDE_WORD
...
Prep work for the next commit.
Signed-off-by: Christian Gmeiner <christian.gmeiner@gmail.com >
Reviewed-by: Rob Clark <robdclark@chromium.org >
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/11321 >
2021-09-21 20:25:31 +00:00
6619877bdf
clover: Don't remove sampler/image uniforms
...
Otherwise we fool nir_shader_gather_info() into telling us there are no
samplers/images.
Signed-off-by: Rob Clark <robdclark@chromium.org >
Reviewed-by: Karol Herbst <kherbst@redhat.com >
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/12733 >
2021-09-05 18:14:08 +00:00
353e632393
clover: add kernel attributes support for SPIR-V
...
Fixes CTS api kernel_attributes
Reviewed-by: Karol Herbst <kherbst@redhat.com >
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/12225 >
2021-08-09 19:16:29 +00:00
b4e5bf0637
clover/nir: Set constant buffer pointer size to host
...
The `argument::size` is supposed to represent the size of a pointer on
the host and not on the device (for which argument::target_size`
exists).
v3: Use `sizeof(buf)` instead of `marg.size`. (Francisco Jerez)
Fixes: 7c6f1d3bf9currojerez@riseup.net >
Signed-off-by: Pierre Moreau <dev@pmoreau.org >
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/10256 >
2021-08-03 16:43:49 +00:00
8af6766062
nir: Move workgroup_size and workgroup_variable_size into common shader_info
...
Move it out the "cs" sub-struct, since these will be used for other
shader stages in the future.
Reviewed-by: Kenneth Graunke <kenneth@whitecape.org >
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/11225 >
2021-06-08 09:23:55 -07:00
430d2206da
compiler: Rename local_size to workgroup_size
...
Acked-by: Emma Anholt <emma@anholt.net >
Acked-by: Alyssa Rosenzweig <alyssa.rosenzweig@collabora.com >
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net >
Acked-by: Timur Kristóf <timur.kristof@gmail.com >
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/11190 >
2021-06-07 22:34:42 +00:00
fa677c8644
nir_lower_readonly_images_to_tex: Support non-CL semantics
...
For non-CL, intrinsic access isn't set, because the image type doesn't
have access qualifier. Instead, the access qualifier is set on the variable.
So, add a mode to this pass which can chase back to the variable in addition
to the intrinsic access. Also, update the variable type and the deref chain
types so everything is consistent, that the tex is accessing a sampler. Note
we can't do this for CL, because void-typed samplers don't exist.
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net >
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/10356 >
2021-04-23 23:16:15 +00:00
29c9731400
nir: Rename nir_lower_cl_images_to_tex, replace 'cl' with 'readonly'
...
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net >
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/10356 >
2021-04-23 23:16:15 +00:00
8027a7ba8a
shader_info: convert textures_used to a bitset.
...
For now keep it a bitset of 1 32-bit dword.
Reviewed-by: Eric Anholt <eric@anholt.net >
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/9456 >
2021-03-10 06:16:09 +10:00
117668b811
nir: Make nir_ssa_def_rewrite_uses take an SSA value
...
This commit replaces the new_src parameter of nir_ssa_def_rewrite_uses()
with an SSA def, removes nir_ssa_def_rewrite_uses_ssa(), and rewrites
all the users as needed.
Reviewed-by: Rhys Perry <pendingchaos02@gmail.com >
Acked-by: Alyssa Rosenzweig <alyssa@collabora.com >
Acked-By: Mike Blumenkrantz <michael.blumenkrantz@gmail.com >
Reviewed-by: Eric Anholt <eric@anholt.net >
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/9383 >
2021-03-08 16:59:55 +00:00
76788353b2
clover/nir: hookup printf (v3)
...
This connects printf up for NIR drivers, it lowers using the NIR
pass where it places the idx to the strings into the output buffer.
It also sets the global buffer header to the nir paths.
v2: remove dead function temps after lowering
v3: move to single string
Acked-by: Jesse Natalie <jenatali@microsoft.com >
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/8254 >
2020-12-29 09:16:52 +10:00
c730ace12b
nir,clover: Drop nir_lower_mem_constant_vars
...
We have a more generic helper now so clover doesn't need quite as many
special paths.
Reviewed-by: Jesse Natalie <jenatali@microsoft.com >
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/7565 >
2020-11-18 04:05:37 +00:00
1a775b71ca
clover/nir: set kernel_image cap
...
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net >
Reviewed-by: Dave Airlie <airlied@redhat.com >
Reviewed-by: Francisco Jerez <currojerez@riseup.net >
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/7069 >
2020-10-20 23:46:42 +02:00
4f24dee22a
clover/nir: Add an image lowering pass
...
Reviewed-by: Dave Airlie <airlied@redhat.com >
Reviewed-by: Karol Herbst <kherbst@redhat.com >
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/7069 >
2020-10-20 23:46:42 +02:00
5e31fad8c9
clover/nir: Calculate sizes of images and samplers properly
...
Clover uses very specific sizes and alignments for images and samplers
to pass various bits of data. We need to add a new size/align helper
for inputs which matches the standard CL size/align for most types but
also has the right size/align for images and samplers.
v2 (Karol): use sizeof(cl_mem) instead of 8 to fix 32 bit runtimes.
Reviewed-by: Dave Airlie <airlied@redhat.com >
Reviewed-by: Karol Herbst <kherbst@redhat.com >
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/7069 >
2020-10-20 23:46:42 +02:00
7e55f0e17d
clover/nir: add a constant folding pass before lowering mem const
...
If we lower mem constants first, then direct array accesses to
constants never get lowered, so do a constant fold pass first to
remove direct const array accesses.
Reviewed-by: Karol Herbst <kherbst@redhat.com >
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/7209 >
2020-10-19 19:32:22 +00:00
54eae33558
clover: Stop leaking NIR shaders
...
Reviewed-by: Karol Herbst <kherbst@redhat.com >
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/7068 >
2020-10-08 16:01:38 +00:00
2fa7c79045
spirv: Move nir_lower_libclc to src/compiler/spirv
...
This puts it in a shared place where everyone can get at it.
Reviewed-by: Jesse Natalie <jenatali@microsoft.com >
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/7034 >
2020-10-07 21:52:04 +00:00
43390a546d
clover: Use core libclc loader
...
v2 (Jason Ekstrand):
- Use the newly added nir_can_find_libclc() helper
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net >
Reviewed-by: Jesse Natalie <jenatali@microsoft.com >
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/7034 >
2020-10-07 21:52:04 +00:00
ee5b46fcfd
clover/spirv: support CL_KERNEL_COMPILE_WORK_GROUP_SIZE
...
Reviewed-by: Serge Martin <edb@sigluy.net >
Reviewed-by: Francisco Jerez <currojerez@riseup.net >
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/4974 >
2020-10-07 13:18:22 +00:00
c04d5e7efa
clover: implements clGetKernelWorkGroupInfo CL_KERNEL_COMPILE_WORK_GROUP_SIZE
...
Reviewed-by: Francisco Jerez <currojerez@riseup.net >
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/4974 >
2020-10-07 13:18:22 +00:00
aadd134081
clover: add CL_KERNEL_ATTRIBUTES for clGetKernelInfo
...
Reviewed-by: Francisco Jerez <currojerez@riseup.net >
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/4974 >
2020-10-07 13:18:22 +00:00
7c6f1d3bf9
clover/nir: extract constant buffer into its own section
...
Fixes test_basic constant_source
Signed-off-by: Karol Herbst <kherbst@redhat.com >
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net >
Reviewed-by: Francisco Jerez <currojerez@riseup.net >
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/6569 >
2020-10-06 17:30:41 +00:00
bc7ed03ef8
clover/nir: Call nir_lower_convert_alu_types
...
Reviewed-by: Jesse Natalie <jenatali@microsoft.com >
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/6945 >
2020-10-01 18:36:53 +00:00
bf80fb7c30
clover/nir: Call the memcpy lowering pass
...
Reviewed-by: Jesse Natalie <jenatali@microsoft.com >.
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/6713 >
2020-09-25 23:48:03 +00:00
f33b417652
clover: handle libclc shader (v3)
...
This works by taking the spirv produced by libclc which contains
a lot of mangled function entrypoints identified with LinkageAttribute decorations.
This patch just sets up clover to load the libclc blob and convert it to
library nir, and support inlining application nir with calls to libclc.
v2: Add a disk cache support for this object, to avoid the spirv parsing
overheads each time. move spirv->nir to lazy instantiation to avoid
the mess with glsl types and constructor ordering.
v3: make disk cache optional
v1-Reviewed-by: Jesse Natalie <jenatali@microsoft.com >
Reviewed-by: Karol Herbst <kherbst@redhat.com >
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/6035 >
2020-09-25 20:09:08 +00:00
796d3fe9e0
clover/nir: Use lower_vars_to_explicit for uniform and global
...
Reviewed-by: Jesse Natalie <jenatali@microsoft.com >
Reviewed-by: Karol Herbst <kherbst@redhat.com >
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/6472 >
2020-09-03 18:02:50 +00:00
8bea5aaa14
clover: Use args.size() to compute new var locations
...
This is better than using num_uniforms as it guarantees what we want: a
mapping from nir_variable to the args vector.
Reviewed-by: Karol Herbst <kherbst@redhat.com >
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/6472 >
2020-09-03 18:02:50 +00:00
bcfeead5f3
clover: Call nir_lower_mem_constant_vars
...
Fixes: 26a4c8f375kherbst@redhat.com >
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/6472 >
2020-09-03 18:02:50 +00:00
526f356633
clover: Use 64-bit offsets for shader_in on 64-bit GPUs
...
This really shouldn't matter as inputs should have logical pointers.
However, nir_builder defaults to building derefs based on the pointer
size in the shader_info. It's easier for now to just be consistent
everywhere.
Reviewed-by: Karol Herbst <kherbst@redhat.com >
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/6379 >
2020-09-01 20:50:04 +00:00
26a4c8f375
clover/nir: Use nir_var_mem_constant for __constant memory
...
Reviewed-by: Eric Anholt <eric@anholt.net >
Reviewed-by: Jesse Natalie <jenatali@microsoft.com >
Reviewed-by: Karol Herbst <kherbst@redhat.com >
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/6379 >
2020-09-01 20:50:04 +00:00
7dc39838ed
clover/nir: use offset for temp memory
...
Signed-off-by: Karol Herbst <kherbst@redhat.com >
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net >
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/6433 >
2020-09-01 18:47:30 +00:00
d421af3a99
clover/nir: Lower function_temp to scratch.
...
Signed-off-by: Karol Herbst <kherbst@redhat.com >
Reviewed-by: Jesse Natalie <jenatali@microsoft.com >
Reviewed-by: Boris Brezillon <boris.brezillon@collabora.com >
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net >
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/6433 >
2020-09-01 18:47:30 +00:00
70cbddc4a7
nir: use enum operator helper for nir_variable_mode and nir_metadata
...
those are used quite a bit
Signed-off-by: Karol Herbst <kherbst@redhat.com >
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net >
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/6520 >
2020-09-01 17:45:08 +00:00
865a2ad086
clover/nir/spirv: Use uniform rather than shader_in for kernel inputs
...
The semantics of inputs for CL are a closer match to the semantics of uniforms for graphics.
Rather than cross-stage data, it's data that every thread sees uniformly.
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net >
Reviewed-by: Karol Herbst <kherbst@redhat.com >
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/6497 >
2020-08-31 19:58:14 +00:00
4fd2a45267
clover/nir: add support for global invocation id offsets
...
v2: create variables only once
Signed-off-by: Karol Herbst <kherbst@redhat.com >
Acked-by: Jesse Natalie <jenatali@microsoft.com >
Acked-by: Jason Ekstrand <jason@jlekstrand.net >
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/5891 >
2020-08-21 22:07:05 +00:00
d3faac7a15
nir: Add options to nir_lower_compute_system_values to control compute ID base lowering
...
If no options are provided, existing intrinsics are used.
If the lowering pass indicates there should be offsets used for global
invocation ID or work group ID, then those instructions are lowered to
include the offset.
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net >
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/5891 >
2020-08-21 22:07:05 +00:00
2e1df6a17f
nir: Move compute system value lowering to a separate pass
...
The actual variable -> intrinsic lowering stays where it is, but
ops which convert one intrinsic to be implemented in terms of
another have moved.
Reviewed-by: Karol Herbst <kherbst@redhat.com >
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net >
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/5891 >
2020-08-21 22:07:05 +00:00
a624faeef9
clover/nir: Register callback for translation messages (v2)
...
This allows us to add SPIR-V to NIR translation failure messages to the
program’s compilation log, which can then be queried by the user.
v2: Replace the if-statement in `debug_function()` with an assert.
Signed-off-by: Pierre Moreau <dev@pmoreau.org >
Reviewed-by: Serge Martin <edb@sigluy.net >
Reviewed-by: Francisco Jerez <currojerez@riseup.net >
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/5038 >
2020-08-20 19:48:12 +00:00
8f7784ee8d
clover/nir: Use the correct address mode for shared
...
Shared memory needs to have 64-bit pointers but we want 32-bit offsets
most of the time. This is exactly what 32bit_offset_as_64bit is for.
Reviewed-by: Karol Herbst <kherbst@redhat.com >
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/6367 >
2020-08-19 18:11:42 +00:00
b2226f7a98
clover/nir: Stop computing the global address format twice
...
Reviewed-by: Karol Herbst <kherbst@redhat.com >
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/6367 >
2020-08-19 18:11:42 +00:00
884d2021d9
clover/nir: Stop setting ubo_addr_format
...
We unconditionally set constant_as_global = true so we should never get
UBO access out of spirv_to_nir.
Reviewed-by: Karol Herbst <kherbst@redhat.com >
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/6367 >
2020-08-19 18:11:42 +00:00
Karol Herbst