Lines Matching full:nir
24 #include "nir.h"
276 clc_lower_images(nir_shader *nir, struct clc_image_lower_context *context) in clc_lower_images() argument
278 nir_foreach_function(func, nir) { in clc_lower_images()
299 nir_foreach_variable_with_modes_safe(var, nir, nir_var_image) { in clc_lower_images()
306 clc_lower_64bit_semantics(nir_shader *nir) in clc_lower_64bit_semantics() argument
308 nir_foreach_function_impl(impl, nir) { in clc_lower_64bit_semantics()
347 clc_lower_nonnormalized_samplers(nir_shader *nir, in clc_lower_nonnormalized_samplers() argument
350 nir_foreach_function(func, nir) { in clc_lower_nonnormalized_samplers()
416 add_kernel_inputs_var(struct clc_dxil_object *dxil, nir_shader *nir, in add_kernel_inputs_var() argument
424 nir_foreach_variable_with_modes(var, nir, nir_var_uniform) in add_kernel_inputs_var()
434 nir_variable_create(nir, nir_var_mem_ubo, in add_kernel_inputs_var()
444 struct nir_shader *nir, unsigned *cbv_id) in add_work_properties_var() argument
452 nir_variable_create(nir, nir_var_mem_ubo, in add_work_properties_var()
461 clc_lower_constant_to_ssbo(nir_shader *nir, in clc_lower_constant_to_ssbo() argument
465 nir_foreach_variable_with_modes(var, nir, nir_var_mem_constant) { in clc_lower_constant_to_ssbo()
473 nir_foreach_function(func, nir) { in clc_lower_constant_to_ssbo()
496 clc_change_variable_mode(nir_shader *nir, nir_variable_mode from, nir_variable_mode to) in clc_change_variable_mode() argument
498 nir_foreach_variable_with_modes(var, nir, from) in clc_change_variable_mode()
501 nir_foreach_function(func, nir) { in clc_change_variable_mode()
583 static bool shader_has_double(nir_shader *nir) in shader_has_double() argument
585 foreach_list_typed(nir_function, func, node, &nir->functions) { in shader_has_double()
631 struct nir_shader *nir; in clc_spirv_to_dxil() local
685 nir = spirv_to_nir(linked_spirv->data, linked_spirv->size / 4, in clc_spirv_to_dxil()
691 if (!nir) { in clc_spirv_to_dxil()
695 nir->info.workgroup_size_variable = true; in clc_spirv_to_dxil()
697 NIR_PASS_V(nir, nir_lower_goto_ifs); in clc_spirv_to_dxil()
698 NIR_PASS_V(nir, nir_opt_dead_cf); in clc_spirv_to_dxil()
714 NIR_PASS(progress, nir, nir_copy_prop); in clc_spirv_to_dxil()
715 NIR_PASS(progress, nir, nir_opt_copy_prop_vars); in clc_spirv_to_dxil()
716 NIR_PASS(progress, nir, nir_opt_deref); in clc_spirv_to_dxil()
717 NIR_PASS(progress, nir, nir_opt_dce); in clc_spirv_to_dxil()
718 NIR_PASS(progress, nir, nir_opt_undef); in clc_spirv_to_dxil()
719 NIR_PASS(progress, nir, nir_opt_constant_folding); in clc_spirv_to_dxil()
720 NIR_PASS(progress, nir, nir_opt_cse); in clc_spirv_to_dxil()
721 NIR_PASS(progress, nir, nir_lower_vars_to_ssa); in clc_spirv_to_dxil()
722 NIR_PASS(progress, nir, nir_opt_algebraic); in clc_spirv_to_dxil()
728 NIR_PASS_V(nir, nir_lower_variable_initializers, nir_var_function_temp); in clc_spirv_to_dxil()
729 NIR_PASS_V(nir, nir_lower_returns); in clc_spirv_to_dxil()
730 NIR_PASS_V(nir, nir_link_shader_functions, clc_libclc_get_clc_shader(lib)); in clc_spirv_to_dxil()
731 NIR_PASS_V(nir, nir_inline_functions); in clc_spirv_to_dxil()
734 nir_remove_non_entrypoints(nir); in clc_spirv_to_dxil()
741 NIR_PASS(progress, nir, nir_copy_prop); in clc_spirv_to_dxil()
742 NIR_PASS(progress, nir, nir_opt_copy_prop_vars); in clc_spirv_to_dxil()
743 NIR_PASS(progress, nir, nir_opt_deref); in clc_spirv_to_dxil()
744 NIR_PASS(progress, nir, nir_opt_dce); in clc_spirv_to_dxil()
745 NIR_PASS(progress, nir, nir_opt_undef); in clc_spirv_to_dxil()
746 NIR_PASS(progress, nir, nir_opt_constant_folding); in clc_spirv_to_dxil()
747 NIR_PASS(progress, nir, nir_opt_cse); in clc_spirv_to_dxil()
748 NIR_PASS(progress, nir, nir_split_var_copies); in clc_spirv_to_dxil()
749 NIR_PASS(progress, nir, nir_lower_var_copies); in clc_spirv_to_dxil()
750 NIR_PASS(progress, nir, nir_lower_vars_to_ssa); in clc_spirv_to_dxil()
751 NIR_PASS(progress, nir, nir_opt_algebraic); in clc_spirv_to_dxil()
752 NIR_PASS(progress, nir, nir_opt_if, nir_opt_if_optimize_phi_true_false); in clc_spirv_to_dxil()
753 NIR_PASS(progress, nir, nir_opt_dead_cf); in clc_spirv_to_dxil()
754 NIR_PASS(progress, nir, nir_opt_remove_phis); in clc_spirv_to_dxil()
755 NIR_PASS(progress, nir, nir_opt_peephole_select, 8, true, true); in clc_spirv_to_dxil()
756 NIR_PASS(progress, nir, nir_lower_vec3_to_vec4, nir_var_mem_generic | nir_var_uniform); in clc_spirv_to_dxil()
757 NIR_PASS(progress, nir, nir_opt_memcpy); in clc_spirv_to_dxil()
761 NIR_PASS_V(nir, nir_scale_fdiv); in clc_spirv_to_dxil()
767 NIR_PASS_V(nir, nir_lower_variable_initializers, ~(nir_var_function_temp | nir_var_shader_temp)); in clc_spirv_to_dxil()
771 assert(nir->scratch_size == 0); in clc_spirv_to_dxil()
772 …NIR_PASS_V(nir, nir_lower_vars_to_explicit_types, nir_var_function_temp, glsl_get_cl_type_size_ali… in clc_spirv_to_dxil()
777 NIR_PASS_V(nir, nir_lower_printf, &printf_options); in clc_spirv_to_dxil()
779 metadata->printf.info_count = nir->printf_info_count; in clc_spirv_to_dxil()
780 metadata->printf.infos = calloc(nir->printf_info_count, sizeof(struct clc_printf_info)); in clc_spirv_to_dxil()
781 for (unsigned i = 0; i < nir->printf_info_count; i++) { in clc_spirv_to_dxil()
782 metadata->printf.infos[i].str = malloc(nir->printf_info[i].string_size); in clc_spirv_to_dxil()
783 …memcpy(metadata->printf.infos[i].str, nir->printf_info[i].strings, nir->printf_info[i].string_size… in clc_spirv_to_dxil()
784 metadata->printf.infos[i].num_args = nir->printf_info[i].num_args; in clc_spirv_to_dxil()
785 metadata->printf.infos[i].arg_sizes = malloc(nir->printf_info[i].num_args * sizeof(unsigned)); in clc_spirv_to_dxil()
786 …memcpy(metadata->printf.infos[i].arg_sizes, nir->printf_info[i].arg_sizes, nir->printf_info[i].num… in clc_spirv_to_dxil()
790 NIR_PASS_V(nir, nir_lower_vars_to_explicit_types, nir_var_uniform, glsl_get_cl_type_size_align); in clc_spirv_to_dxil()
794 nir_foreach_variable_with_modes(var, nir, nir_var_uniform) { in clc_spirv_to_dxil()
825 nir_foreach_image_variable(var, nir) { in clc_spirv_to_dxil()
850 NIR_PASS_V(nir, nir_dedup_inline_samplers); in clc_spirv_to_dxil()
851 NIR_PASS_V(nir, nir_remove_dead_variables, nir_var_uniform | nir_var_mem_ubo | in clc_spirv_to_dxil()
855 nir_foreach_variable_with_modes(var, nir, nir_var_uniform) { in clc_spirv_to_dxil()
877 NIR_PASS_V(nir, nir_lower_readonly_images_to_tex, false); in clc_spirv_to_dxil()
879 NIR_PASS_V(nir, clc_lower_images, &image_lower_context); in clc_spirv_to_dxil()
880 NIR_PASS_V(nir, clc_lower_nonnormalized_samplers, int_sampler_states); in clc_spirv_to_dxil()
881 NIR_PASS_V(nir, nir_lower_samplers); in clc_spirv_to_dxil()
882 NIR_PASS_V(nir, dxil_lower_sample_to_txf_for_integer_tex, in clc_spirv_to_dxil()
885 nir->scratch_size = 0; in clc_spirv_to_dxil()
886 NIR_PASS_V(nir, nir_lower_vars_to_explicit_types, in clc_spirv_to_dxil()
895 NIR_PASS(progress, nir, nir_opt_memcpy); in clc_spirv_to_dxil()
896 NIR_PASS(progress, nir, nir_copy_prop); in clc_spirv_to_dxil()
897 NIR_PASS(progress, nir, nir_opt_copy_prop_vars); in clc_spirv_to_dxil()
898 NIR_PASS(progress, nir, nir_opt_deref); in clc_spirv_to_dxil()
899 NIR_PASS(progress, nir, nir_opt_dce); in clc_spirv_to_dxil()
900 NIR_PASS(progress, nir, nir_split_var_copies); in clc_spirv_to_dxil()
901 NIR_PASS(progress, nir, nir_lower_var_copies); in clc_spirv_to_dxil()
902 NIR_PASS(progress, nir, nir_lower_vars_to_ssa); in clc_spirv_to_dxil()
903 NIR_PASS(progress, nir, nir_opt_constant_folding); in clc_spirv_to_dxil()
904 NIR_PASS(progress, nir, nir_opt_cse); in clc_spirv_to_dxil()
907 NIR_PASS_V(nir, nir_lower_memcpy); in clc_spirv_to_dxil()
910 NIR_PASS_V(nir, dxil_nir_lower_constant_to_temp); in clc_spirv_to_dxil()
913 nir->info.cs.ptr_size = 32; in clc_spirv_to_dxil()
914 NIR_PASS_V(nir, nir_split_struct_vars, nir_var_shader_temp); in clc_spirv_to_dxil()
915 NIR_PASS_V(nir, dxil_nir_flatten_var_arrays, nir_var_shader_temp); in clc_spirv_to_dxil()
916 NIR_PASS_V(nir, dxil_nir_lower_var_bit_size, nir_var_shader_temp, in clc_spirv_to_dxil()
918 nir->info.cs.ptr_size = 64; in clc_spirv_to_dxil()
920 NIR_PASS_V(nir, clc_lower_constant_to_ssbo, out_dxil->kernel, &uav_id); in clc_spirv_to_dxil()
921 NIR_PASS_V(nir, clc_change_variable_mode, nir_var_shader_temp, nir_var_mem_constant); in clc_spirv_to_dxil()
922 NIR_PASS_V(nir, clc_change_variable_mode, nir_var_mem_global, nir_var_mem_ssbo); in clc_spirv_to_dxil()
925 NIR_PASS(has_printf, nir, clc_lower_printf_base, uav_id); in clc_spirv_to_dxil()
928 NIR_PASS_V(nir, dxil_nir_lower_deref_ssbo); in clc_spirv_to_dxil()
930 …NIR_PASS_V(nir, dxil_nir_split_unaligned_loads_stores, nir_var_mem_shared | nir_var_function_temp); in clc_spirv_to_dxil()
932 assert(nir->info.cs.ptr_size == 64); in clc_spirv_to_dxil()
933 NIR_PASS_V(nir, nir_lower_explicit_io, nir_var_mem_ssbo, in clc_spirv_to_dxil()
935 NIR_PASS_V(nir, nir_lower_explicit_io, in clc_spirv_to_dxil()
939 NIR_PASS_V(nir, nir_lower_system_values); in clc_spirv_to_dxil()
945 NIR_PASS_V(nir, nir_lower_compute_system_values, &compute_options); in clc_spirv_to_dxil()
947 NIR_PASS_V(nir, clc_lower_64bit_semantics); in clc_spirv_to_dxil()
949 NIR_PASS_V(nir, nir_opt_deref); in clc_spirv_to_dxil()
950 NIR_PASS_V(nir, nir_lower_vars_to_ssa); in clc_spirv_to_dxil()
955 add_kernel_inputs_var(out_dxil, nir, &cbv_id); in clc_spirv_to_dxil()
957 add_work_properties_var(out_dxil, nir, &cbv_id); in clc_spirv_to_dxil()
959 memcpy(metadata->local_size, nir->info.workgroup_size, in clc_spirv_to_dxil()
961 memcpy(metadata->local_size_hint, nir->info.cs.workgroup_size_hint, in clc_spirv_to_dxil()
966 for (unsigned i = 0; i < ARRAY_SIZE(nir->info.workgroup_size); i++) { in clc_spirv_to_dxil()
968 conf->local_size[i] == nir->info.workgroup_size[i]) in clc_spirv_to_dxil()
971 if (nir->info.workgroup_size[i] && in clc_spirv_to_dxil()
972 nir->info.workgroup_size[i] != conf->local_size[i]) { in clc_spirv_to_dxil()
977 nir->info.workgroup_size[i] = conf->local_size[i]; in clc_spirv_to_dxil()
979 memcpy(metadata->local_size, nir->info.workgroup_size, in clc_spirv_to_dxil()
983 for (unsigned i = 0; i < ARRAY_SIZE(nir->info.workgroup_size); i++) { in clc_spirv_to_dxil()
984 if (nir->info.workgroup_size[i] == 0) in clc_spirv_to_dxil()
985 nir->info.workgroup_size[i] = 1; in clc_spirv_to_dxil()
989 NIR_PASS_V(nir, clc_nir_lower_kernel_input_loads, inputs_var); in clc_spirv_to_dxil()
990 NIR_PASS_V(nir, nir_lower_explicit_io, nir_var_mem_ubo, in clc_spirv_to_dxil()
992 NIR_PASS_V(nir, clc_nir_lower_system_values, work_properties_var); in clc_spirv_to_dxil()
1019 nir->info.shared_size = align(nir->info.shared_size, alignment); in clc_spirv_to_dxil()
1020 metadata->args[i].localptr.sharedmem_offset = nir->info.shared_size; in clc_spirv_to_dxil()
1021 nir->info.shared_size += size; in clc_spirv_to_dxil()
1024 NIR_PASS_V(nir, dxil_nir_lower_loads_stores_to_dxil, &loads_stores_options); in clc_spirv_to_dxil()
1025 NIR_PASS_V(nir, dxil_nir_opt_alu_deref_srcs); in clc_spirv_to_dxil()
1026 NIR_PASS_V(nir, nir_lower_fp16_casts, nir_lower_fp16_all); in clc_spirv_to_dxil()
1027 NIR_PASS_V(nir, nir_lower_convert_alu_types, NULL); in clc_spirv_to_dxil()
1030 NIR_PASS_V(nir, nir_lower_pack); in clc_spirv_to_dxil()
1032 NIR_PASS_V(nir, nir_opt_algebraic); in clc_spirv_to_dxil()
1034 NIR_PASS_V(nir, nir_opt_dce); in clc_spirv_to_dxil()
1036 nir_validate_shader(nir, "Validate before feeding NIR to the DXIL compiler"); in clc_spirv_to_dxil()
1047 metadata->local_mem_size = nir->info.shared_size; in clc_spirv_to_dxil()
1048 metadata->priv_mem_size = nir->scratch_size; in clc_spirv_to_dxil()
1050 /* DXIL double math is too limited compared to what NIR expects. Let's refuse in clc_spirv_to_dxil()
1054 if (shader_has_double(nir)) { in clc_spirv_to_dxil()
1055 clc_error(logger, "NIR shader contains doubles, which we don't support yet"); in clc_spirv_to_dxil()
1063 if (!nir_to_dxil(nir, &opts, logger ? &dxil_logger : NULL, &tmp)) { in clc_spirv_to_dxil()
1068 nir_foreach_variable_with_modes(var, nir, nir_var_mem_ssbo) { in clc_spirv_to_dxil()
1092 ralloc_free(nir); in clc_spirv_to_dxil()