Lines Matching refs:metadata
45 struct clc_dxil_metadata *metadata;
64 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;
404 struct clc_dxil_metadata *metadata = &dxil->metadata;
429 struct clc_dxil_metadata *metadata = &dxil->metadata;
856 struct clc_dxil_metadata *metadata = &out_dxil->metadata;
858 metadata->args = calloc(out_dxil->kernel->num_args,
859 sizeof(*metadata->args));
860 if (!metadata->args) {
953 metadata->printf.info_count = nir->printf_info_count;
954 metadata->printf.infos = calloc(nir->printf_info_count, sizeof(struct clc_printf_info));
956 metadata->printf.infos[i].str = malloc(nir->printf_info[i].string_size);
957 memcpy(metadata->printf.infos[i].str, nir->printf_info[i].strings, nir->printf_info[i].string_size);
958 metadata->printf.infos[i].num_args = nir->printf_info[i].num_args;
959 metadata->printf.infos[i].arg_sizes = malloc(nir->printf_info[i].num_args * sizeof(unsigned));
960 memcpy(metadata->printf.infos[i].arg_sizes, nir->printf_info[i].arg_sizes, nir->printf_info[i].num_args * sizeof(unsigned));
975 // Calculate input offsets/metadata.
984 metadata->args[i].offset = var->data.driver_location;
985 metadata->args[i].size = size;
986 metadata->kernel_inputs_buf_size = MAX2(metadata->kernel_inputs_buf_size,
992 metadata->args[i].globconstptr.buf_id = uav_id++;
1002 metadata->args[i].sampler.sampler_id = var->data.binding = sampler_id++;
1017 metadata->args[i].image.buf_ids[0] = srv_id++;
1020 metadata->args[i].image.buf_ids[0] = uav_id++;
1023 metadata->args[i].image.num_buf_ids = 1;
1024 var->data.binding = metadata->args[i].image.buf_ids[0];
1032 // Fill out inline sampler metadata, now that they've been deduped and dead ones removed
1045 assert(metadata->num_const_samplers < CLC_MAX_SAMPLERS);
1046 metadata->const_samplers[metadata->num_const_samplers].sampler_id = var->data.binding;
1047 metadata->const_samplers[metadata->num_const_samplers].addressing_mode = var->data.sampler.addressing_mode;
1048 metadata->const_samplers[metadata->num_const_samplers].normalized_coords = var->data.sampler.normalized_coordinates;
1049 metadata->const_samplers[metadata->num_const_samplers].filter_mode = var->data.sampler.filter_mode;
1050 metadata->num_const_samplers++;
1056 struct clc_image_lower_context image_lower_context = { metadata, &srv_id, &uav_id };
1076 metadata->printf.uav_id = has_printf ? uav_id++ : -1;
1109 memcpy(metadata->local_size, nir->info.workgroup_size,
1110 sizeof(metadata->local_size));
1111 memcpy(metadata->local_size_hint, nir->info.cs.workgroup_size_hint,
1112 sizeof(metadata->local_size));
1129 memcpy(metadata->local_size, nir->info.workgroup_size,
1130 sizeof(metadata->local_size));
1188 metadata->args[i].localptr.sharedmem_offset = nir->info.shared_size;
1192 metadata->local_mem_size = nir->info.shared_size;
1193 metadata->priv_mem_size = nir->scratch_size;
1219 metadata->consts[metadata->num_consts].data = data;
1220 metadata->consts[metadata->num_consts].size = size;
1221 metadata->consts[metadata->num_consts].uav_id = var->data.binding;
1222 metadata->num_consts++;
1228 metadata->kernel_inputs_cbv_id = inputs_var ? inputs_var->data.binding : 0;
1229 metadata->work_properties_cbv_id = work_properties_var->data.binding;
1230 metadata->num_uavs = uav_id;
1231 metadata->num_srvs = srv_id;
1232 metadata->num_samplers = sampler_id;
1248 for (unsigned i = 0; i < dxil->metadata.num_consts; i++)
1249 free(dxil->metadata.consts[i].data);
1251 for (unsigned i = 0; i < dxil->metadata.printf.info_count; i++) {
1252 free(dxil->metadata.printf.infos[i].arg_sizes);
1253 free(dxil->metadata.printf.infos[i].str);
1255 free(dxil->metadata.printf.infos);