CL_BUILD_PROGRAM_FAILURE on 02MDF (GW-opencl-ati)

Sagittarius Lupus
Joined: 14 May 10
Posts: 5
Credit: 15,399,022
RAC: 3,885
Topic 220609

A batch of tasks for Gravitational Wave search O2 Multi-Directional GPU v2.07 () x86_64-pc-linux-gnu just failed fast on my system which I've been using successfully for other OpenCL projects. In this case, the program fails to compile with a detailed build error that I hope will be useful to the program developers here. 

 

Here's a link to one of the tasks. The portion of the program log containing the build error is this:

 

XLAL Error - XLALGetOpenCLProgramFromSource (/home/jenkins/workspace/workspace/EaH-GW-OpenCL-Testing/SLAVE/LIBC215/TARGET/linux-x86_64/EinsteinAtHome/source/lalsuite/lalpulsar/src/OpenCLutils.c:612): clBuildProgram failed: CL_BUILD_PROGRAM_FAILURE . Error message: /tmp/comgr-793bba/input/CompileSource:175:42: error: passing 'REAL4 *' (aka 'float *') to parameter of type 'const REAL4 *' (aka 'const float *') changes address space of pointer outBSGL[j] = XLALComputeBSGL( twoF[j], twoFPerDet_local, setup ); ^~~~~~~~~~~~~~~~ /tmp/comgr-793bba/input/CompileSource:70:31: note: passing argument to parameter 'twoFPerDet' here const REAL4 twoFPerDet[PULSAR_MAX_DETECTORS], ^ /tmp/comgr-793bba/input/CompileSource:194:46: error: passing 'REAL4 *' (aka 'float *') to parameter of type 'const REAL4 *' (aka 'const float *') changes address space of pointer outBSGLtL[j] = XLALComputeBSGLtL( twoF[j], twoFPerDet_local, maxTwoFSegPerDet_local, setup); ^~~~~~~~~~~~~~~~ /tmp/comgr-793bba/input/CompileSource:97:33: note: passing argument to parameter 'twoFPerDet' here c XLAL Error - XLALGetOpenCLProgramFromSource (/home/jenkins/workspace/workspace/EaH-GW-OpenCL-Testing/SLAVE/LIBC215/TARGET/linux-x86_64/EinsteinAtHome/source/lalsuite/lalpulsar/src/OpenCLutils.c:612): Internal function call failed XLAL Error - MAIN (/home/jenkins/workspace/workspace/EaH-GW-OpenCL-Testing/SLAVE/LIBC215/TARGET/linux-x86_64/EinsteinAtHome/source/lalsuite/lalapps/src/pulsar/GCT/HierarchSearchGCT.c:1276): Check failed: XLALGetOpenCLProgramFromSource ( source, &HierarchSearchGCTProgramm ) == XLAL_SUCCESS XLAL Error - MAIN (/home/jenkins/workspace/workspace/EaH-GW-OpenCL-Testing/SLAVE/LIBC215/TARGET/linux-x86_64/EinsteinAtHome/source/lalsuite/lalapps/src/pulsar/GCT/HierarchSearchGCT.c:1276): Internal function call failed 2020-02-02 08:09:54.7560 (3970521) [CRITICAL]: ERROR: MAIN() returned with error '-1'

 

I have highlighted the messages I believe to be critical to understanding the failure. All of the tasks in this batch failed in exactly the same way.

 

I'm using the ROCm (Radeon Open Compute) drivers for OpenCL on Linux 5.3.7 with a Radeon RX Vega 64, supporting the full OpenCL 2.0 profile. Here is the output of clinfo:

Number of platforms 1 Platform Name AMD Accelerated Parallel Processing Platform Vendor Advanced Micro Devices, Inc. Platform Version OpenCL 2.0 AMD-APP.internal (3052.0) Platform Profile FULL_PROFILE Platform Extensions cl_khr_icd cl_amd_object_metadata cl_amd_event_callback Platform Max metadata object keys (AMD) 8 Platform Extensions function suffix AMDPlatform Name AMD Accelerated Parallel Processing Number of devices 1 Device Name gfx900 Device Vendor Advanced Micro Devices, Inc. Device Vendor ID 0x1002 Device Version OpenCL 2.0 Driver Version 3052.0 (HSA1.1,LC) Device OpenCL C Version OpenCL C 2.0 Device Type GPU Device Board Name (AMD) Vega 10 XL/XT [Radeon RX Vega 56/64] Device Topology (AMD) PCI-E, 0b:00.0 Device Profile FULL_PROFILE Device Available Yes Compiler Available Yes Linker Available Yes Max compute units 64 SIMD per compute unit (AMD) 4 SIMD width (AMD) 16 SIMD instruction width (AMD) 1 Max clock frequency 1630MHz Graphics IP (AMD) 9.0 Device Partition (core) Max number of sub-devices 64 Supported partition types None Supported affinity domains (n/a) Max work item dimensions 3 Max work item sizes 1024x1024x1024 Max work group size 256 Preferred work group size (AMD) 256 Max work group size (AMD) 1024 Preferred work group size multiple 64 Wavefront width (AMD) 64 Preferred / native vector sizes char 4 / 4 short 2 / 2 int 1 / 1 long 1 / 1 half 1 / 1 (cl_khr_fp16) float 1 / 1 double 1 / 1 (cl_khr_fp64) Half-precision Floating-point support (cl_khr_fp16) Denormals No Infinity and NANs No Round to nearest No Round to zero No Round to infinity No IEEE754-2008 fused multiply-add No Support is emulated in software No Single-precision Floating-point support (core) Denormals Yes Infinity and NANs Yes Round to nearest Yes Round to zero Yes Round to infinity Yes IEEE754-2008 fused multiply-add Yes Support is emulated in software No Correctly-rounded divide and sqrt operations Yes Double-precision Floating-point support (cl_khr_fp64) Denormals Yes Infinity and NANs Yes Round to nearest Yes Round to zero Yes Round to infinity Yes IEEE754-2008 fused multiply-add Yes Support is emulated in software No Address bits 64, Little-Endian Global memory size 8573157376 (7.984GiB) Global free memory (AMD) 8372224 (7.984GiB) Global memory channels (AMD) 64 Global memory banks per channel (AMD) 4 Global memory bank width (AMD) 256 bytes Error Correction support No Max memory allocation 7287183769 (6.787GiB) Unified memory for Host and Device No Shared Virtual Memory (SVM) capabilities (core) Coarse-grained buffer sharing Yes Fine-grained buffer sharing Yes Fine-grained system sharing No Atomics No Minimum alignment for any data type 128 bytes Alignment of base address 1024 bits (128 bytes) Preferred alignment for atomics SVM 0 bytes Global 0 bytes Local 0 bytes Max size for global variable 7287183769 (6.787GiB) Preferred total size of global vars 8573157376 (7.984GiB) Global Memory cache type Read/Write Global Memory cache size 16384 (16KiB) Global Memory cache line size 64 bytes Image support No Base address alignment for 2D image buffers 0 bytes Pitch alignment for 2D image buffers 0 pixels Max number of pipe args 16 Max active pipe reservations 16 Max pipe packet size 2992216473 (2.787GiB) Local memory type Local Local memory size 65536 (64KiB) Local memory syze per CU (AMD) 65536 (64KiB) Local memory banks (AMD) 32 Max number of constant args 8 Max constant buffer size 7287183769 (6.787GiB) Preferred constant buffer size (AMD) 16384 (16KiB) Max size of kernel argument 1024 Queue properties (on host) Out-of-order execution No Profiling Yes Queue properties (on device) Out-of-order execution Yes Profiling Yes Preferred size 262144 (256KiB) Max size 8388608 (8MiB) Max queues on device 1 Max events on device 1024 Prefer user sync for interop Yes Number of P2P devices (AMD) 0 P2P devices (AMD) <printDeviceInfo:147: get number of CL_DEVICE_P2P_DEVICES_AMD : error -30> Profiling timer resolution 1ns Profiling timer offset since Epoch (AMD) 0ns (Wed Dec 31 19:00:00 1969) Execution capabilities Run OpenCL kernels Yes Run native kernels No Thread trace supported (AMD) No Number of async queues (AMD) 8 Max real-time compute queues (AMD) 8 Max real-time compute units (AMD) 64 printf() buffer size 4194304 (4MiB) Built-in kernels (n/a) Device Extensions cl_khr_fp64 cl_khr_global_int32_base_atomics cl_khr_global_int32_extended_atomics cl_khr_local_int32_base_atomics cl_khr_local_int32_extended_atomics cl_khr_int64_base_atomics cl_khr_int64_extended_atomics cl_khr_3d_image_writes cl_khr_byte_addressable_store cl_khr_fp16 cl_khr_gl_sharing cl_amd_device_attribute_query cl_amd_media_ops cl_amd_media_ops2 cl_khr_image2d_from_buffer cl_khr_subgroups cl_khr_depth_images cl_amd_copy_buffer_p2p cl_amd_assembly_programNULL platform behavior clGetPlatformInfo(NULL, CL_PLATFORM_NAME, ...) AMD Accelerated Parallel Processing clGetDeviceIDs(NULL, CL_DEVICE_TYPE_ALL, ...) Success [AMD] clCreateContext(NULL, ...) [default] Success [AMD] clCreateContextFromType(NULL, CL_DEVICE_TYPE_DEFAULT) Success (1) Platform Name AMD Accelerated Parallel Processing Device Name gfx900 clCreateContextFromType(NULL, CL_DEVICE_TYPE_CPU) No devices found in platform clCreateContextFromType(NULL, CL_DEVICE_TYPE_GPU) Success (1) Platform Name AMD Accelerated Parallel Processing Device Name gfx900 clCreateContextFromType(NULL, CL_DEVICE_TYPE_ACCELERATOR) No devices found in platform clCreateContextFromType(NULL, CL_DEVICE_TYPE_CUSTOM) No devices found in platform clCreateContextFromType(NULL, CL_DEVICE_TYPE_ALL) Success (1) Platform Name AMD Accelerated Parallel Processing Device Name gfx900ICD loader properties ICD loader Name OpenCL ICD Loader ICD loader Vendor OCL Icd free software ICD loader Version 2.2.12 ICD loader Profile OpenCL 2.2

Please let me know if you need more information. I'm eager to solve this problem.

Sagittarius Lupus
Joined: 14 May 10
Posts: 5
Credit: 15,399,022
RAC: 3,885

Anyone? I have to suspend

Anyone? I have to suspend this project until the problem is solved, since every GPU task is failing immediately.

Keith Myers
Keith Myers
Joined: 11 Feb 11
Posts: 502
Credit: 502,811,794
RAC: 1,054,252

Pretty sure NO BOINC project

Pretty sure NO BOINC project supports RocM drivers.  You must use the stock proprietary AMD drivers with the OpenCL legacy option during install.

BoincStats

Sagittarius Lupus
Joined: 14 May 10
Posts: 5
Credit: 15,399,022
RAC: 3,885

Keith Myers wrote:Pretty sure

Keith Myers wrote:
Pretty sure NO BOINC project supports RocM drivers.  You must use the stock proprietary AMD drivers with the OpenCL legacy option during install.

That's not helpful. All the other GPU projects I've signed up for are working with these drivers. None of those projects have any opinion concerning the OpenCL ICD in use. You do not appear to be affiliated with this project; kindly refrain from telling others what they "must" do.

Besides. If this were a driver problem, I would be seeing the application attempt to execute -- but this is a build problem. It isn't getting that far. The code isn't even compiling. That means I need one of the project developers to respond -- someone who is actually qualified to examine the code.

It may be possible that my system OpenCL headers are at a different version than the code is compatible with, but again, that is a question only the project developers can answer.

Edit: my originally reported task seems to have aged out and is no longer viewable from the link in the OP. Here are my current tasks.

Sagittarius Lupus
Joined: 14 May 10
Posts: 5
Credit: 15,399,022
RAC: 3,885

I've gone ahead and

I've gone ahead and researched the problem myself. Standard disclaimer: I am not a C programmer, and I do not program in OpenCL, but I am literate, and I can break apart this program just enough to see what's going on.

I am about to get very technical, so let's run it from the top.

When you run this application in BOINC on Linux with an AMD GPU, you are running a binary called einstein_O2MDF_2.07_x86_64-pc-linux-gnu__GW-opencl-ati for this particular version of the app. This is a big, opaque thing to me, and I can't see most of what it's doing, but I took one of my jobs out of BOINC to try and run it standalone. The thing runs from a BOINC slot with a command line like this:

../../projects/einstein.phys.uwm.edu/einstein_O2MDF_2.07_x86_64-pc-linux-gnu__GW-opencl-ati --refTime=1177858472.0 --Freq=1045 --FreqBand=0.05 --dFreq=4.035776e-07 --f1dot=-4.547049e-08 --f1dotBand=2.298573e-10 --df1dot=2.558432e-12 --gammaRefine=31 --f2dot=0 --f2dotBand=1.499184e-17 --df2dot=1.356969e-18 --gamma2Refine=51 --computeBSGL --BSGLlogcorr=0 --Fstar0=81.325 --oLGX=0.5,0.5 --nCand1=10000 --SortToplist=7 --recalcToplistStats=1 --outputTiming=GCT.timing --outputTimingDetails=GCT.timing -o GCT.out --printCand1 --semiCohToplist --ephemE=earth --ephemS=sun --segmentList=segments --FstatMethod=ResampBest --FstatMethodRecalc=DemodBest --gridType=3 --skyGridFile={2.321389134249074 -0.808054282417587} --loudestSegOutput --getMaxFperSeg --peakThrF=2.6 --resampTurnOffNarrowbanding=TRUE --DataFiles1='h1_1044.40_O2C02Cl4In0.uisl;l1_1044.40_O2C02Cl4In0.uisl;h1_1044.45_O2C02Cl4In0.4o3E;l1_1044.45_O2C02Cl4In0.4o3E;h1_1044.50_O2C02Cl4In0.KvXE;l1_1044.50_O2C02Cl4In0.KvXE;h1_1044.55_O2C02Cl4In0.pTWG;l1_1044.55_O2C02Cl4In0.pTWG;h1_1044.60_O2C02Cl4In0.4LWe;l1_1044.60_O2C02Cl4In0.4LWe;h1_1044.65_O2C02Cl4In0.6YaM;l1_1044.65_O2C02Cl4In0.6YaM;h1_1044.70_O2C02Cl4In0.QlxB;l1_1044.70_O2C02Cl4In0.QlxB;h1_1044.75_O2C02Cl4In0.I4qK;l1_1044.75_O2C02Cl4In0.I4qK;h1_1044.80_O2C02Cl4In0.CHYs;l1_1044.80_O2C02Cl4In0.CHYs;h1_1044.85_O2C02Cl4In0.iRE-;l1_1044.85_O2C02Cl4In0.iRE-;h1_1044.90_O2C02Cl4In0.ZCHm;l1_1044.90_O2C02Cl4In0.ZCHm;h1_1044.95_O2C02Cl4In0.saW0;l1_1044.95_O2C02Cl4In0.saW0;h1_1045.00_O2C02Cl4In0.W_4R;l1_1045.00_O2C02Cl4In0.W_4R;h1_1045.05_O2C02Cl4In0.6s89;l1_1045.05_O2C02Cl4In0.6s89;h1_1045.10_O2C02Cl4In0.tU8m;l1_1045.10_O2C02Cl4In0.tU8m;h1_1045.15_O2C02Cl4In0.J1U0;l1_1045.15_O2C02Cl4In0.J1U0;h1_1045.20_O2C02Cl4In0.EmvQ;l1_1045.20_O2C02Cl4In0.EmvQ;h1_1045.25_O2C02Cl4In0.c_nX;l1_1045.25_O2C02Cl4In0.c_nX;h1_1045.30_O2C02Cl4In0.0Xzn;l1_1045.30_O2C02Cl4In0.0Xzn;h1_1045.35_O2C02Cl4In0.IEck;l1_1045.35_O2C02Cl4In0.IEck;h1_1045.40_O2C02Cl4In0.bcrl;l1_1045.40_O2C02Cl4In0.bcrl;h1_1045.45_O2C02Cl4In0.9-f1;l1_1045.45_O2C02Cl4In0.9-f1;h1_1045.50_O2C02Cl4In0.RQkR;l1_1045.50_O2C02Cl4In0.RQkR;h1_1045.55_O2C02Cl4In0.EhZ7;l1_1045.55_O2C02Cl4In0.EhZ7;h1_1045.60_O2C02Cl4In0.ekrn;l1_1045.60_O2C02Cl4In0.ekrn;h1_1045.65_O2C02Cl4In0.Fytl;l1_1045.65_O2C02Cl4In0.Fytl' --WUfpops=2.88e+14 --device 0

That is a doozy. Almost all of the command arguments are just the parameters of that particular job and don't matter to the problem -- they're just necessary to get the program running where I can retry several times and debug its execution.

As the standard error stream indicates, very early on the app has to build several OpenCL kernels. One of these kernels has a problem with address space casting. Here's the message from the build log again:

XLAL Error - XLALGetOpenCLProgramFromSource (/home/jenkins/workspace/workspace/EaH-GW-OpenCL-Testing/SLAVE/LIBC215/TARGET/linux-x86_64/EinsteinAtHome/source/lalsuite/lalpulsar/src/OpenCLutils.c:612): clBuildProgram failed: CL_BUILD_PROGRAM_FAILURE . Error message:
/tmp/comgr-8d6b77/input/CompileSource:175:42: error: passing 'REAL4 *' (aka 'float *') to parameter of type 'const REAL4 *' (aka 'const float *') changes address space of pointer
  outBSGL[j] = XLALComputeBSGL( twoF[j], twoFPerDet_local, setup );
                                         ^~~~~~~~~~~~~~~~

I'm not sure whether the preformatted text here is going to render properly after I submit the post, so I've added my own emphasis. Here's what the error means: the compiler is indicating that a parameter called twoFPerDet_local was given the type 'const REAL4 *' in the function definition of XLALComputeBSGL(), but the argument passed to the function for this parameter has a different type: just 'REAL4 *'. But this is only part of the problem -- the real issue is the "changes address space of pointer" bit. More on that shortly.

Since I ran the app standalone, I was able to grab the temporary files it creates on my machine for building the OpenCL kernels, which contain the source code.

Here's the function definition for XLALComputeBSGL():

REAL4
 XLALComputeBSGL ( const REAL4 twoF,
                   const REAL4 twoFPerDet[PULSAR_MAX_DETECTORS],
                   __global const BSGLSetup *setup
                   )
 {
   // meaty, irrelevant stuff
 }

Pay close attention to the definition of the second parameter: const REAL4 twoFPerDet[PULSAR_MAX_DETECTORS]

Now, here's where something very, very subtle happens. Here's the definition of another function -- the one that fails to compile:

__kernel void XLALVectorComputeBSGL ( __global REAL4 *outBSGL,
                                      __global const REAL4 *twoF,
                                      __global const REAL4 *twoFPerDet,
                                      const UINT4 len,
                                      __global const BSGLSetup *setup
                                      )
 {
   UINT4 j = get_global_id(0);
   if (j >= len) return;
   REAL4 twoFPerDet_local[PULSAR_MAX_DETECTORS];
   for ( UINT4 i = 0 ; i < PULSAR_MAX_DETECTORS; i++ )
     {
       twoFPerDet_local[i] = twoFPerDet[j+len*i];
     }
   outBSGL[j] = XLALComputeBSGL( twoF[j], twoFPerDet_local, setup );
 } 

This is where address spaces come in. See the __global qualifiers on all of those variables? We're inside a __kernel function, which has special rules. Namely:

__kernel function arguments declared to be a pointer of a type can point to one of the following address spaces only: __global, __local or __constant. A pointer to address space A can only be assigned to a pointer to the same address space A. Casting a pointer to address space A to a pointer to address space B is illegal.

At some point inside of this function, we declare a local variable and initialize it from a pointer in the __global address space. Then we go and pass that local variable as an argument to another function that has no address space qualifier on the corresponding parameter. The compiler does not like this.

The same thing happens a bit later in another kernel function XLALVectorComputeBSGLtL() when it passes an argument with similar properties to XLALComputeBSGLtL().

Now some details get a little fuzzy to me: this isn't biting everyone, is it? It appears that some compilers are more strict than others. If you're on a system that supports OpenCL 2.0 (like mine), there is a feature in this version that allows programmers to get away with not explicitly qualifying the address space in some situations, but you have to turn it on when you call clBuildProgram() or the default 1.2 compatibility specifications will cause this to blow up. It's not clear to me whether the Einstein GPU apps are supposed to be OpenCL 1.2 or OpenCL 2.0 programs.

What is clear to me is that if the programmers simply reconcile the address space qualifiers for these variables, we'll get past this specific problem on all compilers. I can't say what the prevalence of this issue is because it is doubtful most casual users are technical enough to properly report it let alone notice it, but a bit of googling shows that programmers have been running into it all over the place.

Moderators: consider this a bug report, please. If you had a bug tracker or a git repository somewhere, I'd be happy to submit a patch.

Gary Roberts
Gary Roberts
Moderator
Joined: 9 Feb 05
Posts: 5,035
Credit: 34,660,383,349
RAC: 34,789,964

Sagittarius Lupus

Sagittarius Lupus wrote:
That's not helpful. All the other GPU projects I've signed up for are working with these drivers. None of those projects have any opinion concerning the OpenCL ICD in use. You do not appear to be affiliated with this project; kindly refrain from telling others what they "must" do.

Was it really necessary to do that??  You post about a problem you are having that is clearly well outside the skill set of the bulk of ordinary volunteers, you beg for anyone to respond, and when someone actually tries to be helpful, you reward that with quite some invective.  If you're so confident that you know exactly where the problem lies, then I suggest you desist from deliberately insulting a respondent just trying to help and wait patiently for someone who is qualified to study the problem in depth and is, perhaps, able to do something about it.

If you know anything about how University type research environments operate, you would realise that often the scientists themselves and not professional teams of programmers, are the ones that have to do a lot of the grunt work in developing the very specialised applications needed.  There is often insufficient funding and resources to make that happen quickly.  You haven't received a single 'developer' type response to your problem, after all this time that your message has been up.  Doesn't that tell you something?  It suggests to me that they are already stretched to the limit in supporting what they have already without delving into new and uncharted territory.  They would also be concluding that you could make changes to the way you operate so that the tasks wouldn't end in error.  That would tend to lower the urgency for looking into this particular problem.

I'm not a programmer at all so I have no idea of how to solve your problem.  I would give you much the same response as Keith, except that I would word it this way:-

AMD are on a path to developing ROCm as the future framework under which compute applications will run.  It's under active development and (reading between the lines) it's not ready for prime time yet and it's not surprising that some more 'traditional' applications don't yet work properly with it.  We know for certain that your type of GPU can run the current science apps using the existing OpenCL implementations that (hopefully one day) will be replaced by ROCm when it is fully developed.  If you wish to run these Einstein apps now, and in the absence of any staff response, you may need to consider using the currently supported OpenCL implementations.

Sagittarius Lupus wrote:
Edit: my originally reported task seems to have aged out and is no longer viewable from the link in the OP. Here are my current tasks.

Thanks for providing the new link.

If you ever need to post excerpts like those shown in your first post, please take a look at the toolbar across the top of the message composition window.  At the far right hand end (3rd from the right) is a "Post as plain text" icon which allows you to get rid of the hidden control characters which otherwise disrupt the formatting of what you are trying to show.  So just click that icon before pasting the content you wish to show.  Post the content between [ code ] tags which should preserve the spacing of any columnised data.  You can also add [ size ] and [ font ] tags to change both those things.  The BBCODE HELP below the composition window is also quite good to read.  Here is an example of your latest error message using this technique.  It's tends to be somewhat easier to read :-).

% --- GPS reference time = 1177858472.0000 , GPS data mid time = 1177858472.0000
2020-02-16 15:17:00.1710 (2750439) [normal]: dFreqStack = 4.035776e-07, df1dot = 2.558432e-12, df2dot = 1.356969e-18, df3dot = 0.000000e+00
% --- Setup, N = 17, T = 864000 s, Tobs = 19750204 s, gammaRefine = 31, gamma2Refine = 51, gamma3Refine = 1
XLAL Error - XLALGetOpenCLProgramFromSource (/home/jenkins/workspace/workspace/EaH-GW-OpenCL-Testing/SLAVE/LIBC215/TARGET/linux-x86_64/EinsteinAtHome/source/lalsuite/lalpulsar/src/OpenCLutils.c:612): clBuildProgram failed: CL_BUILD_PROGRAM_FAILURE . Error message:
/tmp/comgr-8d6b77/input/CompileSource:175:42: error: passing 'REAL4 *' (aka 'float *') to parameter of type 'const REAL4 *' (aka 'const float *') changes address space of pointer
outBSGL[j] = XLALComputeBSGL( twoF[j], twoFPerDet_local, setup );
^~~~~~~~~~~~~~~~
/tmp/comgr-8d6b77/input/CompileSource:70:31: note: passing argument to parameter 'twoFPerDet' here
const REAL4 twoFPerDet[PULSAR_MAX_DETECTORS],
^
/tmp/comgr-8d6b77/input/CompileSource:194:46: error: passing 'REAL4 *' (aka 'float *') to parameter of type 'const REAL4 *' (aka 'const float *') changes address space of pointer
outBSGLtL[j] = XLALComputeBSGLtL( twoF[j], twoFPerDet_local, maxTwoFSegPerDet_local, setup);
^~~~~~~~~~~~~~~~
/tmp/comgr-8d6b77/input/CompileSource:97:33: note: passing argument to parameter 'twoFPerDet' here
c
XLAL Error - XLALGetOpenCLProgramFromSource (/home/jenkins/workspace/workspace/EaH-GW-OpenCL-Testing/SLAVE/LIBC215/TARGET/linux-x86_64/EinsteinAtHome/source/lalsuite/lalpulsar/src/OpenCLutils.c:612): Internal function call failed
XLAL Error - MAIN (/home/jenkins/workspace/workspace/EaH-GW-OpenCL-Testing/SLAVE/LIBC215/TARGET/linux-x86_64/EinsteinAtHome/source/lalsuite/lalapps/src/pulsar/GCT/HierarchSearchGCT.c:1276): Check failed: XLALGetOpenCLProgramFromSource ( source, &HierarchSearchGCTProgramm ) == XLAL_SUCCESS
XLAL Error - MAIN (/home/jenkins/workspace/workspace/EaH-GW-OpenCL-Testing/SLAVE/LIBC215/TARGET/linux-x86_64/EinsteinAtHome/source/lalsuite/lalapps/src/pulsar/GCT/HierarchSearchGCT.c:1276): Internal function call failed
2020-02-16 15:17:00.1848 (2750439) [CRITICAL]: ERROR: MAIN() returned with error '-1'
2020-02-16 15:17:00.1848 (2750439) [debug]: resultfile '../../projects/einstein.phys.uwm.edu/h1_1044.40_O2C02Cl4In0__O2MDFV2e_VelaJr1_1045.00Hz_161_2_0' (len 95), current config file: 0
Code-version: %% LAL: 6.19.2.1 (CLEAN 98bbe72a728eb25935e9195dafae691335dabf8c)
%% LALPulsar: 1.17.1.1 (CLEAN 98bbe72a728eb25935e9195dafae691335dabf8c)
%% LALApps: 6.23.0.1 (CLEAN 98bbe72a728eb25935e9195dafae691335dabf8c)

Cheers,
Gary.

Gary Roberts
Gary Roberts
Moderator
Joined: 9 Feb 05
Posts: 5,035
Credit: 34,660,383,349
RAC: 34,789,964

Sagittarius Lupus wrote:I've

Sagittarius Lupus wrote:
I've gone ahead and researched the problem myself. Standard disclaimer: I am not a C programmer, and I do not program in OpenCL, but I am literate, and I can break apart this program just enough to see what's going on.

I had not seen this further response until after I posted my response to your earlier post.

My first response still stands - there's absolutely no reason to be downright rude towards someone just trying to help.

I do understand (in a basic fashion) your explanation of what is causing the app to fail for your setup.  I would guess that the app has been coded with OpenCL 1.2 behaviour in mind and is not necessarily 2.0 compliant.  If your analysis is correct, then perhaps this problem can be resolved fairly easily without causing any regressions for the vast bulk of people still using OpenCL 1.2 implementations.

I suggest you send a PM to Bernd Maschenschalk who will know exactly who is responsible for the code.  Either send him a copy of your analysis or point him towards your analysis message.  If it's easy to solve, I'm sure it will be attended to.

Cheers,
Gary.

San-Fernando-Valley
San-Fernando-Valley
Joined: 16 Mar 16
Posts: 63
Credit: 1,703,282,648
RAC: 5,601,589

Sagittarius Lupus

Sagittarius Lupus wrote:

...  You do not appear to be affiliated with this project; kindly refrain from telling others what they "must" do.

 

-1

San-Fernando-Valley
San-Fernando-Valley
Joined: 16 Mar 16
Posts: 63
Credit: 1,703,282,648
RAC: 5,601,589

Sagittarius Lupus wrote: ...

Sagittarius Lupus wrote:

... Standard disclaimer: I am not a C programmer, and I do not program in OpenCL, but I am literate, and I can break apart this program just enough to see what's going on.   ...

That is nice for you ...

San-Fernando-Valley
San-Fernando-Valley
Joined: 16 Mar 16
Posts: 63
Credit: 1,703,282,648
RAC: 5,601,589

Gary Roberts wrote: ... Was

Gary Roberts wrote:

... Was it really necessary to do that??  You post about a problem you are having that is clearly well outside the skill set of the bulk of ordinary volunteers, you beg for anyone to respond, and when someone actually tries to be helpful, you reward that with quite some invective.  If you're so confident that you know exactly where the problem lies, then I suggest you desist from deliberately insulting a respondent just trying to help and wait patiently for someone who is qualified to study the problem in depth and is, perhaps, able to do something about it. ...

+1

San-Fernando-Valley
San-Fernando-Valley
Joined: 16 Mar 16
Posts: 63
Credit: 1,703,282,648
RAC: 5,601,589

Gary Roberts wrote: My first

Gary Roberts wrote:

My first response still stands - there's absolutely no reason to be downright rude towards someone just trying to help.

+1

Comment viewing options

Select your preferred way to display the comments and click "Save settings" to activate your changes.