[ENHANCEMENT]: `__device__ for_each` that provides a mutable reference rather than a copy of the slot to the callback function

Abierto
  • ¿Algo no está bien?
#639 0 comentarios 0 reacciones 0 asignados Ver en GitHub

Los mantenedores suelen responder en 1 día

Nadie ha tomado este issue todavía.

Evaluación

Dificultad
5/5
Tiempo estimado
Más de una semana
Aptitud para principiantes
35/100
Tipo de issue
Nueva funcionalidad
Claridad
Bastante claro
Estado de actividad
Estancado
Stack tecnológico
cpp
Área
api

Línea de trabajo

Comienza revisando la API de dispositivo for_each actual y la maquinaria de cuco::detail utilizada por la solución alternativa basada en custom-kernel. Define las opciones de copia y referencia mutable, incluidas sus implicaciones de seguridad, y considera cómo debería cubrirse el comportamiento actual de la API.

Escrito por el modelo de indexación a partir del texto del issue.

Descripción

type: feature request
Is your feature request related to a problem? Please describe.

Feature request on behalf of @LutzCle who is implementing a mark join for an internal project.

tl;dr The algorithm requires setting a "(matching) slot has been seen" bit during probing which is explicitly excluded from the hash value generation and key comparison, so it doesn't break the probing sequence if set or unset.

Our initial approach was to use the for_each device API and set the bit for each matching element but unfortunately this doesn't work since the current for_each implementation returns a copy of the matching slot rather than a mutable reference.

Describe the solution you'd like

Provide both copy and mutable reference options for for_each.

Naming ideas: for_each(_reference)/for_each_copy/for_each_mutable, ...

Describe alternatives you've considered

The current workaround consists of basically re-implementing for_each with the proposed feature in a custom kernel which involves using some of the internal cuco::detail machinery - not ideal.

Additional context

I'm not sure if this feature is a foot gun or not so I would like to spark a discussion around the potential dangers.

Lenguaje dominante
Cuda
Estrellas
675
Forks
123
Merge medio
22 h 39 min
PR fusionados (30 d)
8

Preparar el entorno

Abrir en Codespaces

Inicia el contenedor de desarrollo del proyecto en tu navegador, con tu propia cuenta de GitHub.

Primeros pasos

  1. Lee el issue completo y luego la guía de contribución del proyecto.
  2. Comenta en el issue que vas a ocuparte — evita que dos personas hagan lo mismo.
  3. Haz un fork del repositorio y trabaja en una rama.
  4. Abre un pull request que haga referencia al número del issue.

Más de NVIDIA/cuCollections

Todos los issues de NVIDIA/cuCollections

Issues similares

Más issues de Backend & API Design

Recibe los nuevos issues en tu correo

Un resumen breve de issues de GitHub para principiantes.