🥇 Gesamter Bluetooth-Empfang auf SDR mit CUDA? Kein Problem

In letzter Zeit haben mich mehrere Kollegen unabhängig voneinander gefragt, wie man alle Bluetooth-Kanäle gleichzeitig von einem SDR-Empfänger erhalten kann. Der Frequenzbereich erlaubt dies, da es SDRs mit einer Bandbreite von 80 MHz und mehr gibt. Man könnte das natürlich auch auf FPGAs umsetzen, aber die Entwicklungszeit wäre erheblich. Mir ist schon lange bekannt, dass die Umsetzung auf GPUs recht einfach ist, aber so einfach!

Der Bluetooth-Standard definiert die physikalische Schicht in zwei Versionen: Classic und Low Energy. Die Spezifikation ist vorhanden. hierDas Dokument ist furchtbar umfangreich, es ist gefährlich für das Gehirn, es vollständig zu lesen. Glücklicherweise verfügen große Unternehmen, die Messgeräte herstellen, über die Mittel, um anschauliche Dokumente zu diesem Thema zu erstellen. Tektronix und National Instruments, zum Beispiel. Ich habe absolut keine Chance, mit ihnen in Bezug auf die Qualität der Materialpräsentation zu konkurrieren. Interessierte bitte ich, die Links zu studieren.

Alles, was ich über die physikalische Schicht wissen muss, um einen Multikanalfilter zu erstellen, sind die Frequenzrasterschritte und die Modulationsgeschwindigkeit. Diese sind in einer Tabelle in einem der genannten Dokumente zusammengestellt:

Daher müssen wir ein 80-MHz-Band in 79 Filter mit einem Abstimmungsabstand von 1 MHz und gleichzeitig in 40 Filter mit einem Abstimmungsabstand von 2 MHz aufteilen. Die Abtastfrequenzen an den Ausgängen der Filter sollten 1 MHz und 2 MHz betragen.

Somit benötigen wir zwei Filtergitter.

Zunächst wählen wir die Parameter dieser Filter basierend auf den Frequenzbändern von Bluetooth Classic und Bluetooth Low Energy. Wir benötigen deren Impulscharakteristika, um die Belastung auf das Rechenwerk des Filters zu berechnen. Hier sollte sofort erwähnt werden, dass wir die Längen der Impulscharakteristika basierend auf den Anforderungen eines "schnellen" Filteralgorithms gewählt haben. Die grundlegende Idee bleibt dabei unverändert. Auch die Anzahl der Impuls-Filterkoeffizienten sollte nicht zu groß sein, damit der Filter auf vernünftiger Rechenhardware realisierbar ist.

Für Filter mit einem Schritt von 1 MHz wählen wir eine Bandbreite des Tiefpassfilters (die Hälfte der Bandbreite des Bandfilters) von 500 kHz und passen die Impulsantwort auf 480 Abtastwerte an. Für Filter mit einem Schritt von 2 MHz setzen wir diese Parameter auf 1 MHz und 240 Abtastwerte entsprechend. Fenster Typ wählen wir Kaiser. Wir berechnen die Impulsantworten in filterDesigner und exportieren sie im Format C-Header:

Screenshots aus filterDesigner

Die Aufgabe kann direkt gelöst werden: ein entsprechendes Array DDC (Digital Down Converter). Dieser Ansatz ist gut für FPGAs, wo durch Verringerung der Bitanzahl der ersten Berechnungsschritte Einsparungen möglich sind. Auch FPGAs sind die energieeffizienteste Implementierungsart. Jedoch sind die Arbeitsaufwände bei dieser Methode am höchsten.

Bei der Verwendung von Filterbank-Crossovers auf heutzutage populären GPUs eröffnet sich die Möglichkeit, einen ausgeklügelteren Algorithmus zu implementieren: die polyphase Filterbank auf Basis der FFT, die über die CUDA-Bibliothek verfügbar ist. In der internationalen Literatur wird dieser Algorithmus als Polyphase oder WOLA (Weight, Overlap and Add) FFT Filterbank bezeichnet. Meine Unlust am Zeichnen hindert mich daran, eine anschauliche Erklärung selbst zu erstellen. Im Internet gibt es viele Materialien zu diesem Thema, insbesondere eine anschauliche Grafik, die auf Seite 11 zu finden ist (vielen Dank an die geschätzten Autoren). hier Hier ist sie:

Ich werde versuchen, das Verarbeitungsdiagramm in meinen eigenen Worten zu erläutern. Zartbesaitete Gemüter bitte ich, dies nicht zu lesen.

Ich werde versuchen, das Verarbeitungsmuster in meinen eigenen Worten innerhalb meiner methodischen Möglichkeiten zu erklären. Die schnelle Fourier-Transformation (FFT) ist eine Faltung des Eingangssignals mit dem gesamten Spektrum komplexer orthogonaler Harmoniken, die im Intervall der Impulsantwort angeordnet sind. Die Impulsantwort des Filters, mit der das Signal vor dem Eingang der FFT multipliziert wird, wird durch dieses Spektrum der Harmoniken moduliert. Mit anderen Worten, die Hüllkurve der Impulsantworten der resultierenden Comb-Filter wird ausgeklammert. Dann wird das Spektrum der Harmoniken in ein gewisses Maß reduziert, da die Bandbreite des Filters im Vergleich zu einem Filter mit rechteckigem Fenster erweitert wird. In der Abbildung sehen wir eine Reduzierung um das Vierfache. Mit anderen Worten, nach der Erweiterung der Bandbreite mit einem Kaiser-Fenster (bei gleichzeitiger Erhöhung der Dämpfung im Stoppband) benötigen wir nicht mehr alle Filter, sondern nur ein Viertel davon. Die anderen sind überflüssig, ihre Frequenzeigenschaften überlappen sich. Von vier aufeinanderfolgenden FFT-Punkten wählen wir nur den Nullpunkt, wobei die Berechnung die Summe der vier Eingangspunkte ist.
die in einem Zeitintervall, das einem Viertel der Dauer der ursprünglichen FFT entspricht, erfasst wurden.

Wir wählen die Hardware aus, die uns zur Verfügung steht. Es handelt sich um die Eingangsplatine der Firma "Instrumentensysteme" FMC126P. Darüber habe ich bereits in einem früheren Beitrag geschrieben. ArtikelIn den FMC-Slot der Platine wurde ein Submodul derselben Firma mit dem Transceiver AD9371 mit einer Bandbreite von 100 MHz eingesteckt. Der gesamte Datenstrom vom Transceiver kann kontinuierlich an den Computer zur Verarbeitung übertragen werden.

Wir wählen eine Grafikkarte mit dem GPU GTX 1050. (Ich habe gelogen, sie hat uns gewählt: Das war alles, was zur Verfügung stand; sie wurde aus dem Rechner für die Berechnung der Antennen herausgerissen, aber umso erstaunlicher war es, eine funktionierende Kombination zu sehen). Lassen Sie uns zum Software-Teil übergehen.

Leider können wir aufgrund von Lizenzen den vollständigen Code nicht veröffentlichen. Wir können nur die GPU-Kerne zeigen. Der restliche Code ist jedoch nicht besonders interessant.

Hier ist der Kern, der das Signal mit einem Fenster multipliziert und addiert, sowie die Hülle für seinen Aufruf:

__global__ void cuComplexMultiplyWindowKernel(const cuComplex *data, const float *window, size_t windowSize, cuComplex *result) {
    __shared__ cuComplex multiplicationResult[480];
    multiplicationResult[threadIdx.x] = cuComplexMultiplyFloat(data[threadIdx.x + windowSize / 4 * blockIdx.x], window[threadIdx.x]);
    __syncthreads();
    cuComplex sum;
    sum.x = sum.y = 0;
    if (threadIdx.x < windowSize / 4) {
        for(int i = 0; i < 4; i++) {
            sum = cuComplexAdd(sum, multiplicationResult[threadIdx.x + i * windowSize / 4]);
        }
        result[threadIdx.x + windowSize / 4 * blockIdx.x] = sum;
    }
}

cudaError_t cuComplexMultiplyWindow(const cuComplex *data, const float *window, size_t windowSize, cuComplex *result, size_t dataSize, cudaStream_t stream) {
    size_t windowStep = windowSize / 4;
    cuComplexMultiplyWindowKernel<<>>(data, window, windowSize, result);
    return cudaGetLastError();
}

Der Code zur Signalverarbeitung, der diesen Kernel aufruft, spiegelt exakt das Schema des oben abgebildeten Algorithmus wider, daher sehe ich keinen Sinn darin, ihn hier wiederzugeben.

Die Frequenzgänge wurden im Ausgangsspektrum der Kanäle in Echtzeit überprüft. Ein 2450 MHz Signal vom Signal-Generator wurde dem AD9371 zugeführt und die Filterselektion entsprach den Berechnungen.