Changes

← Older edit

GPU610/Cosmosis

28,552 bytes added, 15:10, 13 April 2013

→‎Dynamic Shared Memory (not implemented)

==== Baseline ====

The following profiles were made under the following compiler and computer settings:

<pre>nvcc main.cpp timer.cpp sim\simbody.cu sim\simulation.cu -DWIN32 -O3</pre>

* i5 2500K @ 4.5Ghz

* Nvidia GTX 560Ti

* Raw computations, no graphics drawn to the screen for unbiased results.

* Random position, velocity and mass for each body.

* Brute force algorithm for calculating the forces (O(n^2)).

==== Initial Profiling ====

We initially intended on using the fast-math library provided by CUDA. At first our results were marginally faster than our previous code. Though after some optimizations we discovered that our code actually performed better than the fast-math library. With fast-math, it took 0.451889 seconds to process 1000 bodies for 512 samples, conversely without fast-math we got 0.409581 seconds, which is a considerable improvement.

==== Optimized Code ====

void __global__ SimCalc(BodyArray a)

{

int_fast32_t idx = blockIdx.x * blockDim.x + threadIdx.x;

if (idx < a.size) {

const _T G = 6.67384f * pow(10.0f, -11.0f);

//precompute positions at index

const _T px = a.array[idx].Position.x;

const _T py = a.array[idx].Position.y;

//mass at the index

const _T M_idx = G*a.array[idx].Mass;

a.array[idx].Force = vec2_t();

for (int_fast32_t j(0); j != a.size; ++j) {

if (idx != j) {

_T dx = a.array[j].Position.x - px;

_T dy = a.array[j].Position.y - py;

_T r = sqrt(dx*dx + dy*dy);

_T F = (M_idx*a.array[j].Mass)/(r*r);

a.array[idx].Force.x += F * (dx / r);

a.array[idx].Force.y += F * (dy / r);

}

void __global__ SimTick(BodyArray a, _T dt)

{

int idx = blockIdx.x * blockDim.x + threadIdx.x;

if (idx < a.size)

{

_T mass = a.array[idx].Mass;

a.array[idx].Velocity.x += dt * (a.array[idx].Force.x / mass);

a.array[idx].Velocity.y += dt * (a.array[idx].Force.y / mass);

a.array[idx].Position.x += dt * a.array[idx].Velocity.x;

a.array[idx].Position.y += dt * a.array[idx].Velocity.y;

}

</div>

==== Launch Control ====

We used the following calculations to determine the the number of threads and blocks to launch with:

numThreads_ = prop.maxThreadsPerMultiProcessor / maxBlocks;

numBlocks_ = (bodies_.size() + numThreads_ - 1) / numThreads_;

numThreads_ = (numThreads_ + 1) & ~1;

</div>

==== Dynamic Shared Memory (not implemented) ====

This is the roundabout way we thought of, of how to send in chunks to the kernel so that the kernel can handle shared memory size of no greater than the max shared memory size of the GPU:

CHUNKSIZE = 512;

shared_ = CHUNKSIZE * sizeof(SimBody);

while (chunks > 0)

{

BodyArray ba = { &arr.array[index], CHUNKSIZE };

SimCalc <<< numBlocks_, numThreads_, shared_ >>>(ba);

cudaThreadSynchronize();

SimTick <<< numBlocks_, numThreads_, shared_ >>>(ba, timeStep);

cudaThreadSynchronize();

index += CHUNKSIZE;

--chunks;

}

chunks = arr.size / CHUNKSIZE + 1;

index = 0;

</div>

It handles calculations in chunks so that the kernel can do calculations on body sizes of more than 1175 for devices with compute capabilities of 3.x.

Here is what the shared memory kernels would look like (not implemented because not correct):

void __global__ SimCalc(BodyArray a)

{

int_fast32_t idx = blockIdx.x * blockDim.x + threadIdx.x;

int tid = threadIdx.x;

extern __shared__ SimBody sa[];

if (idx >= a.size)

return;

sa[tid] = a.array[idx];

__syncthreads();

const _T G = 6.67384f * pow(10.0f, -11.0f);

//precompute positions at index

const _T px = sa[tid].Position.x;

const _T py = sa[tid].Position.y;

//mass at the index

const _T M_idx = G*sa[tid].Mass;

sa[tid].Force = vec2_t();

for (int_fast32_t j(0); j != a.size; ++j) {

if (idx != j) {

_T dx = a.array[j].Position.x - px;

_T dy = a.array[j].Position.y - py;

_T r = sqrt(dx*dx + dy*dy);

_T F = (M_idx*a.array[j].Mass)/(r*r);

sa[tid].Force.x += F * (dx / r);

sa[tid].Force.y += F * (dy / r);

}

__syncthreads();

}

a.array[idx] = sa[tid];

}

void __global__ SimTick(BodyArray a, _T dt)

{

int idx = blockIdx.x * blockDim.x + threadIdx.x;

int tid = threadIdx.x;

extern __shared__ SimBody sa[];

if (idx >= a.size)

return;

sa[tid] = a.array[idx];

__syncthreads();

_T mass = sa[tid].Mass;

sa[tid].Velocity.x += dt * (sa[tid].Force.x / mass);

sa[tid].Velocity.y += dt * (sa[tid].Force.y / mass);

sa[tid].Position.x += dt * sa[tid].Velocity.x;

sa[tid].Position.y += dt * sa[tid].Velocity.y;

__syncthreads();

a.array[idx] = sa[tid];

}

</div>

Nbguzman

1

edit

Changes

GPU610/Cosmosis

Navigation menu

Personal tools

Namespaces

Variants

Views

More

Search

Navigation

get involved with CDOT

courses

course projects

links

Tools