Lines Matching defs:metadata
45 struct clc_dxil_metadata *metadata;
65 context->metadata->args[context->metadata_index].
131 while (context->metadata->args[context->metadata_index].image.buf_ids[0] != in_var->data.binding)
251 context->metadata->args[context->metadata_index].image.num_buf_ids = context->num_buf_ids;
807 struct clc_dxil_metadata *metadata = &out_dxil->metadata;
809 metadata->args = calloc(out_dxil->kernel->num_args,
810 sizeof(*metadata->args));
811 if (!metadata->args) {
900 metadata->printf.info_count = nir->printf_info_count;
901 metadata->printf.infos = calloc(nir->printf_info_count, sizeof(struct clc_printf_info));
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));
922 // Calculate input offsets/metadata.
931 metadata->args[i].offset = var->data.driver_location;
932 metadata->args[i].size = size;
933 metadata->kernel_inputs_buf_size = MAX2(metadata->kernel_inputs_buf_size,
937 metadata->args[i].globconstptr.buf_id = uav_id++;
947 metadata->args[i].sampler.sampler_id = var->data.binding = sampler_id++;
963 metadata->args[i].image.buf_ids[0] = srv_id++;
966 metadata->args[i].image.buf_ids[0] = uav_id++;
969 metadata->args[i].image.num_buf_ids = 1;
970 var->data.binding = metadata->args[i].image.buf_ids[0];
973 var->data.driver_location = metadata->kernel_inputs_buf_size;
974 metadata->args[i].offset = metadata->kernel_inputs_buf_size;
975 metadata->args[i].size = 8;
976 metadata->kernel_inputs_buf_size += metadata->args[i].size;
984 // Fill out inline sampler metadata, now that they've been deduped and dead ones removed
997 assert(metadata->num_const_samplers < CLC_MAX_SAMPLERS);
998 metadata->const_samplers[metadata->num_const_samplers].sampler_id = var->data.binding;
999 metadata->const_samplers[metadata->num_const_samplers].addressing_mode = var->data.sampler.addressing_mode;
1000 metadata->const_samplers[metadata->num_const_samplers].normalized_coords = var->data.sampler.normalized_coordinates;
1001 metadata->const_samplers[metadata->num_const_samplers].filter_mode = var->data.sampler.filter_mode;
1002 metadata->num_const_samplers++;
1008 struct clc_image_lower_context image_lower_context = { metadata, &srv_id, &uav_id };
1028 metadata->printf.uav_id = has_printf ? uav_id++ : -1;
1061 memcpy(metadata->local_size, nir->info.workgroup_size,
1062 sizeof(metadata->local_size));
1063 memcpy(metadata->local_size_hint, nir->info.cs.workgroup_size_hint,
1064 sizeof(metadata->local_size));
1081 memcpy(metadata->local_size, nir->info.workgroup_size,
1082 sizeof(metadata->local_size));
1142 metadata->args[i].localptr.sharedmem_offset = nir->info.shared_size;
1146 metadata->local_mem_size = nir->info.shared_size;
1147 metadata->priv_mem_size = nir->scratch_size;
1173 metadata->consts[metadata->num_consts].data = data;
1174 metadata->consts[metadata->num_consts].size = size;
1175 metadata->consts[metadata->num_consts].uav_id = var->data.binding;
1176 metadata->num_consts++;
1182 metadata->kernel_inputs_cbv_id = inputs_var ? inputs_var->data.binding : 0;
1183 metadata->work_properties_cbv_id = work_properties_var->data.binding;
1184 metadata->num_uavs = uav_id;
1185 metadata->num_srvs = srv_id;
1186 metadata->num_samplers = sampler_id;
1202 for (unsigned i = 0; i < dxil->metadata.num_consts; i++)
1203 free(dxil->metadata.consts[i].data);
1205 for (unsigned i = 0; i < dxil->metadata.printf.info_count; i++) {
1206 free(dxil->metadata.printf.infos[i].arg_sizes);
1207 free(dxil->metadata.printf.infos[i].str);
1209 free(dxil->metadata.printf.infos);