Lines Matching full:nir
36 #include <compiler/nir/nir_builder.h>
37 #include <compiler/nir/nir_serialize.h>
180 clover_lower_nir(nir_shader *nir, std::vector<binary::argument> &args, in clover_lower_nir() argument
184 if (nir->constant_data_size) { in clover_lower_nir()
187 constant_var = nir_variable_create(nir, nir_var_uniform, type, in clover_lower_nir()
198 return nir_shader_lower_instructions(nir, in clover_lower_nir()
232 struct disk_cache *clover::nir::create_clc_disk_cache(void) in create_clc_disk_cache()
239 if (!disk_cache_get_function_identifier((void *)clover::nir::create_clc_disk_cache, &ctx)) in create_clc_disk_cache()
248 void clover::nir::check_for_libclc(const device &dev) in check_for_libclc()
254 nir_shader *clover::nir::load_libclc_nir(const device &dev, std::string &r_log) in load_libclc_nir()
272 binary clover::nir::spirv_to_nir(const binary &mod, const device &dev, in spirv_to_nir()
276 std::shared_ptr<nir_shader> nir = dev.clc_nir; in spirv_to_nir() local
277 spirv_options.clc_shader = nir.get(); in spirv_to_nir()
295 nir_shader *nir = spirv_to_nir(data, num_words, nullptr, 0, in spirv_to_nir() local
298 if (!nir) { in spirv_to_nir()
299 r_log += "Translation from SPIR-V to NIR for kernel \"" + sym.name + in spirv_to_nir()
304 nir->info.workgroup_size_variable = sym.reqd_work_group_size[0] == 0; in spirv_to_nir()
305 nir->info.workgroup_size[0] = sym.reqd_work_group_size[0]; in spirv_to_nir()
306 nir->info.workgroup_size[1] = sym.reqd_work_group_size[1]; in spirv_to_nir()
307 nir->info.workgroup_size[2] = sym.reqd_work_group_size[2]; in spirv_to_nir()
308 nir_validate_shader(nir, "clover"); in spirv_to_nir()
312 NIR_PASS_V(nir, nir_lower_variable_initializers, nir_var_function_temp); in spirv_to_nir()
313 NIR_PASS_V(nir, nir_lower_returns); in spirv_to_nir()
314 NIR_PASS_V(nir, nir_link_shader_functions, spirv_options.clc_shader); in spirv_to_nir()
316 NIR_PASS_V(nir, nir_inline_functions); in spirv_to_nir()
317 NIR_PASS_V(nir, nir_copy_prop); in spirv_to_nir()
318 NIR_PASS_V(nir, nir_opt_deref); in spirv_to_nir()
321 nir_remove_non_entrypoints(nir); in spirv_to_nir()
323 nir_validate_shader(nir, "clover after function inlining"); in spirv_to_nir()
325 NIR_PASS_V(nir, nir_lower_variable_initializers, ~nir_var_function_temp); in spirv_to_nir()
330 NIR_PASS_V(nir, nir_lower_printf, &printf_options); in spirv_to_nir()
332 NIR_PASS_V(nir, nir_remove_dead_variables, nir_var_function_temp, NULL); in spirv_to_nir()
335 NIR_PASS_V(nir, nir_split_var_copies); in spirv_to_nir()
336 NIR_PASS_V(nir, nir_opt_copy_prop_vars); in spirv_to_nir()
337 NIR_PASS_V(nir, nir_lower_var_copies); in spirv_to_nir()
338 NIR_PASS_V(nir, nir_lower_vars_to_ssa); in spirv_to_nir()
339 NIR_PASS_V(nir, nir_opt_dce); in spirv_to_nir()
340 NIR_PASS_V(nir, nir_lower_convert_alu_types, NULL); in spirv_to_nir()
343 NIR_PASS_V(nir, nir_lower_alu_to_scalar, in spirv_to_nir()
346 NIR_PASS_V(nir, nir_lower_system_values); in spirv_to_nir()
349 NIR_PASS_V(nir, nir_lower_compute_system_values, &sysval_options); in spirv_to_nir()
352 NIR_PASS_V(nir, nir_opt_constant_folding); in spirv_to_nir()
354 NIR_PASS_V(nir, nir_remove_dead_variables, nir_var_mem_constant, NULL); in spirv_to_nir()
355 NIR_PASS_V(nir, nir_lower_vars_to_explicit_types, nir_var_mem_constant, in spirv_to_nir()
357 if (nir->constant_data_size > 0) { in spirv_to_nir()
358 assert(nir->constant_data == NULL); in spirv_to_nir()
359 nir->constant_data = rzalloc_size(nir, nir->constant_data_size); in spirv_to_nir()
360 nir_gather_explicit_io_initializers(nir, nir->constant_data, in spirv_to_nir()
361 nir->constant_data_size, in spirv_to_nir()
364 NIR_PASS_V(nir, nir_lower_explicit_io, nir_var_mem_constant, in spirv_to_nir()
368 NIR_PASS_V(nir, clover_lower_nir, args, dev.max_block_size().size(), in spirv_to_nir()
371 NIR_PASS_V(nir, clover_nir_add_image_uniforms); in spirv_to_nir()
372 NIR_PASS_V(nir, nir_lower_vars_to_explicit_types, in spirv_to_nir()
374 NIR_PASS_V(nir, nir_lower_vars_to_explicit_types, in spirv_to_nir()
379 NIR_PASS_V(nir, nir_opt_deref); in spirv_to_nir()
380 NIR_PASS_V(nir, nir_lower_readonly_images_to_tex, false); in spirv_to_nir()
381 NIR_PASS_V(nir, nir_lower_cl_images, true, true); in spirv_to_nir()
382 NIR_PASS_V(nir, nir_lower_memcpy); in spirv_to_nir()
385 NIR_PASS_V(nir, nir_lower_explicit_io, nir_var_uniform, in spirv_to_nir()
386 nir->info.cs.ptr_size == 64 ? in spirv_to_nir()
390 NIR_PASS_V(nir, nir_lower_explicit_io, nir_var_mem_constant, in spirv_to_nir()
392 NIR_PASS_V(nir, nir_lower_explicit_io, nir_var_mem_shared, in spirv_to_nir()
395 NIR_PASS_V(nir, nir_lower_explicit_io, nir_var_function_temp, in spirv_to_nir()
398 NIR_PASS_V(nir, nir_lower_explicit_io, nir_var_mem_global, in spirv_to_nir()
403 NIR_PASS_V(nir, nir_remove_dead_variables, nir_var_all, &remove_dead_variables_options); in spirv_to_nir()
406 NIR_PASS_V(nir, nir_lower_int64); in spirv_to_nir()
408 NIR_PASS_V(nir, nir_opt_dce); in spirv_to_nir()
410 if (nir->constant_data_size) { in spirv_to_nir()
411 const char *ptr = reinterpret_cast<const char *>(nir->constant_data); in spirv_to_nir()
415 nir->constant_data_size, in spirv_to_nir()
416 { ptr, ptr + nir->constant_data_size } in spirv_to_nir()
418 nir->constant_data = NULL; in spirv_to_nir()
419 nir->constant_data_size = 0; in spirv_to_nir()
424 unsigned printf_info_count = nir->printf_info_count; in spirv_to_nir()
425 u_printf_info *printf_infos = nir->printf_info; in spirv_to_nir()
431 nir_serialize(&blob, nir, false); in spirv_to_nir()
433 ralloc_free(nir); in spirv_to_nir()
467 binary clover::nir::spirv_to_nir(const binary &mod, const device &dev, std::string &r_log) in spirv_to_nir()