企业官网建设流程全解析

电商网站会员体制怎么做,网站功能方案,小程序开发哪家公司好,沈阳seo推广公司目录 前言 一、POCL仿例列表及功能框图 二、POCL仿例环境 2.1 APP使用的驱动层函数不同 2.2 APP Makefile不同 2.2.1 编译应用层main.cc 2.2.2 链接APP应用程序 2.2.3 执行应用程序 三、POCL在Vortex中的功能 总结 前言 本篇内容继承上一篇Vortex RTLSIM仿真环…目录前言一、POCL仿例列表及功能框图二、POCL仿例环境2.1 APP使用的驱动层函数不同2.2 APP Makefile不同2.2.1 编译应用层main.cc2.2.2 链接APP应用程序2.2.3 执行应用程序三、POCL在Vortex中的功能总结前言本篇内容继承上一篇Vortex RTLSIM仿真环境简介着重描述跑POCL(Portable OpenCL)仿例在环境上的不同点。本系列探索Vortex开源GPGPURISC-V SIMT架构一、POCL仿例列表及功能框图Vortex POCL仿例位于$VORTEX_HOME/tests/opencl共有20个例子。功能框图如下。二、POCL仿例环境2.1 APP使用的驱动层函数不同POCL仿例使用了不同的头文件如下的CL/opencl.h。在紫色粗体字的每个步骤中都使用了cl*的函数这个是POCL(Portable OpenCL)封装的驱动层函数。根据函数的近似功能列出对比表格。之所以是近似是因为功能类似传参却差别很大。POCL仿例VX仿例clGetPlatformIDs/clGetDeviceIDsvx_dev_openclCreateBuffervx_mem_allocclCreateProgramWithSource/clBuildProgram/clCreateKernelMakefile编译RISC-V程序vx_upload_kernel_fileclSetKernelArgvx_upload_bytesclCreateCommandQueueN/AclEnqueueWriteBuffervx_copy_to_devclEnqueueNDRangeKernelvx_startclFinishvx_ready_waitclEnqueueReadBuffervx_copy_from_devcl_device_id device_id NULL;uint8_t *kernel_bin NULL;#define KERNEL_NAME vecaddint main (int argc, char **argv) {............cl_platform_id platform_id;size_t kernel_size;//Getting platform and device informationCL_CHECK(clGetPlatformIDs(1, platform_id, NULL));CL_CHECK(clGetDeviceIDs(platform_id, CL_DEVICE_TYPE_DEFAULT, 1, device_id, NULL));printf(Create context\n);context CL_CHECK2(clCreateContext(NULL, 1, device_id, NULL, NULL, _err));printf(Allocate device buffers\n);size_t nbytes size * sizeof(TYPE);a_memobj CL_CHECK2(clCreateBuffer(context, CL_MEM_READ_ONLY, nbytes, NULL, _err));b_memobj CL_CHECK2(clCreateBuffer(context, CL_MEM_READ_ONLY, nbytes, NULL, _err));c_memobj CL_CHECK2(clCreateBuffer(context, CL_MEM_WRITE_ONLY, nbytes, NULL, _err));printf(Create program from kernel source\n);if (0 !read_kernel_file(kernel.cl, kernel_bin, kernel_size))return -1;program CL_CHECK2(clCreateProgramWithSource(context, 1, (const char**)kernel_bin, kernel_size, _err));//Build programCL_CHECK(clBuildProgram(program, 1, device_id, NULL, NULL, NULL));//Create kernelkernel CL_CHECK2(clCreateKernel(program, KERNEL_NAME, _err));//Set kernel argumentsCL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *)a_memobj));CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *)b_memobj));CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), (void *)c_memobj));//Allocate memories for input arrays and output arrays.std::vectorTYPE h_a(size);std::vectorTYPE h_b(size);std::vectorTYPE h_c(size);//Generate input valuesfor (uint32_t i 0; i size; i) {h_a[i] ComparatorTYPE::generate();h_b[i] ComparatorTYPE::generate();}//Creating command queuecommandQueue CL_CHECK2(clCreateCommandQueue(context, device_id, 0, _err));printf(Upload source buffers\n);CL_CHECK(clEnqueueWriteBuffer(commandQueue, a_memobj, CL_TRUE, 0, nbytes, h_a.data(), 0, NULL, NULL));CL_CHECK(clEnqueueWriteBuffer(commandQueue, b_memobj, CL_TRUE, 0, nbytes, h_b.data(), 0, NULL, NULL));printf(Execute the kernel\n);size_t global_work_size[1] {size};size_t local_work_size[1] {1};auto time_start std::chrono::high_resolution_clock::now();CL_CHECK(clEnqueueNDRangeKernel(commandQueue, kernel, 1, NULL, global_work_size, local_work_size, 0, NULL, NULL));CL_CHECK(clFinish(commandQueue));auto time_end std::chrono::high_resolution_clock::now();double elapsed std::chrono::duration_caststd::chrono::milliseconds(time_end - time_start).count();printf(Elapsed time: %lg ms\n, elapsed);printf(Download destination buffer\n);CL_CHECK(clEnqueueReadBuffer(commandQueue, c_memobj, CL_TRUE, 0, nbytes, h_c.data(), 0, NULL, NULL));......2.2 APP Makefile不同POCL仿例同样用blackbox.sh来跑也是分为4个步骤(详见前一篇)前3步都一样唯一差别是第4步APP目录下的$ROOT_DIR/tests/opencl/vecadd/Makefile内容有很大差别。POCL仿例的Makefile主要执行了3个命令。从第1个步骤开始就出现POCL相关目录如红色字体所示。CONFIGS-DTRACING_ALL -DTRACING_ALL ./ci/blackbox.sh--driverrtlsim --appvecadd --debug1 --cores2 --clusters2 --args-n64 --l2cache --l3cache --warps4 --threads42.2.1 编译应用层main.ccg -stdc17 -Wall -Wextra -Wfatal-errors -Wno-deprecated-declarations -Wno-unused-parameter -Wno-narrowing -pthread -I$TOOLDIR/pocl/include-DTRACING_ALL -DTRACING_ALL -DNUM_CORES2 -DNUM_CLUSTERS2 -DL2_ENABLE -DL3_ENABLE -DNUM_WARPS4 -DNUM_THREADS4 -g -O0 -c $VORTEX_HOME/tests/opencl/vecadd/main.cc-omain.cc.ocp $VORTEX_HOME/tests/opencl/vecadd/kernel.cl kernel.clcp $VORTEX_HOME/tests/opencl/vecadd/common.h common.h2.2.2 链接APP应用程序g -stdc17 -Wall -Wextra -Wfatal-errors -Wno-deprecated-declarations -Wno-unused-parameter -Wno-narrowing -pthread -I$TOOLDIR/pocl/include-DTRACING_ALL -DTRACING_ALL -DNUM_CORES2 -DNUM_CLUSTERS2 -DL2_ENABLE -DL3_ENABLE -DNUM_WARPS4 -DNUM_THREADS4 -g -O0main.cc.o-Wl,-rpath,$TOOLDIR/llvm-vortex/lib -L$ROOT_DIR/runtime-lvortex-L$TOOLDIR/pocl/lib-lOpenCL-ovecadd//同样用到前3步生成的$ROOT_DIR/runtime/libvortex.so回调函数实现C TB//包含动态分配地址控制RAM和verilog TB数据传输//下载RISC-V GPGPU运行程序控制内核开始运行上传运行结果等底层驱动功能2.2.3 执行应用程序这一步差别很大。除了引入POCL的相关目录多了3个跟kernel RISC-V的编译和链接有关的变量如红色/紫色/蓝色字体所示。一般的VX仿例是用独立的命令放在Makefile里面来编译和链接。这个是包在环境变量里面传给APP。看过去是交给APP里面的POCL函数来做这个步骤猜测是2.1列表中的clCreateProgramWithSource/clBuildProgram/clCreateKernel因为如果把$ROOT_DIR/tests/opencl/vecadd/kernel.cl故意改个名称前面几个POCL函数都执行了(可以看到log里面有Allocate device buffers)但是它的下一步就报错Create program from kernel source Failed to load kernel。从这一点可以看出POCL是runtime的编译RISC-V kernel文件。LD_LIBRARY_PATH$TOOLDIR/pocl/lib:$ROOT_DIR/runtime:$TOOLDIR/llvm-vortex/lib:POCL_VORTEX_XLEN32LLVM_PREFIX$TOOLDIR/llvm-vortexPOCL_VORTEX_BINTOOLOBJCOPY$TOOLDIR/llvm-vortex/bin/llvm-objcopy $VORTEX_HOME/kernel/scripts/vxbin.pyPOCL_VORTEX_CFLAGS-marchrv32imaf -mabiilp32f -O3 -mcmodelmedany --sysroot$TOOLDIR/riscv32-gnu-toolchain/riscv32-unknown-elf --gcc-toolchain$TOOLDIR/riscv32-gnu-toolchain -fno-rtti -fno-exceptions -nostartfiles -nostdlib -fdata-sections -ffunction-sections -I$ROOT_DIR/hw -I$VORTEX_HOME/kernel/include -DXLEN_32 -DNDEBUG -DTRACING_ALL -DTRACING_ALL -DNUM_CORES2 -DNUM_CLUSTERS2 -DL2_ENABLE -DL3_ENABLE -DNUM_WARPS4 -DNUM_THREADS4 -Xclang -target-feature -Xclang vortex -Xclang -target-feature -Xclang zicond -mllvm -disable-loop-idiom-all POCL_VORTEX_LDFLAGS-Wl,-Bstatic,--gc-sections,-T$VORTEX_HOME/kernel/scripts/link32.ld,--defsymSTARTUP_ADDR0x80000000$ROOT_DIR/kernel/libvortex.a-L$TOOLDIR/libc32/lib -lm -lc $TOOLDIR/libcrt32/lib/baremetal/libclang_rt.builtins-riscv32.aPOCL_DEBUGall VORTEX_DRIVERrtlsim ./vecadd -n128POCL APP目录下的kernel很简单。#include common.h__kernel void vecadd (__global const TYPE *A,__global const TYPE *B,__global TYPE *C){int gid get_global_id(0);C[gid] A[gid] B[gid];}作为对比VX仿例的kernel如下。#include vx_spawn.h#include common.hvoid kernel_body(kernel_arg_t* __UNIFORM__ arg) {auto src0_ptr reinterpret_castTYPE*(arg-src0_addr);auto src1_ptr reinterpret_castTYPE*(arg-src1_addr);auto dst_ptr reinterpret_castTYPE*(arg-dst_addr);dst_ptr[blockIdx.x] src0_ptr[blockIdx.x] src1_ptr[blockIdx.x];}int main() {kernel_arg_t* arg (kernel_arg_t*)csr_read(VX_CSR_MSCRATCH);return vx_spawn_threads(1, arg-num_points, nullptr, (vx_kernel_func_cb)kernel_body, arg);}三、POCL在Vortex中的功能综合2.2.3的描述可以推断POCL在Vortex中的功能对应前面的框图。PC端首先是中间层适配功能因为标准的opencl函数和Vortex的C TB(参见前一章的回调函数体现在$VORTEX_HOME/runtime/stub/vortex.cpp)差别挺大必须有POCL来作为中间的桥梁这些中间层代码最终要嵌入到APP作为APP的一部分PC端其次是运行时编译功能POCL能通过约定的环境变量来获知如何编译RISC-V kernel文件如2.2.3中标红色/紫色/蓝色的那些变量。它们能提供的信息用什么样的编译器比如是ARM还是RISC-V编译和链接的参数link script启动地址等等DEV端中间层适配功能符合opencl标准的kernel.cl 和VX同样功能的仿例kernel.cpp有所差别所以POCL要提供中间层代码DEV端KERNEL的其他库文件这个通过环境变量约定。比如2.2.3环境变量中的$ROOT_DIR/kernel/libvortex.a主要是6个自定义指令集的功能warp和thread管理总结本文对比了Vortex RTLSIM仿真环境中运行POCL仿例与VX仿例的关键差异。POCL仿例位于$VORTEX_HOME/tests/opencl目录下共20个。主要区别在于1)驱动层使用POCL封装的cl*函数替代VX的vx_*函数2)Makefile流程不同POCL通过环境变量传递RISC-V编译参数由POCL函数完成内核编译3)内核代码结构差异POCL采用标准语法而VX使用特定API。这些差异体现了两种编程接口在编程思想和实现方式上的不同。