From: Dave Airlie Date: Fri, 10 Jun 2022 18:34:15 +0000 (-0500) Subject: nvk: Initial wiring in of the compiler X-Git-Tag: upstream/23.3.3~4555 X-Git-Url: http://review.tizen.org/git/?a=commitdiff_plain;h=c165e5b5b025c203101b5bedcfc07ca660250fa8;p=platform%2Fupstream%2Fmesa.git nvk: Initial wiring in of the compiler Part-of: --- diff --git a/src/nouveau/vulkan/meson.build b/src/nouveau/vulkan/meson.build index a095ab1..46b6bb8 100644 --- a/src/nouveau/vulkan/meson.build +++ b/src/nouveau/vulkan/meson.build @@ -33,6 +33,8 @@ nvk_files = files( 'nvk_private.h', 'nvk_sampler.c', 'nvk_sampler.h', + 'nvk_shader.c', + 'nvk_shader.h', 'nvk_wsi.c', 'nvk_wsi.h' ) diff --git a/src/nouveau/vulkan/nvk_shader.c b/src/nouveau/vulkan/nvk_shader.c new file mode 100644 index 0000000..7216bd3 --- /dev/null +++ b/src/nouveau/vulkan/nvk_shader.c @@ -0,0 +1,186 @@ + +#include "nvk_device.h" +#include "nvk_shader.h" +#include "nvk_physical_device.h" +#include "nvk_pipeline_layout.h" +#include "nvk_nir.h" + +#include "nouveau_bo.h" +#include "vk_shader_module.h" + +#include "nir.h" +#include "compiler/spirv/nir_spirv.h" + +#include "nv50_ir_driver.h" + +static void +shared_var_info(const struct glsl_type *type, unsigned *size, unsigned *align) +{ + assert(glsl_type_is_vector_or_scalar(type)); + + uint32_t comp_size = glsl_type_is_boolean(type) ? 4 : glsl_get_bit_size(type) / 8; + unsigned length = glsl_get_vector_elements(type); + *size = comp_size * length, *align = comp_size; +} + +static inline enum pipe_shader_type +pipe_shader_type_from_mesa(gl_shader_stage stage) +{ + switch (stage) { + case MESA_SHADER_VERTEX: + return PIPE_SHADER_VERTEX; + case MESA_SHADER_TESS_CTRL: + return PIPE_SHADER_TESS_CTRL; + case MESA_SHADER_TESS_EVAL: + return PIPE_SHADER_TESS_EVAL; + case MESA_SHADER_GEOMETRY: + return PIPE_SHADER_GEOMETRY; + case MESA_SHADER_FRAGMENT: + return PIPE_SHADER_FRAGMENT; + case MESA_SHADER_COMPUTE: + case MESA_SHADER_KERNEL: + return PIPE_SHADER_COMPUTE; + default: + unreachable("bad shader stage"); + } +} + +VkResult +nvk_shader_compile_to_nir(struct nvk_device *device, + struct vk_shader_module *module, + const char *entrypoint_name, + gl_shader_stage stage, + const VkSpecializationInfo *spec_info, + const struct nvk_pipeline_layout *layout, + nir_shader **nir_out) +{ + struct nvk_physical_device *pdevice = nvk_device_physical(device); + const nir_shader_compiler_options *nir_options = + nv50_ir_nir_shader_compiler_options(pdevice->dev->chipset, + pipe_shader_type_from_mesa(stage)); + + const struct spirv_to_nir_options spirv_options = { + .caps = + { + }, + .ssbo_addr_format = nir_address_format_64bit_global_32bit_offset, + .ubo_addr_format = nir_address_format_64bit_global_32bit_offset, + .shared_addr_format = nir_address_format_32bit_offset, + }; + + nir_shader *nir; + VkResult result = vk_shader_module_to_nir(&device->vk, module, stage, + entrypoint_name, spec_info, + &spirv_options, nir_options, + NULL, &nir); + if (result != VK_SUCCESS) + return result; + + NIR_PASS(_, nir, nir_lower_global_vars_to_local); + + NIR_PASS(_, nir, nir_split_struct_vars, nir_var_function_temp); + NIR_PASS(_, nir, nir_lower_vars_to_ssa); + + NIR_PASS(_, nir, nvk_nir_lower_descriptors, layout, true); + NIR_PASS(_, nir, nir_lower_system_values); + + nir_lower_compute_system_values_options csv_options = { + }; + NIR_PASS(_, nir, nir_lower_compute_system_values, &csv_options); + + /* Vulkan uses the separate-shader linking model */ + nir->info.separate_shader = true; + + NIR_PASS(_, nir, nir_lower_explicit_io, nir_var_mem_ssbo, + nir_address_format_64bit_global_32bit_offset); + + NIR_PASS(_, nir, nir_lower_explicit_io, nir_var_mem_ubo, + nir_address_format_64bit_global_32bit_offset); + + if (!nir->info.shared_memory_explicit_layout) { + NIR_PASS(_, nir, nir_lower_vars_to_explicit_types, + nir_var_mem_shared, shared_var_info); + } + NIR_PASS(_, nir, nir_lower_explicit_io, nir_var_mem_shared, + nir_address_format_32bit_offset); + + NIR_PASS(_, nir, nvk_nir_lower_descriptors, layout, true); + NIR_PASS(_, nir, nir_copy_prop); + NIR_PASS(_, nir, nir_opt_dce); + + *nir_out = nir; + + return VK_SUCCESS; +} + +VkResult +nvk_compile_nir(struct nvk_physical_device *device, nir_shader *nir, + struct nvk_shader *shader) +{ + struct nv50_ir_prog_info *info; + struct nv50_ir_prog_info_out info_out = {}; + int ret; + + info = CALLOC_STRUCT(nv50_ir_prog_info); + if (!info) + return false; + + info->type = pipe_shader_type_from_mesa(nir->info.stage); + info->target = device->dev->chipset; + info->bin.nir = nir; + + for (unsigned i = 0; i < 3; i++) + shader->cp.block_size[i] = nir->info.workgroup_size[i]; + + info->bin.smemSize = shader->cp.smem_size; + info->dbgFlags = debug_get_num_option("NV50_PROG_DEBUG", 0); + info->optLevel = debug_get_num_option("NV50_PROG_OPTIMIZE", 3); + info->io.auxCBSlot = 15; + info->io.uboInfoBase = 0; + if (nir->info.stage == MESA_SHADER_COMPUTE) { + info->io.auxCBSlot = 1; + info->prop.cp.gridInfoBase = 0; + } + ret = nv50_ir_generate_code(info, &info_out); + if (ret) + return VK_ERROR_UNKNOWN; + + shader->code_ptr = (uint8_t *)info_out.bin.code; + shader->code_size = info_out.bin.codeSize; + + if (info_out.target >= NVISA_GV100_CHIPSET) + shader->num_gprs = MIN2(info_out.bin.maxGPR + 5, 256); //XXX: why? + else + shader->num_gprs = MAX2(4, (info_out.bin.maxGPR + 1)); + shader->cp.smem_size = info_out.bin.smemSize; + shader->num_barriers = info_out.numBarriers; + + if (info_out.bin.tlsSpace) { + assert(info_out.bin.tlsSpace < (1 << 24)); + shader->hdr[0] |= 1 << 26; + shader->hdr[1] |= align(info_out.bin.tlsSpace, 0x10); /* l[] size */ + shader->need_tls = true; + } + + if (info_out.io.globalAccess) + shader->hdr[0] |= 1 << 26; + if (info_out.io.globalAccess & 0x2) + shader->hdr[0] |= 1 << 16; + if (info_out.io.fp64) + shader->hdr[0] |= 1 << 27; + + ralloc_free(nir); + return VK_SUCCESS; +} + +void +nvk_shader_upload(struct nvk_physical_device *physical, struct nvk_shader *shader) +{ + void *ptr; + shader->bo = nouveau_ws_bo_new(physical->dev, shader->code_size, 256, + NOUVEAU_WS_BO_LOCAL | NOUVEAU_WS_BO_MAP); + + ptr = nouveau_ws_bo_map(shader->bo, NOUVEAU_WS_BO_WR); + + memcpy(ptr, shader->code_ptr, shader->code_size); +} diff --git a/src/nouveau/vulkan/nvk_shader.h b/src/nouveau/vulkan/nvk_shader.h new file mode 100644 index 0000000..479c93c --- /dev/null +++ b/src/nouveau/vulkan/nvk_shader.h @@ -0,0 +1,47 @@ +#ifndef NVK_SHADER_H +#define NVK_SHADER_H 1 + +#include "nir.h" + +struct vk_shader_module; +struct nvk_device; +struct nvk_pipeline_layout; +struct nvk_physical_device; + +#define GF100_SHADER_HEADER_SIZE (20 * 4) +#define TU102_SHADER_HEADER_SIZE (32 * 4) +#define NVC0_MAX_SHADER_HEADER_SIZE TU102_SHADER_HEADER_SIZE + +struct nvk_shader { + uint8_t *code_ptr; + uint32_t code_size; + + uint32_t hdr[NVC0_MAX_SHADER_HEADER_SIZE/4]; + bool need_tls; + uint8_t num_gprs; + uint8_t num_barriers; + struct { + uint32_t lmem_size; /* local memory (TGSI PRIVATE resource) size */ + uint32_t smem_size; /* shared memory (TGSI LOCAL resource) size */ + uint32_t block_size[3]; + } cp; + + struct nouveau_ws_bo *bo; +}; + +VkResult +nvk_shader_compile_to_nir(struct nvk_device *device, + struct vk_shader_module *module, + const char *entrypoint_name, + gl_shader_stage stage, + const VkSpecializationInfo *spec_info, + const struct nvk_pipeline_layout *layout, + nir_shader **nir_out); + +VkResult +nvk_compile_nir(struct nvk_physical_device *device, nir_shader *nir, + struct nvk_shader *shader); + +void +nvk_shader_upload(struct nvk_physical_device *physical, struct nvk_shader *shader); +#endif