CUDA заменяет __syncthreads вместо __threadfence ()

Question

Я скопировал приведенный ниже код из руководства NVIDIA, например: для __threadfence().Почему они использовали __threadfence() в приведенном ниже коде.Я думаю, что использование __syncthreads() вместо __threadfence() даст вам тот же результат.

Может кто-нибудь объяснить разницу между звонками __syncthreads() и __threadfence()?

__device__ unsigned int count = 0;
__shared__ bool isLastBlockDone;

__global__ void sum(const float* array, unsigned int N,float* result)
{
    // Each block sums a subset of the input array
    float partialSum = calculatePartialSum(array, N);

    if (threadIdx.x == 0) {
        // Thread 0 of each block stores the partial sum
        // to global memory
        result[blockIdx.x] = partialSum;

        // Thread 0 makes sure its result is visible to
        // all other threads
        __threadfence();

        // Thread 0 of each block signals that it is done
        unsigned int value = atomicInc(&count, gridDim.x);

        // Thread 0 of each block determines if its block is
        // the last block to be done
        isLastBlockDone = (value == (gridDim.x - 1));
    }

    // Synchronize to make sure that each thread reads
    // the correct value of isLastBlockDone
    __syncthreads();

    if (isLastBlockDone) 
    {
        // The last block sums the partial sums
        // stored in result[0 .. gridDim.x-1]
        float totalSum = calculateTotalSum(result);

        if (threadIdx.x == 0)
        {
            // Thread 0 of last block stores total sum
            // to global memory and resets count so that
            // next kernel call works properly
            result[0] = totalSum;
            count = 0;
        }
    }
}

Answer 1

С точки зрения общей памяти __syncthreads() просто сильнее, чем __threadfence(). Что касается глобальной памяти - это две разные вещи.

• __threadfence_block() останавливает текущий поток, пока все записи в общую память не будут видны другим потокам из того же блока. Это предотвращает оптимизацию компилятора путем кэширования записей в общую память в регистрах. не синхронизирует потоки, и не обязательно, чтобы все потоки фактически достигли этой инструкции.
• __threadfence() останавливает текущий поток, пока все записи в общую и глобальную память не будут видны всем другим потокам.
• __syncthreads() должен быть достигнут всеми потоками из блока (например, без расходящихся if операторов) и обеспечивает выполнение кода, предшествующего инструкции до инструкций, следующих за ней, для всех потоков в блок.

В вашем конкретном случае инструкция __threadfence() используется, чтобы убедиться, что записи в глобальный массив result видны всем. __syncthreads() будет просто синхронизировать потоки только в текущем блоке, без принудительной записи в глобальную память для другого блока. Более того, в тот момент кода, который находится внутри ветви if, только один поток выполняет этот код; использование __syncthreads() приведет к неопределенному поведению графического процессора, что, скорее всего, приведет к полной десинхронизации ядра.

Ознакомьтесь со следующими главами Руководства по программированию в CUDA C:

• 3.2.2 «Общая память» - пример умножения матриц
• 5.4.3 «Инструкция по синхронизации»
• B.2.5 "летучий"
• B.5 «Функции забора памяти»