Closed yazun closed 2 years ago
nb. Error is similar for both CPU and GPU generated kernels.
GPU kernel and profile for completeness
........#pragma OPENCL EXTENSION cl_khr_fp64 : 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
#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
#pragma OPENCL EXTENSION cl_khr_fp64 : enable
#pragma OPENCL EXTENSION cl_khr_int64_base_atomics : enable
__kernel void computeGPU(__global uchar *_heap_base, ulong _frame_base, __constant uchar *_constant_region, __local uchar *_local_region, __global int *_atomics)
{
double d_85, d_87, d_86, d_73, d_72, d_75, d_74, d_77, d_76, d_65, d_64, d_67, d_66, d_69, d_68, d_71, d_70, d_57, d_56, d_59, d_58, d_61, d_60, d_63, d_62, d_113, d_112, d_115, d_114, d_117, d_116, d_105, d_104, d_107, d_106, d_109, d_108, d_111, d_110, d_97, d_96, d_99, d_98, d_101, d_100, d_103, d_102, d_89, d_88, d_91, d_90, d_93, d_92, d_95, d_94, d_17, d_19, d_18, d_20, d_23, d_22, d_15, d_48, d_53, d_52, d_55, d_54, d_41, d_43, d_42, d_45, d_44, d_47, d_46, d_33, d_32, d_35, d_34, d_37, d_36, d_38, d_25, d_27, d_26, d_29, d_28, d_31, d_30;
int i_79, i_78, i_49, i_51, i_50, i_21, i_84, i_39, i_118, i_24, i_40, i_10;
long l_11, l_82, l_80, l_81, l_12, l_13;
ulong ul_0, ul_16, ul_14, ul_83, ul_2, ul_1;
__global ulong *_frame = (__global ulong *) &_heap_base[_frame_base];
// BLOCK 0
ul_0 = (ulong) _frame[3];
ul_1 = (ulong) _frame[4];
ul_2 = (ulong) _frame[9];
__private double ul_3[1494001];
__private double ul_4[1494001];
__private double ul_5[1494001];
__private double ul_6[1494001];
__private double ul_7[1494001];
__private double ul_8[1494001];
__private double ul_9[1494001];
// BLOCK 1 MERGES [0 5 ]
i_10 = 0;
for(;i_10 < 165;)
{
// BLOCK 2
l_11 = (long) i_10;
l_12 = l_11 << 3;
l_13 = l_12 + 24L;
ul_14 = ul_0 + l_13;
d_15 = *((__global double *) ul_14);
ul_16 = ul_1 + l_13;
d_17 = *((__global double *) ul_16);
d_18 = d_15 * 0.37699111843077515;
d_19 = native_cos(d_18);
d_20 = native_sin(d_18);
i_21 = get_global_id(1);
// BLOCK 3 MERGES [2 4 ]
d_22 = d_20;
d_23 = d_19;
i_24 = i_21;
for(;i_24 < 1494001;)
{
// BLOCK 4
d_25 = ul_3[i_24];
d_26 = d_22 + d_25;
ul_3[i_24] = d_26;
d_25 = d_26;
d_27 = ul_4[i_24];
d_28 = d_23 + d_27;
ul_4[i_24] = d_28;
d_27 = d_28;
d_29 = ul_5[i_24];
d_30 = fma(d_22, d_22, d_29);
ul_5[i_24] = d_30;
d_29 = d_30;
d_31 = ul_6[i_24];
d_32 = fma(d_23, d_23, d_31);
ul_6[i_24] = d_32;
d_31 = d_32;
d_33 = ul_7[i_24];
d_34 = fma(d_22, d_23, d_33);
ul_7[i_24] = d_34;
d_33 = d_34;
d_35 = ul_8[i_24];
d_36 = fma(d_22, d_17, d_35);
ul_8[i_24] = d_36;
d_35 = d_36;
d_37 = ul_9[i_24];
d_38 = fma(d_23, d_17, d_37);
ul_9[i_24] = d_38;
d_37 = d_38;
i_39 = get_global_size(1);
i_40 = i_39 + i_24;
d_41 = d_15 * 6.283185307179587E-5;
d_42 = native_cos(d_41);
d_43 = d_42 * d_23;
d_44 = native_sin(d_41);
d_45 = d_44 * d_22;
d_46 = d_43 - d_45;
d_47 = d_44 * d_23;
d_48 = fma(d_42, d_22, d_47);
d_22 = d_48;
d_23 = d_46;
i_24 = i_40;
} // B4
// BLOCK 5
i_49 = i_10 + 1;
i_10 = i_49;
} // B5
// BLOCK 6
i_50 = get_global_id(0);
// BLOCK 7 MERGES [6 11 ]
i_51 = i_50;
for(;i_51 < 1494001;)
{
// BLOCK 8
d_52 = ul_6[i_51];
d_53 = ul_5[i_51];
d_54 = ul_7[i_51];
d_55 = ul_7[i_51];
d_56 = ul_4[i_51];
d_57 = ul_4[i_51];
d_58 = ul_5[i_51];
d_59 = ul_3[i_51];
d_60 = ul_3[i_51];
d_61 = ul_6[i_51];
d_62 = ul_3[i_51];
d_63 = ul_4[i_51];
d_64 = ul_7[i_51];
d_65 = d_62 * 2.0;
d_66 = d_65 * d_63;
d_67 = d_52 * d_53;
d_68 = d_54 * d_55;
d_69 = d_67 - d_68;
d_70 = d_69 * 165.0;
d_71 = d_56 * d_57;
d_72 = d_71 * d_58;
d_73 = d_70 - d_72;
d_74 = d_59 * d_60;
d_75 = d_74 * d_61;
d_76 = d_73 - d_75;
d_77 = fma(d_66, d_64, d_76);
i_78 = get_global_size(0);
i_79 = i_78 + i_51;
l_80 = (long) i_51;
l_81 = l_80 << 3;
l_82 = l_81 + 24L;
ul_83 = ul_2 + l_82;
i_84 = isless(d_77, 0.0);
if(i_84 == 1)
{
// BLOCK 9
*((__global double *) ul_83) = 1.0;
} // B9
else
{
// BLOCK 10
d_85 = ul_4[i_51];
d_86 = ul_8[i_51];
d_87 = ul_3[i_51];
d_88 = ul_9[i_51];
d_89 = ul_9[i_51];
d_90 = ul_5[i_51];
d_91 = ul_8[i_51];
d_92 = ul_7[i_51];
d_93 = ul_3[i_51];
d_94 = ul_8[i_51];
d_95 = ul_6[i_51];
d_96 = ul_9[i_51];
d_97 = ul_7[i_51];
d_97 = ul_7[i_51];
d_97 = ul_7[i_51];
d_97 = ul_7[i_51];
d_98 = ul_4[i_51];
d_99 = ul_9[i_51];
d_100 = ul_8[i_51];
d_101 = d_85 * d_86;
d_102 = d_87 * d_88;
d_103 = d_101 - d_102;
d_104 = d_89 * d_90;
d_105 = d_91 * d_92;
d_106 = d_104 - d_105;
d_107 = d_106 * 165.0;
d_108 = fma(d_103, d_93, d_107);
d_109 = d_94 * d_95;
d_110 = d_96 * d_97;
d_111 = d_109 - d_110;
d_112 = d_111 * 165.0;
d_113 = d_98 * d_103;
d_114 = d_112 - d_113;
d_115 = d_114 * d_100;
d_116 = fma(d_108, d_99, d_115);
d_117 = d_116 / d_77;
*((__global double *) ul_83) = d_117;
} // B10
// BLOCK 11 MERGES [9 10 ]
i_118 = i_79;
i_51 = i_118;
} // B11
// BLOCK 12
return;
} // kernel
[TornadoVM-OCL-JNI] ERROR : clCreateKernel -> Returned: -5
[TornadoVM-OCL-JNI] ERROR : clGetKernelInfo -> Returned: 4294967248
[TornadoVM-OCL-JNI] ERROR : clSetKernelArg -> Returned: 4294967248
[TornadoVM-OCL-JNI] ERROR : clSetKernelArg -> Returned: 4294967248
[TornadoVM-OCL-JNI] ERROR : clSetKernelArg -> Returned: 4294967248
[TornadoVM-OCL-JNI] ERROR : clSetKernelArg -> Returned: 4294967248
[TornadoVM-OCL-JNI] ERROR : clSetKernelArg -> Returned: 4294967248
[TornadoVM-OCL-JNI] ERROR : clEnqueueNDRangeKernel -> Returned: -48
[TornadoVM-OCL-JNI] ERROR : clWaitForEvents -> Returned: -58
[TornadoVM-OCL-JNI] ERROR : clGetEventProfilingInfo -> Returned: -58
[TornadoVM-OCL-JNI] ERROR : clGetEventProfilingInfo -> Returned: -58
[TornadoVM-OCL-JNI] ERROR : clGetEventProfilingInfo -> Returned: -58
[TornadoVM-OCL-JNI] ERROR : clGetEventProfilingInfo -> Returned: -58
[TornadoVM-OCL-JNI] ERROR : clGetEventProfilingInfo -> Returned: -58
[TornadoVM-OCL-JNI] ERROR : clGetEventProfilingInfo -> Returned: -58
{
"PeriodSearch": {
"TOTAL_BYTE_CODE_GENERATION": "7627479",
"COPY_IN_TIME": "3906",
"TOTAL_TASK_SCHEDULE_TIME": "1284552037",
"TOTAL_KERNEL_TIME": "0",
"COPY_OUT_TIME": "972347",
"TOTAL_GRAAL_COMPILE_TIME": "88152834",
"TOTAL_DRIVER_COMPILE_TIME": "343693118",
"TOTAL_DISPATCH_DATA_TRANSFERS_TIME": "23902",
"TOTAL_DISPATCH_KERNEL_TIME": "0",
"TOTAL_COPY_IN_SIZE_BYTES": "5376",
"TOTAL_COPY_OUT_SIZE_BYTES": "11952032",
"PeriodSearch.t0": {
"METHOD": "MethodLeastSquareGPU.computeGPU",
"DEVICE_ID": "0:1",
"DEVICE": "Intel(R) Iris(TM) Pro Graphics P580 [0x193a]",
"TOTAL_COPY_IN_SIZE_BYTES": "80",
"TASK_KERNEL_TIME": "0",
"TASK_COMPILE_DRIVER_TIME": "343693118",
"TASK_COMPILE_GRAAL_TIME": "88152834"
}
}
}
There are a few things here:
From the code being generated:
How to parallelize the Java application:
I suggest splitting the task (method) into several tasks. Each method runs a parallel loop. For the sequential part, it will need to run on the host side and call a parallel kernel per iteration. Having smaller methods (compute-kernels) has the consequence of having more space for private memory of large arrays.
Hope this helps.
Thanks for checking and for the hints!
After few iterations it's quite clear that this particular algorithm is a hard case - there too many intra-iteration dependencies between computations to disentangle them. It would be easier if some of the loops could be marked as non-parallelizable in the kernel code, but from what I understand there's no such control - i.e. one cannot specify that only the outer loop is parallelized when we have two loops? We have the code running now, but obviously it does not give proper results due to
private static void computeGPU(double[] cPhArr,double sPhArr[], double[] cTArr, double[] normObsTimes, double[] normObsValues, double[] sumSx, double[] sumCx, double[] sumSx2, double[] sumCx2, double[] sumSxCx, double[] sumSxVal, double[] sumCxVal, double omega, double deltaOmega, double deltaEpsilon)
{
int nObservations = normObsTimes.length;
int nScanFrequencies = cPhArr.length;
// for each set of observation data
for (@Parallel int i = 0; i < nObservations; i++) {
final double obsValue = normObsValues[i];
// for each frequency to test, increment the phase with the phase step
for ( int j = 0; j < nScanFrequencies; j++) {
sumSx[j] += sPhArr[j];
sumCx[j] += cPhArr[j];
sumSx2[j] += sPhArr[j]*sPhArr[j];
sumCx2[j] += cPhArr[j]*cPhArr[j];
sumSxCx[j] += cPhArr[j]*sPhArr[j];
sumSxVal[j] += sPhArr[j]*obsValue;
sumCxVal[j] += cPhArr[j]*obsValue;
}
}
...
TaskSchedule task = new TaskSchedule("PeriodSearch")
.streamIn(sumSx, sumCx, sumSx2, sumCx2, sumSxCx, sumSxVal, sumCxVal)
.task("t0", MethodLeastSquareGPU::computeGPU, cPhArr,sPhArr, cTArr, normObsTimes, normObsValues, sumSx, sumCx, sumSx2, sumCx2, sumSxCx, sumSxVal, sumCxVal, omega, deltaOmega, deltaEpsilon)
.streamOut(amplitudes,sumSx, sumCx, sumSx2, sumCx2, sumSxCx, sumSxVal, sumCxVal);
task.execute();
will not give mathematically valid result. While not surprising I wonder if such cases could be intercepted by the kernel compiler so a warning/error is thrown?
Thank you for the suggestion. Yes I agree TornadoVM should be able to at least, throw the error and bailout. We annotate this for future work.
Meanwhile, keep in mind TornadoVM offers two APIs:
It might be a better fit for your usecase.
Hope this helps.
Yes, porting to Kernel Parallel API will be the next step now. Thanks!
@yazun , is this issue solved? Can this issue be closed?
yes, thanks!
Describe the bug This is continuation of the issue https://github.com/beehive-lab/TornadoVM/issues/152 but with asm ver 9.2 replacing bundled ver 7.2
How To Reproduce Same test as in https://github.com/beehive-lab/TornadoVM/issues/152
The code of the computeGPU function looks like:
This time the kernel is generated but the code fails.
test fails with empty array returned.
I understand
[TornadoVM-OCL-JNI] ERROR : clEnqueueNDRangeKernel -> Returned: -5
isCL_OUT_OF_RESOURCES
and-58 is CL_INVALID_EVENT
?Expected behavior
To run without any error.
Computing system setup (please complete the following information):
tornado --devices
Additional context Normally, nScanFrequencies could be 10-1000x higher but we lowered it for the sake of the test.