🥇Empfang aller Bluetooth-Kanäle gleichzeitig auf SDR mit CUDA? Ganz einfach

In letzter Zeit haben mich meine Kollegen aus der "Branche" unabhängig voneinander gefragt: Wie kann ich mit einem SDR-Empfänger gleichzeitig alle Bluetooth-Kanäle empfangen? Die Bandbreite ist schließlich ausreichend, es gibt SDRs mit einer Ausgangsbande von 80 MHz und mehr. Natürlich lässt sich das auch mit einer FPGA-Lösung umsetzen, aber die Entwicklungszeit wäre ziemlich hoch. Ich wusste schon länger, dass es relativ einfach ist, so etwas auf einer GPU zu realisieren, aber so!

Der Bluetooth-Standard definiert die physikalische Ebene in zwei Versionen: Classic und Low Energy. Es gibt eine Spezifikation, hier. Das Dokument ist schrecklich umfangreich, es ist gefährlich, es vollständig zu lesen. Glücklicherweise verfügen große Unternehmen, die Messgeräte herstellen, über die Mittel, anschauliche Dokumente zu diesem Thema zu erstellen. Tektronix und National Instruments, zum Beispiel. Ich habe keine Chance, qualitativ mit ihnen zu konkurrieren, was die Darstellung des Materials angeht. Ich bitte Interessierte, sich über die Links zu informieren.

Alles, was ich über die physikalische Ebene wissen muss, um einen Mehrkanalfilter zu erstellen, sind der Frequenzrasterabstand und die Modulationsgeschwindigkeit. Diese sind in einer Tabelle in einem der genannten Dokumente zusammengefasst:

Daher müssen wir das 80-MHz-Spektrum in 79 Filter mit einem Abstimmabstand von 1 MHz aufteilen und gleichzeitig in 40 Filter mit einem Abstimmabstand von 2 MHz. Die Abtastraten von den Filterausgängen sollten 1 MHz und 2 MHz betragen, entsprechend.

Wir benötigen also zwei Filtergitter.

Zunächst wählen wir die Parameter dieser Filter basierend auf den Signalen von Bluetooth Classic und Bluetooth Low Energy. Wir benötigen ihre Impulsantworten, um die Last auf das Rechenwerk des Filters zu berechnen. Es sei gleich gesagt, dass wir die Längen der Impulsantworten basierend auf den Anforderungen eines "schnellen" Filteralgorithmus gewählt haben. Das ändert nichts an der Sache. Auch die Anzahl der Koeffizienten der Impulsantwort sollte nicht zu hoch sein, damit der Filter auf einer angemessenen Rechentechnik implementiert werden kann.

Für die Filter mit einem Abstimmabstand von 1 MHz wählen wir eine Bandbreite des Tiefpassfilters (die Hälfte der Bandbreite des Bandfilters) von 500 kHz und passen die Länge der Impulsantwort an 480 Koeffizienten an. Für die Filter mit einem Abstimmabstand von 2 MHz wählen wir diese Parameter auf 1 MHz und 240 Koeffizienten, entsprechend. Der Fenstertyp wird als Kaiser gewählt. Wir berechnen die Impulsantworten in filterDesigner und exportieren sie im C-Header-Format:

Screenshots aus filterDesigner

Man kann die Aufgabe auf direkte Weise lösen: ein entsprechendes Array von DDCs (Digital Down Converter). Dieser Ansatz ist gut für FPGAs, wo Einsparungen durch die Verringerung der Bitbreite der Berechnungen in den ersten Stufen möglich sind. Zudem ist FPGA die energieeffizienteste Implementierung. Aber der Arbeitsaufwand bei dieser Methode ist am höchsten.

Bei der Durchführung von Filterkämmen auf heutzutage beliebten GPUs ergibt sich die Möglichkeit, einen raffinierten Algorithmus zu implementieren: einen polyphasischen Filterkamm basierend auf FFT, der in der CUDA-Bibliothek verfügbar ist. In der ausländischen Literatur wird der Algorithmus als Polyphase oder WOLA (Weight, Overlap and Add) FFT Filterbank bezeichnet. Mein Mangel an künstlerischem Talent hindert mich daran, eine anschauliche Erklärung selbst zu erstellen. Im Internet gibt es viele Materialien zu diesem Thema, besonders anschauliche Diagramme finden sich hier auf Seite 11 (ein enormes Dankeschön an die geschätzten Autoren), hier ist es:

Ich werde versuchen, das Verarbeitungsschema mit meinen eigenen Worten zu erklären. Zartbesaitete Personen bitte nicht weiterlesen.

Ich werde versuchen, das Verarbeitungsschema mit meinen eigenen Worten im Rahmen meiner methodischen Möglichkeiten zu erklären. FFT ist die Faltung des Eingangssignals mit dem gesamten Spektrum komplexer orthogonaler Harmonien, die sich im Intervall der Impulsantwort anordnen. Die Impulsantwort des Filters, mit der das Signal vor dem FFT multipliziert wird, wird durch dieses Spektrum moduliert. Mit anderen Worten, die Hülle der Impulsantworten der resultierenden Filter wird herausgeholt. Danach wird das Spektrum der Harmonien um einen bestimmten Faktor reduziert, aufgrund der Erweiterung der Durchlassbandbreite des Filters im Vergleich zu dem Filter im rechteckigen Fenster. Auf der Abbildung sehen wir eine Reduzierung um den Faktor vier. Mit anderen Worten, nach der Erweiterung des Bandes mit dem Kaiser-Fenster (bei gleichzeitiger Erhöhung der Dämpfung im Verzögerungsband) benötigen wir nicht alle Filter, sondern nur ein Viertel davon. Die übrigen sind überflüssig, ihre Frequenzeigenschaften überschneiden sich. Von vier aufeinander folgenden FFT-Punkten wählen wir nur den nullten Punkt, dessen Berechnung die Summe von vier
Eingangspunkten ist, die in einem Zeitraum genommen werden, der gleich einem Viertel der Dauer der ursprünglichen FFT ist.

Die Hardware wählen wir aus, was zur Hand ist. Es handelt sich um die Eingangsplatine der Firma "Instrumentensysteme" FMC126P. Darüber habe ich bereits in einem vorherigen Beitrag geschrieben. Artikel. In den FMC-Anschluss der Platine wurde ein Submodul desselben Unternehmens mit dem Transceiver AD9371 und einer Bandbreite von 100 MHz eingesetzt. Der gesamte Datenstrom des Transceivers kann kontinuierlich an den Computer zur Verarbeitung übertragen werden.

Wir wählen eine Grafikkarte mit einem GTX 1050 GPU. (Ich habe gelogen, sie hat uns ausgewählt: das war alles, was zur Hand war, sie wurde aus dem Rechner für Antennenkalkulationen herausgerissen, aber desto erstaunlicher war es, eine funktionierende Kombination zu sehen). Lassen Sie uns zum Softwareteil ü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 auch nicht sonderlich interessant.

Hier ist der Kern, der das Signal mit dem Fenster multipliziert und addiert, sowie die Umhüllung 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 Signalverarbeitungs-Code, der diesen Kern aufruft, folgt genau dem Algorithmusdiagramm, das oben abgebildet ist, daher sehe ich keinen Sinn darin, ihn hier anzuführen.

Die Spektren der Kanäle wurden in Echtzeit überprüft. Am Eingang des AD9371 wurde ein Signal von einem Signalgenerator mit 2450 MHz angelegt, die Selektivität der Filter entsprach den Berechnungen.