Zhen Gou Project 1 submission

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" />
@@ -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>

Part1/src/kernel.cu

@@ -9,7 +9,7 @@
 dim3 threadsPerBlock(blockSize);
 
 int numObjects;
-const float planetMass = 3e8;
+const float planetMass = 3e7;
 const __device__ float starMass = 5e10;
 
 const float scene_scale = 2e2; //size of the height map in simulation space
@@ -87,21 +87,55 @@ __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 accelerate(int N, glm::vec4 my_pos, glm::vec4 * their_pos)
 {
-    return glm::vec3(0.0f);
+	int index = (blockIdx.x * blockDim.x) + threadIdx.x;
+	glm::vec3 totalF(0.0f);
+	for(int i = 0;i < N; i ++)
+	{
+		if(i!=index) totalF += universalForce(my_pos, their_pos[i]); //????
+	}
+
+	totalF += universalForce(my_pos,glm::vec4(0.0f,0.0f,0.0f,starMass));
+
+	return totalF;
+}
+
+
+__device__ glm::vec3 universalForce(glm::vec4 pos1,glm::vec4 pos2)
+{
+	glm::vec3 dir = glm::vec3(pos2) - glm::vec3(pos1);
+	float r = glm::length(dir);
+	//to prevent two planets getting too close and blow up the force
+	if(r<1.0f) return glm::vec3(0.0f);
+	glm::vec3 F =(float) G * dir *pos2.w /(r*r*r);
+	return F;
 }
 
 // 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
+	int index =(blockIdx.x * blockDim.x) + threadIdx.x;
+	if(index<N)
+	{
+		acc[index] = accelerate(N,pos[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)
 {
 	// FILL IN HERE
+	int index = (blockIdx.x * blockDim.x) + threadIdx.x;
+	if(index<N)
+	{
+		vel[index] += dt * (acc[index]);
+		pos[index].x += dt * (vel[index].x);
+		pos[index].y += dt * (vel[index].y);
+		pos[index].z += dt * (vel[index].z);
+	}
 }
 
 // Update the vertex buffer object
@@ -179,7 +213,11 @@ 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
+    dim3 fullBlocksPerGrid((int)ceil(float(numObjects)/float(blockSize)));
+
+	updateF<<<fullBlocksPerGrid, blockSize>>>(numObjects,dt,dev_pos,dev_vel,dev_acc);
+	updateS<<<fullBlocksPerGrid, blockSize>>>(numObjects,dt,dev_pos,dev_vel,dev_acc);
+    cudaThreadSynchronize();
 }
 
 void cudaUpdateVBO(float * vbodptr, int width, int height)

Part1/src/kernel.h

@@ -6,13 +6,14 @@
 #include <cuda.h>
 #include <cmath>
 
-#define blockSize 128
+#define blockSize 1024
 #define checkCUDAErrorWithLine(msg) checkCUDAError(msg, __LINE__)
 #define SHARED 0
 
 void checkCUDAError(const char *msg, int line);
 void cudaNBodyUpdateWrapper(float dt);
 void initCuda(int N);
+glm::vec3 universalForce(glm::vec4 pos1,glm::vec4 pos2);
 void cudaUpdatePBO(float4 * pbodptr, int width, int height);
 void cudaUpdateVBO(float * vbodptr, int width, int height);
 #endif

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

Part2/Project2/Project1/Project1.vcxproj

@@ -0,0 +1,78 @@
+<?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>
+  <ItemGroup>
+    <CudaCompile Include="matrix_math.cu" />
+  </ItemGroup>
+  <PropertyGroup Label="Globals">
+    <ProjectGuid>{0B2C1FC2-1FB8-4813-A8B5-BDD9B296214F}</ProjectGuid>
+    <RootNamespace>Project1</RootNamespace>
+  </PropertyGroup>
+  <Import Project="$(VCTargetsPath)\Microsoft.Cpp.Default.props" />
+  <PropertyGroup Condition="'$(Configuration)|$(Platform)'=='Debug|Win32'" Label="Configuration">
+    <ConfigurationType>Application</ConfigurationType>
+    <UseDebugLibraries>true</UseDebugLibraries>
+    <PlatformToolset>v110</PlatformToolset>
+    <CharacterSet>MultiByte</CharacterSet>
+  </PropertyGroup>
+  <PropertyGroup Condition="'$(Configuration)|$(Platform)'=='Release|Win32'" Label="Configuration">
+    <ConfigurationType>Application</ConfigurationType>
+    <UseDebugLibraries>false</UseDebugLibraries>
+    <PlatformToolset>v110</PlatformToolset>
+    <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>
+      <SDLCheck>true</SDLCheck>
+    </ClCompile>
+    <Link>
+      <GenerateDebugInformation>true</GenerateDebugInformation>
+      <AdditionalDependencies>kernel32.lib;user32.lib;gdi32.lib;winspool.lib;comdlg32.lib;advapi32.lib;shell32.lib;ole32.lib;oleaut32.lib;uuid.lib;odbc32.lib;odbccp32.lib;cudart.lib;%(AdditionalDependencies)</AdditionalDependencies>
+    </Link>
+    <CudaCompile>
+      <Include>$(CudaToolkitIncludeDir)</Include>
+    </CudaCompile>
+  </ItemDefinitionGroup>
+  <ItemDefinitionGroup Condition="'$(Configuration)|$(Platform)'=='Release|Win32'">
+    <ClCompile>
+      <WarningLevel>Level3</WarningLevel>
+      <Optimization>MaxSpeed</Optimization>
+      <FunctionLevelLinking>true</FunctionLevelLinking>
+      <IntrinsicFunctions>true</IntrinsicFunctions>
+      <SDLCheck>true</SDLCheck>
+    </ClCompile>
+    <Link>
+      <GenerateDebugInformation>true</GenerateDebugInformation>
+      <EnableCOMDATFolding>true</EnableCOMDATFolding>
+      <OptimizeReferences>true</OptimizeReferences>
+    </Link>
+  </ItemDefinitionGroup>
+  <Import Project="$(VCTargetsPath)\Microsoft.Cpp.targets" />
+  <ImportGroup Label="ExtensionTargets">
+    <Import Project="$(VCTargetsPath)\BuildCustomizations\CUDA 6.5.targets" />
+  </ImportGroup>
+</Project>

Part2/Project2/Project1/Project1.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>

Part2/Project2/Project1/matrix_math.cu

@@ -0,0 +1,134 @@
+#include <cuda.h>
+#include <cuda_runtime.h>
+#include <iostream>
+
+__global__ void dev_matrix_add(int dim, float * A, float * B, float * result)
+{
+	int row = (blockIdx.y * blockDim.y) + threadIdx.y;
+	int col = (blockIdx.x * blockDim.x) + threadIdx.x;
+	if(row < dim && col < dim) result[row * dim + col] = A[row * dim + col] + B[row * dim + col];
+}
+
+__global__ void dev_matrix_sub(int dim, float * A, float * B, float * result)
+{
+	int row = (blockIdx.y * blockDim.y) + threadIdx.y;
+	int col = (blockIdx.x * blockDim.x) + threadIdx.x;
+	if(row < dim && col < dim) result[row * dim + col] = A[row * dim + col] - B[row * dim + col];
+}
+
+__global__ void dev_matrix_mult(int dim, float * A, float * B, float * result)
+{
+	int row = (blockIdx.y * blockDim.y) + threadIdx.y;
+	int col = (blockIdx.x * blockDim.x) + threadIdx.x;
+	if(row >= dim || col >= dim) return;
+	float sum = 0.0f;
+	for (int i = 0; i < dim; i++)
+	{
+		sum += A[row * dim + i] * B[col * dim + i];
+	}
+
+	result[row * dim + col] = sum;
+}
+
+__global__ void dev_initialize(int dim, float * tar, float val)
+{
+	int row = (blockIdx.y * blockDim.y) + threadIdx.y;
+	int col = (blockIdx.x * blockDim.x) + threadIdx.x;
+	if(row >= dim || col >= dim) return;
+	tar[row * dim + col] = val;
+}
+
+void matrix_add(int dim, float * A, float * B, float * result)
+{
+	for(int row = 0; row < dim; row++)
+	{
+		for(int col = 0; col < dim; col++)
+		{
+			result[row * dim + col] = A[row * dim + col] + B[row * dim + col];
+		}
+	}
+}
+
+void matrix_sub(int dim, float * A, float * B, float * result)
+{
+	for(int row = 0; row < dim; row++)
+	{
+		for(int col = 0; col < dim; col++)
+		{
+			result[row * dim + col] = A[row * dim + col] - B[row * dim + col];
+		}
+	}
+}
+
+void matrix_mult(int dim, float * A, float * B, float * result)
+{
+	for(int row = 0; row < dim; row++)
+	{
+		for(int col = 0; col < dim; col++)
+		{
+			float sum = 0.0f;
+			for (int i = 0; i < dim; i++)
+			{
+				sum += A[row * dim + i] * B[col * dim + i];
+			}
+			result[row * dim + col] = sum;
+		}
+	}
+}
+
+int main(int argc, char** argv)
+{
+	float * A, *B, * result, * dev_A, * dev_B, * dev_result;
+	int M(4);
+	int tileWidth = 2;
+	int N = M * M * sizeof(float);
+
+	A = (float*) malloc( N);
+	B = (float*) malloc(N);
+	result = (float*) malloc(N);
+
+	cudaMalloc((void**) & dev_A, N);
+	cudaMalloc((void**) & dev_B, N);
+	cudaMalloc((void**) & dev_result,N);
+
+	//initialize 
+	for(int i = 0;i<M*M;i++)
+	{
+		A[i] = 1.0f;
+		B[i] = 2.0f;
+	}
+
+	cudaMemcpy(dev_A,A,N,cudaMemcpyHostToDevice);
+	cudaMemcpy(dev_B,B,N,cudaMemcpyHostToDevice);
+
+	dim3 gridDim((int)ceil((float)M/(float)tileWidth),(int)ceil((float)M/(float)tileWidth));
+	dim3 blockDim(tileWidth,tileWidth);
+	//dev_initialize<<<gridDim,blockDim>>>(M,dev_A,1.0f);
+	//dev_initialize<<<gridDim,blockDim>>>(M,dev_B,1.0f);
+	dev_matrix_mult<<<gridDim,blockDim>>>(M,dev_A,dev_B,dev_result);
+
+
+	cudaMemcpy(result, dev_result,N,cudaMemcpyDeviceToHost);
+
+	//display matrix
+	std::cout<<"matrix A:"<<std::endl;
+	for(int i = 0;i<M*M;i++)
+	{
+		std::cout<<A[i]<<" ";
+	}
+
+	std::cout<<std::endl<<"matrix B:"<<std::endl;
+	for(int i = 0;i<M*M;i++)
+	{
+		std::cout<<B[i]<<" ";
+	}
+
+	std::cout<<std::endl<<"result:"<<std::endl;
+	for(int i = 0;i<M*M;i++)
+	{
+		std::cout<<result[i]<<" ";
+	}
+	std::cin.get();
+    return 0;
+}
+

Part2/Project2/Project2.sln

@@ -0,0 +1,20 @@
+
+Microsoft Visual Studio Solution File, Format Version 12.00
+# Visual Studio 2012
+Project("{8BC9CEB8-8B4A-11D0-8D11-00A0C91BC942}") = "Project1", "Project1\Project1.vcxproj", "{0B2C1FC2-1FB8-4813-A8B5-BDD9B296214F}"
+EndProject
+Global
+	GlobalSection(SolutionConfigurationPlatforms) = preSolution
+		Debug|Win32 = Debug|Win32
+		Release|Win32 = Release|Win32
+	EndGlobalSection
+	GlobalSection(ProjectConfigurationPlatforms) = postSolution
+		{0B2C1FC2-1FB8-4813-A8B5-BDD9B296214F}.Debug|Win32.ActiveCfg = Debug|Win32
+		{0B2C1FC2-1FB8-4813-A8B5-BDD9B296214F}.Debug|Win32.Build.0 = Debug|Win32
+		{0B2C1FC2-1FB8-4813-A8B5-BDD9B296214F}.Release|Win32.ActiveCfg = Release|Win32
+		{0B2C1FC2-1FB8-4813-A8B5-BDD9B296214F}.Release|Win32.Build.0 = Release|Win32
+	EndGlobalSection
+	GlobalSection(SolutionProperties) = preSolution
+		HideSolutionNode = FALSE
+	EndGlobalSection
+EndGlobal