Lines Matching refs:nir
24 #include "nir.h"
258 clc_lower_images(nir_shader *nir, struct clc_image_lower_context *context)
260 nir_foreach_function(func, nir) {
284 clc_lower_64bit_semantics(nir_shader *nir)
286 nir_foreach_function(func, nir) {
326 clc_lower_nonnormalized_samplers(nir_shader *nir,
329 nir_foreach_function(func, nir) {
398 add_kernel_inputs_var(struct clc_dxil_object *dxil, nir_shader *nir,
406 nir_foreach_variable_with_modes(var, nir, nir_var_uniform)
416 nir_variable_create(nir, nir_var_mem_ubo,
426 struct nir_shader *nir, unsigned *cbv_id)
434 nir_variable_create(nir, nir_var_mem_ubo,
443 clc_lower_constant_to_ssbo(nir_shader *nir,
447 nir_foreach_variable_with_modes(var, nir, nir_var_mem_constant) {
455 nir_foreach_function(func, nir) {
481 clc_lower_global_to_ssbo(nir_shader *nir)
483 nir_foreach_function(func, nir) {
695 static bool shader_has_double(nir_shader *nir)
697 foreach_list_typed(nir_function, func, node, &nir->functions) {
743 struct nir_shader *nir;
792 nir = spirv_to_nir(linked_spirv->data, linked_spirv->size / 4,
798 if (!nir) {
802 nir->info.workgroup_size_variable = true;
804 NIR_PASS_V(nir, nir_lower_goto_ifs);
805 NIR_PASS_V(nir, nir_opt_dead_cf);
821 NIR_PASS(progress, nir, nir_copy_prop);
822 NIR_PASS(progress, nir, nir_opt_copy_prop_vars);
823 NIR_PASS(progress, nir, nir_opt_deref);
824 NIR_PASS(progress, nir, nir_opt_dce);
825 NIR_PASS(progress, nir, nir_opt_undef);
826 NIR_PASS(progress, nir, nir_opt_constant_folding);
827 NIR_PASS(progress, nir, nir_opt_cse);
828 NIR_PASS(progress, nir, nir_lower_vars_to_ssa);
829 NIR_PASS(progress, nir, nir_opt_algebraic);
835 NIR_PASS_V(nir, nir_lower_variable_initializers, nir_var_function_temp);
836 NIR_PASS_V(nir, nir_lower_returns);
837 NIR_PASS_V(nir, nir_lower_libclc, clc_libclc_get_clc_shader(lib));
838 NIR_PASS_V(nir, nir_inline_functions);
841 nir_remove_non_entrypoints(nir);
848 NIR_PASS(progress, nir, nir_copy_prop);
849 NIR_PASS(progress, nir, nir_opt_copy_prop_vars);
850 NIR_PASS(progress, nir, nir_opt_deref);
851 NIR_PASS(progress, nir, nir_opt_dce);
852 NIR_PASS(progress, nir, nir_opt_undef);
853 NIR_PASS(progress, nir, nir_opt_constant_folding);
854 NIR_PASS(progress, nir, nir_opt_cse);
855 NIR_PASS(progress, nir, nir_split_var_copies);
856 NIR_PASS(progress, nir, nir_lower_var_copies);
857 NIR_PASS(progress, nir, nir_lower_vars_to_ssa);
858 NIR_PASS(progress, nir, nir_opt_algebraic);
859 NIR_PASS(progress, nir, nir_opt_if, nir_opt_if_aggressive_last_continue | nir_opt_if_optimize_phi_true_false);
860 NIR_PASS(progress, nir, nir_opt_dead_cf);
861 NIR_PASS(progress, nir, nir_opt_remove_phis);
862 NIR_PASS(progress, nir, nir_opt_peephole_select, 8, true, true);
863 NIR_PASS(progress, nir, nir_lower_vec3_to_vec4, nir_var_mem_generic | nir_var_uniform);
867 NIR_PASS_V(nir, nir_scale_fdiv);
876 nir_foreach_variable_with_modes_safe(var, nir, nir_var_uniform) {
882 exec_node_insert_list_after(exec_list_get_tail(&nir->variables), &inline_samplers_list);
884 NIR_PASS_V(nir, nir_lower_variable_initializers, ~(nir_var_function_temp | nir_var_shader_temp));
887 NIR_PASS_V(nir, dxil_nir_lower_memcpy_deref);
891 assert(nir->scratch_size == 0);
892 NIR_PASS_V(nir, nir_lower_vars_to_explicit_types, nir_var_function_temp, glsl_get_cl_type_size_align);
898 NIR_PASS_V(nir, nir_lower_printf, &printf_options);
900 metadata->printf.info_count = nir->printf_info_count;
901 metadata->printf.infos = calloc(nir->printf_info_count, sizeof(struct clc_printf_info));
902 for (unsigned i = 0; i < nir->printf_info_count; i++) {
903 metadata->printf.infos[i].str = malloc(nir->printf_info[i].string_size);
904 memcpy(metadata->printf.infos[i].str, nir->printf_info[i].strings, nir->printf_info[i].string_size);
905 metadata->printf.infos[i].num_args = nir->printf_info[i].num_args;
906 metadata->printf.infos[i].arg_sizes = malloc(nir->printf_info[i].num_args * sizeof(unsigned));
907 memcpy(metadata->printf.infos[i].arg_sizes, nir->printf_info[i].arg_sizes, nir->printf_info[i].num_args * sizeof(unsigned));
911 NIR_PASS_V(nir, nir_split_var_copies);
912 NIR_PASS_V(nir, nir_opt_copy_prop_vars);
913 NIR_PASS_V(nir, nir_lower_var_copies);
914 NIR_PASS_V(nir, nir_lower_vars_to_ssa);
915 NIR_PASS_V(nir, nir_lower_alu);
916 NIR_PASS_V(nir, nir_opt_dce);
917 NIR_PASS_V(nir, nir_opt_deref);
920 NIR_PASS_V(nir, nir_lower_vars_to_explicit_types, nir_var_uniform, glsl_get_cl_type_size_align);
924 nir_foreach_variable_with_modes(var, nir, nir_var_uniform) {
955 nir_foreach_image_variable(var, nir) {
980 NIR_PASS_V(nir, clc_nir_dedupe_const_samplers);
981 NIR_PASS_V(nir, nir_remove_dead_variables, nir_var_uniform | nir_var_mem_ubo |
985 nir_foreach_variable_with_modes(var, nir, nir_var_uniform) {
1007 NIR_PASS_V(nir, nir_lower_readonly_images_to_tex, false);
1009 NIR_PASS_V(nir, clc_lower_images, &image_lower_context);
1010 NIR_PASS_V(nir, clc_lower_nonnormalized_samplers, int_sampler_states);
1011 NIR_PASS_V(nir, nir_lower_samplers);
1012 NIR_PASS_V(nir, dxil_lower_sample_to_txf_for_integer_tex,
1015 NIR_PASS_V(nir, nir_remove_dead_variables, nir_var_mem_shared | nir_var_function_temp, NULL);
1017 nir->scratch_size = 0;
1018 NIR_PASS_V(nir, nir_lower_vars_to_explicit_types,
1022 NIR_PASS_V(nir, dxil_nir_lower_ubo_to_temp);
1023 NIR_PASS_V(nir, clc_lower_constant_to_ssbo, out_dxil->kernel, &uav_id);
1024 NIR_PASS_V(nir, clc_lower_global_to_ssbo);
1027 NIR_PASS(has_printf, nir, clc_lower_printf_base, uav_id);
1030 NIR_PASS_V(nir, dxil_nir_lower_deref_ssbo);
1032 NIR_PASS_V(nir, split_unaligned_loads_stores);
1034 assert(nir->info.cs.ptr_size == 64);
1035 NIR_PASS_V(nir, nir_lower_explicit_io, nir_var_mem_ssbo,
1037 NIR_PASS_V(nir, nir_lower_explicit_io,
1041 NIR_PASS_V(nir, nir_lower_system_values);
1047 NIR_PASS_V(nir, nir_lower_compute_system_values, &compute_options);
1049 NIR_PASS_V(nir, clc_lower_64bit_semantics);
1051 NIR_PASS_V(nir, nir_opt_deref);
1052 NIR_PASS_V(nir, nir_lower_vars_to_ssa);
1057 add_kernel_inputs_var(out_dxil, nir, &cbv_id);
1059 add_work_properties_var(out_dxil, nir, &cbv_id);
1061 memcpy(metadata->local_size, nir->info.workgroup_size,
1063 memcpy(metadata->local_size_hint, nir->info.cs.workgroup_size_hint,
1068 for (unsigned i = 0; i < ARRAY_SIZE(nir->info.workgroup_size); i++) {
1070 conf->local_size[i] == nir->info.workgroup_size[i])
1073 if (nir->info.workgroup_size[i] &&
1074 nir->info.workgroup_size[i] != conf->local_size[i]) {
1079 nir->info.workgroup_size[i] = conf->local_size[i];
1081 memcpy(metadata->local_size, nir->info.workgroup_size,
1085 for (unsigned i = 0; i < ARRAY_SIZE(nir->info.workgroup_size); i++) {
1086 if (nir->info.workgroup_size[i] == 0)
1087 nir->info.workgroup_size[i] = 1;
1091 NIR_PASS_V(nir, clc_nir_lower_kernel_input_loads, inputs_var);
1092 NIR_PASS_V(nir, split_unaligned_loads_stores);
1093 NIR_PASS_V(nir, nir_lower_explicit_io, nir_var_mem_ubo,
1095 NIR_PASS_V(nir, clc_nir_lower_system_values, work_properties_var);
1096 NIR_PASS_V(nir, dxil_nir_lower_loads_stores_to_dxil);
1097 NIR_PASS_V(nir, dxil_nir_opt_alu_deref_srcs);
1098 NIR_PASS_V(nir, dxil_nir_lower_atomics_to_dxil);
1099 NIR_PASS_V(nir, nir_lower_fp16_casts);
1100 NIR_PASS_V(nir, nir_lower_convert_alu_types, NULL);
1103 NIR_PASS_V(nir, nir_lower_pack);
1105 NIR_PASS_V(nir, nir_opt_algebraic);
1107 NIR_PASS_V(nir, nir_opt_dce);
1109 nir_validate_shader(nir, "Validate before feeding NIR to the DXIL compiler");
1141 nir->info.shared_size = align(nir->info.shared_size, alignment);
1142 metadata->args[i].localptr.sharedmem_offset = nir->info.shared_size;
1143 nir->info.shared_size += size;
1146 metadata->local_mem_size = nir->info.shared_size;
1147 metadata->priv_mem_size = nir->scratch_size;
1153 if (shader_has_double(nir)) {
1159 if (!nir_to_dxil(nir, &opts, &tmp)) {
1164 nir_foreach_variable_with_modes(var, nir, nir_var_mem_ssbo) {
1188 ralloc_free(nir);