CUDA объединяет независимые от потока (??) переменные во время выполнения - PullRequest
1 голос
/ 14 апреля 2011

Ребята прошу прощения, если название сбивает с толку. Я хоть долго и трудно, и не мог придумать надлежащим образом сформулировать вопрос в одной строке. Итак, здесь более подробно. Я делаю базовое вычитание изображения, когда второе изображение было изменено, и мне нужно найти отношение того, сколько изменений было сделано для изображения. для этого я использовал следующий код. Оба изображения имеют размер 128х1024.

for(int i = 0; i < 128; i++)
{
    for(int j = 0; j < 1024; j++)
    {
        den++;
        diff[i * 1024 + j] = orig[i * 1024 + j] - modified[i * 1024 + j];
        if(diff[i * 1024 + j] < error)
        {
            num++;
        }
    }
}
ratio = num/den;

Приведенный выше код прекрасно работает на процессоре, но я хочу попробовать сделать это на CUDA. Для этого я могу настроить CUDA на базовое вычитание изображений (см. Код ниже), но не могу понять, как сделать условный оператор if для получения моего соотношения.

__global__ void calcRatio(float *orig, float *modified, int size, float *result)
{
    int index = threadIdx.x + blockIdx.x * blockDim.x;
    if(index < size)
        result[index] = orig[index] - modified[index];
}

Итак, до этого момента это работает, но я не могу понять, как распараллелить счетчики num и den в каждом потоке, чтобы вычислить соотношение в конце всех выполнений потока. Мне кажется, что счетчики num и den независимы от потоков, так как каждый раз, когда я пытаюсь их использовать, кажется, что они увеличиваются только один раз.

Буду признателен за любую помощь, так как я только начинаю в CUDA, и каждый пример, который я вижу онлайн, никогда не подходит для того, что мне нужно делать.

РЕДАКТИРОВАТЬ: Исправлен мой наивный код. Забыл ввести одно из основных условий в коде. Это был длинный длинный день.

for(int i = 0; i < 128; i++)
{
    for(int j = 0; j < 1024; j++)
    {
        if(modified[i * 1024 + j] < 400.0)  //400.0 threshold value to ignore noise
        {
            den++;  
            diff[i * 1024 + j] = orig[i * 1024 + j] - modified[i * 1024 + j];
            if(diff[i * 1024 + j] < error)
            {
                num++;
            }
        }
    }
}
ratio = num/den;

Ответы [ 2 ]

4 голосов
/ 14 апреля 2011

Операция, которую необходимо использовать для выполнения глобального суммирования во всех потоках, называется «параллельным сокращением». Хотя вы могли бы использовать атомарные операции для этого, я бы не рекомендовал это. В CUDA SDK есть сокращенное ядро ​​и очень хорошая статья, в которой обсуждается техника, которую стоит прочитать.

Если бы я писал код, чтобы делать то, что вы хотите, он, вероятно, выглядел бы так:

template <int blocksize>
__global__ void calcRatio(float *orig, float *modified, int size, float *result, 
                            int *count, const float error)
{
    __shared__ volatile float buff[blocksize];

    int index = threadIdx.x + blockIdx.x * blockDim.x;
    int stride = blockDim.x * gridDim.x;

    int count = 0;
    for(int i=index; i<n; i+=stride) {
        val = orig[index] - modified[index];
        count += (val < error);
        result[index] = val;
    }

    buff[threadIdx.x] = count;
    __syncthreads();


    // Parallel reduction in shared memory using 1 warp
    if (threadId.x < warpSize) {

        for(int i=threadIdx.x + warpSize; i<blocksize; i+= warpSize) {
            buff[threadIdx.x] += buff[i];

        if (threadIdx.x < 16) buff[threadIdx.x] +=buff[threadIdx.x + 16];
        if (threadIdx.x < 8)  buff[threadIdx.x] +=buff[threadIdx.x + 8];
        if (threadIdx.x < 4)  buff[threadIdx.x] +=buff[threadIdx.x + 4];
        if (threadIdx.x < 2)  buff[threadIdx.x] +=buff[threadIdx.x + 2];
        if (threadIdx.x == 0) count[blockIdx.x] = buff[0] + buff[1];
    }
}

Первый раздел выполняет то, что делает ваш последовательный код - вычисляет разницу и локальный поток общее количество элементов, которые меньше ошибки. Примечание. Я написал эту версию так, что каждый поток предназначен для обработки более одной записи входных данных. Это было сделано, чтобы помочь компенсировать вычислительные затраты на параллельное сокращение, которое следует, и идея состоит в том, что вы будете использовать меньше блоков и потоков, чем было во входных наборах данных.

Второй раздел - это само сокращение, выполняемое в общей памяти. По сути, это «древовидная» операция, в которой размер набора локальных промежуточных итогов потока в пределах одного блока потоков сначала суммируется до 32 промежуточных итогов, затем промежуточные итоги объединяются до тех пор, пока не будет получен окончательный промежуточный итог для блока, и что затем сохраняется сумма для блока . Вы получите небольшой список промежуточных итогов, по одному для каждого блока, который вы запустили, который можно скопировать обратно на хост и рассчитать там конечный результат, который вам нужен.

Обратите внимание, что я закодировал это в браузере и не скомпилировал его, могут быть ошибки, но это должно дать представление о том, как будет работать "продвинутая" версия того, что вы пытаетесь сделать.

0 голосов
/ 14 апреля 2011

Знаменатель довольно прост, так как это просто размер.

Числитель более проблематичен, поскольку его значение для данного потока зависит от всех предыдущих значений.Вам придется выполнять эту операцию последовательно.

Вероятно, вам нужно найти atomicAdd.Хотя это очень медленно.

Я думаю, вы найдете этот вопрос актуальным.Ваш номер в основном глобальные данные. CUDA сумма от массива к массиву

Кроме того, вы можете записать результаты проверки на ошибки в массив.Подсчет результатов можно затем распараллелить.Это было бы немного сложно, но я думаю, что что-то вроде этого увеличится: http://tekpool.wordpress.com/2006/09/25/bit-count-parallel-counting-mit-hakmem/

...