How to use extensions of RISC-V in OpenCL code？

[TJU-PanYizhe]

Version

4.0.0

What is your question or problem?

I compiled oneapi-construction-kit on my RISC-V CPU hardware using command below:

cmake -GNinja \
   -Bbuild-riscv \
   -DCA_LLVM_INSTALL_DIR=/usr/lib/llvm-18 \
   -DCA_ENABLE_HOST_IMAGE_SUPPORT=OFF \
   -DCA_CL_ENABLE_ICD_LOADER=ON \
   -DCMAKE_C_COMPILER=clang-18 \
   -DCMAKE_CXX_COMPILER=clang++-18
ninja -C build-riscv install -j32

After set ICD file under /etc/OpenCL/vendors/, clinfo and sycl-ls can both find the RISC-V CPU as an OpenCL device.

Now this RISC-V CPU can ran OpenCL code. Here is one of the OpenCL kernel I tried:

__kernel void hello_kernel(__global const float *a,
	__global const float *b,
	__global float *result)
{
	int gid = get_global_id(0);
 
	result[gid] = a[gid] + b[gid];
}

I compiled the kernel above with ocl and disassembled the output with llvm-objdump using the command:

clc --strip-binary-header -o HelloWorld.o HelloWorld.cl -cl-wfv=never -d riscv64
llvm-objdump --disassemble --triple=riscv64 --mattr="v,c,zbc" HelloWorld.o > HelloWorld.txt

There aren't any RISC-V V Extention Assembly Code in HelloWorld.txt though RVV is supported on my RISC-V CPU.

• How to use any of the RISC-V Extention in OpenCL when run OpenCL code on a RISC-V CPU?
• Is it correct that OpenCL code is translated into LLVM IR and processed by Clang to generate the binary in this scenario that RISC-V CPU is regarded as the OpenCL device?

[hvdijk]

Hi,

• oneAPI Construction Kit by default intentionally does not make use of extensions that are not required to be supported. We follow the hardware baseline agreed between Linux distributions of RV64GC, without V, to ensure that device code will work even if moving to different hardware. We do have options to override that though, which in general allows RVV. I cannot give a concrete example right now but if you have trouble getting it to work, I should be able to go in more detail and show concrete results next week.
• What you write is largely correct. The OpenCL code is translated into LLVM IR. That translation is the part that is done using libclang. Once it has been converted to LLVM IR, further processing is done by oneAPI Construction Kit using libLLVM, not using libclang.

[TJU-PanYizhe]

Thanks for your reply!
Looking forward to your concrete example and results!

We do have options to override that though, which in general allows RVV. I cannot give a concrete example right now but if you have trouble getting it to work, I should be able to go in more detail and show concrete results next week.

[hvdijk]

Having been able to test now: the options that I referenced work without anything further. With the CMake option CA_HOST_TARGET_RISCV64_FEATURES=+v, or in debug builds of oneAPI Construction Kit, with the environment variable CA_HOST_TARGET_FEATURES=+v, RVV instructions will be generated.

With the -cl-wfv=never you specified, vector operations are only generated by LLVM's own vectorization passes. Without -cl-wfv=never, oneAPI Construction Kit's own vectorizer will result in different vectorization. But in both cases, you should see RVV instructions.