Согласно руководство по программированию , __syncthreads() является одновременно барьером для выполнения и ограничением памяти:
ожидает, пока все потоки в блоке потоков достигнут этой точки и все обращения к глобальной и совместно используемой памяти, сделанные этими потоками до __syncthreads(), будут видны всем потокам в блоке .
Функция ограждения памяти (т. Е. «Видимость») «заставляет» все обновления в общей и глобальной памяти быть видимыми для других потоков.
Полагаю, это то, о чем вы спрашиваете. Я не думаю, что делать общие заявления типа «вам не нужно использовать volatile при использовании __syncthreads()» - разумная идея. Это будет зависеть от кода. Но в некоторых ситуациях, например, при классическом параллельном сокращении , использование __syncthreads() на каждом шаге сокращения по всему блоку будет означать, что совместно используемая память, используемая для такого сокращения, не должна быть отмечена volatile.
Поскольку __syncthreads() является одновременно барьером для выполнения и ограничением памяти, существуют определенные утверждения, которые мы можем сделать относительно использования __syncthreads(), которые не будут применимы только к использованию только __threadfence().
Предположим, у меня есть этот код:
__global__ void k(int *data){
...
*data = 1;
__syncthreads();
if (*data == 1){
...}
...
}
В этом случае любой поток в конкретном блоке, выполняющий оператор if, гарантированно увидит *data как 1. Для этого есть два компонента:
__syncthreads() - забор памяти для всего устройства. Это заставляет любой поток, который записал значение, сделать это значение видимым. Это фактически означает, что, поскольку это ограничение памяти для всего устройства, записанное значение, по крайней мере, заполнило кэш-память второго уровня (который является устройством вставки в глобальную память для всего устройства, фактически прокси для глобальной памяти).
__syncthreads() - барьер исполнения (для всего потока). Это заставляет все нити достигать барьера прежде, чем любой сможет продолжить. Такое поведение порядка выполнения означает, что к тому моменту, когда любой поток выполнит вышеуказанный оператор if, гарантия в пункте 1 выше вступит в силу.
Обратите внимание, что здесь есть тонкое различие. Другие потоки в других блоках в других точках кода могут видеть или не видеть значение, записанное другим блоком.
Только когда мы объединили синхронизацию выполнения и ограждение памяти, мы можем быть уверены, что значения, заполненные одним потоком, действительно видны другому потоку. А без использования кооперативных групп CUDA не предоставляет механизма для синхронизации выполнения между отдельными блоками.
__threadfence() само по себе делает значение в конечном итоге видимым, но без понимания относительного порядка выполнения между потоком записи и потоком чтения невозможно дать гарантии только на основе проверки кода.
Аналогично volatile гарантирует что-то похожее на __threadfence() (для потока записи), но также несколько отличается. __threadfence() гарантирует, что записывающий поток в конечном итоге отправит свои данные в L2 (т.е. сделает его видимым). volatile делает что-то подобное, но также гарантирует, что поток чтения не будет читать «устаревшую копию» в L1, но перейдет к L2 (по крайней мере), чтобы извлечь текущее значение, каждый раз, когда в коде происходит чтение этого значения .
Обратите внимание, что никогда не происходит «аннулирование» данных кэша L1, вызванных действиями кода устройства на другом SM. volatile эффективно гарантирует, что нагрузка будет обходить L1. volatile также гарантирует, что магазин перейдет непосредственно к L2. __threadfence() делает что-то похожее на последнее (по крайней мере, из-за того, что поток вышел за пределы __threadfence()), но не дает никаких гарантий о состоянии L1 в других SM или о том, как потоки в других SM будут читать значение.