[Intel GPU] how to know the SIMD width compiled from dpc++ code

[guoyejun]

I'm working on Intel GPU with:

queue q(gpu_selector{}, propList);

And my dpcpp code is normal, looks like:

event add_data(queue q, float *input0, float *input1, float *output)
{
    auto event = q.submit([&](sycl::handler &h) {
        h.parallel_for<class ker_1>(len, [=](cl::sycl::item<1> item) {
            int idx = item.get_id(0);
            output[idx] = input0[idx] + input1[idx];
        });
    });
    return event;
}

I just use the simple way to compile the code:

dpcpp mycode.cpp

My question is, what SIMD width will be used, for example, SIMD-8, SIMD-16 or SIMD-32 for the GPU instruction? How could I query it? thanks.

[bader]

My question is, what SIMD width will be used, for example, SIMD-8, SIMD-16 or SIMD-32 for the GPU instruction? How could I query it?

@PawelJurek, which API exposes this kernel property?

[PawelJurek]

@bader , please correct me if I'm wrong, but it seems that this value can be queried via SYCL_INTEL_sub_group extension

[bader]

You seem to be right. Thanks for pointing this out!

[guoyejun]

thanks, got it.

btw, is it possible to dump the GPU asm, and so we can also know the SIMD width from the kernel asm code.

[PawelJurek]

Yes, this is available on Linux drivers via IGC_ShaderDumpEnable flag. Instructions are here: https://github.com/intel/intel-graphics-compiler/blob/master/documentation/configuration_flags.md

[guoyejun]

thanks!

[guoyejun]

not know how to change the status to be answered

[bader]

This should help - https://docs.github.com/en/discussions/collaborating-with-your-community-using-discussions/participating-in-a-discussion#marking-a-comment-as-an-answer