HeCBench softmax (real run)

{
  "run": {
    "command": "/opt/syclomatic/bin/c2s --in-root=examples/softmax_hecbench --out-root=migrated --cuda-include-path=/opt/cuda-headers/targets/x86_64-linux/include --extra-arg=-Iexamples/softmax_hecbench examples/softmax_hecbench/softmax.cu",
    "returncode": 0,
    "duration_s": 5.851147,
    "stdout": "NOTE: Could not auto-detect compilation database for file 'softmax.cu' in 'examples/softmax_hecbench' or any parent directory.\nParsing: examples/softmax_hecbench/softmax.cu\nAnalyzing: examples/softmax_hecbench/softmax.cu\nMigrating: examples/softmax_hecbench/softmax.cu\nProcessed 1 file(s) in -in-root folder \"examples/softmax_hecbench\"\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": "\n"
  },
  "input_files": [
    "examples/softmax_hecbench/softmax.cu"
  ],
  "in_root": "examples/softmax_hecbench",
  "out_dir": "migrated",
  "migrated_files": [
    "softmax.dp.cpp",
    "softmax_input.hpp"
  ],
  "warnings": [],
  "warning_counts": {},
  "errors": []
}

{
  "path": "migrated/softmax.dp.cpp",
  "content": "// CUDA softmax (the code to be ported): the two kernels of HeCBench softmax-cuda.\n// Adapted from HeCBench (https://github.com/zjin-lcf/HeCBench, src/softmax-cuda/main.cu,\n// commit f0c817aa). Copyright (c) 2020 Argonne National Laboratory, 2020- Oak Ridge National\n// Laboratory (Zheming Jin). BSD 3-Clause License; see LICENSE.HeCBench in this directory.\n// Adaptations: kernels unchanged; deterministic LCG input; the result is written to a file\n// instead of being checked in-process; argument order changed.\n// Usage: softmax <out.bin> [numSlice=8192] [sliceSize=256] [kernel=1] [repeat=1]\n#include <sycl/sycl.hpp>\n#include <dpct/dpct.hpp>\n#include <chrono>\n#include <cstdlib>\n#include <cstdio>\n#include <vector>\n\n#include \"softmax_input.hpp\"\n\n#define BLOCK_SIZE 256\n\n\nvoid softMax (const int numSlice, const int sliceSize,\n              const float* src, float* dest)\n{\n  auto item_ct1 = sycl::ext::oneapi::this_work_item::get_nd_item<3>();\n  int i = item_ct1.get_group(2) * item_ct1.get_local_range(2) +\n          item_ct1.get_local_id(2);\n  if (i >= numSlice) return;\n  float max_ = src[i * sliceSize];\n  for (int j = 0; j < sliceSize; j++) {\n    max_ = sycl::max(max_, src[i * sliceSize + j]);\n  }\n  float sum = 0;\n  for (int j = 0; j < sliceSize; j++) {\n    sum += sycl::native::exp(src[i * sliceSize + j] - max_);\n  }\n  for (int j = 0; j < sliceSize; j++) {\n    dest[i * sliceSize + j] =\n        sycl::native::exp(src[i * sliceSize + j] - max_) / sum;\n  }\n}\n\n\nvoid softMax2 (const int numSlice, const int sliceSize,\n              const float* src, float* dest)\n{\n\n  sycl::group<3> block = sycl::ext::oneapi::this_work_item::get_work_group<3>();\n  sycl::sub_group warp = sycl::ext::oneapi::this_work_item::get_sub_group();\n  int i =\n      sycl::ext::oneapi::this_work_item::get_nd_item<3>().get_group(2) *\n          sycl::ext::oneapi::this_work_item::get_sub_group()\n              .get_group_linear_range() +\n      sycl::ext::oneapi::this_work_item::get_sub_group().get_group_linear_id();\n  if (i >= numSlice) return;\n  float max_ = src[i * sliceSize];\n  for (int j = sycl::ext::oneapi::this_work_item::get_sub_group()\n                   .get_local_linear_id();\n       j < sliceSize; j += sycl::ext::oneapi::this_work_item::get_sub_group()\n                               .get_local_linear_range()) {\n    max_ = sycl::max(max_, src[i * sliceSize + j]);\n  }\n  max_ = sycl::reduce_over_group(\n      sycl::ext::oneapi::this_work_item::get_sub_group(), max_,\n      sycl::maximum<float>{});\n  float sum = 0;\n  for (int j = sycl::ext::oneapi::this_work_item::get_sub_group()\n                   .get_local_linear_id();\n       j < sliceSize; j += sycl::ext::oneapi::this_work_item::get_sub_group()\n                               .get_local_linear_range()) {\n    sum += sycl::native::exp(src[i * sliceSize + j] - max_);\n  }\n  sum = sycl::reduce_over_group(\n      sycl::ext::oneapi::this_work_item::get_sub_group(), sum,\n      sycl::plus<float>{});\n  for (int j = sycl::ext::oneapi::this_work_item::get_sub_group()\n                   .get_local_linear_id();\n       j < sliceSize; j += sycl::ext::oneapi::this_work_item::get_sub_group()\n                               .get_local_linear_range()) {\n    dest[i * sliceSize + j] =\n        sycl::native::exp(src[i * sliceSize + j] - max_) / sum;\n  }\n}\n\nint main(int argc, char *argv[]) {\n  dpct::device_ext &dev_ct1 = dpct::get_current_device();\n  sycl::queue &q_ct1 = dev_ct1.in_order_queue();\n  if (argc < 2) {\n    printf(\"Usage: %s <out.bin> [number of slices] [slice size] [kernel 0|1] [repeat]\\n\", argv[0]);\n    return 1;\n  }\n  int numSlice = argc > 2 ? atoi(argv[2]) : 8192;\n  int sliceSize = argc > 3 ? atoi(argv[3]) : 256;\n  int kernel = argc > 4 ? atoi(argv[4]) : 1;\n  int repeat = argc > 5 ? atoi(argv[5]) : 1;\n  int numElem = numSlice * sliceSize;\n\n  std::vector<float> input(numElem), output_gpu(numElem);\n  fill_input(input.data(), numElem);\n\n  float *d_input, *d_output;\n  d_input = sycl::malloc_device<float>(numElem, q_ct1);\n  d_output = sycl::malloc_device<float>(numElem, q_ct1);\n  q_ct1.memcpy(d_input, input.data(), sizeof(float) * numElem).wait();\n\n  dev_ct1.queues_wait_and_throw();\n  auto start = std::chrono::steady_clock::now();\n  if (kernel == 1) {\n    dpct::dim3 grids((numSlice + BLOCK_SIZE / 32 - 1) / (BLOCK_SIZE / 32));\n    dpct::dim3 blocks(BLOCK_SIZE);\n    for (int n = 0; n < repeat; n++) {\n      q_ct1.parallel_for(sycl::nd_range<3>(grids * blocks, blocks),\n                         [=](sycl::nd_item<3> item_ct1)\n                             [[sycl::reqd_sub_group_size(32)]] {\n                               softMax2(numSlice, sliceSize, d_input, d_output);\n                             });\n    }\n  } else {\n    dpct::dim3 grids((numSlice + BLOCK_SIZE - 1) / BLOCK_SIZE);\n    dpct::dim3 blocks(BLOCK_SIZE);\n    for (int n = 0; n < repeat; n++) {\n      q_ct1.parallel_for(sycl::nd_range<3>(grids * blocks, blocks),\n                         [=](sycl::nd_item<3> item_ct1) {\n                           softMax(numSlice, sliceSize, d_input, d_output);\n                         });\n    }\n  }\n  dev_ct1.queues_wait_and_throw();\n  auto end = std::chrono::steady_clock::now();\n  auto time = std::chrono::duration_cast<std::chrono::nanoseconds>(end - start).count();\n  printf(\"kernel_ms=%f\\n\", (time * 1e-6f) / repeat);\n\n  q_ct1.memcpy(output_gpu.data(), d_output, sizeof(float) * numElem).wait();\n  FILE* f = fopen(argv[1], \"wb\");\n  if (!f || fwrite(output_gpu.data(), sizeof(float), numElem, f) != size_t(numElem)) return 1;\n  fclose(f);\n  dpct::dpct_free(d_input, q_ct1);\n  dpct::dpct_free(d_output, q_ct1);\n  return 0;\n}\n"
}

{
  "path": "migrated/softmax_input.hpp",
  "content": "// Deterministic input shared by all softmax builds (replaces HeCBench's srand(2)/rand() % 13,\n// whose sequence depends on the C library). Values are integers in [0, 12], as in HeCBench.\n#pragma once\n\n#include <cstdint>\n\ninline void fill_input(float* x, int n) {\n  uint32_t s = 2u;\n  for (int i = 0; i < n; i++) {\n    s = s * 1664525u + 1013904223u;  // Numerical Recipes LCG\n    x[i] = static_cast<float>((s >> 16) % 13u);\n  }\n}\n"
}

// SYCL softmax: the two kernels of HeCBench softmax-cuda.
// Adapted from HeCBench (https://github.com/zjin-lcf/HeCBench, src/softmax-cuda/main.cu,
// commit f0c817aa). Copyright (c) 2020 Argonne National Laboratory, 2020- Oak Ridge National
// Laboratory (Zheming Jin). BSD 3-Clause License; see LICENSE.HeCBench in this directory.
// Usage: softmax <out.bin> [numSlice=8192] [sliceSize=256] [kernel=1] [repeat=1]
#include <sycl/sycl.hpp>
#include <chrono>
#include <cstdlib>
#include <cstdio>
#include <vector>

#include "softmax_input.hpp"

constexpr int BLOCK_SIZE = 256;

// One work-item per slice.
static void softMax(sycl::nd_item<1> it, const int numSlice, const int sliceSize,
                    const float* src, float* dest) {
  int i = static_cast<int>(it.get_global_id(0));
  if (i >= numSlice) return;
  const float* in = src + i * sliceSize;
  float* out = dest + i * sliceSize;
  float max_ = in[0];
  for (int j = 0; j < sliceSize; j++) max_ = sycl::max(max_, in[j]);
  float sum = 0;
  for (int j = 0; j < sliceSize; j++) sum += sycl::exp(in[j] - max_);
  for (int j = 0; j < sliceSize; j++) out[j] = sycl::exp(in[j] - max_) / sum;
}

// One sub-group per slice; independent of the sub-group size (grid-stride over slices).
static void softMax2(sycl::nd_item<1> it, const int numSlice, const int sliceSize,
                     const float* src, float* dest) {
  sycl::sub_group sg = it.get_sub_group();
  const int lane = static_cast<int>(sg.get_local_linear_id());
  const int width = static_cast<int>(sg.get_local_linear_range());
  const int sgPerGroup = static_cast<int>(sg.get_group_linear_range());
  const int totalSg = static_cast<int>(it.get_group_range(0)) * sgPerGroup;
  for (int i = static_cast<int>(it.get_group(0)) * sgPerGroup +
               static_cast<int>(sg.get_group_linear_id());
       i < numSlice; i += totalSg) {
    const float* in = src + i * sliceSize;
    float* out = dest + i * sliceSize;
    float max_ = in[0];
    for (int j = lane; j < sliceSize; j += width) max_ = sycl::max(max_, in[j]);
    max_ = sycl::reduce_over_group(sg, max_, sycl::maximum<float>{});
    float sum = 0;
    for (int j = lane; j < sliceSize; j += width) sum += sycl::exp(in[j] - max_);
    sum = sycl::reduce_over_group(sg, sum, sycl::plus<float>{});
    for (int j = lane; j < sliceSize; j += width) out[j] = sycl::exp(in[j] - max_) / sum;
  }
}

int main(int argc, char* argv[]) {
  if (argc < 2) {
    printf("Usage: %s <out.bin> [number of slices] [slice size] [kernel 0|1] [repeat]\n", argv[0]);
    return 1;
  }
  int numSlice = argc > 2 ? atoi(argv[2]) : 8192;
  int sliceSize = argc > 3 ? atoi(argv[3]) : 256;
  int kernel = argc > 4 ? atoi(argv[4]) : 1;
  int repeat = argc > 5 ? atoi(argv[5]) : 1;
  int numElem = numSlice * sliceSize;

  std::vector<float> input(numElem), output_gpu(numElem);
  fill_input(input.data(), numElem);

  sycl::queue q{sycl::property::queue::in_order()};
  float* d_input = sycl::malloc_device<float>(numElem, q);
  float* d_output = sycl::malloc_device<float>(numElem, q);
  q.memcpy(d_input, input.data(), sizeof(float) * numElem).wait();

  auto start = std::chrono::steady_clock::now();
  if (kernel == 1) {
    // Assume at least 8 sub-groups per work-group worth of slices; the grid-stride
    // loop keeps the result correct for any sub-group size.
    size_t groups = (numSlice + BLOCK_SIZE / 32 - 1) / (BLOCK_SIZE / 32);
    for (int n = 0; n < repeat; n++) {
      q.parallel_for(sycl::nd_range<1>(groups * BLOCK_SIZE, BLOCK_SIZE),
                     [=](sycl::nd_item<1> it) {
                       softMax2(it, numSlice, sliceSize, d_input, d_output);
                     });
    }
  } else {
    size_t groups = (numSlice + BLOCK_SIZE - 1) / BLOCK_SIZE;
    for (int n = 0; n < repeat; n++) {
      q.parallel_for(sycl::nd_range<1>(groups * BLOCK_SIZE, BLOCK_SIZE),
                     [=](sycl::nd_item<1> it) {
                       softMax(it, numSlice, sliceSize, d_input, d_output);
                     });
    }
  }
  q.wait_and_throw();
  auto end = std::chrono::steady_clock::now();
  auto time = std::chrono::duration_cast<std::chrono::nanoseconds>(end - start).count();
  printf("kernel_ms=%f\n", (time * 1e-6f) / repeat);

  q.memcpy(output_gpu.data(), d_output, sizeof(float) * numElem).wait();
  FILE* f = fopen(argv[1], "wb");
  if (!f || fwrite(output_gpu.data(), sizeof(float), numElem, f) != size_t(numElem)) return 1;
  fclose(f);
  sycl::free(d_input, q);
  sycl::free(d_output, q);
  return 0;
}

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

{
  "run": {
    "command": "/opt/dpcpp/bin/icpx -fsycl -O2 -Iexamples/softmax_hecbench migrated/softmax.dp.cpp -o build/candidate",
    "returncode": 0,
    "duration_s": 4.383667,
    "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": "array",
  "cases": [
    {
      "input": null,
      "reference": {
        "command": "build/reference verify_out/case0_reference.out 2048 256",
        "returncode": 0,
        "duration_s": 0.012049,
        "stdout": "",
        "stderr": "",
        "cwd": ".",
        "output_file": "verify_out/case0_reference.out"
      },
      "candidate": {
        "command": "build/candidate verify_out/case0_candidate.out 2048 256",
        "returncode": 0,
        "duration_s": 0.272008,
        "stdout": "kernel_ms=200.125824\n",
        "stderr": "",
        "cwd": ".",
        "output_file": "verify_out/case0_candidate.out"
      },
      "reference_format": "f32",
      "candidate_format": "f32",
      "mode": "array",
      "stats": {
        "n_reference": 524288,
        "n_candidate": 524288,
        "atol": 1e-06,
        "rtol": 0.0,
        "max_ulp_allowed": null,
        "passed": true,
        "max_abs_err": 9.685754776000977e-08,
        "rms_err": 1.2343314380908246e-08,
        "mean_abs_err": 4.604832123386105e-09,
        "max_rel_err": 2.598395370445306e-06,
        "max_ulp_f32": 41,
        "median_ulp_f32": 14.0,
        "exact_matches": 4144,
        "mismatches": 0,
        "nan_mismatches": 0,
        "reference_range": [
          1.2936880011693574e-07,
          0.056325364857912064
        ],
        "first_mismatches": []
      },
      "passed": true
    }
  ],
  "work_dir": "verify_out",
  "passed_cases": 1,
  "total_cases": 1,
  "worst_max_abs_err": 9.685754776000977e-08
}

{
  "command": "build/candidate verify_out/bench_out 2048 256",
  "runs": 5,
  "warmup": 1,
  "device_discovery": {
    "sycl_ls": {
      "command": "/opt/dpcpp/bin/sycl-ls",
      "returncode": 0,
      "duration_s": 0.03496,
      "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.2721300769999857,
        "median": 0.2756937129999528,
        "p95": 0.28072772819999725,
        "max": 0.28166232000000946,
        "mean": 0.2763203885999815,
        "stdev": 0.003477090220821803
      },
      "samples_s": [
        0.281662,
        0.275126,
        0.275694,
        0.27213,
        0.276989
      ]
    }
  ]
}

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