Skip to content

Explore kernel sizes of refine_*.cu #922

@divyegala

Description

@divyegala
  1. refine_device_float_float.cu.o: ~45 MB
  2. refine_device_int8_t_float.cu.o: ~26 MB
  3. refine_device_uint8_t_float.cu.o: ~26 MB
  4. refine_device_half_float.cu.o: ~24 MB

refine_device.cuh seemingly instantiates no kernels, but it is likely instantiating kernels from calling into detail of other files. There may be an opportunity here to consolidate the kernel instantiations to just once.

  1. refine and ivf_flat::extend both call into build_index_kernel
    build_index_kernel<T, IdxT, LabelT, true>
    <<<grid_dim, block_dim, 0, stream>>>(new_labels.data(),
    dataset,
    candidate_idx,
    refinement_index->data_ptrs().data_handle(),
    refinement_index->inds_ptrs().data_handle(),
    list_sizes_ptr,
    n_queries * n_candidates,
    refinement_index->dim(),
    refinement_index->veclen());
  2. refine and ivf_flat::search both call into ivfflat_interleaved_scan kernels
  3. refine, ivf_flat::search, and ivf_pq::search all call into postprocess_neighbors_kernel
    postprocess_neighbors_kernel(IdxT* neighbors_out, // [n_queries, topk]

Metadata

Metadata

Assignees

Labels

No labels
No labels

Type

Projects

Status

Done

Milestone

No milestone

Relationships

None yet

Development

No branches or pull requests

Issue actions