CIS565-Fall-2014 · drerucha · Sep 22, 2014 · Sep 22, 2014 · Sep 22, 2014 · Sep 22, 2014
diff --git a/Part1/PROJ_WIN/CIS565_PROJ_1.sln b/Part1/PROJ_WIN/CIS565_PROJ_1.sln
@@ -3,6 +3,8 @@ Microsoft Visual Studio Solution File, Format Version 11.00
 # Visual Studio 2010
 Project("{8BC9CEB8-8B4A-11D0-8D11-00A0C91BC942}") = "CIS565_PROJ_1", "cIS565_PROJ_1\CIS565_PROJ_1.vcxproj", "{D7BEFF7A-4902-4B7E-922B-B0417A66864C}"
 EndProject
+Project("{8BC9CEB8-8B4A-11D0-8D11-00A0C91BC942}") = "CIS565_PROJ_1_PART_2", "..\..\Part2\CIS565_PROJ_1_PART_2\CIS565_PROJ_1_PART_2.vcxproj", "{B5E14F8D-EF1A-4763-808E-A148FEF3CCD2}"
+EndProject
 Global
 	GlobalSection(SolutionConfigurationPlatforms) = preSolution
 		Debug|Win32 = Debug|Win32
@@ -13,6 +15,10 @@ Global
 		{D7BEFF7A-4902-4B7E-922B-B0417A66864C}.Debug|Win32.Build.0 = Debug|Win32
 		{D7BEFF7A-4902-4B7E-922B-B0417A66864C}.Release|Win32.ActiveCfg = Release|Win32
 		{D7BEFF7A-4902-4B7E-922B-B0417A66864C}.Release|Win32.Build.0 = Release|Win32
+		{B5E14F8D-EF1A-4763-808E-A148FEF3CCD2}.Debug|Win32.ActiveCfg = Debug|Win32
+		{B5E14F8D-EF1A-4763-808E-A148FEF3CCD2}.Debug|Win32.Build.0 = Debug|Win32
+		{B5E14F8D-EF1A-4763-808E-A148FEF3CCD2}.Release|Win32.ActiveCfg = Release|Win32
+		{B5E14F8D-EF1A-4763-808E-A148FEF3CCD2}.Release|Win32.Build.0 = Release|Win32
 	EndGlobalSection
 	GlobalSection(SolutionProperties) = preSolution
 		HideSolutionNode = FALSE

diff --git a/Part1/PROJ_WIN/CIS565_PROJ_1.suo b/Part1/PROJ_WIN/CIS565_PROJ_1.suo
diff --git a/Part1/PROJ_WIN/CIS565_PROJ_1/CIS565_PROJ_1.vcxproj b/Part1/PROJ_WIN/CIS565_PROJ_1/CIS565_PROJ_1.vcxproj
@@ -30,7 +30,7 @@
   </PropertyGroup>
   <Import Project="$(VCTargetsPath)\Microsoft.Cpp.props" />
   <ImportGroup Label="ExtensionSettings">
-    <Import Project="$(VCTargetsPath)\BuildCustomizations\CUDA 5.5.props" />
+    <Import Project="$(VCTargetsPath)\BuildCustomizations\CUDA 6.5.props" />
   </ImportGroup>
   <ImportGroup Label="PropertySheets" Condition="'$(Configuration)|$(Platform)'=='Debug|Win32'">
     <Import Project="$(UserRootDir)\Microsoft.Cpp.$(Platform).user.props" Condition="exists('$(UserRootDir)\Microsoft.Cpp.$(Platform).user.props')" Label="LocalAppDataPlatform" />
@@ -62,7 +62,7 @@
       <GPUDebugInfo>true</GPUDebugInfo>
       <GenerateLineInfo>true</GenerateLineInfo>
       <HostDebugInfo>true</HostDebugInfo>
-      <CodeGeneration>compute_10,sm_10;compute_20,sm_20;compute_30,sm_30</CodeGeneration>
+      <CodeGeneration>compute_20,sm_20;compute_20,sm_20;compute_30,sm_30</CodeGeneration>
     </CudaCompile>
   </ItemDefinitionGroup>
   <ItemDefinitionGroup Condition="'$(Configuration)|$(Platform)'=='Release|Win32'">
@@ -95,7 +95,7 @@
   <ItemGroup>
     <CudaCompile Include="..\..\src\kernel.cu">
       <FileType>Document</FileType>
-      <CodeGeneration Condition="'$(Configuration)|$(Platform)'=='Debug|Win32'">compute_10,sm_10;compute_20,sm_20</CodeGeneration>
+      <CodeGeneration Condition="'$(Configuration)|$(Platform)'=='Debug|Win32'">compute_20,sm_20;compute_20,sm_20</CodeGeneration>
     </CudaCompile>
   </ItemGroup>
   <ItemGroup>
@@ -114,6 +114,6 @@
   </ItemGroup>
   <Import Project="$(VCTargetsPath)\Microsoft.Cpp.targets" />
   <ImportGroup Label="ExtensionTargets">
-    <Import Project="$(VCTargetsPath)\BuildCustomizations\CUDA 5.5.targets" />
+    <Import Project="$(VCTargetsPath)\BuildCustomizations\CUDA 6.5.targets" />
   </ImportGroup>
 </Project>
diff --git a/Part1/PROJ_WIN/CIS565_PROJ_1/vc100.pdb b/Part1/PROJ_WIN/CIS565_PROJ_1/vc100.pdb
diff --git a/Part1/PROJ_WIN/src/kernel.cu.deps b/Part1/PROJ_WIN/src/kernel.cu.deps
diff --git a/Part1/src/kernel.cu b/Part1/src/kernel.cu
@@ -27,8 +27,9 @@ void checkCUDAError(const char *msg, int line = -1)
         {
             fprintf(stderr, "Line %d: ", line);
         }
-        fprintf(stderr, "Cuda error: %s: %s.\n", msg, cudaGetErrorString( err) ); 
-        exit(EXIT_FAILURE); 
+        fprintf(stderr, "Cuda error: %s: %s.\n", msg, cudaGetErrorString( err) );
+
+        exit(EXIT_FAILURE);
     }
 } 
 
@@ -83,25 +84,105 @@ __global__ void generateCircularVelArray(int time, int N, glm::vec3 * arr, glm::
     }
 }
 
+// Helper function for accelerate().
+// Compute the gravitational acceleration body2 imposes on body1.
+__device__
+glm::vec3 computeGravitationalAccelerationBetweenTwoBodies( glm::vec4 body1, glm::vec4 body2 )
+{
+	// NOTE: F = (G * m_a * m_b) / (r_ab ^ 2)
+	//		 F = ma => a = F / m.
+	//		 So, a = (G * m_a * m_b) / (m_a * r_ab ^ 2) => a = (G * m_b) / (r_ab ^ 2)
+
+	// NOTE: The the 4th component (w) of each body is that body's mass.
+
+	glm::vec3 r( body2.x - body1.x, body2.y - body1.y, body2.z - body1.z );
+
+	// Threshold value determined through trial and error so that planets aren't flung out into the ether.
+	// Also, this avoids potential divide by zeros.
+	if ( glm::length( r ) < 0.2f ) {
+		return glm::vec3( 0.0f, 0.0f, 0.0f );
+	}
+
+	float numerator = ( float )( G * body2.w );
+	float denominator = ( float )( glm::length( r ) * glm::length( r ) );
+
+	return ( numerator / denominator ) * glm::normalize( r );
+}
+
 // TODO: Core force calc kernel global memory
 //		 HINT : You may want to write a helper function that will help you 
 //              calculate the acceleration contribution of a single body.
 //		 REMEMBER : F = (G * m_a * m_b) / (r_ab ^ 2)
-__device__  glm::vec3 accelerate(int N, glm::vec4 my_pos, glm::vec4 * their_pos)
+// N: Number of planets.
+// my_pos: Position of current planet.
+// their_pos: Array of planet positions.
+__device__
+glm::vec3 accelerate( int N, glm::vec4 my_pos, glm::vec4 *their_pos )
 {
-    return glm::vec3(0.0f);
+	// Danny was here.
+
+	glm::vec3 acceleration( 0.0f, 0.0f, 0.0f );
+
+	// Compute acceleration caused by the star at the center of the "solar system".
+	acceleration += computeGravitationalAccelerationBetweenTwoBodies( my_pos, glm::vec4( 0.0f, 0.0f, 0.0f, starMass ) );
+
+	// Compute acceleration caused by the other N-1 planets.
+	for ( unsigned int i = 0; i < N; ++i ) {
+		acceleration += computeGravitationalAccelerationBetweenTwoBodies( my_pos, their_pos[i] );
+	}
+
+	return acceleration;
 }
 
 // TODO : update the acceleration of each body
-__global__ void updateF(int N, float dt, glm::vec4 * pos, glm::vec3 * vel, glm::vec3 * acc)
+// N: Number of planets.
+// dt: Time step.
+// pos: Array of planet positions.
+// vel: Array of planet velocities.
+// acc: array of planet accelerations.
+__global__
+void updateF( int N, float dt, glm::vec4 * pos, glm::vec3 * vel, glm::vec3 * acc )
 {
-	// FILL IN HERE
+	// Danny was here.
+
+	// NOTE: Blocks are a one-dimensional array in the grid (they only have an x-component).
+	// NOTE: Threads are a one dimensional array in the current block (they only have an x-component).
+
+	int planet_index = threadIdx.x + ( blockIdx.x * blockDim.x );
+
+	// If number of planets do not fit perfectly inside the number of blocks specified,
+	// then some thread indices will not correlate to a planet, so check bounds.
+	if ( planet_index < N ) {
+		acc[planet_index] = accelerate( N, pos[planet_index], pos );
+	}
 }
 
 // TODO : update velocity and position using a simple Euler integration scheme
-__global__ void updateS(int N, float dt, glm::vec4 * pos, glm::vec3 * vel, glm::vec3 * acc)
+// N: Number of planets.
+// dt: Time step.
+// pos: Array of planet positions.
+// vel: Array of planet velocities.
+// acc: array of planet accelerations.
+__global__
+void updateS( int N, float dt, glm::vec4 *pos, glm::vec3 *vel, glm::vec3 *acc )
 {
-	// FILL IN HERE
+	// Danny was here.
+
+	// NOTE: Blocks are a one-dimensional array in the grid (they only have an x-component).
+	// NOTE: Threads are a one dimensional array in the current block (they only have an x-component).
+
+	int planet_index = threadIdx.x + ( blockIdx.x * blockDim.x );
+
+	// If number of planets do not fit perfectly inside the number of blocks specified,
+	// then some thread indices will not correlate to a planet, so check bounds.
+	if ( planet_index < N ) {
+		vel[planet_index] += ( acc[planet_index] * dt );
+
+		// Must separate into x-, y-, and z-components since pos is an array of glm::vec4s while vel is an array of glm::vec3s.
+		pos[planet_index].x += ( vel[planet_index].x * dt );
+		pos[planet_index].y += ( vel[planet_index].y * dt );
+		pos[planet_index].z += ( vel[planet_index].z * dt );
+	}
 }
 
 // Update the vertex buffer object
@@ -179,7 +260,21 @@ void initCuda(int N)
 // TODO : Using the functions you wrote above, write a function that calls the CUDA kernels to update a single sim step
 void cudaNBodyUpdateWrapper(float dt)
 {
-	// FILL IN HERE
+	// Danny was here.
+
+	// Compute grid dimensions to pass into kernel.
+	// Block dimensions are already given in threadsPerBlock.
+	dim3 blocks_per_grid( ( int )ceil( ( float )numObjects / ( float )blockSize ) );
+
+	// dev_pos, dev_vel, and dev_acc are arrays of vectors that have already been allocated on the device.
+	updateF<<<blocks_per_grid, threadsPerBlock>>>( numObjects, dt, dev_pos, dev_vel, dev_acc );
+	checkCUDAErrorWithLine("Kernel failed!"); // Because initCuda() does this.
+
+	// dev_pos, dev_vel, and dev_acc are arrays of vectors that have already been allocated on the device.
+	updateS<<<blocks_per_grid, threadsPerBlock>>>(numObjects, dt, dev_pos, dev_vel, dev_acc);
+	checkCUDAErrorWithLine("Kernel failed!"); // Because initCuda() does this.
+
+	cudaThreadSynchronize(); // Because cudaThreadSynchronize() and cudaUpdatePBO() do this.
 }
 
 void cudaUpdateVBO(float * vbodptr, int width, int height)
@@ -194,6 +289,4 @@ void cudaUpdatePBO(float4 * pbodptr, int width, int height)
     dim3 fullBlocksPerGrid((int)ceil(float(width*height)/float(blockSize)));
     sendToPBO<<<fullBlocksPerGrid, blockSize, blockSize*sizeof(glm::vec4)>>>(numObjects, dev_pos, pbodptr, width, height, scene_scale);
     cudaThreadSynchronize();
-}
-
-
+}
diff --git a/Part1/src/main.cpp b/Part1/src/main.cpp
@@ -62,7 +62,7 @@ void runCuda()
     // execute the kernel
     cudaNBodyUpdateWrapper(DT);
 #if VISUALIZE == 1
-    cudaUpdatePBO(dptr, field_width, field_height);
+    //cudaUpdatePBO(dptr, field_width, field_height);
     cudaUpdateVBO(dptrvert, field_width, field_height);
 #endif
     // unmap buffer object

diff --git a/Part2/CIS565_PROJ_1_PART_2/CIS565_PROJ_1_PART_2.vcxproj b/Part2/CIS565_PROJ_1_PART_2/CIS565_PROJ_1_PART_2.vcxproj
@@ -0,0 +1,77 @@
+<?xml version="1.0" encoding="utf-8"?>
+<Project DefaultTargets="Build" ToolsVersion="4.0" xmlns="http://schemas.microsoft.com/developer/msbuild/2003">
+  <ItemGroup Label="ProjectConfigurations">
+    <ProjectConfiguration Include="Debug|Win32">
+      <Configuration>Debug</Configuration>
+      <Platform>Win32</Platform>
+    </ProjectConfiguration>
+    <ProjectConfiguration Include="Release|Win32">
+      <Configuration>Release</Configuration>
+      <Platform>Win32</Platform>
+    </ProjectConfiguration>
+  </ItemGroup>
+  <PropertyGroup Label="Globals">
+    <ProjectGuid>{B5E14F8D-EF1A-4763-808E-A148FEF3CCD2}</ProjectGuid>
+    <RootNamespace>CIS565_PROJ_1_PART_2</RootNamespace>
+  </PropertyGroup>
+  <Import Project="$(VCTargetsPath)\Microsoft.Cpp.Default.props" />
+  <PropertyGroup Condition="'$(Configuration)|$(Platform)'=='Debug|Win32'" Label="Configuration">
+    <ConfigurationType>Application</ConfigurationType>
+    <UseDebugLibraries>true</UseDebugLibraries>
+    <CharacterSet>MultiByte</CharacterSet>
+  </PropertyGroup>
+  <PropertyGroup Condition="'$(Configuration)|$(Platform)'=='Release|Win32'" Label="Configuration">
+    <ConfigurationType>Application</ConfigurationType>
+    <UseDebugLibraries>false</UseDebugLibraries>
+    <WholeProgramOptimization>true</WholeProgramOptimization>
+    <CharacterSet>MultiByte</CharacterSet>
+  </PropertyGroup>
+  <Import Project="$(VCTargetsPath)\Microsoft.Cpp.props" />
+  <ImportGroup Label="ExtensionSettings">
+    <Import Project="$(VCTargetsPath)\BuildCustomizations\CUDA 6.5.props" />
+  </ImportGroup>
+  <ImportGroup Label="PropertySheets" Condition="'$(Configuration)|$(Platform)'=='Debug|Win32'">
+    <Import Project="$(UserRootDir)\Microsoft.Cpp.$(Platform).user.props" Condition="exists('$(UserRootDir)\Microsoft.Cpp.$(Platform).user.props')" Label="LocalAppDataPlatform" />
+  </ImportGroup>
+  <ImportGroup Label="PropertySheets" Condition="'$(Configuration)|$(Platform)'=='Release|Win32'">
+    <Import Project="$(UserRootDir)\Microsoft.Cpp.$(Platform).user.props" Condition="exists('$(UserRootDir)\Microsoft.Cpp.$(Platform).user.props')" Label="LocalAppDataPlatform" />
+  </ImportGroup>
+  <PropertyGroup Label="UserMacros" />
+  <PropertyGroup />
+  <ItemDefinitionGroup Condition="'$(Configuration)|$(Platform)'=='Debug|Win32'">
+    <ClCompile>
+      <WarningLevel>Level3</WarningLevel>
+      <Optimization>Disabled</Optimization>
+    </ClCompile>
+    <Link>
+      <GenerateDebugInformation>true</GenerateDebugInformation>
+    </Link>
+  </ItemDefinitionGroup>
+  <ItemDefinitionGroup Condition="'$(Configuration)|$(Platform)'=='Release|Win32'">
+    <ClCompile>
+      <WarningLevel>Level3</WarningLevel>
+      <Optimization>MaxSpeed</Optimization>
+      <FunctionLevelLinking>true</FunctionLevelLinking>
+      <IntrinsicFunctions>true</IntrinsicFunctions>
+    </ClCompile>
+    <Link>
+      <GenerateDebugInformation>true</GenerateDebugInformation>
+      <EnableCOMDATFolding>true</EnableCOMDATFolding>
+      <OptimizeReferences>true</OptimizeReferences>
+      <AdditionalDependencies>cudart.lib;%(AdditionalDependencies)</AdditionalDependencies>
+    </Link>
+    <CudaCompile>
+      <CompileOut>$(ProjectDir)$(Platform)/$(Configuration)/%(Filename)%(Extension).obj</CompileOut>
+      <Include>$(CudaToolkitIncludeDir);%(Include)</Include>
+    </CudaCompile>
+  </ItemDefinitionGroup>
+  <ItemGroup>
+    <CudaCompile Include="matrix_math.cu">
+      <FileType>Document</FileType>
+    </CudaCompile>
+  </ItemGroup>
+  <Import Project="$(VCTargetsPath)\Microsoft.Cpp.targets" />
+  <ImportGroup Label="ExtensionTargets">
+    <Import Project="$(VCTargetsPath)\BuildCustomizations\CUDA 6.5.targets" />
+  </ImportGroup>
+</Project>
diff --git a/Part2/CIS565_PROJ_1_PART_2/CIS565_PROJ_1_PART_2.vcxproj.filters b/Part2/CIS565_PROJ_1_PART_2/CIS565_PROJ_1_PART_2.vcxproj.filters
@@ -0,0 +1,22 @@
+<?xml version="1.0" encoding="utf-8"?>
+<Project ToolsVersion="4.0" xmlns="http://schemas.microsoft.com/developer/msbuild/2003">
+  <ItemGroup>
+    <Filter Include="Source Files">
+      <UniqueIdentifier>{4FC737F1-C7A5-4376-A066-2A32D752A2FF}</UniqueIdentifier>
+      <Extensions>cpp;c;cc;cxx;def;odl;idl;hpj;bat;asm;asmx</Extensions>
+    </Filter>
+    <Filter Include="Header Files">
+      <UniqueIdentifier>{93995380-89BD-4b04-88EB-625FBE52EBFB}</UniqueIdentifier>
+      <Extensions>h;hpp;hxx;hm;inl;inc;xsd</Extensions>
+    </Filter>
+    <Filter Include="Resource Files">
+      <UniqueIdentifier>{67DA6AB6-F800-4c08-8B7A-83BB121AAD01}</UniqueIdentifier>
+      <Extensions>rc;ico;cur;bmp;dlg;rc2;rct;bin;rgs;gif;jpg;jpeg;jpe;resx;tiff;tif;png;wav;mfcribbon-ms</Extensions>
+    </Filter>
+  </ItemGroup>
+  <ItemGroup>
+    <CudaCompile Include="matrix_math.cu">
+      <Filter>Source Files</Filter>
+    </CudaCompile>
+  </ItemGroup>
+</Project>