Reduction function throws an instance of 'cl::sycl::nd_range_error'

[Wanzizhu]

When i use reduction function, it raised a 'cl::sycl::nd_range_error' when input array >=409600 bytes.
Below is the detailed error log:

terminate called after throwing an instance of 'cl::sycl::nd_range_error'                                                                                                             
 what():  Global_work_size not evenly divisible by local_work_size. Non-uniform work-groups are not allowed by default.  
 Underlying OpenCL 2.x implementation supports this feature and to enable it, build device program with -cl-std=CL2.0 -54 (CL_INVALID_WORK_GROUP_SIZE)  
 Aborted (core dumped)

but actually my global_work_size = num_work_group * work_group_size
Below is the code:

#include <algorithm>
#include <iostream>
#include <random>
#include <CL/sycl.hpp>

#define GET_VARIABLE_NAME(Variable) (#Variable)

#define SHOW_DEVICE_INFO(dev, prop) \
  std::cout << GET_VARIABLE_NAME(prop) << ": " << dev.get_info<cl::sycl::info::device::prop>() \
            << std::endl;

int device_info(cl::sycl::device dev) {
  const cl::sycl::platform plat = dev.get_platform();
  std::cout << "platform: " << plat.get_info<cl::sycl::info::platform::name>() << std::endl;
  SHOW_DEVICE_INFO(dev, name);
  SHOW_DEVICE_INFO(dev, vendor);
  SHOW_DEVICE_INFO(dev, driver_version);
  std::cout << std::endl;

  return 0;
}

template <typename T>
void RandomInit(T *arr, size_t nelems) {
  std::random_device rd;
  std::mt19937 gen(rd());
  std::uniform_real_distribution<T> rg(0.f, 1.f);
  std::generate(arr, arr + nelems, [&] { return rg(gen); });
}

template <typename IN_T>
float reduce_algorithm(cl::sycl::queue &q, IN_T in, size_t nelem) {
  auto max_group_sz = 256;
  auto num_gz = (nelem + max_group_sz - 1) / max_group_sz;
  cl::sycl::range<1> local(max_group_sz);
  cl::sycl::range<1> global(num_gz * max_group_sz);
  std::array<float, 1> result;
  result[0] = 0.0f;
  {
    cl::sycl::buffer<float> res_buf(result.data(), 1);
    printf("local: %zu, glocal: %zu\n", local.size(), global.size());
    q.submit([&](cl::sycl::handler &cgh) {
      auto res_acc = res_buf.template get_access<cl::sycl::access::mode::read_write>(cgh);
      cgh.parallel_for(cl::sycl::nd_range<1>(global, local),
          cl::sycl::ONEAPI::reduction(res_acc, 0.0f, cl::sycl::ONEAPI::plus<float>()),
          [=](cl::sycl::nd_item<1> it, auto &res_acc) {
            int lid = it.get_global_id(0);
            if (lid >= nelem) return;
            res_acc += in[lid];
          });
    });
  }
  return result[0];
}

int main(int argc, char **argv) {
  cl::sycl::queue q;
  device_info(q.get_device());
  const size_t nelems = 100 * 1024;

  auto *p = reinterpret_cast<float *>(sycl::aligned_alloc_host(4096, sizeof(float) * nelems, q));
  auto *in = reinterpret_cast<float *>(sycl::aligned_alloc_device(4096, sizeof(float) * nelems, q));
  float sum;

  RandomInit(p, nelems);
  q.memcpy(in, p, sizeof(float) * nelems);
  sum = reduce_algorithm<float *>(q, in, nelems);
  return 0;
}

[v-klochkov]

I cannot reproduce the error. Can you please give a) more complete test, b) what OS you used, and how new the compiler is c) how you compile and run the code.

Below is how I tried reproducing it on Linux/Ubuntu using today/fresh compiler:

14:20:47 > clang++ --version
clang version 12.0.0 (https://github.com/intel/llvm 3f9546528d907fd61fa58379865bda85487c43a8)
Target: x86_64-unknown-linux-gnu

14:21:52 > clang++ -fsycl -fsycl-unnamed-lambda wg_size_bug.cpp -o bug.exe
14:22:16 > env SYCL_DEVICE_FILTER=opencl:cpu ./bug.exe 
local: 256, glocal: 102656
finished
14:22:37 > env SYCL_DEVICE_FILTER=opencl:gpu ./bug.exe 
local: 256, glocal: 102656
finished
14:22:43 > env SYCL_DEVICE_FILTER=level_zero:gpu ./bug.exe 
local: 256, glocal: 102656
finished
14:22:49 > 

14:22:49 > cat wg_size_bug.cpp 
#include <CL/sycl.hpp>

// Customer reported a bug for # of bytes >= 409600
const size_t nelem = 409604 / sizeof(float);
//const size_t nelem = 409600 / sizeof(float);

int main() {
  auto max_group_sz = 256;
  auto num_gz = (nelem + max_group_sz - 1)/max_group_sz;
  cl::sycl::range<1> local(max_group_sz);
  cl::sycl::range<1> global(num_gz * max_group_sz);
  std::array<float, 1> result;
  std::vector<float> in_array(nelem);
  result[0] = 0.0f;
  cl::sycl::queue q;
  {
    cl::sycl::buffer<float> res_buf(result.data(), 1);
    cl::sycl::buffer<float> in_buf(in_array.data(), nelem);
    printf("local: %zu, glocal: %zu\n", local.size(), global.size());
    q.submit([&](cl::sycl::handler& cgh){
      auto res_acc = res_buf.template get_access<cl::sycl::access::mode::read_write>(cgh);
      auto in = in_buf.template get_access<cl::sycl::access::mode::read>(cgh);
      cgh.parallel_for(cl::sycl::nd_range<1>(global, local),
                        cl::sycl::ONEAPI::reduction(res_acc, 0.0f, cl::sycl::ONEAPI::plus<float>()),
                        [=](cl::sycl::nd_item<1> it, auto &res_acc) {
                          int lid = it.get_global_id(0);
                          if(lid >= nelem) return;
                          res_acc += in[lid];
                        });    
    });
  }
  std::cout << "finished\n";
}

[Wanzizhu]

Sorry for the unclear test. I already updated the test.Below are platform details.
os: Ubuntu 20.04.2 LTS
compiler: Intel(R) oneAPI DPC++ Compiler 2021.1 (2020.10.0.1113)
platform: Intel(R) Level-Zero
name: Intel(R) Graphics [0x3e98]
vendor: Intel(R) Corporation
driver_version: 1.0.18912

[v-klochkov]

With the complete test and the most recent compiler from intel/llvm I cannot reproduce the error.
It is fixed by this patch: #2859 (Merged to intel/llvm on Dec05'2020). The fix was NOT included into oneAPI 2021.1, but is included into 2021.2

[v-klochkov]

Closing the issue as fixed in 2021.2, and on github.com/intel/llvm by #2859

[Wanzizhu]

Thank you for your reply。