Negli ultimi tempi, i colleghi del "settore" hanno iniziato a chiedermi in modo indipendente: come ottenere da un solo ricevitore SDR tutti i canali Bluetooth contemporaneamente? La banda lo consente, abbiamo SDR con larghezza di banda di 80 MHz o più. Certo, si può fare anche su FPGA, ma il tempo per una tale sviluppazione sarebbe piuttosto lungo. Sapevo da tempo che farlo su GPU è abbastanza semplice, ma così!
Lo standard Bluetooth definisce il livello fisico in due versioni: Classic e Low Energy. La specifica è . Il documento è enormemente grande, leggerlo per intero è pericoloso per il cervello. Fortunatamente, grandi aziende che producono strumenti di misura hanno le risorse per creare documenti visivi sull'argomento. e , per esempio. Non ho assolutamente possibilità di competere con loro per la qualità della presentazione del materiale. Invito gli interessati a studiare tramite i link.
Tutto ciò che devo sapere sul livello fisico per creare un filtro multicanale è la griglia delle frequenze e la velocità di modulazione. Sono riassunti in una tabella in uno dei documenti citati:
Pertanto, dobbiamo suddividere una larghezza di banda di 80 MHz in 79 filtri con passo di regolazione di 1 MHz e, contemporaneamente, in 40 filtri con passo di regolazione di 2 MHz. Le frequenze di campionamento dagli output dei filtri devono essere 1 MHz e 2 MHz, rispettivamente.
Di conseguenza, abbiamo bisogno di due griglie di filtri.
Per iniziare, scegliamo i parametri di questi filtri in base alle bande dei segnali Bluetooth Classic e Bluetooth Low Energy. Abbiamo bisogno delle loro caratteristiche impulsive per calcolare il carico sul dispositivo di calcolo del filtro. È opportuno precisare subito che le lunghezze delle caratteristiche impulsive sono state scelte in base ai requisiti dell'algoritmo di filtraggio "veloce". Il concetto rimane invariato. Inoltre, il numero di coefficienti delle caratteristiche impulsive non deve essere troppo elevato affinché il filtro possa essere realizzato su hardware di calcolo ragionevole.
Per i filtri con passo di 1 MHz, scegliamo una larghezza di banda del filtro passa-basso (metà della larghezza di banda del filtro a banda) di 500 kHz, e la lunghezza della caratteristica impulsiva sarà di 480 campionamenti. Per i filtri con passo di 2 MHz, questi parametri saranno di 1 MHz e 240 campionamenti, rispettivamente. Scegliamo il tipo di finestra di Kaiser. Calcoliamo le caratteristiche impulsive in filterDesigner e le esportiamo in formato C-header:
Screenshot da filterDesigner
È possibile risolvere il problema in modo diretto: costruendo un array DDC corrispondente al numero di filtri (). Questo approccio è vantaggioso per le FPGA, dove è possibile risparmiare riducendo la bit-width dei calcolatori delle prime fasi. Inoltre, le FPGA sono il modo più energeticamente efficiente di implementazione. Ma i costi di lavoro per questo metodo sono i più elevati.
Nell'esecuzione della griglia di filtri su GPU attualmente popolari, emerge la possibilità di implementare un algoritmo più interessante: una griglia di filtri polifase basata su FFT, disponibile su CUDA dalla libreria. Nella letteratura straniera, l'algoritmo è chiamato Polyphase o WOLA (Weight, Overlap and Add) FFT Filterbank. La mia pigrizia nel disegnare non mi consente di fornire una spiegazione visiva da solo. In rete ci sono molti materiali sull'argomento, in particolare un grafico molto chiaro è stato realizzato a pagina 11 (un grande grazie agli autori rispettabili), ecco il grafico:
Cercherò di spiegare il circuito di elaborazione con parole mie. Coloro che sono facilmente impressionabili sono pregati di non leggere.
Cercherò di spiegare il circuito di elaborazione con parole mie, nei limiti delle mie capacità metodologiche. L'FFT è la convoluzione del segnale in ingresso con l'intero spettro delle armoniche complesse ortogonali, che si inseriscono nell'intervallo della risposta all'impulso. La risposta all'impulso del filtro, sulla quale viene moltiplicato il segnale prima di entrare nell'FFT, è modulata da questo spettro di armoniche. In altre parole, l'inviluppo delle risposte all'impulso dei filtri risultanti dalla griglia viene portato fuori dai calcoli. Successivamente, lo spettro delle armoniche viene sottoposto a sotto campionamento in qualche modo, a causa dell'espansione della larghezza di banda del filtro rispetto al filtro in finestra rettangolare. Nella figura vediamo un sotto campionamento di quattro. In altre parole, dopo aver espanso la larghezza di banda con una finestra di Kaiser (con un aumento simultaneo dell'attenuazione nella banda di ritardo), non abbiamo più bisogno di tutti i filtri, ma solo di un quarto di essi. Gli altri sono ridondanti, le loro caratteristiche di frequenza si sovrappongono. Delle quattro posizioni di FFT consecutive, si sceglie solo la zero, il cui calcolo consiste nella somma delle quattro
posizioni di ingresso, prese a intervallo pari a un quarto della durata dell'FFT originale.
Sceglieremo l'hardware che abbiamo a disposizione. Questa è la scheda di input della società "Strumentazioni Sistemi" FMC126P. Ne ho già parlato in un precedente articolo. . Nel connettore FMC della scheda è inserito un sottocomponente della stessa azienda con trasmettitore AD9371 con una larghezza di banda di 100 MHz. L'intero flusso dal trasmettitore può essere continuamente inviato al computer per l'elaborazione.
Scegliamo una scheda video con GPU GTX 1050. (Ho mentito, è stata lei a scegliere noi: era tutto ciò che avevamo a disposizione, strappata da un calcolatore per il calcolo delle antenne, ma è stato sorprendente vedere un comb generator funzionante). Passiamo alla parte software.
Sfortunatamente, a causa delle licenze, non possiamo pubblicare il codice completo. Possiamo mostrare solo i kernel GPU. Comunque, il resto del codice non è particolarmente interessante.
Ecco il kernel che esegue la moltiplicazione del segnale per la finestra e la somma, e il wrapper per la sua chiamata:
__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();
}
Il codice di elaborazione del segnale che chiama questo kernel ripete esattamente lo schema dell'algoritmo mostrato nella figura sopra, quindi non vedo il senso di riportarlo qui.
I comb generator sono stati verificati sullo spettro di uscita dei canali in tempo reale. Al ricevitore AD9371 è stato fornito un segnale del generatore di segnali a 2450 MHz, la selettività dei filtri corrispondeva a quella calcolata.
In programma: adattamento del software alla scheda XRTX e realizzazione della ricerca dei pacchetti, se a qualcuno interessa o se ci sarà tempo libero.
Tutto il lavoro sul software è stato svolto da , gloria a lui!
Fonte: habr.com
