Binglu Du Pull Request for Project-1

Part1/src/kernel.cu

@@ -10,14 +10,17 @@ dim3 threadsPerBlock(blockSize);
 
 int numObjects;
 const float planetMass = 3e8;
-const __device__ float starMass = 5e10;
+const __device__ float starMass = 5e12;
 
 const float scene_scale = 2e2; //size of the height map in simulation space
 
 glm::vec4 * dev_pos;
 glm::vec3 * dev_vel;
 glm::vec3 * dev_acc;
 
+
+
+
 void checkCUDAError(const char *msg, int line = -1)
 {
     cudaError_t err = cudaGetLastError();
@@ -87,21 +90,99 @@ __global__ void generateCircularVelArray(int time, int N, glm::vec3 * arr, glm::
 //		 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 accelerateone( glm::vec4 a_pos, glm::vec4 b_pos)
+{
+	glm::vec3 a_ab = glm::vec3(0.0f) ;
+
+	glm::vec3 r_a(a_pos.x, a_pos.y, a_pos.z);
+	glm::vec3 r_b(b_pos.x, b_pos.y, b_pos.z);
+	glm::vec3 r_ab = r_b - r_a;
+	float r = sqrt(r_ab.x*r_ab.x + r_ab.y*r_ab.y)+ EPSILON;
+
+	a_ab.x = (float)G * b_pos.w/(r*r*r)*r_ab.x;
+	a_ab.y = (float)G * b_pos.w/(r*r*r)*r_ab.y;
+	a_ab.z = (float)G * b_pos.w/(r*r*r)*r_ab.z;
+
+	 return a_ab;
+}
+
+
+// global memory
 __device__  glm::vec3 accelerate(int N, glm::vec4 my_pos, glm::vec4 * their_pos)
 {
-    return glm::vec3(0.0f);
+	//calculate the accelaration to center star at first
+	glm::vec3 returnAcc = accelerateone(my_pos, glm::vec4(0,0,0,starMass)) ;
+
+	for(int index=0; index<N; index++){
+
+			returnAcc += accelerateone(my_pos,their_pos[index]);
+
+	}
+    return returnAcc;
 }
 
+
+
+//shared memory
+/*__device__ glm::vec3 accelerate(int N, glm::vec4 my_pos, glm::vec4 * their_pos)
+{
+    glm::vec3 returnAcc = accelerateone(my_pos, glm::vec4(0,0,0,starMass));
+
+	__shared__ glm::vec4 sharedPositions[blockSize];
+	int positionsFullBlocks = (int)ceil((float)N /(float)blockSize);
+	for(int i = 0; i < positionsFullBlocks; ++i)
+	{
+		int index = threadIdx.x + i * blockSize; // index on global memory
+		if(index < N)
+		{
+			sharedPositions[threadIdx.x] = their_pos[index];		
+		}
+		__syncthreads();
+
+		for(int j = 0; j < blockSize && j + i * blockSize < N; ++j) 
+		{
+			returnAcc += accelerateone(my_pos, sharedPositions[j]);
+		}
+
+		__syncthreads();
+
+	}
+    return returnAcc;
+}*/
+
+
 // TODO : update the acceleration of each body
 __global__ void updateF(int N, float dt, glm::vec4 * pos, glm::vec3 * vel, glm::vec3 * acc)
 {
 	// FILL IN HERE
+	// f = ma
+	int index = threadIdx.x + (blockIdx.x * blockDim.x);
+
+
+	if(index < N){
+		acc[index] = accelerate(N, pos[index], pos);
+	}
+
+	//acc[index] += accelerate(N, *pos, 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)
 {
-	// FILL IN HERE
+	//f(t+dt) = f_prime(t)*dt + f(t)
+	int index = threadIdx.x + (blockIdx.x * blockDim.x);
+	if(index<N){
+		//Eular Intergration
+		vel[index].x += acc[index].x*dt;
+		vel[index].y += acc[index].y*dt;
+		vel[index].z += acc[index].z*dt;
+
+		//Eular Intergration
+		pos[index].x += (vel[index]).x*dt;
+		pos[index].y += (vel[index]).y*dt;
+		pos[index].z += (vel[index]).z*dt;
+	}
 }
 
 // Update the vertex buffer object
@@ -137,14 +218,15 @@ __global__ void sendToPBO(int N, glm::vec4 * pos, float4 * pbo, int width, int h
     float c_scale_h = height / s_scale;
 
     glm::vec3 color(0.05, 0.15, 0.3);
-    glm::vec3 acc = accelerate(N, glm::vec4((x-w2)/c_scale_w,(y-h2)/c_scale_h,0,1), pos);
-
+	glm::vec3 acc(0,0,0);
+	//glm::vec3 acc = accelerate(N, glm::vec4((x-w2)/c_scale_w,(y-h2)/c_scale_h,0,1), pos);
+	//glm::vec3 acc = dev_acc[index];
     if(x<width && y<height)
     {
-        float mag = sqrt(sqrt(acc.x*acc.x + acc.y*acc.y + acc.z*acc.z));
+		float mag = sqrt(sqrt(acc.x*acc.x + acc.y*acc.y + acc.z*acc.z));
 
 		// Each thread writes one pixel location in the texture (textel)
-        pbo[index].w = (mag < 1.0f) ? mag : 1.0f;
+		pbo[index].w = (mag < 1.0f) ? mag : 1.0f;
     }
 }
 
@@ -180,6 +262,16 @@ void initCuda(int N)
 void cudaNBodyUpdateWrapper(float dt)
 {
 	// FILL IN HERE
+	// execution configuration <<< Dgrid, Dblock, Nsharedmemory >>> 
+	//dim3 fullBlocksPerGrid((int)ceil(float(width*height)/float(blockSize)));
+
+    dim3 fullBlocksPerGrid((int)ceil(float(numObjects)/float(blockSize)));
+	//printf("test print");
+	//dim3 fullBlocksPerGrid(1);
+	updateF<<<fullBlocksPerGrid, blockSize, blockSize*sizeof(glm::vec4)>>>(numObjects, dt, dev_pos, dev_vel, dev_acc);
+	//utilityCore::printVec4(*dev_pos);
+	updateS<<<fullBlocksPerGrid, blockSize, blockSize*sizeof(glm::vec4)>>>(numObjects, dt, dev_pos, dev_vel, dev_acc);
+
 }
 
 void cudaUpdateVBO(float * vbodptr, int width, int height)

Part1/src/main.cpp

@@ -4,7 +4,7 @@
 
 #include "main.h"
 
-#define N_FOR_VIS 5000
+#define N_FOR_VIS 2000
 #define DT 0.2
 #define VISUALIZE 1
 //-------------------------------
@@ -16,6 +16,7 @@ int main(int argc, char** argv)
     // Launch CUDA/GL
 
     init(argc, argv);
+
 
     cudaGLSetGLDevice(0);
     initPBO(&pbo);

Part2/CUDAMatricMath/CUDAMatricMath.sln

@@ -0,0 +1,20 @@
+
+Microsoft Visual Studio Solution File, Format Version 11.00
+# Visual Studio 2010
+Project("{8BC9CEB8-8B4A-11D0-8D11-00A0C91BC942}") = "CUDAMatricMath", "CUDAMatricMath\CUDAMatricMath.vcxproj", "{615B7285-9C29-45D0-98E0-DBFAD80CF01E}"
+EndProject
+Global
+	GlobalSection(SolutionConfigurationPlatforms) = preSolution
+		Debug|Win32 = Debug|Win32
+		Release|Win32 = Release|Win32
+	EndGlobalSection
+	GlobalSection(ProjectConfigurationPlatforms) = postSolution
+		{615B7285-9C29-45D0-98E0-DBFAD80CF01E}.Debug|Win32.ActiveCfg = Debug|Win32
+		{615B7285-9C29-45D0-98E0-DBFAD80CF01E}.Debug|Win32.Build.0 = Debug|Win32
+		{615B7285-9C29-45D0-98E0-DBFAD80CF01E}.Release|Win32.ActiveCfg = Release|Win32
+		{615B7285-9C29-45D0-98E0-DBFAD80CF01E}.Release|Win32.Build.0 = Release|Win32
+	EndGlobalSection
+	GlobalSection(SolutionProperties) = preSolution
+		HideSolutionNode = FALSE
+	EndGlobalSection
+EndGlobal

Part2/CUDAMatricMath/CUDAMatricMath/CUDAMatricMath.vcxproj

@@ -0,0 +1,85 @@
+<?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>{615B7285-9C29-45D0-98E0-DBFAD80CF01E}</ProjectGuid>
+    <Keyword>Win32Proj</Keyword>
+    <RootNamespace>CUDAMatricMath</RootNamespace>
+  </PropertyGroup>
+  <Import Project="$(VCTargetsPath)\Microsoft.Cpp.Default.props" />
+  <PropertyGroup Condition="'$(Configuration)|$(Platform)'=='Debug|Win32'" Label="Configuration">
+    <ConfigurationType>Application</ConfigurationType>
+    <UseDebugLibraries>true</UseDebugLibraries>
+    <CharacterSet>Unicode</CharacterSet>
+  </PropertyGroup>
+  <PropertyGroup Condition="'$(Configuration)|$(Platform)'=='Release|Win32'" Label="Configuration">
+    <ConfigurationType>Application</ConfigurationType>
+    <UseDebugLibraries>false</UseDebugLibraries>
+    <WholeProgramOptimization>true</WholeProgramOptimization>
+    <CharacterSet>Unicode</CharacterSet>
+  </PropertyGroup>
+  <Import Project="$(VCTargetsPath)\Microsoft.Cpp.props" />
+  <ImportGroup Label="ExtensionSettings">
+    <Import Project="$(VCTargetsPath)\BuildCustomizations\CUDA 5.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 Condition="'$(Configuration)|$(Platform)'=='Debug|Win32'">
+    <LinkIncremental>true</LinkIncremental>
+  </PropertyGroup>
+  <PropertyGroup Condition="'$(Configuration)|$(Platform)'=='Release|Win32'">
+    <LinkIncremental>false</LinkIncremental>
+  </PropertyGroup>
+  <ItemDefinitionGroup Condition="'$(Configuration)|$(Platform)'=='Debug|Win32'">
+    <ClCompile>
+      <PrecompiledHeader>
+      </PrecompiledHeader>
+      <WarningLevel>Level3</WarningLevel>
+      <Optimization>Disabled</Optimization>
+      <PreprocessorDefinitions>WIN32;_DEBUG;_CONSOLE;%(PreprocessorDefinitions)</PreprocessorDefinitions>
+    </ClCompile>
+    <Link>
+      <SubSystem>Console</SubSystem>
+      <GenerateDebugInformation>true</GenerateDebugInformation>
+      <AdditionalDependencies>cudart.lib;%(AdditionalDependencies)</AdditionalDependencies>
+    </Link>
+  </ItemDefinitionGroup>
+  <ItemDefinitionGroup Condition="'$(Configuration)|$(Platform)'=='Release|Win32'">
+    <ClCompile>
+      <WarningLevel>Level3</WarningLevel>
+      <PrecompiledHeader>
+      </PrecompiledHeader>
+      <Optimization>MaxSpeed</Optimization>
+      <FunctionLevelLinking>true</FunctionLevelLinking>
+      <IntrinsicFunctions>true</IntrinsicFunctions>
+      <PreprocessorDefinitions>WIN32;NDEBUG;_CONSOLE;%(PreprocessorDefinitions)</PreprocessorDefinitions>
+    </ClCompile>
+    <Link>
+      <SubSystem>Console</SubSystem>
+      <GenerateDebugInformation>true</GenerateDebugInformation>
+      <EnableCOMDATFolding>true</EnableCOMDATFolding>
+      <OptimizeReferences>true</OptimizeReferences>
+    </Link>
+  </ItemDefinitionGroup>
+  <ItemGroup>
+    <CudaCompile Include="matrix_math.cu" />
+  </ItemGroup>
+  <Import Project="$(VCTargetsPath)\Microsoft.Cpp.targets" />
+  <ImportGroup Label="ExtensionTargets">
+    <Import Project="$(VCTargetsPath)\BuildCustomizations\CUDA 5.5.targets" />
+  </ImportGroup>
+</Project>

Part2/CUDAMatricMath/CUDAMatricMath/CUDAMatricMath.vcxproj.filters

@@ -0,0 +1,20 @@
+<?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" />
+  </ItemGroup>
+</Project>