CUGIP_GLOBAL void
pushKernelMultiLevelGlobalSync(
TGraph aGraph,
ParallelQueueView<int> aVertices,
int *aLevelStarts,
int aLevelCount,
int aCurrentLevel,
device_ptr<int> aLastProcessedLevel,
device_flag_view aPushSuccessfulFlag,
cub::GridBarrier barrier)
{
uint blockId = __mul24(blockIdx.y, gridDim.x) + blockIdx.x;
int levelStart = aLevelStarts[aCurrentLevel - 1];
int levelEnd = aLevelStarts[aCurrentLevel];
do {
__syncthreads();
barrier.Sync();
int index = levelStart + blockId * blockDim.x + threadIdx.x;
while (index < levelEnd) {
pushImplementation(aGraph, aVertices, index, aCurrentLevel, aPushSuccessfulFlag);
index += blockDim.x * gridDim.x;
}
__syncthreads();
barrier.Sync();
--aCurrentLevel;
levelStart = aLevelStarts[aCurrentLevel - 1];
levelEnd = aLevelStarts[aCurrentLevel];
__syncthreads();
} while (aCurrentLevel > 0 && (levelEnd - levelStart) <= TPolicy::MULTI_LEVEL_GLOBAL_SYNC_LIMIT);
if (threadIdx.x == 0 && blockIdx.x == 0) {
aLastProcessedLevel.assign_device(aCurrentLevel);
}
}
CUGIP_GLOBAL void
pushKernel(
TGraph aGraph,
ParallelQueueView<int> aVertices,
int aLevelStart,
int aLevelEnd,
int aCurrentLevel,
device_flag_view aPushSuccessfulFlag)
{
uint blockId = __mul24(blockIdx.y, gridDim.x) + blockIdx.x;
int index = aLevelStart + blockId * blockDim.x + threadIdx.x;
while (index < aLevelEnd) {
pushImplementation(aGraph, aVertices, index, aCurrentLevel, aPushSuccessfulFlag);
index += blockDim.x * gridDim.x;
}
}
void Object::push(lua_State *L)
{
SLB_DEBUG_CALL;
SLB_DEBUG(3, "(L %p) Object::push(%p) [%s]", L, this, typeid(*this).name());
int top = lua_gettop(L);
// get the table of objects (top+1)
lua_getfield(L, LUA_REGISTRYINDEX, objectsTable_name);
if (lua_isnil(L, -1))
{
lua_pop(L, 1); // remove nil
initialize(L);
lua_getfield(L, LUA_REGISTRYINDEX, objectsTable_name);
}
// get Reference's table (top+2)
lua_getfield(L, LUA_REGISTRYINDEX, refTable_name);
// find the object with our id
lua_pushlightuserdata(L, (void*) this );
lua_rawget(L, top + 1);
if (lua_isnil(L,-1))
{
lua_pop(L, 1); // remove nil
int objpos = lua_gettop(L) + 1; // position of the user pushImplementation
pushImplementation(L); // value (0)
SLB_DEBUG(5, "\t-new object");
// check, pushImplementation should return only one value
if (lua_gettop(L) != objpos)
{
luaL_error(L, "Error on Object::push the current stack "
"has %d elments and should have only one.",
lua_gettop(L) - objpos - 1);
lua_settop(L, top);
return;
}
// make objects[key] = value
lua_pushlightuserdata(L, (void*) this); // key
lua_pushvalue(L, objpos); // copy of value (0)
lua_rawset(L, top+1); // objects[key] = value
lua_pushvalue(L, objpos); // key
// create a reference of this object...
Object **objptr = reinterpret_cast<Object**>(
lua_newuserdata(L, sizeof(Object*)));
*objptr = this; // copy a pointer
ref(); // manual reference
// get or crate a metatable for the reference of this
if(luaL_newmetatable(L, "SLB_ObjectPtr_GC"))
{
// create the __gc function
lua_pushcfunction(L, GC_callback);
lua_setfield(L, -2, "__gc"); // set the gc func of the obj
}
lua_setmetatable(L, -2); // metatable of reference object
// make references[obj] = reference
lua_rawset(L, top+2);
// check object at top
assert( lua_gettop(L) == objpos);
}
else
{
SLB_DEBUG(6, "\t-object exists");
}
// replace top+1 with the object
lua_replace(L, top+1);
lua_settop(L, top+1);
}
