linnaea/beignet

Beignet is an open source implementation of the OpenCL specification - a generic compute oriented API. Here is Beignet Source Code Mirror in github- This is a publish-only repository and all pull requests are ignored. Please follow https://wiki.freedesktop.org/www/Software/Beignet/ for any of your improvements

beignet's Issues

IF/ENDIF fix does not work on large nested if/endif block

Current IF/ENDIF fix only handles the innermost block, this will still trigger the IF/ENDIF out of range error if the block is enclosed in another IF/ENDIF block.

beignet/backend/src/backend/gen_insn_selection.cpp

Lines 2696 to 2708 in 6553d2c

	if (this->ctx.getIFENDIFFix() &&
	!this->block->insnList.empty() &&
	this->block->insnList.size() % 1000 == 0 &&
	this->block->endifLabel.value() != 0) {
	this->curr.flag = 0;
	this->curr.subFlag = 1;
	ir::LabelIndex jip = this->block->endifLabel;
	this->ENDIF(GenRegister::immd(0), jip);
	this->push();
	this->curr.predicate = GEN_PREDICATE_NORMAL;
	this->IF(GenRegister::immd(0), jip, jip);
	this->pop();
	}

This might be the cause of some assertions on hashcat.

Removing the assertion on gen_program.cpp as a workaround.

ASSERTION FAILED: The bit sizes of src and the dst is not identical on Haswell GT2 Mobile

Tracking for upstream GitLab 81.

Issue seems to be that structs with 3 elements that looks like vec3 being converted to vec3, while CL spec says vec3 should take 4 elements' space rather than 3.

typedef struct __attribute__ ((packed)) {
    uint cols;
    uint diagl;
    uint diagr;
} start;

#define DEPTH 9
#define WG_SIZE 64

typedef short int_fast16_t;
typedef ushort uint_fast16_t;
typedef uint uint_fast32_t;
#define UINT_FAST32_MAX UINT_MAX

// for convenience
#define L (get_local_id(0))
#define G (get_global_id(0))

kernel void solve(__global const start* in_starts, __global uint* out_cnt) {
    uint num = 0;
    uint_fast32_t cols[WG_SIZE][DEPTH+1], posibs[WG_SIZE][DEPTH+1];
    uint_fast32_t diagl[WG_SIZE][DEPTH+1], diagr[WG_SIZE][DEPTH+1];
    int_fast16_t d = 1;
    cols[L][d] = in_starts[G].cols;
    diagl[L][d] = in_starts[G].diagl;
    diagr[L][d] = in_starts[G].diagr;
    uint_fast32_t posib = ( diagl[L][d] | diagr[L][d]);
    uint_fast32_t bit = ~posib & (posib + 1);
    uint_fast32_t new_diagr = (bit >> 1) | diagr[L][d];
    bit |= cols[L][d];
    uint_fast32_t new_posib = (bit | new_diagr);

    if (new_posib != UINT_FAST32_MAX) {
      cols[L][d] = bit;
    } else {
        num += bit == UINT_FAST32_MAX;
    }
    posib = posibs[L][d];
    out_cnt[G] += num;
}

Build fails with LLVM/Clang 8/9/10

The code builds fine with LLVM/Clang 7, but not with version 8/9/10:

FAILED: backend/src/libgbeinterp.so 
: && /usr/bin/clang++ -fPIC -D_FORTIFY_SOURCE=2 -march=x86-64 -mtune=generic -O2 -pipe -fno-plt  -funroll-loops -fstrict-aliasing -msse2 -msse3 -mssse3 -msse4.1 -fPIC -Wall -Wno-invalid-offsetof -I/usr/include -D_GNU_SOURCE -D__STDC_CONSTANT_MACROS -D__STDC_FORMAT_MACROS -D__STDC_LIMIT_MACROS -DGBE_DEBUG_MEMORY=0  -DGBE_COMPILER_AVAILABLE=1 -fvisibility=hidden -O2 -DNDEBUG -DGBE_DEBUG=0  -fuse-ld=lld -Wl,-O1,--sort-common,--as-needed,-z,relro,-z,now -Wl,-Bsymbolic -Wl,--no-undefined -L/usr/lib -shared -Wl,-soname,libgbeinterp.so -o backend/src/libgbeinterp.so backend/src/CMakeFiles/gbeinterp.dir/gbe_bin_interpreter.cpp.o   && :
ld.lld: error: undefined symbol: llvm::DisableABIBreakingChecks
>>> referenced by gbe_bin_interpreter.cpp
>>>               backend/src/CMakeFiles/gbeinterp.dir/gbe_bin_interpreter.cpp.o:(llvm::VerifyDisableABIBreakingChecks)
clang-10: error: linker command failed with exit code 1 (use -v to see invocation)

Full build log

Emulation needed for intrinsics

• fshl/fshr (LLVM 8)
• ssub_sat/usub_sat (LLVM 9)

Hint: here