Without the toolkit: softmax ships wrong

{
  "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.685694,
    "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 <algorithm>
#include <chrono>
#include <cstdlib>
#include <cstdio>
#include <vector>

#include "softmax_input.hpp"

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

  // The sub-group size is chosen by the runtime. Size the grid for the smallest
  // supported sub-group so that every slice gets a sub-group; surplus ones exit early.
  size_t min_sg = BLOCK_SIZE;
  for (auto s : q.get_device().get_info<sycl::info::device::sub_group_sizes>())
    min_sg = std::min<size_t>(min_sg, s);
  const size_t max_sg_per_group = BLOCK_SIZE / min_sg;

  auto start = std::chrono::steady_clock::now();
  if (kernel == 1) {
    const size_t groups = (numSlice + max_sg_per_group - 1) / max_sg_per_group;
    for (int n = 0; n < repeat; n++) {
      q.parallel_for(
          sycl::nd_range<1>(groups * BLOCK_SIZE, BLOCK_SIZE),
          [=](sycl::nd_item<1> it) {
            sycl::sub_group sg = it.get_sub_group();
            const int i = int(it.get_group(0) * sg.get_group_range()[0] +
                              sg.get_group_id()[0]);
            if (i >= numSlice) return;
            const int lane = int(sg.get_local_id()[0]);
            const int step = int(sg.get_local_range()[0]);
            const float *row = d_input + size_t(i) * sliceSize;
            float *out = d_output + size_t(i) * sliceSize;
            float max_ = row[0];
            for (int j = lane; j < sliceSize; j += step)
              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 += step)
              sum += sycl::exp(row[j] - max_);
            sum = sycl::reduce_over_group(sg, sum, sycl::plus<float>{});
            for (int j = lane; j < sliceSize; j += step)
              out[j] = sycl::exp(row[j] - max_) / sum;
          });
    }
  } else {
    const 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) {
            const int i = int(it.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": 5.610994,
    "stdout": "",
    "stderr": ""
  },
  "output": "build/candidate",
  "dpct_include_added": null,
  "diagnostics": [],
  "counts": {
    "error": 0,
    "warning": 0,
    "note": 0
  }
}

{
  "model_summary": "Ported to SYCL 2020 with no dpct headers or helpers. Compiled cleanly with add_dpct_include=false. I did not run the binary or compare its output with the reference.\n\n- **Sub-group size:** softMax2 uses `reduce_over_group` on the sub-group, and the sub-group size is not hard-coded to 32. The grid is sized for the device's smallest supported sub-group size, so surplus sub-groups exit early. That is wasteful when the runtime picks a larger sub-group.\n- **Precision:** `sycl::exp` replaces the `native::exp` that c2s generated, to stay within the 1e-6 tolerance.\n- **Launch:** both kernels use 1-D `nd_range` launches."
}

{
  "type": "evaluation",
  "source_hash": "b0bdf6470571f73f",
  "compiled": true,
  "verify_status": "fail",
  "verify_summary": "FAIL: 0/1 case(s) match; first failure: 332679 of 524288 element(s) outside tolerance",
  "helpers": 0,
  "advisories": []
}