Changes

GPU610/Cosmosis

16,378 bytes added, 14:09, 13 April 2013

no edit summary

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

numThreads_ = (numThreads_ + 1) & ~1;

</div>

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

This is the what we thought of 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

CDOT Wiki β

Changes

GPU610/Cosmosis

CDOT Wiki ^β