intel / intel/llvm

Read Image Crash in Kernel Execution with Level Zero Backend

Open
#13,084 1 comment 0 reactions 1 assignee Claimed by @maksimsab View on GitHub
bug confirmed
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

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.