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