Off load computational stuff?

[bradmesserle]

I am trying to learn tornado vm..

So for example i am trying to wrap a computational method..

 /**
 	*  Computes sin(x) - x, where |x| < 1/16.
 	*  Use a Remez polynomial approximation.
 	*  @param x a number smaller than 1/16
 	*  @return sin(x) - x
 	*/
	public static double polySine(final double x) {
   	 
   	 
   	  double x2 = x*x;

     	double p = 2.7553817452272217E-6;
     	p = p * x2 + -1.9841269659586505E-4;
     	p = p * x2 + 0.008333333333329196;
     	p = p * x2 + -0.16666666666666666;
     	p = p * x2 * x;

     	return p;
   	 
	}

public static void polySineTvm(final double x, final Double result) {
     	 final TaskSchedule t0 = new TaskSchedule("s0")
     			 .task("t0",TornadoFastMath::polySine,x)
     			 .streamOut(result);
     	 
   	 t0.execute();
   			 
	}

When I execute the tvm version I am getting the following error:

Exception in thread "main" java.lang.IndexOutOfBoundsException: Index: 1, Size: 0
    at java.util.ArrayList.rangeCheck(ArrayList.java:659)
    at java.util.ArrayList.get(ArrayList.java:435)
    at uk.ac.manchester.tornado.runtime.graph.TornadoExecutionContext.getObjectState(TornadoExecutionContext.java:215)
    at uk.ac.manchester.tornado.runtime.tasks.TornadoTaskSchedule.streamOutInner(TornadoTaskSchedule.java:645)
    at uk.ac.manchester.tornado.api.TaskSchedule.streamOut(TaskSchedule.java:253)
    at org.hipparchus.util.TornadoFastMath.polySinceTvm(TornadoFastMath.java:40)
    at org.hipparchus.runnable.TornadoFastMathRunnable.main(TornadoFastMathRunnable.java:26)

Am I misunderstanding the use case for using tvm?

Thank You!

[jjfumero]

Thank you for trying TornadoVM. TornadoVM is mainly designed for Array Programming (https://en.wikipedia.org/wiki/Array_programming) and inputs/outputs are passed as parameters to the expressions to be offloaded. TornadoVM currently cannot return a value from a sequential expression. Additionally, we cannot use wrappers. Instead, you can use an array double[1] and store the result in that variable.

So I tuned the code as follows (note that you are running only with 1 thread on the accelerator):

import uk.ac.manchester.tornado.api.TaskSchedule;

public class Foo {

	public static void polySine(final double x, double[] result) {
        double x2 = x*x;
        double p = 2.7553817452272217E-6;
     	p = p * x2 + -1.9841269659586505E-4;
     	p = p * x2 + 0.008333333333329196;
     	p = p * x2 + -0.16666666666666666;
     	p = p * x2 * x;
        result[0] = p;
	 }

	public static void polySineTvm(final double x, double[] result) {
     	 final TaskSchedule t0 = new TaskSchedule("s0")
     			 .task("t0", Foo::polySine, x, result)
     			 .streamOut(result);
   	 t0.execute();
   }  


    public static void main(String[] args) {
		double[] r = new double[1];
		polySineTvm(2.2132, r);
	}
}

I executed this program, and you will end up with something like this. TornadoVM compiler use GraalJIT for specializing the code. So we end-up with the following code:

$ tornado --debug --printKernel Foo
#pragma OPENCL EXTENSION cl_khr_fp64 : enable  
#pragma OPENCL EXTENSION cl_khr_int64_base_atomics : enable  
__kernel void polySine(__global uchar *_heap_base, ulong _frame_base, __constant uchar *_constant_region, __local uchar *_local_region, __global int *_atomics)
{
  ulong ul_0, ul_1; 

  __global ulong *_frame = (__global ulong *) &_heap_base[_frame_base];


  // BLOCK 0
  ul_0  =  (ulong) _frame[4];
  ul_1  =  ul_0 + 24L;
  *((__global double *) ul_1)  =  -1.4123909221508955;
  return;
}  //  kernel

Task info: s0.t0
	Backend           : OpenCL
	Device            : GeForce GTX 1050 CL_DEVICE_TYPE_GPU (available)
	Dims              : 0
	Global work offset: [0]
	Global work size  : [1]
	Local  work size  : [1, 1, 1]
	Number of workgroups  : [1]

Let us know if you have more questions.

[bradmesserle]

Thank you..This clarifies a lot!

But my setup is getting wrong results when off loading to the gpu..
The cpu seems to get the correct output.. but my gpu is like 2x the number it should i think..
but your post seems to be correct also.. Intel gpu cant add? :)

My testing results
Thanks.

tornado -Ds0.t0.device=0:1 -Dtornado.profiler=True -cp target/tvm-math.jar org.hipparchus.runnable.TornadoFastMathRunnable
1767036421380
General CPU
Result = -1.4123909221508955
457133
{
"s0": {
"TOTAL_BYTE_CODE_GENERATION": "41095319",
"COPY_OUT_TIME": "7885",
"TOTAL_DRIVER_COMPILE_TIME": "96585581",
"TOTAL_TASK_SCHEDULE_TIME": "227225881",
"TOTAL_GRAAL_COMPILE_TIME": "13388759",
"TOTAL_KERNEL_TIME": "7138",
"s0.t0": {
"METHOD": "TornadoFastMath.polySine",
"DEVICE_ID": "0:1",
"DEVICE": "Intel(R) UHD Graphics 630",
"TASK_COPY_OUT_SIZE_BYTES": "32",
"TASK_KERNEL_TIME": "7138",
"TASK_COMPILE_DRIVER_TIME": "96585581",
"TASK_COMPILE_GRAAL_TIME": "13388759"
}
}
}

Total Time227225881
GPU
1.589061484E-314
660883264

[jjfumero]

I just checked on my Intel HD Graphics, and it seems to work as well.

#pragma OPENCL EXTENSION cl_khr_int64_base_atomics : enable  
__kernel void polySine(__global uchar *_heap_base, ulong _frame_base, __constant uchar *_constant_region, __local uchar *_local_region, __global int *_atomics)
{
  ulong ul_0, ul_1; 

  __global ulong *_frame = (__global ulong *) &_heap_base[_frame_base];


  // BLOCK 0
  ul_0  =  (ulong) _frame[4];
  ul_1  =  ul_0 + 24L;
  *((__global double *) ul_1)  =  -1.4123909221508955;
  return;
}  //  kernel

Task info: s0.t0
	Backend           : OpenCL
	Device            : Intel(R) Graphics [0x591b] CL_DEVICE_TYPE_GPU (available)
	Dims              : 0
	Global work offset: [0]
	Global work size  : [1]
	Local  work size  : [1, 1, 1]
	Number of workgroups  : [1]

Can you run with --debug --printKernel? Feel free to open an issue. Which OpenCL driver are you using? I remember not long ago we encountered a few bugs in one of the drivers. We have a curated and maintained list of the drivers that we have tested with TornadoVM here

[bradmesserle]

Looks like the right number is in the open cl part..

I guess i don't know what driver versions etc. i am running.. should i do a git pull and rebuild everything?

thx

tornado -Ds0.t0.device=0:1 -Dtornado.profiler=True -cp target/tvm-math.jar org.hipparchus.runnable.TornadoFastMathRunnable --debug --printKernel
#pragma OPENCL EXTENSION cl_khr_fp64 : enable
#pragma OPENCL EXTENSION cl_khr_int64_base_atomics : enable
__kernel void lookupBufferAddress(__global uchar *_heap_base, ulong _frame_base, __constant uchar *_constant_region, __local uchar *_local_region, __global int *_atomics)
{

__global ulong *_frame = (__global ulong *) &_heap_base[_frame_base];

// BLOCK 0
_frame[0] = (ulong) _heap_base;
} // kernel

__kernel void lookupBufferAddress(__global uchar *_heap_base, ulong _frame_base, __constant uchar *_constant_region, __local uchar *_local_region, __global int *_atomics)
{

__global ulong *_frame = (__global ulong *) &_heap_base[_frame_base];

// BLOCK 0
_frame[0] = (ulong) _heap_base;
} // kernel

__kernel void polySine(__global uchar *_heap_base, ulong _frame_base, __constant uchar *_constant_region, __local uchar *_local_region, __global int *_atomics)
{
ulong ul_0, ul_1;

__global ulong *_frame = (__global ulong *) &_heap_base[_frame_base];

// BLOCK 0
ul_0 = (ulong) _frame[4];
ul_1 = ul_0 + 24L;
*((__global double *) ul_1) = -1.4123909221508955;
return;
} // kernel

task info: s0.t0
platform : Apple
device : Intel(R) UHD Graphics 630 CL_DEVICE_TYPE_GPU (available)
dims : 0
global work offset: [0]
global work size : [1]
local work size : [1, 1, 1]

{
"s0": {
"TOTAL_BYTE_CODE_GENERATION": "41787852",
"COPY_OUT_TIME": "7802",
"TOTAL_DRIVER_COMPILE_TIME": "704371",
"TOTAL_TASK_SCHEDULE_TIME": "130521775",
"TOTAL_GRAAL_COMPILE_TIME": "13539962",
"TOTAL_KERNEL_TIME": "6972",
"s0.t0": {
"METHOD": "TornadoFastMath.polySine",
"DEVICE_ID": "0:1",
"DEVICE": "Intel(R) UHD Graphics 630",
"TASK_COPY_OUT_SIZE_BYTES": "32",
"TASK_KERNEL_TIME": "6972",
"TASK_COMPILE_DRIVER_TIME": "704371",
"TASK_COMPILE_GRAAL_TIME": "13539962"
}
}
}

GPU
1.589061484E-314
562378803

[jjfumero]

Could you please open an issue with the test case you are running attached?
Thank you

[bradmesserle]

#67
My issue might be related to an existing issue...

[jjfumero]

If possible, could you please provide the code you are running? I can't reproduce the error.

[bradmesserle]

what do you think about creating an fpga aws image where everything works.. what do you think?

thx

[jjfumero]

Hi @bradmesserle , sorry for the delay. I missed this message.

Yes, that should work as well. But keep in mind the compute you are running is too simple to offset the cost of data transfers (from the CPU to the {GPU|FPGA}. You might want to increase the computation.

[bradmesserle]

yea. i know this is just a test and learning exp making sure i know how to use the framework and my setup is correct.

Thanks!