KhronosGroup / KhronosGroup/SYCL-Docs

Iteration Size Limitation?

Open
#386 1 comment 0 reactions 0 assignees View on GitHub
enhancement
Dominant language
JavaScript
Stars
158
Forks
80
Avg merge
7d 6h
Merged PRs (30d)
5

Description

# Context

No implementation with a GPU backend allow: `Q.parallel_for(std::numerical_limits::max(), ...); `,
when the same code works fine with OpenMP:
```
#pragma omp target team distribute parallel for
for (size_t i = 0; i < std::numerical_limits::max(), i++)
{}
```

The SYCL spec said that this code is valid and should be executed. We don't have a concept of `maximum iteration space`.

## TLDR

I propose that:
- We add a CTS test for `Q.parallel_for(std::numerical_limits::max(), ...); `
- We add a `max_global_work_groups` to specify the number of maximum work-group into the `nd_range`. [Intel Extension](https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/experimental/sycl_ext_oneapi_max_work_group_query.md).

# Too Much Info and Text

## SYCL

The current SYCL behavior can explained as in sycl is natural to implement `parrallel_for` as a direct kernel submission using the native backend. And some GPU native backends have a restriction on the number of work-item that can be submitted. For example, both L0 and CUDA have a maximum [WorkGrounpCount, and WorkGroupSize](https://spec.oneapi.io/level-zero/latest/core/api.html#ze-device-compute-properties-t). OpenCL have also a context of maxinun`global_work_size`:
```
global_work_size

Points to an array of work_dim unsigned values that describe the number of global work-items in work_dim dimensions that will execute the kernel function. The total number of global work-items is computed as global_work_size[0] *...* global_work_size[work_dim - 1].

The values specified in global_work_size cannot exceed the range given by the sizeof(size_t) for the device on which the kernel execution will be enqueued. The sizeof(size_t) for a device can be determined using CL_DEVICE_ADDRESS_BITS in the table of OpenCL Device Queries for [clGetDeviceInfo](https://registry.khronos.org/OpenCL/sdk/1.0/docs/man/xhtml/clGetDeviceInfo.html). If, for example, CL_DEVICE_ADDRESS_BITS = 32, i.e. the device uses a 32-bit address space, size_t is a 32-bit unsigned integer and global_work_size values must be in the range 1 .. 2^32 - 1. Values outside this range return a CL_OUT_OF_RESOURCES error.
```

## OpenMP

In OpenMP, the "traditional" code gen will add an internal loop inside the kernel submitted so the `iteration space` are independent of the number of work-item. Just to be clear, the code previous user code will look like something like:

```
__global__ void foo(){
int gindex = threadIdx.x + blockIdx.x * blockDim.x;
// Some trip-count + offset comoutation
for (size_t k=N; k=M )
__user_kernel()
}

int main() {
foo<<>>();
return 0;
}
```

# Problem and Solution

So in short, the current implementation cannot submit a kernel with a large iteration space. 2 Main Solutions:

1/ Implementers should just fix their buggy implementation. We should add a test in the CTF
2/ We should add a concept of `maximum group count` to SYCL. Similar to this [Intel Extension](https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/experimental/sycl_ext_oneapi_max_work_group_query.md).

## 1/ Not your problem

I like 1/, as I like to think of SYCL as a High-Level Abstraction model. As in OpenMP, people should not be concerned about those low-level trivial detail. But, I heard that people like performance.

## 2/ Handling large problem sizes add overhead

The argument in favor of 2/, is that some native-backends have limitations. Any effort to avoid that limitation will have some overhead. We want SYCL to be a thin-abstraction 0 overhead layer. Hence our goal as SYCL is just to standardize general queries. Then people can query for those limitations and deal with them appropriately.

## 3/ Middle ground

We can allow abbriraty iteration space for `range` and add a query for the maximum number of work-group for`nd_range`. This sound like a reasonable tradeoff.

Contributor guide

No contributing guide indexed for this repository

Research direction

No implementation file is named; begin with the SYCL specification sections for queue::parallel_for, range, and nd_range, then inspect CTS coverage for large iteration spaces. Compare the proposed implementation fix with the cited Intel maximum-work-group extension. Done means an agreed specification direction and a CTS test covering the size_t::max case.

Written by the indexing model from the issue text.

Assessment

Tech stack
cpp
Domain
backend-api-design, documentation
Issue type
Feature
Difficulty
5/5
Estimated time
Over a week
Activity status
Stale
Clarity
Needs clarification
Newbie friendliness
25/100

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.