diff --git a/task1/.gitignore b/task1/.gitignore new file mode 100644 index 0000000..9b23ed1 --- /dev/null +++ b/task1/.gitignore @@ -0,0 +1,4 @@ +bin/ +results/*.out +results/*.err +results/*.csv diff --git a/task1/README.md b/task1/README.md new file mode 100644 index 0000000..2755df1 --- /dev/null +++ b/task1/README.md @@ -0,0 +1,188 @@ +# Задание 1: CUDA-реализация LINPACK-подобного теста + +В папке лежит готовый каркас под первое задание: + +- `src/main.cu` — собственная CUDA-реализация решения плотной СЛАУ методом Якоби. +- `scripts/build.sh` — сборка программы через `nvcc`. +- `scripts/run_cuda.slurm` — пакетный запуск собственной CUDA-версии. +- `scripts/run_intel_linpack.slurm` — пакетный запуск стандартного Intel LINPACK на CPU. + +Программа генерирует строго диагонально доминирующую матрицу `A`, заранее известный вектор решения `x_true`, правую часть `b = A * x_true`, после чего решает систему методом Якоби на GPU. В выводе печатаются: + +- размер матрицы `N`; +- лучшее время решения в миллисекундах; +- число итераций; +- норма невязки `||Ax - b||_inf`; +- ошибка `||x - x_true||_inf`; +- LINPACK-like производительность в GFLOPS. + +## Что сделать на СКЦ + +### 1. Передать папку на кластер + +Если алиас `polytech` уже прописан в `~/.ssh/config`, достаточно: + +```bash +scp -r task1 polytech:~/supercomputers/ +``` + +### 2. Подключиться + +```bash +ssh polytech +cd ~/supercomputers/task1 +``` + +### 3. Запустить собственную CUDA-реализацию + +```bash +sbatch scripts/run_cuda.slurm +``` + +Сразу после отправки Slurm вернёт `job id`. Дальше: + +```bash +squeue -u tm3u21 +sacct -j --format=JobID,JobName,Partition,State,Start,End,Elapsed,NNodes,AllocTRES%40,NodeList,ExitCode +``` + +В текущей конфигурации СКЦ в `tornado-k40` GPU выбирается самим разделом, поэтому в `slurm`-скрипте не используется `--gres=gpu:1`. Если снова появится ошибка про `gres`, значит её не надо добавлять вручную. + +После завершения посмотри: + +```bash +less results/task1-cuda-.out +cat results/task1-cuda-.csv +``` + +Если на кластере нужна другая GPU-архитектура, можно пересобрать так: + +```bash +CUDA_ARCH=sm_70 ./scripts/build.sh +``` + +По умолчанию в `build.sh` стоит `sm_35`, потому что пример ориентирован на `tornado-k40`. + +### 4. Запустить стандартный Intel LINPACK + +```bash +sbatch scripts/run_intel_linpack.slurm +``` + +Проверка статуса и итогов: + +```bash +sacct -j --format=JobID,JobName,Partition,State,Start,End,Elapsed,NNodes,AllocTRES%40,NodeList,ExitCode +less results/task1-intel-linpack-.out +``` + +Если каталог с Intel LINPACK на кластере другой, отправь задание так: + +```bash +sbatch --export=ALL,LINPACK_DIR=/linux/share/mkl/benchmarks/linpack scripts/run_intel_linpack.slurm +``` + +## Что нужно собрать для отчёта + +Ниже последовательность, которая даст все обязательные материалы для отчёта и скриншотов. + +### Шаг 1. Скрин входа с логином + +На login-узле выполни: + +```bash +whoami +hostname +date +``` + +Сделай скрин терминала. На нём должен быть виден логин `tm3u21`. + +### Шаг 2. Скрин конфигурации узла и GPU + +После завершения CUDA-задачи открой: + +```bash +less results/task1-cuda-.out +``` + +В начале файла уже будут: + +- `whoami`, `hostname`, `date`; +- `scontrol show job ...`; +- `scontrol show node ...`; +- `lscpu`; +- `nvidia-smi`. + +Сделай отдельные скрины с этой информацией. + +### Шаг 3. Скрин времени выполнения и числа узлов + +Выполни: + +```bash +sacct -j , --format=JobID,JobName,Partition,State,Elapsed,NNodes,AllocTRES%40,NodeList,ExitCode +``` + +На этом скрине будут: + +- время выполнения; +- количество узлов; +- список узлов; +- тип выделенных ресурсов. + +### Шаг 4. Вынести численные результаты + +Для собственной программы значения бери из файла: + +```bash +cat results/task1-cuda-.csv +``` + +Для Intel LINPACK значения времени и GFLOPS бери из: + +```bash +less results/task1-intel-linpack-.out +``` + +Ищи секцию `Performance Summary`. + +## Какие картинки ожидает `report/report.tex` + +В отчёте уже подготовлены следующие пути: + +- `report/img/task1-login.png` +- `report/img/task1-cuda-node.png` +- `report/img/task1-cuda-run.png` +- `report/img/task1-cuda-sacct.png` +- `report/img/task1-intel-run.png` +- `report/img/task1-intel-sacct.png` + +Просто положи туда свои скриншоты с этими именами. + +## Если нужно поменять размеры задач + +Собственная программа сейчас запускается на: + +- `1000` +- `1500` +- `2000` +- `2500` +- `3000` +- `3500` + +Это задаётся параметрами в `scripts/run_cuda.slurm`: + +```bash +--start 1000 --step 500 --count 6 +``` + +Если удобнее задать точный набор размеров, используй: + +```bash +./bin/linpack_cuda --sizes 1000,2000,3000 +``` + +## Ограничения + +Этот код я здесь локально не компилировал, потому что в окружении нет гарантированно настроенного CUDA toolchain и GPU. Поэтому первый реальный прогон лучше делать сразу на СКЦ; если что-то упадёт по модулю, архитектуре GPU или пути к Intel LINPACK, пришли ошибку, и я быстро подправлю. diff --git a/task1/results/.gitkeep b/task1/results/.gitkeep new file mode 100644 index 0000000..8b13789 --- /dev/null +++ b/task1/results/.gitkeep @@ -0,0 +1 @@ + diff --git a/task1/scripts/build.sh b/task1/scripts/build.sh new file mode 100755 index 0000000..cf0d9eb --- /dev/null +++ b/task1/scripts/build.sh @@ -0,0 +1,24 @@ +#!/usr/bin/env bash +set -euo pipefail + +ROOT_DIR="$(cd "$(dirname "$0")/.." && pwd)" +cd "$ROOT_DIR" + +mkdir -p bin results + +CUDA_ARCH="${CUDA_ARCH:-sm_35}" + +module purge +module load compiler/gcc/11 +module load nvidia/cuda/11.6u2 + +nvcc \ + -ccbin g++ \ + -O3 \ + -std=c++14 \ + -lineinfo \ + -arch="${CUDA_ARCH}" \ + -o bin/linpack_cuda \ + src/main.cu + +echo "Built: $ROOT_DIR/bin/linpack_cuda" diff --git a/task1/scripts/run_cuda.slurm b/task1/scripts/run_cuda.slurm new file mode 100755 index 0000000..5f07d3a --- /dev/null +++ b/task1/scripts/run_cuda.slurm @@ -0,0 +1,54 @@ +#!/usr/bin/env bash +#SBATCH --job-name=task1-cuda +#SBATCH --partition=tornado-k40 +#SBATCH --nodes=1 +#SBATCH --ntasks=1 +#SBATCH --time=00:20:00 +#SBATCH --output=results/%x-%j.out +#SBATCH --error=results/%x-%j.err + +set -euo pipefail + +ROOT_DIR="$(cd "$(dirname "$0")/.." && pwd)" +cd "$ROOT_DIR" + +mkdir -p results bin + +./scripts/build.sh + +echo "===== account info =====" +whoami +hostname +date + +echo +echo "===== slurm info =====" +echo "SLURM_JOB_ID=${SLURM_JOB_ID:-unknown}" +echo "SLURM_JOB_NAME=${SLURM_JOB_NAME:-unknown}" +echo "SLURM_JOB_PARTITION=${SLURM_JOB_PARTITION:-unknown}" +echo "SLURM_JOB_NUM_NODES=${SLURM_JOB_NUM_NODES:-unknown}" +echo "SLURM_NODELIST=${SLURM_NODELIST:-unknown}" +echo "CUDA_VISIBLE_DEVICES=${CUDA_VISIBLE_DEVICES:-unset}" +scontrol show job "${SLURM_JOB_ID}" || true + +echo +echo "===== node config =====" +lscpu | sed -n '1,20p' +if [ -n "${SLURMD_NODENAME:-}" ]; then + scontrol show node "${SLURMD_NODENAME}" || true +fi +nvidia-smi -L || true +nvidia-smi || true + +echo +echo "===== benchmark =====" +./bin/linpack_cuda \ + --start 1000 \ + --step 500 \ + --count 6 \ + --eps 1e-6 \ + --max-iters 15000 \ + --threads 256 \ + --repeat 3 \ + --warmup 1 \ + --csv "results/task1-cuda-${SLURM_JOB_ID}.csv" diff --git a/task1/scripts/run_intel_linpack.slurm b/task1/scripts/run_intel_linpack.slurm new file mode 100755 index 0000000..18f837e --- /dev/null +++ b/task1/scripts/run_intel_linpack.slurm @@ -0,0 +1,56 @@ +#!/usr/bin/env bash +#SBATCH --job-name=task1-intel-linpack +#SBATCH --partition=tornado +#SBATCH --nodes=1 +#SBATCH --ntasks=1 +#SBATCH --cpus-per-task=56 +#SBATCH --time=00:20:00 +#SBATCH --output=results/%x-%j.out +#SBATCH --error=results/%x-%j.err + +set -euo pipefail + +ROOT_DIR="$(cd "$(dirname "$0")/.." && pwd)" +cd "$ROOT_DIR" + +mkdir -p results + +LINPACK_DIR="${LINPACK_DIR:-/linux/share/mkl/benchmarks/linpack}" +LINPACK_INPUT="${LINPACK_INPUT:-lininput_xeon64}" + +if [ ! -x "${LINPACK_DIR}/xlinpack_xeon64" ]; then + echo "Intel LINPACK binary not found: ${LINPACK_DIR}/xlinpack_xeon64" + echo "If the path differs on the cluster, submit with:" + echo "sbatch --export=ALL,LINPACK_DIR=/path/to/linpack scripts/run_intel_linpack.slurm" + exit 1 +fi + +echo "===== account info =====" +whoami +hostname +date + +echo +echo "===== slurm info =====" +echo "SLURM_JOB_ID=${SLURM_JOB_ID:-unknown}" +echo "SLURM_JOB_NAME=${SLURM_JOB_NAME:-unknown}" +echo "SLURM_JOB_PARTITION=${SLURM_JOB_PARTITION:-unknown}" +echo "SLURM_JOB_NUM_NODES=${SLURM_JOB_NUM_NODES:-unknown}" +echo "SLURM_NODELIST=${SLURM_NODELIST:-unknown}" +echo "OMP_NUM_THREADS=${SLURM_CPUS_PER_TASK:-56}" +scontrol show job "${SLURM_JOB_ID}" || true + +echo +echo "===== node config =====" +lscpu | sed -n '1,20p' +if [ -n "${SLURMD_NODENAME:-}" ]; then + scontrol show node "${SLURMD_NODENAME}" || true +fi + +echo +echo "===== intel linpack =====" +export OMP_NUM_THREADS="${SLURM_CPUS_PER_TASK:-56}" +export MKL_NUM_THREADS="${SLURM_CPUS_PER_TASK:-56}" + +cd "${LINPACK_DIR}" +./xlinpack_xeon64 "${LINPACK_INPUT}" diff --git a/task1/src/main.cu b/task1/src/main.cu new file mode 100644 index 0000000..8e5fe27 --- /dev/null +++ b/task1/src/main.cu @@ -0,0 +1,559 @@ +#include + +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +#define CUDA_CHECK(call) \ + do { \ + cudaError_t err__ = (call); \ + if (err__ != cudaSuccess) { \ + std::cerr << "CUDA error at " << __FILE__ << ":" << __LINE__ \ + << " -> " << cudaGetErrorString(err__) << std::endl; \ + std::exit(EXIT_FAILURE); \ + } \ + } while (0) + +struct Options { + int start = 1000; + int step = 500; + int count = 6; + std::vector sizes; + int threads = 256; + int max_iters = 10000; + int repeat = 3; + int warmup = 1; + unsigned int seed = 42U; + double eps = 1e-6; + std::string csv_path; +}; + +struct Metrics { + double elapsed_ms = std::numeric_limits::max(); + int iterations = 0; + double residual_inf = std::numeric_limits::infinity(); + double x_error_inf = std::numeric_limits::infinity(); + double gflops = 0.0; + bool converged = false; +}; + +__global__ void jacobi_iteration(const double *a, + const double *b, + const double *x_in, + double *x_out, + int n, + double eps, + int *converged) { + const int row = blockIdx.x * blockDim.x + threadIdx.x; + if (row >= n) { + return; + } + + const int row_offset = row * n; + double sum = 0.0; + for (int col = 0; col < n; ++col) { + if (col != row) { + sum += a[row_offset + col] * x_in[col]; + } + } + + const double next = (b[row] - sum) / a[row_offset + row]; + x_out[row] = next; + if (fabs(next - x_in[row]) > eps) { + atomicAnd(converged, 0); + } +} + +static void print_usage(const char *program) { + std::cout + << "Usage: " << program << " [options]\n" + << "Options:\n" + << " --start N First matrix size (default: 1000)\n" + << " --step N Size increment (default: 500)\n" + << " --count N Number of tests (default: 6)\n" + << " --sizes a,b,c Comma-separated matrix sizes\n" + << " --threads N Threads per block (default: 256)\n" + << " --max-iters N Max Jacobi iterations (default: 10000)\n" + << " --eps X Convergence epsilon (default: 1e-6)\n" + << " --repeat N Timed repetitions, best is kept (default: 3)\n" + << " --warmup N Warmup repetitions (default: 1)\n" + << " --seed N RNG seed (default: 42)\n" + << " --csv PATH Write CSV summary to PATH\n" + << " --help Print this help\n"; +} + +static bool parse_int_arg(const std::string &text, int &out) { + try { + size_t pos = 0; + const int parsed = std::stoi(text, &pos); + if (pos != text.size()) { + return false; + } + out = parsed; + return true; + } catch (...) { + return false; + } +} + +static bool parse_uint_arg(const std::string &text, unsigned int &out) { + try { + size_t pos = 0; + const unsigned long parsed = std::stoul(text, &pos); + if (pos != text.size()) { + return false; + } + out = static_cast(parsed); + return true; + } catch (...) { + return false; + } +} + +static bool parse_double_arg(const std::string &text, double &out) { + try { + size_t pos = 0; + const double parsed = std::stod(text, &pos); + if (pos != text.size()) { + return false; + } + out = parsed; + return true; + } catch (...) { + return false; + } +} + +static bool parse_sizes_arg(const std::string &text, std::vector &sizes) { + std::stringstream ss(text); + std::string token; + std::vector parsed; + + while (std::getline(ss, token, ',')) { + token.erase( + std::remove_if(token.begin(), + token.end(), + [](unsigned char c) { return std::isspace(c) != 0; }), + token.end()); + if (token.empty()) { + continue; + } + + int value = 0; + if (!parse_int_arg(token, value) || value <= 0) { + return false; + } + parsed.push_back(value); + } + + if (parsed.empty()) { + return false; + } + + sizes = parsed; + return true; +} + +static bool parse_options(int argc, char **argv, Options &options) { + for (int i = 1; i < argc; ++i) { + const std::string arg = argv[i]; + auto require_value = [&](const char *name) -> const char * { + if (i + 1 >= argc) { + std::cerr << "Missing value for " << name << '\n'; + return nullptr; + } + return argv[++i]; + }; + + if (arg == "--help") { + print_usage(argv[0]); + return false; + } + if (arg == "--start") { + const char *v = require_value("--start"); + if (!v || !parse_int_arg(v, options.start) || options.start <= 0) { + std::cerr << "Invalid --start value\n"; + return false; + } + continue; + } + if (arg == "--step") { + const char *v = require_value("--step"); + if (!v || !parse_int_arg(v, options.step) || options.step <= 0) { + std::cerr << "Invalid --step value\n"; + return false; + } + continue; + } + if (arg == "--count") { + const char *v = require_value("--count"); + if (!v || !parse_int_arg(v, options.count) || options.count <= 0) { + std::cerr << "Invalid --count value\n"; + return false; + } + continue; + } + if (arg == "--threads") { + const char *v = require_value("--threads"); + if (!v || !parse_int_arg(v, options.threads) || options.threads <= 0 || + options.threads > 1024) { + std::cerr << "Invalid --threads value\n"; + return false; + } + continue; + } + if (arg == "--max-iters") { + const char *v = require_value("--max-iters"); + if (!v || !parse_int_arg(v, options.max_iters) || + options.max_iters <= 0) { + std::cerr << "Invalid --max-iters value\n"; + return false; + } + continue; + } + if (arg == "--eps") { + const char *v = require_value("--eps"); + if (!v || !parse_double_arg(v, options.eps) || options.eps <= 0.0) { + std::cerr << "Invalid --eps value\n"; + return false; + } + continue; + } + if (arg == "--repeat") { + const char *v = require_value("--repeat"); + if (!v || !parse_int_arg(v, options.repeat) || options.repeat <= 0) { + std::cerr << "Invalid --repeat value\n"; + return false; + } + continue; + } + if (arg == "--warmup") { + const char *v = require_value("--warmup"); + if (!v || !parse_int_arg(v, options.warmup) || options.warmup < 0) { + std::cerr << "Invalid --warmup value\n"; + return false; + } + continue; + } + if (arg == "--seed") { + const char *v = require_value("--seed"); + if (!v || !parse_uint_arg(v, options.seed)) { + std::cerr << "Invalid --seed value\n"; + return false; + } + continue; + } + if (arg == "--sizes") { + const char *v = require_value("--sizes"); + if (!v || !parse_sizes_arg(v, options.sizes)) { + std::cerr << "Invalid --sizes value\n"; + return false; + } + continue; + } + if (arg == "--csv") { + const char *v = require_value("--csv"); + if (!v) { + return false; + } + options.csv_path = v; + continue; + } + + std::cerr << "Unknown option: " << arg << '\n'; + return false; + } + + if (options.sizes.empty()) { + options.sizes.reserve(static_cast(options.count)); + for (int i = 0; i < options.count; ++i) { + options.sizes.push_back(options.start + i * options.step); + } + } + + return true; +} + +static double next_random_value(uint32_t &state) { + state = state * 1664525U + 1013904223U; + const double normalized = + static_cast(state & 0x00FFFFFFU) / + static_cast(0x00FFFFFFU); + return normalized * 2.0 - 1.0; +} + +static void build_system(int n, + unsigned int seed, + std::vector &a, + std::vector &x_true, + std::vector &b) { + const size_t nn = static_cast(n) * static_cast(n); + a.assign(nn, 0.0); + x_true.assign(static_cast(n), 0.0); + b.assign(static_cast(n), 0.0); + + uint32_t state = seed; + for (int i = 0; i < n; ++i) { + x_true[static_cast(i)] = next_random_value(state); + } + + for (int row = 0; row < n; ++row) { + double off_diag_sum = 0.0; + const size_t row_offset = static_cast(row) * static_cast(n); + for (int col = 0; col < n; ++col) { + if (col == row) { + continue; + } + const double value = next_random_value(state); + a[row_offset + static_cast(col)] = value; + off_diag_sum += std::fabs(value); + } + a[row_offset + static_cast(row)] = + off_diag_sum + 2.0 + std::fabs(next_random_value(state)); + } + + for (int row = 0; row < n; ++row) { + const size_t row_offset = static_cast(row) * static_cast(n); + double sum = 0.0; + for (int col = 0; col < n; ++col) { + sum += a[row_offset + static_cast(col)] * + x_true[static_cast(col)]; + } + b[static_cast(row)] = sum; + } +} + +static double compute_residual_inf(const std::vector &a, + const std::vector &x, + const std::vector &b, + int n) { + double max_residual = 0.0; + for (int row = 0; row < n; ++row) { + const size_t row_offset = static_cast(row) * static_cast(n); + double sum = 0.0; + for (int col = 0; col < n; ++col) { + sum += a[row_offset + static_cast(col)] * + x[static_cast(col)]; + } + max_residual = + std::max(max_residual, std::fabs(sum - b[static_cast(row)])); + } + return max_residual; +} + +static double compute_x_error_inf(const std::vector &x, + const std::vector &x_true) { + double max_error = 0.0; + for (size_t i = 0; i < x.size(); ++i) { + max_error = std::max(max_error, std::fabs(x[i] - x_true[i])); + } + return max_error; +} + +static Metrics solve_once(const Options &options, + int n, + const std::vector &a, + const std::vector &b, + const std::vector &x_true, + double *d_a, + double *d_b, + double *d_x_old, + double *d_x_new, + int *d_converged) { + Metrics metrics; + std::vector x(static_cast(n), 0.0); + + CUDA_CHECK(cudaMemset(d_x_old, 0, static_cast(n) * sizeof(double))); + CUDA_CHECK(cudaMemset(d_x_new, 0, static_cast(n) * sizeof(double))); + + cudaEvent_t start = nullptr; + cudaEvent_t stop = nullptr; + CUDA_CHECK(cudaEventCreate(&start)); + CUDA_CHECK(cudaEventCreate(&stop)); + CUDA_CHECK(cudaEventRecord(start, 0)); + + const int block = options.threads; + const int grid = (n + block - 1) / block; + + int h_converged = 0; + int iterations = 0; + while (iterations < options.max_iters) { + h_converged = 1; + CUDA_CHECK(cudaMemcpy( + d_converged, &h_converged, sizeof(int), cudaMemcpyHostToDevice)); + + jacobi_iteration<<>>( + d_a, d_b, d_x_old, d_x_new, n, options.eps, d_converged); + CUDA_CHECK(cudaGetLastError()); + + CUDA_CHECK(cudaMemcpy( + &h_converged, d_converged, sizeof(int), cudaMemcpyDeviceToHost)); + + std::swap(d_x_old, d_x_new); + ++iterations; + + if (h_converged == 1) { + metrics.converged = true; + break; + } + } + + CUDA_CHECK(cudaEventRecord(stop, 0)); + CUDA_CHECK(cudaEventSynchronize(stop)); + + float elapsed_ms = 0.0f; + CUDA_CHECK(cudaEventElapsedTime(&elapsed_ms, start, stop)); + CUDA_CHECK(cudaEventDestroy(start)); + CUDA_CHECK(cudaEventDestroy(stop)); + + CUDA_CHECK(cudaMemcpy(x.data(), + d_x_old, + static_cast(n) * sizeof(double), + cudaMemcpyDeviceToHost)); + + metrics.elapsed_ms = static_cast(elapsed_ms); + metrics.iterations = iterations; + metrics.residual_inf = compute_residual_inf(a, x, b, n); + metrics.x_error_inf = compute_x_error_inf(x, x_true); + + const double seconds = metrics.elapsed_ms / 1000.0; + const double flops = + (2.0 / 3.0) * static_cast(n) * static_cast(n) * + static_cast(n); + metrics.gflops = (seconds > 0.0) ? (flops / seconds) / 1e9 : 0.0; + return metrics; +} + +static Metrics benchmark_size(const Options &options, + int n, + const std::vector &a, + const std::vector &b, + const std::vector &x_true) { + const size_t matrix_bytes = + static_cast(n) * static_cast(n) * sizeof(double); + const size_t vector_bytes = static_cast(n) * sizeof(double); + + double *d_a = nullptr; + double *d_b = nullptr; + double *d_x_old = nullptr; + double *d_x_new = nullptr; + int *d_converged = nullptr; + + CUDA_CHECK(cudaMalloc(reinterpret_cast(&d_a), matrix_bytes)); + CUDA_CHECK(cudaMalloc(reinterpret_cast(&d_b), vector_bytes)); + CUDA_CHECK(cudaMalloc(reinterpret_cast(&d_x_old), vector_bytes)); + CUDA_CHECK(cudaMalloc(reinterpret_cast(&d_x_new), vector_bytes)); + CUDA_CHECK(cudaMalloc(reinterpret_cast(&d_converged), sizeof(int))); + + CUDA_CHECK(cudaMemcpy( + d_a, a.data(), matrix_bytes, cudaMemcpyHostToDevice)); + CUDA_CHECK(cudaMemcpy( + d_b, b.data(), vector_bytes, cudaMemcpyHostToDevice)); + + for (int w = 0; w < options.warmup; ++w) { + (void)solve_once(options, n, a, b, x_true, d_a, d_b, d_x_old, d_x_new, + d_converged); + } + + Metrics best; + for (int run = 0; run < options.repeat; ++run) { + Metrics current = solve_once( + options, n, a, b, x_true, d_a, d_b, d_x_old, d_x_new, d_converged); + if (current.elapsed_ms < best.elapsed_ms) { + best = current; + } + } + + CUDA_CHECK(cudaFree(d_a)); + CUDA_CHECK(cudaFree(d_b)); + CUDA_CHECK(cudaFree(d_x_old)); + CUDA_CHECK(cudaFree(d_x_new)); + CUDA_CHECK(cudaFree(d_converged)); + + return best; +} + +int main(int argc, char **argv) { + Options options; + if (!parse_options(argc, argv, options)) { + return 0; + } + + int device = 0; + CUDA_CHECK(cudaGetDevice(&device)); + cudaDeviceProp prop{}; + CUDA_CHECK(cudaGetDeviceProperties(&prop, device)); + + std::ofstream csv_file; + if (!options.csv_path.empty()) { + csv_file.open(options.csv_path.c_str(), std::ios::out | std::ios::trunc); + if (!csv_file) { + std::cerr << "Failed to open CSV file: " << options.csv_path << '\n'; + return 1; + } + csv_file + << "n,elapsed_ms,iterations,residual_inf,x_error_inf,gflops,converged\n"; + } + + std::cout << "CUDA Jacobi LINPACK-like benchmark\n"; + std::cout << "device = " << prop.name << ", compute capability = " + << prop.major << '.' << prop.minor << ", threads = " + << options.threads << ", repeat = " << options.repeat + << ", warmup = " << options.warmup + << ", eps = " << std::scientific << options.eps << std::defaultfloat + << "\n\n"; + + std::cout << std::left << std::setw(8) << "N" << std::setw(14) << "Time(ms)" + << std::setw(12) << "Iter" << std::setw(18) << "ResidualInf" + << std::setw(18) << "XerrInf" << std::setw(14) << "GFLOPS" + << "Status\n"; + + for (size_t i = 0; i < options.sizes.size(); ++i) { + const int n = options.sizes[i]; + if (n <= 0) { + continue; + } + + std::vector a; + std::vector x_true; + std::vector b; + build_system( + n, options.seed + static_cast(n), a, x_true, b); + + Metrics metrics = benchmark_size(options, n, a, b, x_true); + + std::cout << std::left << std::setw(8) << n << std::setw(14) + << std::fixed << std::setprecision(4) << metrics.elapsed_ms + << std::setw(12) << metrics.iterations << std::setw(18) + << std::scientific << std::setprecision(3) + << metrics.residual_inf << std::setw(18) << metrics.x_error_inf + << std::setw(14) << std::fixed << std::setprecision(3) + << metrics.gflops + << (metrics.converged ? "converged" : "max_iters") << '\n'; + + if (csv_file) { + csv_file << n << ',' << std::fixed << std::setprecision(6) + << metrics.elapsed_ms << ',' << metrics.iterations << ',' + << std::scientific << std::setprecision(8) + << metrics.residual_inf << ',' << metrics.x_error_inf << ',' + << std::fixed << std::setprecision(6) << metrics.gflops + << ',' << (metrics.converged ? 1 : 0) << '\n'; + } + } + + if (csv_file) { + csv_file.close(); + } + + return 0; +}