With the toolkit: softmax verified

{
  "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": 7.666796,
    "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.
// Usage: softmax <out.bin> [numSlice=8192] [sliceSize=256] [kernel=1] [repeat=1]
#include <sycl/sycl.hpp>
#include <algorithm>
#include <chrono>
#include <cstdlib>
#include <cstdio>
#include <vector>

#include "softmax_input.hpp"

#define BLOCK_SIZE 256

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();

  // Sub-group size is device dependent: launch enough work-groups for the
  // largest supported sub-group size (fewest sub-groups per work-group).
  size_t maxSg = 1;
  for (size_t s : q.get_device().get_info<sycl::info::device::sub_group_sizes>())
    maxSg = std::max(maxSg, s);
  size_t minSgPerGroup = std::max<size_t>(1, BLOCK_SIZE / maxSg);

  auto start = std::chrono::steady_clock::now();
  if (kernel == 1) {
    size_t groups = (numSlice + minSgPerGroup - 1) / minSgPerGroup;
    for (int n = 0; n < repeat; n++) {
      q.parallel_for(
          sycl::nd_range<1>(groups * BLOCK_SIZE, BLOCK_SIZE),
          [=](sycl::nd_item<1> item) {
            sycl::sub_group sg = item.get_sub_group();
            const int i = item.get_group(0) * sg.get_group_linear_range() +
                          sg.get_group_linear_id();
            if (i >= numSlice) return;
            const float *row = d_input + (size_t)i * sliceSize;
            float *out = d_output + (size_t)i * sliceSize;
            const int lane = sg.get_local_linear_id();
            const int width = sg.get_local_linear_range();
            float max_ = row[0];
            for (int j = lane; j < sliceSize; j += width)
              max_ = sycl::max(max_, row[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(row[j] - max_);
            sum = sycl::reduce_over_group(sg, sum, sycl::plus<float>{});
            for (int j = lane; j < sliceSize; j += width)
              out[j] = sycl::exp(row[j] - max_) / sum;
          });
    }
  } 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> item) {
            const int i = item.get_global_id(0);
            if (i >= numSlice) return;
            const float *row = d_input + (size_t)i * sliceSize;
            float *out = d_output + (size_t)i * sliceSize;
            float max_ = row[0];
            for (int j = 0; j < sliceSize; j++) max_ = sycl::max(max_, row[j]);
            float sum = 0;
            for (int j = 0; j < sliceSize; j++) sum += sycl::exp(row[j] - max_);
            for (int j = 0; j < sliceSize; j++)
              out[j] = sycl::exp(row[j] - max_) / sum;
          });
    }
  }
  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": 6.004739,
    "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.025054,
        "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.362439,
        "stdout": "kernel_ms=295.356537\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.2332124610661183e-08,
        "mean_abs_err": 4.604463579296544e-09,
        "max_rel_err": 2.598395370445306e-06,
        "max_ulp_f32": 41,
        "median_ulp_f32": 14.0,
        "exact_matches": 3909,
        "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.039294,
      "stdout": "[opencl:cpu][opencl:0] Intel(R) OpenCL, Intel(R) Xeon(R) Processor @ 2.80GHz 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.80GHz",
        "description": "Intel(R) OpenCL, Intel(R) Xeon(R) Processor @ 2.80GHz OpenCL 3.0 (Build 0) [2026.21.7.0.24_160000]",
        "selector": "opencl:0"
      },
      "status": "ok",
      "wall_time_s": {
        "n": 5,
        "min": 0.2605348889999277,
        "median": 0.29769088299985924,
        "p95": 0.3158420973999455,
        "max": 0.3169648339999185,
        "mean": 0.2904872989999149,
        "stdev": 0.025933533907642182
      },
      "samples_s": [
        0.260535,
        0.297691,
        0.311351,
        0.265895,
        0.316965
      ]
    }
  ]
}

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

{
  "type": "evaluation",
  "source_hash": "7151056c15044628",
  "compiled": true,
  "verify_status": "ok",
  "verify_summary": "PASS: 1/1 case(s) match (worst max abs err 9.69e-08, atol 1e-06)",
  "helpers": 0,
  "advisories": []
}