|
|
@@ -8,11 +8,201 @@ from collections import defaultdict
|
|
|
|
|
|
|
|
|
# Get the OpenCL kernel
|
|
|
-kernel_path = os.path.join(
|
|
|
- os.path.dirname(__file__),
|
|
|
- "kernels/brute-force.cl"
|
|
|
-)
|
|
|
-kernel = open(kernel_path, "r").read()
|
|
|
+kernel = """
|
|
|
+#include <pyopencl-complex.h>
|
|
|
+
|
|
|
+/*
|
|
|
+ * Returns the nth number where a given digit
|
|
|
+ * is cleared in the binary representation of the number
|
|
|
+ */
|
|
|
+static int nth_cleared(int n, int target)
|
|
|
+{
|
|
|
+ int mask = (1 << target) - 1;
|
|
|
+ int not_mask = ~mask;
|
|
|
+
|
|
|
+ return (n & mask) | ((n & not_mask) << 1);
|
|
|
+}
|
|
|
+
|
|
|
+///////////////////////////////////////////////
|
|
|
+// KERNELS
|
|
|
+///////////////////////////////////////////////
|
|
|
+
|
|
|
+/*
|
|
|
+ * Applies a single qubit gate to the register.
|
|
|
+ * The gate matrix must be given in the form:
|
|
|
+ *
|
|
|
+ * A B
|
|
|
+ * C D
|
|
|
+ */
|
|
|
+__kernel void apply_gate(
|
|
|
+ __global cfloat_t *amplitudes,
|
|
|
+ int target,
|
|
|
+ cfloat_t A,
|
|
|
+ cfloat_t B,
|
|
|
+ cfloat_t C,
|
|
|
+ cfloat_t D)
|
|
|
+{
|
|
|
+ int const global_id = get_global_id(0);
|
|
|
+
|
|
|
+ int const zero_state = nth_cleared(global_id, target);
|
|
|
+
|
|
|
+ // int const zero_state = state & (~(1 << target)); // Could just be state
|
|
|
+ int const one_state = zero_state | (1 << target);
|
|
|
+
|
|
|
+ cfloat_t const zero_amp = amplitudes[zero_state];
|
|
|
+ cfloat_t const one_amp = amplitudes[one_state];
|
|
|
+
|
|
|
+ amplitudes[zero_state] = cfloat_add(cfloat_mul(A, zero_amp), cfloat_mul(B, one_amp));
|
|
|
+ amplitudes[one_state] = cfloat_add(cfloat_mul(D, one_amp), cfloat_mul(C, zero_amp));
|
|
|
+}
|
|
|
+
|
|
|
+/*
|
|
|
+ * Applies a controlled single qubit gate to the register.
|
|
|
+ */
|
|
|
+__kernel void apply_controlled_gate(
|
|
|
+ __global cfloat_t *amplitudes,
|
|
|
+ int control,
|
|
|
+ int target,
|
|
|
+ cfloat_t A,
|
|
|
+ cfloat_t B,
|
|
|
+ cfloat_t C,
|
|
|
+ cfloat_t D)
|
|
|
+{
|
|
|
+ int const global_id = get_global_id(0);
|
|
|
+ int const zero_state = nth_cleared(global_id, target);
|
|
|
+ int const one_state = zero_state | (1 << target); // Set the target bit
|
|
|
+
|
|
|
+ int const control_val_zero = (((1 << control) & zero_state) > 0) ? 1 : 0;
|
|
|
+ int const control_val_one = (((1 << control) & one_state) > 0) ? 1 : 0;
|
|
|
+
|
|
|
+ cfloat_t const zero_amp = amplitudes[zero_state];
|
|
|
+ cfloat_t const one_amp = amplitudes[one_state];
|
|
|
+
|
|
|
+ if (control_val_zero == 1)
|
|
|
+ {
|
|
|
+ amplitudes[zero_state] = cfloat_add(cfloat_mul(A, zero_amp), cfloat_mul(B, one_amp));
|
|
|
+ }
|
|
|
+
|
|
|
+ if (control_val_one == 1)
|
|
|
+ {
|
|
|
+ amplitudes[one_state] = cfloat_add(cfloat_mul(D, one_amp), cfloat_mul(C, zero_amp));
|
|
|
+ }
|
|
|
+
|
|
|
+
|
|
|
+}
|
|
|
+
|
|
|
+/*
|
|
|
+ * Applies a controlled-controlled single qubit gate to the register.
|
|
|
+ * NOT MIGRATED
|
|
|
+ */
|
|
|
+__kernel void apply_controlled_controlled_gate(
|
|
|
+ __global cfloat_t *const amplitudes,
|
|
|
+ __global cfloat_t *amps,
|
|
|
+ int control1,
|
|
|
+ int control2,
|
|
|
+ int target,
|
|
|
+ cfloat_t A,
|
|
|
+ cfloat_t B,
|
|
|
+ cfloat_t C,
|
|
|
+ cfloat_t D)
|
|
|
+{
|
|
|
+ int const state = get_global_id(0);
|
|
|
+ cfloat_t const amp = amplitudes[state];
|
|
|
+
|
|
|
+ int const zero_state = state & (~(1 << target));
|
|
|
+ int const one_state = state | (1 << target);
|
|
|
+
|
|
|
+ int const bit_val = (((1 << target) & state) > 0) ? 1 : 0;
|
|
|
+ int const control1_val = (((1 << control1) & state) > 0) ? 1 : 0;
|
|
|
+ int const control2_val = (((1 << control2) & state) > 0) ? 1 : 0;
|
|
|
+
|
|
|
+ if (control1_val == 0 || control2_val == 0)
|
|
|
+ {
|
|
|
+ // Control is 0, don't apply gate
|
|
|
+ amps[state] = amp;
|
|
|
+ }
|
|
|
+ else
|
|
|
+ {
|
|
|
+ // control is 1, apply gate.
|
|
|
+ if (bit_val == 0)
|
|
|
+ {
|
|
|
+ // Bitval = 0
|
|
|
+ amps[state] = cfloat_add(cfloat_mul(A, amp), cfloat_mul(B, amplitudes[one_state]));
|
|
|
+ }
|
|
|
+ else
|
|
|
+ {
|
|
|
+ amps[state] = cfloat_add(cfloat_mul(D, amp), cfloat_mul(C, amplitudes[zero_state]));
|
|
|
+ }
|
|
|
+ }
|
|
|
+}
|
|
|
+
|
|
|
+/*
|
|
|
+ * Swaps the states of two qubits in the register
|
|
|
+ * NOT MIGRATED
|
|
|
+ */
|
|
|
+// __kernel void swap(
|
|
|
+// __global cfloat_t *const amplitudes,
|
|
|
+// __global cfloat_t *amps,
|
|
|
+// int first_qubit,
|
|
|
+// int second_qubit)
|
|
|
+// {
|
|
|
+// int const state = get_global_id(0);
|
|
|
+
|
|
|
+// int const first_bit_mask = 1 << first_qubit;
|
|
|
+// int const second_bit_mask = 1 << second_qubit;
|
|
|
+
|
|
|
+// int const new_second_bit = ((state & first_bit_mask) >> first_qubit) << second_qubit;
|
|
|
+// int const new_first_bit = ((state & second_bit_mask) >> second_qubit) << first_qubit;
|
|
|
+
|
|
|
+// int const new_state = (state & !first_bit_mask & !second_bit_mask) | new_first_bit | new_second_bit;
|
|
|
+
|
|
|
+// amps[new_state] = amplitudes[state];
|
|
|
+// }
|
|
|
+
|
|
|
+
|
|
|
+/**
|
|
|
+ * Get a single amplitude
|
|
|
+ */
|
|
|
+__kernel void get_single_amplitude(
|
|
|
+ __global cfloat_t *const amplitudes,
|
|
|
+ __global cfloat_t *out,
|
|
|
+ int i)
|
|
|
+{
|
|
|
+ out[0] = amplitudes[i];
|
|
|
+}
|
|
|
+
|
|
|
+/**
|
|
|
+ * Calculates The Probabilities Of A State Vector
|
|
|
+ */
|
|
|
+__kernel void calculate_probabilities(
|
|
|
+ __global cfloat_t *const amplitudes,
|
|
|
+ __global float *probabilities)
|
|
|
+{
|
|
|
+ int const state = get_global_id(0);
|
|
|
+ cfloat_t amp = amplitudes[state];
|
|
|
+
|
|
|
+ probabilities[state] = cfloat_abs(cfloat_mul(amp, amp));
|
|
|
+}
|
|
|
+
|
|
|
+/**
|
|
|
+ * Initializes a register to the value 1|0..100...0>
|
|
|
+ * ^ target
|
|
|
+ */
|
|
|
+__kernel void initialize_register(
|
|
|
+ __global cfloat_t *amplitudes,
|
|
|
+ int const target)
|
|
|
+{
|
|
|
+ int const state = get_global_id(0);
|
|
|
+ if (state == target)
|
|
|
+ {
|
|
|
+ amplitudes[state] = cfloat_new(1, 0);
|
|
|
+ }
|
|
|
+ else
|
|
|
+ {
|
|
|
+ amplitudes[state] = cfloat_new(0, 0);
|
|
|
+ }
|
|
|
+}
|
|
|
+"""
|
|
|
|
|
|
|
|
|
class Backend:
|
Adam Kelly