maki/openairinterface5g
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

Implement CUDA Noise and Phase Noise Kernel

Browse Source 
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.
This commit is contained in:
Nika Ghaderi
2026-02-09 09:29:36 +00:00
committed by Jaroslava Fiedlerova
parent 97a14c3960
commit 257b296e41
3 changed files with 263 additions and 27 deletions
Download Patch File Download Diff File Expand all files Collapse all files
openair1/SIMULATION/TOOLS/CMakeLists.txt
+1
View File
@@ -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)
openair1/SIMULATION/TOOLS/oai_cuda.h
+72 -27
View File
@@ -24,47 +24,92 @@
#include <stdint.h>
#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 <curand_kernel.h>
__device__ float2 complex_mul(float2 a, float2 b);
#include <curand_kernel.h>
__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
}
openair1/SIMULATION/TOOLS/phase_noise.cu
+190
View File
@@ -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 <stdio.h>
#include <stdlib.h>
#include <string.h>
#include <cuda_runtime.h>
#include <curand_kernel.h>
#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<<<numBlocks, threadsPerBlock>>>(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<<<blocks, threads>>>((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"