🥇 Receber todo o Bluetooth de uma vez no SDR com CUDA? Fácil

Recentemente, colegas do "setor" começaram a me perguntar, independentemente, como obter todos os canais Bluetooth ao mesmo tempo de um receptor SDR. A largura de banda permite, há SDRs com largura de banda de saída de 80 MHz ou mais. Claro, é possível fazer isso em FPGA, mas o tempo de desenvolvimento seria bastante longo. Eu já sabia há muito tempo que fazer isso em GPU é bem simples, mas a esse ponto!

O padrão Bluetooth define o nível físico em duas versões: Classic e Low Energy. A especificação existe aqui. O documento é imenso, lê-lo na íntegra é perigoso para a mente. Felizmente, grandes empresas que fabricam equipamentos de medição têm recursos para criar documentos informativos sobre o tema. Tektronix e National Instruments, por exemplo. Eu não tenho chances de concorrência com eles em termos de apresentação de material. Interessados, por favor, explorem os links.

Tudo o que preciso saber sobre o nível físico para criar um filtro multicanal é a malha de frequências e a taxa de modulação. Elas estão resumidas em uma tabela em um dos documentos mencionados:

Portanto, precisamos dividir a faixa de 80 MHz em 79 filtros com uma configuração de 1 MHz e, ao mesmo tempo, em 40 filtros com uma configuração de 2 MHz. As frequências de amostragem das saídas dos filtros devem ser 1 MHz e 2 MHz, respectivamente.

Assim, precisamos de duas grades de filtros.

Para começar, vamos escolher os parâmetros desses filtros com base nas faixas de sinais do Bluetooth Classic e do Bluetooth Low Energy. Precisamos de suas características de impulso para calcular a carga no dispositivo computacional do filtro. Aqui, deve-se esclarecer que as extensões das características de impulso foram escolhidas com base nas exigências de um algoritmo de filtragem "rápido". A essência não muda. E o número de coeficientes da característica de impulso não deve ser muito grande, para que o filtro possa ser implementado em hardware computacional razoável.

Para filtros com uma configuração de 1 MHz, vamos escolher uma largura de banda de um filtro passa-baixa (metade da largura de banda do filtro de banda) de 500 kHz, e adaptar a extensão da característica de impulso para 480 coeficientes. Para filtros com uma configuração de 2 MHz, vamos escolher esses parâmetros como 1 MHz e 240 coeficientes, respectivamente. O tipo de janela será de Kaiser. Vamos calcular as características de impulso no filterDesigner e exportá-las no formato C-header:

Capturas de tela do filterDesigner

É possível resolver o problema de forma direta: construir um array DDC correspondente ao número de filtros (Conversor Digital de Abaixamento). Essa abordagem é boa para FPGAs, onde se pode economizar reduzindo a precisão dos processadores de primeira ordem. Além disso, FPGAs são a forma mais energeticamente eficiente de implementação. Mas os custos de trabalho com esse método são os mais altos.

Ao realizar a grelha de filtros em GPUs populares atualmente, surge a possibilidade de implementar um algoritmo mais sofisticado: a grelha de filtros polifásica baseada em FFT, que está disponível na biblioteca CUDA. Na literatura estrangeira, o algoritmo é chamado de Polyphase ou WOLA (Weight, Overlap and Add) FFT Filterbank. A falta de ânimo para desenhar não me permite realizar uma explicação visual por conta própria. Há muitos materiais sobre o tema na internet, especialmente um gráfico ilustrativo foi feito aqui na página 11 (muito obrigado aos respeitáveis autores), aqui está ele:

Eu tentarei explicar o esquema de processamento com minhas próprias palavras. Os sensíveis devem evitar a leitura.

Tentarei explicar o esquema de processamento com minhas próprias palavras dentro das minhas possibilidades metodológicas. A FFT é a convolução do sinal de entrada com todo o espectro de harmônicos ortogonais complexos, que se ajustam ao intervalo da resposta ao impulso. A resposta ao impulso do filtro, sobre a qual o sinal é multiplicado antes da entrada da FFT, é modulada por esse espectro de harmônicos. Em outras palavras, o contorno das respostas ao impulso dos filtros resultantes da grelha é retirado da equação. Em seguida, o espectro de harmônicos é subamostrado em um certo número de vezes, devido à expansão da largura de banda do filtro em relação ao filtro com janela retangular. Na figura, vemos uma subamostragem em quatro. Em outras palavras, após a expansão da largura da janela de Kaiser (com o aumento simultâneo da atenuação na faixa de atraso), já não precisamos de todos os filtros, apenas de um quarto deles. Os outros são redundantes, suas características de frequência se sobrepõem. Das quatro pontos de FFT consecutivas, escolhemos apenas a zero, cujo cálculo é a soma das quatro
pontos de entrada, tomados a um intervalo igual a um quarto da duração da FFT original.

Escolheremos o hardware que temos à mão. Esta é a placa de entrada da empresa "Instrumental Systems" FMC126P. Já escrevi sobre isso em uma de minhas anteriores. artigoO módulo da mesma empresa com o transceptor AD9371 com largura de banda de 100 MHz está inserido na porta FMC da placa. Todo o fluxo do transceptor pode ser transmitido continuamente para o computador para processamento.

Vamos escolher uma placa de vídeo com GPU GTX 1050. (Mentira, foi ela que nos escolheu: era tudo que estava à mão, foi arrancada do computador para calcular as antenas, mas foi surpreendente ver uma grelha funcionando). Agora vamos para a parte de software.

Infelizmente, devido a licenças, não podemos publicar o código completo. Podemos mostrar apenas os núcleos do GPU. No entanto, o restante do código não é tão interessante.

Aqui está o núcleo que realiza a multiplicação do sinal por uma janela e a soma, e uma camada para sua chamada:

__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();
}

O código de processamento do sinal que chama esse núcleo repete exatamente o esquema do algoritmo mostrado na figura acima, por isso não vejo sentido em apresentá-lo aqui.

As grelhas foram verificadas no espectro de saída dos canais em tempo real. Um sinal do gerador de sinais de 2450 MHz foi enviado para a entrada do AD9371, e a seletividade dos filtros correspondia ao calculado.