Lines Matching full:nir
35 #include <compiler/nir/nir_builder.h>
36 #include <compiler/nir/nir_serialize.h>
345 clover_lower_nir(nir_shader *nir, std::vector<binary::argument> &args, in clover_lower_nir() argument
349 if (nir->constant_data_size) { in clover_lower_nir()
352 constant_var = nir_variable_create(nir, nir_var_uniform, type, in clover_lower_nir()
363 return nir_shader_lower_instructions(nir, in clover_lower_nir()
397 struct disk_cache *clover::nir::create_clc_disk_cache(void) in create_clc_disk_cache()
404 if (!disk_cache_get_function_identifier((void *)clover::nir::create_clc_disk_cache, &ctx)) in create_clc_disk_cache()
413 void clover::nir::check_for_libclc(const device &dev) in check_for_libclc()
419 nir_shader *clover::nir::load_libclc_nir(const device &dev, std::string &r_log) in load_libclc_nir()
436 binary clover::nir::spirv_to_nir(const binary &mod, const device &dev, in spirv_to_nir()
440 std::shared_ptr<nir_shader> nir = dev.clc_nir; in spirv_to_nir() local
441 spirv_options.clc_shader = nir.get(); in spirv_to_nir()
459 nir_shader *nir = spirv_to_nir(data, num_words, nullptr, 0, in spirv_to_nir() local
462 if (!nir) { in spirv_to_nir()
463 r_log += "Translation from SPIR-V to NIR for kernel \"" + sym.name + in spirv_to_nir()
468 nir->info.workgroup_size_variable = sym.reqd_work_group_size[0] == 0; in spirv_to_nir()
469 nir->info.workgroup_size[0] = sym.reqd_work_group_size[0]; in spirv_to_nir()
470 nir->info.workgroup_size[1] = sym.reqd_work_group_size[1]; in spirv_to_nir()
471 nir->info.workgroup_size[2] = sym.reqd_work_group_size[2]; in spirv_to_nir()
472 nir_validate_shader(nir, "clover"); in spirv_to_nir()
476 NIR_PASS_V(nir, nir_lower_variable_initializers, nir_var_function_temp); in spirv_to_nir()
477 NIR_PASS_V(nir, nir_lower_returns); in spirv_to_nir()
478 NIR_PASS_V(nir, nir_lower_libclc, spirv_options.clc_shader); in spirv_to_nir()
480 NIR_PASS_V(nir, nir_inline_functions); in spirv_to_nir()
481 NIR_PASS_V(nir, nir_copy_prop); in spirv_to_nir()
482 NIR_PASS_V(nir, nir_opt_deref); in spirv_to_nir()
485 nir_remove_non_entrypoints(nir); in spirv_to_nir()
487 nir_validate_shader(nir, "clover after function inlining"); in spirv_to_nir()
489 NIR_PASS_V(nir, nir_lower_variable_initializers, ~nir_var_function_temp); in spirv_to_nir()
495 NIR_PASS_V(nir, nir_lower_printf, &printf_options); in spirv_to_nir()
497 NIR_PASS_V(nir, nir_remove_dead_variables, nir_var_function_temp, NULL); in spirv_to_nir()
500 NIR_PASS_V(nir, nir_split_var_copies); in spirv_to_nir()
501 NIR_PASS_V(nir, nir_opt_copy_prop_vars); in spirv_to_nir()
502 NIR_PASS_V(nir, nir_lower_var_copies); in spirv_to_nir()
503 NIR_PASS_V(nir, nir_lower_vars_to_ssa); in spirv_to_nir()
504 NIR_PASS_V(nir, nir_opt_dce); in spirv_to_nir()
505 NIR_PASS_V(nir, nir_lower_convert_alu_types, NULL); in spirv_to_nir()
508 NIR_PASS_V(nir, nir_lower_alu_to_scalar, in spirv_to_nir()
511 NIR_PASS_V(nir, nir_lower_system_values); in spirv_to_nir()
514 NIR_PASS_V(nir, nir_lower_compute_system_values, &sysval_options); in spirv_to_nir()
517 NIR_PASS_V(nir, nir_opt_constant_folding); in spirv_to_nir()
519 NIR_PASS_V(nir, nir_remove_dead_variables, nir_var_mem_constant, NULL); in spirv_to_nir()
520 NIR_PASS_V(nir, nir_lower_vars_to_explicit_types, nir_var_mem_constant, in spirv_to_nir()
522 if (nir->constant_data_size > 0) { in spirv_to_nir()
523 assert(nir->constant_data == NULL); in spirv_to_nir()
524 nir->constant_data = rzalloc_size(nir, nir->constant_data_size); in spirv_to_nir()
525 nir_gather_explicit_io_initializers(nir, nir->constant_data, in spirv_to_nir()
526 nir->constant_data_size, in spirv_to_nir()
529 NIR_PASS_V(nir, nir_lower_explicit_io, nir_var_mem_constant, in spirv_to_nir()
533 NIR_PASS_V(nir, clover_lower_nir, args, dev.max_block_size().size(), in spirv_to_nir()
536 NIR_PASS_V(nir, clover_nir_add_image_uniforms); in spirv_to_nir()
537 NIR_PASS_V(nir, nir_lower_vars_to_explicit_types, in spirv_to_nir()
539 NIR_PASS_V(nir, nir_lower_vars_to_explicit_types, in spirv_to_nir()
544 NIR_PASS_V(nir, nir_opt_deref); in spirv_to_nir()
545 NIR_PASS_V(nir, nir_lower_readonly_images_to_tex, false); in spirv_to_nir()
546 NIR_PASS_V(nir, clover_nir_lower_images); in spirv_to_nir()
547 NIR_PASS_V(nir, nir_lower_memcpy); in spirv_to_nir()
550 NIR_PASS_V(nir, nir_lower_explicit_io, nir_var_uniform, in spirv_to_nir()
551 nir->info.cs.ptr_size == 64 ? in spirv_to_nir()
555 NIR_PASS_V(nir, nir_lower_explicit_io, nir_var_mem_constant, in spirv_to_nir()
557 NIR_PASS_V(nir, nir_lower_explicit_io, nir_var_mem_shared, in spirv_to_nir()
560 NIR_PASS_V(nir, nir_lower_explicit_io, nir_var_function_temp, in spirv_to_nir()
563 NIR_PASS_V(nir, nir_lower_explicit_io, nir_var_mem_global, in spirv_to_nir()
569 NIR_PASS_V(nir, nir_remove_dead_variables, nir_var_all, &remove_dead_variables_options); in spirv_to_nir()
572 NIR_PASS_V(nir, nir_lower_int64); in spirv_to_nir()
574 NIR_PASS_V(nir, nir_opt_dce); in spirv_to_nir()
576 if (nir->constant_data_size) { in spirv_to_nir()
577 const char *ptr = reinterpret_cast<const char *>(nir->constant_data); in spirv_to_nir()
581 nir->constant_data_size, in spirv_to_nir()
582 { ptr, ptr + nir->constant_data_size } in spirv_to_nir()
584 nir->constant_data = NULL; in spirv_to_nir()
585 nir->constant_data_size = 0; in spirv_to_nir()
590 unsigned printf_info_count = nir->printf_info_count; in spirv_to_nir()
591 nir_printf_info *printf_infos = nir->printf_info; in spirv_to_nir()
597 nir_serialize(&blob, nir, false); in spirv_to_nir()
599 ralloc_free(nir); in spirv_to_nir()
633 binary clover::nir::spirv_to_nir(const binary &mod, const device &dev, std::string &r_log) in spirv_to_nir()