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