__device__ int org_trifort_gc_malloc(int size){ int space_size = (int) m_Local[1]; int ret = org_trifort_gc_malloc_no_fail(size); int end = ret + ((size + 16) >> 4); if(end >= space_size){ return -1; } return ret; }
/** * Identical to org_trifort_gc_malloc_no_fail, but checks if the returned * address is actually valid! */ __device__ int org_trifort_gc_malloc( int size ) { int const & objectMemSizeDiv16 = (int) m_Local[1]; int const address = org_trifort_gc_malloc_no_fail( size ); /** * int const end = address + size/16 + 1; // this looks very wrong! * if ( end >= objectMemSizeDiv16 ) * 0 <= size <= 15 -> 1 * This is floor( size/16 ) + 1. It is stricter than this new version, so * it shouldn't have led to a segfault, it was just confusing and lax */ if ( address + padNumberTo( size, MallocAlignBytes ) / MallocAlignBytes > objectMemSizeDiv16 ) { assert( false && "Garbage collector ran out of memory!" ); // this needs to be checked by the caller! Else a segfault before the // allocated memory might occur ... */ return -1; } return address; }