From 2f4fd11c93e19efa982d6c0bc4b40f5b68b93881 Mon Sep 17 00:00:00 2001 From: Hansung Kim Date: Tue, 31 Oct 2023 23:32:52 -0700 Subject: [PATCH] Add vecadd-loop This is the same kernel as vecadd but repeated in a for-loop many times so that the runtime overhead at the startup is amortized. --- tests/opencl/vecadd-loop/.gitignore | 2 + tests/opencl/vecadd-loop/Makefile | 72 +++++ tests/opencl/vecadd-loop/README | 0 .../vecadd-loop/kernel.alll1hit.loop1000.cl | 12 + tests/opencl/vecadd-loop/kernel.cl | 13 + tests/opencl/vecadd-loop/kernel.cl.loop | 9 + tests/opencl/vecadd-loop/kernel.pocl | Bin 0 -> 15483 bytes tests/opencl/vecadd-loop/main.cc | 250 ++++++++++++++++++ 8 files changed, 358 insertions(+) create mode 100644 tests/opencl/vecadd-loop/.gitignore create mode 100644 tests/opencl/vecadd-loop/Makefile create mode 100644 tests/opencl/vecadd-loop/README create mode 100644 tests/opencl/vecadd-loop/kernel.alll1hit.loop1000.cl create mode 100644 tests/opencl/vecadd-loop/kernel.cl create mode 100644 tests/opencl/vecadd-loop/kernel.cl.loop create mode 100644 tests/opencl/vecadd-loop/kernel.pocl create mode 100644 tests/opencl/vecadd-loop/main.cc diff --git a/tests/opencl/vecadd-loop/.gitignore b/tests/opencl/vecadd-loop/.gitignore new file mode 100644 index 00000000..c2f4a18b --- /dev/null +++ b/tests/opencl/vecadd-loop/.gitignore @@ -0,0 +1,2 @@ +vecadd-loop +*.ll diff --git a/tests/opencl/vecadd-loop/Makefile b/tests/opencl/vecadd-loop/Makefile new file mode 100644 index 00000000..bed44ce1 --- /dev/null +++ b/tests/opencl/vecadd-loop/Makefile @@ -0,0 +1,72 @@ +XLEN ?= 32 + +LLVM_PREFIX ?= /opt/llvm-riscv +RISCV_TOOLCHAIN_PATH ?= /opt/riscv-gnu-toolchain +SYSROOT ?= $(RISCV_TOOLCHAIN_PATH)/riscv32-unknown-elf +POCL_CC_PATH ?= /opt/pocl/compiler +POCL_RT_PATH ?= /opt/pocl/runtime + +OPTS ?= -n64 + +VORTEX_DRV_PATH ?= $(realpath ../../../driver) +VORTEX_RT_PATH ?= $(realpath ../../../runtime) + +K_LLCFLAGS += "-O3 -march=riscv32 -target-abi=ilp32f -mcpu=generic-rv32 -mattr=+m,+f -mattr=+vortex -float-abi=hard -code-model=small" +K_CFLAGS += "-v -O3 --sysroot=$(SYSROOT) --gcc-toolchain=$(RISCV_TOOLCHAIN_PATH) -march=rv32imf -mabi=ilp32f -Xclang -target-feature -Xclang +vortex -I$(VORTEX_RT_PATH)/include -fno-rtti -fno-exceptions -ffreestanding -nostartfiles -fdata-sections -ffunction-sections" +K_LDFLAGS += "-Wl,-Bstatic,-T$(VORTEX_RT_PATH)/linker/vx_link$(XLEN).ld -Wl,--gc-sections $(VORTEX_RT_PATH)/libvortexrt.a -lm" + +CXXFLAGS += -std=c++11 -Wall -Wextra -Wfatal-errors + +CXXFLAGS += -Wno-deprecated-declarations -Wno-unused-parameter -Wno-narrowing + +CXXFLAGS += -I$(POCL_RT_PATH)/include + +LDFLAGS += -L$(POCL_RT_PATH)/lib -L$(VORTEX_DRV_PATH)/stub -lOpenCL -lvortex + +# Debugigng +ifdef DEBUG + CXXFLAGS += -g -O0 +else + CXXFLAGS += -O2 -DNDEBUG +endif + +PROJECT = vecadd-loop + +SRCS = main.cc + +all: $(PROJECT) kernel.pocl + +kernel.pocl: kernel.cl + LLVM_PREFIX=$(LLVM_PREFIX) POCL_DEBUG=all LD_LIBRARY_PATH=$(LLVM_PREFIX)/lib:$(POCL_CC_PATH)/lib $(POCL_CC_PATH)/bin/poclcc -CFLAGS $(K_CFLAGS) -LDFLAGS $(K_LDFLAGS) -o kernel.pocl kernel.cl + +$(PROJECT): $(SRCS) + $(CXX) $(CXXFLAGS) $^ $(LDFLAGS) -o $@ + +run-fpga: $(PROJECT) kernel.pocl + LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/fpga:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS) + +run-asesim: $(PROJECT) kernel.pocl + LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/asesim:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS) + +run-vlsim: $(PROJECT) kernel.pocl + LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/vlsim:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS) + +run-simx: $(PROJECT) kernel.pocl + LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/simx:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS) + cp -f args.bin $(PROJECT).args.bin + +run-rtlsim: $(PROJECT) kernel.pocl + LD_LIBRARY_PATH=$(POCL_RT_PATH)/lib:$(VORTEX_DRV_PATH)/rtlsim:$(LD_LIBRARY_PATH) ./$(PROJECT) $(OPTS) + +.depend: $(SRCS) + $(CXX) $(CXXFLAGS) -MM $^ > .depend; + +clean: + rm -rf $(PROJECT) *.o .depend + +clean-all: clean + rm -rf *.pocl *.dump + +ifneq ($(MAKECMDGOALS),clean) + -include .depend +endif diff --git a/tests/opencl/vecadd-loop/README b/tests/opencl/vecadd-loop/README new file mode 100644 index 00000000..e69de29b diff --git a/tests/opencl/vecadd-loop/kernel.alll1hit.loop1000.cl b/tests/opencl/vecadd-loop/kernel.alll1hit.loop1000.cl new file mode 100644 index 00000000..a610b183 --- /dev/null +++ b/tests/opencl/vecadd-loop/kernel.alll1hit.loop1000.cl @@ -0,0 +1,12 @@ +__kernel void vecadd (__global const float *A, + __global const float *B, + __global float *C) +{ + int gid = get_global_id(0); + float sum = 0.; + for (int i = 0; i < 1000; i++) { + int addr = gid + (i % 2); + sum += A[addr] + B[addr]; + } + C[gid] = sum; +} diff --git a/tests/opencl/vecadd-loop/kernel.cl b/tests/opencl/vecadd-loop/kernel.cl new file mode 100644 index 00000000..535b2cf8 --- /dev/null +++ b/tests/opencl/vecadd-loop/kernel.cl @@ -0,0 +1,13 @@ +__kernel void vecadd_loop (__global volatile const float *A, + __global volatile const float *B, + __global volatile float *C) +{ + int gid = get_global_id(0); + float sum = 0.; + for (int i = 0; i < 500; i++) { + // int addr = gid + (i % 2); + int addr = gid; + C[addr] += A[addr] + B[addr]; + } + // C[gid] = sum; +} diff --git a/tests/opencl/vecadd-loop/kernel.cl.loop b/tests/opencl/vecadd-loop/kernel.cl.loop new file mode 100644 index 00000000..9a4c2882 --- /dev/null +++ b/tests/opencl/vecadd-loop/kernel.cl.loop @@ -0,0 +1,9 @@ +__kernel void vecadd_loop (__global volatile const float *A, + __global volatile const float *B, + __global volatile float *C) +{ + int gid = get_global_id(0); + for (int i = 0; i < 100; i++) { + C[gid] = A[gid] + B[gid]; + } +} diff --git a/tests/opencl/vecadd-loop/kernel.pocl b/tests/opencl/vecadd-loop/kernel.pocl new file mode 100644 index 0000000000000000000000000000000000000000..8422ab307437ba9e2b5226f12294016c7a333df1 GIT binary patch literal 15483 zcmeHN4^&fEn!hhEc`rx^!Jvd1(ia{AQj3950a2R}1gwn4>P+oUcTI>OR+I>oKX$fF zc!7tOBiOjgbh|T9wPR;zJF&4=XJ^}BM5&Dy|5T~tu}~11u5D>ipdD)Vy8$ZGKeIh& z&+OTq!}DBz$-Vcx_q*TsyWjotyw(y+asH;Q=#b~lC5OD)M`?r*0YwZ2KPz(=lvo#N z9)y2WYZ76elyHQ~=E}zQi4bzuwnuH2y^cObC`vApq0LA%wT+o*V+HgPOUVoT3W;fzuMTm@B26NUTIoOu`ErqxRG~(i*{$MS^z1)tt zL+u2)UnFpg+6CIWqi85v=W|LC8e#!q?6O3i&rKag@;$N0O*JAd55>VFr%-o4-9jOr zp%JYXqBT|-6aqONxBYR95cNJUyz8u;h1^FEBIl{uXtQlL>b-UF%gsoNoV!%W{kn>9 zw$Fn4Sx~R~VzZ3|9Tmv^P!amNzPM<3vrT|Fup=2WWsL!Q+Kf&@2=+q(AIe38e(cMA zVc=ssk)KK+f7x8d=bUvAX^F!DpJ_#}3nFHMO$DzdY6UMLd8uG*C|2XM)KOoLbI*oh z0slE^7#I7A0y&WY?E@b=Yis`?7Ye{{o+t7AxDDG%D>{JvOb*3qe9ku_U~OP)@C^z2 z|CxU*ZV<13C!> za^Xoo)KBPi+mF9v9``m2?@k?m%6jjziLkEuu;v69D_%{ou}K6QqbJzt)kImG(Kj?t z;~R=azOq=^SXlxZa|bOVgcLLaf4zlq2Ln71bpmmnnEJH9~f9q%KH(J0c+I3th; zfgSOh=L(@u7VzDHNZ_^{_LLKRqpf=zX{lHO+Ysxj4O&k0zVN~r8RsDIzT51%iTy-? zp9ruG0k$EydEmDMJItrYb`oG-48?-3af~kl{3fUQyUx9_NK0RV`9Qb1u%}|%5t-9Y zaEC;QdwCgh^QV!!u1jDM(?T-m20>=OM`*9y)VZlE0ylp}XQ3z|nR|!OcF}t7Wf4KX zx{SbjyQwau6)Q-R9pAfJ`ZVF@X}z}YG}?$rq&2}{RK zsaCvB;G|`+Hj32lsux(+UPtcJr$|fokkqnnL?S1L{MvVi1x{fovdr#9?*3tc+tE*G z-*yPxeGY*oH!{R2l_BlQNYX+*(%_BZ5x1e?4+u zuF$(jXu|o{{m5B2d(`>MMuhj5b9gw^xtAi`H%FjcLHi&^JL{+$EYk}6=rHWxHsITN zf=y^6*!cv~lAj9JG}7IeTUx#!$=5xrhnO=eSJZ<%O|Xfd!M=*lGF6f z%}qDaJ)eAVZmLgucL_(eK|5e?flb5r1Y#%~t0p$v;A|NDpldS%UVo1sL*SQG+7DcV z@%QnlEDpKYmYaYw7yA-M#>F1MsBy6`W7N6WLl})N_SYCqF7`JVtuFRej5aJUD2r>t z@_;Q^9?*s50X2zg5?2SSRT-WkK19?Er zF^~t;9Rqnl<1vs2G#vwZK=gFK-1ILHIqu)GxJ$MS$JSRT-YE?Z~o+rku0}Sq4v2&dP^? zpCV*=fu`ic9ey{}uY;Ht4(rv>C83eKt=*k9R*eU&%==@I0vy@M=85#zZC9=B?ih&wHUIbgp+KCnM@?@VjY$Vdb4kT`9E3v)Li9T zYp)r&`>ncP48Ynx7?YEOt@`_=IxBU>G-hkK!B?e0D z8#uNQy@glqTEZ!WZR4iAI(#ne6Ts)XpFrIy#L{n|{S@?vvR?!UWxo&*%6<|M%6G!Cl>G=0N`Oyy?m(QaMeZB#A?IEZff4i_cJ75UdSV`D4T7_k`qBycnv&r z2@47m%7bB#JO^cl3^{`a{anm)B!?W>k|`88_rxH#@aM?l2o`d`%92@Zl_I(Dw%_SE z>F=GL;{Rq$hn$0eoP&UzgCHM(9OMDWK_G9MJaf3X<&alofqh5(kmm{6mmtrsr@rLk zb||#_c0%9f$c=gPJS_lUfW4uA+|F=II8XFC;=bsur$^;GXzaT?T-@6V?qHT!i~B3_ z?~q{ZKioFPM(+Up<(KkW(L~)Z|AeNWX$$;6wm=I)wn_DG&qTH-w& z@FUfp{nwP4Tl#NZ7op(ss_DPe7MQlcU)cg;8M0l98=pirflt9#ya%Vea3p_#8uw51 z*$=O#{uv5nD}iTni(}NoYYJ3YTw*Ry_)Fa~zugVqdjy`TqZ#gXQ$gd@_n{MSNUi2l zb8&G&abmuuVJ1SE2G8|};4KtZEnlLEdF1go3xD*dBkO}#+w$a40 zGx-LY{&~TM4)Nh|#*kglyiFuO>QE%aCp0Df%=N@>ro98k-N7IPFWFFXbZ!u{p;qXO zKHz{qy^hwV@R@l#8-g7N@?>jSq($NVVh=-~QWBZ9b?IVte}n!IqZ98wr8v+LPP$S~ zwJelrM4B9#FU%$gdI5$+s}%_S8{nLOh#U38L&(s>s2eraZ>yrKoA^=%z1_tBtca|x z;1`$+#@sIE%T4@31#dM$EA3I9u!Eoy zoly3vl~-GpSCQ(vI_jETbw;m3l~r_QJZ%HPX1=gVNqf{8v+CZ`u(*}C!cgBYZc<(7QC@3Rc1&6Z4CNzKS6iY^ z*j1f`228@tO1 zwTE2xV2cg3DhCPXqqe9EZAxsw){swN7IW0qZtT{NK=X-NTIlLS;NSY~uzWFH20YQw zRYe^5WxMX4XJ1l*e?jDWn{ogaZjZXI26s%Z+Qg}49>LlEv?mnYbRdKZL&3%^ zni<48+YE%319$P>lRI?H8c7x*msSI#6!g>0WO`TUKN;p6;Z%DJb3B~N2cB+IcIf9E z;G#P1sxBhxh&=-;kF?D{X;*gJv&DwQP#N(~=KNZf!{_r)QpE2NmCbNOsPkR2^P{&o2+V`0kFqeFSW<6Db9fsvZFF@NUDmu!(vLzvmPv2hq*_U)>O@=AxmIPN zMil7A&GC^-It)u1$R)n2hH@+YGu;-fX^FJAse$p+yCT{=~zw1B{IY1m7LzhTi5ZWR<68=t^!KJR&L=V&#ejl1ciQX zl=d}Adv&2bjPw?)f{Bm1)~14RazU+xU~p}cdWDMBX)eQ}7r6|dF=k>RCLZEMRrBh+ z+fAX^0;(qiJ3}3Hy)Eh)LWLDpM-7^(wY6ND&zs>gECGXzomkvw$f$>}-iG!We04Kl z3HktYA@VibpXilIZ+WCQjM5(>g$Q<>WR8es0won-P9lxdP5U>g9#lrO1H|Sy&5Te9qOk3 z*e&fdN=G$@9*ADBO9qL6=H~26?|ak0SuXE_F2f?9VNtzd5v;n?u+-=J{ZsLDwUsY5 zft&DdhQFl@T|HV`4X(sq+xCzDX_S86ExqB9z9o}>ZpD8mW*iLu%oy6^0rP6W5luL_ zR*cpH0m?>2+dJs0cs}ySir;_}xr{wz##m-rhc|t!1{k{-TnV##V7-Sxmdp-{0M;d`7MM2%ia53iV9|x}`mxu;gAf`81q1EqtkmPCZ0d=F(;GwTO>g z9ogm0=pdKuA(tK_(hqP;_Hc_|G^8H`<{N!B<%KqU8o<^7CKRsS2*#9!_6-J)3`#$D zN#C+cKOdCdL7{JHLjTwTcJ9@aRaHEkHtW-XY_Lo?uXr#KLe&gHMX(=BRk@6zUJ@Bp znw-n9_J-ZN&u6vB;x{Z#O3vKXq~3o>^A@o)Kj!W9;&wUS^_7+)bLslBP1_5~&7~U) z${#f6Z`xFB)g~7jmm8S8f{KD-NFSnL0 zTCiZ_rt+ev@)Ip3TNYH5l$IA%ZM1H*KDD5@xME8pv@z+a8y+gurX+7jPD-||lF}@xX=(YE^z`%st!7?80{_2n;$%Lv-fXcHl$8~f1#+p9En7EPOUetj zmNQuSHJi3riwhXT0>c9jFl#gQdc=yoni`Fa6&sRjGzwOHAxWgEm2F#t)`-`n!bfAZ zR1XdsNCEY(#Gr7rX}xIm z^9{n-cy{5!e_y!oXzFW=jwVNFkGkB)SJ$8Uoomna?lad;bveK?4O-ZV;Y!w~2^1b3Vh(!Rw9xQrJeD3MEAxWy2ln!MPz zW(J{u0h9}u&&THgK4((0(TuMEKhq!tFYi!vPDj&fRR^2{2swHfJ?^))WNX2A)@mu) zT2{_floXrGHx(C5qCV4r=F}BDzEHA3M{RJKPoFDk6g zB@S;mFSE-+vf30OVLLxt3}3_AnDERtMP*89CcH>%)3w`(HAyQN$y&+Z86(|7uddw_ z6ud19EfGJ~$ytMzh_7f?SWTblJywHkciVigN!aWlgpyQYXrwP$m}@(u&_R<1nmD77 zoMkBMXwyB?e85qiiHnn-S>^X`?B(T ztQ?A1d0T!q+`ROBURXcNqaQnv#QRp-xn6lh^3PY;Q)(L5Jx55qAF~@}Q~Gzam(Pzq z$D(%{CH@;1XTOG43r!u?_m}3+Ir~(kCRrgW?wS(?uT%a)-! z$zvb0rzWR0Ewk?}x!B z7&-T#NFg*g?P%MZDXGzPoHVV_Q+maQd>AV5m>Gnc;L~?;ue$S~Njl?XQMU0Hl7UAx zv?xD?O*2yE^O`w9(xT{>>Yj{O>{sV+|4l9$7{&o6q~K5j)~?TuWJ)) zR+D1IRCueVH(v7$Pf}-9${NQPm`pHNJCn;o>!mucEFw7w*ihg2 zqU|s7aM z%q{q-iGE5YHMKlwh|ImPmk5b@az$%^F|LE_I7X73)b;2Qy^Q*~qIlkQM$F}50eH*| zLc0LJlSjUTx`sl>$+u03m_nMBRM|Y)e)h!b=vqu1rtme)qs-D@@2mk@#piPIir!t1 ztQ4E}GQgHj)V%X?fw?8$$h9OF2G!Yj+Cv<*Y0);7BVj<_>`*z1n&+mu;scb42FfTJ zkW-5()2#+Jp&;Y(KvOI<$#$Ar@{6Ra6}dsqJ(^Ww+g{BAnW@WGRHN3hQj_VJI*rsc z7cjeYUi&zI7EY4He*;JoJ(WMRCz}HNnfu@IM=`;lv){*`*eU#p`=R`4|BL)NKEa=# z{O|IoZdJ`uNj|gur*+9qhaPN?91lQ(knsRCNhec)PQV%bY-7fM;Ur6C(GFx2DF!#$ z1c&jj8-xmgV-En%B=?x`d5OQD)jfD^L#+iIUnn{-U_$w)^9E- z-C9t*e#$GX@2#&aDc!uWwB#x459m>7E-Rl%^L}XdR=TOoQlU-$I&;gnV6Dg3_Z@Gn zrf58Ur*HgT(Da=?eCKZZP9NM)f;*D>={tSXclz*Mf$2Ma(|7u&@AOUI>6^aOH+`pX r`c5CbubjTqH+`pX{I})w-@hdbv`pXW`&)A7^jfAZ@K?0}zSH+#PPV9* literal 0 HcmV?d00001 diff --git a/tests/opencl/vecadd-loop/main.cc b/tests/opencl/vecadd-loop/main.cc new file mode 100644 index 00000000..f0eec447 --- /dev/null +++ b/tests/opencl/vecadd-loop/main.cc @@ -0,0 +1,250 @@ +#include +#include +#include +#include +#include +#include +#include +#include + +#define KERNEL_NAME "vecadd_loop" + +#define CL_CHECK(_expr) \ + do { \ + cl_int _err = _expr; \ + if (_err == CL_SUCCESS) \ + break; \ + printf("OpenCL Error: '%s' returned %d!\n", #_expr, (int)_err); \ + cleanup(); \ + exit(-1); \ + } while (0) + +#define CL_CHECK2(_expr) \ + ({ \ + cl_int _err = CL_INVALID_VALUE; \ + decltype(_expr) _ret = _expr; \ + if (_err != CL_SUCCESS) { \ + printf("OpenCL Error: '%s' returned %d!\n", #_expr, (int)_err); \ + cleanup(); \ + exit(-1); \ + } \ + _ret; \ + }) + +static int read_kernel_file(const char* filename, uint8_t** data, size_t* size) { + if (nullptr == filename || nullptr == data || 0 == size) + return -1; + + FILE* fp = fopen(filename, "r"); + if (NULL == fp) { + fprintf(stderr, "Failed to load kernel."); + return -1; + } + fseek(fp , 0 , SEEK_END); + long fsize = ftell(fp); + rewind(fp); + + *data = (uint8_t*)malloc(fsize); + *size = fread(*data, 1, fsize, fp); + + fclose(fp); + + return 0; +} + +static int write_operand_file(const char* filename, void* data, size_t size) { + if (nullptr == filename || nullptr == data || 0 == size) + return -1; + + FILE* fp = fopen(filename, "wb"); + if (NULL == fp) { + fprintf(stderr, "Failed to write operand data.\n"); + return -1; + } + + size_t wsize = fwrite(data, size, 1, fp); + if (wsize != 1) { + fprintf(stderr, "Failed to write operand data.\n"); + return -1; + } + + return 0; +} + +static bool almost_equal(float a, float b, int ulp = 4) { + union fi_t { int i; float f; }; + fi_t fa, fb; + fa.f = a; + fb.f = b; + return std::abs(fa.i - fb.i) <= ulp; +} + +cl_device_id device_id = NULL; +cl_context context = NULL; +cl_command_queue commandQueue = NULL; +cl_program program = NULL; +cl_kernel kernel = NULL; +cl_mem a_memobj = NULL; +cl_mem b_memobj = NULL; +cl_mem c_memobj = NULL; +float *h_a = NULL; +float *h_b = NULL; +float *h_c = NULL; +uint8_t *kernel_bin = NULL; + +static void cleanup() { + if (commandQueue) clReleaseCommandQueue(commandQueue); + if (kernel) clReleaseKernel(kernel); + if (program) clReleaseProgram(program); + if (a_memobj) clReleaseMemObject(a_memobj); + if (b_memobj) clReleaseMemObject(b_memobj); + if (c_memobj) clReleaseMemObject(c_memobj); + if (context) clReleaseContext(context); + if (device_id) clReleaseDevice(device_id); + + if (kernel_bin) free(kernel_bin); + if (h_a) free(h_a); + if (h_b) free(h_b); + if (h_c) free(h_c); +} + +int size = 64; + +static void show_usage() { + printf("Usage: [-n size] [-h: help]\n"); +} + +static void parse_args(int argc, char **argv) { + int c; + while ((c = getopt(argc, argv, "n:h?")) != -1) { + switch (c) { + case 'n': + size = atoi(optarg); + break; + case 'h': + case '?': { + show_usage(); + exit(0); + } break; + default: + show_usage(); + exit(-1); + } + } + + printf("Workload size=%d\n", size); +} + +int main (int argc, char **argv) { + // parse command arguments + parse_args(argc, argv); + + cl_platform_id platform_id; + size_t kernel_size; + cl_int binary_status; + + // read kernel binary from file + if (0 != read_kernel_file("kernel.pocl", &kernel_bin, &kernel_size)) + return -1; + + // Getting platform and device information + CL_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(float); + 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"); + program = CL_CHECK2(clCreateProgramWithBinary( + context, 1, &device_id, &kernel_size, (const uint8_t**)&kernel_bin, &binary_status, &_err)); + if (program == NULL) { + cleanup(); + return -1; + } + + // Build program + CL_CHECK(clBuildProgram(program, 1, &device_id, NULL, NULL, NULL)); + + // Create kernel + kernel = CL_CHECK2(clCreateKernel(program, KERNEL_NAME, &_err)); + + // Set kernel arguments + // NOTE(hansung): clSetKernelArg doesn't seem to incur any device-specific + // operation + CL_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. + h_a = (float*)malloc(nbytes); + h_b = (float*)malloc(nbytes); + h_c = (float*)malloc(nbytes); + + // Initialize values for array members. + for (int i = 0; i < size; ++i) { + h_a[i] = sinf(i)*sinf(i); + h_b[i] = cosf(i)*cosf(i); + h_c[i] = 0xdeadbeef; + //printf("*** [%d]: h_a=%f, h_b=%f\n", i, h_a[i], h_b[i]); + } + + // NOTE(hansung): Dump operand buffer to a file + if (write_operand_file("vecadd.input.a.bin", h_a, nbytes) != 0) + return EXIT_FAILURE; + if (write_operand_file("vecadd.input.b.bin", h_b, nbytes) != 0) + return EXIT_FAILURE; + + // Creating command queue + // NOTE(hansung): The 3rd properties arg is a bit-field, where fields like + // CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE can be set. With value of 0, + // nothing is set and the commands in the queue will be completed in-order. + // See OpenCL 1.2 spec, section 5.1 + commandQueue = CL_CHECK2(clCreateCommandQueue( + context, device_id, 0 /* command-queue properties */, &_err)); + + printf("Upload source buffers\n"); + CL_CHECK(clEnqueueWriteBuffer(commandQueue, a_memobj, CL_TRUE, 0, nbytes, h_a, 0, NULL, NULL)); + CL_CHECK(clEnqueueWriteBuffer(commandQueue, b_memobj, CL_TRUE, 0, nbytes, h_b, 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)); + // NOTE(hansung): clFinish blocks until all kernels in the command queue are + // finished. This seems to be what actually kicks off kernel execution. + CL_CHECK(clFinish(commandQueue)); + auto time_end = std::chrono::high_resolution_clock::now(); + double elapsed = std::chrono::duration_cast(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, 0, NULL, NULL)); + + printf("Verify result\n"); + int errors = 0; + for (int i = 0; i < size; ++i) { + float ref = h_a[i] + h_b[i]; + if (!almost_equal(h_c[i], ref)) { + if (errors < 100) + printf("*** error: [%d] expected=%f, actual=%f, a=%f, b=%f\n", i, ref, h_c[i], h_a[i], h_b[i]); + ++errors; + } + } + if (0 == errors) { + printf("PASSED!\n"); + } else { + printf("FAILED! - %d errors\n", errors); + } + + // Clean up + cleanup(); + + return errors; +}