static
char *GetTimedURLFromPlaylistItem( intf_thread_t *p_intf,
playlist_item_t *p_current_item )
{
#ifdef CMML_INTF_USE_TIMED_URIS
char *psz_url = NULL;
char *psz_return_value = NULL;
char *psz_seconds = NULL;
int i_seconds;
psz_url = XURL_GetWithoutFragment( p_current_item->input->psz_uri );
/* Get current time as a string */
if( XURL_IsFileURL( psz_url ) == VLC_TRUE )
psz_url = xstrcat( psz_url, "#" );
else
psz_url = xstrcat( psz_url, "?" );
/* jump back to 2 seconds before where we are now */
i_seconds = GetCurrentTimeInSeconds( p_intf->p_sys->p_input ) - 2;
psz_seconds = GetTimedURIFragmentForTime( i_seconds < 0 ? 0 : i_seconds );
if( psz_seconds )
{
psz_url = xstrcat( psz_url, psz_seconds );
free( psz_seconds );
psz_return_value = psz_url;
}
return psz_return_value;
#else
void *p;
/* Suppress warning messages about unused functions */
p = GetTimedURIFragmentForTime; /* unused */
p = GetCurrentTimeInSeconds; /* unused */
return strdup( p_current_item->input.psz_uri );
#endif
}
void Tracing() {
// Initialize d_tetrahedralLinks
err = cudaMalloc((void **)&d_tetrahedralLinks, sizeof(int) * globalNumOfCells * 4);
if (err) lcs::Error("Fail to create a buffer for device tetrahedralLinks");
err = cudaMemcpy(d_tetrahedralLinks, tetrahedralLinks, sizeof(int) * globalNumOfCells * 4, cudaMemcpyHostToDevice);
if (err) lcs::Error("Fail to enqueue copy for d_tetrahedralLinks");
// Initialize initial active particle data
InitializeInitialActiveParticles();
err = cudaMalloc((void **)&d_activeParticles, sizeof(int) * numOfInitialActiveParticles);
if (err) lcs::Error("Fail to create d_activeParticles");
err = cudaMalloc((void **)&d_exitCells, sizeof(int) * numOfInitialActiveParticles);
if (err) lcs::Error("Fail to create d_exitCells");
err = cudaMemcpy(d_exitCells, exitCells, sizeof(int) * numOfInitialActiveParticles, cudaMemcpyHostToDevice);
if (err) lcs::Error("Fail to initialize d_exitCells");
// Initialize velocity data
double *velocities[2];
int currStartVIndex = 1;
InitializeVelocityData(velocities);
// Create some dynamic device arrays
err = cudaMalloc((void **)&d_activeBlockList, sizeof(int) * numOfInterestingBlocks);
if (err) lcs::Error("Fail to create a buffer for device activeBlockList");
err = cudaMalloc((void **)&d_startOffsetInParticles, sizeof(int) * (numOfInterestingBlocks + 1));
if (err) lcs::Error("Fail to create a buffer for device startOffsetInParticles");
err = cudaMalloc((void **)&d_blockedActiveParticles, sizeof(int) * numOfInitialActiveParticles);
if (err) lcs::Error("Fail to create a buffer for device blockedAciveParticles");
err = cudaMalloc((void **)&d_blockedCellLocations, sizeof(int) * numOfInitialActiveParticles);
if (err) lcs::Error("Fail to create a buffer for device blockedCellLocations");
// Initialize activeParticles
int *activeParticles = new int [numOfInitialActiveParticles];
for (int i = 0; i < numOfInitialActiveParticles; i++)
activeParticles[i] = i;
// Create cellLocations
int *cellLocations = new int [numOfInitialActiveParticles];
// Create blockLocations
int *blockLocations = new int [numOfInitialActiveParticles];
// Create activeBlockIDs and activeBlockList
int *activeBlockIDs = new int [numOfInterestingBlocks]; // An interesting block has the ID if and only if it has particles
int *activeBlockIDList = new int [numOfInterestingBlocks];
// Create startOffsetInParticle
int *startOffsetInParticle = new int [numOfInterestingBlocks + 1];
// Create blockedActiveParticleIDList and topOfActiveBlocks
int *blockedActiveParticleIDList = new int [numOfInitialActiveParticles];
int *topOfActiveBlocks = new int [numOfInterestingBlocks];
// Create blockedCellLocations
int *blockedCellLocations = new int [numOfInitialActiveParticles];
// Create work arrays
int *countParticlesInInterestingBlocks = new int [numOfInterestingBlocks];
int *marks = new int [numOfInterestingBlocks];
memset(marks, 255, sizeof(int) * numOfInterestingBlocks);
int markCount = 0;
// Some start setting
double currTime = 0;
double interval = configure->GetTimeInterval();
// Create general squeezed arrays
int *squeezedStage = new int [numOfInitialActiveParticles];
int *squeezedExitCells = new int [numOfInitialActiveParticles];
// Create RK4-specific squeezed arrays
double *squeezedLastPositionForRK4, *squeezedK1ForRK4, *squeezedK2ForRK4, *squeezedK3ForRK4;
switch (lcs::ParticleRecord::GetDataType()) {
case lcs::ParticleRecord::RK4: {
squeezedLastPositionForRK4 = new double [numOfInitialActiveParticles * 3];
squeezedK1ForRK4 = new double [numOfInitialActiveParticles * 3];
squeezedK2ForRK4 = new double [numOfInitialActiveParticles * 3];
squeezedK3ForRK4 = new double [numOfInitialActiveParticles * 3];
}break;
}
// Main loop for blocked tracing
numOfTimePoints = configure->GetNumOfTimePoints();
double startTime = GetCurrentTimeInSeconds();
/// DEBUG ///
kernelSum = 0;
int numOfKernelCalls = 0;
/// DEBUG ///
FILE *tracer = fopen("tracer.txt", "w");
for (int frameIdx = 0; frameIdx + 1 < numOfTimePoints; frameIdx++, currTime += interval) {
printf("*********Tracing between frame %d and frame %d*********\n", frameIdx, frameIdx + 1);
printf("\n");
/// DEBUG ///
int startTime;
startTime = clock();
currStartVIndex = 1 - currStartVIndex;
int lastNumOfActiveParticles = 0;
for (int i = 0; i < numOfInitialActiveParticles; i++) {
if (exitCells[i] < -1) exitCells[i] = -(exitCells[i] + 2);
if (exitCells[i] != -1) activeParticles[lastNumOfActiveParticles++] = i;
}
/// DEBUG ///
printf("lastNumOfActiveParticles = %d\n", lastNumOfActiveParticles);
//std::random_shuffle(activeParticles, activeParticles + lastNumOfActiveParticles);
err = cudaMemcpy(d_activeParticles, activeParticles, sizeof(int) * lastNumOfActiveParticles, cudaMemcpyHostToDevice);
if (err) lcs::Error("Fail to initialize d_activeParticles");
// Load end velocities
LoadVelocities(velocities[1 - currStartVIndex], d_velocities[1 - currStartVIndex], (frameIdx + 1) % numOfFrames);
// Naive tracing
int kernelStart = clock();
LaunchGPUForNaiveTracing(d_vertexPositions,
d_velocities[currStartVIndex],
d_velocities[1 - currStartVIndex],
d_tetrahedralConnectivities,
d_tetrahedralLinks,
d_stage,
d_lastPositionForRK4,
d_k1ForRK4,
d_k2ForRK4,
d_k3ForRK4,
d_pastTimes,
d_exitCells,
currTime, currTime + interval, configure->GetTimeStep(),
configure->GetEpsilon(),
d_activeParticles,
lastNumOfActiveParticles
);
int kernelEnd = clock();
kernelSum += (kernelEnd - kernelStart) * 1.0 / CLOCKS_PER_SEC;
numOfKernelCalls++;
cudaMemcpy(exitCells, d_exitCells, sizeof(int) * numOfInitialActiveParticles, cudaMemcpyDeviceToHost);
int endTime = clock();
printf("This interval cost %lf sec.\n", (double)(endTime - startTime) / CLOCKS_PER_SEC);
printf("\n");
}
fclose(tracer);
/// DEBUG ///
printf("kernelSum = %lf\n", kernelSum);
printf("numOfKernelCalls = %d\n", numOfKernelCalls);
/// DEBUG ///
double endTime = GetCurrentTimeInSeconds();
printf("The total tracing time is %lf sec.\n", (double)(endTime - startTime));
printf("\n");
// Release work arrays
delete [] activeParticles;
delete [] cellLocations;
delete [] blockLocations;
delete [] countParticlesInInterestingBlocks;
delete [] marks;
delete [] activeBlockIDs;
delete [] activeBlockIDList;
delete [] startOffsetInParticle;
}
fcAPI fcTime fcGetTime()
{
return GetCurrentTimeInSeconds();
}
