abs undefined within device block

[mloubout]

Hi

I am one of the Devito developers and we ran into an issue when abs is used within a device block. Is it something with the code generated?
MFE:

#include <math.h>
#include <stdlib.h>
#include <omp.h>


int simplesum(int argc, char const * argv[]) {
    const int N = 1000;
    int array[N];
    for (int i=0; i<N; i++){
        array[i] = i;
    }
    #pragma omp target
    {
        #pragma omp parallel for
        for (int i=-N; i<N; i++){
            array[i] += array[(int)abs(i)];
        }
    }
    return 0;
}

Error:

$:aompcc -O3 -g -fPIC -Wall -Wno-unused-result -Wno-unused-variable -ffast-math -shared debug.c -lm -o test
lld: error: undefined symbol: abs
>>> referenced by debug.c:0 (/app/debug.c:0)
>>>               /tmp/debug-f98ea8-gfx906-02f80a.o:(__omp_offloading_40_17c02e4_simplesum_l12)
>>> referenced by debug.c:0 (/app/debug.c:0)
>>>               /tmp/debug-f98ea8-gfx906-02f80a.o:(__omp_offloading_40_17c02e4_simplesum_l12)
>>> referenced by debug.c:0 (/app/debug.c:0)
>>>               /tmp/debug-f98ea8-gfx906-02f80a.o:(__omp_offloading_40_17c02e4_simplesum_l12)
>>> referenced 1 more times
clang-14: error: amdgcn-link command failed with exit code 1 (use -v to see invocation)
ERROR:  The following command failed with return code 1.
        /opt/rocm-5.1.3/llvm/bin/clang  -g  -O3 -target x86_64-pc-linux-gnu -fopenmp -fopenmp-targets=amdgcn-amd-amdhsa -Xopenmp-target=amdgcn-amd-amdhsa -march=gfx906  -fPIC -Wall -Wno-unused-result -Wno-unused-variable -ffast-math -shared -lm  debug.c -o /app/test

[ronlieb]

test case aomp/test/smoke-fails/issue-435

[gregrodgers]

This needs to be resolved in 16.0-2

[doru1004]

This example compiles fine with latest clang++ but not with clang. Will investigate.

[doru1004]

The fix is waiting to be merged, the updated test is here: #438

[gregrodgers]

First, your example has a simple serious bug which will cause a memory fault once we fix the problem with missing integer abs. Since array does not have negative indices, Change this line.

array[i] += array[(int)abs(i)];

to

array[abs(i)] += array[(int)abs(i)];

Ok, now regarding this missing abs() in c.

Sorry, that this is taking so long to fix. It turns out to be a problem with our header library which needs to work for c and c++. I don't know why they only defined abs, labs, and llabs only for cpluplus in the header. I also don't understand why the c++ header definitions are so complicated. This patch will fix your problem and will be in 16.0-3. Just apply this to the installation of AOMP or ROCm compiler

--- /work/grodgers/git/aomp16.0/llvm-project/clang/lib/Headers/__clang_hip_math.h       2022-11-14 09:16:15.790306136 -0600
+++ lib/clang/16.0.0/include/__clang_hip_math.h 2022-11-14 08:54:21.462471385 -0600
@@ -182,6 +182,13 @@
   long long __sgn = __x >> (sizeof(long long) * CHAR_BIT - 1);
   return (__x ^ __sgn) - __sgn;
 }
+#else
+__DEVICE__
+int abs(int __x)               { return (__x < 0) ? -__x : __x; }
+__DEVICE__
+long labs(long __x)            { return (__x < 0) ? -__x : __x; }
+__DEVICE__
+long long llabs(long long __x) { return (__x < 0) ? -__x : __x; }
 #endif

 __DEVICE__
 

Also, your example has a simple serious bug. You need to add abs(i)

[gregrodgers]

@doru1004 I will fix our smoke-fails test and submit gerrit review for 16.0-3

[mloubout]

your example has a simple serious bug which will cause a memory fault

My apologies for that, I merely made that example to trigger the compilation error and did not check whether it was valid to run it or not. Thank you for fixing it.

Will give it a try with 16.0-3, thanks.

[doru1004]

@gregrodgers @mloubout the fix for the test itself is here: #438