THU-DSP-LAB / ventus-gpgpu

GPGPU processor supporting RISCV-V extension, developed with Chisel HDL
Mulan Permissive Software License, Version 2
633 stars 73 forks source link

Cannot use the released version "VENTUS-2.0.2" of software part #37

Closed fxzjshm closed 9 months ago

fxzjshm commented 9 months ago

I've downloaded software part from https://opengpgpu.org.cn/html/web/project/release/index.html, while trying to run some example OpenCL program these problems are encountered:

1. Path used by pocl is absolute:

chmod: cannot access '/home/kl/kl_ws/tsinghua_ws/ventus_ws/llvm-project/install/bin/../../assemble.sh': No such file or directory
sh: 1: /home/kl/kl_ws/tsinghua_ws/ventus_ws/llvm-project/install/bin/../../assemble.sh: not found

which relates to

      std::string assembler_path = CLANG;
    assembler_path = assembler_path.substr(0,assembler_path.length()-6);
      assembler_path += "/../../assemble.sh";

https://github.com/THU-DSP-LAB/pocl/blob/414145738992eccdb830c55d4af0501917326ae3/lib/CL/devices/ventus/pocl_ventus.cc#L689

which seems caused by https://github.com/THU-DSP-LAB/pocl/blob/414145738992eccdb830c55d4af0501917326ae3/config.h.in.cmake#L160 then cannot be used without manually putting the script there. Moreover, assemble.sh is not included in the package.

2. object.riscv not found

nm: 'object.riscv': No such file
llvm-objdump: error: 'object.riscv': No such file or directory

not sure what the situation is, could you please explain how to get object.riscv?

3. poclcc gives "still have references to IRs - can't release LLVM context !"

➜  /tmp cat vector_add.cl
 // ACL kernel for adding two input vectors
__kernel void vectorAdd(__global float *x, 
                        __global float *y, 
                        __global float *z)
{
    // get index of the work item
    int index = get_global_id(0);

    // add the vector elements
    z[index] = x[index] + y[index];
}
➜  /tmp poclcc vector_add.cl
[INFO]: [HW DRIVER] in [FILE] ventus.cpp,[LINE]25,[fn] vt_dev_open: vt_dev_open : hello world from ventus.cpp
spike device initialize: allocating local memory: to allocate at 0x70000000 with 268435456 bytes 
spike device initialize: allocating pc source memory: to allocate at 0x80000000 with 268435456 bytes 
### options: -DPOCL_DEVICE_ADDRESS_BITS=32 -D__USE_CLANG_OPENCL_C_H -xcl -Dinline= -I. -cl-kernel-arg-info  -D__ENDIAN_LITTLE__=1 -DCL_DEVICE_MAX_GLOBAL_VARIABLE_SIZE=0 -D__OPENCL_VERSION__=200 -cl-std=CL2.0 -D__OPENCL_C_VERSION__=200 -Dcl_khr_fp64=1 -D__opencl_c_generic_address_space=1 -D__opencl_c_named_address_space_builtins=1 -cl-ext=-all,+cl_khr_fp64,+__opencl_c_generic_address_space,+__opencl_c_named_address_space_builtins -fno-builtin -triple=riscv32 -target-cpu ventus-gpgpu user_options: 
### Triple: riscv32, CPU: ventus-gpgpu
still have references to IRs - can't release LLVM context !
[1]    62536 IOT instruction (core dumped)  poclcc vector_add.cl

also not sure what's happening...

Jules-Kong commented 9 months ago

@fxzjshm Thanks for your feedback. We have fixed the bugs in the release version.

You can re-download it from https://opengpgpu.org.cn/html/web/project/release/index.html, and follow the steps in https://opengpgpu.org.cn/html/web/project/getstarted/index.html to use the toolchain.
Please note that the new environment variable "VENTUS_INSTALL_PREFIX" has been added to the current version.

If you encounter new problems, you are welcome to continue to provide feedback. Thank you!

fxzjshm commented 9 months ago

Confirm fixed by the new release, thanks!