cuda_course_task_quantum (Задания)

Лихогруд Николайn.lihogrud@gmail.comЗаданиеОднокубитовая квантовая операция Операция применения квантового вентиля к одномуиз кубитов квантового регистра Квантовый вентиль задается унитарной матрицейразмера 2 × 2 Квантовый регистр из кубитов задается квантовойсуперпозицией его состояний На выходе операции – новая суперпозиция состоянийрегистраУнитарная матрица Квадратная матрица с комплексными элементами, результатумножения которой на эрмитово сопряжённую равен единичнойматрице× 2+ 2= + + 1 022 = 0 1 + Квантовая суперпозиция Обозначим состояние квантового регистра, при котором кубитынаходятся в состояниях | >, = 0.

. − 1, как булевый вектор :−1 2 = , = 0 или 1 = 0 1 … −1 ∶=0 Квантовая суперпозиция состояний квантового регистра из кубитзадается 2 комплексными числами, такими, что2 −1| |2 = 1, − амплитуда состояния =0 | |2 - вероятность при измерении получить состояние Квантовая операция Вход: = { } ∈ 2×2 - квантовый вентиль (унитарная матрица) = {0 , 1 , … , 2 −1} – суперпозиция состояний регистра из N кубитов – индекс кубита, к которому нужно применить вентиль Выход ′ = {′ , ′ , … , ′ 0′ 101 ……−12 −1} – новая суперпозиция состояний:= 0 ∗ 01 …0…−1 + 1 ∗ 01…1…−1Пример Регистр из одного кубита:′ 000 ∗ 0 + 01 ∗ 1==×′10 ∗ 0 + 11 ∗ 11Пример Регистр из двух кубитов, = 1:′ 0000 ∗ 00 + 01 ∗ 01′ 0110 ∗ 00 + 11 ∗ 01= ∗ + ∗′100010011110 ∗ 10 + 11 ∗ 11′11Пример Регистр из двух кубитов, = 0:′ 0000 ∗ 00 + 01 ∗ 10′ 0100 ∗ 01 + 01 ∗ 11= ∗ + ∗′101000111010 ∗ 01 + 11 ∗ 11′11Пример Регистр из трех кубитов, = 1:000 00 ∗ 000 + 01 ∗ 010001 00 ∗ 001 + 01 ∗ 01101010 ∗ 000 + 11 ∗ 01001110 ∗ 001 + 11 ∗ 011100 = 00 ∗ 100 + 01 ∗ 11010100 ∗ 101 + 01 ∗ 11111010 ∗ 100 + 11 ∗ 11011110 ∗ 101 + 11 ∗ 111Суть задания Предоставляется C++ реализация: Парсинга параметров Генерации, проверки корректности, вывода суперпозициисостояний квантового регистра и квантового вентиля(унитарной матрицы) Вспомогательных типов для упрощения CUDA-ядер Требуется реализовать однокубитовую квантовую операциюопераций на CPU и GPUОписание программы• Параметры программы:-n <number of qubits> - число кубитов в регистре-i <qubit index>- номер кубита, над которым требуетсяпровести операцию-check- флаг проверки результата-v- флаг вывода в консоль регистров иквантового вентиля• Вывод программы:• Время работы на GPU• Время работы на CPU, если включен флаг проверки результатаИспользуемые типы данныхв definitions.h:typedef double ComplexAmplitudePart; Тип мнимой/вещественной частей комплексной амплитуды.Вынесена вотдельный typedef для быстрого переключения между float/doubletypedef std::complex<ComplexAmplitudePart> ComplexAmplitude; Комплексная амплитудаtypedef std::vector<ComplexAmplitude> QuantumSuperposition; Суперпозиция состояний квантового регистраtypedef ComplexAmplitude QuantumGate[2][2]; Квантовый вентильВспомогательные функции В definitions.h и helpers.cpp:void intializeRandomQuantumRegisterStatesSuperposition(QuantumSuperposition &quantumRegisterStatesSuperposition); Инициализация квантового регистра.

Сумма квадратов амплитуд равнаединице.void intializeRandomQuantumGate(QuantumGate &quantumGate); Инициализация квантового вентиля - унитарной комплексной матрицыvoid printQuantumRegisterStatesSuperposition(constQuantumSuperposition &quantumRegisterStatesSuperposition); Вывести амплитуды состояний квантовой суперпозицииВспомогательные типы для GPU В kernel.cu:Комплексное число с операторами, работающими на GPUstruct ComplexAmplitudeGPU {ComplexAmplitudePart real;ComplexAmplitudePart imag;__device__ ComplexAmplitudeGPU() {}__device__ ComplexAmplitudeGPU(ComplexAmplitudePart real,ComplexAmplitudePart imag): real(real), imag(imag) {}__device__ ComplexAmplitudeGPU operator*(ComplexAmplitudeGPU other);__device__ ComplexAmplitudeGPU operator+(ComplexAmplitudeGPU other);};Вспомогательные типы для GPU В kernel.cu:Класс для раздельного хранения вещественных и мнимых частей, см.лекцию по глобальной памятиstruct QuantumSuperpositionGPU {int numberOfStates;ComplexAmplitudePart *real;ComplexAmplitudePart *imag;__device__ ComplexAmplitudeGPU operator[] (uint64_t index);__device__ void setAmplitudeAtIndex(uint64_t index,ComplexAmplitudeGPU amplitude);};Требуется реализоватьvoid applyQuantumGateToQuantumSuperpositionOnCPU(QuantumSuperposition &quantumSuperposition,QuantumGate quantumGate, int qubitIndex); Применить на CPU вентиль quantumGate к кубиту с индексомqubitIndex в квантовом регистре quantumSuperposition Использовать OpenMP, т.е.

накинуть #pragma omp parallel forна внешний цикл*Результат применения операции записывается в тот же векторТребуется реализоватьvoid applyQuantumGateToQuantumSuperpositionOnGPU(QuantumSuperposition &quantumSuperposition,QuantumGate quantumGate, int qubitIndex); Применить на GPU вентиль quantumGate к кубиту с индексомqubitIndex в квантовом регистре quantumSuperposition Функция выделяет память на GPU, копирует на GPU квантовуюсуперпозицию и вентиль, запускает ядро, копирует результат с GPU,замеряет время выполнения копирований и счета ядра*Результат применения операции записывается в тот же векторТребования к CUDA-части Квантовый вентиль располагаем в константной памяти__constant__ ComplexAmplitudeGPU quantumGateOnGPU[2][2]; Каждая нить рассчитывает амплитуды двух связанныхсостояний.

Общее число нитей 2−1 Блоки грида линейные, по 512 нитей - dim3(512) Размеры грида рассчитывать в зависимости от параметраустройства int cudaDeviceProp::maxGridSize[3] Если число блоков превышает максимальную ширину грида,то высота грида будет больше одного блокаПереход от CPU к GPU На CPU – цикл из 2−1 итераций, на каждой итерациирассчитываются амплитуды двух связанных состояний регистра Для перехода к GPU-реализации нужно распределить итерациипо нитям Ядро будет состоять из вычисления индекса итерации,которую должна выполнить нить, и тела цикла из CPUреализацииТесты• Продемонстрировать результатытестирования в виде графика (для double):Тесты• Продемонстрировать результатытестирования в виде графика (для float):000000000001000010000011000100000101000110000111001000001001001010001011001100001101001110001111010000010001010010010011010100010101010110010111011000011001011010011011011100011101011110011111Независимость блоков• Пусть блок состоит из 2 нитей• Например, = 3, в блоке 8 нитей• Расстояние между отрезками суперпозиции,рассчитываемыми блоком, уменьшается с увеличениеминдекса кубита• При ≥ − ( + 1) блок считает непрерывный отрезоксуперпозиции• - число кубитов, – индекс кубита, каждая нитьсчитает две амплитуды,• Отрезки разных блоков не пересекаются000000000000000000000000000001000001000001000001000010000010000010000010000011000011000011000011000100000100000100000100000101000101000101000101000110000110000110000110000111000111000111000111001000001000001000001000001001001001001001001001001010001010001010001010001011001011001011001011001100001100001100001100001101001101001101001101001110001110001110001110001111001111001111001111010000010000010000010000010001010001010001010001010010010010010010010010010011010011010011010011010100010100010100010100010101010101010101010101010110010110010110010110010111010111010111010111011000011000011000011000011001011001011001011001011010011010011010011010011011011011011011011011011100011100011100011100011101011101011101011101011110011110011110011110011111011111011111011111Суть задания Последовательно применить квантовый вентиль к кубитам синдексами : − log 2 _ + 1≤ ≤ −1 _ – число нитей в блоке, _ = 2 − +1≤ ≤ −1 Последовательно применить квантовый вентиль кпоследним + 1 кубитамСуть задания _ – число нитей в блоке,_ = 2 , − +1≤ ≤ −1 На каждой итерации блоки работают с непересекающимися отрезкамиквантовой суперпозиции Нет необходимости в синхронизации между блоками в концеитерации, каждый блок работает отдельно от других Можем хранить промежуточные результаты в общей памяти Нужна синхронизация внутри блокаОписание программы• Параметры программы:-n <number of qubits> - число кубитов в регистре-block <block size>- размер блока нитей, степень двойки-check- флаг проверки результата-v- флаг вывода в консоль регистров иквантового вентиля• Вывод программы:• Время работы на GPU с общей памятью и без• Время работы на CPU, если включен флаг проверки результатаТребуется реализовать Последовательное применение вентиля к последним + 1кубитам на CPU CUDA-ядро, последовательно применяющее вентиль к последним + 1 кубитам CUDA-ядро, последовательно применяющее вентиль к последним + 1 кубитам.

Промежуточные результаты сохранять в общейпамяти CPU-функцию для управления памятью устройства, запуска двухядер и замера времени их выполненияТестыKepler, 25 кубитовВремя работы, мс62,5Без общей памятиС общей памятью5048,537,536,72527,920,112,522,017,218,011,30блок 1024,doubleблок 512, double блок 1024, float блок 512, floatТесты На Fermi разница в производительности с общей памятью и безнеё меньше, чем на Kepler, т.к. промежуточные результатыполностью влезают в кеш Лично у меня на Fermi нет ускорения на float и есть небольшое наdouble На Kepler ускорение больше, т.к. на нем запросы в глобальнуюпамять не кешируются (см.

в cuda-programming-guide) На Fermi можно замедлить вариант без общей памяти, указав -Xptxas -dlcm=cg при компиляции (выключает кеширование)The end.