цифровая интеграция; CUDA-разработка

Мне нужен совет о том, как действовать и использовать вычислительную мощность устройства CUDA для численного интегрирования функции. Ниже приведена некоторая информация о моем устройстве (не имеет значения)

Аппаратное обеспечение

 Geforce GTX470; Compute Capability 2.0

описание проблемы

У меня есть функция вроде

g(x) = x * f(x, a, b, c)

Что мне нужно интегрировать как указано уравнение

Теперь я уже написал функцию интегрирования, которая просто берет g(x), разбивает интервал на N подинтервалов, вычисляет результат для отдельного подинтервала, а затем я суммирую его на процессоре. Для завершения я привожу ниже пример кода.

__device__ float function(float x, float a, float b, float c) {
   // do some complex calculation
   return result;
}
__global__ void kernel(float *d_arr, float a, float b, float c, int N) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    float x = (float)idx / (float)N;

    if (idx < N)  {
       d_arr[idx] = x * function(x, a, b, c);
    }
}

Приведенный выше код предназначен только для демонстрационных целей, на самом деле я использую метод Ромберга для интеграции моего g(x), но идея та же. Моя настоящая проблема возникает из-за того, что у меня нет только одного набора значений (a, b, c), у меня есть несколько значений этого набора.

У меня есть 2D-массив в памяти устройства, точно (3, 1024) 3 строки, 1024 столбца. Каждый столбец представляет отдельный набор, для которого необходимо выполнить функцию интегрирования.

Проблема возникает, когда я должен решить, буду ли я выполнять блок потоков, такой как 1024, имея в виду, что один поток эквивалентен одной функции интеграции. В этом случае функция, которую я написал выше, бесполезна. Поскольку я хочу выполнить параллельную интеграцию для всех наборов значений, мне нужно написать функцию интеграции, которая может выполнять интеграцию последовательно. Например:

__global__ void kernel(float *d_arr, float a, float b, float c, int N) {
    
   int idx = blockIdx.x * blockDim.x + threadIdx.x;
   float sum = 0;
   for (int i = 0; i < N; i++) {
      float x = (float)i / (float) N;
      sum += x * function(x, a, b, c);
   } 
    d_arr[idx] = sum;
}

Итак, вы понимаете мою точку зрения? Вариант А кажется лучше, но я не могу его использовать, потому что я не знаю, как я могу сделать несколько интегралов, а затем распределить каждый интеграл по N потокам.

Как бы вы это сделали? Можете ли вы предложить мне, как я могу добиться как нескольких интегралов, так и того, что каждый интеграл может быть распределен по N потокам? Есть ли лучший способ сделать это.

С нетерпением жду вашего совета.


performance cuda numerical-integration
person fahad    schedule 14.06.2013    source источник
comment
Потоки в одном блоке потока могут обмениваться информацией через общую память. Вы изучали, как это может помочь в вашем случае использования?   -  person njuffa    schedule 14.06.2013
comment
@njuffa Итак, я предполагаю, что вы имеете в виду, что я генерирую столько потоков, сколько подинтервалов я хочу иметь, и все эти блоки потоков имеют один и тот же набор значений, который я могу хранить в общей памяти.   -  person fahad    schedule 15.06.2013

Ответы (1)


arrow_upward
1
arrow_downward

Если я правильно понимаю вашу проблему, вы хотите выполнить численное интегрирование с несколькими (1024) наборами входных данных (a, b, c), и для каждого интеграла вам нужно N подынтервалов. Назовем количество наборов входов M.

Если N достаточно велико (скажем,> 10000), первый вставленный вами образец ядра может быть достаточно хорошим (вызывая его M раз для разных наборов входных данных). Будет ли он использовать всю доступную пропускную способность устройства, зависит от того, насколько сложна ваша функция.

Я не понял, что именно вы делаете с массивом d_arr[]? Обычно для численного интегрирования вы хотели бы суммировать его. Верно? Вы подводите итоги по процессору? Рассмотрите возможность использования atomicAdd (особенно, если вы собираетесь запускать ядро ​​на вычислительной мощности GPU 3.0 и выше) или параллельного сканирования, если считаете, что atomicAdd недостаточно быстр.

Если N мало, лучше запускать N*M потоков в одном ядре.

В вашем случае, когда M = 1024, вы можете обрабатывать каждый блок одним набором входных данных (т. Е. Установить blockSize = 1024) и передавать входные данные (a, b, c) в виде массивов в ядро ​​- что-то вроде этого:

__global__ void kernel(float *d_arr, float *a_array, float *b_array, float *c_array, int totalThreads, int N) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    float x = (float) blockIdx.x / (float) N;
    float a = a_array[threadIdx.x];
    float b = b_array[threadIdx.x];
    float c = c_array[threadIdx.x];

    if (idx < totalThreads)  {
       // what happen to this array?
       d_arr[idx] = x * function(x, a, b, c);
    }
}

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

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

(в этом примере предполагается, что N ‹= 1024, но можно разбить ядро, чтобы воспользоваться преимуществами этого подхода, даже если это не так)

__global__ void kernel(float *d_arr, float *a_array, float *b_array, float *c_array, int totalThreads) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;

    float x = (float)threadIdx.x / (float) blockDim.x;  // N = blockDim.x

    float a = a_array[blockIdx.x];  // every thread in block accesses same memory location
    float b = b_array[blockIdx.x];
    float c = c_array[blockIdx.x];

    // d_arr has 'M' elements containing the integral for each input set.
    if (idx < totalThreads)  
    {
       atomicAdd(&d_arr[blockIdx.x], x * function(x, a, b, c));
    }
}

В приведенном выше ядре есть a_array, b_array и c_array, выделенные в постоянной памяти. это будет быстрее, так как каждый поток в блоке будет обращаться к одному и тому же местоположению. В качестве примера я также заменил ваши записи d_arr на atomicAdd.

person Ankan    schedule 16.06.2013