CLStatus CLThreadForMsgLoop::Run(void *pContext)
{
if(m_pCoordinator == 0)
return CLStatus(-1, 0);
CLEvent event;
CLThreadInitialFinishedNotifier notifier(&event);
SLExecutiveInitialParameter para;//打包通知变量和运行参数
para.pContext = pContext;
para.pNotifier = ¬ifier;
CLStatus s = m_pCoordinator->Run(¶);
if(!s.IsSuccess())
{
CLLogger::WriteLogMesg("In CLThreadForMsgLoop::Run(), m_pThread->Run error", 0);
m_bWaitForDeath = false;
m_pCoordinator = 0;
return CLStatus(-1, 0);
}
CLStatus s1 = event.Wait();
if(!s1.IsSuccess())
CLLogger::WriteLogMesg("In CLThreadForMsgLoop::Run(), event.Wait error", 0);
if(notifier.IsInitialSuccess())
return CLStatus(0,0);
else
return CLStatus(-1,0);
}
void CLCommand::SetWaitList(cl_uint num_events_in_wait_list,
const cl_event* event_wait_list) {
wait_events_.reserve(num_events_in_wait_list);
for (cl_uint i = 0; i < num_events_in_wait_list; i++) {
CLEvent* e = event_wait_list[i]->c_obj;
e->Retain();
wait_events_.push_back(e);
}
}
void CTLord::OnTimer()
{ // 틱 // 초당 1회 // m_pController가 호출
TRACE( "CTLord.OnTimer: %d\n", time( NULL ) );
CTElection* pElection = static_cast<CTElection*>( GetElection() );
pElection->OnTimer();
CLEvent* pEvent = static_cast<CLEvent*>( GetEvent() );
pEvent->OnTimer();
CTLordSkill* pSkills = static_cast<CTLordSkill*>( GetSkills() );
pSkills->OnTimer();
}
static void* TestThreadForCLEvent2(void *arg)
{
CLEvent *p = (CLEvent *)arg;
for(int i = 0; i < count; i++)
{
EXPECT_TRUE((p->Set()).IsSuccess());
}
return 0;
}
void CCLord::CreateColleagues( void )
{
// 생성
CCElection * pElection = new CCElection( this );
CLEvent* pEvent = new CLEvent( this );
CLordSkill* pSkills = new CLordSkill( this );
// 스크립트 로드
pEvent->Initialize( "lordevent.inc" );
pSkills->Initialize( "lordskill.txt" );
// 설정
m_pElection = pElection;
m_pEvent = pEvent;
m_pSkills = pSkills;
}
/**Generates hit data for viewing rays, from the constructed BVH
* @param err Error info
* @return Result of the operation: Success or failure
*/
Result BVHManager::generateContacts(Camera& cam,Errata& err)
{
//Uploading camera state to GPU
_deviceCamera.reset(new CLBuffer(_context,sizeof(struct Camera),&cam,CLBufferFlags::ReadOnly));
//Allocating memory for primary contacts
size_t contactBufSize = cam.resX * cam.resY * sizeof(struct Contact);
if (_primaryContactsArray)
_primaryContactsArray->resize(contactBufSize);
else
_primaryContactsArray.reset(new CLBuffer(_context,contactBufSize,CLBufferFlags::ReadWrite));
//Getting the launch parameters
size_t processors, warp;
if (Success != _context.getMaximalLaunchExecParams(*_contactGenerateKernel,processors,warp,err))
return Error;
//Setting kernel args
try
{
SET_KERNEL_ARGS((*_contactGenerateKernel),_deviceCamera->getCLMem(),_bvhNodes->getCLMem(),_bvhLeavesCount,_scene.getDeviceSceneData(),_primaryContactsArray->getCLMem());
}
catch (CLInterfaceException e)
{
err = Errata(e);
return Error;
}
CLEvent evt;
evt.reset();
cl_uint totalWorkItems = closestMultipleTo(cam.resX * cam.resY,warp);
CLKernelWorkDimension globalDim(1,totalWorkItems);
CLKernelWorkDimension localDim(1,warp);
CLKernelExecuteParams contactsKernelExecParams(&globalDim,&localDim,&evt);
//Execute kernel
if (Success != _context.enqueueKernel((*_contactGenerateKernel),contactsKernelExecParams,err))
return Error;
//Flush and make sure that calculation is over before returning control
if (Success != _context.flushQueue(err))
return Error;
if (Success != evt.wait(err))
return Error;
return Success;
}
int main_ss()
{
CLEvent *pEvent = (CLEvent *) new CLEvent();
CLCoordinator *pCoordinator = new CLRegularCoordinator();
//CLExecutiveFunctionProvider * printer = new CLParaPrinter();
CLMyFunction *myfunc = new CLMyFunction();
CLExecutive *pThread = new CLThread(pCoordinator,true);
pCoordinator->SetExecObjects(pThread,myfunc);
cout <<"Status Code: "<<pCoordinator->Run((void*)pEvent).m_clReturnCode<<endl;
pEvent->Wait();
cout << "main thread: i received signal." << endl;
pCoordinator->WaitForDeath();
//pCoordinator->WaitForDeath();
return 0;
}
BOOL CLController::Restore( void )
{ // 모든 군주 시스템 협력 객체들의 복원
CQuery* pQuery = CreateQuery();
if( !pQuery )
return FALSE;
CTElection* pElection = static_cast<CTElection*>( m_pLord->GetElection() );
CLEvent* pEvent = static_cast<CLEvent*>( m_pLord->GetEvent() );
CTLordSkill* pSkills = static_cast<CTLordSkill*>( m_pLord->GetSkills() );
if( m_pLord->Restore( pQuery ) && pElection->Restore( pQuery )
&& pEvent->Restore( pQuery ) && pSkills->Restore( pQuery ) )
{
SAFE_DELETE( pQuery );
return TRUE;
}
SAFE_DELETE( pQuery );
return FALSE;
}
/**Generates contacts for rays and fills the contacts array
* @param rays The rays to be traced - A device memory buffer object that contains rays as struct Ray
* @param contact The target device memory that will contain the result - For each ray rays[i] the buffer will contain
* data about its closest intersection with object in the scene as contacts[i]. The contact data will be stored
* as struct Contact
* @param rayCount The number of rays to trace
* @param err Error info
* @return Result of the operation: Success or failure
*/
Result BVHManager::generateContacts(OpenCLUtils::CLBuffer& rays,OpenCLUtils::CLBuffer& contacts, const unsigned int rayCount, Common::Errata& err)
{
//Setting kernel args
try
{
SET_KERNEL_ARGS((*_contactGenerateKernel2),rays.getCLMem(),rayCount,_bvhNodes->getCLMem(),_bvhLeavesCount,_scene.getDeviceSceneData(),contacts.getCLMem());
}
catch (CLInterfaceException e)
{
err = Errata(e);
return Error;
}
//Getting the launch parameters
size_t processors, warp;
if (Success != _context.getMaximalLaunchExecParams(*_contactGenerateKernel2,processors,warp,err))
return Error;
CLEvent evt;
evt.reset();
cl_uint totalWorkItems = closestMultipleTo(rayCount,warp);
CLKernelWorkDimension globalDim(1,totalWorkItems);
CLKernelWorkDimension localDim(1,warp);
CLKernelExecuteParams contactsKernelExecParams(&globalDim,&localDim,&evt);
//Execute kernel
if (Success != _context.enqueueKernel((*_contactGenerateKernel2),contactsKernelExecParams,err))
return Error;
//Flush and make sure that calculation is over before returning control
if (Success != _context.flushQueue(err))
return Error;
if (Success != evt.wait(err))
return Error;
return Success;
}
void CTLord::CreateColleagues( void )
{ // 협력 객체의 생성 및 초기화
// 생성
m_pController = new CLController( this );
m_pController->CreateDbHandler( MIN( 1 ) );
CTElection * pElection = new CTElection( this );
CLEvent* pEvent = new CLEvent( this );
CLordSkill* pSkills = new CTLordSkill( this );
// 스크립트
pElection->Initialize( "election.inc" );
pEvent->Initialize( "lordevent.inc" );
pSkills->Initialize( "lordskill.txt" );
// 설정
m_pElection = pElection;
m_pEvent = pEvent;
m_pSkills = pSkills;
// 데이터베이스로부터 이전 상태 복원
if( !RestoreAll() )
{
election::OutputDebugString( "CTLord.CreateColleagues" );
exit( -1 );
}
}
bool CLCommand::IsExecutable() {
if (wait_events_complete_)
return true;
for (vector<CLEvent*>::iterator it = wait_events_.begin();
it != wait_events_.end();
++it) {
if (!(*it)->IsComplete())
return false;
}
for (vector<CLEvent*>::iterator it = wait_events_.begin();
it != wait_events_.end();
++it) {
CLEvent* event = *it;
if (event->IsError())
wait_events_good_ = false;
event->Release();
}
wait_events_.clear();
wait_events_complete_ = true;
return true;
}
TEST(CLProtocolDataPoster, STL_PostProtocolData_Features_Test)
{
CLProtocolDataPoster *pd = new CLProtocolDataPoster;
CLEvent e;
CLSTLqueue *q = new CLSTLqueue;
EXPECT_TRUE(pd->SetParameters(new CLDataPosterBySTLqueue(q), new CLDataPostResultNotifier(false), &e).IsSuccess());
CLIOVectors *piov = new CLIOVectors;
long i = 32;
EXPECT_TRUE(piov->PushBack((char *)(&i), 8).IsSuccess());
CLStatus s1 = pd->PostProtocolData(piov);
EXPECT_TRUE(s1.IsSuccess());
EXPECT_TRUE(e.Wait().IsSuccess());
long j = 0;
CLIOVectors iov1;
EXPECT_TRUE(iov1.PushBack((char *)(&j), 8).IsSuccess());
EXPECT_TRUE(q->PopData(iov1).IsSuccess());
EXPECT_EQ(j, 32);
delete q;
}
void GPURayTracer::RenderFrame(const Camera &camera) {
PROFILE_TIMER_START("frame");
frame_stopwatch_.Restart();
CLEvent ev;
PROFILE_TIMER_START("GPU");
PROFILE_TIMER_START("GPU init tracing");
/*__kernel void InitTracingFrame(__global TracingState *tracing_states,
float16 view_proj_inv,
float lod_voxel_size)*/
init_frame_kernel_->SetBufferArg(0, tracing_states_.get());
init_frame_kernel_->SetFloat16Arg(1, camera.ViewProjectionMatrix().Inverse());
init_frame_kernel_->SetFloatArg(2, lod_voxel_size_);
init_frame_kernel_->Run2D(frame_width_, frame_height_, NULL, &ev);
ev.WaitFor();
PROFILE_TIMER_COMMIT("GPU init tracing");
PROFILE_TIMER_STOP();
cache_manager_->StartRequestSession();
int iterations;
for (iterations = 0;; ++iterations) {
PROFILE_TIMER_START("GPU");
PROFILE_TIMER_START("GPU tracing passes");
/*__kernel void RaytracingPass(__global int *faults_and_hits, // out
__global TracingState *tracing_states, // in-out
__global uint *node_links,
__global uint *far_pointers,
int root_node_index)*/
raytracing_pass_kernel_->SetBufferArg(0, faults_and_hits_.get());
raytracing_pass_kernel_->SetBufferArg(1, tracing_states_.get());
raytracing_pass_kernel_->SetBufferArg(2, cache_manager_->data()->ChannelByIndex(0)->cl_buffer());
raytracing_pass_kernel_->SetBufferArg(3, cache_manager_->data()->far_pointers_buffer());
raytracing_pass_kernel_->SetIntArg(4, cache_manager_->root_node_index());
const CLBuffer *b = NULL;
if (cache_manager_->data()->ChannelByName("normals"))
b = cache_manager_->data()->ChannelByName("normals")->cl_buffer();
raytracing_pass_kernel_->SetBufferArg(5, b);
b = NULL;
if (cache_manager_->data()->ChannelByName("colors"))
b = cache_manager_->data()->ChannelByName("colors")->cl_buffer();
raytracing_pass_kernel_->SetBufferArg(6, b);
raytracing_pass_kernel_->Run2D(frame_width_, frame_height_, NULL, &ev);
ev.WaitFor();
PROFILE_TIMER_STOP();
PROFILE_TIMER_STOP();
// always give it at least one iteration of loading
if (iterations > 0 && frame_stopwatch_.TimeSinceRestart() > frame_time_limit_)
break;
feedback_extractor_->Process(*faults_and_hits_, cache_manager_->data()->nodes_in_block());
cache_manager_->StartRequestTransaction();
PROFILE_TIMER_START("CPU mark as used");
const vector<int> &hits = feedback_extractor_->hit_blocks();
for (int i = 0; i < (int)hits.size(); ++i) {
cache_manager_->MarkBlockAsUsed(hits[i]);
}
const vector<int> &dup_hits = feedback_extractor_->duplicate_hit_blocks();
for (int i = 0; i < (int)dup_hits.size(); ++i) {
cache_manager_->MarkParentAsUsed(dup_hits[i]);
}
PROFILE_TIMER_STOP();
PROFILE_VALUE_COMMIT_SET("hits per pass", hits.size());
PROFILE_VALUE_COMMIT_SET("hits dup per pass", dup_hits.size());
PROFILE_TIMER_START("CPU request blocks");
const vector<int> &faults = feedback_extractor_->fault_blocks();
for (int i = 0; i < (int)faults.size(); ++i) {
if (cache_manager_->TransactionFilledCache())
break;
cache_manager_->RequestBlock(faults[i]);
}
PROFILE_TIMER_STOP();
PROFILE_VALUE_ADD("faults", faults.size());
cache_manager_->EndRequestTransaction();
if (faults.empty() || cache_manager_->SessionFilledCache())
break;
}
PROFILE_TIMER_COMMIT("CPU request blocks");
PROFILE_TIMER_COMMIT("CPU mark as used");
PROFILE_TIMER_COMMIT("GPU tracing passes");
PROFILE_VALUE_COMMIT_SET("tracing iterations", iterations);
PROFILE_VALUE_COMMIT_SET("cache too small", cache_manager_->SessionFilledCache() ? 1 : 0);
PROFILE_VALUE_COMMIT("faults");
PROFILE_VALUE_COMMIT("uploaded blocks");
PROFILE_VALUE_COMMIT("updated far pointers");
PROFILE_VALUE_COMMIT("updated pointers");
cache_manager_->EndRequestSession();
PROFILE_TIMER_START("GPU");
PROFILE_TIMER_START("GPU finish tracing");
/*__kernel void FinishTracingFrame(__global uchar4 *result_colors,
__global TracingState *tracing_states,
__constant float4 background_color)*/
finish_frame_kernel_->SetBufferArg(0, out_image_.get());
finish_frame_kernel_->SetBufferArg(1, tracing_states_.get());
finish_frame_kernel_->SetFloat4Arg(2, fvec3(0, 216/255., 1), 1);
const CLBuffer *b = NULL;
if (cache_manager_->data()->ChannelByName("normals"))
b = cache_manager_->data()->ChannelByName("normals")->cl_buffer();
finish_frame_kernel_->SetBufferArg(3, b);
b = NULL;
if (cache_manager_->data()->ChannelByName("colors"))
b = cache_manager_->data()->ChannelByName("colors")->cl_buffer();
finish_frame_kernel_->SetBufferArg(4, b);
finish_frame_kernel_->Run2D(frame_width_, frame_height_, NULL, &ev);
ev.WaitFor();
PROFILE_TIMER_COMMIT("GPU finish tracing");
PROFILE_TIMER_STOP();
PROFILE_TIMER_COMMIT("frame");
// this part kinda breaks encapsulation
// the excuse for this is that it's just profiling code, not code that "matters" :)
PROFILE_TIMER_COMMIT("GPU sort");
PROFILE_TIMER_COMMIT("GPU scan");
PROFILE_TIMER_COMMIT("GPU/BUS compr/size");
PROFILE_TIMER_COMMIT("BUS read faults");
PROFILE_TIMER_COMMIT("BUS read hits");
PROFILE_TIMER_COMMIT("HDD? load block");
PROFILE_TIMER_COMMIT("CPU process links");
PROFILE_TIMER_COMMIT("BUS upload block");
PROFILE_TIMER_COMMIT("BUS upload far ptr");
PROFILE_TIMER_COMMIT("BUS upload ptr");
PROFILE_TIMER_COMMIT("BUS unload pointer");
PROFILE_TIMER_COMMIT("GPU");
PROFILE_TIMER_COMMIT("BUS");
PROFILE_TIMER_COMMIT("HDD");
}
/**Constructs BVH acceleration structure, according to associated scene state
* @param err Error info
* @return Result of the operation: Success or failure
*/
Result BVHManager::construct(Common::Errata& err)
{
try
{
SET_KERNEL_ARGS((*_mortonCalcKernel),_bvhNodes->getCLMem(),_sortedMortonCodes->getCLMem(),_scene.getDeviceSceneData());
SET_KERNEL_ARGS((*_radixTreeBuildKernel),_bvhNodes->getCLMem(),_sortedMortonCodes->getCLMem(),_bvhLeavesCount);
SET_KERNEL_ARGS((*_bbCalcKernel),_bvhNodes->getCLMem(),_nodeVisitCounters->getCLMem(), _bvhLeavesCount);
}
catch (CLInterfaceException e)
{
err = Errata(e);
return Error;
}
CLEvent evt;
evt.reset();
size_t processors, warp;
if (Success != _context.getMaximalLaunchExecParams(*_mortonCalcKernel,processors,warp,err))
return Error;
CL_UINT worksize = closestMultipleTo(_bvhLeavesCount,warp);
CLKernelWorkDimension globalDim(1,worksize);
CLKernelWorkDimension localDim(1,warp);
CLKernelExecuteParams mortonKernelExecParams(&globalDim,&localDim,&evt);
//1. Morton Codes
//Execute kernel
if (Success != _context.enqueueKernel((*_mortonCalcKernel),mortonKernelExecParams,err))
return Error;
//Flush and make sure that calculation is over before returning control
if (Success != _context.flushQueue(err))
return Error;
if (Success != evt.wait(err))
return Error;
//2.Sort the leaves by Morton Codes
if (Success != _bitonicSorter->sort(_sortedMortonCodes->getCLMem(),_mortonBufferItems,err))
return Error;
//3. Build Radix Tree
worksize = closestMultipleTo(_bvhLeavesCount-1,warp);
CLKernelWorkDimension globalDim_Radix(1,worksize);
CLKernelWorkDimension localDim_Radix(1,warp);
CLKernelExecuteParams radixKernelExecParams(&globalDim_Radix,&localDim_Radix,&evt);
//Execute kernel
if (Success != _context.enqueueKernel((*_radixTreeBuildKernel),radixKernelExecParams,err))
return Error;
//Flush and make sure that calculation is over before returning control
if (Success != _context.flushQueue(err))
return Error;
if (Success != evt.wait(err))
return Error;
worksize = closestMultipleTo(_bvhLeavesCount,warp);
CLKernelExecuteParams boundingBoxKernelExecParams(new CLKernelWorkDimension(1,worksize),new CLKernelWorkDimension(1,warp,&evt));
//Execute kernel
if (Success != _context.enqueueKernel((*_bbCalcKernel),boundingBoxKernelExecParams,err))
return Error;
//Flush and make sure that calculation is over before returning control
if (Success != _context.flushQueue(err))
return Error;
if (Success != evt.wait(err))
return Error;
return Success;
}
