fdo-mirrors/mesa
1
0
Fork 0
mirror of https://gitlab.freedesktop.org/mesa/mesa.git synced 2025-12-30 23:00:11 +01:00
Code Issues Projects Releases Packages Wiki Activity Actions 1

nouveau/nir: Add support for pre-GF100 images and ssbos.

Browse source 
We have to allocate them slots in the global file.

Reviewed-by: Karol Herbst <kherbst@redhat.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/15949>
This commit is contained in:
Emma Anholt 2022-04-24 12:44:23 -07:00 • committed by Marge Bot
parent 75f0127d78
commit c228cb3889
1 changed files with 24 additions and 0 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

24
src/gallium/drivers/nouveau/codegen/nv50_ir_from_nir.cpp
View file

@ -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)