6 апреля
This presentation is the property of its rightful owner.
Sponsored Links
1 / 45

Использование CUDA в расчете динамики пучка PowerPoint PPT Presentation


  • 122 Views
  • Uploaded on
  • Presentation posted in: General

6 апреля 2 010. Использование CUDA в расчете динамики пучка. С.Б. Ворожцов , В.Л. Смирнов , Е.Е. Перепелкин Дубна, ОИЯИ. http://parallel-compute.ru. Циклотрон. Постановка задачи Численные методы Программная реализация на CUDA Результаты. http://cbda.jinr.ru

Download Presentation

Использование CUDA в расчете динамики пучка

An Image/Link below is provided (as is) to download presentation

Download Policy: Content on the Website is provided to you AS IS for your information and personal use and may not be sold / licensed / shared on other websites without getting consent from its author.While downloading, if for some reason you are not able to download a presentation, the publisher may have deleted the file from their server.


- - - - - - - - - - - - - - - - - - - - - - - - - - E N D - - - - - - - - - - - - - - - - - - - - - - - - - -

Presentation Transcript


Cuda

6 апреля 2010

Использование CUDA в расчете динамики пучка

С.Б. Ворожцов, В.Л. Смирнов, Е.Е. Перепелкин

Дубна, ОИЯИ

http://parallel-compute.ru


Cuda

Циклотрон

  • Постановка задачи

  • Численные методы

  • Программная реализация на CUDA

  • Результаты

http://cbda.jinr.ru

CBDA: Cyclotron Beam Dynamic Analysis code


Cuda

Постановка задачи


Cuda

Линия инжекции

Инфлектор

Дуант

ЭСД

Магнитный сектор

Компьютерная модель циклотрона


Cuda

Инфлектор

Электрическое поле

Линза

Магнитное поле

Аксиальный канал

Магнитное поле

Области задания карт полей


Cuda

Ресурсоемкое моделирование

  • Необходимость рассмотреть не менее 5 различных конфигураций центральной зоны;

  • Необходимость ускорять различные ионы;

  • Сложная геометрическая структура;

  • Учет пространственного заряда;

Одна итерация требует ~

несколько дней расчетов


Cuda

Уравнения движения


Cuda

Пространственный заряд

PIC метод

PP метод


Cuda

Численные методы


Cuda

Уравнение движения

Уравнение движения из постановки задачи

можно представить в упрощенном виде, дополнив его вторым уравнением для определения координат частиц


Cuda

Пример решения ОДУ

Рассмотрим решение обыкновенного дифференциального уравнения (ОДУ) методом Рунге -Кутта

Задача Коши


Cuda

Метод Рунге - Кутта


Cuda

Метод Рунге - Кутта


Cuda

Метод Рунге - Кутта


Cuda

Решение краевой задачи

При поиске коэффициентов Фурье используется алгоритм БПФ

(Быстрого Преобразования Фурье)


Cuda

Lz

Lx

Lz

Lx

Ly

Lx

Lz

Ly

Ly

Задание области для краевой задачи


Cuda

Ячейка 7

Ячейка 8

Узел

Ячейка 6

Ячейка 5

Ячейка 3

Ячейка 2

Ячейка 1

Раздача плотности заряда


Cuda

B

D

tn+1

A

tn

C

Потери частиц

Если точка D принадлежит треугольнику ABC, тогда

Условие пересечения

где εΔ – допустимое отклонение от поверхности


Cuda

Программная реализация на CUDA


Cuda

Функции ядра

  • Track ( карты полей, координаты и скорости частиц )

    • метод Рунге-Кутта

  • Losses ( геометрия установки, координаты частиц )

    • проверка пересечений с геометрией

  • Rho ( координаты частиц )

    • раздача заряда в узлы сетки

  • FFT ( функцияплотности заряда или потенциал)

    • БПФ по базисным функциям sin(πn/N)

  • PoissonSolver ( Фурье коэффициенты )

    • решение краевой задачи

  • E_SC ( потенциал электрического поля )

    • поиск электрического поля


Global void track

__global__ void Track ( )

  • Много входных параметров. Использование типа переменной __constant__ для неизменных параметров:

    • __device__ __constant__ float d_float[200];

    • __device__ __constant__ int d_int[80];

  • Каждой частице соответствует нить:

    • int n = threadIdx.x+blockIdx.x*blockDim.x;

  • Количество “if, goto, for” необходимо максимально сократить


If goto for

Инфлектор

Электрическое поле

Линза

Магнитное поле

Аксиальный канал

Магнитное поле

Проблема количества “if, goto, for”


Global void losses

__global__ void Losses ( )

  • Нити одного блока копируют вершины треугольников из globalв sharedпамять.

  • Синхронизация нитей после копирования треугольников __syncthreads()

  • Каждой частице соответствует номер нити:

    • int n = threadIdx.x+blockIdx.x*blockDim.x;

  • Проверка условия пересечения частицей c номеромn,загруженныхвsharedпамять,треугольников

  • Для каждого блока геометрии есть своя функция Losses


Global void rho

__global__ void Rho

  • Каждая частица с номером n = threadIdx.x + blockIdx.x*blockDim.x дает свой вклад, в окружающие ее узлы. Для этого по координатам частицы определяется какой ячейки она принадлежит

  • Одна частица может дать вклад в 8 ближайших узлов. Таким образом, каждая нить заполняет свои 16 ячеек в общем массиве вклада: 8 – номеров узлов и 8 – значений вклада.

  • Далее производится сложение этих вкладов для каждого узла.


Global fft

__global__ FFT ( )

  • Действительное БПФ по базисным функциям sin(πn/N);

  • 3D преобразование состоит из трех последовательных 1D БПФпо осям: X, Y, Z соответственно

  • int n = threadIdx.x+blockIdx.x*blockDim.x;

    k=(int)(n/(NY+1));

    j=n-k*(NY+1);

    m=j*(NX+1)+k*(NX+1)*(NY+1);

    FFT_X[i+1]=Rho[i+m];

n = j + k*(NY+1)

Массив данных для функции Rho трех переменных

NY

NZ


Global poissonsolver

__global__ PoissonSolver ( )

  • Номер нити

    int n = threadIdx.x+blockIdx.x*blockDim.x;

  • Каждая нить находит значение коэффициентов Фурье PhiF потенциала Phi

    PhiFind(i,j,k) = -RhoFind(i,j,k) / ( kxi2 + kyj2 + kzk2 )

    В узле с номером:

    ind(i,j,k)=i+j*(NX+1)+k*(NX+1)*(NY+1),

    где

    k=(int)(n/(NX+1)*(NY+1));

    j=(int)(n-k*(NX+1)*(NY+1))/(NX+1);

    i=n-j*(NX+1)-k*(NX+1)*(NY+1);

  • RhoF – коэффициенты Фурье для функции плотности заряда Rho.


Global e sc

φn + ( NX + 1 )( NY + 1 )

φn + ( NX + 1 )

φn

φn - 1

φn - ( NX + 1 )

φn + 1

φn - ( NX + 1 )( NY + 1 )

__global__ E_SC ( )

  • Вычисление электрического поля в узле с номером

    int n = threadIdx.x+blockIdx.x*blockDim.x+st_ind


Cuda

Результаты


Cuda

Аксиальная инжекция пучка


Cuda

Процесс банчировки пучка


Cuda

Ускорение в циклотроне


Cuda

Анимация


Cuda

Анимация


Cuda

Анимация


Cuda

Потери частиц


Cuda

Ускорение банчей


Cuda

«Земля»

φRF = 15°

φRF = 13°

Без постов

Дуант

φRF = 10°

φRF = 28°

С постами

Оптимизация центральной области

F = ZURF - WGAP


Cuda

S0

S1

S2

S3

S4

Выбор оптимальной конфигурации


Cuda

Распределение ускоряющего поля


8800gtx

Производительность на 8800GTX

*Размер сетки: 25 x 25 x 25. Число частиц: 100,000 треугольников: 2054

**CPU с частотой 2.4 ГГц


Cpu geforce 8800gtx

Сравнение CPU и GeForce 8800GTX

*CPU с частотой 2.4 ГГц


Cpu tesla c1060

Сравнение CPU с Tesla C1060

БЕЗ пространственного заряда


Cpu tesla c10601

Сравнение CPU с Tesla C1060

С пространственным зарядом


Cuda

Эффект пространственного заряда

I ~ 0

Потери 24%

I = 4 мА

Потери 94%


Cuda

Заключение

  • Очень дешевая технология в сравнении с CPU;

  • Увеличение производительности на 1.5 – 2 порядкадает шанс проведения моделирования ресурсоемких физических моделей;

  • Требует аккуратного программирования.


  • Login