Mesa (main): nouveau/nir: Add support for pre-GF100 images and ssbos.
GitLab Mirror
gitlab-mirror at kemper.freedesktop.org
Fri Apr 29 23:13:23 UTC 2022
Module: Mesa
Branch: main
Commit: c228cb388956ce3acd8cec71e4de25b6e0a738dd
URL: http://cgit.freedesktop.org/mesa/mesa/commit/?id=c228cb388956ce3acd8cec71e4de25b6e0a738dd
Author: Emma Anholt <emma at anholt.net>
Date: Sun Apr 24 12:44:23 2022 -0700
nouveau/nir: Add support for pre-GF100 images and ssbos.
We have to allocate them slots in the global file.
Reviewed-by: Karol Herbst <kherbst at redhat.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/15949>
---
.../drivers/nouveau/codegen/nv50_ir_from_nir.cpp | 24 ++++++++++++++++++++++
1 file changed, 24 insertions(+)
diff --git a/src/gallium/drivers/nouveau/codegen/nv50_ir_from_nir.cpp b/src/gallium/drivers/nouveau/codegen/nv50_ir_from_nir.cpp
index 6d9ba1ccca5..0cd0bfab8d5 100644
--- a/src/gallium/drivers/nouveau/codegen/nv50_ir_from_nir.cpp
+++ b/src/gallium/drivers/nouveau/codegen/nv50_ir_from_nir.cpp
@@ -852,6 +852,7 @@ uint32_t
Converter::getIndirect(nir_intrinsic_instr *insn, uint8_t s, uint8_t c, Value *&indirect, bool isScalar)
{
int32_t idx = nir_intrinsic_base(insn) + getIndirect(&insn->src[s], c, indirect);
+
if (indirect && !isScalar)
indirect = mkOp2v(OP_SHL, TYPE_U32, getSSA(4, FILE_ADDRESS), indirect, loadImm(NULL, 4));
return idx;
@@ -1311,6 +1312,23 @@ Converter::parseNIR()
info->prop.cp.numThreads[1] = nir->info.workgroup_size[1];
info->prop.cp.numThreads[2] = nir->info.workgroup_size[2];
info_out->bin.smemSize = std::max(info_out->bin.smemSize, nir->info.shared_size);
+
+ if (info->target < NVISA_GF100_CHIPSET) {
+ int gmemSlot = 0;
+
+ for (unsigned i = 0; i < nir->info.num_ssbos; i++) {
+ info_out->prop.cp.gmem[gmemSlot++] = {.valid = 1, .image = 0, .slot = i};
+ assert(gmemSlot < 16);
+ }
+ nir_foreach_image_variable(var, nir) {
+ int image_count = glsl_type_get_image_count(var->type);
+ for (int i = 0; i < image_count; i++) {
+ info_out->prop.cp.gmem[gmemSlot++] = {.valid = 1, .image = 1, .slot = var->data.binding + i};
+ assert(gmemSlot < 16);
+ }
+ }
+ }
+
break;
case Program::TYPE_FRAGMENT:
info_out->prop.fp.earlyFragTests = nir->info.fs.early_fragment_tests;
@@ -2270,6 +2288,12 @@ Converter::visit(nir_intrinsic_instr *insn)
else
location = getIndirect(&insn->src[0], 0, indirect);
+ /* Pre-GF100, SSBOs and images are in the same HW file, managed by
+ * prop.cp.gmem. images are located after SSBOs.
+ */
+ if (info->target < NVISA_GF100_CHIPSET)
+ location += nir->info.num_ssbos;
+
// coords
if (opInfo.num_srcs >= 2)
for (unsigned int i = 0u; i < argCount; ++i)
More information about the mesa-commit
mailing list