intel / intel/llvm

sycl_ext_intel_grf_size large grf not working for sycl_explicit_simd kernel

Open
#12,704 3 comments 0 reactions 0 assignees View on GitHub
bug confirmed esimd
Dominant language
LLVM
Stars
1.5k
Forks
854
Avg merge
3d 17h
Merged PRs (30d)
137

Description

**Describe the bug**
Setting the kernel_properties to include grf_size<256> for an ESIMD kernel does not change the register file size.

**To Reproduce**
```c++
#include
#include
#include
int main() {
printf("SYCL_EXT_INTEL_GRF_SIZE: %d\n", SYCL_EXT_INTEL_GRF_SIZE);
sycl::queue q(sycl::gpu_selector_v);
sycl::ext::oneapi::experimental::properties kernel_properties{sycl::ext::intel::experimental::grf_size<256>};
q.submit([&](sycl::handler &h) {
h.parallel_for(
sycl::nd_range{
sycl::range<3>(1, 1, 1),
sycl::range<3>(1, 1, 1)
},
kernel_properties,
[=](sycl::nd_item<3> item)
[[intel::sycl_explicit_simd]]
{
}
);
}).wait();
return 0;
}
```

Save file as `no_large_grf_reproducer.cpp`

Export environment variables to keep assembly code
```bash
export IGC_ShaderDumpEnable=1
export IGC_DumpToCustomDir=assembly
```

Compile with
```bash
icpx -std=c++20 -fsycl -fsycl-targets=intel_gpu_acm_g10 -Xsycl-target-frontend -O3 -o no_large_grf_reproducer src/no_large_grf_reproducer.cpp -lsycl
```

Look at the assembly code (.asm file in the `assembly` directory) and note the .thread_config numGRF value. This should be 256. This value is correctly set when not using [[intel::sycl_explicit_simd]]. It appears to also not work correctly when targeting intel_gpu_pvc.

**Environment (please complete the following information):**

- OS: Ubuntu 23.10, kernel 6.5.0-17-generic
- Target device and vendor: Intel A770 GPU
- DPC++ version: icpx --version: Intel(R) oneAPI DPC++/C++ Compiler 2024.0.2 (2024.0.2.20231213)
- Dependencies version: I'm not sure what this is referring to.

Contributor guide

Open the contributing guide

Assessment

This issue has not been assessed yet.

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.