search

cpc / openasip

Open Application-Specific Instruction Set processor tools (OpenASIP)
http://openasip.org
Other
140 stars 42 forks source link

Assertion `isRegLoc()' when using long in OpenCL kernels #179

Open nrother opened 1 year ago

nrother commented 1 year ago

Compiling the following OpenCL kernel code for a TTA target fails with an assertion:

__kernel void test(long a, __global int *b)
{
    long mul = (a * (long)10);
    int x_in = (int)(mul / (1 << 30));
    *b = x_in;
}
llvm-tce: /path/llvm14/include/llvm/CodeGen/CallingConvLower.h:150: llvm::Register llvm::CCValAssign::getLocReg() const: Assertion `isRegLoc()' failed.

This seems to be related to the usage of long.

pjaaskel commented 1 year ago

Likely due to it since in OpenCL C long is 64b always while in C we can (and do) define it to 32b for 32b TTAs. There's experimental 64b support in TCE which you could try out, but prepare for issues.