From 257b296e412422c9668a04e39debb03951dfb864 Mon Sep 17 00:00:00 2001 From: Nika Ghaderi Date: Wed, 24 Sep 2025 17:34:09 +0000 Subject: [PATCH] Implement CUDA Noise and Phase Noise Kernel This commit adds the CUDA implementation for adding Additive White Gaussian Noise (AWGN) and phase noise to the signal. The code is located in phase_noise.cu. - The `add_noise_and_phase_noise_kernel` uses the NVIDIA cuRAND library for efficient parallel random number generation. - Helper functions are included to create, initialize, and destroy the cuRAND states required by the kernel. - Placeholder for a C-callable wrapper, `add_noise_cuda`, is also provided. --- openair1/SIMULATION/TOOLS/CMakeLists.txt | 1 + openair1/SIMULATION/TOOLS/oai_cuda.h | 99 ++++++++---- openair1/SIMULATION/TOOLS/phase_noise.cu | 190 +++++++++++++++++++++++ 3 files changed, 263 insertions(+), 27 deletions(-) create mode 100644 openair1/SIMULATION/TOOLS/phase_noise.cu diff --git a/openair1/SIMULATION/TOOLS/CMakeLists.txt b/openair1/SIMULATION/TOOLS/CMakeLists.txt index 1a7b5e8f81..a1937dfa64 100644 --- a/openair1/SIMULATION/TOOLS/CMakeLists.txt +++ b/openair1/SIMULATION/TOOLS/CMakeLists.txt @@ -22,6 +22,7 @@ if(CUDA_ENABLE) add_library(oai_cuda_lib STATIC multipath_channel.cu + phase_noise.cu ) target_link_libraries(oai_cuda_lib PUBLIC CUDA::toolkit) diff --git a/openair1/SIMULATION/TOOLS/oai_cuda.h b/openair1/SIMULATION/TOOLS/oai_cuda.h index f711e26ea8..9a590d0f8e 100644 --- a/openair1/SIMULATION/TOOLS/oai_cuda.h +++ b/openair1/SIMULATION/TOOLS/oai_cuda.h @@ -24,47 +24,92 @@ #include - #ifdef __NVCC__ +typedef struct complex16 { + int16_t r; + int16_t i; +} c16_t; #else - #include "PHY/TOOLS/tools_defs.h" +#include "PHY/TOOLS/tools_defs.h" #endif #ifdef __NVCC__ - #include - __device__ float2 complex_mul(float2 a, float2 b); +#include +__device__ float2 complex_mul(float2 a, float2 b); - __global__ void multipath_channel_kernel( - const float2* __restrict__ d_channel_coeffs, - const float* __restrict__ tx_sig, - float2* __restrict__ rx_sig, - int num_samples, - int channel_length, - int nb_tx, - int nb_rx); +__global__ void multipath_channel_kernel(const float2 *__restrict__ d_channel_coeffs, + const float *__restrict__ tx_sig, + float2 *__restrict__ rx_sig, + int num_samples, + int channel_length, + int nb_tx, + int nb_rx); + +__global__ void add_noise_and_phase_noise_kernel(const float2 *__restrict__ r_sig, + short2 *__restrict__ output_sig, + curandState_t *states, + int num_samples, + float sigma, + float pn_std_dev, + bool apply_phase_noise); #endif // __NVCC__ - #ifdef __cplusplus extern "C" { #endif - void multipath_channel_cuda( - float **rx_sig_re, float **rx_sig_im, - int nb_tx, int nb_rx, int channel_length, - uint32_t length, uint64_t channel_offset, - float *h_channel_coeffs, - void *d_tx_sig, void *d_rx_sig, - void *d_channel_coeffs, - void *h_tx_sig_pinned - ); +void multipath_channel_cuda(float **rx_sig_re, + float **rx_sig_im, + int nb_tx, + int nb_rx, + int channel_length, + uint32_t length, + uint64_t channel_offset, + float *h_channel_coeffs, + void *d_tx_sig, + void *d_rx_sig, + void *d_channel_coeffs, + void *h_tx_sig_pinned); - void interleave_channel_output_cuda(float **rx_sig_re, - float **rx_sig_im, - void **output_interleaved, - int nb_rx, - int num_samples); +void add_noise_cuda(const float **r_re, + const float **r_im, + c16_t **output_signal, + int num_samples, + int nb_rx, + float sigma2, + double ts, + int slot_offset, + int delay, + uint16_t pdu_bit_map, + uint16_t ptrs_bit_map, + void *d_r_sig, + void *d_output_sig, + void *d_curand_states, + void *h_r_sig_pinned, + void *h_output_sig_pinned); + +/** + * \brief Interleaves separate real and imaginary signal buffers into a complex format. + * + * This utility function takes separate planar buffers for the real (I) and imaginary (Q) + * components of a signal and combines them into a single interleaved buffer of float2 + * structs, which is a common format for CUDA processing. + * + * \param[in] rx_sig_re Array of pointers to the real parts of the input signals. + * \param[in] rx_sig_im Array of pointers to the imaginary parts of the input signals. + * \param[out] output_interleaved Pointer to an array of pointers where the interleaved output will be stored. + * \param[in] nb_rx The number of receive antennas (i.e., the number of signals). + * \param[in] num_samples The number of samples per signal. + * + * \note The caller is responsible for allocating memory for the `output_interleaved` buffer. It should + * point to an array of `nb_rx` pointers, where each of those pointers is allocated with a size of + * `num_samples * sizeof(float2)`. + */ +void interleave_channel_output_cuda(float **rx_sig_re, float **rx_sig_im, void **output_interleaved, int nb_rx, int num_samples); + +void *create_and_init_curand_states_cuda(int num_elements, unsigned long long seed); +void destroy_curand_states_cuda(void *d_curand_states); #ifdef __cplusplus } diff --git a/openair1/SIMULATION/TOOLS/phase_noise.cu b/openair1/SIMULATION/TOOLS/phase_noise.cu new file mode 100644 index 0000000000..3bca71fdf0 --- /dev/null +++ b/openair1/SIMULATION/TOOLS/phase_noise.cu @@ -0,0 +1,190 @@ +/* + * Licensed to the OpenAirInterface (OAI) Software Alliance under one or more + * contributor license agreements. See the NOTICE file distributed with + * this work for additional information regarding copyright ownership. + * The OpenAirInterface Software Alliance licenses this file to You under + * the OAI Public License, Version 1.1 (the "License"); you may not use this file + * except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.openairinterface.org/?page_id=698 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + *------------------------------------------------------------------------------- + * For more information about the OpenAirInterface (OAI) Software Alliance: + * contact@openairinterface.org + */ + +#include +#include +#include +#include +#include +#include "oai_cuda.h" + +#define CHECK_CUDA(val) checkCuda((val), __FILE__, __LINE__) +static void checkCuda(cudaError_t result, const char *file, int line) +{ + if (result != cudaSuccess) { + fprintf(stderr, "CUDA Error at %s:%d: %s\n", file, line, cudaGetErrorString(result)); + exit(1); + } +} + +__device__ __forceinline__ float2 complex_mul(float2 a, float2 b) +{ + return make_float2(a.x * b.x - a.y * b.y, a.x * b.y + a.y * b.x); +} + +__global__ void init_curand_states_kernel(curandState_t *states, unsigned long long seed, int num_elements) +{ + int idx = blockIdx.x * blockDim.x + threadIdx.x; + if (idx < num_elements) { + curand_init(seed, idx, 0, &states[idx]); + } +} + +__global__ void add_noise_and_phase_noise_kernel(const float2 *__restrict__ r_sig, + short2 *__restrict__ output_sig, + curandState_t *states, + int num_samples, + float sigma, + float pn_std_dev, + bool apply_phase_noise) +{ + const int i = blockIdx.x * blockDim.x + threadIdx.x; + const int ii = blockIdx.y; + + if (i >= num_samples) + return; + + // Each thread handles one sample and its corresponding cuRAND state + curandState_t local_state = states[ii * num_samples + i]; + + float2 noisy_signal = r_sig[ii * num_samples + i]; + + float2 awgn = curand_normal2(&local_state); + noisy_signal.x += awgn.x * sigma; + noisy_signal.y += awgn.y * sigma; + + float2 final_signal = noisy_signal; + + if (apply_phase_noise) { + float phase_error = curand_normal(&local_state) * pn_std_dev; + float cos_phi, sin_phi; + __sincosf(phase_error, &sin_phi, &cos_phi); + float2 phase_rot = make_float2(cos_phi, sin_phi); + final_signal = complex_mul(noisy_signal, phase_rot); + } + + states[ii * num_samples + i] = local_state; + output_sig[ii * num_samples + i] = make_short2((short)fmaxf(-32768.0f, fminf(32767.0f, final_signal.x)), + (short)fmaxf(-32768.0f, fminf(32767.0f, final_signal.y))); +} + +extern "C" { + +void add_noise_cuda(const float **r_re, + const float **r_im, + c16_t **output_signal, + int num_samples, + int nb_rx, + float sigma2, + double ts, + int slot_offset, + int delay, + uint16_t pdu_bit_map, + uint16_t ptrs_bit_map, + void *d_r_sig_void, + void *d_output_sig_void, + void *d_curand_states_void, + void *h_r_sig_void, + void *h_output_temp_void) +{ + float2 *d_r_sig = (float2 *)d_r_sig_void; + short2 *d_output_sig = (short2 *)d_output_sig_void; + curandState_t *d_curand_states = (curandState_t *)d_curand_states_void; + + float2 *kernel_input_ptr; + +#if defined(USE_UNIFIED_MEMORY) + for (int ii = 0; ii < nb_rx; ii++) { + for (int i = 0; i < num_samples; i++) { + d_r_sig[ii * num_samples + i] = make_float2(r_re[ii][i], r_im[ii][i]); + } + } + kernel_input_ptr = d_r_sig; +#elif defined(USE_ATS_MEMORY) + float2 *h_r_sig = (float2 *)h_r_sig_void; + for (int ii = 0; ii < nb_rx; ii++) { + for (int i = 0; i < num_samples; i++) { + h_r_sig[ii * num_samples + i] = make_float2(r_re[ii][i], r_im[ii][i]); + } + } + kernel_input_ptr = h_r_sig; +#else // EXPLICIT COPY + float2 *h_r_sig = (float2 *)h_r_sig_void; + for (int ii = 0; ii < nb_rx; ii++) { + for (int i = 0; i < num_samples; i++) { + h_r_sig[ii * num_samples + i] = make_float2(r_re[ii][i], r_im[ii][i]); + } + } + CHECK_CUDA(cudaMemcpy(d_r_sig, h_r_sig, nb_rx * num_samples * sizeof(float2), cudaMemcpyHostToDevice)); + kernel_input_ptr = d_r_sig; +#endif + + dim3 threadsPerBlock(256, 1); + dim3 numBlocks((num_samples + threadsPerBlock.x - 1) / threadsPerBlock.x, nb_rx); + float pn_variance = 1e-5f * 2.0f * 3.1415926535f * 300.0f * (float)ts; + bool apply_phase_noise = (pdu_bit_map & ptrs_bit_map); + add_noise_and_phase_noise_kernel<<>>(kernel_input_ptr, + d_output_sig, + d_curand_states, + num_samples, + sqrtf(sigma2 / 2.0f), + sqrtf(pn_variance), + apply_phase_noise); + +#if defined(USE_UNIFIED_MEMORY) + CHECK_CUDA(cudaDeviceSynchronize()); + for (int ii = 0; ii < nb_rx; ii++) { + for (int i = 0; i < num_samples; i++) { + short2 result = d_output_sig[ii * num_samples + i]; + output_signal[ii][i + slot_offset + delay].r = result.x; + output_signal[ii][i + slot_offset + delay].i = result.y; + } + } +#else + short2 *h_output_temp = (short2 *)h_output_temp_void; + CHECK_CUDA(cudaMemcpy(h_output_temp, d_output_sig, nb_rx * num_samples * sizeof(short2), cudaMemcpyDeviceToHost)); + CHECK_CUDA(cudaDeviceSynchronize()); + for (int ii = 0; ii < nb_rx; ii++) { + memcpy(output_signal[ii] + slot_offset + delay, h_output_temp + ii * num_samples, num_samples * sizeof(short2)); + } +#endif +} + +void *create_and_init_curand_states_cuda(int num_elements, unsigned long long seed) +{ + void *d_states_void; + CHECK_CUDA(cudaMalloc(&d_states_void, num_elements * sizeof(curandState_t))); + + int threads = 256; + int blocks = (num_elements + threads - 1) / threads; + init_curand_states_kernel<<>>((curandState_t *)d_states_void, seed, num_elements); + CHECK_CUDA(cudaDeviceSynchronize()); + + return d_states_void; +} + +void destroy_curand_states_cuda(void *d_curand_states) +{ + if (d_curand_states) { + CHECK_CUDA(cudaFree(d_curand_states)); + } +} +} // extern "C"