parallel_for_work_item does not work inside lambdas
- Dominant language
- LLVM
- Stars
- 1.5k
- Forks
- 854
- Avg merge
- 3d 17h
- Merged PRs (30d)
- 137
Description
I have a simple test which fills buffer with local ids. If parallel_for_work_item is invoked directly everything is fine. But if it is located inside lambda, only the first element of range is executed, even if range is set explicitly.
The test:
```
#include
void FillLambda(cl::sycl::queue &queue, std::vector &vec,
size_t workGroupSize) {
using namespace cl::sycl;
auto buf = buffer{vec};
auto nWorkGroups = vec.size() / workGroupSize;
queue.submit([&](handler &h) {
auto global = buf.template get_access(h);
auto lambda = [=](group<1> g) {
g.parallel_for_work_item([=](h_item<1> it) {
global[it.get_global_id()[0]] = it.get_local_id()[0];
});
};
h.parallel_for_work_group(
range<1>{nWorkGroups}, range<1>{workGroupSize},
[=](group<1> g) { lambda(g); });
});
queue.wait();
buf.template get_access();
}
void FillWithoutLambda(cl::sycl::queue &queue, std::vector &vec,
size_t workGroupSize) {
using namespace cl::sycl;
auto buf = buffer{vec};
auto nWorkGroups = vec.size() / workGroupSize;
queue.submit([&](handler &h) {
auto global = buf.template get_access(h);
h.parallel_for_work_group(
range<1>{nWorkGroups}, range<1>{workGroupSize}, [=](group<1> g) {
g.parallel_for_work_item([=](h_item<1> it) {
global[it.get_global_id()[0]] = it.get_local_id()[0];
});
});
});
queue.wait();
buf.template get_access();
}
int main() {
auto size = 16;
auto workGroupSize = 4;
auto nWorkGroups = size / workGroupSize;
auto emptyVec = std::vector(size);
auto lambdaVec = emptyVec;
auto withoutLambdaVec = emptyVec;
auto GPUSelector = cl::sycl::gpu_selector{};
auto queue = cl::sycl::queue{GPUSelector};
FillLambda(queue, lambdaVec, workGroupSize);
FillWithoutLambda(queue, withoutLambdaVec, workGroupSize);
std::sort(lambdaVec.begin(), lambdaVec.end());
std::sort(withoutLambdaVec.begin(), withoutLambdaVec.end());
for (auto i = 0; i != size; ++i)
std::cout << "Expected: " << i / nWorkGroups
<< "; Computed without lambda: " << withoutLambdaVec[i]
<< "; Computed with lambda: " << lambdaVec[i] << std::endl;
}
```
`clang++ -fsycl LambdaTest.cpp`
**Output:**
```
Expected: 0; Computed without lambda: 0; Computed with lambda: 0
Expected: 0; Computed without lambda: 0; Computed with lambda: 0
Expected: 0; Computed without lambda: 0; Computed with lambda: 0
Expected: 0; Computed without lambda: 0; Computed with lambda: 0
Expected: 1; Computed without lambda: 1; Computed with lambda: 0
Expected: 1; Computed without lambda: 1; Computed with lambda: 0
Expected: 1; Computed without lambda: 1; Computed with lambda: 0
Expected: 1; Computed without lambda: 1; Computed with lambda: 0
Expected: 2; Computed without lambda: 2; Computed with lambda: 0
Expected: 2; Computed without lambda: 2; Computed with lambda: 0
Expected: 2; Computed without lambda: 2; Computed with lambda: 0
Expected: 2; Computed without lambda: 2; Computed with lambda: 0
Expected: 3; Computed without lambda: 3; Computed with lambda: 0
Expected: 3; Computed without lambda: 3; Computed with lambda: 0
Expected: 3; Computed without lambda: 3; Computed with lambda: 0
Expected: 3; Computed without lambda: 3; Computed with lambda: 0
```
Contributor guide
Assessment
This issue has not been assessed yet.