🥇¿Recepción de todo Bluetooth a la vez en SDR con CUDA? Fácil

Recientemente, colegas del "sector" me han preguntado de manera independiente: ¿cómo obtener de un receptor SDR todos los canales Bluetooth al mismo tiempo? La banda lo permite, hay SDR con un ancho de banda de salida de 80 MHz o más. Por supuesto, se podría hacer esto en FPGA, pero el tiempo de desarrollo sería bastante largo. Hace tiempo supe que hacerlo en GPU es bastante simple, ¡pero de esta manera!

El estándar Bluetooth define el nivel físico en dos versiones: Classic y Low Energy. Hay especificaciones aquí. El documento es extremadamente extenso, leerlo en su totalidad puede ser perjudicial para el cerebro. Afortunadamente, grandes empresas que producen equipos de medición tienen los recursos para crear documentos visuales sobre el tema. Tektronix y National Instruments, por ejemplo. No tengo ninguna posibilidad de competir con ellos en la calidad de presentación del material. A los interesados les pido que estudien a través de los enlaces.

Todo lo que necesito saber sobre el nivel físico para crear un filtro multicanal es el paso de la cuadrícula de frecuencias y la velocidad de modulación. Se resumen en una tabla en uno de los documentos mencionados:

Por lo tanto, necesitamos dividir la banda de 80 MHz en 79 filtros con un ajuste de 1 MHz y, al mismo tiempo, en 40 filtros con un ajuste de 2 MHz. Las frecuencias de muestreo de las salidas de los filtros deben ser de 1 MHz y 2 MHz, respectivamente.

Por lo tanto, necesitamos dos peines de filtros.

Para comenzar, elijamos los parámetros de estos filtros basándonos en los anchos de banda de las señales Bluetooth Classic y Bluetooth Low Energy. Necesitamos sus características de impulso para calcular la carga en el dispositivo computacional del filtro. Aquí vale la pena aclarar que hemos seleccionado las longitudes de las características de impulso basándonos en los requisitos de un algoritmo de filtrado "rápido". La esencia no cambia. Y el número de coeficientes de la característica de impulso no debe ser demasiado grande para que el filtro sea realizable en un hardware computacional razonable.

Para los filtros con un ajuste de 1 MHz, seleccionaremos un ancho de banda de paso bajo (la mitad del ancho de banda del filtro de banda) de 500 kHz, y ajustaremos la longitud de la característica de impulso a 480 muestras. Para los filtros con un ajuste de 2 MHz, elegiremos estos parámetros como 1 MHz y 240 muestras, respectivamente. El tipo de ventana será Kaiser. Calcularemos las características de impulso en filterDesigner y las exportaremos en formato de encabezado C:

Capturas de pantalla de filterDesigner

Se puede abordar el problema de forma directa: construir un arreglo DDC correspondiente al número de filtros (Convertidor Digital Descendente). Este enfoque es bueno para FPGAs, donde se puede economizar reduciendo la precisión de los procesadores de las primeras etapas. Además, las FPGAs son el método más energéticamente eficiente de implementar esto. Pero el esfuerzo requerido para este método es el más alto.

Al ejecutar un peinado de filtros en las GPU populares actualmente, surge la oportunidad de implementar un algoritmo más sofisticado: un peinado polifásico de filtros basado en FFT, que está disponible en CUDA a partir de la biblioteca. En la literatura extranjera, el algoritmo se llama Peine Polifásico o WOLA (Weight, Overlap and Add) FFT Filterbank. Mi pereza para dibujar no me permite realizar una explicación visual por mí mismo. Hay muchos materiales sobre el tema en la red, especialmente un gráfico visual realizado aquí en la página 11 (muchas gracias a los respetables autores), aquí está:

Intentaré explicar el esquema de procesamiento con mis propias palabras. A los que son de nervios débiles, les pido que no lean.

Intentaré explicar el esquema de procesamiento con mis propias palabras dentro de mis capacidades metodológicas. La FFT es la convolución de la señal de entrada con todo el espectro de armónicos complejos ortogonales que se ajustan al intervalo de la respuesta al impulso. La respuesta al impulso del filtro, multiplicada por la señal antes de entrar en la FFT, es modulada por este espectro de armónicos. En otras palabras, el sobre de las respuestas al impulso de los filtros resultantes del peinado se saca de la ecuación. Luego, el espectro de armónicos se muestrea en un número determinado de veces, debido a la ampliación del ancho de banda del filtro en relación con el filtro de ventana rectangular. En la ilustración vemos un muestreo de cuatro. En otras palabras, después de ampliar el ancho usando la ventana de Kaiser (con un aumento paralelo de la atenuación en la banda de retardo), solo necesitamos una cuarta parte de los filtros. Los demás son redundantes, sus características de frecuencia se superponen. De cuatro puntos de la FFT que van en secuencia, elegimos solo el cero, cuyo cálculo es la suma de cuatro
puntos de entrada, tomados en un intervalo igual a un cuarto de la duración de la FFT original.

Elegiremos el hardware que tengamos a mano. Esta es la placa de entrada de la empresa "Instrumental Systems" FMC126P. Ya he escrito sobre ella en uno de mis anteriores el artículoEn el conector FMC de la placa se ha insertado un submódulo de la misma empresa con un transceptor AD9371 que tiene un ancho de banda de 100 MHz. Todo el flujo del transceptor puede ser transmitido continuamente a una computadora para su procesamiento.

Elijamos una tarjeta gráfica con GPU GTX 1050. (Mintió, fue ella quien nos eligió: era todo lo que teníamos a mano, fue extraída de un procesador para calcular antenas, pero fue aún más sorprendente ver un combiner operativo). Pasemos a la parte de software.

Desafortunadamente, debido a licencias, no podemos publicar el código completo. Solo podemos mostrar los núcleos de la GPU. Sin embargo, el resto del código no es tan interesante.

Aquí está el núcleo que realiza la multiplicación de la señal por la ventana y la suma, y el envoltorio para su invocación:

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

El código de procesamiento de señal que llama a este núcleo reproduce exactamente el esquema del algoritmo mostrado en la figura anterior, por lo que no veo sentido en incluirlo aquí.

Los combinadores fueron probados en el espectro de salida de los canales en tiempo real. A la entrada de AD9371 se aplicó una señal de generador de 2450 MHz, la selectividad de los filtros correspondía al cálculo.