Crush revealing on AMD gpu

I have a function:

void Core::SingleLine3(handler& H, accessor<Particle, 1> Particles, accessor<Particle, 1> eParticles, int Z,
                                                                           accessor<double, 1> idMap,
                                                                           accessor<double, 1> occupancyPdfMap,
                                                                          accessor<double, 1> bindingMap,
                                                                           accessor<double, 2> components,
                                                                           accessor<double, 1> biggs){
    double ScatFuncFitParam[101][16] = {...
    };
    sycl::stream out(1064,64,H);
    
    H.parallel_for(Particles.size(), [=](id<1> idx){
        int index = idx.get(0);
        double pi = 3.14159265458979323846;
        double twoPi = 2*pi;
        double photonEn = 0.04;
        double elMas_c2 = 0.510999;
        double e0m = photonEn/elMas_c2;
        double epsilon0Local = 1./(1. + 2.*e0m);
        double epsilon0Sq = epsilon0Local*epsilon0Local;
        double alpha1 = -sycl::log(epsilon0Local);
        double alpha2 = 0.5*(1. - epsilon0Sq);
        double wlPhoton = 3.0996e-08;
        double epsilon;
        double epsilonSq;
        double oneCosT;
        double sinT2;
        double gReject;
        oneapi::dpl::uniform_real_distribution<double> distr;
        oneapi::dpl::minstd_rand engine3(13,index);
        do{
            if(alpha1/(alpha1 + alpha2) > distr(engine3)){
                epsilon = sycl::exp(-alpha1*distr(engine3));
                epsilonSq = epsilon*epsilon;
            }
            else{
                epsilonSq = epsilon0Sq + (1. - epsilon0Sq)*distr(engine3);
                epsilon = sycl::sqrt(epsilonSq);
            }
            oneCosT = (1. - epsilon)/(epsilon*e0m);
            sinT2 = oneCosT*(2. - oneCosT);
            double x = sycl::sqrt(oneCosT/2.)*10/wlPhoton;
            double value = Z;
            if(x <= ScatFuncFitParam[Z][3]){
                double lgq = sycl::log(x)/sycl::log(10.);
                if(lgq < ScatFuncFitParam[Z][1]) {
                      value = ScatFuncFitParam[Z][4] + lgq*ScatFuncFitParam[Z][5];
                    }
                    else if(lgq >= ScatFuncFitParam[Z][1] && lgq < ScatFuncFitParam[Z][2]) {
                      value = ScatFuncFitParam[Z][6] + lgq*(ScatFuncFitParam[Z][7] + lgq*(ScatFuncFitParam[Z][8]
                                            + lgq*(ScatFuncFitParam[Z][9] + lgq*ScatFuncFitParam[Z][10])));
                    }
                    else {
                      value = ScatFuncFitParam[Z][11] + lgq*(ScatFuncFitParam[Z][12] + lgq*(ScatFuncFitParam[Z][13]
                                            + lgq*(ScatFuncFitParam[Z][14] + lgq*ScatFuncFitParam[Z][15])));
                    }
                value = sycl::exp(value*sycl::log(10.));
            }
            double scatteringFunction = value;
            gReject = (1. - epsilon*sinT2/(1. + epsilonSq))*scatteringFunction;
        }while(gReject < distr(engine3)*Z);

        double cosTheta = 1. - oneCosT;
        double sinTheta = sycl::sqrt(sinT2);
        double phi = twoPi*distr(engine3);
        double dirx = sinTheta*std::cos(phi);
        double diry = sinTheta*std::sin(phi);
        double dirz = cosTheta;

        // Doppler broadening -  Method based on:
        // Y. Namito, S. Ban and H. Hirayama,
        // "Implementation of the Doppler Broadening of a Compton-Scattered Photon
        // into the EGS4 Code", NIM A 349, pp. 489-494, 1994

        // Maximum number of sampling iterations
        int maxDopplerIterations = 1000;
        double bindingE = 0.;
        double photonEoriginal = epsilon*photonEn;
        double photonE = -1.;
        double iteration = 0;
        double eMax = photonEn;
        double param = photonEn;
        do{
            ++iteration;
            //SelectRandomShell()
            // Select shell based on shell occupancy
            int shellIndex = 0;
            double ran = distr(engine3);
            // Binary search the shell with probability less or equal random
            int nShellsLocal = idMap.size();
            int upperBound = nShellsLocal;

            while(shellIndex <= upperBound){
                int midShell = (shellIndex + upperBound)/2;
                if(ran < occupancyPdfMap[midShell]) upperBound = midShell - 1;
                else shellIndex = midShell + 1;
            }

            if(shellIndex >= nShellsLocal) shellIndex = nShellsLocal - 1;
            //
            //BindingEnergy()
            bindingE = bindingMap[shellIndex];
            eMax = photonEn - bindingE;
        
            // Randomly sample bound electron momentum
            // (memento: the data set is in Atomic Units)

            int si = -3;
            int i = 0;
            do{
                if(components[shellIndex][i] == 9 || i == (sizeof(components[shellIndex])/sizeof(components[shellIndex][0]) - 1)) si = i;
                else ++i;
            }while(si == -3);
            //if(si!= 8)            out << si << sycl::endl;
            double value = 0.;
            double x = distr(engine3);
            std::size_t lowerBound = 0;
            //upperBound = components[shellIndex].size() - 1;
            upperBound = si - 1;
            while (lowerBound <= upperBound){
                std::size_t midBin((lowerBound + upperBound) / 2);

                if(x < (components[shellIndex])[midBin]) upperBound = midBin - 1;
                else lowerBound = midBin + 1;
            }
            int bin = upperBound;
            double y;
            if(bin == 0){
                //int nBins = int(components[shellIndex].size() - 1);
                int nBins = si - 1;
                double value = 0.;
                if(x < (components[shellIndex])[0]){
                    value = 0.;
                }
                else if(bin < nBins){
                    double e1 = (components[shellIndex])[bin];
                    double e2 = (components[shellIndex])[bin+1];
                    double d1 = biggs[bin];
                    double d2 = biggs[bin+1];
                    value = d1 + (d2 - d1)*(x - e1)/(e2 - e1);
                }
                else{
                     value = biggs[nBins];
                }
                y = value;
            }
            else{
                int nBins = int(biggs.size() - 1);
                double value = 0.;
                if(x < (components[shellIndex])[0]){
                    value = 0.;
                }
                else if(bin < nBins){
                    double e1 = (components[shellIndex])[bin];
                    double e2 = (components[shellIndex])[bin+1];
                    double d1 = biggs[bin];
                    double d2 = biggs[bin+1];
                    // Check of e1, e2, d1 and d2 values to avoid doubleing-point errors when estimating the
                    // interpolated value below
                    if((d1 > 0.) && (d2 > 0.) && (e1 > 0.) && (e2 > 0.)){
                        value = std::log10(d1)+(std::log10(d2/d1)/std::log10(e2/e1)*std::log10(x/e1));
                        value = std::pow(10.,value);
                    }
                    else value = 0.;
                }
                else value = biggs[nBins];
                y = value;
            }

            double pSample = y;
            // Rescale from atomic units
            double fine_structure_const = 0.00729735;
            double pDoppler = pSample*fine_structure_const;
            double pDoppler2 = pDoppler*pDoppler;
            double var2 = 1. + oneCosT*e0m;
            double var3 = var2*var2 - pDoppler2;
            double var4 = var2 - pDoppler2*cosTheta;
            double var = var4*var4 - var3 + pDoppler2*var3;
            if(var > 0.){
                double varSqrt = sycl::sqrt(var);
                double scale = photonEn/var3;
                    // Random select either root
                if(distr(engine3) < 0.5){ photonE = (var4 - varSqrt)*scale; }
                else {                    photonE = (var4 + varSqrt)*scale; }
            }
            else{
                photonE = -1.;
            }
         }while(iteration <= maxDopplerIterations && photonE > eMax);
         
        // End of recalculation of photon energy with Doppler broadening
        // Revert to original if maximum number of iterations threshold
        // has been reached
        if(iteration >= maxDopplerIterations){
            photonE = photonEoriginal;
            bindingE = 0.;
        }

        if(bindingE > 1.5){  out << bindingE << sycl::endl;}
        // Update G4VParticleChange for the scattered photon
        double photonEnergy1 = photonE;
        Particles[index].position.Set(photonEn, photonEnergy1, bindingE);

        if(photonEnergy1 > 0.){
            // Kinematics of the scattered electron
            double eKineticEnergy = photonEn - photonEnergy1 - bindingE;
            // protection against negative final energy: no e- is created

            if(eKineticEnergy > 0.0){
                double eTotalEnergy = eKineticEnergy + elMas_c2;
                double electronE = photonEn*(1. - epsilon) + elMas_c2;
                double electronP2 = electronE*electronE - elMas_c2*elMas_c2;
                double sinThetaE = -1.;
                double cosThetaE = 0.;
                if(electronP2 > 0){
                    cosThetaE = (eTotalEnergy + photonEnergy1 )*(1. - epsilon)/sycl::sqrt(electronP2);
                    sinThetaE = -1.*sycl::sqrt(1. - cosThetaE*cosThetaE);
                }
                double x = sinThetaE*sycl::cos(phi);
                double y = sinThetaE*sycl::sin(phi);
                double z = cosThetaE;
                double ox = Particles[index].momentum.normPx();
                double oy = Particles[index].momentum.normPy();
                double oz = Particles[index].momentum.normPz();
                double up = ox*ox + oy*oy;
                double pMod = sycl::sqrt((eKineticEnergy + elMas_c2)*(eKineticEnergy + elMas_c2) - elMas_c2*elMas_c2);
                if(up > 0){
                    up = sycl::sqrt(up);
                    double px = (ox*oz*x - oy*y)/up + ox*z;
                    double py = (oy*oz*x + ox*y)/up + oy*z;
                    double pz = -up*x + oz*z; 
                    px = pMod*px;
                    py = pMod*py;
                    pz = pMod*pz;
                    eParticles[index].momentum.Set(px,py,pz);
                }
                else if(up < 0.){ 
                    x = pMod*x;
                    y = pMod*y;
                    z = pMod*z;
                    eParticles[index].momentum.Set(-x, y, -z);
                }
                else{
                    x = pMod*x;
                    y = pMod*y;
                    z = pMod*z;
                    eParticles[index].momentum.Set(x, y, z);
                }
            }
            //momentum residual photon
                
            double x = dirx;
            double y = diry;
            double z = dirz;
            double ox = Particles[index].momentum.normPx();
            double oy = Particles[index].momentum.normPy();
            double oz = Particles[index].momentum.normPz();
            double up = ox*ox + oy*oy;
            double pMod = photonEnergy1;

            if(up > 0){
                up = sycl::sqrt(up);
                double px = (ox*oz*x - oy*y)/up + ox*z;
                double py = (oy*oz*x + ox*y)/up + oy*z;
                double pz = -up*x + oz*z; 
                px = pMod*px;
                py = pMod*py;
                pz = pMod*pz;
                Particles[index].momentum.Set(px,py,pz);
            }
            else if(up < 0){ 
                x = pMod*x;
                y = pMod*y;
                z = pMod*z;
                Particles[index].momentum.Set(-x, y, -z);
            }
            else{
                x = pMod*x;
                y = pMod*y;
                z = pMod*z;
                Particles[index].momentum.Set(x, y, z);
            }
        }
        else{
            // photon absorbed
            photonEnergy1 = 0.;
            Particles[index].momentum.Set(0,0,0);
            eParticles[index].momentum.Set(0,0,0);
        }
    });
}

and Error

lld: error: no libcall available for flog10
lld: error: no libcall available for flog10
lld: error: no libcall available for flog10
lld: error: no libcall available for flog10
LLVM ERROR: Cannot select: 0x55eb2117b600: f64 = fpow nsz arcp contract afn reassoc ConstantFP:f64<1.000000e+01>, 0x55eb20e8f020, tpt4/core/src/core.cpp:1240:33 @[ intel/oneapi/compiler/2025.0/bin/compiler/../../include/sycl/handler.hpp:405:7 @[ intel/oneapi/compiler/2025.0/bin/compiler/../../include/sycl/handler.hpp:1639:5 ] ]
 intel/oneapi/compiler/2025.0/bin/compiler/../../include/sycl/handler.hpp:405:7 @[ intel/oneapi/compiler/2025.0/bin/compiler/../../include/sycl/handler.hpp:1639:5 ] ]
        0x55eb216c4df0: i32 = undef
        0x55eb216c4df0: i32 = undef
In function: _ZTSN4sycl3_V16detail18RoundedRangeKernelINS0_4itemILi1ELb1EEELi1EZN4Core11SingleLine3ERNS0_7handlerENS0_8accessorI8ParticleLi1ELNS0_6access4modeE1026ELNSA_6targetE2014ELNSA_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEEESI_iNS8_IdLi1ELSB_1026ELSC_2014ELSD_0ESH_EESJ_SJ_NS8_IdLi2ELSB_1026ELSC_2014ELSD_0ESH_EESJ_EUlNS0_2idILi1EEEE_EE
PLEASE submit a bug report to https://software.intel.com/en-us/support/priority-support and include the crash backtrace.
Stack dump:
0.	Program arguments: /home/aleksei/intel/oneapi/compiler/2025.0/bin/compiler/lld -flavor gnu -m elf64_amdgpu --no-undefined -shared -plugin-opt=-amdgpu-internalize-symbols -plugin-opt=-vector-library=SVML -plugin-opt=mcpu=gfx1102 -plugin-opt=fintel-libirc-allowed -plugin-opt=-disable-hir-generate-mkl-call -plugin-opt=-intel-abi-compatible=true -plugin-opt=-x86-enable-unaligned-vector-move=true -plugin-opt=-alt-math-library=svml --whole-archive -o /tmp/main-gfx1102-c82396-86ad7a.out /tmp/main-gfx1102-acd46c-f72cc5.o --no-whole-archive
1.	Running pass 'CallGraph Pass Manager' on module 'ld-temp.o'.
2.	Running pass 'AMDGPU DAG->DAG Pattern Instruction Selection' on function ''
 #0 0x000055eb1c9caba7 llvm::sys::PrintStackTrace(llvm::raw_ostream&, int) (/home/aleksei/intel/oneapi/compiler/2025.0/bin/compiler/lld+0x3902ba7)
 #1 0x000055eb1c9c9904 llvm::sys::RunSignalHandlers() (/home/aleksei/intel/oneapi/compiler/2025.0/bin/compiler/lld+0x3901904)
 #2 0x000055eb1c9cb3cf SignalHandler(int) Signals.cpp:0:0
 #3 0x00007fe086242520 (/lib/x86_64-linux-gnu/libc.so.6+0x42520)
 #4 0x00007fe0862969fc __pthread_kill_implementation ./nptl/pthread_kill.c:44:76
 #5 0x00007fe0862969fc __pthread_kill_internal ./nptl/pthread_kill.c:78:10
 #6 0x00007fe0862969fc pthread_kill ./nptl/pthread_kill.c:89:10
 #7 0x00007fe086242476 gsignal ./signal/../sysdeps/posix/raise.c:27:6
 #8 0x00007fe0862287f3 abort ./stdlib/abort.c:81:7
 #9 0x000055eb1c1a83c7 llvm::report_fatal_error(llvm::Twine const&, bool) (/home/aleksei/intel/oneapi/compiler/2025.0/bin/compiler/lld+0x30e03c7)
#10 0x000055eb1be31dbf llvm::SelectionDAGISel::CannotYetSelect(llvm::SDNode*) (/home/aleksei/intel/oneapi/compiler/2025.0/bin/compiler/lld+0x2d69dbf)
#11 0x000055eb1be3c5b0 llvm::SelectionDAGISel::SelectCodeCommon(llvm::SDNode*, unsigned char const*, unsigned int) (/home/aleksei/intel/oneapi/compiler/2025.0/bin/compiler/lld+0x2d745b0)
#12 0x000055eb1cf13a6a AMDGPUDAGToDAGISel::Select(llvm::SDNode*) AMDGPUISelDAGToDAG.cpp:0:0
#13 0x000055eb1be515fa llvm::SelectionDAGISel::DoInstructionSelection() (/home/aleksei/intel/oneapi/compiler/2025.0/bin/compiler/lld+0x2d895fa)
#14 0x000055eb1bbd2d62 llvm::SelectionDAGISel::CodeGenAndEmitDAG() (/home/aleksei/intel/oneapi/compiler/2025.0/bin/compiler/lld+0x2b0ad62)
#15 0x000055eb1bc341bc llvm::SelectionDAGISel::SelectAllBasicBlocks(llvm::Function const&) (/home/aleksei/intel/oneapi/compiler/2025.0/bin/compiler/lld+0x2b6c1bc)
#16 0x000055eb1c32025f llvm::SelectionDAGISel::runOnMachineFunction(llvm::MachineFunction&) (/home/aleksei/intel/oneapi/compiler/2025.0/bin/compiler/lld+0x325825f)
#17 0x000055eb1cf1240b AMDGPUDAGToDAGISel::runOnMachineFunction(llvm::MachineFunction&) AMDGPUISelDAGToDAG.cpp:0:0
#18 0x000055eb1c31e166 llvm::SelectionDAGISelLegacy::runOnMachineFunction(llvm::MachineFunction&) (/home/aleksei/intel/oneapi/compiler/2025.0/bin/compiler/lld+0x3256166)
#19 0x000055eb1bd5c44b llvm::MachineFunctionPass::runOnFunction(llvm::Function&) (/home/aleksei/intel/oneapi/compiler/2025.0/bin/compiler/lld+0x2c9444b)
#20 0x000055eb1bd3d485 llvm::FPPassManager::runOnFunction(llvm::Function&) (/home/aleksei/intel/oneapi/compiler/2025.0/bin/compiler/lld+0x2c75485)
#21 0x000055eb1e1ca2b2 (anonymous namespace)::CGPassManager::runOnModule(llvm::Module&) CallGraphSCCPass.cpp:0:0
#22 0x000055eb1bd36e56 llvm::legacy::PassManagerImpl::run(llvm::Module&) (/home/aleksei/intel/oneapi/compiler/2025.0/bin/compiler/lld+0x2c6ee56)
#23 0x000055eb1d37ccee codegen(llvm::lto::Config const&, llvm::TargetMachine*, std::__1::function<llvm::Expected<std::__1::unique_ptr<llvm::CachedFileStream, std::__1::default_delete<llvm::CachedFileStream>>> (unsigned int, llvm::Twine const&)>, unsigned int, llvm::Module&, llvm::ModuleSummaryIndex const&) LTOBackend.cpp:0:0
#24 0x000055eb1d37c33c llvm::lto::backend(llvm::lto::Config const&, std::__1::function<llvm::Expected<std::__1::unique_ptr<llvm::CachedFileStream, std::__1::default_delete<llvm::CachedFileStream>>> (unsigned int, llvm::Twine const&)>, unsigned int, llvm::Module&, llvm::ModuleSummaryIndex&) (/home/aleksei/intel/oneapi/compiler/2025.0/bin/compiler/lld+0x42b433c)
#25 0x000055eb1d36a1f1 llvm::lto::LTO::runRegularLTO(std::__1::function<llvm::Expected<std::__1::unique_ptr<llvm::CachedFileStream, std::__1::default_delete<llvm::CachedFileStream>>> (unsigned int, llvm::Twine const&)>) (/home/aleksei/intel/oneapi/compiler/2025.0/bin/compiler/lld+0x42a21f1)
#26 0x000055eb1d369a7f llvm::lto::LTO::run(std::__1::function<llvm::Expected<std::__1::unique_ptr<llvm::CachedFileStream, std::__1::default_delete<llvm::CachedFileStream>>> (unsigned int, llvm::Twine const&)>, std::__1::function<llvm::Expected<std::__1::function<llvm::Expected<std::__1::unique_ptr<llvm::CachedFileStream, std::__1::default_delete<llvm::CachedFileStream>>> (unsigned int, llvm::Twine const&)>> (unsigned int, llvm::StringRef, llvm::Twine const&)>) (/home/aleksei/intel/oneapi/compiler/2025.0/bin/compiler/lld+0x42a1a7f)
#27 0x000055eb1cb5fded lld::elf::BitcodeCompiler::compile() (/home/aleksei/intel/oneapi/compiler/2025.0/bin/compiler/lld+0x3a97ded)
#28 0x000055eb1cab3cfc void lld::elf::LinkerDriver::link<llvm::object::ELFType<(llvm::endianness)1, true>>(llvm::opt::InputArgList&) (/home/aleksei/intel/oneapi/compiler/2025.0/bin/compiler/lld+0x39ebcfc)
#29 0x000055eb1caa0f22 lld::elf::LinkerDriver::linkerMain(llvm::ArrayRef<char const*>) (/home/aleksei/intel/oneapi/compiler/2025.0/bin/compiler/lld+0x39d8f22)
#30 0x000055eb1ca9f5e4 lld::elf::link(llvm::ArrayRef<char const*>, llvm::raw_ostream&, llvm::raw_ostream&, bool, bool) (/home/aleksei/intel/oneapi/compiler/2025.0/bin/compiler/lld+0x39d75e4)
#31 0x000055eb1c9ccb97 lld::unsafeLldMain(llvm::ArrayRef<char const*>, llvm::raw_ostream&, llvm::raw_ostream&, llvm::ArrayRef<lld::DriverDef>, bool) (/home/aleksei/intel/oneapi/compiler/2025.0/bin/compiler/lld+0x3904b97)
#32 0x000055eb1c6c8abf lld_main(int, char**, llvm::ToolContext const&) (/home/aleksei/intel/oneapi/compiler/2025.0/bin/compiler/lld+0x3600abf)
#33 0x000055eb1c6c8629 main (/home/aleksei/intel/oneapi/compiler/2025.0/bin/compiler/lld+0x3600629)
#34 0x00007fe086229d90 __libc_start_call_main ./csu/../sysdeps/nptl/libc_start_call_main.h:58:16
#35 0x00007fe086229e40 call_init ./csu/../csu/libc-start.c:128:20
#36 0x00007fe086229e40 __libc_start_main ./csu/../csu/libc-start.c:379:5
#37 0x000055eb1c982c3e _start /export/project/tools-build/toolchain-cross/src/glibc/csu/../sysdeps/x86_64/start.S:122:0
llvm-foreach: Aborted (core dumped)
icpx: error: amdgcn-link command failed with exit code 254 (use -v to see invocation)
Intel(R) oneAPI DPC++/C++ Compiler 2025.0.0 (2025.0.0.20241008)
Target: x86_64-unknown-linux-gnu
Thread model: posix
InstalledDir: /home/aleksei/intel/oneapi/compiler/2025.0/bin/compiler
Configuration file: /home/aleksei/intel/oneapi/compiler/2025.0/bin/compiler/../icpx.cfg
icpx: note: diagnostic msg: Error generating preprocessed source(s).
make[2]: *** [CMakeFiles/test.dir/build.make:111: test] Error 1
make[1]: *** [CMakeFiles/Makefile2:434: CMakeFiles/test.dir/all] Error 2
make: *** [Makefile:91: all] Error 2

while i do standard compilation with icpx -fsycl the code works well, but when i try to build with AMD support, in compilation flow errors occurs. Furthermore i cannot trace it while commenting: the problem in second do block as i understand because if i comment all his content the error is vanished. Then an idea that may be 2D matrix is somehow wrong commenting everythin below int si = -3;, didn’t resolve compilation. And at finest when i only comment if(components[shell Index][i] == 9 etc line error hides.

Found the problem here:

if((d1 > 0.) && (d2 > 0.) && (e1 > 0.) && (e2 > 0.)){
                        value = std::log10(d1)+(std::log10(d2/d1)/std::log10(e2/e1)*std::log10(x/e1));
                        value = std::pow(10.,value);
                    }

Forgot that on gpu std:: doesn’t works.

You could try using sycl::log10 and sycl::pow as they are functions in the SYCL spec. They will try to find an optimized function for those, and if it can’t I think it will fall back to the std:: version on CPU. See the SYCL reference for some information on what is defined in the spec. Math functions — SYCL Reference documentation