CUDAを使用してSDRですべてのBluetoothを一度に受信しますか？簡単

最近、「ワークショップ」の同僚が独自に私に尋ね始めました。1つのSDRレシーバーからすべてのBluetoothチャネルを同時に取得する方法は？帯域幅が許せば、80MHz以上の出力帯域幅を持つSDRがあります。もちろん、FPGAでも実行できますが、開発時間はかなり長くなります。 GPUでこれを行うのは非常に簡単であることを私は長い間知っていましたが、そうです！

Bluetooth標準では、物理レイヤーをクラシックと低エネルギーの2つのバージョンで定義しています。仕様はこちらです。文書はひどく大きく、全体を読むことは脳にとって危険です。幸いなことに、大規模な計測会社には、トピックに関する視覚的なドキュメントを作成する手段があります。たとえば、TektronixとNationalInstruments。資料の提示の質の面で彼らと競争する機会は絶対にありません。興味のある方は、リンクをたどってください。

マルチチャネルフィルターを作成するために物理レイヤーについて知る必要があるのは、周波数グリッドのステップと変調率だけです。それらは、次のいずれかのドキュメントにまとめられています。

, 80 79 1 , , 40 2 . 1 2 , .

, .

Bluetooth Classic Bluetooth Low Energy. , . , "" . . , .

1 ( ) 500 , 480 . 2 1 240 , . . filterDesigner -header:

: DDC (Digital Down Converter). , . — - . .

GPU : , CUDA . Polyphase or WOLA (Weight, Overlap and Add) FFT Filterbank. . , 11 ( ), :

, . " " FMC126P. . FMC AD9371 100 . .

- GPU GTX 1050. (, : , , ). .

, - . GPU. , .

, , :

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

, , , , .

. AD9371 2450 , .

: XRTX , - .

gaudima, !