From 321f2e9d1d3942e56da98c87b7df806d31210348 Mon Sep 17 00:00:00 2001 From: Tomas Jusko Date: Fri, 28 Jan 2022 21:47:19 +0100 Subject: Added stacked traces class and GPU combine algorithms --- pyecsca/sca/__init__.py | 1 + pyecsca/sca/stacked_trace/__init__.py | 3 - pyecsca/sca/stacked_trace/stacked_trace.py | 194 ---------------------- pyecsca/sca/stacked_traces/__init__.py | 1 + pyecsca/sca/stacked_traces/stacked_traces.py | 233 +++++++++++++++++++++++++++ 5 files changed, 235 insertions(+), 197 deletions(-) delete mode 100644 pyecsca/sca/stacked_trace/__init__.py delete mode 100644 pyecsca/sca/stacked_trace/stacked_trace.py create mode 100644 pyecsca/sca/stacked_traces/__init__.py create mode 100644 pyecsca/sca/stacked_traces/stacked_traces.py diff --git a/pyecsca/sca/__init__.py b/pyecsca/sca/__init__.py index 1aae9d9..5b359b8 100644 --- a/pyecsca/sca/__init__.py +++ b/pyecsca/sca/__init__.py @@ -5,3 +5,4 @@ from .scope import * from .target import * from .trace import * from .trace_set import * +from .stacked_traces import * diff --git a/pyecsca/sca/stacked_trace/__init__.py b/pyecsca/sca/stacked_trace/__init__.py deleted file mode 100644 index 0e77c18..0000000 --- a/pyecsca/sca/stacked_trace/__init__.py +++ /dev/null @@ -1,3 +0,0 @@ - - -from .stacked_trace import * \ No newline at end of file diff --git a/pyecsca/sca/stacked_trace/stacked_trace.py b/pyecsca/sca/stacked_trace/stacked_trace.py deleted file mode 100644 index 9ab74b1..0000000 --- a/pyecsca/sca/stacked_trace/stacked_trace.py +++ /dev/null @@ -1,194 +0,0 @@ -from audioop import avg -from numba import cuda, float32 -import numpy as np -from public import public -from typing import Any, Iterable, Mapping, MutableSequence, Optional -from math import ceil, sqrt - -TPB = 128 - - -@public -class StackedTraces: - """Samples of multiple traces and metadata""" - - meta: Mapping[str, Any] - traces: np.ndarray - - def __init__( - self, traces: np.ndarray, - meta: Mapping[str, Any] = None) -> None: - if meta is None: - meta = dict() - self.meta = meta - self.traces = traces - - @classmethod - def fromarray(cls, traces: MutableSequence[np.ndarray], - meta: Mapping[str, Any] = None) -> 'StackedTraces': - min_samples = min(map(len, traces)) - for i, t in enumerate(traces): - traces[i] = t[:min_samples] - stacked = np.stack(traces) - return cls(stacked, meta) - - @classmethod - def fromtraceset(cls, traceset) -> 'StackedTraces': - traces = [t.samples for t in traceset] - return cls.fromarray(traces) - - def __len__(self): - return self.traces.shape[0] - - def __getitem__(self, index): - return self.traces - - def __iter__(self): - yield from self.traces - - -class GPUTraceManager: - @staticmethod - def average(traces: StackedTraces) -> np.ndarray: - samples = traces.traces - samples_global = cuda.to_device(samples) - device_result = cuda.device_array(samples.shape[1]) - - tpb = TPB - bpg = (samples.size + (tpb - 1)) // tpb - - gpu_average[bpg, tpb](samples_global, device_result) - res = device_result.copy_to_host() - return res - - def conditional_average(traces: StackedTraces) -> np.ndarray: - raise NotImplementedError - - def standard_deviation(traces: StackedTraces) -> np.ndarray: - samples = traces.traces - samples_global = cuda.to_device(samples) - device_result = cuda.device_array(samples.shape[1]) - - tpb = TPB - bpg = (samples.size + (tpb - 1)) // tpb - - gpu_std_dev[bpg, tpb](samples_global, device_result) - res = device_result.copy_to_host() - return res - - -@cuda.jit -def gpu_average(samples: np.ndarray, result: np.ndarray): - col = cuda.grid(1) - - if col >= samples.shape[1]: - return - - acc = 0. - for row in range(samples.shape[0]): - acc += samples[row, col] - result[col] = acc / samples.shape[0] - - -@cuda.jit() -def gpu_std_dev(samples: np.ndarray, result: np.ndarray): - col = cuda.grid(1) - - if col >= samples.shape[1]: - return - - avg = 0. - for row in range(samples.shape[0]): - avg += samples[row, col] - avg /= samples.shape[0] - - var = 0. - for row in range(samples.shape[0]): - current = samples[row, col] - avg - var += current * current - result[col] = sqrt(var / samples.shape[0]) - - -@cuda.jit() -def gpu_variance(samples: np.ndarray, result: np.ndarray): - col = cuda.grid(1) - - if col >= samples.shape[1]: - return - - avg = 0. - for row in range(samples.shape[0]): - avg += samples[row, col] - avg /= samples.shape[0] - - var = 0. - for row in range(samples.shape[0]): - current = samples[row, col] - avg - var += current * current - result[col] = var / samples.shape[0] - - -@cuda.jit() -def gpu_avg_var(samples: np.ndarray, result_avg: np.ndarray, - result_var: np.ndarray): - col = cuda.grid(1) - - if col >= samples.shape[1]: - return - - avg = 0. - for row in range(samples.shape[0]): - avg += samples[row, col] - avg /= samples.shape[0] - - var = 0. - for row in range(samples.shape[0]): - current = samples[row, col] - avg - var += current * current - result_avg[col] = avg - result_var[col] = var - - -@cuda.jit() -def gpu_add(samples: np.ndarray, result: np.ndarray): - col = cuda.grid(1) - - if col >= samples.shape[1]: - return - - res = 0. - for row in range(samples.shape[0]): - res += samples[row, col] - result[col] = res - - -@cuda.jit() -def gpu_subtract(samples_one: np.ndarray, samples_other: np.ndarray, - result: np.ndarray): - col = cuda.grid(1) - - if col >= samples_one.shape[1]: - return - - result[col] = samples_one[col] - samples_other[col] - - -def test_average(): - samples = np.random.rand(4 * TPB, 8 * TPB) - ts = StackedTraces.fromarray(np.array(samples)) - res = GPUTraceManager.average(ts) - check_res = samples.sum(0) / ts.traces.shape[0] - print(all(check_res == res)) - - -def test_standard_deviation(): - samples: np.ndarray = np.random.rand(4 * TPB, 8 * TPB) - ts = StackedTraces.fromarray(np.array(samples)) - res = GPUTraceManager.standard_deviation(ts) - check_res = samples.std(0, dtype=samples.dtype) - print(all(np.isclose(res, check_res))) - - -if __name__ == '__main__': - test_average() - test_standard_deviation() diff --git a/pyecsca/sca/stacked_traces/__init__.py b/pyecsca/sca/stacked_traces/__init__.py new file mode 100644 index 0000000..090c00c --- /dev/null +++ b/pyecsca/sca/stacked_traces/__init__.py @@ -0,0 +1 @@ +from .stacked_traces import * \ No newline at end of file diff --git a/pyecsca/sca/stacked_traces/stacked_traces.py b/pyecsca/sca/stacked_traces/stacked_traces.py new file mode 100644 index 0000000..7af96e3 --- /dev/null +++ b/pyecsca/sca/stacked_traces/stacked_traces.py @@ -0,0 +1,233 @@ +from numba import cuda, float32 +import numpy as np +from public import public +from typing import Any, Iterable, Mapping, MutableSequence, Optional +from math import ceil, sqrt + +TPB = 128 + + +@public +class StackedTraces: + """Samples of multiple traces and metadata""" + + meta: Mapping[str, Any] + samples: np.ndarray + + def __init__( + self, samples: np.ndarray, + meta: Mapping[str, Any] = None) -> None: + if meta is None: + meta = dict() + self.meta = meta + self.samples = samples + + @classmethod + def fromarray(cls, traces: MutableSequence[np.ndarray], + meta: Mapping[str, Any] = None) -> 'StackedTraces': + min_samples = min(map(len, traces)) + for i, t in enumerate(traces): + traces[i] = t[:min_samples] + stacked = np.stack(traces) + return cls(stacked, meta) + + @classmethod + def fromtraceset(cls, traceset) -> 'StackedTraces': + traces = [t.samples for t in traceset] + return cls.fromarray(traces) + + def __len__(self): + return self.traces.shape[0] + + def __getitem__(self, index): + return self.traces + + def __iter__(self): + yield from self.traces + + +@public +class GPUTraceManager: + @staticmethod + def average(traces: StackedTraces) -> np.ndarray: + samples = traces.samples + samples_global = cuda.to_device(samples) + device_result = cuda.device_array(samples.shape[1]) + + tpb = TPB + bpg = (samples.size + (tpb - 1)) // tpb + + gpu_average[bpg, tpb](samples_global, device_result) + res = device_result.copy_to_host() + return res + + @staticmethod + def conditional_average(traces: StackedTraces) -> np.ndarray: + raise NotImplementedError + + @staticmethod + def standard_deviation(traces: StackedTraces) -> np.ndarray: + samples = traces.samples + samples_global = cuda.to_device(samples) + device_result = cuda.device_array(samples.shape[1]) + + tpb = TPB + bpg = (samples.size + (tpb - 1)) // tpb + + gpu_std_dev[bpg, tpb](samples_global, device_result) + res = device_result.copy_to_host() + return res + + +@cuda.jit(device=True) +def _gpu_average(col: int, samples: np.ndarray, result: np.ndarray): + # col = cuda.grid(1) + + # if col >= samples.shape[1]: + # return + + acc = 0. + for row in range(samples.shape[0]): + acc += samples[row, col] + result[col] = acc / samples.shape[0] + + +@cuda.jit +def gpu_average(samples: np.ndarray, result: np.ndarray): + col = cuda.grid(1) + + if col >= samples.shape[1]: + return + + _gpu_average(col, samples, result) + + +@cuda.jit(device=True) +def _gpu_var_from_avg(col: int, samples: np.ndarray, averages: np.ndarray, result: np.ndarray): + var = 0. + for row in range(samples.shape[0]): + current = samples[row, col] - averages[col] + var += current * current + result[col] = var / samples.shape[0] + + +@cuda.jit(device=True) +def _gpu_variance(col: int, samples: np.ndarray, result: np.ndarray): + # col = cuda.grid(1) + + # if col >= samples.shape[1]: + # return + + # avg = 0. + # for row in range(samples.shape[0]): + # avg += samples[row, col] + # avg /= samples.shape[0] + + _gpu_average(col, samples, result) + _gpu_var_from_avg(col, samples, result, result) + # var = 0. + # for row in range(samples.shape[0]): + # current = samples[row, col] - result[col] + # var += current * current + # result[col] = var / samples.shape[0] + + +@cuda.jit +def gpu_std_dev(samples: np.ndarray, result: np.ndarray): + col = cuda.grid(1) + + if col >= samples.shape[1]: + return + + # avg = 0. + # for row in range(samples.shape[0]): + # avg += samples[row, col] + # avg /= samples.shape[0] + + # var = 0. + # for row in range(samples.shape[0]): + # current = samples[row, col] - result[col] + # var += current * current + # result[col] = sqrt(var / samples.shape[0]) + + _gpu_variance(col, samples, result) + + result[col] = sqrt(result[col]) + + +@cuda.jit +def gpu_variance(samples: np.ndarray, result: np.ndarray): + col = cuda.grid(1) + + if col >= samples.shape[1]: + return + + _gpu_variance(col, samples, result) + + +@cuda.jit +def gpu_avg_var(samples: np.ndarray, result_avg: np.ndarray, + result_var: np.ndarray): + col = cuda.grid(1) + + if col >= samples.shape[1]: + return + + # avg = 0. + # for row in range(samples.shape[0]): + # avg += samples[row, col] + # avg /= samples.shape[0] + + _gpu_average(samples, result_avg) + _gpu_var_from_avg(col, samples, result_avg, result_var) + + # var = 0. + # for row in range(samples.shape[0]): + # current = samples[row, col] - result_avg[col] + # var += current * current + # result_var[col] = var / samples.shape[0] + + +@cuda.jit +def gpu_add(samples: np.ndarray, result: np.ndarray): + col = cuda.grid(1) + + if col >= samples.shape[1]: + return + + res = 0. + for row in range(samples.shape[0]): + res += samples[row, col] + result[col] = res + + +@cuda.jit +def gpu_subtract(samples_one: np.ndarray, samples_other: np.ndarray, + result: np.ndarray): + col = cuda.grid(1) + + if col >= samples_one.shape[1]: + return + + result[col] = samples_one[col] - samples_other[col] + + +def test_average(): + samples = np.random.rand(4 * TPB, 8 * TPB) + ts = StackedTraces.fromarray(np.array(samples)) + res = GPUTraceManager.average(ts) + check_res = samples.sum(0) / ts.traces.shape[0] + print(all(check_res == res)) + + +def test_standard_deviation(): + samples: np.ndarray = np.random.rand(4 * TPB, 8 * TPB) + ts = StackedTraces.fromarray(np.array(samples)) + res = GPUTraceManager.standard_deviation(ts) + check_res = samples.std(0, dtype=samples.dtype) + print(all(np.isclose(res, check_res))) + + +if __name__ == '__main__': + test_average() + test_standard_deviation() -- cgit v1.3.1