From 615ce29c84e5e33bbab2817712c06209cf2d83e3 Mon Sep 17 00:00:00 2001 From: Juan Fumero Date: Tue, 12 Dec 2023 12:24:44 +0100 Subject: [PATCH 1/6] Rename module to beehive-levelzero-jni --- pom.xml | 6 +++--- scripts/compileAndRun.sh | 4 ++-- scripts/copies.sh | 2 +- scripts/events.sh | 2 +- scripts/fences.sh | 2 +- scripts/kernelTimers.sh | 2 +- scripts/largeBuffers.sh | 2 +- scripts/run.sh | 2 +- scripts/transfersTimers.sh | 2 +- 9 files changed, 12 insertions(+), 12 deletions(-) diff --git a/pom.xml b/pom.xml index b5bc15b..c4afd8c 100644 --- a/pom.xml +++ b/pom.xml @@ -4,9 +4,9 @@ xsi:schemaLocation="http://maven.apache.org/POM/4.0.0 http://maven.apache.org/xsd/maven-4.0.0.xsd"> 4.0.0 - uk.ac.manchester.levelzero - levelzero - 0.1.0 + beehive-lab + beehive-levelzero-jni + 0.1.2 levelzero https://github.com/beehive-lab/levelzero-jni diff --git a/scripts/compileAndRun.sh b/scripts/compileAndRun.sh index 4b3964e..6259799 100755 --- a/scripts/compileAndRun.sh +++ b/scripts/compileAndRun.sh @@ -1,7 +1,7 @@ -mvn clean package +mvn clean install clang -cc1 -triple spir copy_data.cl -O0 -finclude-default-header -emit-llvm-bc -o opencl-copy.bc llvm-spirv opencl-copy.bc -o opencl-copy.spv cp opencl-copy.spv /tmp/opencl-copy.spv -java -Djava.library.path=./levelZeroLib/build -cp target/levelzero-0.1.0.jar uk.ac.manchester.tornado.drivers.spirv.levelzero.samples.TestLevelZero +java -Djava.library.path=./levelZeroLib/build -cp target/beehive-levelzero-jni-0.1.2.jar uk.ac.manchester.tornado.drivers.spirv.levelzero.samples.TestLevelZero diff --git a/scripts/copies.sh b/scripts/copies.sh index de90580..3c92816 100755 --- a/scripts/copies.sh +++ b/scripts/copies.sh @@ -1 +1 @@ -java -Djava.library.path=./levelZeroLib/build -cp target/levelzero-0.1.0.jar uk.ac.manchester.tornado.drivers.spirv.levelzero.samples.TestCopies +java -Djava.library.path=./levelZeroLib/build -cp target/beehive-levelzero-jni-0.1.2.jar uk.ac.manchester.tornado.drivers.spirv.levelzero.samples.TestCopies diff --git a/scripts/events.sh b/scripts/events.sh index 9127f99..796d6a1 100755 --- a/scripts/events.sh +++ b/scripts/events.sh @@ -1,3 +1,3 @@ java -Djava.library.path=./levelZeroLib/build \ - -cp target/levelzero-0.1.0.jar \ + -cp target/beehive-levelzero-jni-0.1.2.jar \ uk.ac.manchester.tornado.drivers.spirv.levelzero.samples.TestWithEvents diff --git a/scripts/fences.sh b/scripts/fences.sh index b5a9fd9..583db71 100755 --- a/scripts/fences.sh +++ b/scripts/fences.sh @@ -1,4 +1,4 @@ #!/bin/bash ./scripts/compileSPIRVKernelCopy.sh -java -Djava.library.path=./levelZeroLib/build -cp target/levelzero-0.1.0.jar uk.ac.manchester.tornado.drivers.spirv.levelzero.samples.TestFences +java -Djava.library.path=./levelZeroLib/build -cp target/beehive-levelzero-jni-0.1.2.jar uk.ac.manchester.tornado.drivers.spirv.levelzero.samples.TestFences diff --git a/scripts/kernelTimers.sh b/scripts/kernelTimers.sh index 6ca2037..c06c1e2 100755 --- a/scripts/kernelTimers.sh +++ b/scripts/kernelTimers.sh @@ -1,4 +1,4 @@ #!/bin/bash ./scripts/compileSPIRVKernelCopy.sh -java -Djava.library.path=./levelZeroLib/build -cp target/levelzero-0.1.0.jar uk.ac.manchester.tornado.drivers.spirv.levelzero.samples.TestKernelTimer +java -Djava.library.path=./levelZeroLib/build -cp target/beehive-levelzero-jni-0.1.2.jar uk.ac.manchester.tornado.drivers.spirv.levelzero.samples.TestKernelTimer diff --git a/scripts/largeBuffers.sh b/scripts/largeBuffers.sh index 899fac2..b74de2b 100755 --- a/scripts/largeBuffers.sh +++ b/scripts/largeBuffers.sh @@ -1 +1 @@ -java -Djava.library.path=./levelZeroLib/build -cp target/levelzero-0.1.0.jar uk.ac.manchester.tornado.drivers.spirv.levelzero.samples.TestLargeBuffer +java -Djava.library.path=./levelZeroLib/build -cp target/beehive-levelzero-jni-0.1.2.jar uk.ac.manchester.tornado.drivers.spirv.levelzero.samples.TestLargeBuffer diff --git a/scripts/run.sh b/scripts/run.sh index 485e208..4b973b5 100755 --- a/scripts/run.sh +++ b/scripts/run.sh @@ -2,4 +2,4 @@ clang -cc1 -triple spir copy_data.cl -O0 -finclude-default-header -emit-llvm-bc llvm-spirv opencl-copy.bc -o opencl-copy.spv mv opencl-copy.spv /tmp/opencl-copy.spv -java -Djava.library.path=./levelZeroLib/build -cp target/levelzero-0.1.0.jar uk.ac.manchester.tornado.drivers.spirv.levelzero.samples.TestLevelZero +java -Djava.library.path=./levelZeroLib/build -cp target/beehive-levelzero-jni-0.1.2.jar uk.ac.manchester.tornado.drivers.spirv.levelzero.samples.TestLevelZero diff --git a/scripts/transfersTimers.sh b/scripts/transfersTimers.sh index 8db9429..d57965f 100755 --- a/scripts/transfersTimers.sh +++ b/scripts/transfersTimers.sh @@ -1 +1 @@ -java -Djava.library.path=./levelZeroLib/build -cp target/levelzero-0.1.0.jar uk.ac.manchester.tornado.drivers.spirv.levelzero.samples.TestTransferTimers +java -Djava.library.path=./levelZeroLib/build -cp target/beehive-levelzero-jni-0.1.2.jar uk.ac.manchester.tornado.drivers.spirv.levelzero.samples.TestTransferTimers From 53406064afc84226212cf6a83ec2b08d23c12df3 Mon Sep 17 00:00:00 2001 From: Juan Fumero Date: Tue, 12 Dec 2023 12:29:36 +0100 Subject: [PATCH 2/6] Update for 0.1.2 version --- CHANGELOG.md | 9 +++++++++ README.md | 2 +- 2 files changed, 10 insertions(+), 1 deletion(-) diff --git a/CHANGELOG.md b/CHANGELOG.md index 06cfe52..ffba324 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -1,3 +1,12 @@ +## Beehive LevelZero-JNI 0.1.2 +12/12/23 + +- Multiple SPIR-V Devices Fixed: [commit](https://github.com/beehive-lab/levelzero-jni/commit/fe20b18c9623b4d0533ee50d878b266ecdce46dc) +- Support for data send/receive using Panama off-heap buffers +- Use `-O2` optimization: [commit](https://github.com/beehive-lab/levelzero-jni/commit/721b8aed7ac4e419843b3029be99c11267eeb32c) +- Fix `string` release in JNI code: [commit](https://github.com/beehive-lab/levelzero-jni/commit/3c6d463ebafbf9d2de7be128f79483ff28c5ace6) +- Fix Device properties: [commit](https://github.com/beehive-lab/levelzero-jni/commit/83c2e032197e2f8a13d895d2b75f72693424bd7b) + ## LevelZero-JNI 0.1.1 10/03/2022 diff --git a/README.md b/README.md index 72dbd4c..e2fa3de 100644 --- a/README.md +++ b/README.md @@ -1,4 +1,4 @@ -# LevelZero JNI +# Beehive LevelZero JNI Baremetal GPU and FPGA programming for Java using the [LevelZero API](https://spec.oneapi.io/level-zero/latest/index.html). From a98d947725e251ba3683272f30c7a08abd899d63 Mon Sep 17 00:00:00 2001 From: otabuzzman Date: Sun, 10 Mar 2024 16:11:12 +0100 Subject: [PATCH 3/6] Save work on Windows support --- .gitignore | 4 +- CHANGELOG.md | 7 +- README.md | 131 +++++++++++++----- copyData.cl | 4 + copyData.spv | Bin 0 -> 1204 bytes copyLong.cl | 4 + copyLong.spv | Bin 0 -> 1228 bytes copyTest.cl | 17 +++ copyTest.spv | Bin 0 -> 2108 bytes copy_data.cl | 5 - levelZeroLib/CMakeLists.txt | 57 +++++++- levelZeroLib/src/levelZeroBuffer.cpp | 15 ++ levelZeroLib/src/levelZeroCommandList.cpp | 19 +++ levelZeroLib/src/levelZeroContext.cpp | 123 ++++++++++++++++ levelZeroLib/src/levelZeroDescriptors.cpp | 46 +++++- levelZeroLib/src/levelZeroDevice.cpp | 20 +++ levelZeroLib/src/levelZeroDriver.cpp | 16 +++ levelZeroLib/src/levelZeroModule.cpp | 4 + levelZeroLib/src/levelZeroTimeStampKernel.cpp | 17 ++- lookUpBufferAddress.cl | 3 + lookUpBufferAddress.spv | Bin 0 -> 1096 bytes scripts/compileAndRun.cmd | 3 + scripts/compileAndRun.sh | 10 +- scripts/compileSPIRVKernelCopy.sh | 6 +- scripts/copies.cmd | 1 + scripts/copies.sh | 2 + scripts/events.cmd | 1 + scripts/events.sh | 6 +- scripts/fences.cmd | 4 + scripts/fences.sh | 4 +- scripts/kernelTimers.cmd | 1 + scripts/kernelTimers.sh | 4 +- scripts/largeBuffers.cmd | 1 + scripts/largeBuffers.sh | 2 + scripts/run.cmd | 1 + scripts/run.sh | 9 +- scripts/transfersTimers.cmd | 1 + scripts/transfersTimers.sh | 2 + .../levelzero/samples/SimulationLKBuffer.java | 51 +++---- .../spirv/levelzero/samples/TestFences.java | 19 ++- .../levelzero/samples/TestKernelTimer.java | 26 ++-- .../levelzero/samples/TestLevelZero.java | 19 ++- .../samples/TestLevelZeroDedicatedMemory.java | 19 ++- .../TestLevelZeroDedicatedMemoryLong.java | 19 ++- .../samples/TestLookUpBufferAddress.java | 20 ++- 45 files changed, 567 insertions(+), 156 deletions(-) create mode 100644 copyData.cl create mode 100644 copyData.spv create mode 100644 copyLong.cl create mode 100644 copyLong.spv create mode 100644 copyTest.cl create mode 100644 copyTest.spv delete mode 100644 copy_data.cl create mode 100644 lookUpBufferAddress.cl create mode 100644 lookUpBufferAddress.spv create mode 100644 scripts/compileAndRun.cmd create mode 100644 scripts/copies.cmd create mode 100644 scripts/events.cmd create mode 100644 scripts/fences.cmd create mode 100644 scripts/kernelTimers.cmd create mode 100644 scripts/largeBuffers.cmd create mode 100644 scripts/run.cmd create mode 100644 scripts/transfersTimers.cmd diff --git a/.gitignore b/.gitignore index 3934d96..445176e 100644 --- a/.gitignore +++ b/.gitignore @@ -1,6 +1,6 @@ levelZeroLib/.idea/ levelZeroLib/build/ levelzero.iml -opencl-copy.bc +*.bc +*.log target/ - diff --git a/CHANGELOG.md b/CHANGELOG.md index ffba324..0604143 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -1,4 +1,4 @@ -## Beehive LevelZero-JNI 0.1.2 +## Beehive LevelZero-JNI 0.1.2 12/12/23 - Multiple SPIR-V Devices Fixed: [commit](https://github.com/beehive-lab/levelzero-jni/commit/fe20b18c9623b4d0533ee50d878b266ecdce46dc) @@ -20,8 +20,9 @@ ## LevelZero-JNI 0.1.0 03/12/2021 -- Initial prototype of LevelZero JNI. +- Initial prototype of LevelZero JNI - It covers a subset of the Intel Level-Zero 1.1.2 Spec (Feb 2021) - C++ Wrapper for JNI provided - Java Interface provided -- Set of examples and documentation \ No newline at end of file +- Set of examples and documentation + diff --git a/README.md b/README.md index e2fa3de..43e272b 100644 --- a/README.md +++ b/README.md @@ -1,67 +1,106 @@ -# Beehive LevelZero JNI +# Beehive LevelZero JNI -Baremetal GPU and FPGA programming for Java using the [LevelZero API](https://spec.oneapi.io/level-zero/latest/index.html). +Baremetal GPU and FPGA programming for Java using Intel's [Level Zero API](https://spec.oneapi.io/level-zero/latest/index.html). This project is a Java Native Interface (JNI) binding for Intel's Level Zero. This library is designed to be as closed as possible to the Level Zero API for C++. Subset of Level Zero 1.4.0 supported (Level Zero May 2022 version) -This project is a Java Native Interface (JNI) binding for Intel's Level Zero. This library is as designed to be as closed as possible to the LevelZero API for C++. +## Compilation & configuration of the JNI Level Zero API +### 1) Compile Level Zero API -Subset of LevelZero 1.4.0 supported (Level Zero May 2022 version) +#### Linux +Note: Using tag `v1.4.1` from `level-zero` which implements Level Zero Specification 1.2. -## Compilation & Configuration of the JNI Level-Zero API +```bash +git clone https://github.com/oneapi-src/level-zero.git +cd level-zero +git checkout tags/v1.4.1 +mkdir build +cd build +cmake .. +cmake --build . --config Release +cmake --build . --config Release --target package +``` -### 1) Install Level-Zero +#### Windows +Configuration: +- Lenovo IdeaPad Gaming 3 15IHU6 +- Windows 11 +- VS Community 2022 + + components C++, Git, Spectre mitigated libraries +- CMake 3.26.3, Maven 3.9.1, JDK 21 -Note: Using commit `551dd5810a3cea7a7e26ac4441da31878e804b53` from `level-zero` +Run commands in _x64 Native Tools Command Prompt for VS 2022_. +```cmd +git clone https://github.com/oneapi-src/level-zero +cd level-zero +md build +cd build +cmake .. +cmake --build . --config Release +``` + +Note: Check for extisting Level Zero API libraries (e.g. `ze_tracing_layer.dll`) in `c:\windows\system32` if `zello_world.exe` fails. -```bash -$ git clone https://github.com/oneapi-src/level-zero -$ mkdir build -$ cd build -$ cmake .. -$ cmake --build . --config Release -$ cmake --build . --config Release --target package +``` +rem check +.\bin\Release\zello_world.exe ``` +### 2) Compile Level Zero JNI native code -### 2) Compile JNI Native Code +Set the paths to the directory of Level Zero installation. Here are examples: -Set the paths to the directory of Level-Zero installation. Here's an example: +#### Linux ```bash -$ scl enable devtoolset-9 bash # << Only for CentOS -$ export CPLUS_INCLUDE_PATH=/include:$CPLUS_INCLUDE_PATH -$ export LD_LIBRARY_PATH=/build/lib:$LD_LIBRARY_PATH -$ export ZE_SHARED_LOADER="/build/lib/libze_loader.so" -$ cd levelZeroLib -$ mkdir build -$ cd build -$ cmake .. -$ make +scl enable devtoolset-9 bash # << Only for CentOS +git clone https://github.com/beehive-lab/levelzero-jni +export ZE_SHARED_LOADER="/build/lib/libze_loader.so" +export CPLUS_INCLUDE_PATH=/include:$CPLUS_INCLUDE_PATH +export C_INCLUDE_PATH=/include:$CPLUS_INCLUDE_PATH +cd levelzero-jni/levelZeroLib +mkdir build +cd build +cmake .. +make ``` -##### 2.1 Obtain the LLVM-SPIRV Compiler +#### Windows + +Note: Run commands in _x64 Native Tools Command Prompt for VS 2022_ + +```cmd +git clone https://github.com/otabuzzman/levelzero-jni +set ZE_SHARED_LOADER=%USERPROFILE%\lab\level-zero\build\lib\release\ze_loader.lib +set CPLUS_INCLUDE_PATH=%USERPROFILE%\lab\level-zero\include +set C_INCLUDE_PATH=%USERPROFILE%\lab\level-zero\include +cd levelzero-jni\levelZeroLib +md build +cd build +cmake .. +cmake --build . --config Release +``` -In case you want to compile kernels from OpenCL C to SPIR-V and use the Level Zero library, you need to download the `llvm-spirv` compiler. -The implementation we are currently using is the `intel/llvm`: [https://github.com/intel/llvm](https://github.com/intel/llvm). +#### Obtain a SPIR-V compiler (Linux and Windwos) +In case you want to compile kernels from OpenCL C to SPIR-V and use the Level Zero API, you need to download the `llvm-spirv` compiler. The implementation we are currently using is [Intel LLVM](https://github.com/intel/llvm). -### 3) Compile & Run a Java test +### 3) Compile and run a Java test +#### Linux ```bash -$ mvn clean package -$ ./scripts/run.sh ## < This script compiles an OpenCL C program to SPIR-V using the llvm-spirv compiler (see 2.1) +mvn clean package +./scripts/run.sh ## < This script compiles an OpenCL C program to SPIR-V using the llvm-spirv compiler (see 2.1) ``` The OpenCL C kernel provided for this example is as follows: - ```c -__kernel void copydata(__global int* input, __global int* output) { +__kernel void copyData(__global int* input, __global int* output) { uint idx = get_global_id(0); output[idx] = input[idx]; } @@ -70,11 +109,28 @@ __kernel void copydata(__global int* input, __global int* output) { To compile from OpenCL C to SPIR-V: ```bash -$ clang -cc1 -triple spir copy_data.cl -O0 -finclude-default-header -emit-llvm-bc -o opencl-copy.bc -$ llvm-spirv opencl-copy.bc -o opencl-copy.spv -$ mv opencl-copy.spv /tmp/opencl-copy.spv +clang -cc1 -triple spir copyData.cl -O0 -finclude-default-header -emit-llvm-bc -o copyData.bc +llvm-spirv copyData.bc -o copyData.spv +``` + +#### Windows + +```cmd +mvn clean package + +rem copyData.spv file expected in CWD +.\scripts\run.cmd + +# more tests +.\scripts\copies.cmd +.\scripts\events.cmd +.\scripts\fences.cmd +.\scripts\kernelTimers.cmd +.\scripts\transferTimers.cmd +.\scripts\largeBuffers.cmd ``` + ## License This project is developed at [The University of Manchester](https://www.manchester.ac.uk/), and it is fully open source under the [MIT](https://github.com/beehive-lab/levelzero-jni/blob/master/LICENSE) license. @@ -82,5 +138,4 @@ This project is developed at [The University of Manchester](https://www.manchest ## Acknowledgments -The work was partially funded by the EU Horizon 2020 [Elegant 957286](https://www.elegant-h2020.eu/) project, and Intel Coorporation (https://www.intel.it/content/www/it/it/homepage.html). - +The work was partially funded by the EU Horizon 2020 [Elegant 957286](https://www.elegant-h2020.eu/) project, and [Intel Coorporation](https://www.intel.it/content/www/it/it/homepage.html). diff --git a/copyData.cl b/copyData.cl new file mode 100644 index 0000000..b5cdf19 --- /dev/null +++ b/copyData.cl @@ -0,0 +1,4 @@ +__kernel void copyData(__global int* input, __global int* output) { + uint idx = get_global_id(0); + output[idx] = input[idx]; +} diff --git a/copyData.spv b/copyData.spv new file mode 100644 index 0000000000000000000000000000000000000000..7b9fcb84d65d823b16a3e5f2a4bbe2af17fcade8 GIT binary patch literal 1204 zcma)++e#ck5Qb}aXFSF8F^L*Cgr}X?QgSYEnBu3>)P~Nv$jPt(Q%_^q(2C?q;JcI{OZ^7V?8w0 zb)2s)|BB6u#cWqHD9XuoTBkA{M@M-b+2x;>Fid9s{$n|+p8A*LQC9czBrERItd~C( zgR~wMc~A7pi)~)|HD6>4@{RIxT)Uh|jAq5Sc8*xnFA6{8bv1GESaUY{$LVlbO?g-T z)l8kSCR`(L_*-2g+mJm-vrNj(u`QieRXQP|@qG?_C(akyo^&|dzObJyu_bYt9ZKGF ze)gScE^6qXXfABdOP@qLXFu?TI}&uc?CD6teNszZh}lbs*-41mM~LYeV$SCAmiT!G zT?yZgFn5j*-4$jI&II02g!h79kI152$XR-WR~t;t&xV`b!2eww)UzAn z@SUnFTH&?_85`-gOD{;fLo= z2%C3CjE{Qw@rCvD+mz5R^kZ*HsPX_uL42v#si{=hAaz$*NcT=W*BG1eTh7;Hv#{z zXIeUt3kQm-zq;!0s_L3f`*7N^UDK)hN$;qE}pObJG9vMPe`;WyQF1F_A4u=i|ybVs-IJ_%W->i93%~TN8hj42NaS zyY!w@J7ZP2Lfr60w??)mdyu55l&fP?Iw{L!LO|o&34F)S7ui?ou(my6Ut3p|(ffDF zNBU>pUDZX-hO;AQ(?4_WYUa!bzHml@I;-dOL-tIAnNx_F3z&B&o|%N0d4!nW!hCx3 zcwPDI!M=nKtR=H+fWA#(c7Pc4zA*RFBY0IO><@nGrDyUFWX(!*h94njQJQm)mm3WK zMZ*mT{3qfdpILDZ->GgTmsxsyLH4U6*dI8!gBfv;$IQrM?qo*bE8RgXeB42A_=tzs z`z}>U-EZ9^eqw$r7Ms0WmGF@dKfW-Zy4NM#7y7Z;6LP%&mpF)Lf63vT8mM_of}dDo z@T0dS)Z{tvV}FtSkDLF19L}j347ca~t5-Oa_;+A;ISb!`-SY3iUhxj!q|bGCy%+WS J5y!dbd;|8{U?%_o literal 0 HcmV?d00001 diff --git a/copyTest.cl b/copyTest.cl new file mode 100644 index 0000000..771b508 --- /dev/null +++ b/copyTest.cl @@ -0,0 +1,17 @@ +__kernel void copyTest(__global uchar *_heap_base) +{ + int i_8, i_7, i_1, i_2; + ulong ul_0, ul_6; + long l_3, l_5, l_4; + + __global ulong *_frame = (__global ulong *) &_heap_base[0]; + + ul_0 = (ulong) _frame[3]; + i_1 = get_global_id(0); + i_2 = i_1; + l_3 = (long) i_2; + l_4 = l_3 << 3; // Long buffer + l_5 = l_4 + 16L; // Randomly starting in position 16 + ul_6 = ul_0 + l_5; + *((__global int *) ul_6) = 555; +} diff --git a/copyTest.spv b/copyTest.spv new file mode 100644 index 0000000000000000000000000000000000000000..27d8f84a4e4f59a9b2c0648fdb2f4208aee4f940 GIT binary patch literal 2108 zcma)+T~8B16o#j4OI1KbMEtgZii#rN6j4Jou^}-rCf;urT0@hyt=pDB;)TD!U*(VT z#>DrT-Gif7ZZbLNeb2|9nKNgik%g)&m)u1+>R!0~n{p$rEEQQPE%~(7bO4AM+U_owUS(gFelU{aLr4Q41MZ)9R9B<_~a@VA!@w$(@ z?)6Qt-|+f|kGUy5{*)!3eXSLDM>xr{9wb@)_PNF#fx;QwR2v)|NE>#f)R;-`lE?e{)BEiXE3EjyjlWY5`LR`wubq zB*fg45OYrqQ#bb{#N3k*b5BCdJqa=Q#4z=7Pr$FVFKT5E?2A3W)K1wK^MTFyxAKMO zO~A%oAf8-s=4bJD#bGo4tFqUG=d`eq54@lSH?Ku3{kW$EPb~lBKb!%)arkma?rYH} zc+UBQ5IS!XHa_AxV>oQKPxM7CY{uiu`NJE>ztA6QVKY8cChFz=z(y>+0OLdLwx`F^ zu^InEzQi(<*q9ySZKlo7;-AWn&G>IcJntzs=8SlH4`+TBzaow^{FmnqC&qZoyC#kp z?BwOn!I_`sT^EPV_%AvW^bIX+C&J9uR>6@Aoi_kKdhttBR@bsQs~5I#&Un9=v*(gK O%E3JDYB7`8EXQvR+KW~I literal 0 HcmV?d00001 diff --git a/copy_data.cl b/copy_data.cl deleted file mode 100644 index 308de85..0000000 --- a/copy_data.cl +++ /dev/null @@ -1,5 +0,0 @@ -__kernel void copydata(__global int* input, __global int* output) { - uint idx = get_global_id(0); - output[idx] = input[idx]; -} - diff --git a/levelZeroLib/CMakeLists.txt b/levelZeroLib/CMakeLists.txt index 2100ccf..f1478bf 100644 --- a/levelZeroLib/CMakeLists.txt +++ b/levelZeroLib/CMakeLists.txt @@ -7,24 +7,68 @@ find_package(JNI REQUIRED) # message("- ZE_SHARED_LOADER found") #else() # message("ERROR: ZE_SHARED_LOADER is not set") +#if(MSVC) +# message("Set this export to PATH-TO\level-zero\build\lib\release\ze_loader.lib") +#else() # message("Set this export to PATH-TO/level-zero/build/lib/libze_loader.so") +#endif() # return() #endif() # set the C++14 standard set(CMAKE_CXX_STANDARD 14) -set(GCC_INTEL_LEVEL0 "-std=c++14 -O2 -fpermissive -rdynamic -fPIC -lstdc++ -fno-stack-protector") +if(MSVC) + string(REPLACE "/MDd" "/MTd" CMAKE_CXX_FLAGS_DEBUG "${CMAKE_CXX_FLAGS_DEBUG}") + string(REPLACE "/MD" "/MT" CMAKE_CXX_FLAGS_RELEASE "${CMAKE_CXX_FLAGS_RELEASE}") -# This should point to: "PATH-TO/level-zero/build/lib/libze_loader.so" -set(ZE_LOADER $ENV{ZE_SHARED_LOADER}) + # enable multi-process compilation, not supported by clang-cl + if(NOT CMAKE_CXX_COMPILER_ID STREQUAL Clang) + set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} /MP") + endif() + + # enable exceptions handling + set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} /EHsc") + + # enable creation of PDB files for Release Builds + set(CMAKE_CXX_FLAGS_RELEASE "${CMAKE_CXX_FLAGS_RELEASE} /Zi") + set(CMAKE_SHARED_LINKER_FLAGS_RELEASE "${CMAKE_SHARED_LINKER_FLAGS_RELEASE} /DEBUG /OPT:REF /OPT:ICF") -add_definitions(${GCC_INTEL_LEVEL0}) + # enable CET shadow stack + set(CMAKE_SHARED_LINKER_FLAGS "${CMAKE_SHARED_LINKER_FLAGS} /CETCOMPAT") + + #Use of sccache with MSVC requires workaround of replacing /Zi with /Z7 + #https://github.com/mozilla/sccache + if(USE_Z7) #sccache + string(REPLACE "/Zi" "/Z7" CMAKE_CXX_FLAGS_DEBUG "${CMAKE_CXX_FLAGS_DEBUG}") + string(REPLACE "/Zi" "/Z7" CMAKE_C_FLAGS_DEBUG "${CMAKE_C_FLAGS_DEBUG}") + string(REPLACE "/Zi" "/Z7" CMAKE_CXX_FLAGS_RELEASE "${CMAKE_CXX_FLAGS_RELEASE}") + string(REPLACE "/Zi" "/Z7" CMAKE_C_FLAGS_RELEASE "${CMAKE_C_FLAGS_RELEASE}") + string(REPLACE "/Zi" "/Z7" CMAKE_CXX_FLAGS_RELWITHDEBINFO "${CMAKE_CXX_FLAGS_RELWITHDEBINFO}") + string(REPLACE "/Zi" "/Z7" CMAKE_C_FLAGS_RELWITHDEBINFO "${CMAKE_C_FLAGS_RELWITHDEBINFO}") + endif() +else() + set(GCC_INTEL_LEVEL0 "-std=c++14") + set(GCC_INTEL_LEVEL0 "${GCC_INTEL_LEVEL0} -O2") + set(GCC_INTEL_LEVEL0 "${GCC_INTEL_LEVEL0} -fpermissive") + set(GCC_INTEL_LEVEL0 "${GCC_INTEL_LEVEL0} -fno-stack-protector") + set(GCC_INTEL_LEVEL0 "${GCC_INTEL_LEVEL0} -fPIC") + set(GCC_INTEL_LEVEL0 "${GCC_INTEL_LEVEL0} -rdynamic") + set(GCC_INTEL_LEVEL0 "${GCC_INTEL_LEVEL0} -lstdc++") + add_definitions(${GCC_INTEL_LEVEL0}) +endif() + +# This should point to: +# Linos: "PATH-TO/level-zero/build/lib/libze_loader.so" +# Winos: "PATH-TO\level-zero\build\lib\release\ze_loader.lib" +set(ZE_LOADER $ENV{ZE_SHARED_LOADER}) include_directories( ./ ${JNI_INCLUDE_DIRS} ${OPENCL_INCLUDE_DIRS} + $ENV{CPLUS_INCLUDE_PATH} + $ENV{C_INCLUDE_PATH} ) file(GLOB_RECURSE "src/*.cpp") @@ -36,10 +80,11 @@ add_library(tornado-levelzero SHARED src/levelZeroModule.cpp src/levelZeroKernel.cpp src/levelZeroCommandList.cpp - src/levelZeroCommandQueue.cpp + src/levelZeroCommandQueue.cpp src/levelZeroFence.cpp src/levelZeroTimeStampKernel.cpp - src/levelZeroDescriptors.cpp) + src/levelZeroDescriptors.cpp + ) target_link_libraries(tornado-levelzero ${OpenCL_LIBRARIES} ${JNI_LIB_DIRS}) target_link_libraries(tornado-levelzero ${ZE_LOADER}) diff --git a/levelZeroLib/src/levelZeroBuffer.cpp b/levelZeroLib/src/levelZeroBuffer.cpp index fbc54e9..0ad8d52 100644 --- a/levelZeroLib/src/levelZeroBuffer.cpp +++ b/levelZeroLib/src/levelZeroBuffer.cpp @@ -140,7 +140,11 @@ JNIEXPORT void JNICALL Java_uk_ac_manchester_tornado_drivers_spirv_levelzero_Lev jbyte *offHeapByteArray = static_cast(buffer); jbyte *arrayByte = env->GetByteArrayElements(javaArray, 0); int size = env->GetArrayLength(javaArray); +#ifdef _WIN32 + _memccpy(offHeapByteArray, arrayByte, 0, size); +#else memccpy(offHeapByteArray, arrayByte, 0, size); +#endif } /* @@ -169,10 +173,17 @@ JNIEXPORT void JNICALL Java_uk_ac_manchester_tornado_drivers_spirv_levelzero_Lev jfieldID fieldPointer = env->GetFieldID(klass, "ptrBuffer", "J"); jlong ptr = env->GetLongField(javaBufferObject, fieldPointer); +#ifdef _WIN32 + int64_t *buffer = nullptr; + if (ptr != -1) { + buffer = reinterpret_cast(ptr); + } +#else long *buffer = nullptr; if (ptr != -1) { buffer = reinterpret_cast(ptr); } +#endif for (int i = 0; i < size; i++) { buffer[i] = value; } @@ -217,7 +228,11 @@ JNIEXPORT void JNICALL Java_uk_ac_manchester_tornado_drivers_spirv_levelzero_Lev jlong *offHeapByteArray = static_cast(buffer); jlong *arrayByte = env->GetLongArrayElements(array, 0); int size = env->GetArrayLength(array); +#ifdef _WIN32 + _memccpy(offHeapByteArray, arrayByte, 0, size); +#else memccpy(offHeapByteArray, arrayByte, 0, size); +#endif } /* diff --git a/levelZeroLib/src/levelZeroCommandList.cpp b/levelZeroLib/src/levelZeroCommandList.cpp index 19c1555..d18d2a6 100644 --- a/levelZeroLib/src/levelZeroCommandList.cpp +++ b/levelZeroLib/src/levelZeroCommandList.cpp @@ -58,7 +58,11 @@ JNIEXPORT jint JNICALL Java_uk_ac_manchester_tornado_drivers_spirv_levelzero_Lev if (javaSignalEvent != nullptr) { jclass signalEventClass = env->GetObjectClass(javaSignalEvent); jfieldID fieldSignal = env->GetFieldID(signalEventClass, "ptrZeEventHandle", "J"); +#ifdef _WIN32 + int64_t eventSignalPtr = env->GetLongField(javaSignalEvent, fieldSignal); +#else long eventSignalPtr = env->GetLongField(javaSignalEvent, fieldSignal); +#endif signalEvent = reinterpret_cast(eventSignalPtr); } @@ -470,7 +474,11 @@ JNIEXPORT jint JNICALL Java_uk_ac_manchester_tornado_drivers_spirv_levelzero_Lev if (javaEventHandler != nullptr) { jclass signalEventClass = env->GetObjectClass(javaEventHandler); jfieldID fieldSignal = env->GetFieldID(signalEventClass, "ptrZeEventHandle", "J"); +#ifdef _WIN32 + int64_t eventSignalPtr = env->GetLongField(javaEventHandler, fieldSignal); +#else long eventSignalPtr = env->GetLongField(javaEventHandler, fieldSignal); +#endif events = reinterpret_cast(eventSignalPtr); } @@ -515,6 +523,7 @@ JNIEXPORT jint JNICALL Java_uk_ac_manchester_tornado_drivers_spirv_levelzero_Lev } ze_result_t result = zeCommandListAppendWriteGlobalTimestamp(commandList, (uint64_t *) timestampBuffer ,nullptr, numWaitEvents, nullptr); + LOG_ZE_JNI("zeCommandListAppendWriteGlobalTimestamp", result); return result; } @@ -531,8 +540,13 @@ JNIEXPORT jint JNICALL Java_uk_ac_manchester_tornado_drivers_spirv_levelzero_Lev jclass classLevelZeroIntegerBuffer = env->GetObjectClass(levelZeroBufferInteger); jfieldID fieldBufferPtr = env->GetFieldID(classLevelZeroIntegerBuffer, "ptrBuffer", "J"); +#ifdef _WIN32 + int64_t bufferPtr = env->GetLongField(levelZeroBufferInteger, fieldBufferPtr); + const int64_t *ptr = reinterpret_cast(bufferPtr); +#else long bufferPtr = env->GetLongField(levelZeroBufferInteger, fieldBufferPtr); const void *ptr = reinterpret_cast(bufferPtr); +#endif ze_result_t result = zeCommandListAppendMemoryPrefetch(commandList, ptr, size); LOG_ZE_JNI("zeCommandListAppendMemoryPrefetch", result); @@ -555,8 +569,13 @@ JNIEXPORT jint JNICALL Java_uk_ac_manchester_tornado_drivers_spirv_levelzero_Lev jclass classLevelZeroIntegerBuffer = env->GetObjectClass(levelZeroBufferInteger); jfieldID fieldBufferPtr = env->GetFieldID(classLevelZeroIntegerBuffer, "ptrBuffer", "J"); +#ifdef _WIN32 + int64_t bufferPtr = env->GetLongField(levelZeroBufferInteger, fieldBufferPtr); + const int64_t *ptr = reinterpret_cast(bufferPtr); +#else long bufferPtr = env->GetLongField(levelZeroBufferInteger, fieldBufferPtr); const void *ptr = reinterpret_cast(bufferPtr); +#endif ze_result_t result = zeCommandListAppendMemAdvise(commandList, deviceHandle, ptr, size, memoryAdvice); LOG_ZE_JNI("zeCommandListAppendMemAdvise", result); diff --git a/levelZeroLib/src/levelZeroContext.cpp b/levelZeroLib/src/levelZeroContext.cpp index 722f023..553c337 100644 --- a/levelZeroLib/src/levelZeroContext.cpp +++ b/levelZeroLib/src/levelZeroContext.cpp @@ -51,7 +51,11 @@ JNIEXPORT jint JNICALL Java_uk_ac_manchester_tornado_drivers_spirv_levelzero_Lev } jfieldID fieldDescriptorPointer = env->GetFieldID(descriptorClass, "nativePointer", "J"); +#ifdef _WIN32 + int64_t valuePointerDescriptor = env->GetLongField(descriptorObject, fieldDescriptorPointer); +#else long valuePointerDescriptor = env->GetLongField(descriptorObject, fieldDescriptorPointer); +#endif ze_context_desc_t contextDesc = {}; ze_context_desc_t *contextDescPtr; @@ -67,7 +71,11 @@ JNIEXPORT jint JNICALL Java_uk_ac_manchester_tornado_drivers_spirv_levelzero_Lev contextJavaArray[0] = reinterpret_cast(context); env->ReleaseLongArrayElements(contextArray, contextJavaArray, 0); +#ifdef _WIN32 + valuePointerDescriptor = reinterpret_cast(&(contextDesc)); +#else valuePointerDescriptor = reinterpret_cast(&(contextDesc)); +#endif env->SetLongField(descriptorObject, fieldDescriptorPointer, valuePointerDescriptor); return result; @@ -98,7 +106,11 @@ JNIEXPORT jint JNICALL Java_uk_ac_manchester_tornado_drivers_spirv_levelzero_Lev ze_command_queue_desc_t cmdQueueDesc = {}; jclass commandDescriptorClass = env->GetObjectClass(javaCommandQueueDescriptor); field = env->GetFieldID(commandDescriptorClass, "ptrZeCommandDescriptor", "J"); +#ifdef _WIN32 + int64_t ptrZeCommandDescriptor = env->GetLongField(javaCommandQueueDescriptor, field); +#else long ptrZeCommandDescriptor = env->GetLongField(javaCommandQueueDescriptor, field); +#endif if (ptrZeCommandDescriptor != -1) { ze_command_queue_desc_t *cmdQueueDescPtr = reinterpret_cast(ptrZeCommandDescriptor); cmdQueueDesc = *cmdQueueDescPtr; @@ -108,10 +120,18 @@ JNIEXPORT jint JNICALL Java_uk_ac_manchester_tornado_drivers_spirv_levelzero_Lev int type = env->GetIntField(javaCommandQueueDescriptor, field); field = env->GetFieldID(commandDescriptorClass, "ordinal", "J"); +#ifdef _WIN32 + int64_t ordinal = env->GetLongField(javaCommandQueueDescriptor, field); +#else long ordinal = env->GetLongField(javaCommandQueueDescriptor, field); +#endif field = env->GetFieldID(commandDescriptorClass, "index", "J"); +#ifdef _WIN32 + int64_t index = env->GetLongField(javaCommandQueueDescriptor, field); +#else int index = env->GetLongField(javaCommandQueueDescriptor, field); +#endif field = env->GetFieldID(commandDescriptorClass, "mode", "I"); int mode = env->GetIntField(javaCommandQueueDescriptor, field); @@ -185,7 +205,11 @@ JNIEXPORT jint JNICALL Java_uk_ac_manchester_tornado_drivers_spirv_levelzero_Lev ze_command_list_desc_t cmdListDesc = {}; jclass commanddescriptorClass = env->GetObjectClass(javaCommandListDescriptor); jfieldID field = env->GetFieldID(commanddescriptorClass, "ptrZeCommandListDescriptor", "J"); +#ifdef _WIN32 + int64_t ptrZeCommandDescriptor = env->GetLongField(javaCommandListDescriptor, field); +#else long ptrZeCommandDescriptor = env->GetLongField(javaCommandListDescriptor, field); +#endif if (ptrZeCommandDescriptor != -1) { ze_command_list_desc_t *cmdListDescPtr = reinterpret_cast(ptrZeCommandDescriptor); cmdListDesc = *cmdListDescPtr; @@ -195,7 +219,11 @@ JNIEXPORT jint JNICALL Java_uk_ac_manchester_tornado_drivers_spirv_levelzero_Lev int type = env->GetIntField(javaCommandListDescriptor, field); field = env->GetFieldID(commanddescriptorClass, "commandQueueGroupOrdinal", "J"); +#ifdef _WIN32 + int64_t ordinal = env->GetLongField(javaCommandListDescriptor, field); +#else long ordinal = env->GetLongField(javaCommandListDescriptor, field); +#endif cmdListDesc.stype = static_cast(type); cmdListDesc.commandQueueGroupOrdinal = ordinal; @@ -250,7 +278,11 @@ JNIEXPORT jint JNICALL Java_uk_ac_manchester_tornado_drivers_spirv_levelzero_Lev ze_command_queue_desc_t commandQueueDesc = {}; jclass commanddescriptorClass = env->GetObjectClass(javaCommandQueueDescriptor); field = env->GetFieldID(commanddescriptorClass, "ptrZeCommandDescriptor", "J"); +#ifdef _WIN32 + int64_t ptrZeCommandDescriptor = env->GetLongField(javaCommandQueueDescriptor, field); +#else long ptrZeCommandDescriptor = env->GetLongField(javaCommandQueueDescriptor, field); +#endif if (ptrZeCommandDescriptor != -1) { ze_command_queue_desc_t *cmdQueueDescPtr = reinterpret_cast(ptrZeCommandDescriptor); commandQueueDesc = *cmdQueueDescPtr; @@ -260,7 +292,11 @@ JNIEXPORT jint JNICALL Java_uk_ac_manchester_tornado_drivers_spirv_levelzero_Lev int type = env->GetIntField(javaCommandQueueDescriptor, field); field = env->GetFieldID(commanddescriptorClass, "ordinal", "J"); +#ifdef _WIN32 + int64_t ordinal = env->GetLongField(javaCommandQueueDescriptor, field); +#else long ordinal = env->GetLongField(javaCommandQueueDescriptor, field); +#endif field = env->GetFieldID(commanddescriptorClass, "index", "J"); int index = env->GetIntField(javaCommandQueueDescriptor, field); @@ -331,12 +367,21 @@ JNIEXPORT jint JNICALL Java_uk_ac_manchester_tornado_drivers_spirv_levelzero_Lev jfieldID fieldTypeDeviceDesc = env->GetFieldID(javaDeviceMemAllocDescClass, "stype", "I"); int typeDeviceDesc = env->GetIntField(javaDeviceMemAllocDesc, fieldTypeDeviceDesc); jfieldID fieldFlagsDeviceDesc = env->GetFieldID(javaDeviceMemAllocDescClass, "flags", "J"); +#ifdef _WIN32 + int64_t flagDeviceDesc = env->GetLongField(javaDeviceMemAllocDesc, fieldFlagsDeviceDesc); + jfieldID fieldOrdinalDeviceDesc = env->GetFieldID(javaDeviceMemAllocDescClass, "ordinal", "J"); + int64_t ordinalDeviceDesc = env->GetLongField(javaDeviceMemAllocDesc, fieldOrdinalDeviceDesc); + + jfieldID fieldPNextMemAlloc = env->GetFieldID(javaDeviceMemAllocDescClass, "pNext", "J"); + int64_t pnextDeviceAlloc = env->GetLongField(javaDeviceMemAllocDesc, fieldPNextMemAlloc); +#else long flagDeviceDesc = env->GetLongField(javaDeviceMemAllocDesc, fieldFlagsDeviceDesc); jfieldID fieldOrdinalDeviceDesc = env->GetFieldID(javaDeviceMemAllocDescClass, "ordinal", "J"); long ordinalDeviceDesc = env->GetLongField(javaDeviceMemAllocDesc, fieldOrdinalDeviceDesc); jfieldID fieldPNextMemAlloc = env->GetFieldID(javaDeviceMemAllocDescClass, "pNext", "J"); ulong pnextDeviceAlloc = env->GetLongField(javaDeviceMemAllocDesc, fieldPNextMemAlloc); +#endif ze_device_mem_alloc_desc_t deviceDesc = {}; @@ -351,10 +396,17 @@ JNIEXPORT jint JNICALL Java_uk_ac_manchester_tornado_drivers_spirv_levelzero_Lev jfieldID fieldTypeHostDesc = env->GetFieldID(javaHostMemAllocDescClass, "stype", "I"); int typeHostDesc = env->GetIntField(javaHostMemAllocDesc, fieldTypeHostDesc); jfieldID fieldFlagsHostDesc = env->GetFieldID(javaHostMemAllocDescClass, "flags", "J"); +#ifdef _WIN32 + int64_t flagsHostDesc = env->GetLongField(javaHostMemAllocDesc, fieldFlagsHostDesc); + + jfieldID fieldPNextHostAlloc = env->GetFieldID(javaDeviceMemAllocDescClass, "pNext", "J"); + int64_t pnextHostAlloc = env->GetLongField(javaDeviceMemAllocDesc, fieldPNextHostAlloc); +#else long flagsHostDesc = env->GetLongField(javaHostMemAllocDesc, fieldFlagsHostDesc); jfieldID fieldPNextHostAlloc = env->GetFieldID(javaDeviceMemAllocDescClass, "pNext", "J"); ulong pnextHostAlloc = env->GetLongField(javaDeviceMemAllocDesc, fieldPNextHostAlloc); +#endif ze_host_mem_alloc_desc_t hostDesc; hostDesc.stype = static_cast(typeHostDesc); @@ -391,12 +443,21 @@ JNIEXPORT jint JNICALL Java_uk_ac_manchester_tornado_drivers_spirv_levelzero_Lev jfieldID fieldTypeDeviceDesc = env->GetFieldID(javaDeviceMemAllocDescClass, "stype", "I"); int typeDeviceDesc = env->GetIntField(javaDeviceMemAllocDesc, fieldTypeDeviceDesc); jfieldID fieldFlagsDeviceDesc = env->GetFieldID(javaDeviceMemAllocDescClass, "flags", "J"); +#ifdef _WIN32 + int64_t flagDeviceDesc = env->GetLongField(javaDeviceMemAllocDesc, fieldFlagsDeviceDesc); + jfieldID fieldOrdinalDeviceDesc = env->GetFieldID(javaDeviceMemAllocDescClass, "ordinal", "J"); + int64_t ordinalDeviceDesc = env->GetLongField(javaDeviceMemAllocDesc, fieldOrdinalDeviceDesc); + + jfieldID fieldPNextMemAlloc = env->GetFieldID(javaDeviceMemAllocDescClass, "pNext", "J"); + int64_t pnextDeviceAlloc = env->GetLongField(javaDeviceMemAllocDesc, fieldPNextMemAlloc); +#else long flagDeviceDesc = env->GetLongField(javaDeviceMemAllocDesc, fieldFlagsDeviceDesc); jfieldID fieldOrdinalDeviceDesc = env->GetFieldID(javaDeviceMemAllocDescClass, "ordinal", "J"); long ordinalDeviceDesc = env->GetLongField(javaDeviceMemAllocDesc, fieldOrdinalDeviceDesc); jfieldID fieldPNextMemAlloc = env->GetFieldID(javaDeviceMemAllocDescClass, "pNext", "J"); ulong pnextDeviceAlloc = env->GetLongField(javaDeviceMemAllocDesc, fieldPNextMemAlloc); +#endif ze_device_mem_alloc_desc_t deviceDesc = {}; deviceDesc.stype = static_cast(typeDeviceDesc); @@ -411,10 +472,17 @@ JNIEXPORT jint JNICALL Java_uk_ac_manchester_tornado_drivers_spirv_levelzero_Lev jfieldID fieldTypeHostDesc = env->GetFieldID(javaHostMemAllocDescClass, "stype", "I"); int typeHostDesc = env->GetIntField(javaHostMemAllocDesc, fieldTypeHostDesc); jfieldID fieldFlagsHostDesc = env->GetFieldID(javaHostMemAllocDescClass, "flags", "J"); +#ifdef _WIN32 + int64_t flagsHostDesc = env->GetLongField(javaHostMemAllocDesc, fieldFlagsHostDesc); + + jfieldID fieldPNextHostAlloc = env->GetFieldID(javaHostMemAllocDescClass, "pNext", "J"); + int64_t pnextHostAlloc = env->GetLongField(javaHostMemAllocDesc, fieldPNextHostAlloc); +#else long flagsHostDesc = env->GetLongField(javaHostMemAllocDesc, fieldFlagsHostDesc); jfieldID fieldPNextHostAlloc = env->GetFieldID(javaHostMemAllocDescClass, "pNext", "J"); ulong pnextHostAlloc = env->GetLongField(javaHostMemAllocDesc, fieldPNextHostAlloc); +#endif ze_host_mem_alloc_desc_t hostDesc; hostDesc.stype = static_cast(typeHostDesc); @@ -452,12 +520,21 @@ JNIEXPORT jint JNICALL Java_uk_ac_manchester_tornado_drivers_spirv_levelzero_Lev jfieldID fieldTypeDeviceDesc = env->GetFieldID(javaDeviceMemAllocDescClass, "stype", "I"); int typeDeviceDesc = env->GetIntField(javaDeviceMemAllocDesc, fieldTypeDeviceDesc); jfieldID fieldFlagsDeviceDesc = env->GetFieldID(javaDeviceMemAllocDescClass, "flags", "J"); +#ifdef _WIN32 + int64_t flagDeviceDesc = env->GetLongField(javaDeviceMemAllocDesc, fieldFlagsDeviceDesc); + jfieldID fieldOrdinalDeviceDesc = env->GetFieldID(javaDeviceMemAllocDescClass, "ordinal", "J"); + int64_t ordinalDeviceDesc = env->GetLongField(javaDeviceMemAllocDesc, fieldOrdinalDeviceDesc); + + jfieldID fieldPNextMemAlloc = env->GetFieldID(javaDeviceMemAllocDescClass, "pNext", "J"); + int64_t pnextDeviceAlloc = env->GetLongField(javaDeviceMemAllocDesc, fieldPNextMemAlloc); +#else long flagDeviceDesc = env->GetLongField(javaDeviceMemAllocDesc, fieldFlagsDeviceDesc); jfieldID fieldOrdinalDeviceDesc = env->GetFieldID(javaDeviceMemAllocDescClass, "ordinal", "J"); long ordinalDeviceDesc = env->GetLongField(javaDeviceMemAllocDesc, fieldOrdinalDeviceDesc); jfieldID fieldPNextMemAlloc = env->GetFieldID(javaDeviceMemAllocDescClass, "pNext", "J"); ulong pnextDeviceAlloc = env->GetLongField(javaDeviceMemAllocDesc, fieldPNextMemAlloc); +#endif ze_device_mem_alloc_desc_t deviceDesc = {}; deviceDesc.stype = static_cast(typeDeviceDesc); @@ -474,7 +551,11 @@ JNIEXPORT jint JNICALL Java_uk_ac_manchester_tornado_drivers_spirv_levelzero_Lev // Set Buffer Pointer and attributes jfieldID fieldBufferSize = env->GetFieldID(javaBufferClass, "size", "J"); jfieldID alignmentField = env->GetFieldID(javaBufferClass, "alignment", "J"); +#ifdef _WIN32 + env->SetLongField(javaLevelZeroBuffer, fieldBuffer, reinterpret_cast(buffer)); +#else env->SetLongField(javaLevelZeroBuffer, fieldBuffer, reinterpret_cast(buffer)); +#endif env->SetLongField(javaLevelZeroBuffer, fieldBufferSize, allocSize); env->SetLongField(javaLevelZeroBuffer, alignmentField, alignment); @@ -499,12 +580,21 @@ JNIEXPORT jint JNICALL Java_uk_ac_manchester_tornado_drivers_spirv_levelzero_Lev jfieldID fieldTypeDeviceDesc = env->GetFieldID(javaDeviceMemAllocDescClass, "stype", "I"); int typeDeviceDesc = env->GetIntField(javaDeviceMemAllocDesc, fieldTypeDeviceDesc); jfieldID fieldFlagsDeviceDesc = env->GetFieldID(javaDeviceMemAllocDescClass, "flags", "J"); +#ifdef _WIN32 + int64_t flagDeviceDesc = env->GetLongField(javaDeviceMemAllocDesc, fieldFlagsDeviceDesc); + jfieldID fieldOrdinalDeviceDesc = env->GetFieldID(javaDeviceMemAllocDescClass, "ordinal", "J"); + int64_t ordinalDeviceDesc = env->GetLongField(javaDeviceMemAllocDesc, fieldOrdinalDeviceDesc); + + jfieldID fieldPNextMemAlloc = env->GetFieldID(javaDeviceMemAllocDescClass, "pNext", "J"); + int64_t pnextDeviceAlloc = env->GetLongField(javaDeviceMemAllocDesc, fieldPNextMemAlloc); +#else long flagDeviceDesc = env->GetLongField(javaDeviceMemAllocDesc, fieldFlagsDeviceDesc); jfieldID fieldOrdinalDeviceDesc = env->GetFieldID(javaDeviceMemAllocDescClass, "ordinal", "J"); long ordinalDeviceDesc = env->GetLongField(javaDeviceMemAllocDesc, fieldOrdinalDeviceDesc); jfieldID fieldPNextMemAlloc = env->GetFieldID(javaDeviceMemAllocDescClass, "pNext", "J"); ulong pnextDeviceAlloc = env->GetLongField(javaDeviceMemAllocDesc, fieldPNextMemAlloc); +#endif ze_device_mem_alloc_desc_t deviceDesc = {}; @@ -522,7 +612,11 @@ JNIEXPORT jint JNICALL Java_uk_ac_manchester_tornado_drivers_spirv_levelzero_Lev // Set Buffer Pointer and attributes jfieldID fieldBufferSize = env->GetFieldID(javaBufferClass, "size", "J"); jfieldID alignmentField = env->GetFieldID(javaBufferClass, "alignment", "J"); +#ifdef _WIN32 + env->SetLongField(javaLevelZeroBuffer, fieldBuffer, reinterpret_cast(buffer)); +#else env->SetLongField(javaLevelZeroBuffer, fieldBuffer, reinterpret_cast(buffer)); +#endif env->SetLongField(javaLevelZeroBuffer, fieldBufferSize, allocSize); env->SetLongField(javaLevelZeroBuffer, alignmentField, alignment); @@ -713,7 +807,11 @@ JNIEXPORT jint JNICALL Java_uk_ac_manchester_tornado_drivers_spirv_levelzero_Lev field = env->GetFieldID(klassEventPoolDescriptor, "flags", "I"); jint flags = env->GetIntField(javaEventPoolDescriptor, field); field = env->GetFieldID(klassEventPoolDescriptor, "pNext", "J"); +#ifdef _WIN32 + jlong pNext = env->GetLongField(javaEventPoolDescriptor, field); +#else jint pNext = env->GetLongField(javaEventPoolDescriptor, field); +#endif ze_event_pool_desc_t eventPoolDescriptor = {}; eventPoolDescriptor.stype = static_cast(stype); @@ -750,7 +848,11 @@ JNIEXPORT jint JNICALL Java_uk_ac_manchester_tornado_drivers_spirv_levelzero_Lev ze_event_pool_handle_t eventPool; jclass eventPoolClass = env->GetObjectClass(javaEventPoolHandler); jfieldID field = env->GetFieldID(eventPoolClass, "ptrZeEventPoolHandle", "J"); +#ifdef _WIN32 + int64_t ptrEventPool = env->GetLongField(javaEventPoolHandler, field); +#else long ptrEventPool = env->GetLongField(javaEventPoolHandler, field); +#endif if (ptrEventPool != -1) { eventPool = reinterpret_cast(ptrEventPool); } else { @@ -760,6 +862,15 @@ JNIEXPORT jint JNICALL Java_uk_ac_manchester_tornado_drivers_spirv_levelzero_Lev jclass klassEventDesc = env->GetObjectClass(javaEventDescriptor); field = env->GetFieldID(klassEventDesc, "stype", "I"); +#ifdef _WIN32 + int64_t stype = env->GetLongField(javaEventDescriptor, field); + field = env->GetFieldID(klassEventDesc, "index", "J"); + jlong index = env->GetLongField(javaEventDescriptor, field); + field = env->GetFieldID(klassEventDesc, "signal", "I"); + jlong signal = env->GetLongField(javaEventDescriptor, field); + field = env->GetFieldID(klassEventDesc, "wait", "I"); + jlong wait = env->GetLongField(javaEventDescriptor, field); +#else int stype = env->GetLongField(javaEventDescriptor, field); field = env->GetFieldID(klassEventDesc, "index", "J"); jint index = env->GetLongField(javaEventDescriptor, field); @@ -767,6 +878,7 @@ JNIEXPORT jint JNICALL Java_uk_ac_manchester_tornado_drivers_spirv_levelzero_Lev jint signal = env->GetLongField(javaEventDescriptor, field); field = env->GetFieldID(klassEventDesc, "wait", "I"); jint wait = env->GetLongField(javaEventDescriptor, field); +#endif ze_event_desc_t eventDesc = {}; eventDesc.stype = static_cast(stype); @@ -862,10 +974,17 @@ JNIEXPORT jint JNICALL Java_uk_ac_manchester_tornado_drivers_spirv_levelzero_Lev jfieldID fieldTypeDeviceDesc = env->GetFieldID(javaHostDescClass, "stype", "I"); int typeDeviceDesc = env->GetIntField(javaHostMemAllocDesc, fieldTypeDeviceDesc); jfieldID fieldFlagsDeviceDesc = env->GetFieldID(javaHostDescClass, "flags", "J"); +#ifdef _WIN32 + int64_t flagDeviceDesc = env->GetLongField(javaHostMemAllocDesc, fieldFlagsDeviceDesc); + + jfieldID fieldPNextMemAlloc = env->GetFieldID(javaHostDescClass, "pNext", "J"); + int64_t pnextHostAlloc = env->GetLongField(javaHostMemAllocDesc, fieldPNextMemAlloc); +#else long flagDeviceDesc = env->GetLongField(javaHostMemAllocDesc, fieldFlagsDeviceDesc); jfieldID fieldPNextMemAlloc = env->GetFieldID(javaHostDescClass, "pNext", "J"); ulong pnextHostAlloc = env->GetLongField(javaHostMemAllocDesc, fieldPNextMemAlloc); +#endif ze_host_mem_alloc_desc_t hostDescriptor = {}; hostDescriptor.stype = static_cast(typeDeviceDesc); @@ -874,7 +993,11 @@ JNIEXPORT jint JNICALL Java_uk_ac_manchester_tornado_drivers_spirv_levelzero_Lev if (pnextHostAlloc != -1) { +#ifdef _WIN32 + hostDescriptor.pNext = reinterpret_cast(pnextHostAlloc); +#else hostDescriptor.pNext = reinterpret_cast(pnextHostAlloc); +#endif } ze_result_t result = zeMemAllocHost(context, &hostDescriptor, allocSize, alignment, (void**) &buffer); diff --git a/levelZeroLib/src/levelZeroDescriptors.cpp b/levelZeroLib/src/levelZeroDescriptors.cpp index 5745603..d0223bb 100644 --- a/levelZeroLib/src/levelZeroDescriptors.cpp +++ b/levelZeroLib/src/levelZeroDescriptors.cpp @@ -49,7 +49,11 @@ JNIEXPORT void JNICALL Java_uk_ac_manchester_tornado_drivers_spirv_levelzero_ZeD ordinal }; +#ifdef _WIN32 + int64_t ptrToStruct = reinterpret_cast(descriptor); +#else ulong ptrToStruct = reinterpret_cast(descriptor); +#endif jfieldID fieldSelfPTr = env->GetFieldID(classDescriptor, "selfPtr", "J"); env->SetLongField(thisObject, fieldSelfPTr, ptrToStruct); } @@ -73,7 +77,11 @@ JNIEXPORT void JNICALL Java_uk_ac_manchester_tornado_drivers_spirv_levelzero_ZeR flags, }; +#ifdef _WIN32 + int64_t ptrToStruct = reinterpret_cast(descriptor); +#else ulong ptrToStruct = reinterpret_cast(descriptor); +#endif jfieldID fieldSelfPTr = env->GetFieldID(classDescriptor, "selfPtr", "J"); env->SetLongField(thisObject, fieldSelfPTr, ptrToStruct); } @@ -98,7 +106,11 @@ JNIEXPORT void JNICALL Java_uk_ac_manchester_tornado_drivers_spirv_levelzero_ZeR flags }; +#ifdef _WIN32 + int64_t ptrToStruct = reinterpret_cast(exceedCapacity); +#else ulong ptrToStruct = reinterpret_cast(exceedCapacity); +#endif jfieldID fieldSelfPTr = env->GetFieldID(classDescriptor, "selfPtr", "J"); env->SetLongField(thisObject, fieldSelfPTr, ptrToStruct); } @@ -126,7 +138,11 @@ JNIEXPORT void JNICALL Java_uk_ac_manchester_tornado_drivers_spirv_levelzero_ZeR count, }; +#ifdef _WIN32 + int64_t ptrToStruct = reinterpret_cast(descriptor); +#else ulong ptrToStruct = reinterpret_cast(descriptor); +#endif jfieldID fieldSelfPTr = env->GetFieldID(classDescriptor, "selfPtr", "J"); env->SetLongField(thisObject, fieldSelfPTr, ptrToStruct); } @@ -154,7 +170,11 @@ JNIEXPORT void JNICALL Java_uk_ac_manchester_tornado_drivers_spirv_levelzero_ZeR flags, }; +#ifdef _WIN32 + int64_t ptrToStruct = reinterpret_cast(descriptor); +#else ulong ptrToStruct = reinterpret_cast(descriptor); +#endif jfieldID fieldSelfPTr = env->GetFieldID(classDescriptor, "selfPtr", "J"); env->SetLongField(thisObject, fieldSelfPTr, ptrToStruct); } @@ -178,7 +198,11 @@ JNIEXPORT void JNICALL Java_uk_ac_manchester_tornado_drivers_spirv_levelzero_ZeR flags, }; +#ifdef _WIN32 + int64_t ptrToStruct = reinterpret_cast(descriptor); +#else ulong ptrToStruct = reinterpret_cast(descriptor); +#endif jfieldID fieldSelfPTr = env->GetFieldID(classDescriptor, "selfPtr", "J"); env->SetLongField(thisObject, fieldSelfPTr, ptrToStruct); } @@ -207,7 +231,11 @@ JNIEXPORT void JNICALL Java_uk_ac_manchester_tornado_drivers_spirv_levelzero_ZeR pKernelName, }; +#ifdef _WIN32 + int64_t ptrToStruct = reinterpret_cast(descriptor); +#else ulong ptrToStruct = reinterpret_cast(descriptor); +#endif jfieldID fieldSelfPTr = env->GetFieldID(classDescriptor, "selfPtr", "J"); env->SetLongField(thisObject, fieldSelfPTr, ptrToStruct); env->ReleaseStringUTFChars(javaString, pKernelName); @@ -248,7 +276,11 @@ JNIEXPORT void JNICALL Java_uk_ac_manchester_tornado_drivers_spirv_levelzero_ZeR static_cast(priority) }; +#ifdef _WIN32 + int64_t ptrToStruct = reinterpret_cast(descriptor); +#else ulong ptrToStruct = reinterpret_cast(descriptor); +#endif jfieldID fieldSelfPTr = env->GetFieldID(classDescriptor, "selfPtr", "J"); env->SetLongField(thisObject, fieldSelfPTr, ptrToStruct); } @@ -270,7 +302,11 @@ JNIEXPORT void JNICALL Java_uk_ac_manchester_tornado_drivers_spirv_levelzero_ZeR }; +#ifdef _WIN32 + int64_t ptrToStruct = reinterpret_cast(descriptor); +#else ulong ptrToStruct = reinterpret_cast(descriptor); +#endif jfieldID fieldSelfPTr = env->GetFieldID(classDescriptor, "selfPtr", "J"); env->SetLongField(thisObject, fieldSelfPTr, ptrToStruct); } @@ -302,7 +338,11 @@ JNIEXPORT void JNICALL Java_uk_ac_manchester_tornado_drivers_spirv_levelzero_ZeR wait, }; +#ifdef _WIN32 + int64_t ptrToStruct = reinterpret_cast(descriptor); +#else ulong ptrToStruct = reinterpret_cast(descriptor); +#endif jfieldID fieldSelfPTr = env->GetFieldID(classDescriptor, "selfPtr", "J"); env->SetLongField(thisObject, fieldSelfPTr, ptrToStruct); } @@ -326,7 +366,11 @@ JNIEXPORT void JNICALL Java_uk_ac_manchester_tornado_drivers_spirv_levelzero_ZeR flags, }; +#ifdef _WIN32 + int64_t ptrToStruct = reinterpret_cast(descriptor); +#else ulong ptrToStruct = reinterpret_cast(descriptor); +#endif jfieldID fieldSelfPTr = env->GetFieldID(classDescriptor, "selfPtr", "J"); env->SetLongField(thisObject, fieldSelfPTr, ptrToStruct); -} \ No newline at end of file +} diff --git a/levelZeroLib/src/levelZeroDevice.cpp b/levelZeroLib/src/levelZeroDevice.cpp index db906bc..71a92e4 100644 --- a/levelZeroLib/src/levelZeroDevice.cpp +++ b/levelZeroLib/src/levelZeroDevice.cpp @@ -260,7 +260,11 @@ JNIEXPORT jint JNICALL Java_uk_ac_manchester_tornado_drivers_spirv_levelzero_Lev if (javaMemoryPropertiesArray != nullptr) { // set the values back to Java +#ifdef _WIN32 + for (int64_t i = 0; i < memoryCount; i++) { +#else for (int i = 0; i < memoryCount; i++) { +#endif jobject javaMemoryProperty = static_cast(env->GetObjectArrayElement(javaMemoryPropertiesArray, i)); jclass descriptionClass = env->GetObjectClass(javaMemoryProperty); jfieldID field = env->GetFieldID(descriptionClass, "type", "I"); @@ -366,7 +370,11 @@ JNIEXPORT jint JNICALL Java_uk_ac_manchester_tornado_drivers_spirv_levelzero_Lev if (javaCachePropertiesArray != nullptr) { // set the values back to Java +#ifdef _WIN32 + for (int64_t i = 0; i < cacheCount; i++) { +#else for (int i = 0; i < cacheCount; i++) { +#endif jobject javaMemoryProperty = static_cast(env->GetObjectArrayElement(javaCachePropertiesArray, i)); jclass descriptionClass = env->GetObjectClass(javaMemoryProperty); jfieldID field = env->GetFieldID(descriptionClass, "type", "I"); @@ -402,10 +410,18 @@ JNIEXPORT jint JNICALL Java_uk_ac_manchester_tornado_drivers_spirv_levelzero_Lev jfieldID fieldType = env->GetFieldID(descriptionClass, "stype", "I"); ze_structure_type_t type = static_cast(env->GetIntField(javaDeviceModuleProperties, fieldType)); jfieldID fieldPNext = env->GetFieldID(descriptionClass, "pNext", "J"); +#ifdef _WIN32 + int64_t pNext = static_cast(env->GetLongField(javaDeviceModuleProperties, fieldPNext)); +#else long pNext = static_cast(env->GetLongField(javaDeviceModuleProperties, fieldPNext)); +#endif void* ptrNext = nullptr; if (pNext != -1) { +#ifdef _WIN32 + ptrNext = reinterpret_cast(pNext); +#else ptrNext = reinterpret_cast(pNext); +#endif } ze_device_module_properties_t deviceModuleProperties; @@ -475,7 +491,11 @@ JNIEXPORT jint JNICALL Java_uk_ac_manchester_tornado_drivers_spirv_levelzero_Lev LOG_ZE_JNI("zeDeviceGetCommandQueueGroupProperties", result); if (javaQueuePropertiesArray != nullptr) { +#ifdef _WIN32 + for (int64_t i = 0; i < numQueueGroups; i++) { +#else for (int i = 0; i < numQueueGroups; i++) { +#endif jobject javaMemoryProperty = static_cast(env->GetObjectArrayElement(javaQueuePropertiesArray, i)); jclass descriptionClass = env->GetObjectClass(javaMemoryProperty); jfieldID field = env->GetFieldID(descriptionClass, "type", "I"); diff --git a/levelZeroLib/src/levelZeroDriver.cpp b/levelZeroLib/src/levelZeroDriver.cpp index d3d5604..c6a1cfd 100644 --- a/levelZeroLib/src/levelZeroDriver.cpp +++ b/levelZeroLib/src/levelZeroDriver.cpp @@ -124,7 +124,11 @@ JNIEXPORT jint JNICALL Java_uk_ac_manchester_tornado_drivers_spirv_levelzero_Lev numDevices[0] = deviceCountNumber; } else { // Update object javaDeviceHandler +#ifdef _WIN32 + for (int64_t i = 0; i < deviceCountNumber; i++) { +#else for (int i = 0; i < deviceCountNumber; i++) { +#endif deviceHandlerArray[i] = reinterpret_cast(*(&(deviceHandler) + i)); } env->SetIntField(javaDeviceHandler, fieldNumDevices, deviceCountNumber); @@ -147,7 +151,11 @@ JNIEXPORT jint JNICALL Java_uk_ac_manchester_tornado_drivers_spirv_levelzero_Lev jclass descriptionClass = env->GetObjectClass(javaDriverProperties); jfieldID fieldDescriptorPointer = env->GetFieldID(descriptionClass, "nativePointer", "J"); +#ifdef _WIN32 + int64_t valuePointerDescriptor = env->GetLongField(javaDriverProperties, fieldDescriptorPointer); +#else long valuePointerDescriptor = env->GetLongField(javaDriverProperties, fieldDescriptorPointer); +#endif jfieldID fieldDescriptorType = env->GetFieldID(descriptionClass, "type", "I"); auto type = static_cast(env->GetIntField(javaDriverProperties, fieldDescriptorType)); @@ -164,7 +172,11 @@ JNIEXPORT jint JNICALL Java_uk_ac_manchester_tornado_drivers_spirv_levelzero_Lev ze_result_t result = zeDriverGetProperties(driver, &driverProperties); LOG_ZE_JNI("zeDriverGetProperties", result); +#ifdef _WIN32 + valuePointerDescriptor = reinterpret_cast(&(driverProperties)); +#else valuePointerDescriptor = reinterpret_cast(&(driverProperties)); +#endif env->SetLongField(javaDriverProperties, fieldDescriptorPointer, valuePointerDescriptor); jfieldID field = env->GetFieldID(descriptionClass, "uuid", "[I"); @@ -201,7 +213,11 @@ JNIEXPORT jint JNICALL Java_uk_ac_manchester_tornado_drivers_spirv_levelzero_Lev ze_result_t result = zeDriverGetApiVersion(driver, &version); LOG_ZE_JNI("zeDriverGetApiVersion", result); +#ifdef _WIN32 + valueAPIVersion = reinterpret_cast(&version); +#else valueAPIVersion = reinterpret_cast(&version); +#endif env->SetIntField(javaAPIVersion, fieldAPIVersionPtr, valueAPIVersion); std::stringstream ss; diff --git a/levelZeroLib/src/levelZeroModule.cpp b/levelZeroLib/src/levelZeroModule.cpp index 9ed6179..e80d6cf 100644 --- a/levelZeroLib/src/levelZeroModule.cpp +++ b/levelZeroLib/src/levelZeroModule.cpp @@ -76,7 +76,11 @@ JNIEXPORT jint JNICALL Java_uk_ac_manchester_tornado_drivers_spirv_levelzero_Lev int type = env->GetIntField(javaKernelDesc, fieldType); jfieldID fieldFlags = env->GetFieldID(javaKernelDescClass, "flags", "J"); +#ifdef _WIN32 + int64_t flags = env->GetLongField(javaKernelDesc, fieldFlags); +#else long flags = env->GetLongField(javaKernelDesc, fieldFlags); +#endif jfieldID fieldName = env->GetFieldID(javaKernelDescClass, "kernelName", "Ljava/lang/String;"); jstring javaStringName = static_cast(env->GetObjectField(javaKernelDesc, fieldName)); diff --git a/levelZeroLib/src/levelZeroTimeStampKernel.cpp b/levelZeroLib/src/levelZeroTimeStampKernel.cpp index 7acd55c..8168d92 100644 --- a/levelZeroLib/src/levelZeroTimeStampKernel.cpp +++ b/levelZeroLib/src/levelZeroTimeStampKernel.cpp @@ -50,16 +50,31 @@ JNIEXPORT jint JNICALL Java_uk_ac_manchester_tornado_drivers_spirv_levelzero_ZeK // Set fields for Java Event Handle jclass eventHandleClass = env->GetObjectClass(javaZeTimeKernelStampResult); jfieldID field = env->GetFieldID(eventHandleClass, "globalKernelStart", "J"); +#ifdef _WIN32 + env->SetLongField(javaZeTimeKernelStampResult, field, reinterpret_cast(kernelTsResults->global.kernelStart)); +#else env->SetLongField(javaZeTimeKernelStampResult, field, reinterpret_cast(kernelTsResults->global.kernelStart)); - +#endif field = env->GetFieldID(eventHandleClass, "globalKernelEnd", "J"); +#ifdef _WIN32 + env->SetLongField(javaZeTimeKernelStampResult, field, reinterpret_cast(kernelTsResults->global.kernelEnd)); +#else env->SetLongField(javaZeTimeKernelStampResult, field, reinterpret_cast(kernelTsResults->global.kernelEnd)); +#endif field = env->GetFieldID(eventHandleClass, "contextKernelStart", "J"); +#ifdef _WIN32 + env->SetLongField(javaZeTimeKernelStampResult, field, reinterpret_cast(kernelTsResults->context.kernelStart)); +#else env->SetLongField(javaZeTimeKernelStampResult, field, reinterpret_cast(kernelTsResults->context.kernelStart)); +#endif field = env->GetFieldID(eventHandleClass, "contextKernelEnd", "J"); +#ifdef _WIN32 + env->SetLongField(javaZeTimeKernelStampResult, field, reinterpret_cast(kernelTsResults->context.kernelEnd)); +#else env->SetLongField(javaZeTimeKernelStampResult, field, reinterpret_cast(kernelTsResults->context.kernelEnd)); +#endif return 1; } \ No newline at end of file diff --git a/lookUpBufferAddress.cl b/lookUpBufferAddress.cl new file mode 100644 index 0000000..cb79492 --- /dev/null +++ b/lookUpBufferAddress.cl @@ -0,0 +1,3 @@ +__kernel void lookUp(__global long *heap, __global long* output) { + output[get_global_id(0)] = (ulong) heap; +} diff --git a/lookUpBufferAddress.spv b/lookUpBufferAddress.spv new file mode 100644 index 0000000000000000000000000000000000000000..688537ecb2fdcffbe809d4e46b0270002c59ee6b GIT binary patch literal 1096 zcma)*TT8=05QVpGS}(1ww_5L3d=NqU=mk+k!BP;!2cJu93kE|HVk-J4{Jp*io-fHl zd~;!$IWx0oc1|{x^1+NTCDSxDvuCokWXi^oi>>CimfL!64`u7hmSi8INpO4LdrQVf zJe4$7R=g!n&TVS@VfgwKnUt@hnu=o4S4kKaV}`@G$cx{HH&f3~20`BspIv_tyoVz< z@xovrYSlTW@rTPX3&JlGSIk+yDH~3cXqsfWB|Hg|_%qE{HB(mKbI0R2kFCgNXR;`L z))gDMzAvM`RkI=M#+Fc7^H7eBc#jcoW3AFZ03&p3gCS`^qpowbYTLW>MDm2QTNC{Dm;R;H2M)JkB?? zxet7&x&?f^8+^{8*S9KoD{#;Q&FBNB2bzJe#7U0xaMF($d3bE>Ql06us!hF(Pj!O% z+IfyT?!Y&P!`2Z;uZ}kJ5a(WxLioH-a;!VW;NXqILykLv=|eqt`H?DT+Ey3-OL9eB Zcnek+-T}Jo$e!yyXt}G6p45@IoWIsPRYm{+ literal 0 HcmV?d00001 diff --git a/scripts/compileAndRun.cmd b/scripts/compileAndRun.cmd new file mode 100644 index 0000000..78e0071 --- /dev/null +++ b/scripts/compileAndRun.cmd @@ -0,0 +1,3 @@ +mvn clean install + +java -Djava.library.path=./levelZeroLib/build/Release -cp target/beehive-levelzero-jni-0.1.2.jar uk.ac.manchester.tornado.drivers.spirv.levelzero.samples.TestLevelZero copyData.spv diff --git a/scripts/compileAndRun.sh b/scripts/compileAndRun.sh index 6259799..693812f 100755 --- a/scripts/compileAndRun.sh +++ b/scripts/compileAndRun.sh @@ -1,7 +1,7 @@ -mvn clean install -clang -cc1 -triple spir copy_data.cl -O0 -finclude-default-header -emit-llvm-bc -o opencl-copy.bc -llvm-spirv opencl-copy.bc -o opencl-copy.spv -cp opencl-copy.spv /tmp/opencl-copy.spv +#!/bin/bash -java -Djava.library.path=./levelZeroLib/build -cp target/beehive-levelzero-jni-0.1.2.jar uk.ac.manchester.tornado.drivers.spirv.levelzero.samples.TestLevelZero +mvn clean install +clang -cc1 -triple spir copyData.cl -O0 -finclude-default-header -emit-llvm-bc -o copyData.bc +llvm-spirv copyData.bc -o copyData.spv +java -Djava.library.path=./levelZeroLib/build -cp target/beehive-levelzero-jni-0.1.2.jar uk.ac.manchester.tornado.drivers.spirv.levelzero.samples.TestLevelZero copyData.spv diff --git a/scripts/compileSPIRVKernelCopy.sh b/scripts/compileSPIRVKernelCopy.sh index 69e91ec..061caff 100755 --- a/scripts/compileSPIRVKernelCopy.sh +++ b/scripts/compileSPIRVKernelCopy.sh @@ -1,6 +1,4 @@ #!/bin/bash -clang -cc1 -triple spir copy_data.cl -O0 -finclude-default-header -emit-llvm-bc -o opencl-copy.bc -llvm-spirv opencl-copy.bc -o opencl-copy.spv -mv opencl-copy.spv /tmp/opencl-copy.spv - +clang -cc1 -triple spir copyData.cl -O0 -finclude-default-header -emit-llvm-bc -o copyData.bc +llvm-spirv copyData.bc -o copyData.spv diff --git a/scripts/copies.cmd b/scripts/copies.cmd new file mode 100644 index 0000000..74fd063 --- /dev/null +++ b/scripts/copies.cmd @@ -0,0 +1 @@ +java -Djava.library.path=./levelZeroLib/build/Release -cp target/beehive-levelzero-jni-0.1.2.jar uk.ac.manchester.tornado.drivers.spirv.levelzero.samples.TestCopies diff --git a/scripts/copies.sh b/scripts/copies.sh index 3c92816..f438af1 100755 --- a/scripts/copies.sh +++ b/scripts/copies.sh @@ -1 +1,3 @@ +#!/bin/bash + java -Djava.library.path=./levelZeroLib/build -cp target/beehive-levelzero-jni-0.1.2.jar uk.ac.manchester.tornado.drivers.spirv.levelzero.samples.TestCopies diff --git a/scripts/events.cmd b/scripts/events.cmd new file mode 100644 index 0000000..66c3c14 --- /dev/null +++ b/scripts/events.cmd @@ -0,0 +1 @@ +java -Djava.library.path=./levelZeroLib/build/Release -cp target/beehive-levelzero-jni-0.1.2.jar uk.ac.manchester.tornado.drivers.spirv.levelzero.samples.TestWithEvents diff --git a/scripts/events.sh b/scripts/events.sh index 796d6a1..1bd878f 100755 --- a/scripts/events.sh +++ b/scripts/events.sh @@ -1,3 +1,3 @@ -java -Djava.library.path=./levelZeroLib/build \ - -cp target/beehive-levelzero-jni-0.1.2.jar \ - uk.ac.manchester.tornado.drivers.spirv.levelzero.samples.TestWithEvents +#!/bin/bash + +java -Djava.library.path=./levelZeroLib/build -cp target/beehive-levelzero-jni-0.1.2.jar uk.ac.manchester.tornado.drivers.spirv.levelzero.samples.TestWithEvents diff --git a/scripts/fences.cmd b/scripts/fences.cmd new file mode 100644 index 0000000..c6032e2 --- /dev/null +++ b/scripts/fences.cmd @@ -0,0 +1,4 @@ +#!/bin/bash + +./scripts/compileSPIRVKernelCopy.sh +java -Djava.library.path=./levelZeroLib/build/Release -cp target/beehive-levelzero-jni-0.1.2.jar uk.ac.manchester.tornado.drivers.spirv.levelzero.samples.TestFences copyData.spv diff --git a/scripts/fences.sh b/scripts/fences.sh index 583db71..b2dece3 100755 --- a/scripts/fences.sh +++ b/scripts/fences.sh @@ -1,4 +1,4 @@ -#!/bin/bash +#!/bin/bash ./scripts/compileSPIRVKernelCopy.sh -java -Djava.library.path=./levelZeroLib/build -cp target/beehive-levelzero-jni-0.1.2.jar uk.ac.manchester.tornado.drivers.spirv.levelzero.samples.TestFences +java -Djava.library.path=./levelZeroLib/build -cp target/beehive-levelzero-jni-0.1.2.jar uk.ac.manchester.tornado.drivers.spirv.levelzero.samples.TestFences copyData.spv diff --git a/scripts/kernelTimers.cmd b/scripts/kernelTimers.cmd new file mode 100644 index 0000000..775df9e --- /dev/null +++ b/scripts/kernelTimers.cmd @@ -0,0 +1 @@ +java -Djava.library.path=./levelZeroLib/build/Release -cp target/beehive-levelzero-jni-0.1.2.jar uk.ac.manchester.tornado.drivers.spirv.levelzero.samples.TestKernelTimer copyData.spv diff --git a/scripts/kernelTimers.sh b/scripts/kernelTimers.sh index c06c1e2..43a5653 100755 --- a/scripts/kernelTimers.sh +++ b/scripts/kernelTimers.sh @@ -1,4 +1,4 @@ -#!/bin/bash +#!/bin/bash ./scripts/compileSPIRVKernelCopy.sh -java -Djava.library.path=./levelZeroLib/build -cp target/beehive-levelzero-jni-0.1.2.jar uk.ac.manchester.tornado.drivers.spirv.levelzero.samples.TestKernelTimer +java -Djava.library.path=./levelZeroLib/build -cp target/beehive-levelzero-jni-0.1.2.jar uk.ac.manchester.tornado.drivers.spirv.levelzero.samples.TestKernelTimer copyData.spv diff --git a/scripts/largeBuffers.cmd b/scripts/largeBuffers.cmd new file mode 100644 index 0000000..05e593c --- /dev/null +++ b/scripts/largeBuffers.cmd @@ -0,0 +1 @@ +java -Djava.library.path=./levelZeroLib/build/Release -cp target/beehive-levelzero-jni-0.1.2.jar uk.ac.manchester.tornado.drivers.spirv.levelzero.samples.TestLargeBuffer diff --git a/scripts/largeBuffers.sh b/scripts/largeBuffers.sh index b74de2b..944bcb2 100755 --- a/scripts/largeBuffers.sh +++ b/scripts/largeBuffers.sh @@ -1 +1,3 @@ +#!/bin/bash + java -Djava.library.path=./levelZeroLib/build -cp target/beehive-levelzero-jni-0.1.2.jar uk.ac.manchester.tornado.drivers.spirv.levelzero.samples.TestLargeBuffer diff --git a/scripts/run.cmd b/scripts/run.cmd new file mode 100644 index 0000000..748cf88 --- /dev/null +++ b/scripts/run.cmd @@ -0,0 +1 @@ +java -Djava.library.path=./levelZeroLib/build/Release -cp target/beehive-levelzero-jni-0.1.2.jar uk.ac.manchester.tornado.drivers.spirv.levelzero.samples.TestLevelZero copyData.spv diff --git a/scripts/run.sh b/scripts/run.sh index 4b973b5..9e76b54 100755 --- a/scripts/run.sh +++ b/scripts/run.sh @@ -1,5 +1,6 @@ -clang -cc1 -triple spir copy_data.cl -O0 -finclude-default-header -emit-llvm-bc -o opencl-copy.bc -llvm-spirv opencl-copy.bc -o opencl-copy.spv -mv opencl-copy.spv /tmp/opencl-copy.spv +#!/bin/bash -java -Djava.library.path=./levelZeroLib/build -cp target/beehive-levelzero-jni-0.1.2.jar uk.ac.manchester.tornado.drivers.spirv.levelzero.samples.TestLevelZero +clang -cc1 -triple spir copyData.cl -O0 -finclude-default-header -emit-llvm-bc -o copyData.bc +llvm-spirv copyData.bc -o copyData.spv + +java -Djava.library.path=./levelZeroLib/build -cp target/beehive-levelzero-jni-0.1.2.jar uk.ac.manchester.tornado.drivers.spirv.levelzero.samples.TestLevelZero copyData.spv diff --git a/scripts/transfersTimers.cmd b/scripts/transfersTimers.cmd new file mode 100644 index 0000000..78a3c7a --- /dev/null +++ b/scripts/transfersTimers.cmd @@ -0,0 +1 @@ +java -Djava.library.path=./levelZeroLib/build/Release -cp target/beehive-levelzero-jni-0.1.2.jar uk.ac.manchester.tornado.drivers.spirv.levelzero.samples.TestTransferTimers diff --git a/scripts/transfersTimers.sh b/scripts/transfersTimers.sh index d57965f..9e1d761 100755 --- a/scripts/transfersTimers.sh +++ b/scripts/transfersTimers.sh @@ -1 +1,3 @@ +#!/bin/bash + java -Djava.library.path=./levelZeroLib/build -cp target/beehive-levelzero-jni-0.1.2.jar uk.ac.manchester.tornado.drivers.spirv.levelzero.samples.TestTransferTimers diff --git a/src/main/java/uk/ac/manchester/tornado/drivers/spirv/levelzero/samples/SimulationLKBuffer.java b/src/main/java/uk/ac/manchester/tornado/drivers/spirv/levelzero/samples/SimulationLKBuffer.java index 0545f1a..4e1ac0c 100644 --- a/src/main/java/uk/ac/manchester/tornado/drivers/spirv/levelzero/samples/SimulationLKBuffer.java +++ b/src/main/java/uk/ac/manchester/tornado/drivers/spirv/levelzero/samples/SimulationLKBuffer.java @@ -46,39 +46,40 @@ import uk.ac.manchester.tornado.drivers.spirv.levelzero.utils.LevelZeroUtils; /** - * How to run? + * Kernels to test: * * - * __kernel void lookUp(__global uchar *heap) { - * __global ulong *_frame = (__global ulong *) &heap[0]; - * *((__global long *) &heap[get_global_id(0)]) = (ulong) _frame; - * } - * + * __kernel void lookUp(__global uchar *heap) { + * __global ulong *_frame = (__global ulong *) &heap[0]; + * *((__global long *) &heap[get_global_id(0)]) = (ulong) _frame; + * } * * * - * __kernel void copyTest(__global uchar *_heap_base) - * { - * int i_8, i_7, i_1, i_2; - * ulong ul_0, ul_6; - * long l_3, l_5, l_4; + * __kernel void copyTest(__global uchar *_heap_base) + * { + * int i_8, i_7, i_1, i_2; + * ulong ul_0, ul_6; + * long l_3, l_5, l_4; * - * __global ulong *_frame = (__global ulong *) &_heap_base[0]; + * __global ulong *_frame = (__global ulong *) &_heap_base[0]; * - * ul_0 = (ulong) _frame[3]; - * i_1 = get_global_id(0); - * i_2 = i_1; - * l_3 = (long) i_2; - * l_4 = l_3 << 3; // Long buffer - * l_5 = l_4 + 16L; // Randomly starting in position 16 - * ul_6 = ul_0 + l_5; - * *((__global int *) ul_6) = 555; - * } + * ul_0 = (ulong) _frame[3]; + * i_1 = get_global_id(0); + * i_2 = i_1; + * l_3 = (long) i_2; + * l_4 = l_3 << 3; // Long buffer + * l_5 = l_4 + 16L; // Randomly starting in position 16 + * ul_6 = ul_0 + l_5; + * *((__global int *) ul_6) = 555; + * } * * * + * How to run? + * * - * $ tornado uk.ac.manchester.tornado.drivers.spirv.levelzero.samples.SimulationLKBuffer + * $ tornado uk.ac.manchester.tornado.drivers.spirv.levelzero.samples.SimulationLKBuffer lookUpBufferAddress.spv copyTest.spv * */ public class SimulationLKBuffer { @@ -87,6 +88,7 @@ public class SimulationLKBuffer { // private static final int DEVICE_HEAP_SIZE = 128 * 8; private static final int DEVICE_HEAP_SIZE = 1000000000; // 1GB private static long[] stack; + private static String[] args; private static void dispatchCopyKernel(LevelZeroCommandList commandList, LevelZeroCommandQueue commandQueue, LevelZeroKernel levelZeroKernel, long[] output, int bufferSize, ByteBuffer stack) { ZeKernelHandle kernel = levelZeroKernel.getKernelHandle(); @@ -155,7 +157,7 @@ private static void simulateLookUpBufferAddress(LevelZeroContext context, LevelZ int result = context.zeMemAllocDevice(context.getDefaultContextPtr(), deviceMemAllocDesc, DEVICE_HEAP_SIZE, 1, device.getDeviceHandlerPtr(), deviceHeapBuffer); LevelZeroUtils.errorLog("zeMemAllocDevice", result); - LevelZeroKernel levelZeroKernel = LevelZeroUtils.compileSPIRVKernel(device, context, "lookUp", "/tmp/lookUpBufferAddress.spv"); + LevelZeroKernel levelZeroKernel = LevelZeroUtils.compileSPIRVKernel(device, context, "lookUp", args[0]); LevelZeroUtils.dispatchLookUpBuffer(commandList, commandQueue, levelZeroKernel, deviceHeapBuffer, output, bufferSize); result = commandList.zeCommandListReset(commandList.getCommandListHandlerPtr()); @@ -175,7 +177,7 @@ private static void simulateLookUpBufferAddress(LevelZeroContext context, LevelZ result = commandList.zeCommandListAppendBarrier(commandList.getCommandListHandlerPtr(), null, 0, null); LevelZeroUtils.errorLog("zeCommandListAppendBarrier", result); - LevelZeroKernel kernelCopy = LevelZeroUtils.compileSPIRVKernel(device, context, "copyTest", "/tmp/example.spv"); + LevelZeroKernel kernelCopy = LevelZeroUtils.compileSPIRVKernel(device, context, "copyTest", args[1]); long[] output2 = new long[128]; dispatchCopyKernel(commandList, commandQueue, kernelCopy, output2, 128 * Sizeof.LONG.getNumBytes(), stack); @@ -199,6 +201,7 @@ private static void simulateLookUpBufferAddress(LevelZeroContext context, LevelZ * @param args */ public static void main(String[] args) { + SimulationLKBuffer.args = args; LevelZeroDriver driver = new LevelZeroDriver(); LevelZeroContext context = LevelZeroUtils.zeInitContext(driver); LevelZeroDevice device = LevelZeroUtils.zeGetDevices(context, driver); diff --git a/src/main/java/uk/ac/manchester/tornado/drivers/spirv/levelzero/samples/TestFences.java b/src/main/java/uk/ac/manchester/tornado/drivers/spirv/levelzero/samples/TestFences.java index 9c7da18..de6fa97 100644 --- a/src/main/java/uk/ac/manchester/tornado/drivers/spirv/levelzero/samples/TestFences.java +++ b/src/main/java/uk/ac/manchester/tornado/drivers/spirv/levelzero/samples/TestFences.java @@ -71,9 +71,9 @@ * Kernel to test: * * - * __kernel void copydata(__global int* input, __global int* output) { - * uint idx = get_global_id(0); - * output[idx] = input[idx]; + * __kernel void copyData(__global int* input, __global int* output) { + * uint idx = get_global_id(0); + * output[idx] = input[idx]; * } * * @@ -81,17 +81,16 @@ * To compile to SPIR-V: * * - * $ clang -cc1 -triple spir opencl-copy.cl -O0 -finclude-default-header -emit-llvm-bc -o opencl-copy.bc - * $ llvm-spirv opencl-copy.bc -o opencl-copy.spv - * $ mv opencl-copy.spv /tmp + * $ clang -cc1 -triple spir copyData.cl -O0 -finclude-default-header -emit-llvm-bc -o copyData.bc + * $ llvm-spirv copyData.bc -o copyData.spv * * + * * How to run? * * - * tornado uk.ac.manchester.tornado.drivers.spirv.levelzero.samples.TestLevelZero + * tornado uk.ac.manchester.tornado.drivers.spirv.levelzero.samples.TestFences copyData.spv * - * */ public class TestFences { @@ -227,7 +226,7 @@ public static void main(String[] args) throws IOException { moduleDesc.setFormat(ZeModuleFormat.ZE_MODULE_FORMAT_IL_SPIRV); moduleDesc.setBuildFlags(""); - result = context.zeModuleCreate(context.getDefaultContextPtr(), device.getDeviceHandlerPtr(), moduleDesc, module, buildLog, "/tmp/opencl-copy.spv"); + result = context.zeModuleCreate(context.getDefaultContextPtr(), device.getDeviceHandlerPtr(), moduleDesc, module, buildLog, args[0]); LevelZeroUtils.errorLog("zeModuleCreate", result); if (result != ZeResult.ZE_RESULT_SUCCESS) { @@ -248,7 +247,7 @@ public static void main(String[] args) throws IOException { ZeKernelDescriptor kernelDesc = new ZeKernelDescriptor(); ZeKernelHandle kernel = new ZeKernelHandle(); - kernelDesc.setKernelName("copydata"); + kernelDesc.setKernelName("copyData"); result = levelZeroModule.zeKernelCreate(module.getPtrZeModuleHandle(), kernelDesc, kernel); LevelZeroUtils.errorLog("zeKernelCreate", result); diff --git a/src/main/java/uk/ac/manchester/tornado/drivers/spirv/levelzero/samples/TestKernelTimer.java b/src/main/java/uk/ac/manchester/tornado/drivers/spirv/levelzero/samples/TestKernelTimer.java index c838117..06ec64b 100644 --- a/src/main/java/uk/ac/manchester/tornado/drivers/spirv/levelzero/samples/TestKernelTimer.java +++ b/src/main/java/uk/ac/manchester/tornado/drivers/spirv/levelzero/samples/TestKernelTimer.java @@ -76,25 +76,25 @@ * Kernel to test: * * - * __kernel void copydata(__global int* input, __global int* output) { - * uint idx = get_global_id(0); - * output[idx] = input[idx]; - * } + * __kernel void copyData(__global int* input, __global int* output) { + * uint idx = get_global_id(0); + * output[idx] = input[idx]; + * } * - *

- *

+ * + * * To compile to SPIR-V: * * - * $ clang -cc1 -triple spir opencl-copy.cl -O0 -finclude-default-header -emit-llvm-bc -o opencl-copy.bc - * $ llvm-spirv opencl-copy.bc -o opencl-copy.spv - * $ mv opencl-copy.spv /tmp + * $ clang -cc1 -triple spir copyData.cl -O0 -finclude-default-header -emit-llvm-bc -o copyData.bc + * $ llvm-spirv copyData.bc -o copyData.spv * - *

+ * + * * How to run? * * - * tornado uk.ac.manchester.tornado.drivers.spirv.levelzero.samples.TestKernelTimer + * tornado uk.ac.manchester.tornado.drivers.spirv.levelzero.samples.TestKernelTimer copyData.spv * */ public class TestKernelTimer { @@ -241,7 +241,7 @@ public static void main(String[] args) throws IOException { moduleDesc.setFormat(ZeModuleFormat.ZE_MODULE_FORMAT_IL_SPIRV); moduleDesc.setBuildFlags(""); - result = context.zeModuleCreate(context.getDefaultContextPtr(), device.getDeviceHandlerPtr(), moduleDesc, module, buildLog, "/tmp/opencl-copy.spv"); + result = context.zeModuleCreate(context.getDefaultContextPtr(), device.getDeviceHandlerPtr(), moduleDesc, module, buildLog, args[0]); LevelZeroUtils.errorLog("zeModuleCreate", result); if (result != ZeResult.ZE_RESULT_SUCCESS) { @@ -262,7 +262,7 @@ public static void main(String[] args) throws IOException { ZeKernelDescriptor kernelDesc = new ZeKernelDescriptor(); ZeKernelHandle kernel = new ZeKernelHandle(); - kernelDesc.setKernelName("copydata"); + kernelDesc.setKernelName("copyData"); result = levelZeroModule.zeKernelCreate(module.getPtrZeModuleHandle(), kernelDesc, kernel); LevelZeroUtils.errorLog("zeKernelCreate", result); diff --git a/src/main/java/uk/ac/manchester/tornado/drivers/spirv/levelzero/samples/TestLevelZero.java b/src/main/java/uk/ac/manchester/tornado/drivers/spirv/levelzero/samples/TestLevelZero.java index b6e0bdd..5cd8c35 100644 --- a/src/main/java/uk/ac/manchester/tornado/drivers/spirv/levelzero/samples/TestLevelZero.java +++ b/src/main/java/uk/ac/manchester/tornado/drivers/spirv/levelzero/samples/TestLevelZero.java @@ -75,9 +75,9 @@ * Kernel to test: * * - * __kernel void copydata(__global int* input, __global int* output) { - * uint idx = get_global_id(0); - * output[idx] = input[idx]; + * __kernel void copyData(__global int* input, __global int* output) { + * uint idx = get_global_id(0); + * output[idx] = input[idx]; * } * * @@ -85,17 +85,16 @@ * To compile to SPIR-V: * * - * $ clang -cc1 -triple spir opencl-copy.cl -O0 -finclude-default-header -emit-llvm-bc -o opencl-copy.bc - * $ llvm-spirv opencl-copy.bc -o opencl-copy.spv - * $ cp opencl-copy.spv /tmp + * $ clang -cc1 -triple spir copyData.cl -O0 -finclude-default-header -emit-llvm-bc -o copyData.bc + * $ llvm-spirv copyData.bc -o copyData.spv * * + * * How to run? * * - * tornado uk.ac.manchester.tornado.drivers.spirv.levelzero.samples.TestLevelZero + * tornado uk.ac.manchester.tornado.drivers.spirv.levelzero.samples.TestLevelZero copyData.spv * - * */ public class TestLevelZero { @@ -306,7 +305,7 @@ public static void main(String[] args) throws IOException { System.gc(); - result = context.zeModuleCreate(context.getDefaultContextPtr(), device.getDeviceHandlerPtr(), moduleDesc, module, buildLog, "/tmp/opencl-copy.spv"); + result = context.zeModuleCreate(context.getDefaultContextPtr(), device.getDeviceHandlerPtr(), moduleDesc, module, buildLog, args[0]); LevelZeroUtils.errorLog("zeModuleCreate", result); System.gc(); @@ -333,7 +332,7 @@ public static void main(String[] args) throws IOException { ZeKernelDescriptor kernelDesc = new ZeKernelDescriptor(); ZeKernelHandle kernel = new ZeKernelHandle(); - kernelDesc.setKernelName("copydata"); + kernelDesc.setKernelName("copyData"); result = levelZeroModule.zeKernelCreate(module.getPtrZeModuleHandle(), kernelDesc, kernel); LevelZeroUtils.errorLog("zeKernelCreate", result); diff --git a/src/main/java/uk/ac/manchester/tornado/drivers/spirv/levelzero/samples/TestLevelZeroDedicatedMemory.java b/src/main/java/uk/ac/manchester/tornado/drivers/spirv/levelzero/samples/TestLevelZeroDedicatedMemory.java index c8fe537..1fbb807 100644 --- a/src/main/java/uk/ac/manchester/tornado/drivers/spirv/levelzero/samples/TestLevelZeroDedicatedMemory.java +++ b/src/main/java/uk/ac/manchester/tornado/drivers/spirv/levelzero/samples/TestLevelZeroDedicatedMemory.java @@ -54,9 +54,9 @@ * Kernel to test: * * - * __kernel void copydata(__global int* input, __global int* output) { - * uint idx = get_global_id(0); - * output[idx] = input[idx]; + * __kernel void copyData(__global int* input, __global int* output) { + * uint idx = get_global_id(0); + * output[idx] = input[idx]; * } * * @@ -64,17 +64,16 @@ * To compile to SPIR-V: * * - * $ clang -cc1 -triple spir copydata.cl -O0 -finclude-default-header -emit-llvm-bc -o copydata.bc - * $ llvm-spirv copydata.bc -o copydata.spv - * $ cp opencl-copy.spv /tmp/copydata.spv + * $ clang -cc1 -triple spir copyData.cl -O0 -finclude-default-header -emit-llvm-bc -o copyData.bc + * $ llvm-spirv copyData.bc -o copyData.spv * * + * * How to run? * * - * $ tornado uk.ac.manchester.tornado.drivers.spirv.levelzero.samples.TestLevelZeroDedicatedMemory + * $ tornado uk.ac.manchester.tornado.drivers.spirv.levelzero.samples.TestLevelZeroDedicatedMemory copyData.spv * - * */ public class TestLevelZeroDedicatedMemory { @@ -120,7 +119,7 @@ public static void main(String[] args) throws IOException { moduleDesc.setFormat(ZeModuleFormat.ZE_MODULE_FORMAT_IL_SPIRV); moduleDesc.setBuildFlags(""); - result = context.zeModuleCreate(context.getDefaultContextPtr(), device.getDeviceHandlerPtr(), moduleDesc, module, buildLog, "/tmp/copydata.spv"); + result = context.zeModuleCreate(context.getDefaultContextPtr(), device.getDeviceHandlerPtr(), moduleDesc, module, buildLog, args[0]); LevelZeroUtils.errorLog("zeModuleCreate", result); if (result != ZeResult.ZE_RESULT_SUCCESS) { @@ -142,7 +141,7 @@ public static void main(String[] args) throws IOException { ZeKernelDescriptor kernelDesc = new ZeKernelDescriptor(); ZeKernelHandle kernel = new ZeKernelHandle(); - kernelDesc.setKernelName("copydata"); + kernelDesc.setKernelName("copyData"); result = levelZeroModule.zeKernelCreate(module.getPtrZeModuleHandle(), kernelDesc, kernel); LevelZeroUtils.errorLog("zeKernelCreate", result); diff --git a/src/main/java/uk/ac/manchester/tornado/drivers/spirv/levelzero/samples/TestLevelZeroDedicatedMemoryLong.java b/src/main/java/uk/ac/manchester/tornado/drivers/spirv/levelzero/samples/TestLevelZeroDedicatedMemoryLong.java index 1558b6f..e4fb46b 100644 --- a/src/main/java/uk/ac/manchester/tornado/drivers/spirv/levelzero/samples/TestLevelZeroDedicatedMemoryLong.java +++ b/src/main/java/uk/ac/manchester/tornado/drivers/spirv/levelzero/samples/TestLevelZeroDedicatedMemoryLong.java @@ -54,9 +54,9 @@ * Kernel to test: * * - * __kernel void copydata(__global long* input, __global long* output) { - * uint idx = get_global_id(0); - * output[idx] = input[idx]; + * __kernel void copyData(__global long* input, __global long* output) { + * uint idx = get_global_id(0); + * output[idx] = input[idx]; * } * * @@ -64,17 +64,16 @@ * To compile to SPIR-V: * * - * $ clang -cc1 -triple spir opencl-copy.cl -O0 -finclude-default-header -emit-llvm-bc -o opencl-copy.bc - * $ llvm-spirv opencl-copy.bc -o opencl-copy.spv - * $ cp opencl-copy.spv /tmp/copyLong.spv + * $ clang -cc1 -triple spir copyData.cl -O0 -finclude-default-header -emit-llvm-bc -o copyData.bc + * $ llvm-spirv copyData.bc -o copyData.spv * * + * * How to run? * * - * $ tornado uk.ac.manchester.tornado.drivers.spirv.levelzero.samples.TestLevelZeroDedicatedMemoryLong + * $ tornado uk.ac.manchester.tornado.drivers.spirv.levelzero.samples.TestLevelZeroDedicatedMemoryLong copyData.spv * - * */ public class TestLevelZeroDedicatedMemoryLong { @@ -120,7 +119,7 @@ public static void main(String[] args) throws IOException { moduleDesc.setFormat(ZeModuleFormat.ZE_MODULE_FORMAT_IL_SPIRV); moduleDesc.setBuildFlags(""); - result = context.zeModuleCreate(context.getDefaultContextPtr(), device.getDeviceHandlerPtr(), moduleDesc, module, buildLog, "/tmp/copyLong.spv"); + result = context.zeModuleCreate(context.getDefaultContextPtr(), device.getDeviceHandlerPtr(), moduleDesc, module, buildLog, args[0]); LevelZeroUtils.errorLog("zeModuleCreate", result); if (result != ZeResult.ZE_RESULT_SUCCESS) { @@ -142,7 +141,7 @@ public static void main(String[] args) throws IOException { ZeKernelDescriptor kernelDesc = new ZeKernelDescriptor(); ZeKernelHandle kernel = new ZeKernelHandle(); - kernelDesc.setKernelName("copydata"); + kernelDesc.setKernelName("copyData"); result = levelZeroModule.zeKernelCreate(module.getPtrZeModuleHandle(), kernelDesc, kernel); LevelZeroUtils.errorLog("zeKernelCreate", result); diff --git a/src/main/java/uk/ac/manchester/tornado/drivers/spirv/levelzero/samples/TestLookUpBufferAddress.java b/src/main/java/uk/ac/manchester/tornado/drivers/spirv/levelzero/samples/TestLookUpBufferAddress.java index 143b632..6972080 100644 --- a/src/main/java/uk/ac/manchester/tornado/drivers/spirv/levelzero/samples/TestLookUpBufferAddress.java +++ b/src/main/java/uk/ac/manchester/tornado/drivers/spirv/levelzero/samples/TestLookUpBufferAddress.java @@ -44,19 +44,32 @@ import uk.ac.manchester.tornado.drivers.spirv.levelzero.utils.LevelZeroUtils; /** - * How to run? + * Kernel to test: * * * __kernel void lookUp(__global long *heap, __global long* output) { * output[get_global_id(0)] = (ulong) heap; * } + * } * * + * + * To compile to SPIR-V: + * + * + * $ clang -cc1 -triple spir lookUpBufferAddress.cl -O0 -finclude-default-header -emit-llvm-bc -o lookUpBufferAddress.bc + * $ llvm-spirv lookUpBufferAddress.bc -o lookUpBufferAddress.spv + * + * + * + * How to run? + * * - * $ tornado uk.ac.manchester.tornado.drivers.spirv.levelzero.samples.TestLookUpBufferAddress + * $ tornado uk.ac.manchester.tornado.drivers.spirv.levelzero.samples.TestLookUpBufferAddress lookUpBufferAddress.spv * */ public class TestLookUpBufferAddress { + private static String[] args; private static void dispatchLookUpBuffer(LevelZeroCommandList commandList, LevelZeroCommandQueue commandQueue, LevelZeroKernel levelZeroKernel, LevelZeroByteBuffer deviceBuffer, LevelZeroByteBuffer bufferB, long[] output, int bufferSize) { @@ -140,7 +153,7 @@ private static void testLookUpBufferAddress(LevelZeroContext context, LevelZeroD result = commandList.zeCommandListAppendBarrier(commandList.getCommandListHandlerPtr(), null, 0, null); LevelZeroUtils.errorLog("zeCommandListAppendBarrier", result); - LevelZeroKernel levelZeroKernel = LevelZeroUtils.compileSPIRVKernel(device, context, "lookUp", "/home/juan/manchester/tornado/tornado/assembly/src/bin/spirv/lookUpBufferAddress.spv"); + LevelZeroKernel levelZeroKernel = LevelZeroUtils.compileSPIRVKernel(device, context, "lookUp", args[0]); dispatchLookUpBuffer(commandList, commandQueue, levelZeroKernel, deviceBuffer, bufferB, output, bufferSize); result = commandList.zeCommandListReset(commandList.getCommandListHandlerPtr()); @@ -165,6 +178,7 @@ private static void testLookUpBufferAddress(LevelZeroContext context, LevelZeroD * @param args */ public static void main(String[] args) { + TestLookUpBufferAddress.args = args; LevelZeroDriver driver = new LevelZeroDriver(); LevelZeroContext context = LevelZeroUtils.zeInitContext(driver); LevelZeroDevice device = LevelZeroUtils.zeGetDevices(context, driver); From 17e47f186e5b179852144aa65258e00ea1d449eb Mon Sep 17 00:00:00 2001 From: otabuzzman Date: Thu, 14 Mar 2024 20:07:50 +0100 Subject: [PATCH 4/6] Save work on Windows support --- CHANGELOG.md | 5 +++++ 1 file changed, 5 insertions(+) diff --git a/CHANGELOG.md b/CHANGELOG.md index 0604143..2ba0647 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -1,3 +1,8 @@ +## Beehive LevelZero-JNI 0.1.3 +14/03/2024 + +- Add Windwos installer support + ## Beehive LevelZero-JNI 0.1.2 12/12/23 From 6e70a0a91c4769ac40a4757945bdaf7d791bd598 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?J=C3=BCrgen=20Schuck?= Date: Thu, 21 Mar 2024 21:48:19 +0100 Subject: [PATCH 5/6] Update instructions on referencing DLLs via PATH --- README.md | 17 ++++++++++++----- 1 file changed, 12 insertions(+), 5 deletions(-) diff --git a/README.md b/README.md index 43e272b..a1a5f41 100644 --- a/README.md +++ b/README.md @@ -40,15 +40,13 @@ md build cd build cmake .. cmake --build . --config Release -``` - -Note: Check for extisting Level Zero API libraries (e.g. `ze_tracing_layer.dll`) in `c:\windows\system32` if `zello_world.exe` fails. -``` rem check .\bin\Release\zello_world.exe ``` +Note: If `zello_world.exe` fails, search for existing Level Zero API DLLs (file names start with `ze_`, e.g. `ze_tracing_layer.dll`) in `c:\windows\system32` and move them to another folder. + ### 2) Compile Level Zero JNI native code Set the paths to the directory of Level Zero installation. Here are examples: @@ -70,13 +68,17 @@ make #### Windows -Note: Run commands in _x64 Native Tools Command Prompt for VS 2022_ +Note: Run commands in _x64 Native Tools Command Prompt for VS 2022_. ```cmd git clone https://github.com/otabuzzman/levelzero-jni set ZE_SHARED_LOADER=%USERPROFILE%\lab\level-zero\build\lib\release\ze_loader.lib set CPLUS_INCLUDE_PATH=%USERPROFILE%\lab\level-zero\include set C_INCLUDE_PATH=%USERPROFILE%\lab\level-zero\include + +rem add the folder with Intel's Level Zero API DLLs to PATH +set PATH=%USERPROFILE%\lab\level-zero\build\bin\release;%PATH% + cd levelzero-jni\levelZeroLib md build cd build @@ -115,9 +117,14 @@ llvm-spirv copyData.bc -o copyData.spv #### Windows +Note: Java programs that use levelzero-jni are based on DLLs, which are provided by Intel's Level Zero API. For these programs to find these DLLs, the PATH environment variable must contain the folder that contains the DLLs. + ```cmd mvn clean package +rem add the folder with Intel's Level Zero API DLLs to PATH +set PATH=%USERPROFILE%\lab\level-zero\build\bin\release;%PATH% + rem copyData.spv file expected in CWD .\scripts\run.cmd From 314b34bdfd244e6e783972ae81365782545d1bc5 Mon Sep 17 00:00:00 2001 From: Juan Fumero Date: Tue, 26 Mar 2024 13:48:34 +0100 Subject: [PATCH 6/6] Update CHANGELOG.md Co-authored-by: Thanos Stratikopoulos <34061419+stratika@users.noreply.github.com> --- CHANGELOG.md | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/CHANGELOG.md b/CHANGELOG.md index 2ba0647..3b31d8c 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -1,7 +1,7 @@ ## Beehive LevelZero-JNI 0.1.3 14/03/2024 -- Add Windwos installer support +- Add Windows installer support ## Beehive LevelZero-JNI 0.1.2 12/12/23