64-tap FIR filter (real run)

{
  "run": {
    "command": "/opt/syclomatic/bin/c2s --in-root=examples/fir --out-root=migrated --cuda-include-path=/opt/cuda-headers/targets/x86_64-linux/include --extra-arg=-Iexamples/common --extra-arg=-Iexamples/fir examples/fir/fir.cu",
    "returncode": 0,
    "duration_s": 2.233351,
    "stdout": "NOTE: Could not auto-detect compilation database for file 'fir.cu' in 'examples/fir' or any parent directory.\nParsing: examples/fir/fir.cu\nAnalyzing: examples/fir/fir.cu\nMigrating: examples/fir/fir.cu\nProcessed 1 file(s) in -in-root folder \"examples/fir\"\n\nSee Diagnostics Reference to resolve warnings and complete the migration:\nhttps://oneapi-src.github.io/SYCLomatic/dev_guide/reference/diagnostics-reference.html\n",
    "stderr": "examples/fir/fir.cu:36:3: warning: DPCT1113:0: Consider replacing sycl::nd_item::barrier(sycl::access::fence_space::local_space) with sycl::nd_item::barrier() if function \"fir_kernel\" is called in a multidimensional kernel.\n   36 |   __syncthreads();\n      |   ^\nexamples/fir/fir.cu:14:9: warning: DPCT1009:1: SYCL reports errors using exceptions and does not use error codes. Please replace the \"get_error_string_dummy(...)\" with a real error-handling function.\n   14 | #define CUDA_CHECK(call)                                                        \\\n      |         ^\nexamples/fir/fir.cu:70:16: warning: DPCT1010:2: SYCL uses exceptions to report errors and does not use the error codes. The cudaGetLastError function call was replaced with 0. You need to rewrite this code.\n   70 |     CUDA_CHECK(cudaGetLastError());\n      |                ^\nexamples/fir/fir.cu:28:20: warning: DPCT1101:3: 'BLOCK + kTaps - 1' expression was replaced with a value. Modify the code to use the original expression, provided in comments, if it is correct.\n   28 |   __shared__ float tile[BLOCK + kTaps - 1];\n      |                    ^\n\n"
  },
  "input_files": [
    "examples/fir/fir.cu"
  ],
  "in_root": "examples/fir",
  "out_dir": "migrated",
  "migrated_files": [
    "fir.dp.cpp",
    "fir_coeffs.hpp"
  ],
  "warnings": [
    {
      "code": "DPCT1113",
      "id": 0,
      "message": "Consider replacing sycl::nd_item::barrier(sycl::access::fence_space::local_space) with sycl::nd_item::barrier() if function \"fir_kernel\" is called in a multidimensional kernel.",
      "file": "examples/fir/fir.cu",
      "line": 36,
      "column": 3,
      "migrated_file": "fir.dp.cpp",
      "migrated_line": 41
    },
    {
      "code": "DPCT1009",
      "id": 1,
      "message": "SYCL reports errors using exceptions and does not use error codes. Please replace the \"get_error_string_dummy(...)\" with a real error-handling function.",
      "file": "examples/fir/fir.cu",
      "line": 14,
      "column": 9,
      "migrated_file": "fir.dp.cpp",
      "migrated_line": 16
    },
    {
      "code": "DPCT1010",
      "id": 2,
      "message": "SYCL uses exceptions to report errors and does not use the error codes. The cudaGetLastError function call was replaced with 0. You need to rewrite this code.",
      "file": "examples/fir/fir.cu",
      "line": 70,
      "column": 16,
      "migrated_file": "fir.dp.cpp",
      "migrated_line": 110
    },
    {
      "code": "DPCT1101",
      "id": 3,
      "message": "'BLOCK + kTaps - 1' expression was replaced with a value. Modify the code to use the original expression, provided in comments, if it is correct.",
      "file": "examples/fir/fir.cu",
      "line": 28,
      "column": 20,
      "migrated_file": "fir.dp.cpp",
      "migrated_line": 90
    }
  ],
  "warning_counts": {
    "DPCT1113": 1,
    "DPCT1009": 1,
    "DPCT1010": 1,
    "DPCT1101": 1
  },
  "errors": []
}

{
  "path": "migrated/fir.dp.cpp",
  "content": "// CUDA version of the 64-tap FIR (the code to be ported).\n// Coefficients live in __constant__ memory; each thread computes one output sample of\n// one channel; the input tile + history is staged through shared memory.\n// Usage: fir <in.wav> <out.wav> [cutoff_hz=2000]\n#include <sycl/sycl.hpp>\n#include <dpct/dpct.hpp>\n#include <cstdio>\n#include <cstdlib>\n\n#include \"fir_coeffs.hpp\"\n#include \"wav.hpp\"\n\n#define BLOCK 256\n\n/*\nDPCT1009:1: SYCL reports errors using exceptions and does not use error codes.\nPlease replace the \"get_error_string_dummy(...)\" with a real error-handling\nfunction.\n*/\n#define CUDA_CHECK(call)                                                       \\\n  do {                                                                         \\\n    dpct::err0 err_ = (call);                                                  \\\n                                                                               \\\n  } while (0)\n\ninline dpct::constant_memory<float, 1> c_taps(kTaps);\n\n// x and y are one de-interleaved channel of length n.\nvoid fir_kernel(const float* __restrict__ x, float* __restrict__ y, int n,\n                float const *c_taps, float *tile) {\n  auto item_ct1 = sycl::ext::oneapi::this_work_item::get_nd_item<3>();\n\n  const int base = item_ct1.get_group(2) * item_ct1.get_local_range(2);\n  const int t = item_ct1.get_local_id(2);\n  // Load history (kTaps-1 samples before the block) and the block itself.\n  for (int i = t; i < BLOCK + kTaps - 1; i += item_ct1.get_local_range(2)) {\n    const int src = base - (kTaps - 1) + i;\n    tile[i] = (src >= 0 && src < n) ? x[src] : 0.0f;\n  }\n  /*\n  DPCT1113:0: Consider replacing\n  sycl::nd_item::barrier(sycl::access::fence_space::local_space) with\n  sycl::nd_item::barrier() if function \"fir_kernel\" is called in a\n  multidimensional kernel.\n  */\n  item_ct1.barrier(sycl::access::fence_space::local_space);\n  const int idx = base + t;\n  if (idx >= n) return;\n  float acc = 0.0f;\n  for (int k = 0; k < kTaps; ++k) acc += c_taps[k] * tile[t + (kTaps - 1) - k];\n  y[idx] = acc;\n}\n\nint main(int argc, char **argv) try {\n  dpct::device_ext &dev_ct1 = dpct::get_current_device();\n  sycl::queue &q_ct1 = dev_ct1.in_order_queue();\n  if (argc < 3) {\n    std::fprintf(stderr, \"usage: %s in.wav out.wav [cutoff_hz]\\n\", argv[0]);\n    return 2;\n  }\n  const double cutoff = argc > 3 ? std::strtod(argv[3], nullptr) : 2000.0;\n  WavData wav;\n  if (!wav_read(argv[1], wav)) {\n    std::fprintf(stderr, \"cannot read %s\\n\", argv[1]);\n    return 1;\n  }\n  float h[kTaps];\n  make_lowpass(h, cutoff, wav.sample_rate);\n  CUDA_CHECK(\n      DPCT_CHECK_ERROR(q_ct1.memcpy(c_taps.get_ptr(), h, sizeof(h)).wait()));\n\n  const int ch = wav.channels;\n  const int n = static_cast<int>(wav.frames());\n  std::vector<float> chan(n), result(n), out(wav.samples.size());\n  float *d_x = nullptr, *d_y = nullptr;\n  CUDA_CHECK(DPCT_CHECK_ERROR(d_x = sycl::malloc_device<float>(n, q_ct1)));\n  CUDA_CHECK(DPCT_CHECK_ERROR(d_y = sycl::malloc_device<float>(n, q_ct1)));\n\n  for (int c = 0; c < ch; ++c) {\n    for (int i = 0; i < n; ++i) chan[i] = wav.samples[i * ch + c];\n    CUDA_CHECK(DPCT_CHECK_ERROR(\n        q_ct1.memcpy(d_x, chan.data(), n * sizeof(float)).wait()));\n    {\n      c_taps.init();\n\n      q_ct1.submit([&](sycl::handler &cgh) {\n        auto c_taps_ptr_ct1 = c_taps.get_ptr();\n\n        /*\n        DPCT1101:3: 'BLOCK + kTaps - 1' expression was replaced with a value.\n        Modify the code to use the original expression, provided in comments,\n        if it is correct.\n        */\n        sycl::local_accessor<float, 1> tile_acc_ct1(\n            sycl::range<1>(319 /*BLOCK + kTaps - 1*/), cgh);\n\n        cgh.parallel_for(\n            sycl::nd_range<3>(sycl::range<3>(1, 1, (n + BLOCK - 1) / BLOCK) *\n                                  sycl::range<3>(1, 1, BLOCK),\n                              sycl::range<3>(1, 1, BLOCK)),\n            [=](sycl::nd_item<3> item_ct1) {\n              fir_kernel(\n                  d_x, d_y, n, c_taps_ptr_ct1,\n                  tile_acc_ct1.get_multi_ptr<sycl::access::decorated::no>()\n                      .get());\n            });\n      });\n    }\n    /*\n    DPCT1010:2: SYCL uses exceptions to report errors and does not use the error\n    codes. The cudaGetLastError function call was replaced with 0. You need to\n    rewrite this code.\n    */\n    CUDA_CHECK(0);\n    CUDA_CHECK(DPCT_CHECK_ERROR(\n        q_ct1.memcpy(result.data(), d_y, n * sizeof(float)).wait()));\n    for (int i = 0; i < n; ++i) out[i * ch + c] = result[i];\n  }\n  CUDA_CHECK(DPCT_CHECK_ERROR(dpct::dpct_free(d_x, q_ct1)));\n  CUDA_CHECK(DPCT_CHECK_ERROR(dpct::dpct_free(d_y, q_ct1)));\n  if (!wav_write_f32(argv[2], out.data(), out.size(), ch, wav.sample_rate)) return 1;\n  return 0;\n}\ncatch (sycl::exception const &exc) {\n  std::cerr << exc.what() << \"Exception caught at file:\" << __FILE__\n            << \", line:\" << __LINE__ << std::endl;\n  std::exit(1);\n}\n"
}

{
  "path": "migrated/fir_coeffs.hpp",
  "content": "// Shared host-side coefficient design for the FIR example (Hann-windowed sinc low-pass).\n// Computed in double and rounded once to float so every build gets identical taps.\n#pragma once\n\n#include <cmath>\n\nconstexpr int kTaps = 64;\n\ninline void make_lowpass(float* h, double cutoff_hz, int sample_rate) {\n  const double pi = 3.14159265358979323846;\n  const double fc = cutoff_hz / sample_rate;  // normalised cutoff (cycles/sample)\n  const double mid = (kTaps - 1) / 2.0;\n  double taps[kTaps];\n  double sum = 0.0;\n  for (int k = 0; k < kTaps; ++k) {\n    const double t = k - mid;\n    const double sinc = (t == 0.0) ? 2.0 * fc : std::sin(2.0 * pi * fc * t) / (pi * t);\n    const double window = 0.5 - 0.5 * std::cos(2.0 * pi * k / (kTaps - 1));\n    taps[k] = sinc * window;\n    sum += taps[k];\n  }\n  for (int k = 0; k < kTaps; ++k) h[k] = static_cast<float>(taps[k] / sum);  // unity DC gain\n}\n"
}

// SYCL 2020 version of the 64-tap FIR.
// Taps are captured by value in the kernel (passed as kernel arguments, i.e. constant
// parameter storage); each work-item computes one output sample of one channel; the input
// tile + history is staged through work-group local memory.
// Usage: fir <in.wav> <out.wav> [cutoff_hz=2000]
#include <sycl/sycl.hpp>
#include <array>
#include <cstdio>
#include <cstdlib>
#include <iostream>
#include <vector>

#include "fir_coeffs.hpp"
#include "wav.hpp"

constexpr int kBlock = 256;
constexpr int kTile = kBlock + kTaps - 1;

int main(int argc, char **argv) try {
  if (argc < 3) {
    std::fprintf(stderr, "usage: %s in.wav out.wav [cutoff_hz]\n", argv[0]);
    return 2;
  }
  const double cutoff = argc > 3 ? std::strtod(argv[3], nullptr) : 2000.0;
  WavData wav;
  if (!wav_read(argv[1], wav)) {
    std::fprintf(stderr, "cannot read %s\n", argv[1]);
    return 1;
  }
  std::array<float, kTaps> taps;
  make_lowpass(taps.data(), cutoff, wav.sample_rate);

  sycl::queue q{sycl::property::queue::in_order()};

  const int ch = wav.channels;
  const int n = static_cast<int>(wav.frames());
  std::vector<float> chan(n), result(n), out(wav.samples.size());
  float *d_x = sycl::malloc_device<float>(n, q);
  float *d_y = sycl::malloc_device<float>(n, q);

  const size_t groups = (static_cast<size_t>(n) + kBlock - 1) / kBlock;

  for (int c = 0; c < ch; ++c) {
    for (int i = 0; i < n; ++i) chan[i] = wav.samples[i * ch + c];
    q.memcpy(d_x, chan.data(), n * sizeof(float));
    q.submit([&](sycl::handler &cgh) {
      sycl::local_accessor<float, 1> tile(sycl::range<1>(kTile), cgh);
      const float *x = d_x;
      float *y = d_y;
      cgh.parallel_for(
          sycl::nd_range<1>(groups * kBlock, kBlock),
          [=](sycl::nd_item<1> it) {
            const int base = static_cast<int>(it.get_group(0)) * kBlock;
            const int t = static_cast<int>(it.get_local_id(0));
            for (int i = t; i < kTile; i += kBlock) {
              const int src = base - (kTaps - 1) + i;
              tile[i] = (src >= 0 && src < n) ? x[src] : 0.0f;
            }
            sycl::group_barrier(it.get_group());
            const int idx = base + t;
            if (idx >= n) return;
            float acc = 0.0f;
            for (int k = 0; k < kTaps; ++k)
              acc += taps[k] * tile[t + (kTaps - 1) - k];
            y[idx] = acc;
          });
    });
    q.memcpy(result.data(), d_y, n * sizeof(float)).wait();
    for (int i = 0; i < n; ++i) out[i * ch + c] = result[i];
  }
  sycl::free(d_x, q);
  sycl::free(d_y, q);
  if (!wav_write_f32(argv[2], out.data(), out.size(), ch, wav.sample_rate)) return 1;
  return 0;
} catch (sycl::exception const &e) {
  std::cerr << "SYCL exception: " << e.what() << std::endl;
  return 1;
}

{
  "path": "migrated/fir.dp.cpp"
}

{
  "run": {
    "command": "/opt/dpcpp/bin/icpx -fsycl -O2 -Iexamples/common -Iexamples/fir migrated/fir.dp.cpp -o build/candidate",
    "returncode": 0,
    "duration_s": 5.066484,
    "stdout": "",
    "stderr": ""
  },
  "output": "build/candidate",
  "dpct_include_added": null,
  "diagnostics": [],
  "counts": {
    "error": 0,
    "warning": 0,
    "note": 0
  }
}

{
  "path": "migrated",
  "files_scanned": 2,
  "total": 0,
  "by_category": {},
  "by_symbol": {},
  "by_file": {},
  "occurrences": [],
  "occurrences_truncated": 0,
  "suggestions": {},
  "unresolved_dpct_comments": [],
  "advisories": []
}

{
  "mode": "audio",
  "cases": [
    {
      "input": "examples/data/sweep_noise_48k.wav",
      "reference": {
        "command": "build/reference examples/data/sweep_noise_48k.wav verify_out/case0_reference.wav",
        "returncode": 0,
        "duration_s": 0.008299,
        "stdout": "",
        "stderr": "",
        "cwd": ".",
        "output_file": "verify_out/case0_reference.wav"
      },
      "candidate": {
        "command": "build/candidate examples/data/sweep_noise_48k.wav verify_out/case0_candidate.wav",
        "returncode": 0,
        "duration_s": 0.246474,
        "stdout": "",
        "stderr": "",
        "cwd": ".",
        "output_file": "verify_out/case0_candidate.wav"
      },
      "reference_format": "wav",
      "candidate_format": "wav",
      "mode": "audio",
      "stats": {
        "threshold_dbfs": -96.0,
        "reference_frames": 96000,
        "candidate_frames": 96000,
        "reference_channels": 1,
        "candidate_channels": 1,
        "reference_sample_rate": 48000,
        "candidate_sample_rate": 48000,
        "passed": true,
        "bit_exact": false,
        "residual_peak": 1.1920928955078125e-07,
        "residual_peak_dbfs": -138.47379800543135,
        "residual_rms_dbfs": -155.56082682898696,
        "reference_rms_dbfs": -10.983389376767594,
        "null_depth_db": 144.57743745221936,
        "per_channel": [
          {
            "channel": 0,
            "residual_peak_dbfs": -138.47379800543135,
            "residual_rms_dbfs": -155.56082682898696,
            "peak_frame": 604
          }
        ]
      },
      "passed": true
    },
    {
      "input": "examples/data/sweep_noise_stereo_44k.wav",
      "reference": {
        "command": "build/reference examples/data/sweep_noise_stereo_44k.wav verify_out/case1_reference.wav",
        "returncode": 0,
        "duration_s": 0.008478,
        "stdout": "",
        "stderr": "",
        "cwd": ".",
        "output_file": "verify_out/case1_reference.wav"
      },
      "candidate": {
        "command": "build/candidate examples/data/sweep_noise_stereo_44k.wav verify_out/case1_candidate.wav",
        "returncode": 0,
        "duration_s": 0.227987,
        "stdout": "",
        "stderr": "",
        "cwd": ".",
        "output_file": "verify_out/case1_candidate.wav"
      },
      "reference_format": "wav",
      "candidate_format": "wav",
      "mode": "audio",
      "stats": {
        "threshold_dbfs": -96.0,
        "reference_frames": 44100,
        "candidate_frames": 44100,
        "reference_channels": 2,
        "candidate_channels": 2,
        "reference_sample_rate": 44100,
        "candidate_sample_rate": 44100,
        "passed": true,
        "bit_exact": false,
        "residual_peak": 1.7881393432617188e-07,
        "residual_peak_dbfs": -134.95197282431772,
        "residual_rms_dbfs": -155.6051398509516,
        "reference_rms_dbfs": -10.941526370582011,
        "null_depth_db": 144.66361348036958,
        "per_channel": [
          {
            "channel": 0,
            "residual_peak_dbfs": -134.95197282431772,
            "residual_rms_dbfs": -155.5692483331675,
            "peak_frame": 5773
          },
          {
            "channel": 1,
            "residual_peak_dbfs": -138.47379800543135,
            "residual_rms_dbfs": -155.6413304615059,
            "peak_frame": 22507
          }
        ]
      },
      "passed": true
    }
  ],
  "work_dir": "verify_out",
  "passed_cases": 2,
  "total_cases": 2,
  "worst_residual_peak_dbfs": -134.95197282431772
}

{
  "command": "build/candidate examples/data/sweep_noise_48k.wav verify_out/bench_out",
  "runs": 5,
  "warmup": 1,
  "device_discovery": {
    "sycl_ls": {
      "command": "/opt/dpcpp/bin/sycl-ls",
      "returncode": 0,
      "duration_s": 0.04091,
      "stdout": "[opencl:cpu][opencl:0] Intel(R) OpenCL, Intel(R) Xeon(R) Processor @ 2.10GHz OpenCL 3.0 (Build 0) [2026.21.7.0.24_160000]\n",
      "stderr": ""
    }
  },
  "results": [
    {
      "selector": "opencl:0",
      "device": {
        "backend": "opencl",
        "device_type": "cpu",
        "index": 0,
        "platform": "Intel(R) OpenCL",
        "name": "Intel(R) Xeon(R) Processor @ 2.10GHz",
        "description": "Intel(R) OpenCL, Intel(R) Xeon(R) Processor @ 2.10GHz OpenCL 3.0 (Build 0) [2026.21.7.0.24_160000]",
        "selector": "opencl:0"
      },
      "status": "ok",
      "wall_time_s": {
        "n": 5,
        "min": 0.21268019299998286,
        "median": 0.2237716370000271,
        "p95": 0.26450542000001176,
        "max": 0.27016712900001494,
        "mean": 0.23341669000000137,
        "stdev": 0.023260744771388144
      },
      "samples_s": [
        0.241859,
        0.270167,
        0.218606,
        0.223772,
        0.21268
      ]
    }
  ]
}

{
  "path": ".",
  "files_scanned": 2,
  "total": 0,
  "by_category": {},
  "by_symbol": {},
  "by_file": {},
  "occurrences": [],
  "occurrences_truncated": 0,
  "suggestions": {},
  "unresolved_dpct_comments": [],
  "advisories": []
}