atomicité de mise à jour de mémoire GPU

J'ai une mémoire globale principale (gpu_mem), avec une variable (gpu_mem_offset) pour suivre le décalage actuel de cette mémoire globale où un thread mettra à jour son contenu. Le script suivant garantira-t-il l'atomicité du contenu de la mémoire et des valeurs de décalage?atomicité de mise à jour de mémoire GPU

__global__ void kernel(int *gpu_mem, int *gpu_mem_offset) 
{ 
    gpu_mem[(*gpu_mem_offset)++] = some_value; 
}

Mise à jour

Je réalise rapidement que le script ci-dessus est stupide, mais je ne décrirais mon intention: comment garantir l'atomicité à la fois du contenu de la mémoire et les valeurs de décalage? Ou, comment les verrouiller tous les deux en même temps? Peut-être pas réalisable?

Source

2016-03-07 Hailiang Zhang

En un mot, non, chaque fois que vous exécutez plus d'un seul thread – talonmies

La seule façon d'assurer une mise à jour cohérente à la fois contre et tableau dans cet exemple est comme ceci:

__global__ void kernel(int *gpu_mem, int *gpu_mem_offset) 
{ 
    int offset = atomicAdd(gpu_mem_offset, 1); 
    gpu_mem[offset] = some_value; 
}

à savoir si vous avez besoin des mises à jour atomiques, utilisez alors une intrinsèque atomique. C'est ce qu'ils sont pour. Ici, l'accès atomique à gpu_mem_offset garantit que chaque thread obtient une valeur unique du décalage. L'écriture est alors garantie, car chaque thread accède à un index unique.

Source

2016-03-07 20:47:30 talonmies

C'est tellement évident. Désolé que je sois coincé. –

Gardez à l'esprit que cette approche ne garantit que la cohérence de l'écriture. Si plusieurs threads font la même chose, ils * écriront * à des cellules différentes. Toutefois, cela échouera si vous essayez de lire à partir du même tableau avec des threads différents, par ex. lorsque vous essayez de mettre en œuvre la file d'attente FIFO de vol de travail. Pour ce faire, vous devrez en faire plus ... – CygnusX1

atomicité de mise à jour de mémoire GPU

Répondre

Questions connexes