PortBLAS: core aborted on build with AdaptiveCPP on Intel GPUs

[chsasank]

I am building PortBLAS and benchmaking it for diff GPUs - see #1417. On Intel's the build crashes.

$ echo $CC $CXX $ACPP_TARGETS
clang-16 acpp generic

cmake -GNinja -DTUNING_TARGET=INTEL_GPU -DCMAKE_BUILD_TYPE=Release       -DSYCL_COMPILER=adaptivecpp -DACPP_TARGETS=$ACPP_TARGETS       -DBLAS_ENABLE_BENCHMARK=ON ..

(I removed lines from SYCL.cmake that disables benchmarks for adaptivecpp in PortBLAS)

Run benchmark:

$ cat << EOF > params.csv
n,n,1024,1024,1024,1,0
n,n,2048,2048,2048,1,0
n,n,4096,4096,4096,1,0
EOF
./benchmark/portblas/bench_gemm --csv-param params.csv --benchmark_out=../results.json \
    --benchmark_out_format=json --benchmark_format=console

Result:

Device vendor: pci:32902
Device name: Intel(R) Arc(TM) A770 Graphics
Device type: gpu
2024-03-25T17:22:51+05:30
Running ./benchmark/portblas/bench_gemm
Run on (12 X 4400 MHz CPU s)
CPU Caches:
  L1 Data 48 KiB (x6)
  L1 Instruction 32 KiB (x6)
  L2 Unified 1280 KiB (x6)
  L3 Unified 18432 KiB (x1)
Load Average: 6.80, 7.69, 6.19
***WARNING*** CPU scaling is enabled, the benchmark real time measurements may be noisy and will incur extra overhead.
[AdaptiveCpp Warning] kernel_cache: This application run has resulted in new binaries being JIT-compiled. This indicates that the runtime optimization process has not yet reached peak performance. You may want to run the application again until this warning no longer appears to achieve optimal performance.
terminate called after throwing an instance of 'hipsycl::sycl::exception'
  what():  Operation not profiled: No execution start timestamp available
Aborted (core dumped)

[chsasank]

Can confirm that acpp-info can detect the Intel GPU:

$ acpp-info -l
=================Backend information===================
Loaded backend 0: OpenMP
  Found device: hipSYCL OpenMP host device
Loaded backend 1: Level Zero
  Found device: Intel(R) Arc(TM) A770 Graphics

[chsasank]

Can confirm that things work fine with dpcpp compiler on the same system/repo.

[illuhad]

what(): Operation not profiled: No execution start timestamp available

Likely the library attempts to use the SYCL queue profiling functionality (which is an optional feature according to the SYCL specification). The OpenCL and Level Zero backends do not currently implement this functionality. This is currently a known limitation, and one that is in line with the SYCL specification.

So I think we can close this here.

[chsasank]

Ah that's why portBLAS folks disabled benchmarking on AdaptiveCPP. Rest of the functionality should work fine?

[illuhad]

I think so. Compiler support in the generic SSCP compiler is on par between all the backends, and on the runtime side, the only additional limitation that OpenCL and L0 backends have is 2D and 3D data transfers due to limitations in the OpenCL/L0 USM API. I don't think that portBLAS uses that though.