Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Project 1: Gabriel Naghi #18

Open
wants to merge 45 commits into
base: master
Choose a base branch
from
Open
Changes from 1 commit
Commits
Show all changes
45 commits
Select commit Hold shift + click to select a range
8ffe708
brute force kernel and helper implementation
Sep 5, 2016
6d5e5f2
step simulation impl
Sep 5, 2016
b274078
update for Moore100 computers
Sep 5, 2016
1f2907c
= != ==
gabenaghi Sep 6, 2016
941f9d7
remove incorrect averaging, add relative COM velocity calc. Looks pre…
Sep 7, 2016
312b489
fixed rule3 and update README
Sep 7, 2016
48872be
initial part 2.1 stuff
Sep 7, 2016
834ab74
kernComputeIndices impl
Sep 8, 2016
5077f67
kernIdentifyCellStartEnd impl
Sep 8, 2016
fbf50cf
started writing getNeighbors. Yet i know its crap
Sep 8, 2016
cb23e69
this commit actually has getNeighbors;
gabenaghi Sep 8, 2016
745582c
inital work on getNeighbors as described by @krupkad
Sep 8, 2016
4b82963
does getNeighbors imple, some updateScattered impl
Sep 8, 2016
7f93c65
add rules calculation to updateScattered
Sep 8, 2016
970c8a9
finish updateScattered impl
Sep 8, 2016
9559bbf
finished uniformGrid impl. does not build
Sep 8, 2016
2690585
fix build errors
Sep 9, 2016
823bea9
Revert "fixed rule3 and update README"
Sep 9, 2016
b4e890d
whoops. Start debugging 2.1 functionality
Sep 9, 2016
1f746fb
lots of int/float casting confusing. build status unknown
Sep 9, 2016
c0e138d
initial implementation work for coherent search
Sep 9, 2016
50ee99f
a couple fixes. More dbug
Sep 11, 2016
ee95cc1
a couple minor fixes, more debug. Boids seem to be gravitating to cer…
Sep 11, 2016
23c925d
fix wrong boid index for rule2
Sep 11, 2016
85c4174
add second pos buffer
Sep 12, 2016
f80c0b3
initial coherent implementation. Doesnt work properly
Sep 12, 2016
70183a1
implementation complete. Now need to clean up, readme, and profile
Sep 12, 2016
c2b4c50
readme edits
Sep 12, 2016
7c0df7a
some profiling features
Sep 12, 2016
da98383
needs work
gabenaghi Sep 13, 2016
6923df4
boids everywhere
gabenaghi Sep 13, 2016
ad8c79d
Update README.md
gabenaghi Sep 13, 2016
2c4040e
add simulation image
gabenaghi Sep 13, 2016
1e5ff00
Update README.md
gabenaghi Sep 13, 2016
2bc2bc4
Merge branch 'master' of https://github.com/gabenaghi/Project1-CUDA-F…
gabenaghi Sep 13, 2016
e82f6ba
Update README.md
gabenaghi Sep 13, 2016
c4de626
performace analysis
Sep 13, 2016
f4b9963
Merge branch 'master' of https://github.com/gabenaghi/Project1-CUDA-F…
gabenaghi Sep 13, 2016
c696eff
README.md: analysis questions
Sep 13, 2016
a007436
add blocksize images
gabenaghi Sep 13, 2016
ea2c2b0
block size analysis and bugs
Sep 13, 2016
6ed7c83
add blocksize images
gabenaghi Sep 13, 2016
cf4e457
Merge branch 'master' of https://github.com/gabenaghi/Project1-CUDA-F…
gabenaghi Sep 13, 2016
2739d49
update readme with scattered threadperblock images
Sep 14, 2016
588e6e1
Update README.md
gabenaghi Sep 14, 2016
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
Prev Previous commit
Next Next commit
lots of int/float casting confusing. build status unknown
Gabriel Naghi authored and Gabriel Naghi committed Sep 9, 2016
commit 1f746fbf9354b27cde19b8c7cd2e1172d6b4d89b
40 changes: 21 additions & 19 deletions src/kernel.cu
Original file line number Diff line number Diff line change
@@ -352,13 +352,6 @@ __device__ int gridIndex3Dto1D(int x, int y, int z, int gridResolution) {
return x + y * gridResolution + z * gridResolution * gridResolution;
}

__device__ int vec3ToGridIndex(glm::vec3 pos, float inverseCellWidth, int gridResolution)
{
return gridIndex3Dto1D(
(int)(pos.x * inverseCellWidth),
(int)(pos.y * inverseCellWidth),
(int)(pos.z * inverseCellWidth), gridResolution);
}

__global__ void kernComputeIndices(int N, int gridResolution,
glm::vec3 gridMin, float inverseCellWidth,
@@ -420,51 +413,60 @@ __global__ void kernIdentifyCellStartEnd(int N, int *particleGridIndices,
}
}


__device__ glm::vec3 posToFloat3DIndex (glm::vec3 pos, glm::vec3 gridMin, float inverseCellWidth)
{
//to zero-index everything, must subtract off minimum value
//NOTE these are still floats!!
return glm::vec3(((pos.x - gridMin.x) * inverseCellWidth),
((pos.y - gridMin.y) * inverseCellWidth),
((pos.z - gridMin.z) * inverseCellWidth));
}


__device__ int getNeighbors(glm::vec3 pos, float inverseCellWidth,
float cellWidth, int gridResolution, int * neighbors)
float cellWidth, int gridResolution, glm::vec3 gridMin, int * neighbors)
{

float halfWidth = cellWidth * 0.5;
glm::vec3 myGridPos = glm::vec3((int)(pos.x * inverseCellWidth),
(int)(pos.y * inverseCellWidth),
(int)(pos.z * inverseCellWidth));
glm::vec3 myFloatGridPos = posToFloat3DIndex (pos, gridMin, inverseCellWidth);

glm::vec3 gridStart = glm::vec3( 0, 0, 0 );
glm::vec3 gridEnd = glm::vec3( 0, 0, 0 );

//if adding a halfwidth results in the same tile, then they are in
if (((pos.x + halfWidth) * inverseCellWidth) == myGridPos.x)
if ((int)((pos.x - gridMin.x + halfWidth) * inverseCellWidth) == (int) myGridPos.x)
gridStart.x = -1 ;
else
gridEnd.x = 1 ;

if ((int)((pos.y + halfWidth) * inverseCellWidth) == myGridPos.y)
if ((int)((pos.y - gridMin.y + halfWidth) * inverseCellWidth) == (int) myGridPos.y)
gridStart.y = -1 ;
else
gridEnd.y = 1 ;

if ((int)((pos.z + halfWidth) * inverseCellWidth) == myGridPos.z)
if ((int)((pos.z - gridMin.z + halfWidth) * inverseCellWidth) == (int) myGridPos.z)
gridStart.z = -1 ;
else
gridEnd.z = 1 ;

//our cell is always a "neighbor"
neighbors[0] = gridIndex3Dto1D(myGridPos.x, myGridPos.y, myGridPos.z, gridResolution);
neighbors[0] = gridIndex3Dto1D((int) myGridPos.x, (int) myGridPos.y, (int) myGridPos.z, gridResolution);

//calculate which cells are adjacent to me and put them in the buffer
int neighborCnt = 1; //self index is already in buffer, so start at 1

for (int i = myGridPos.x + gridStart.x; i < myGridPos.x + gridEnd.x; ++i)
for (int i = (int) myGridPos.x + (int) gridStart.x; i < (int) myGridPos.x + (int) gridEnd.x; ++i)
{
if (i < 0 || i >= gridResolution)
continue;

for (int j = myGridPos.y + gridStart.y; j < myGridPos.y + gridEnd.y; ++j)
for (int j = (int) myGridPos.y + (int) gridStart.y; j < (int) myGridPos.y + (int) gridEnd.y; ++j)
{
if (j < 0 || j >= gridResolution)
continue;

for (int k = myGridPos.z + gridStart.z; k < myGridPos.z + gridEnd.z; ++k)
for (int k = (int) myGridPos.z + (int) gridStart.z; k < (int) myGridPos.z + (int) gridEnd.z; ++k)
{
if (k < 0 || k >= gridResolution)
continue;
@@ -501,7 +503,7 @@ __global__ void kernUpdateVelNeighborSearchScattered(
//and its closest relevant neighbors
int neighbors[9];
int neighborCnt = getNeighbors(pos[myBoidIndex],
inverseCellWidth, cellWidth, gridResolution, neighbors);
inverseCellWidth, cellWidth, gridResolution, gridMin, neighbors);

glm::vec3 centerOfMass = glm::vec3(0.0f, 0.0f, 0.0f); //rule 1
glm::vec3 keepAway = glm::vec3(0.0f, 0.0f, 0.0f); //rule 2