diff --git a/openair1/PHY/TOOLS/tests/CMakeLists.txt b/openair1/PHY/TOOLS/tests/CMakeLists.txt index 3326ef0c8d..91b1f10023 100644 --- a/openair1/PHY/TOOLS/tests/CMakeLists.txt +++ b/openair1/PHY/TOOLS/tests/CMakeLists.txt @@ -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) diff --git a/openair1/PHY/TOOLS/tests/benchmark_channel_pipeline.cpp b/openair1/PHY/TOOLS/tests/benchmark_channel_pipeline.cpp new file mode 100644 index 0000000000..681955aaa6 --- /dev/null +++ b/openair1/PHY/TOOLS/tests/benchmark_channel_pipeline.cpp @@ -0,0 +1,239 @@ +/* + * SPDX-License-Identifier: LicenseRef-CSSL-1.0 + */ + +#include +#include +#include +#include +#include +#include +#include "oai_cuda.h" +#include "common/config/config_userapi.h" +#include +#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 input(nb_tx); + for (int i = 0; i < nb_tx; ++i) { + input[i] = new c16_t[num_input_samples]; + } + + std::vector output(nb_rx); + for (int i = 0; i < nb_rx; ++i) { + output[i] = new c16_t[num_samples]; + } + + std::vector 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 input(nb_tx); + for (int i = 0; i < nb_tx; ++i) { + input[i] = new c16_t[num_input_samples]; + } + + std::vector output(nb_rx); + for (int i = 0; i < nb_rx; ++i) { + output[i] = new c16_t[num_samples]; + } + + std::vector 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 input(nb_tx); + for (int i = 0; i < nb_tx; ++i) { + input[i] = new c16_t[num_input_samples]; + } + + std::vector output(nb_rx); + for (int i = 0; i < nb_rx; ++i) { + output[i] = new c16_t[num_samples]; + } + + std::vector 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; +} diff --git a/openair1/PHY/TOOLS/tests/test_channel_pipeline.cpp b/openair1/PHY/TOOLS/tests/test_channel_pipeline.cpp new file mode 100644 index 0000000000..e4b8d2cdf9 --- /dev/null +++ b/openair1/PHY/TOOLS/tests/test_channel_pipeline.cpp @@ -0,0 +1,277 @@ +/* + * SPDX-License-Identifier: LicenseRef-CSSL-1.0 + */ + +#include +#include +#include +#include +#include +#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> { + 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 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 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 output_cpu(nb_rx); + std::vector 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 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 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 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 output_cpu(nb_rx); + std::vector output_cpu2(nb_rx); + std::vector output_tpool(nb_rx); + std::vector 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 &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(); +} diff --git a/openair1/PHY/TOOLS/tests/test_channel_pipeline_tools.c b/openair1/PHY/TOOLS/tests/test_channel_pipeline_tools.c new file mode 100644 index 0000000000..34e7af3fae --- /dev/null +++ b/openair1/PHY/TOOLS/tests/test_channel_pipeline_tools.c @@ -0,0 +1,84 @@ +/* + * SPDX-License-Identifier: LicenseRef-CSSL-1.0 + */ + +#include +#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; + } +} diff --git a/openair1/PHY/TOOLS/tests/test_channel_pipeline_tools.h b/openair1/PHY/TOOLS/tests/test_channel_pipeline_tools.h new file mode 100644 index 0000000000..9057c6582e --- /dev/null +++ b/openair1/PHY/TOOLS/tests/test_channel_pipeline_tools.h @@ -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 diff --git a/openair1/SIMULATION/TOOLS/CMakeLists.txt b/openair1/SIMULATION/TOOLS/CMakeLists.txt index ebf2430528..4ea2d2373c 100644 --- a/openair1/SIMULATION/TOOLS/CMakeLists.txt +++ b/openair1/SIMULATION/TOOLS/CMakeLists.txt @@ -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() diff --git a/openair1/SIMULATION/TOOLS/channel_pipeline.c b/openair1/SIMULATION/TOOLS/channel_pipeline.c new file mode 100644 index 0000000000..069222e2f4 --- /dev/null +++ b/openair1/SIMULATION/TOOLS/channel_pipeline.c @@ -0,0 +1,185 @@ +/* + * SPDX-License-Identifier: LicenseRef-CSSL-1.0 + */ + +#include +#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(); +} diff --git a/openair1/SIMULATION/TOOLS/channel_pipeline.h b/openair1/SIMULATION/TOOLS/channel_pipeline.h new file mode 100644 index 0000000000..fc8d9e29c0 --- /dev/null +++ b/openair1/SIMULATION/TOOLS/channel_pipeline.h @@ -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 diff --git a/openair1/SIMULATION/TOOLS/channel_pipeline_v2.cu b/openair1/SIMULATION/TOOLS/channel_pipeline_v2.cu new file mode 100644 index 0000000000..5c270e670f --- /dev/null +++ b/openair1/SIMULATION/TOOLS/channel_pipeline_v2.cu @@ -0,0 +1,210 @@ +/* + * SPDX-License-Identifier: LicenseRef-CSSL-1.0 + */ + +#include "oai_cuda.h" +#include +#include +#include +#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<<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)); +} diff --git a/openair1/SIMULATION/TOOLS/multipath_channel.cu b/openair1/SIMULATION/TOOLS/multipath_channel.cu index ec0e032b73..b8704c0da0 100644 --- a/openair1/SIMULATION/TOOLS/multipath_channel.cu +++ b/openair1/SIMULATION/TOOLS/multipath_channel.cu @@ -5,6 +5,7 @@ #include #include #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) diff --git a/openair1/SIMULATION/TOOLS/oai_cuda.h b/openair1/SIMULATION/TOOLS/oai_cuda.h index 03d2ea15b8..3dfd79a067 100644 --- a/openair1/SIMULATION/TOOLS/oai_cuda.h +++ b/openair1/SIMULATION/TOOLS/oai_cuda.h @@ -8,10 +8,7 @@ #include #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