Lines Matching defs:mod
452 struct dxil_module mod;
514 const struct dxil_func *func = dxil_get_function(&ctx->mod,
520 const struct dxil_value *opcode = dxil_module_get_int32_const(&ctx->mod, intr);
529 return dxil_emit_call(&ctx->mod, func, args, ARRAY_SIZE(args));
537 const struct dxil_func *func = dxil_get_function(&ctx->mod, "dx.op.binary", overload);
541 const struct dxil_value *opcode = dxil_module_get_int32_const(&ctx->mod, intr);
551 return dxil_emit_call(&ctx->mod, func, args, ARRAY_SIZE(args));
561 const struct dxil_func *func = dxil_get_function(&ctx->mod, "dx.op.tertiary", overload);
565 const struct dxil_value *opcode = dxil_module_get_int32_const(&ctx->mod, intr);
576 return dxil_emit_call(&ctx->mod, func, args, ARRAY_SIZE(args));
587 const struct dxil_func *func = dxil_get_function(&ctx->mod, "dx.op.quaternary", overload);
591 const struct dxil_value *opcode = dxil_module_get_int32_const(&ctx->mod, intr);
603 return dxil_emit_call(&ctx->mod, func, args, ARRAY_SIZE(args));
609 const struct dxil_func *func = dxil_get_function(&ctx->mod, "dx.op.threadId", DXIL_I32);
613 const struct dxil_value *opcode = dxil_module_get_int32_const(&ctx->mod,
623 return dxil_emit_call(&ctx->mod, func, args, ARRAY_SIZE(args));
630 const struct dxil_func *func = dxil_get_function(&ctx->mod, "dx.op.threadIdInGroup", DXIL_I32);
635 const struct dxil_value *opcode = dxil_module_get_int32_const(&ctx->mod,
645 return dxil_emit_call(&ctx->mod, func, args, ARRAY_SIZE(args));
651 const struct dxil_func *func = dxil_get_function(&ctx->mod, "dx.op.flattenedThreadIdInGroup", DXIL_I32);
656 const struct dxil_value *opcode = dxil_module_get_int32_const(&ctx->mod,
665 return dxil_emit_call(&ctx->mod, func, args, ARRAY_SIZE(args));
671 const struct dxil_func *func = dxil_get_function(&ctx->mod, "dx.op.groupId", DXIL_I32);
676 const struct dxil_value *opcode = dxil_module_get_int32_const(&ctx->mod,
686 return dxil_emit_call(&ctx->mod, func, args, ARRAY_SIZE(args));
695 const struct dxil_func *func = dxil_get_function(&ctx->mod, "dx.op.bufferLoad", overload);
699 const struct dxil_value *opcode = dxil_module_get_int32_const(&ctx->mod,
703 return dxil_emit_call(&ctx->mod, func, args, ARRAY_SIZE(args));
714 const struct dxil_func *func = dxil_get_function(&ctx->mod, "dx.op.bufferStore", overload);
719 const struct dxil_value *opcode = dxil_module_get_int32_const(&ctx->mod,
727 return dxil_emit_call_void(&ctx->mod, func,
737 const struct dxil_func *func = dxil_get_function(&ctx->mod, "dx.op.textureLoad", overload);
740 const struct dxil_type *int_type = dxil_module_get_int_type(&ctx->mod, 32);
741 const struct dxil_value *int_undef = dxil_module_get_undef(&ctx->mod, int_type);
743 const struct dxil_value *opcode = dxil_module_get_int32_const(&ctx->mod,
750 return dxil_emit_call(&ctx->mod, func, args, ARRAY_SIZE(args));
761 const struct dxil_func *func = dxil_get_function(&ctx->mod, "dx.op.textureStore", overload);
766 const struct dxil_value *opcode = dxil_module_get_int32_const(&ctx->mod,
774 return dxil_emit_call_void(&ctx->mod, func,
785 const struct dxil_func *func = dxil_get_function(&ctx->mod, "dx.op.atomicBinOp", DXIL_I32);
791 dxil_module_get_int32_const(&ctx->mod, DXIL_INTR_ATOMIC_BINOP);
793 dxil_module_get_int32_const(&ctx->mod, atomic_op);
799 return dxil_emit_call(&ctx->mod, func, args, ARRAY_SIZE(args));
810 dxil_get_function(&ctx->mod, "dx.op.atomicCompareExchange", DXIL_I32);
816 dxil_module_get_int32_const(&ctx->mod, DXIL_INTR_ATOMIC_CMPXCHG);
821 return dxil_emit_call(&ctx->mod, func, args, ARRAY_SIZE(args));
831 const struct dxil_value *opcode = dxil_module_get_int32_const(&ctx->mod, DXIL_INTR_CREATE_HANDLE);
832 const struct dxil_value *resource_class_value = dxil_module_get_int8_const(&ctx->mod, resource_class);
833 const struct dxil_value *resource_range_id_value = dxil_module_get_int32_const(&ctx->mod, resource_range_id);
834 const struct dxil_value *non_uniform_resource_index_value = dxil_module_get_int1_const(&ctx->mod, non_uniform_resource_index);
848 dxil_get_function(&ctx->mod, "dx.op.createHandle", DXIL_NONE);
853 return dxil_emit_call(&ctx->mod, func, args, ARRAY_SIZE(args));
864 const struct dxil_value *resource_range_index_value = dxil_module_get_int32_const(&ctx->mod, resource_range_index);
880 if (ctx->mod.minor_validator >= 6) {
901 if (ctx->mod.minor_validator >= 6 && ctx->num_uavs > 8)
902 ctx->mod.feats.use_64uavs = 1;
943 unsigned resource_element_size = ctx->mod.minor_validator >= 6 ?
978 const struct dxil_type *res_type_as_type = dxil_module_get_res_type(&ctx->mod, res_kind, comp_type, false /* readwrite */);
981 res_type_as_type = dxil_module_get_array_type(&ctx->mod, res_type_as_type, count);
983 const struct dxil_mdnode *srv_meta = emit_srv_metadata(&ctx->mod, res_type_as_type, var->name,
992 ctx->mod.raw_and_structured_buffers = true;
1006 const struct dxil_type *struct_type = dxil_module_get_res_type(&ctx->mod,
1012 dxil_module_get_array_type(&ctx->mod, struct_type, size);
1018 emit_uav_metadata(&ctx->mod, array_type,
1026 if (ctx->mod.minor_validator < 6 &&
1028 ctx->mod.feats.use_64uavs = 1;
1031 ctx->mod.raw_and_structured_buffers = true;
1042 const struct dxil_type *res_type = dxil_module_get_res_type(&ctx->mod, res_kind, comp_type, true /* readwrite */);
1043 res_type = dxil_module_get_array_type(&ctx->mod, res_type, count);
1044 const struct dxil_mdnode *uav_meta = emit_uav_metadata(&ctx->mod, res_type, name,
1051 if (ctx->mod.minor_validator < 6 &&
1053 ctx->mod.feats.use_64uavs = 1;
1057 ctx->mod.raw_and_structured_buffers = true;
1058 if (ctx->mod.shader_kind != DXIL_PIXEL_SHADER &&
1059 ctx->mod.shader_kind != DXIL_COMPUTE_SHADER)
1060 ctx->mod.feats.uavs_at_every_stage = true;
1174 const_vals[i] = dxil_module_get_int32_const(&ctx->mod, const_ints[i]);
1176 const struct dxil_type *elt_type = dxil_module_get_int_type(&ctx->mod, 32);
1180 dxil_module_get_array_type(&ctx->mod, elt_type, num_members);
1184 dxil_module_get_array_const(&ctx->mod, type, const_vals);
1188 const struct dxil_value *gvar = dxil_add_global_ptr_var(&ctx->mod, var->name, type,
1209 const struct dxil_type *float32 = dxil_module_get_float_type(&ctx->mod, 32);
1210 const struct dxil_type *array_type = dxil_module_get_array_type(&ctx->mod, float32, size);
1211 const struct dxil_type *buffer_type = dxil_module_get_struct_type(&ctx->mod, name,
1214 const struct dxil_type *final_type = count != 1 ? dxil_module_get_array_type(&ctx->mod, buffer_type, count) : buffer_type;
1216 const struct dxil_mdnode *cbv_meta = emit_cbv_metadata(&ctx->mod, final_type,
1257 const struct dxil_type *int32_type = dxil_module_get_int_type(&ctx->mod, 32);
1258 const struct dxil_type *sampler_type = dxil_module_get_struct_type(&ctx->mod, "struct.SamplerState", &int32_type, 1);
1261 sampler_type = dxil_module_get_array_type(&ctx->mod, sampler_type, count);
1263 const struct dxil_mdnode *sampler_meta = emit_sampler_metadata(&ctx->mod, sampler_type, var, &layout);
1284 unsigned resource_element_size = ctx->mod.minor_validator >= 6 ?
1355 gs_state_nodes[0] = dxil_get_metadata_int32(&ctx->mod, dxil_get_input_primitive(s->info.gs.input_primitive));
1356 gs_state_nodes[1] = dxil_get_metadata_int32(&ctx->mod, s->info.gs.vertices_out);
1357 gs_state_nodes[2] = dxil_get_metadata_int32(&ctx->mod, MAX2(s->info.gs.active_stream_mask, 1));
1358 gs_state_nodes[3] = dxil_get_metadata_int32(&ctx->mod, dxil_get_primitive_topology(s->info.gs.output_primitive));
1359 gs_state_nodes[4] = dxil_get_metadata_int32(&ctx->mod, s->info.gs.invocations);
1366 return dxil_get_metadata_node(&ctx->mod, gs_state_nodes, ARRAY_SIZE(gs_state_nodes));
1415 hs_state_nodes[0] = dxil_get_metadata_func(&ctx->mod, ctx->tess_ctrl_patch_constant_func_def->func);
1416 hs_state_nodes[1] = dxil_get_metadata_int32(&ctx->mod, ctx->tess_input_control_point_count);
1417 hs_state_nodes[2] = dxil_get_metadata_int32(&ctx->mod, ctx->shader->info.tess.tcs_vertices_out);
1418 hs_state_nodes[3] = dxil_get_metadata_int32(&ctx->mod, get_tessellator_domain(ctx->shader->info.tess._primitive_mode));
1419 hs_state_nodes[4] = dxil_get_metadata_int32(&ctx->mod, get_tessellator_partitioning(ctx->shader->info.tess.spacing));
1420 hs_state_nodes[5] = dxil_get_metadata_int32(&ctx->mod, get_tessellator_output_primitive(&ctx->shader->info));
1421 hs_state_nodes[6] = dxil_get_metadata_float32(&ctx->mod, 64.0f);
1423 return dxil_get_metadata_node(&ctx->mod, hs_state_nodes, ARRAY_SIZE(hs_state_nodes));
1431 ds_state_nodes[0] = dxil_get_metadata_int32(&ctx->mod, get_tessellator_domain(ctx->shader->info.tess._primitive_mode));
1432 ds_state_nodes[1] = dxil_get_metadata_int32(&ctx->mod, ctx->shader->info.tess.tcs_vertices_out);
1434 return dxil_get_metadata_node(&ctx->mod, ds_state_nodes, ARRAY_SIZE(ds_state_nodes));
1441 const struct dxil_mdnode *threads_x = dxil_get_metadata_int32(&ctx->mod, MAX2(s->info.workgroup_size[0], 1));
1442 const struct dxil_mdnode *threads_y = dxil_get_metadata_int32(&ctx->mod, MAX2(s->info.workgroup_size[1], 1));
1443 const struct dxil_mdnode *threads_z = dxil_get_metadata_int32(&ctx->mod, MAX2(s->info.workgroup_size[2], 1));
1448 return dxil_get_metadata_node(&ctx->mod, threads_nodes, ARRAY_SIZE(threads_nodes));
1460 if (ctx->mod.feats.doubles)
1465 if (ctx->mod.raw_and_structured_buffers)
1467 if (ctx->mod.feats.min_precision)
1469 if (ctx->mod.feats.dx11_1_double_extensions)
1471 if (ctx->mod.feats.array_layer_from_vs_or_ds)
1473 if (ctx->mod.feats.inner_coverage)
1475 if (ctx->mod.feats.typed_uav_load_additional_formats)
1477 if (ctx->mod.feats.use_64uavs)
1479 if (ctx->mod.feats.uavs_at_every_stage)
1481 if (ctx->mod.feats.cs_4x_raw_sb)
1483 if (ctx->mod.feats.wave_ops)
1485 if (ctx->mod.feats.int64_ops)
1487 if (ctx->mod.feats.barycentrics)
1489 if (ctx->mod.feats.stencil_ref)
1491 if (ctx->mod.feats.native_low_precision)
1510 const struct dxil_mdnode *func_md = dxil_get_metadata_func(&ctx->mod, func);
1511 const struct dxil_mdnode *name_md = dxil_get_metadata_string(&ctx->mod, truncated_name);
1519 return dxil_get_metadata_node(&ctx->mod, nodes,
1534 resources_nodes[0] = dxil_get_metadata_node(&ctx->mod, ARRAY_AND_SIZE(ctx->srv_metadata_nodes));
1539 resources_nodes[1] = dxil_get_metadata_node(&ctx->mod, ARRAY_AND_SIZE(ctx->uav_metadata_nodes));
1544 resources_nodes[2] = dxil_get_metadata_node(&ctx->mod, ARRAY_AND_SIZE(ctx->cbv_metadata_nodes));
1549 resources_nodes[3] = dxil_get_metadata_node(&ctx->mod, ARRAY_AND_SIZE(ctx->sampler_metadata_nodes));
1556 dxil_get_metadata_node(&ctx->mod, resources_nodes, ARRAY_SIZE(resources_nodes)): NULL;
1563 const struct dxil_mdnode *tag_node = dxil_get_metadata_int32(&ctx->mod, tag);
1577 assert(ctx->mod.major_version == 6);
1579 unsigned dxilMinor = ctx->mod.minor_version;
1580 unsigned valMajor = ctx->mod.major_validator;
1581 unsigned valMinor = ctx->mod.minor_validator;
1582 if (!emit_llvm_ident(&ctx->mod) ||
1583 !emit_named_version(&ctx->mod, "dx.version", dxilMajor, dxilMinor) ||
1584 !emit_named_version(&ctx->mod, "dx.valver", valMajor, valMinor) ||
1585 !emit_dx_shader_model(&ctx->mod))
1595 const struct dxil_mdnode *main_entrypoint = dxil_get_metadata_func(&ctx->mod, main_func);
1596 const struct dxil_mdnode *node27 = dxil_get_metadata_node(&ctx->mod, NULL, 0);
1598 const struct dxil_mdnode *node4 = dxil_get_metadata_int32(&ctx->mod, 0);
1602 const struct dxil_mdnode *node28 = dxil_get_metadata_node(&ctx->mod, nodes_4_27_27,
1605 const struct dxil_mdnode *node29 = dxil_get_metadata_node(&ctx->mod, &node28, 1);
1607 const struct dxil_mdnode *node3 = dxil_get_metadata_int32(&ctx->mod, 1);
1611 const struct dxil_mdnode *main_type_annotation = dxil_get_metadata_node(&ctx->mod, main_type_annotation_nodes,
1614 if (ctx->mod.shader_kind == DXIL_GEOMETRY_SHADER) {
1617 } else if (ctx->mod.shader_kind == DXIL_HULL_SHADER) {
1628 } else if (ctx->mod.shader_kind == DXIL_DOMAIN_SHADER) {
1631 } else if (ctx->mod.shader_kind == DXIL_COMPUTE_SHADER) {
1638 if (!emit_tag(ctx, DXIL_SHADER_TAG_FLAGS, dxil_get_metadata_int64(&ctx->mod, flags)))
1643 shader_properties = dxil_get_metadata_node(&ctx->mod, ctx->shader_property_nodes,
1651 entry_func_impl->function->name, get_signatures(&ctx->mod), resources_node, shader_properties);
1657 dxil_add_metadata_named_node(&ctx->mod, "dx.resources",
1662 return dxil_add_metadata_named_node(&ctx->mod, "dx.typeAnnotations",
1665 dxil_add_metadata_named_node(&ctx->mod, "dx.entryPoints",
1673 const struct dxil_type *type = dxil_module_get_int_type(&ctx->mod, bit_size);
1677 return dxil_emit_cast(&ctx->mod, DXIL_CAST_BITCAST, type, value);
1684 const struct dxil_type *type = dxil_module_get_float_type(&ctx->mod, bit_size);
1688 return dxil_emit_cast(&ctx->mod, DXIL_CAST_BITCAST, type, value);
1703 value = dxil_emit_cast(&ctx->mod, DXIL_CAST_BITCAST, expect_type, value);
1724 ctx->mod.feats.doubles = true;
1730 ctx->mod.feats.native_low_precision = true;
1732 ctx->mod.feats.int64_ops = true;
1772 assert(bit_size != 64 || ctx->mod.feats.int64_ops);
1773 const struct dxil_type *expect_type = dxil_module_get_int_type(&ctx->mod, bit_size);
1783 assert(nir_src_bit_size(*src) != 64 || ctx->mod.feats.doubles);
1784 if (dxil_value_type_equal_to(value, dxil_module_get_float_type(&ctx->mod, bit_size)))
1791 return dxil_emit_cast(&ctx->mod, DXIL_CAST_TRUNC,
1792 dxil_module_get_int_type(&ctx->mod, 1), value);
1834 const struct dxil_value *v = dxil_emit_binop(&ctx->mod, opcode, op0, op1, flags);
1850 dxil_module_get_int_type(&ctx->mod, op0_bit_size);
1853 op1 = dxil_emit_cast(&ctx->mod, cast_op, type, op1);
1857 dxil_emit_binop(&ctx->mod, opcode, op0, op1, 0);
1869 const struct dxil_value *v = dxil_emit_cmp(&ctx->mod, pred, op0, op1);
1958 return dxil_module_get_int_type(&ctx->mod, dst_bits);
1961 return dxil_module_get_float_type(&ctx->mod, dst_bits);
1989 ctx->mod.feats.dx11_1_double_extensions = true;
1994 ctx->mod.feats.dx11_1_double_extensions = true;
2000 const struct dxil_value *v = dxil_emit_cast(&ctx->mod, opcode, type,
2111 const struct dxil_value *compare_width = dxil_emit_cmp(&ctx->mod, DXIL_ICMP_SGE,
2112 width, dxil_module_get_int32_const(&ctx->mod, 32));
2113 v = dxil_emit_select(&ctx->mod, compare_width, insert, v);
2127 const struct dxil_value *v = dxil_emit_select(&ctx->mod, sel, val_true, val_false);
2140 struct dxil_module *m = &ctx->mod;
2156 struct dxil_module *m = &ctx->mod;
2172 struct dxil_module *m = &ctx->mod;
2180 ctx->mod.feats.doubles = 1;
2189 const struct dxil_value *zero = dxil_module_get_float_const(&ctx->mod, 0.0f);
2197 val = dxil_emit_binop(&ctx->mod, DXIL_BINOP_LSHR, val,
2198 dxil_module_get_int32_const(&ctx->mod, 16), 0);
2203 const struct dxil_func *func = dxil_get_function(&ctx->mod,
2209 const struct dxil_value *opcode = dxil_module_get_int32_const(&ctx->mod, DXIL_INTR_LEGACY_F16TOF32);
2218 const struct dxil_value *v = dxil_emit_call(&ctx->mod, func, args, ARRAY_SIZE(args));
2228 const struct dxil_func *func = dxil_get_function(&ctx->mod,
2234 const struct dxil_value *opcode = dxil_module_get_int32_const(&ctx->mod, DXIL_INTR_LEGACY_F32TOF16);
2243 const struct dxil_value *v = dxil_emit_call(&ctx->mod, func, args, ARRAY_SIZE(args));
2249 const struct dxil_value *v_high = dxil_emit_call(&ctx->mod, func, args, ARRAY_SIZE(args));
2253 v_high = dxil_emit_binop(&ctx->mod, DXIL_BINOP_SHL, v_high,
2254 dxil_module_get_int32_const(&ctx->mod, 16), 0);
2258 v = dxil_emit_binop(&ctx->mod, DXIL_BINOP_OR, v, v_high, 0);
2287 const struct dxil_func *func = dxil_get_function(&ctx->mod, "dx.op.makeDouble", DXIL_F64);
2291 const struct dxil_value *opcode = dxil_module_get_int32_const(&ctx->mod, DXIL_INTR_MAKE_DOUBLE);
2303 const struct dxil_value *v = dxil_emit_call(&ctx->mod, func, args, ARRAY_SIZE(args));
2313 const struct dxil_func *func = dxil_get_function(&ctx->mod, "dx.op.splitDouble", DXIL_F64);
2317 const struct dxil_value *opcode = dxil_module_get_int32_const(&ctx->mod, DXIL_INTR_SPLIT_DOUBLE);
2328 const struct dxil_value *v = dxil_emit_call(&ctx->mod, func, args, ARRAY_SIZE(args));
2332 const struct dxil_value *hi = dxil_emit_extractval(&ctx->mod, v, 0);
2333 const struct dxil_value *lo = dxil_emit_extractval(&ctx->mod, v, 1);
2392 ctx->mod.feats.dx11_1_double_extensions = 1;
2410 const struct dxil_value *negative_one = dxil_module_get_int_const(&ctx->mod, val, bit_size);
2445 const struct dxil_value *one = dxil_module_get_float_const(&ctx->mod, 1.0f);
2464 ctx->mod.feats.dx11_1_double_extensions = 1;
2518 const struct dxil_value *opcode = dxil_module_get_int32_const(&ctx->mod, DXIL_INTR_CBUFFER_LOAD_LEGACY);
2526 const struct dxil_func *func = dxil_get_function(&ctx->mod, "dx.op.cbufferLoadLegacy", overload);
2529 return dxil_emit_call(&ctx->mod, func, args, ARRAY_SIZE(args));
2552 func = dxil_get_function(&ctx->mod, "dx.op.barrier", DXIL_NONE);
2556 opcode = dxil_module_get_int32_const(&ctx->mod, DXIL_INTR_BARRIER);
2560 mode = dxil_module_get_int32_const(&ctx->mod, flags);
2566 return dxil_emit_call_void(&ctx->mod, func,
2627 const struct dxil_value *idx = dxil_module_get_int32_const(&ctx->mod, i);
2651 *idx = dxil_module_get_int32_const(&ctx->mod, i);
2688 const struct dxil_value *idx = dxil_module_get_int32_const(&ctx->mod, i);
2706 dxil_get_function(&ctx->mod, name, DXIL_I32);
2711 dxil_module_get_int32_const(&ctx->mod, dxil_intr);
2717 return dxil_emit_call(&ctx->mod, func, args, ARRAY_SIZE(args));
2738 if (ctx->mod.info.has_per_sample_input) {
2739 value = dxil_emit_binop(&ctx->mod, DXIL_BINOP_AND, value,
2740 dxil_emit_binop(&ctx->mod, DXIL_BINOP_SHL,
2741 dxil_module_get_int32_const(&ctx->mod, 1),
2754 dxil_get_function(&ctx->mod, "dx.op.domainLocation", DXIL_F32);
2759 dxil_module_get_int32_const(&ctx->mod, DXIL_INTR_DOMAIN_LOCATION);
2767 const struct dxil_value *component = dxil_module_get_int32_const(&ctx->mod, component_idx);
2774 dxil_emit_call(&ctx->mod, func, args, ARRAY_SIZE(args));
2779 const struct dxil_value *value = dxil_module_get_float_const(&ctx->mod, 0.0f);
2807 const struct dxil_value *zero = dxil_module_get_int32_const(&ctx->mod, 0);
2812 return dxil_emit_gep_inbounds(&ctx->mod, ops, ARRAY_SIZE(ops));
2897 const struct dxil_value *int32_undef = get_int32_undef(&ctx->mod);
2926 dxil_emit_extractval(&ctx->mod, load, i);
2953 const struct dxil_value *int32_undef = get_int32_undef(&ctx->mod);
2966 dxil_module_get_int8_const(&ctx->mod, (1u << num_components) - 1);
2986 const struct dxil_value *int32_undef = get_int32_undef(&ctx->mod);
3010 zero = dxil_module_get_int32_const(&ctx->mod, 0);
3024 ptr = dxil_emit_gep_inbounds(&ctx->mod, ops, ARRAY_SIZE(ops));
3033 return dxil_emit_store(&ctx->mod, value, ptr, 4, false);
3039 if (!dxil_emit_atomicrmw(&ctx->mod, mask, ptr, DXIL_RMWOP_AND, false,
3044 if (!dxil_emit_atomicrmw(&ctx->mod, value, ptr, DXIL_RMWOP_OR, false,
3063 zero = dxil_module_get_int32_const(&ctx->mod, 0);
3074 ptr = dxil_emit_gep_inbounds(&ctx->mod, ops, ARRAY_SIZE(ops));
3082 return dxil_emit_store(&ctx->mod, value, ptr, 4, false);
3095 offset = dxil_module_get_int32_const(&ctx->mod, const_offset->i32 >> 4);
3098 const struct dxil_value *c4 = dxil_module_get_int32_const(&ctx->mod, 4);
3102 offset = dxil_emit_binop(&ctx->mod, DXIL_BINOP_ASHR, offset_src, c4, 0);
3111 const struct dxil_value *retval = dxil_emit_extractval(&ctx->mod, agg, i);
3137 dxil_emit_extractval(&ctx->mod, agg, i));
3160 ctx->mod.shader_kind == DXIL_HULL_SHADER);
3162 ctx->mod.shader_kind == DXIL_HULL_SHADER;
3165 const struct dxil_func *func = dxil_get_function(&ctx->mod, is_patch_constant ?
3172 const struct dxil_value *opcode = dxil_module_get_int32_const(&ctx->mod, is_patch_constant ?
3174 const struct dxil_value *output_id = dxil_module_get_int32_const(&ctx->mod, nir_intrinsic_base(intr));
3188 col = dxil_module_get_int8_const(&ctx->mod, 0);
3199 if (ctx->mod.minor_validator >= 5) {
3201 &ctx->mod.patch_consts[nir_intrinsic_base(intr)] :
3202 &ctx->mod.outputs[nir_intrinsic_base(intr)];
3219 &ctx->mod.psv_patch_consts[nir_intrinsic_base(intr)] :
3220 &ctx->mod.psv_outputs[nir_intrinsic_base(intr)];
3228 row = dxil_module_get_int32_const(&ctx->mod, i + base_component);
3230 col = dxil_module_get_int8_const(&ctx->mod, i + base_component);
3238 success &= dxil_emit_call_void(&ctx->mod, func, args, ARRAY_SIZE(args));
3249 if (ctx->mod.shader_kind == DXIL_PIXEL_SHADER &&
3258 bool is_patch_constant = (ctx->mod.shader_kind == DXIL_DOMAIN_SHADER &&
3260 (ctx->mod.shader_kind == DXIL_HULL_SHADER &&
3269 if (ctx->mod.minor_validator >= 6)
3270 ctx->mod.feats.barycentrics = 1;
3282 const struct dxil_value *opcode = dxil_module_get_int32_const(&ctx->mod, opcode_val);
3286 const struct dxil_value *input_id = dxil_module_get_int32_const(&ctx->mod,
3289 ctx->mod.input_mappings[nir_intrinsic_base(intr)]);
3302 vertex_id = dxil_module_get_int8_const(&ctx->mod, ctx->opts->provoking_vertex);
3304 const struct dxil_type *int32_type = dxil_module_get_int_type(&ctx->mod, 32);
3308 vertex_id = dxil_module_get_undef(&ctx->mod, int32_type);
3325 comp = dxil_module_get_int8_const(&ctx->mod, 0);
3332 const struct dxil_func *func = dxil_get_function(&ctx->mod, func_name, overload);
3341 if (ctx->mod.minor_validator >= 5 &&
3345 &ctx->mod.patch_consts[nir_intrinsic_base(intr)] :
3346 &ctx->mod.inputs[ctx->mod.input_mappings[nir_intrinsic_base(intr)]];
3357 &ctx->mod.psv_patch_consts[nir_intrinsic_base(intr)] :
3358 &ctx->mod.psv_inputs[ctx->mod.input_mappings[nir_intrinsic_base(intr)]];
3365 row = dxil_module_get_int32_const(&ctx->mod, i + base_component);
3367 comp = dxil_module_get_int8_const(&ctx->mod, i + base_component);
3377 const struct dxil_value *retval = dxil_emit_call(&ctx->mod, func, args, num_args);
3403 const struct dxil_value *offset_16 = dxil_emit_binop(&ctx->mod,
3404 DXIL_BINOP_MUL, float_offset, dxil_module_get_float_const(&ctx->mod, 16.0f), 0);
3405 args[i + 4] = dxil_emit_cast(&ctx->mod, DXIL_CAST_FPTOSI,
3406 dxil_module_get_int_type(&ctx->mod, 32), offset_16);
3413 args[4] = args[5] = dxil_module_get_int32_const(&ctx->mod, 0);
3429 args[0] = dxil_module_get_int32_const(&ctx->mod, opcode_val);
3430 args[1] = dxil_module_get_int32_const(&ctx->mod, nir_intrinsic_base(intr));
3433 const struct dxil_func *func = dxil_get_function(&ctx->mod, func_name, DXIL_F32);
3442 if (ctx->mod.minor_validator >= 5) {
3444 &ctx->mod.inputs[ctx->mod.input_mappings[nir_intrinsic_base(intr)]];
3453 &ctx->mod.psv_inputs[ctx->mod.input_mappings[nir_intrinsic_base(intr)]];
3459 args[3] = dxil_module_get_int8_const(&ctx->mod, i + base_component);
3461 const struct dxil_value *retval = dxil_emit_call(&ctx->mod, func, args, num_args);
3485 dxil_emit_load(&ctx->mod, ptr, 4, false);
3506 zero = dxil_module_get_int32_const(&ctx->mod, 0);
3517 ptr = dxil_emit_gep_inbounds(&ctx->mod, ops, ARRAY_SIZE(ops));
3521 retval = dxil_emit_load(&ctx->mod, ptr, align, false);
3542 zero = dxil_module_get_int32_const(&ctx->mod, 0);
3553 ptr = dxil_emit_gep_inbounds(&ctx->mod, ops, ARRAY_SIZE(ops));
3557 retval = dxil_emit_load(&ctx->mod, ptr, align, false);
3568 const struct dxil_value *opcode = dxil_module_get_int32_const(&ctx->mod, DXIL_INTR_DISCARD);
3577 const struct dxil_func *func = dxil_get_function(&ctx->mod, "dx.op.discard", DXIL_NONE);
3581 return dxil_emit_call_void(&ctx->mod, func, args, ARRAY_SIZE(args));
3597 const struct dxil_value *value = dxil_module_get_int1_const(&ctx->mod, true);
3604 const struct dxil_value *opcode = dxil_module_get_int32_const(&ctx->mod, DXIL_INTR_EMIT_STREAM);
3605 const struct dxil_value *stream_id = dxil_module_get_int8_const(&ctx->mod, nir_intrinsic_stream_id(intr));
3614 const struct dxil_func *func = dxil_get_function(&ctx->mod, "dx.op.emitStream", DXIL_NONE);
3618 return dxil_emit_call_void(&ctx->mod, func, args, ARRAY_SIZE(args));
3624 const struct dxil_value *opcode = dxil_module_get_int32_const(&ctx->mod, DXIL_INTR_CUT_STREAM);
3625 const struct dxil_value *stream_id = dxil_module_get_int8_const(&ctx->mod, nir_intrinsic_stream_id(intr));
3634 const struct dxil_func *func = dxil_get_function(&ctx->mod, "dx.op.cutStream", DXIL_NONE);
3638 return dxil_emit_call_void(&ctx->mod, func, args, ARRAY_SIZE(args));
3654 const struct dxil_value *int32_undef = get_int32_undef(&ctx->mod);
3690 dxil_module_get_int8_const(&ctx->mod, (1u << num_components) - 1);
3714 const struct dxil_value *int32_undef = get_int32_undef(&ctx->mod);
3750 const struct dxil_value *component = dxil_emit_extractval(&ctx->mod, load_result, i);
3761 ctx->mod.feats.typed_uav_load_additional_formats = true;
3781 const struct dxil_value *int32_undef = get_int32_undef(&ctx->mod);
3827 const struct dxil_value *int32_undef = get_int32_undef(&ctx->mod);
3873 const struct dxil_func *func = dxil_get_function(&ctx->mod, "dx.op.getDimensions", DXIL_NONE);
3878 dxil_module_get_int32_const(&ctx->mod, DXIL_INTR_TEXTURE_SIZE),
3883 return dxil_emit_call(&ctx->mod, func, args, ARRAY_SIZE(args));
3906 const struct dxil_value *retval = dxil_emit_extractval(&ctx->mod, dimensions, i);
3930 &ctx->mod, dxil_module_get_int_type(&ctx->mod, 32))
3937 const struct dxil_value *retval = dxil_emit_extractval(&ctx->mod, dimensions, 0);
3956 const struct dxil_value *int32_undef = get_int32_undef(&ctx->mod);
3988 const struct dxil_value *int32_undef = get_int32_undef(&ctx->mod);
4014 zero = dxil_module_get_int32_const(&ctx->mod, 0);
4025 ptr = dxil_emit_gep_inbounds(&ctx->mod, ops, ARRAY_SIZE(ops));
4033 retval = dxil_emit_atomicrmw(&ctx->mod, value, ptr, op, false,
4050 zero = dxil_module_get_int32_const(&ctx->mod, 0);
4061 ptr = dxil_emit_gep_inbounds(&ctx->mod, ops, ARRAY_SIZE(ops));
4070 retval = dxil_emit_cmpxchg(&ctx->mod, cmpval, newval, ptr, false,
4090 const struct dxil_value *index_value = dxil_module_get_int32_const(&ctx->mod, binding);
4099 index_value = dxil_emit_binop(&ctx->mod, DXIL_BINOP_ADD, index_value, offset, 0);
4105 store_dest(ctx, &intr->dest, 1, dxil_module_get_int32_const(&ctx->mod, 0), nir_type_uint32);
4159 const struct dxil_func *func = dxil_get_function(&ctx->mod, "dx.op.renderTargetGetSamplePosition", DXIL_NONE);
4163 const struct dxil_value *opcode = dxil_module_get_int32_const(&ctx->mod, DXIL_INTR_RENDER_TARGET_GET_SAMPLE_POSITION);
4174 const struct dxil_value *v = dxil_emit_call(&ctx->mod, func, args, ARRAY_SIZE(args));
4180 const struct dxil_value *coord = dxil_emit_binop(&ctx->mod, DXIL_BINOP_ADD,
4181 dxil_emit_extractval(&ctx->mod, v, i),
4182 dxil_module_get_float_const(&ctx->mod, 0.5f), 0);
4191 const struct dxil_value *layer_id = dxil_module_get_int32_const(&ctx->mod, 0);
4200 assert(ctx->mod.info.has_per_sample_input ||
4203 if (ctx->mod.info.has_per_sample_input)
4207 store_dest_value(ctx, &intr->dest, 0, dxil_module_get_int32_const(&ctx->mod, 0));
4249 switch (ctx->mod.shader_kind) {
4416 value = dxil_module_get_int1_const(&ctx->mod,
4420 ctx->mod.feats.native_low_precision = true;
4421 value = dxil_module_get_int16_const(&ctx->mod,
4425 value = dxil_module_get_int32_const(&ctx->mod,
4429 ctx->mod.feats.int64_ops = true;
4430 value = dxil_module_get_int64_const(&ctx->mod,
4474 binding = dxil_module_get_int32_const(&ctx->mod, binding_val);
4482 offset = dxil_emit_binop(&ctx->mod, DXIL_BINOP_MUL, offset,
4483 dxil_module_get_int32_const(&ctx->mod, glsl_get_aoa_size(instr->type)), 0);
4487 binding = dxil_emit_binop(&ctx->mod, DXIL_BINOP_ADD, base, offset, 0);
4531 return dxil_emit_branch(&ctx->mod, cond, true_block, false_block);
4538 return dxil_emit_branch(&ctx->mod, NULL, block, -1);
4565 const struct dxil_type *type = dxil_module_get_int_type(&ctx->mod,
4572 struct dxil_instr *phi = vphi->comp[i] = dxil_emit_phi(&ctx->mod, type);
4635 const struct dxil_func *func = dxil_get_function(&ctx->mod, "dx.op.sample", params->overload);
4640 dxil_module_get_int32_const(&ctx->mod, DXIL_INTR_SAMPLE),
4647 return dxil_emit_call(&ctx->mod, func, args, ARRAY_SIZE(args));
4653 const struct dxil_func *func = dxil_get_function(&ctx->mod, "dx.op.sampleBias", params->overload);
4660 dxil_module_get_int32_const(&ctx->mod, DXIL_INTR_SAMPLE_BIAS),
4667 return dxil_emit_call(&ctx->mod, func, args, ARRAY_SIZE(args));
4673 const struct dxil_func *func = dxil_get_function(&ctx->mod, "dx.op.sampleLevel", params->overload);
4680 dxil_module_get_int32_const(&ctx->mod, DXIL_INTR_SAMPLE_LEVEL),
4687 return dxil_emit_call(&ctx->mod, func, args, ARRAY_SIZE(args));
4697 if (ctx->mod.shader_kind == DXIL_PIXEL_SHADER) {
4698 func = dxil_get_function(&ctx->mod, "dx.op.sampleCmp", DXIL_F32);
4702 func = dxil_get_function(&ctx->mod, "dx.op.sampleCmpLevelZero", DXIL_F32);
4711 dxil_module_get_int32_const(&ctx->mod, opcode),
4718 return dxil_emit_call(&ctx->mod, func, args, numparam);
4724 const struct dxil_func *func = dxil_get_function(&ctx->mod, "dx.op.sampleGrad", params->overload);
4729 dxil_module_get_int32_const(&ctx->mod, DXIL_INTR_SAMPLE_GRAD),
4738 return dxil_emit_call(&ctx->mod, func, args, ARRAY_SIZE(args));
4744 const struct dxil_func *func = dxil_get_function(&ctx->mod, "dx.op.textureLoad", params->overload);
4749 params->lod_or_sample = dxil_module_get_undef(&ctx->mod, dxil_module_get_int_type(&ctx->mod, 32));
4752 dxil_module_get_int32_const(&ctx->mod, DXIL_INTR_TEXTURE_LOAD),
4758 return dxil_emit_call(&ctx->mod, func, args, ARRAY_SIZE(args));
4764 const struct dxil_func *func = dxil_get_function(&ctx->mod, "dx.op.calculateLOD", DXIL_F32);
4769 dxil_module_get_int32_const(&ctx->mod, DXIL_INTR_TEXTURE_LOD),
4775 dxil_module_get_int1_const(&ctx->mod, clamped ? 1 : 0)
4778 return dxil_emit_call(&ctx->mod, func, args, ARRAY_SIZE(args));
4784 const struct dxil_func *func = dxil_get_function(&ctx->mod,
4790 dxil_module_get_int32_const(&ctx->mod, params->cmp ?
4800 dxil_module_get_int32_const(&ctx->mod, component),
4804 return dxil_emit_call(&ctx->mod, func, args, ARRAY_SIZE(args) - (params->cmp ? 0 : 1));
4817 const struct dxil_type *int_type = dxil_module_get_int_type(&ctx->mod, 32);
4818 const struct dxil_type *float_type = dxil_module_get_float_type(&ctx->mod, 32);
4819 const struct dxil_value *int_undef = dxil_module_get_undef(&ctx->mod, int_type);
4820 const struct dxil_value *float_undef = dxil_module_get_undef(&ctx->mod, float_type);
4914 dxil_emit_binop(&ctx->mod, DXIL_BINOP_ADD,
4916 dxil_module_get_int32_const(&ctx->mod, instr->texture_index), 0),
4924 dxil_emit_binop(&ctx->mod, DXIL_BINOP_ADD,
4926 dxil_module_get_int32_const(&ctx->mod, instr->sampler_index), 0),
4960 } else if (ctx->mod.shader_kind == DXIL_PIXEL_SHADER) {
4964 params.lod_or_sample = dxil_module_get_float_const(&ctx->mod, 0);
5003 params.lod_or_sample = dxil_module_get_int_const(&ctx->mod, 0, 32);
5005 const struct dxil_value *retval = dxil_emit_extractval(&ctx->mod, sample, 3);
5018 const struct dxil_value *retval = dxil_emit_extractval(&ctx->mod, sample, i);
5029 store_ssa_def(ctx, &undef->def, i, dxil_module_get_int32_const(&ctx->mod, 0));
5063 assert(block->index < ctx->mod.cur_emitting_func->num_basic_block_ids);
5064 ctx->mod.cur_emitting_func->basic_block_ids[block->index] = ctx->mod.cur_emitting_func->curr_block;
5215 const struct dxil_value *dummy = dxil_module_get_int_const(&ctx->mod, 0, bitsize);
5267 const struct dxil_type *int32 = dxil_module_get_int_type(&ctx->mod, 32);
5268 const struct dxil_value *array_length = dxil_module_get_int32_const(&ctx->mod, size / sizeof(uint32_t));
5273 &ctx->mod, int32, size / sizeof(uint32_t));
5277 ctx->scratchvars = dxil_emit_alloca(&ctx->mod, type, int32, array_length, 4);
5329 const struct dxil_type *void_type = dxil_module_get_void_type(&ctx->mod);
5330 const struct dxil_type *func_type = dxil_module_add_function_type(&ctx->mod, void_type, NULL, 0);
5331 struct dxil_func_def *func_def = dxil_add_function_def(&ctx->mod, func->name, func_type, impl->num_blocks);
5366 if (!dxil_emit_ret_void(&ctx->mod))
5438 type = dxil_module_get_array_type(&ctx->mod,
5439 dxil_module_get_int_type(&ctx->mod, 32),
5441 ctx->sharedvars = dxil_add_global_ptr_var(&ctx->mod, "shared", type,
5499 ctx->mod.info.has_per_sample_input =
5501 if (!ctx->mod.info.has_per_sample_input && ctx->shader->info.stage == MESA_SHADER_FRAGMENT) {
5504 ctx->mod.info.has_per_sample_input = true;
5510 unsigned input_clip_size = ctx->mod.shader_kind == DXIL_PIXEL_SHADER ?
5512 preprocess_signatures(&ctx->mod, ctx->shader, input_clip_size);
5522 ctx->mod.feats.stencil_ref = true;
5529 ctx->mod.feats.array_layer_from_vs_or_ds = true;
5532 if (ctx->mod.feats.native_low_precision)
5533 ctx->mod.minor_version = MAX2(ctx->mod.minor_version, 2);
5536 dxil_emit_module(&ctx->mod);
5631 unsigned resource_element_size = ctx->mod.minor_validator >= 6 ?
5636 state->state.psv1.shader_stage = (uint8_t)ctx->mod.shader_kind;
5637 state->state.psv1.sig_input_elements = (uint8_t)ctx->mod.num_sig_inputs;
5638 state->state.psv1.sig_output_elements = (uint8_t)ctx->mod.num_sig_outputs;
5639 state->state.psv1.sig_patch_const_or_prim_elements = (uint8_t)ctx->mod.num_sig_patch_consts;
5641 switch (ctx->mod.shader_kind) {
5643 state->state.psv1.psv0.vs.output_position_present = ctx->mod.info.has_out_position;
5647 state->state.psv1.psv0.ps.depth_output = ctx->mod.info.has_out_depth;
5649 ctx->mod.info.has_per_sample_input;
5661 state->state.psv1.psv0.gs.output_position_present = ctx->mod.info.has_out_position;
5668 state->state.psv1.sig_patch_const_or_prim_vectors = ctx->mod.num_psv_patch_consts;
5673 state->state.psv1.psv0.ds.output_position_present = ctx->mod.info.has_out_position;
5674 state->state.psv1.sig_patch_const_or_prim_vectors = ctx->mod.num_psv_patch_consts;
5846 dxil_module_init(&ctx->mod, ctx->ralloc_ctx);
5847 ctx->mod.shader_kind = get_dxil_shader_kind(s);
5848 ctx->mod.major_version = 6;
5849 ctx->mod.minor_version = 1;
5850 ctx->mod.major_validator = validator_version >> 16;
5851 ctx->mod.minor_validator = validator_version & 0xffff;
5874 if (ctx->mod.shader_kind == DXIL_HULL_SHADER)
5877 if (ctx->mod.shader_kind == DXIL_HULL_SHADER ||
5878 ctx->mod.shader_kind == DXIL_DOMAIN_SHADER) {
5904 assert(ctx->mod.major_version == 6 && ctx->mod.minor_version >= 1);
5905 if ((ctx->mod.major_version << 16 | ctx->mod.minor_version) > opts->shader_model_max) {
5911 assert(ctx->mod.major_validator == 1);
5912 if (!dxil_validator_can_validate_shader_model(ctx->mod.minor_version, ctx->mod.minor_validator)) {
5920 dxil_dump_module(dumper, &ctx->mod);
5929 if (!dxil_container_add_features(&container, &ctx->mod.feats)) {
5937 ctx->mod.num_sig_inputs,
5938 ctx->mod.inputs,
5939 ctx->mod.minor_validator >= 7)) {
5947 ctx->mod.num_sig_outputs,
5948 ctx->mod.outputs,
5949 ctx->mod.minor_validator >= 7)) {
5955 if ((ctx->mod.shader_kind == DXIL_HULL_SHADER ||
5956 ctx->mod.shader_kind == DXIL_DOMAIN_SHADER) &&
5959 ctx->mod.num_sig_patch_consts,
5960 ctx->mod.patch_consts,
5961 ctx->mod.minor_validator >= 7)) {
5971 if (!dxil_container_add_state_validation(&container,&ctx->mod,
5978 if (!dxil_container_add_module(&container, &ctx->mod)) {
5995 get_shader_kind_str(ctx->mod.shader_kind), shader_id++);
6005 dxil_module_release(&ctx->mod);