Read Image Crash in Kernel Execution with Level Zero Backend
- Dominant language
- LLVM
- Stars
- 1.5k
- Forks
- 854
- Avg merge
- 3d 17h
- Merged PRs (30d)
- 137
Description
### Describe the bug
When reading a unrom_int8 or unsigned_int8 image in sycl, it always crash in kernel.
I've tried w/ and w/o sampler. No matter data port or sampler, it both fails
### To reproduce
```
#include
#include
#include
using namespace sycl;
class Scaling
{
public:
Scaling(
accessor inputImg,
sampler smplr,
accessor debugInfo) :
m_inputImage(inputImg),
m_sampler(smplr),
m_debugInfo(debugInfo)
{
}
void operator()(sycl::item<2> idx) const
{
float2 coord(0.5, 1);
float4 pixel = m_inputImage.read(coord,m_sampler);
if (idx[0] == 1 && idx[1] == 1)
{
m_debugInfo[0] = pixel[2];
m_debugInfo[1] = pixel[3];
}
m_debugInfo[2] = 9;
}
accessor m_inputImage;
sampler m_sampler;
accessor m_debugInfo;
};
int main() {
std::cout << "start process" << std::endl;
std::vector debug_info(100);
//prepare input image data
uint8_t tmpIn[16];
for (int i = 0; i < 16; i++) {
tmpIn[i] = i * 4;
}
tmpIn[12] = 8;
tmpIn[13] = 64;
tmpIn[14] = 128;
tmpIn[15] = 255;
for (int i = 0; i < 4; ++i)
{
std::cout << "input: " << int(tmpIn[i * 4 + 0]) << " " << int(tmpIn[i * 4 + 1]) << " " << int(tmpIn[i * 4 + 2]) << " " << int(tmpIn[i * 4 + 3]) << std::endl;
}
try {
image<2> srcImage(tmpIn, image_channel_order::rgba, image_channel_type::unorm_int8, range<2>(2, 2));
buffer buffer(debug_info.data(), debug_info.size());
std::cout << "image load done" << std::endl;
queue q;
std::cout << "Running on "
<< q.get_device().get_info()
<< "\n";
q.submit([&](handler& cgh) {
auto image_acc = srcImage.get_access(cgh);
auto debugBuffer = buffer.get_access(cgh);
sampler smplr(coordinate_normalization_mode::normalized, addressing_mode::clamp, filtering_mode::linear);
std::cout << "accessor ready" << std::endl;
cgh.parallel_for(range<2>(2, 2), Scaling(image_acc, smplr, debugBuffer));
}).wait();
}
catch (exception e)
{
std::cout << "SYCL exception caught: " << e.what() << std::endl;
}
std::cout << "gpu process done" << std::endl;
for (int i = 0; i < 4; ++i)
{
std::cout << "input: " << int(tmpIn[i * 4 + 0]) << " " << int(tmpIn[i * 4 + 1]) << " " << int(tmpIn[i * 4 + 2]) << " " << int(tmpIn[i * 4 + 3]) << std::endl;
}
std::cout << "debugInfo: " << (debug_info[0]) << " " << (debug_info[1]) << " " << (debug_info[2]) << std::endl;
std::cout << "all done" << std::endl;
return 0;
}
```
**How to compile**
goto onAPI root folder
open cmd in the folder
`./setvars.bat`
`.\compiler\2024.0\bin\icpx.exe -fsycl -fsycl-targets=spir64_gen -Xsycl-target-backend "-device dg2" readImage.cpp -o readImage.exe`
**How to run**
goto onAPI root folder
open cmd in the folder
`./setvars.bat`
`readImage.exe`
**Expected Result**
It is expected that the console output the image pixel. But it outputs all 0.
### Environment
OS : Windows
Target Device : Intel GPU (DG2 or MTL)
Intel(R) oneAPI DPC++/C++ Compiler 2024.0.2 (2024.0.2.20231213)
Target: x86_64-pc-windows-msvc
Thread model: posix
[ext_oneapi_level_zero:gpu:0] Intel(R) Level-Zero, Intel(R) Arc(TM) A370M Graphics 1.3 [1.3.0]
Platform [#4]:
Version : 1.3
Name : Intel(R) Level-Zero
Vendor : Intel(R) Corporation
Devices : 1
Device [#0]:
Type : gpu
Version : 1.3
Name : Intel(R) Arc(TM) A370M Graphics
Vendor : Intel(R) Corporation
Driver : 1.3.0
Aspects : gpu fp16 online_compiler online_linker queue_profiling usm_device_allocations usm_host_allocations usm_shared_allocations ext_intel_pci_address ext_intel_gpu_eu_count ext_intel_gpu_eu_simd_width ext_intel_gpu_slices ext_intel_gpu_subslices_per_slice ext_intel_gpu_eu_count_per_subslice atomic64 ext_intel_device_info_uuid ext_intel_gpu_hw_threads_per_eu ext_intel_device_id ext_intel_memory_clock_rate ext_intel_memory_bus_width ext_intel_legacy_image
info::device::sub_group_sizes: 8 16 32
default_selector() : gpu, Intel(R) Level-Zero, Intel(R) Arc(TM) A370M Graphics 1.3 [1.3.0]
accelerator_selector() : acc, Intel(R) FPGA Emulation Platform for OpenCL(TM), Intel(R) FPGA Emulation Device OpenCL 1.2 [2023.16.12.0.12_195853.xmain-hotfix]
cpu_selector() : cpu, Intel(R) OpenCL, 11th Gen Intel(R) Core(TM) i9-11900H @ 2.50GHz OpenCL 3.0 (Build 0) [2023.16.12.0.12_195853.xmain-hotfix]
gpu_selector() : gpu, Intel(R) Level-Zero, Intel(R) Arc(TM) A370M Graphics 1.3 [1.3.0]
custom_selector(gpu) : gpu, Intel(R) Level-Zero, Intel(R) Arc(TM) A370M Graphics 1.3 [1.3.0]
custom_selector(cpu) : cpu, Intel(R) OpenCL, 11th Gen Intel(R) Core(TM) i9-11900H @ 2.50GHz OpenCL 3.0 (Build 0) [2023.16.12.0.12_195853.xmain-hotfix]
custom_selector(acc) : acc, Intel(R) FPGA Emulation Platform for OpenCL(TM), Intel(R) FPGA Emulation Device OpenCL 1.2 [2023.16.12.0.12_195853.xmain-hotfix]
### Additional context
_No response_
Contributor guide
Assessment
This issue has not been assessed yet.