yigid/openairinterface5g
1
0
Fork 0
mirror of https://gitlab.eurecom.fr/oai/openairinterface5g.git synced 2026-07-13 04:30:28 +00:00
Code Issues Packages Projects Releases Wiki Activity

Acceleration of channel convolution and noise generation

Browse Source 
Added two versions of channel convolution and noise generation:
 - accelerated via threadpool
 - accelerated using CUDA

The main difference between this and previous versions of channel
convolution implementations is that these functions take a real-world
approach to input/output where both could be split unevenly over a set of
buffers, e.g. ring buffers used in vrtsim.

The CUDA-accelerated version only works on systems with unified memory, e.g. NVidia
DGX Spark or GH
This commit is contained in:
Bartosz Podrygajlo
2026-02-25 12:00:56 +01:00
parent c7799886cf
commit 47ba5d3ca5
11 changed files with 1099 additions and 5 deletions
Download Patch File Download Diff File Expand all files Collapse all files

11
openair1/PHY/TOOLS/tests/CMakeLists.txt
View File

@@ -42,6 +42,17 @@ add_dependencies(tests test_sse_intrinsics)
add_test(NAME test_sse_intrinsics
COMMAND ./test_sse_intrinsics)
add_executable(benchmark_channel_pipeline benchmark_channel_pipeline.cpp test_channel_pipeline_tools.c)
target_link_libraries(benchmark_channel_pipeline PRIVATE UTIL SIMU PHY_COMMON LOG CONFIG_LIB shlib_loader m channel_pipeline benchmark::benchmark thread-pool)
add_executable(test_channel_pipeline test_channel_pipeline.cpp test_channel_pipeline_tools.c)
target_link_libraries(test_channel_pipeline PRIVATE UTIL SIMU PHY_COMMON LOG CONFIG_LIB shlib_loader m channel_pipeline GTest::gtest thread-pool)
if (CUDA_ENABLE)
target_compile_definitions(test_channel_pipeline PRIVATE CUDA_ENABLE)
target_compile_definitions(benchmark_channel_pipeline PRIVATE CUDA_ENABLE)
endif()
if(CUDA_ENABLE)
add_executable(test_multipath test_multipath.c)
target_link_libraries(test_multipath PRIVATE UTIL SIMU LOG CONFIG_LIB shlib_loader m oai_cuda_lib)

239
openair1/PHY/TOOLS/tests/benchmark_channel_pipeline.cpp Normal file
View File

@@ -0,0 +1,239 @@
/*
* SPDX-License-Identifier: LicenseRef-CSSL-1.0
*/
#include <cstdio>
#include <cstdlib>
#include <cmath>
#include <vector>
#include <time.h>
#include <getopt.h>
#include "oai_cuda.h"
#include "common/config/config_userapi.h"
#include <memory>
#include "benchmark/benchmark.h"
#include "test_channel_pipeline_tools.h"
#include "channel_pipeline.h"
extern "C" {
#include "openair1/SIMULATION/TOOLS/sim.h"
}
configmodule_interface_t *uniqCfg = NULL;
extern "C" void exit_function(const char *file, const char *function, const int line, const char *s, const int assert)
{
fprintf(stderr, "FATAL: %s at %s:%s:%d\n", s, file, function, line);
exit(EXIT_FAILURE);
}
#ifdef CUDA_ENABLE
static void BM_channel_convolution_gpu(benchmark::State &state)
{
int nb_rx = state.range(0);
int nb_tx = state.range(1);
int num_samples = state.range(2);
int channel_length = 16;
size_t num_input_samples = num_samples + channel_length - 1;
std::vector<c16_t *> input(nb_tx);
for (int i = 0; i < nb_tx; ++i) {
input[i] = new c16_t[num_input_samples];
}
std::vector<c16_t *> output(nb_rx);
for (int i = 0; i < nb_rx; ++i) {
output[i] = new c16_t[num_samples];
}
std::vector<cf_t *> channel(nb_rx * nb_tx);
for (int i = 0; i < nb_rx * nb_tx; ++i) {
channel[i] = new cf_t[channel_length];
}
for (int i = 0; i < nb_rx * nb_tx; ++i) {
generate_random_signal_float(channel[i], channel_length);
}
void *gpu_context = cuda_channel_pipeline_init(61440 * 4);
for (int aatx = 0; aatx < nb_tx; aatx++) {
generate_random_signal(input[aatx], num_input_samples);
}
size_t total_samples = 0;
for (auto _ : state) {
cuda_channel_pipeline(gpu_context,
(const cf_t **)channel.data(),
(const c16_t **)input.data(),
nullptr,
num_input_samples,
output.data(),
nullptr,
num_samples,
num_samples,
channel_length,
nb_tx,
nb_rx,
0.0f);
total_samples += num_samples;
}
state.counters["MSPS"] = benchmark::Counter(total_samples / 1000000.f, benchmark::Counter::kIsRate);
cuda_channel_pipeline_shutdown(gpu_context);
}
#endif
static void BM_channel_convolution_cpu(benchmark::State &state)
{
int nb_rx = state.range(0);
int nb_tx = state.range(1);
int num_samples = state.range(2);
int channel_length = 16;
size_t num_input_samples = num_samples + channel_length - 1;
std::vector<c16_t *> input(nb_tx);
for (int i = 0; i < nb_tx; ++i) {
input[i] = new c16_t[num_input_samples];
}
std::vector<c16_t *> output(nb_rx);
for (int i = 0; i < nb_rx; ++i) {
output[i] = new c16_t[num_samples];
}
std::vector<cf_t *> channel(nb_rx * nb_tx);
for (int i = 0; i < nb_rx * nb_tx; ++i) {
channel[i] = new cf_t[channel_length];
}
for (int i = 0; i < nb_rx * nb_tx; ++i) {
generate_random_signal_float(channel[i], channel_length);
}
for (int aatx = 0; aatx < nb_tx; aatx++) {
generate_random_signal(input[aatx], num_input_samples);
}
size_t total_samples = 0;
for (auto _ : state) {
channel_convolution_cpu((const cf_t **)channel.data(),
(const c16_t **)input.data(),
nullptr,
num_input_samples,
output.data(),
nullptr,
num_samples,
num_samples,
channel_length,
nb_tx,
nb_rx);
total_samples += num_samples;
}
state.counters["MSPS"] = benchmark::Counter(total_samples / 1000000.0f, benchmark::Counter::kIsRate);
for (int i = 0; i < nb_tx; ++i)
delete[] input[i];
for (int i = 0; i < nb_rx * nb_tx; ++i)
delete[] channel[i];
for (int i = 0; i < nb_rx; ++i) {
delete[] output[i];
}
}
static void BM_channel_convolution_tpool(benchmark::State &state)
{
int nb_rx = state.range(0);
int nb_tx = state.range(1);
int num_samples = state.range(2);
int channel_length = 16;
size_t num_input_samples = num_samples + channel_length - 1;
std::vector<c16_t *> input(nb_tx);
for (int i = 0; i < nb_tx; ++i) {
input[i] = new c16_t[num_input_samples];
}
std::vector<c16_t *> output(nb_rx);
for (int i = 0; i < nb_rx; ++i) {
output[i] = new c16_t[num_samples];
}
std::vector<cf_t *> channel(nb_rx * nb_tx);
for (int i = 0; i < nb_rx * nb_tx; ++i) {
channel[i] = new cf_t[channel_length];
}
for (int i = 0; i < nb_rx * nb_tx; ++i) {
generate_random_signal_float(channel[i], channel_length);
}
for (int aatx = 0; aatx < nb_tx; aatx++) {
generate_random_signal(input[aatx], num_input_samples);
}
void *tpool = init_tpool(16);
channel_pipeline_init(0.0f);
size_t total_samples = 0;
for (auto _ : state) {
channel_pipeline(tpool,
(const cf_t **)channel.data(),
(const c16_t **)input.data(),
nullptr,
num_input_samples,
output.data(),
nullptr,
num_samples,
num_samples,
channel_length,
nb_tx,
nb_rx,
0.0f);
total_samples += num_samples;
}
state.counters["MSPS"] = benchmark::Counter(total_samples / 1000000.f, benchmark::Counter::kIsRate);
channel_pipeline_shutdown();
destroy_tpool(tpool);
for (int i = 0; i < nb_tx; ++i)
delete[] input[i];
for (int i = 0; i < nb_rx * nb_tx; ++i)
delete[] channel[i];
for (int i = 0; i < nb_rx; ++i) {
delete[] output[i];
}
}
#ifdef CUDA_ENABLE
BENCHMARK(BM_channel_convolution_gpu)
->ArgsProduct({
{1, 2, 4, 16, 64}, // nb_rx
{1, 2, 4, 16, 64}, // nb_tx
{61440}, // num_samples
})
->Iterations(100);
#endif
BENCHMARK(BM_channel_convolution_cpu)
->ArgsProduct({
{1, 2, 4}, // nb_rx
{1, 2, 4}, // nb_tx
{61440}, // num_samples
})
->Iterations(50);
BENCHMARK(BM_channel_convolution_tpool)
->ArgsProduct({
{1, 2, 4, 8, 16}, // nb_rx
{1, 2, 4, 8, 16}, // nb_tx
{61440}, // num_samples
})
->Iterations(50);
int main(int argc, char **argv)
{
logInit();
randominit();
benchmark::Initialize(&argc, argv);
benchmark::RunSpecifiedBenchmarks();
return 0;
}

277
openair1/PHY/TOOLS/tests/test_channel_pipeline.cpp Normal file
View File

@@ -0,0 +1,277 @@
/*
* SPDX-License-Identifier: LicenseRef-CSSL-1.0
*/
#include <gtest/gtest.h>
#include <vector>
#include <tuple>
#include <cmath>
#include <cstdlib>
#include "oai_cuda.h"
#include "test_channel_pipeline_tools.h"
#include "channel_pipeline.h"
extern "C" {
#include "openair1/SIMULATION/TOOLS/sim.h"
}
configmodule_interface_t *uniqCfg = NULL;
extern "C" void exit_function(const char *file, const char *function, const int line, const char *s, const int assert)
{
fprintf(stderr, "FATAL: %s at %s:%s:%d\n", s, file, function, line);
exit(EXIT_FAILURE);
}
class ChannelConvolutionTest : public ::testing::TestWithParam<std::tuple<int, int, int>> {
protected:
void SetUp() override
{
#ifdef CUDA_ENABLE
gpu_context = cuda_channel_pipeline_init(614400 * 4);
#endif
tpool = init_tpool(8);
channel_pipeline_init(0.0f);
}
void TearDown() override
{
#ifdef CUDA_ENABLE
cuda_channel_pipeline_shutdown(gpu_context);
#endif
destroy_tpool(tpool);
channel_pipeline_shutdown();
}
void *gpu_context = nullptr;
void *tpool = nullptr;
};
#ifdef CUDA_ENABLE
TEST_P(ChannelConvolutionTest, CompareCpuGpu)
{
int nb_rx = std::get<0>(GetParam());
int nb_tx = std::get<1>(GetParam());
int num_samples = std::get<2>(GetParam());
int channel_length = 16;
// The input buffer must be padded at the beginning to handle the convolution history.
size_t num_input_samples = num_samples + channel_length - 1;
std::vector<c16_t *> input(nb_tx);
for (int i = 0; i < nb_tx; ++i) {
input[i] = new c16_t[num_input_samples];
generate_random_signal(input[i], num_input_samples);
}
std::vector<cf_t *> channel(nb_rx * nb_tx);
for (int i = 0; i < nb_rx * nb_tx; ++i) {
channel[i] = new cf_t[channel_length];
generate_random_signal_float(channel[i], channel_length);
}
std::vector<c16_t *> output_cpu(nb_rx);
std::vector<c16_t *> output_gpu(nb_rx);
for (int i = 0; i < nb_rx; ++i) {
output_cpu[i] = new c16_t[num_samples];
output_gpu[i] = new c16_t[num_samples];
memset(output_cpu[i], 0, num_samples * sizeof(c16_t));
memset(output_gpu[i], 0, num_samples * sizeof(c16_t));
}
// Run CPU implementation
channel_convolution_cpu((const cf_t **)channel.data(),
(const c16_t **)input.data(),
nullptr,
num_input_samples,
output_cpu.data(),
nullptr,
num_samples,
num_samples,
channel_length,
nb_tx,
nb_rx);
// Run GPU implementation
cuda_channel_pipeline(gpu_context,
(const cf_t **)channel.data(),
(const c16_t **)input.data(),
nullptr,
num_input_samples,
output_gpu.data(),
nullptr,
num_samples,
num_samples,
channel_length,
nb_tx,
nb_rx,
0.0f);
// Compare results
for (int r = 0; r < nb_rx; ++r) {
for (int i = 0; i < num_samples; ++i) {
EXPECT_LE(std::abs(output_cpu[r][i].r - output_gpu[r][i].r), 1) << "Real part mismatch at rx=" << r << " sample=" << i;
EXPECT_LE(std::abs(output_cpu[r][i].i - output_gpu[r][i].i), 1) << "Imag part mismatch at rx=" << r << " sample=" << i;
}
}
// Cleanup
for (int i = 0; i < nb_tx; ++i)
delete[] input[i];
for (int i = 0; i < nb_rx * nb_tx; ++i)
delete[] channel[i];
for (int i = 0; i < nb_rx; ++i) {
delete[] output_cpu[i];
delete[] output_gpu[i];
}
}
#endif
TEST_P(ChannelConvolutionTest, CompareCpuTpool)
{
int nb_rx = std::get<0>(GetParam());
int nb_tx = std::get<1>(GetParam());
int num_samples = std::get<2>(GetParam());
int channel_length = 16;
// The input buffer must be padded at the beginning to handle the convolution history.
size_t num_input_samples = num_samples + channel_length - 1;
size_t num_input_samples_input0 = num_input_samples;
size_t num_input_samples_input1 = 0;
if (num_samples > 5000) {
num_input_samples_input1 = 5000;
num_input_samples_input0 = num_input_samples - num_input_samples_input1;
}
std::vector<c16_t *> input(nb_tx);
for (int i = 0; i < nb_tx; ++i) {
input[i] = new c16_t[num_input_samples_input0];
generate_random_signal(input[i], num_input_samples_input0);
}
std::vector<c16_t *> input2(nb_tx);
if (num_input_samples_input1 > 0) {
for (int i = 0; i < nb_tx; i++) {
input2[i] = new c16_t[num_input_samples_input1];
generate_random_signal(input2[i], num_input_samples_input1);
}
}
std::vector<cf_t *> channel(nb_rx * nb_tx);
for (int i = 0; i < nb_rx * nb_tx; ++i) {
channel[i] = new cf_t[channel_length];
generate_random_signal_float(channel[i], channel_length);
}
std::vector<c16_t *> output_cpu(nb_rx);
std::vector<c16_t *> output_cpu2(nb_rx);
std::vector<c16_t *> output_tpool(nb_rx);
std::vector<c16_t *> output_tpool2(nb_rx);
size_t num_output_samples_output_0 = num_samples;
size_t num_output_samples_output_1 = 0;
if (num_samples > 10000) {
num_output_samples_output_1 = 10000;
num_output_samples_output_0 = num_samples - num_output_samples_output_1;
}
for (int i = 0; i < nb_rx; ++i) {
output_cpu[i] = new c16_t[num_output_samples_output_0];
output_tpool[i] = new c16_t[num_output_samples_output_0];
memset(output_cpu[i], 0, num_output_samples_output_0 * sizeof(c16_t));
memset(output_tpool[i], 0, num_output_samples_output_0 * sizeof(c16_t));
}
if (num_output_samples_output_1 > 0) {
for (int i = 0; i < nb_rx; ++i) {
output_cpu2[i] = new c16_t[num_output_samples_output_1];
output_tpool2[i] = new c16_t[num_output_samples_output_1];
memset(output_cpu2[i], 0, num_output_samples_output_1 * sizeof(c16_t));
memset(output_tpool2[i], 0, num_output_samples_output_1 * sizeof(c16_t));
}
}
// Run CPU implementation
channel_convolution_cpu((const cf_t **)channel.data(),
(const c16_t **)input.data(),
num_input_samples_input1 > 0 ? (const c16_t **)input2.data() : nullptr,
num_input_samples_input0,
output_cpu.data(),
num_output_samples_output_1 > 0 ? output_cpu2.data() : nullptr,
num_output_samples_output_0,
num_samples,
channel_length,
nb_tx,
nb_rx);
// Run tpool implementation
channel_pipeline(tpool,
(const cf_t **)channel.data(),
(const c16_t **)input.data(),
num_input_samples_input1 > 0 ? (const c16_t **)input2.data() : nullptr,
num_input_samples_input0,
output_tpool.data(),
num_output_samples_output_1 > 0 ? output_tpool2.data() : nullptr,
num_output_samples_output_0,
num_samples,
channel_length,
nb_tx,
nb_rx,
0.0f);
// Compare results
for (int r = 0; r < nb_rx; ++r) {
for (uint i = 0; i < num_output_samples_output_0; ++i) {
EXPECT_LE(std::abs(output_cpu[r][i].r - output_tpool[r][i].r), 1) << "Real part mismatch at rx=" << r << " sample=" << i;
EXPECT_LE(std::abs(output_cpu[r][i].i - output_tpool[r][i].i), 1) << "Imag part mismatch at rx=" << r << " sample=" << i;
}
}
if (num_output_samples_output_1 > 0) {
for (int r = 0; r < nb_rx; ++r) {
for (uint i = 0; i < num_output_samples_output_1; ++i) {
EXPECT_LE(std::abs(output_cpu2[r][i].r - output_tpool2[r][i].r), 1) << "Real part mismatch at rx=" << r << " sample=" << i;
EXPECT_LE(std::abs(output_cpu2[r][i].i - output_tpool2[r][i].i), 1) << "Imag part mismatch at rx=" << r << " sample=" << i;
}
}
}
// Cleanup
for (int i = 0; i < nb_tx; ++i)
delete[] input[i];
if (num_input_samples_input1 > 0)
for (int i = 0; i < nb_tx; i++)
delete[] input2[i];
for (int i = 0; i < nb_rx * nb_tx; ++i)
delete[] channel[i];
for (int i = 0; i < nb_rx; ++i) {
delete[] output_cpu[i];
delete[] output_tpool[i];
}
if (num_output_samples_output_1 > 0) {
for (int i = 0; i < nb_rx; ++i) {
delete[] output_cpu2[i];
delete[] output_tpool2[i];
}
}
}
INSTANTIATE_TEST_SUITE_P(ChannelConvolutionTests,
ChannelConvolutionTest,
::testing::Combine(::testing::Values(1, 2, 4),
::testing::Values(1, 2, 4),
::testing::Values(100, 1024, 6000, 614400)),
[](const ::testing::TestParamInfo<ChannelConvolutionTest::ParamType> &info) {
int rx = std::get<0>(info.param);
int tx = std::get<1>(info.param);
int samples = std::get<2>(info.param);
std::ostringstream name;
name << "Rx" << rx << "_Tx" << tx << "_Samples" << samples;
return name.str();
});
int main(int argc, char **argv)
{
logInit();
randominit();
testing::InitGoogleTest(&argc, argv);
return RUN_ALL_TESTS();
}

84
openair1/PHY/TOOLS/tests/test_channel_pipeline_tools.c Normal file
View File

@@ -0,0 +1,84 @@
/*
* SPDX-License-Identifier: LicenseRef-CSSL-1.0
*/
#include <stdlib.h>
#include "test_channel_pipeline_tools.h"
#include "channel_pipeline.h"
#include "thread-pool.h"
void *init_tpool(int num_threads)
{
tpool_t *tpool = malloc(sizeof(*tpool));
initFloatingCoresTpool(num_threads, tpool, false, NULL);
return tpool;
}
void destroy_tpool(void *tpool)
{
abortTpool((tpool_t *)tpool);
free(tpool);
}
void channel_convolution_cpu(const cf_t **channel,
const c16_t **tx_sig0,
const c16_t **tx_sig1,
int num_samples_tx_sig0,
c16_t **rx_sig0,
c16_t **rx_sig1,
int num_samples_rx_sig0,
int num_samples,
int channel_length,
int nb_tx,
int nb_rx)
{
for (int rx_ant = 0; rx_ant < nb_rx; rx_ant++) {
for (int i = 0; i < num_samples; i++) {
float rx_r = 0.0f;
float rx_i = 0.0f;
for (int tx_ant = 0; tx_ant < nb_tx; tx_ant++) {
for (int l = 0; l < channel_length; l++) {
int idx = i + (channel_length - 1) - l;
c16_t in_sample;
if (idx < num_samples_tx_sig0) {
in_sample = tx_sig0[tx_ant][idx];
} else {
in_sample = tx_sig1[tx_ant][idx - num_samples_tx_sig0];
}
cf_t tx_sample = {(float)in_sample.r, (float)in_sample.i};
int chan_idx = tx_ant + nb_tx * rx_ant;
cf_t ch = channel[chan_idx][l];
rx_r += tx_sample.r * ch.r - tx_sample.i * ch.i;
rx_i += tx_sample.r * ch.i + tx_sample.i * ch.r;
}
}
if (i < num_samples_rx_sig0) {
rx_sig0[rx_ant][i].r = (int16_t)rx_r;
rx_sig0[rx_ant][i].i = (int16_t)rx_i;
} else {
rx_sig1[rx_ant][i - num_samples_rx_sig0].r = (int16_t)rx_r;
rx_sig1[rx_ant][i - num_samples_rx_sig0].i = (int16_t)rx_i;
}
}
}
}
void generate_random_signal(c16_t *sig, int num_samples)
{
for (int i = 0; i < num_samples; i++) {
sig[i].r = (rand() % 2000) - 1000;
sig[i].i = (rand() % 2000) - 1000;
}
}
void generate_random_signal_float(cf_t *sig, int num_samples)
{
for (int i = 0; i < num_samples; i++) {
sig[i].r = (rand() % 2000) - 1000;
sig[i].i = (rand() % 2000) - 1000;
}
}

31
openair1/PHY/TOOLS/tests/test_channel_pipeline_tools.h Normal file
View File

@@ -0,0 +1,31 @@
/*
* SPDX-License-Identifier: LicenseRef-CSSL-1.0
*/
#pragma once
#include "common/platform_types.h"
#ifdef __cplusplus
extern "C" {
#endif
void channel_convolution_cpu(const cf_t **channel,
const c16_t **tx_sig0,
const c16_t **tx_sig1,
int num_samples_tx_sig0,
c16_t **rx_sig0,
c16_t **rx_sig1,
int num_samples_rx_sig0,
int num_samples,
int channel_length,
int nb_tx,
int nb_rx);
void generate_random_signal(c16_t *sig, int num_samples);
void generate_random_signal_float(cf_t *sig, int num_samples);
void *init_tpool(int num_threads);
void destroy_tpool(void *tpool);
#ifdef __cplusplus
}
#endif

12
openair1/SIMULATION/TOOLS/CMakeLists.txt
View File

@@ -1,8 +1,17 @@
# SPDX-License-Identifier: LicenseRef-CSSL-1.0
add_library(noise_device noise_device.c)
target_link_libraries(noise_device PRIVATE log_headers)
target_link_libraries(noise_device PRIVATE log_headers SIMU)
target_include_directories(noise_device PUBLIC ./)
set(CHANNEL_PIPELINE_SOURCES channel_pipeline.c)
if (CUDA_ENABLE)
list(APPEND CHANNEL_PIPELINE_SOURCES channel_pipeline_v2.cu)
endif()
add_library(channel_pipeline ${CHANNEL_PIPELINE_SOURCES})
target_link_libraries(channel_pipeline PUBLIC thread-pool noise_device)
if (CUDA_ENABLE)
target_link_libraries(channel_pipeline PUBLIC CUDA::toolkit)
endif()
if(CUDA_ENABLE)
add_library(oai_cuda_lib STATIC
@@ -16,4 +25,5 @@ if(CUDA_ENABLE)
target_include_directories(oai_cuda_lib PUBLIC
${CMAKE_CURRENT_SOURCE_DIR}
)
target_link_libraries(channel_pipeline PUBLIC oai_cuda_lib)
endif()

185
openair1/SIMULATION/TOOLS/channel_pipeline.c Normal file
View File

@@ -0,0 +1,185 @@
/*
* SPDX-License-Identifier: LicenseRef-CSSL-1.0
*/
#include <math.h>
#include "utils.h"
#include "channel_pipeline.h"
#include "task_ans.h"
#include "noise_device.h"
#include "thread-pool.h"
typedef struct {
const c16_t **tx_sig0;
const c16_t **tx_sig1;
int num_samples_tx_sig0;
int nb_tx;
int num_samples_rx_sig0;
int num_samples;
int channel_length;
int num_jobs;
task_ans_t *task_ans;
float noise_power;
} job_common_args_t;
typedef struct {
job_common_args_t *common;
const cf_t **channel;
c16_t *rx_sig0;
c16_t *rx_sig1;
int job_index;
} job_args_t;
void do_convolution_and_noise(int nb_tx,
int channel_length,
const cf_t **channel,
const c16_t **tx_sig0,
const c16_t **tx_sig1,
int num_samples_tx_sig0,
c16_t *rx_sig0,
c16_t *rx_sig1,
int num_samples_rx_sig0,
int num_samples,
int job_index,
int num_jobs,
float noise_power)
{
const int batch_size = 4000;
cf_t work_buffer[batch_size] __attribute__((aligned(64)));
for (int batch_start = job_index * batch_size; batch_start < num_samples; batch_start += batch_size * num_jobs) {
int batch_end = min(batch_start + batch_size, num_samples);
if (noise_power > 0.0f) {
get_noise_vector((float *)work_buffer, batch_size * 2);
}
for (int i = batch_start; i < batch_end; i++) {
float rx_r = 0.0f;
float rx_i = 0.0f;
for (int tx_ant = 0; tx_ant < nb_tx; tx_ant++) {
for (int l = 0; l < channel_length; l++) {
int idx = i + (channel_length - 1) - l;
c16_t in_sample;
if (idx < num_samples_tx_sig0) {
in_sample = tx_sig0[tx_ant][idx];
} else {
in_sample = tx_sig1[tx_ant][idx - num_samples_tx_sig0];
}
cf_t tx_sample = {(float)in_sample.r, (float)in_sample.i};
cf_t ch = channel[tx_ant][l];
rx_r += tx_sample.r * ch.r - tx_sample.i * ch.i;
rx_i += tx_sample.r * ch.i + tx_sample.i * ch.r;
}
}
if (noise_power > 0.0f) {
work_buffer[i - batch_start].r += rx_r;
work_buffer[i - batch_start].i += rx_i;
} else {
work_buffer[i - batch_start].r = rx_r;
work_buffer[i - batch_start].i = rx_i;
}
}
for (int i = batch_start; i < batch_end; i++) {
if (i < num_samples_rx_sig0) {
rx_sig0[i].r = work_buffer[i - batch_start].r;
rx_sig0[i].i = work_buffer[i - batch_start].i;
} else {
rx_sig1[i - num_samples_rx_sig0].r = work_buffer[i - batch_start].r;
rx_sig1[i - num_samples_rx_sig0].i = work_buffer[i - batch_start].i;
}
}
}
}
void channel_job(void *args)
{
job_args_t *job_args = (job_args_t *)args;
job_common_args_t *common = job_args->common;
do_convolution_and_noise(common->nb_tx,
common->channel_length,
job_args->channel,
common->tx_sig0,
common->tx_sig1,
common->num_samples_tx_sig0,
job_args->rx_sig0,
job_args->rx_sig1,
common->num_samples_rx_sig0,
common->num_samples,
job_args->job_index,
common->num_jobs,
common->noise_power);
completed_task_ans(common->task_ans);
}
void channel_pipeline(void *tpool,
const cf_t **channel,
const c16_t **tx_sig0,
const c16_t **tx_sig1,
int num_samples_tx_sig0,
c16_t **rx_sig0,
c16_t **rx_sig1,
int num_samples_rx_sig0,
int num_samples,
int channel_length,
int nb_tx,
int nb_rx,
float noise_power)
{
AssertFatal(nb_tx > 0, "nb_tx must be positive (%d)\n", nb_tx);
AssertFatal(nb_rx > 0, "nb_rx must be positive (%d)\n", nb_rx);
AssertFatal(nb_tx <= 64, "Number of TX antennas is too large (%d)\n", nb_tx);
AssertFatal(nb_rx <= 64, "Number of RX antennas is too large (%d)\n", nb_rx);
tpool_t *thread_pool = (tpool_t *)tpool;
size_t num_tpool_threads = thread_pool->len_thr;
job_common_args_t common_args;
common_args.tx_sig0 = tx_sig0;
if (tx_sig1) {
common_args.tx_sig1 = tx_sig1;
} else {
common_args.tx_sig1 = NULL;
}
// At least 1 job per RX antenna. Attempt to saturate the threadpool
size_t num_jobs_per_rx_antenna = (num_tpool_threads + nb_rx - 1) / nb_rx;
common_args.num_samples_tx_sig0 = num_samples_tx_sig0;
common_args.num_samples_rx_sig0 = num_samples_rx_sig0;
common_args.num_samples = num_samples;
common_args.channel_length = channel_length;
common_args.nb_tx = nb_tx;
common_args.num_jobs = num_jobs_per_rx_antenna;
common_args.noise_power = noise_power;
task_ans_t task_ans;
init_task_ans(&task_ans, nb_rx * num_jobs_per_rx_antenna);
common_args.task_ans = &task_ans;
job_args_t job_args_array[nb_rx][num_jobs_per_rx_antenna];
for (int aarx = 0; aarx < nb_rx; aarx++) {
for (int job_idx = 0; job_idx < num_jobs_per_rx_antenna; job_idx++) {
job_args_t *job_args = &job_args_array[aarx][job_idx];
job_args->common = &common_args;
job_args->channel = &channel[nb_tx * aarx];
job_args->rx_sig0 = rx_sig0[aarx];
if (rx_sig1) {
job_args->rx_sig1 = rx_sig1[aarx];
} else {
job_args->rx_sig1 = NULL;
}
job_args->job_index = job_idx;
task_t task;
task.args = job_args;
task.func = channel_job;
pushTpool(tpool, task);
}
}
join_task_ans(&task_ans);
}
void channel_pipeline_init(float noise_power) {
init_noise_device(noise_power);
}
void channel_pipeline_shutdown(void) {
free_noise_device();
}

49
openair1/SIMULATION/TOOLS/channel_pipeline.h Normal file
View File

@@ -0,0 +1,49 @@
/*
* SPDX-License-Identifier: LicenseRef-CSSL-1.0
*/
#ifndef _CHANNEL_CONVOLUTION_H_
#define _CHANNEL_CONVOLUTION_H_
#ifdef __cplusplus
extern "C" {
#endif
#include "common/platform_types.h"
#ifdef CUDA_ENABLE
void *cuda_channel_pipeline_init(int max_samples);
void cuda_channel_pipeline_shutdown(void *context_handle);
void cuda_channel_pipeline(void *context_handle,
const cf_t **channel,
const c16_t **tx_sig0,
const c16_t **tx_sig1,
int num_samples_tx_sig0,
c16_t **rx_sig0,
c16_t **rx_sig1,
int num_samples_rx_sig0,
int num_samples,
int channel_length,
int nb_tx,
int nb_rx,
float noise_power);
#endif
void channel_pipeline_init(float noise_power);
void channel_pipeline_shutdown(void);
void channel_pipeline(void *tpool,
const cf_t **channel,
const c16_t **tx_sig0,
const c16_t **tx_sig1,
int num_samples_tx_sig0,
c16_t **rx_sig0,
c16_t **rx_sig1,
int num_samples_rx_sig0,
int num_samples,
int channel_length,
int nb_tx,
int nb_rx,
float noise_power);
#ifdef __cplusplus
}
#endif
#endif

210
openair1/SIMULATION/TOOLS/channel_pipeline_v2.cu Normal file
View File

@@ -0,0 +1,210 @@
/*
* SPDX-License-Identifier: LicenseRef-CSSL-1.0
*/
#include "oai_cuda.h"
#include <cstdint>
#include <cstdio>
#include <cuda_runtime.h>
#include "common/platform_types.h"
#include "common/utils/assertions.h"
#define CHECK_CUDA(val) \
{ \
if (val != cudaSuccess) { \
fprintf(stderr, "CUDA Error at %s:%d: %s\n", __FILE__, __LINE__, cudaGetErrorString(val)); \
exit(EXIT_FAILURE); \
} \
}
__device__ __forceinline__ cf_t complex_mul(cf_t a, cf_t b)
{
return cf_t{a.r * b.r - a.i * b.i, a.r * b.i + a.i * b.r};
}
__device__ __forceinline__ cf_t complex_add(cf_t a, cf_t b)
{
return cf_t{a.r + b.r, a.i + b.i};
}
/**
* @brief CUDA kernel for performing multipath channel convolution and noise generation
*
* This kernel simulates the effect of a multipath channel by convolving the transmitted
* signal with the channel impulse response. It supports multiple transmit and receive
* antennas (MIMO). The input signal can be split across two buffers (tx_sig0 and tx_sig1),
* and the output is similarly written to two buffers (rx_sig0 and rx_sig1).
*
* @param channel Pointer to the channel impulse response coefficients for each link.
* @param tx_sig0 Pointer to the first part of the transmitted signal buffers.
* @param tx_sig1 Pointer to the second part of the transmitted signal buffers (optional).
* @param num_samples_tx_sig0 Number of samples in the first transmit buffer NOTE: The total input
samples provided to the kernel should be equal to num_samples + channel_length - 1. They
can be arbitrarily split between tx_sig0 and tx_sig1 via this parameter
* @param rx_sig0 Pointer to the first part of the received signal buffers.
* @param rx_sig1 Pointer to the second part of the received signal buffers (optional).
* @param num_samples_rx_sig0 Number of samples in the first receive buffer. NOTE: The total output
samples privided to the kernel should be equal to num_samples. They can be arbitrarily split
between rx_sig0 and rx_sig1 via this parameter
* @param num_samples Total number of output samples to compute.
* @param channel_length Number of taps in the multipath channel.
* @param nb_tx Number of transmit antennas.
* @param nb_rx Number of receive antennas.
*/
__global__ void channel_convolution_and_noise(const cf_t **__restrict__ channel,
const c16_t **__restrict__ tx_sig0,
const c16_t **__restrict__ tx_sig1,
int num_samples_tx_sig0,
c16_t **__restrict__ rx_sig0,
c16_t **__restrict__ rx_sig1,
int num_samples_rx_sig0,
int num_samples,
int channel_length,
int nb_tx,
int nb_rx,
float noise_power,
curandState_t *curand_states)
{
extern __shared__ cf_t tx_sig[];
const int i = blockIdx.x * blockDim.x + threadIdx.x;
const int rx_ant_idx = blockIdx.y;
const int padding_len = channel_length - 1;
cf_t rx_tmp = cf_t{0.0f, 0.0f};
for (int aatx = 0; aatx < nb_tx; aatx++) {
const int tid = threadIdx.x;
const int block_start_idx = blockIdx.x * blockDim.x;
const int shared_mem_size = blockDim.x + padding_len;
for (int k = tid; k < shared_mem_size; k += blockDim.x) {
int load_idx = block_start_idx + k;
if (load_idx < num_samples + padding_len) {
c16_t val;
if (load_idx < num_samples_tx_sig0) {
val = tx_sig0[aatx][load_idx];
} else {
val = tx_sig1[aatx][load_idx - num_samples_tx_sig0];
}
tx_sig[k] = cf_t{(float)val.r, (float)val.i};
} else {
tx_sig[k] = cf_t{0.0f, 0.0f};
}
}
__syncthreads();
if (i < num_samples) {
for (int l = 0; l < channel_length; l++) {
cf_t tx_sample = tx_sig[tid + (channel_length - 1) - l];
int chan_link_idx = aatx + (rx_ant_idx * nb_tx);
cf_t chan_weight = channel[chan_link_idx][l];
rx_tmp = complex_add(rx_tmp, complex_mul(tx_sample, chan_weight));
}
}
__syncthreads();
}
if (i < num_samples) {
if (noise_power > 0.0f) {
curandState_t local_state = curand_states[rx_ant_idx * num_samples + i];
float2 awgn = curand_normal2(&local_state);
if (i < num_samples_rx_sig0) {
rx_sig0[rx_ant_idx][i].r = rx_tmp.r + awgn.x * noise_power;
rx_sig0[rx_ant_idx][i].i = rx_tmp.i + awgn.y * noise_power;
} else {
rx_sig1[rx_ant_idx][i - num_samples_rx_sig0].r = rx_tmp.r + awgn.x * noise_power;
rx_sig1[rx_ant_idx][i - num_samples_rx_sig0].i = rx_tmp.i + awgn.y * noise_power;
}
curand_states[rx_ant_idx * num_samples + i] = local_state;
} else {
if (i < num_samples_rx_sig0) {
rx_sig0[rx_ant_idx][i].r = rx_tmp.r;
rx_sig0[rx_ant_idx][i].i = rx_tmp.i;
} else {
rx_sig1[rx_ant_idx][i - num_samples_rx_sig0].r = rx_tmp.r;
rx_sig1[rx_ant_idx][i - num_samples_rx_sig0].i = rx_tmp.i;
}
}
}
}
struct GpuContext {
cudaStream_t stream;
curandState_t *curand_states;
size_t curand_states_size;
};
extern "C" void *cuda_channel_pipeline_init(int max_samples)
{
GpuContext *ctx = new GpuContext();
int dev = 0;
struct cudaDeviceProp prop;
CHECK_CUDA(cudaGetDeviceProperties(&prop, dev));
int pageable;
int integrated;
cudaDeviceGetAttribute(&pageable, cudaDevAttrPageableMemoryAccess, dev);
cudaDeviceGetAttribute(&integrated, cudaDevAttrIntegrated,dev);
if (!(pageable && integrated)) {
return NULL;
}
CHECK_CUDA(cudaStreamCreate(&ctx->stream));
ctx->curand_states = (curandState_t *)create_and_init_curand_states_cuda(max_samples, time(NULL));
ctx->curand_states_size = max_samples;
return (void *)ctx;
}
extern "C" void cuda_channel_pipeline_shutdown(void *context_handle)
{
if (context_handle == nullptr) {
return;
}
GpuContext *ctx = (GpuContext *)context_handle;
CHECK_CUDA(cudaStreamDestroy(ctx->stream));
destroy_curand_states_cuda((void *)ctx->curand_states);
delete ctx;
}
extern "C" void cuda_channel_pipeline(void *context_handle,
const cf_t **channel,
const c16_t **tx_sig0,
const c16_t **tx_sig1,
int num_samples_tx_sig0,
c16_t **rx_sig0,
c16_t **rx_sig1,
int num_samples_rx_sig0,
int num_samples,
int channel_length,
int nb_tx,
int nb_rx,
float noise_power)
{
AssertFatal(context_handle, "No context handle provided\n");
AssertFatal(channel, "No channel provided\n");
AssertFatal(tx_sig0, "No tx_sig0 provided\n");
AssertFatal(num_samples_tx_sig0 == (num_samples + channel_length - 1) || tx_sig1, "No tx_sig1 provided\n");
AssertFatal(rx_sig0, "No rx_sig0 provided\n");
AssertFatal(num_samples_rx_sig0 == num_samples || rx_sig1, "No rx_sig1 provided\n");
GpuContext *ctx = (GpuContext *)context_handle;
if (num_samples * nb_rx > ctx->curand_states_size) {
destroy_curand_states_cuda((void *)ctx->curand_states);
ctx->curand_states_size = num_samples * nb_rx;
ctx->curand_states = (curandState_t *)create_and_init_curand_states_cuda(ctx->curand_states_size, time(NULL));
}
dim3 threadsPerBlock(512, 1);
dim3 numBlocks((num_samples + threadsPerBlock.x - 1) / threadsPerBlock.x, nb_rx);
size_t sharedMemSize = (threadsPerBlock.x + channel_length - 1) * sizeof(cf_t);
channel_convolution_and_noise<<<numBlocks, threadsPerBlock, sharedMemSize, ctx->stream>>>(channel,
tx_sig0,
tx_sig1,
num_samples_tx_sig0,
rx_sig0,
rx_sig1,
num_samples_rx_sig0,
num_samples,
channel_length,
nb_tx,
nb_rx,
noise_power,
ctx->curand_states);
CHECK_CUDA(cudaStreamSynchronize(ctx->stream));
}

1
openair1/SIMULATION/TOOLS/multipath_channel.cu
View File

@@ -5,6 +5,7 @@
#include <stdio.h>
#include <cuda_runtime.h>
#include "oai_cuda.h"
#include "common/platform_types.h"
#define CHECK_CUDA(val) checkCuda((val), #val, __FILE__, __LINE__)
static void checkCuda(cudaError_t result, const char *const func, const char *const file, const int line)

5
openair1/SIMULATION/TOOLS/oai_cuda.h
View File

@@ -8,10 +8,7 @@
#include <stdint.h>
#ifdef __NVCC__
typedef struct complex16 {
int16_t r;
int16_t i;
} c16_t;
#include "common/platform_types.h"
#else
#include "PHY/TOOLS/tools_defs.h"
#endif