🥇¿Recibir todos los Bluetooth a la vez en SDR con CUDA? Fácil

Recientemente, mis colegas en la “tienda” comenzaron a preguntarme de forma independiente: ¿cómo obtener todos los canales Bluetooth simultáneamente desde un receptor SDR? El ancho de banda lo permite, existen SDR con un ancho de banda de salida de 80 MHz o más. Por supuesto, puedes hacer esto en una FPGA, pero ese tiempo de desarrollo será bastante largo. Sé desde hace mucho tiempo que hacer esto en una GPU es bastante simple, ¡pero hacerlo así!

El estándar Bluetooth define una capa física en dos versiones: Classic y Low Energy. Especificación disponible aquí. El documento es terriblemente extenso, leerlo completo es peligroso para el cerebro. Afortunadamente, las grandes empresas de tecnología de medición tienen los medios para crear documentos visuales sobre un tema. Tektronix и National Instruments, Por ejemplo. No tengo absolutamente ninguna posibilidad de competir con ellos en cuanto a la calidad de presentación del material. Si está interesado, siga los enlaces.

Todo lo que necesito saber sobre la capa física para crear un filtro multicanal es el espaciado de la cuadrícula de frecuencias y la tasa de modulación. Están tabulados en uno de los siguientes documentos:

Por tanto, necesitamos cortar la banda de 80 MHz en 79 filtros con un paso de sintonía de 1 MHz y, al mismo tiempo, en 40 filtros con un paso de sintonía de 2 MHz. Las frecuencias de muestreo de las salidas del filtro deben ser de 1 MHz y 2 MHz, respectivamente.

Por tanto, necesitamos dos bancos de filtros.

Primero, seleccionemos los parámetros de estos filtros según las bandas de señal de Bluetooth Classic y Bluetooth Low Energy. Necesitamos sus respuestas de impulso para calcular la carga en el dispositivo informático del filtro. Aquí vale la pena mencionar de inmediato que elegimos la duración de la respuesta al impulso en función de los requisitos de un algoritmo de filtrado "rápido". La esencia no cambia a partir de esto. Y el número de coeficientes de respuesta al impulso no debe ser demasiado grande para que el filtro se implemente en equipos informáticos sensatos.

Para filtros con un paso de 1 MHz, seleccionamos un ancho de banda de filtro de paso bajo (la mitad del ancho de banda de un filtro de paso de banda) de 500 kHz y ajustamos la longitud de la respuesta de impulso a 480 derivaciones. Para filtros con un paso de 2 MHz, seleccionaremos estos parámetros como 1 MHz y 240 taps, respectivamente. Elegimos el tipo de ventana Kaiser. Calculemos las respuestas de impulso en filterDesigner y carguemoslas en formato de encabezado C:

Capturas de pantalla de filterDesigner

Puede resolver el problema mediante fuerza bruta: cree una matriz DDC correspondiente a la cantidad de filtros (Convertidor descendente digital). Este enfoque es bueno para FPGA, donde es posible ahorrar reduciendo la capacidad de bits de las computadoras de la primera etapa. Además, FPGA es el método de implementación más eficiente energéticamente. Pero los costos laborales con este método son los más altos.

Cuando se ejecuta un banco de filtros en las GPU populares de hoy en día, es posible implementar un algoritmo más sofisticado: un banco de filtros polifásico basado en FFT, que en CUDA está disponible en la biblioteca. En la literatura extranjera, el algoritmo se llama Polyphase o WOLA (Weight, Overlap and Add) FFT Filterbank. La pereza en el dibujo no me da la oportunidad de completar una explicación visual por mi cuenta. Hay mucho material sobre el tema en Internet, se hace un gráfico particularmente claro. aquí en la página 11 (muchas gracias a los respetados autores), aquí está:

Intentaré explicar el esquema de procesamiento con mis propias palabras. Por favor, no lo lea para los débiles de corazón.

Intentaré explicar el esquema de procesamiento con mis propias palabras dentro del marco de mis capacidades metodológicas. La FFT es la convolución de la señal de entrada con todo el espectro de armónicos ortogonales complejos que caben dentro del intervalo de respuesta al impulso. La respuesta al impulso del filtro, por el cual se multiplica la señal antes de la entrada FFT, es modulada por este espectro armónico. En otras palabras, la envolvente de respuesta al impulso de los filtros de peine resultantes está entre corchetes. Además, el espectro armónico se adelgaza un cierto número de veces, debido a la expansión del ancho de banda del filtro con respecto al filtro en la ventana rectangular. En la imagen vemos un adelgazamiento de cuatro. En otras palabras, después de expandir la banda con la ventana de Kaiser (con un aumento simultáneo de la atenuación en la banda de parada), ya no necesitamos todos los filtros, sino solo una cuarta parte de ellos. El resto son redundantes, sus características de frecuencia se superponen. De los cuatro puntos FFT seguidos, seleccionamos solo el cero, cuyo cálculo es la suma de cuatro
puntos de entrada tomados después de un tiempo igual a un cuarto de la duración de la FFT original.

Elegiremos la plancha que tengamos a mano. Esta es la placa de entrada de Instrumental Systems FMC126P. Ya escribí sobre esto en uno anterior. статье. En el conector FMC de la placa se inserta un submódulo de la misma empresa con un transceptor AD9371 con banda de 100 MHz. Todo el flujo del transceptor se puede transmitir continuamente a una computadora para su procesamiento.

Elijamos una tarjeta de video con GPU GTX 1050. (Mentí, fue la que nos eligió: era todo lo que tenía a mano, lo arrancaron de la computadora para calcular las antenas, pero fue aún más sorprendente ver un peine que funcione). Pasemos a la parte del software.

Lamentablemente, debido a licencias, no podemos publicar el código completo. Sólo podemos mostrar núcleos de GPU. Sin embargo, el resto del código no es particularmente interesante.

Aquí está el núcleo que realiza la multiplicación y suma de ventanas de señales, y un contenedor para llamarlo:

__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<<<dataSize / windowStep - 3, windowSize, 1024, stream>>>(data, window, windowSize, result);
    return cudaGetLastError();
}

El código de procesamiento de señales que llama a este núcleo repite exactamente el diagrama del algoritmo que se muestra en la figura anterior, por lo que no veo ninguna razón para presentarlo aquí.

Los peines se probaron en el espectro de salida de los canales en tiempo real. Se suministró una señal del generador de señal de 9371 MHz a la entrada AD2450 y la selectividad de los filtros correspondió a la calculada.