Ultimamente, i colleghi del "settore" hanno iniziato a chiedermi indipendentemente: come si possono ricevere simultaneamente tutti i canali Bluetooth da un solo ricevitore SDR? La banda lo consente, ci sono SDR con una larghezza di banda in uscita di 80 MHz o più. Certo, si può fare con un FPGA, ma il tempo necessario per tale sviluppo sarebbe piuttosto lungo. Sapevo già da tempo che farlo su GPU è abbastanza semplice, ma non così!
Lo standard Bluetooth definisce il livello fisico in due versioni: Classic e Low Energy. La specifica è disponibile . Il documento è estremamente voluminoso, leggerlo per intero è dannoso per il cervello. Fortunatamente, grandi aziende che producono attrezzature di misura dispongono delle risorse per creare documenti visivi sul tema. e , ad esempio. Non ho alcuna possibilità di competere con loro per qualità di presentazione del materiale. Invito chi è interessato a studiare tramite i link.
Tutto ciò che devo sapere sul livello fisico per creare un filtro multicanale sono il passo della griglia di frequenze e la velocità di modulazione. Sono riassunti in una tabella in uno dei documenti menzionati:
Pertanto, dobbiamo suddividere una banda di 80 MHz in 79 filtri con intervallo di regolazione di 1 MHz e, contemporaneamente, in 40 filtri con intervallo di regolazione di 2 MHz. Le frequenze di campionamento in uscita dai filtri devono essere rispettivamente di 1 MHz e 2 MHz.
Pertanto, abbiamo bisogno di due griglie di filtri.
Iniziamo a scegliere i parametri di questi filtri in base alle bande di segnale di Bluetooth Classic e Bluetooth Low Energy. Abbiamo bisogno delle loro caratteristiche impulsive per calcolare il carico sul dispositivo di calcolo del filtro. È importante sottolineare che le lunghezze delle risposte impulsive sono state selezionate in base ai requisiti dell'algoritmo di filtraggio 'veloce'. La sostanza non cambia. Inoltre, il numero di coefficienti della risposta impulsiva non dovrebbe essere troppo elevato affinché il filtro possa essere implementato su un hardware di calcolo ragionevole.
Per i filtri con passo di 1 MHz, selezioniamo una banda passante del filtro passa-basso (metà della banda passante del filtro a banda) di 500 kHz, e regoleremo la lunghezza della risposta all'impulso a 480 campioni. Per i filtri con passo di 2 MHz, questi parametri saranno 1 MHz e 240 campioni, rispettivamente. Scegliamo il tipo di finestra Kaiser. Calcoliamo le risposte all'impulso in filterDesigner e le esportiamo nel formato C-header:
Screenshot da filterDesigner
È possibile affrontare il problema in modo diretto: costruire un array DDC corrispondente al numero di filtri (). Questo approccio è vantaggioso per le FPGA, dove è possibile risparmiare riducendo la dimensione dei calcolatori di primo livello. Inoltre, le FPGA rappresentano il modo più energeticamente efficiente di implementazione. Tuttavia, il lavoro richiesto con questo metodo è il più elevato.
Con l'implementazione di filtri a pettine su GPU attuali, emerge la possibilità di realizzare un algoritmo più sofisticato: il filtro a pettine polifasico basato su FFT, che è disponibile su CUDA tramite una libreria. Nella letteratura estera, l'algoritmo è chiamato Polyphase o WOLA (Weight, Overlap and Add) FFT Filterbank. La mia pigrizia nei disegni non mi permette di fornire una spiegazione visiva autonoma. In rete ci sono molti materiali sull'argomento, in particolare un grafico esplicativo è presente a pagina 11 (un enorme ringraziamento agli autorevoli autori), eccolo:
Cercherò di spiegare lo schema di elaborazione con parole mie. Le persone sensibili sono pregate di non leggere.
Proverò a spiegare lo schema di elaborazione con parole mie, secondo le mie capacità metodologiche. L'FFT è una convoluzione del segnale in ingresso con tutto lo spettro delle armoniche complesse ortogonali, che si collocano nell'intervallo della risposta all'impulso. La risposta all'impulso del filtro, sul quale viene moltiplicato il segnale prima dell'ingresso nell'FFT, è modulata da questo spettro di armoniche. In altre parole, l'inviluppo delle risposte all'impulso dei filtri risultanti della griglia viene rimosso. Successivamente, lo spettro di armoniche viene decimato di un certo numero di volte, a causa dell'espansione della banda passante del filtro rispetto al filtro a finestra rettangolare. Nella figura vediamo una decimazione di quattro. In altre parole, dopo l'espansione della banda con la 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. Tra quattro punti dell'FFT, consecutivi, scegliamo solo il punto zero, il cui calcolo consiste nella somma dei quattro
punti di ingresso, presi a un intervallo pari a un quarto della durata dell'FFT originale.
Selezioniamo l'hardware che abbiamo a disposizione. Si tratta della scheda di input dell'azienda "Strumentali Sistemi" FMC126P. Di essa ho già parlato in un precedente . Nel connettore FMC della scheda è inserito un sottocomponente della stessa azienda con trasmettitore AD9371 con una banda di 100 MHz. L'intero flusso dal trasmettitore può essere trasmesso continuamente al computer per l'elaborazione.
Scegliamo una scheda video con GPU GTX 1050. (Ho mentito, è stata lei a sceglierci: è tutto ciò che avevamo a disposizione, strappata da un calcolatore per il calcolo delle antenne, ma è stato sorprendente vedere un sistema operativo funzionante). Passiamo alla parte software.
Sfortunatamente, a causa delle licenze non possiamo pubblicare il codice completo. Possiamo mostrare solo i nuclei GPU. Tuttavia, il resto del codice non è particolarmente interessante.
Ecco un nucleo che esegue la moltiplicazione del segnale per la finestra e l'addizione, con un 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 richiama questo kernel segue esattamente lo schema dell'algoritmo mostrato nell'immagine sopra, quindi non vedo il motivo di riportarlo qui.
Le griglie sono state verificate sullo spettro di uscita dei canali in tempo reale. Al ingresso dell'AD9371 è stato fornito un segnale da un generatore a 2450 MHz, la selettività dei filtri corrispondeva a quella calcolata.
In programma: adattamento del software sulla scheda XRTX e implementazione della ricerca dei pacchetti, se a qualcuno interessa o se ci sarà tempo libero.
Tutto il lavoro relativo al software è stato svolto da , gloria a lui!
Fonte: habr.com
