Enable locality test (#247)

Enables the locality test by using an IGC intrinsic instead of an OpenCL
function for `get_global_size`. IGC does not appear to support the
OpenCL function `get_global_size`, and SPIRV builtins are not currently
an option. This PR uses intel/llvm#12383 and then simply removes the XPU skip from `test_locality`.

Closes #154

python/test/unit/language/test_core.py

@@ -1884,8 +1884,6 @@ def histogram_kernel(x_ptr, z_ptr, M: tl.constexpr, N: tl.constexpr):
 @pytest.mark.parametrize("N", [512, 1024, 2048])
 @pytest.mark.parametrize("num_pid_n", [2, 4])
 def test_locality(op, BLOCK_N, N, num_pid_n, device):
-    if is_xpu(device):
-        pytest.skip("FIXME: Does run correctly on XPU")
 
     @triton.jit
     def kernel(X, Y, N, BLOCK_M: tl.constexpr, BLOCK_N: tl.constexpr):