Optimized GEMM & GEMV for Intel platforms

[gongzg]

This PR implemented optimized GEMV and GEMM kernels for Intel Gen Graphics. For the GEMM function, we force the profile to always choose image based GEMM implementation as we found for the real work load, the image based kernels always get better performance. If we use the default tuning mechanism, isaac may choose different implementations which are much slower.

[ptillet]

Hi,

Thanks a lot for the PR. I went through it quite quickly and here are some comments. Let's make sure it eventually gets integrated at just a tiny maintenance cost.

• I don't see the GEMV using any Intel OpenCL extension. The improvements for GEMV should probably be analyzed in more depth to see what's wrong with the current template. Upgrading the current template would massively simplify the PR (no need to check for GEMV) and changes would carry out to row/col-wise reductions and any input shape.
• Isaac should not depend on pthread and OpenCL at compile-time. If there are missing dynamic loads for OpenCL, they should be added in driver/dispatch.cpp. Currently this breaks compilation on NVidia machines, and those that don't have OpenCL installed in the system's default path.
• So the new .json profile for Broadwell seems to suggest that the Image kernel beats Isaac for every single shape. This is a bit surprising, especially considering that this kernel requires a copy from the OpenCL buffers to images. It also means that the improved buffer kernels are never used. I want to check what's going on on my Broadwell machine.

[gongzg]

• GEMV kernel really doesn't use any Intel specific extension. There are some vector load/store consideration to make it more efficient. Not sure whether this change will bring benefit for other platforms.
• You are correct. The pthread dependency should be from clBLAS, I will fix it.
• Actually, this json file is mainly for SKL, but also is good for BDW. And the fact is that the image based GEMM kernel is only faster for some relatively large matrix, for small matrixes, the generic kernel may be faster. But once we enable other kernels during the auto-tuning phase and use that json file. We found during runtime, the isaac may choose the generic kernel rather than the image based kernel for some even relatively large matrix and thus get much worse performance than we force it to image based gemm kernel.

[ptillet]

So I could try the branch today at work, and saw indeed some good 20-30% improvements in some cases for GEMM. Good job :)
However, I get some crashes in the tests/benchmarks, and substantial decrease in performance for some shapes.

For GEMV, there is only one single case in the benchmarks where the Intel kernel provides gains:
(M, N)=(1024,60000), A_trans = 1
This corresponds to a case where the matrix is too large to fit in the eDRAM. Which is a bit strange since the two algorithms seem to be the same. I'll investigate this in more details.

[gongzg]

@ptillet Could you let me know how to reproduce the crashes? Which platform are you working on? @listenlink will take care of those crashes.

As those GEMV kernels, we found Intel GEMV wins for many cases. @listenlink could you share some cases which GEMV wins in both BDW and SKL. Thanks.

BTW, we found the master branch isaac cause random fail with clcaffe. If you build clcaffe with isaac and enable the intel spatial engine then run the clcaffe's test suite. It may crash some time or get incorrect result. But if we run those crashing or fail case directly, it could pass. Do you have time to look into this issue as well?

[ptillet]

For GEMM:
Can you see any HW reason why the image kernels would be faster than the buffer kernels? On NVidia Hardware, the compiler will use the read-only texture-cache when the pointers are restrict. I would be surprised if there was no instruction in the ISA to use the texture cache while fetching from global memory.
Also, there are indeed some problems with the prediction mechanism, hence the branch labels-cache branch which obviates the prediction mechanism. It is not in production yet, but it's definitely useful for testing new kernels.

For GEMV:
ISAAC's GEMV kernel also uses vector loads and vector stores. Keep also in mind that the .json file on isaac-master was generated for my platform, which is a BDW GT3e with 2x DDR3 sticks. If you have SKL , DDR4 or quad-channel, the results could be very different. I can't benchmark Skylake unfortunately; I have no machine.

For the crashes:
I get crashes at the end of test-blas-2 and test-blas-3, with "The supplied parameters for this template are invalid". I will get the error code tomorrow.
./bench/bench-blas hangs on GEMM. I will try with hangcheck disabled tomorrow, but it's a bad sign since no kernel should take long enough to activate the watchdog.
Do you have the guarantee that the subgroup size will indeed be what you assume it to be in the kernels?

For the bug:
Can you open another issue for the error with caffe? A commit SHA1 that doesn't cause the error would also be useful. The test-suite is pretty strict for the kernels now, but after all it is not impossible that things get messed up with queues at some point.

[listenlink]

Hi,
For the crashes,
that's because the hardware support the double datatype test while the broadwell.json file don't have the "float64" profiles, the crashes will also occur on the original master branch.

For GEMV perf,
because the original benchmark is vs. the commit on Jan last year, when rebase the latest code, i also find it will not get benefit against the current reduce_2d kernel.

[gongzg]

@ptillet
For GEMM:
There is indeed texture only cache for image for Intel Gen graphics. And the tiling mode of image works very well with the sub group blocking read extension which could get better IO efficiency than the buffer based implementation. This is the major reason why image is better.

For GEMV, @listenlink already explained that we are comparing with the orignal implemtnation. But as you may know, to open source our contribution, we came from a long way and you already did some nice improvement thus we can't see too much difference now. What's your suggestion for this case?

[ptillet]

So at this point it's probably useful that I summarize my comments on the PR:

1 - There shouldn't be any static dependency for Isaac -- not even OpenCL. I use libdl to load OpenCL and/or CUDA at runtime (depending on what's installed), see driver/dispatch.cpp for details. The pthread dependency is also not necessary.

2 - Actually I have not touched GEMV kernels in a very long time, although I may have re-ran the auto-tuner. At this point it seems like the maintenance troubles associated with intel-specific gemv kernels are not worth it. However I think your expertise could be useful to improve the current GEMV template if you wish, but it already uses vectorized load/store. Is there any way to have global atomic add on GEN ? (i.e., is it supported by the hardware and is there any way to access it via OpenCL without a repeated use of atomic_inc).

3 - I think the GEMM kernels are very valuable, as they yield some good improvements for some input shapes. Here are my worries, though:

• Is there any guarantee that the sub-group size will always be set by the compiler to the value required by the kernel?
• I think it's time for Isaac to find a way to differentiate between different GEN architectures. It seems like the best profiles for BDW and SKL are different.
• the file generation/gemm.cpp is becoming quite big, perhaps having intel-specific definitions somewhere else (gemm/intel.cpp) could be good.

Overall, I think the PR could be simplified to just a few file:
runtime/profiles.cpp
jit/gemm.h
jit/gemm/intel.cpp (new file)
driver/dispatch.h|cpp
driver/device.h|cpp (for checking BDW vs SKL)
python/src/bind/kernels.cpp

I think that changes in the auto-tuner should not be needed, but I may be wrong

[gongzg]

@ptillet Thanks for the comments, we are working on that. One question which I forgot to anwser,
"Do you have the guarantee that the subgroup size will indeed be what you assume it to be in the kernels?"

There is a sub group extension function attribution to make sure the kernel will be built with specified sub group size:

__attribute__((intel_reqd_sub_group_size(16)))
__kernel void foo() {
...
}