CMake example.

NVIDIA-developer-blog · Jun 9, 2017 · 030fb1a · 030fb1a · enemy1205 · Nov 17, 2022
1 parent 33b83df
commit 030fb1a
Show file tree

Hide file tree

Showing 6 changed files with 357 additions and 0 deletions.
diff --git a/posts/cmake/CMakeLists.txt b/posts/cmake/CMakeLists.txt
@@ -0,0 +1,47 @@
+
+cmake_minimum_required(VERSION 3.8 FATAL_ERROR)
+project(cmake_and_cuda LANGUAGES CXX CUDA)
+
+include(CTest)
+
+add_library(particles STATIC
+  particle.cu
+  particle.h
+  v3.cu
+  v3.h
+  )
+
+# Request that particles be built with -std=c++11
+# As this is a public compile feature anything that links to particles
+# will also build with -std=c++11
+target_compile_features(particles PUBLIC cxx_std_11)
+
+# We need to explicitly state that we need all CUDA files in the particle
+# library to be built with -dc as the member functions could be called by
+# other libraries and executables
+set_target_properties( particles
+                       PROPERTIES CUDA_SEPARABLE_COMPILATION ON)
+
+if(BUILD_TESTING)
+
+  add_executable(particle_test test.cu)
+
+  set_target_properties(particle_test PROPERTIES CUDA_SEPARABLE_COMPILATION ON)
+  target_link_libraries(particle_test PRIVATE particles)
+
+  add_test(NAME particles_10k COMMAND particle_test 10000 )
+  add_test(NAME particles_256k COMMAND particle_test 256000 )
+
+  if(APPLE)
+    # We need to add the default path to the driver (libcuda.dylib) as an rpath,
+    # so that the static cuda runtime can find it at runtime.
+    target_link_libraries(particle_test
+                          PRIVATE
+                          "-Wl,-rpath,${CMAKE_CUDA_IMPLICIT_LINK_DIRECTORIES}"
+                          )
+  endif()
+
+endif()
+
+
+
diff --git a/posts/cmake/particle.cu b/posts/cmake/particle.cu
@@ -0,0 +1,57 @@
+/* Copyright (c) 1993-2015, NVIDIA CORPORATION. All rights reserved.
+ *
+ * Redistribution and use in source and binary forms, with or without
+ * modification, are permitted provided that the following conditions
+ * are met:
+ *  * Redistributions of source code must retain the above copyright
+ *    notice, this list of conditions and the following disclaimer.
+ *  * Redistributions in binary form must reproduce the above copyright
+ *    notice, this list of conditions and the following disclaimer in the
+ *    documentation and/or other materials provided with the distribution.
+ *  * Neither the name of NVIDIA CORPORATION nor the names of its
+ *    contributors may be used to endorse or promote products derived
+ *    from this software without specific prior written permission.
+ *
+ * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS ``AS IS'' AND ANY
+ * EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+ * IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR
+ * PURPOSE ARE DISCLAIMED.  IN NO EVENT SHALL THE COPYRIGHT OWNER OR
+ * CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL,
+ * EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO,
+ * PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR
+ * PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY
+ * OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
+ * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
+ * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+ */
+#include "particle.h"
+
+particle::particle() : 	position(), velocity(), totalDistance(1,0,0)
+{
+}
+
+__device__ __host__
+void particle::advance(float d)
+{
+  velocity.normalize();
+  float dx = d * velocity.x;
+  position.x += dx;
+  totalDistance.x += dx;
+  float dy = d * velocity.y;
+  position.y += dy;
+  totalDistance.y += dy;
+  float dz = d * velocity.z;
+  position.z += dz;
+  totalDistance.z += dz;
+// #if __CUDA_ARCH__
+//   int idx = threadIdx.x + blockIdx.x*blockDim.x;
+//   if(idx == 0)
+//   {
+//     printf("totalDistance: %f\n", totalDistance.x );
+//   }
+// #endif
+  velocity.scramble();
+}
+
+const v3& particle::getTotalDistance() const
+{	return totalDistance; }
diff --git a/posts/cmake/particle.h b/posts/cmake/particle.h
@@ -0,0 +1,46 @@
+/* Copyright (c) 1993-2015, NVIDIA CORPORATION. All rights reserved.
+ *
+ * Redistribution and use in source and binary forms, with or without
+ * modification, are permitted provided that the following conditions
+ * are met:
+ *  * Redistributions of source code must retain the above copyright
+ *    notice, this list of conditions and the following disclaimer.
+ *  * Redistributions in binary form must reproduce the above copyright
+ *    notice, this list of conditions and the following disclaimer in the
+ *    documentation and/or other materials provided with the distribution.
+ *  * Neither the name of NVIDIA CORPORATION nor the names of its
+ *    contributors may be used to endorse or promote products derived
+ *    from this software without specific prior written permission.
+ *
+ * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS ``AS IS'' AND ANY
+ * EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+ * IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR
+ * PURPOSE ARE DISCLAIMED.  IN NO EVENT SHALL THE COPYRIGHT OWNER OR
+ * CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL,
+ * EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO,
+ * PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR
+ * PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY
+ * OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
+ * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
+ * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+ */
+#ifndef __particle_h__
+#define __particle_h__
+
+#include "v3.h"
+
+class particle
+{
+	private:
+		v3 position;
+		v3 velocity;
+		v3 totalDistance;
+
+	public:
+		particle();
+		__host__ __device__ void advance(float dist);
+		const v3& getTotalDistance() const;
+
+};
+
+#endif
diff --git a/posts/cmake/test.cu b/posts/cmake/test.cu
@@ -0,0 +1,103 @@
+/* Copyright (c) 1993-2015, NVIDIA CORPORATION. All rights reserved.
+ *
+ * Redistribution and use in source and binary forms, with or without
+ * modification, are permitted provided that the following conditions
+ * are met:
+ *  * Redistributions of source code must retain the above copyright
+ *    notice, this list of conditions and the following disclaimer.
+ *  * Redistributions in binary form must reproduce the above copyright
+ *    notice, this list of conditions and the following disclaimer in the
+ *    documentation and/or other materials provided with the distribution.
+ *  * Neither the name of NVIDIA CORPORATION nor the names of its
+ *    contributors may be used to endorse or promote products derived
+ *    from this software without specific prior written permission.
+ *
+ * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS ``AS IS'' AND ANY
+ * EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+ * IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR
+ * PURPOSE ARE DISCLAIMED.  IN NO EVENT SHALL THE COPYRIGHT OWNER OR
+ * CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL,
+ * EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO,
+ * PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR
+ * PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY
+ * OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
+ * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
+ * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+ */
+#include "particle.h"
+#include <stdlib.h>
+#include <stdio.h>
+
+__global__ void advanceParticles(float dt, particle * pArray, int nParticles)
+{
+	int idx = threadIdx.x + blockIdx.x*blockDim.x;
+	if(idx < nParticles)
+	{
+		pArray[idx].advance(dt);
+	}
+}
+
+int main(int argc, char ** argv)
+{
+	cudaError_t error;
+	int n = 1000000;
+	if(argc > 1)	{ n = atoi(argv[1]);}     // Number of particles
+	if(argc > 2)	{	srand(atoi(argv[2])); } // Random seed
+
+	error = cudaGetLastError();
+	if (error != cudaSuccess)
+  	{
+  	printf("0 %s\n",cudaGetErrorString(error));
+  	exit(1);
+  	}
+
+	particle * pArray = new particle[n];
+	particle * devPArray = NULL;
+	cudaMalloc(&devPArray, n*sizeof(particle));
+	cudaDeviceSynchronize(); error = cudaGetLastError();
+	if (error != cudaSuccess)
+  	{
+  	printf("1 %s\n",cudaGetErrorString(error));
+  	exit(1);
+  	}
+
+	cudaMemcpy(devPArray, pArray, n*sizeof(particle), cudaMemcpyHostToDevice);
+	cudaDeviceSynchronize(); error = cudaGetLastError();
+	if (error != cudaSuccess)
+  	{
+  	printf("2 %s\n",cudaGetErrorString(error));
+  	exit(1);
+  	}
+
+	for(int i=0; i<100; i++)
+	{
+		float dt = (float)rand()/(float) RAND_MAX; // Random distance each step
+		advanceParticles<<< 1 +  n/256, 256>>>(dt, devPArray, n);
+		error = cudaGetLastError();
+		if (error != cudaSuccess)
+    	{
+    	printf("3 %s\n",cudaGetErrorString(error));
+    	exit(1);
+    	}
+
+		cudaDeviceSynchronize();
+	}
+	cudaMemcpy(pArray, devPArray, n*sizeof(particle), cudaMemcpyDeviceToHost);
+
+	v3 totalDistance(0,0,0);
+	v3 temp;
+	for(int i=0; i<n; i++)
+	{
+		temp = pArray[i].getTotalDistance();
+		totalDistance.x += temp.x;
+		totalDistance.y += temp.y;
+		totalDistance.z += temp.z;
+	}
+	float avgX = totalDistance.x /(float)n;
+	float avgY = totalDistance.y /(float)n;
+	float avgZ = totalDistance.z /(float)n;
+	float avgNorm = sqrt(avgX*avgX + avgY*avgY + avgZ*avgZ);
+	printf(	"Moved %d particles 100 steps. Average distance traveled is |(%f, %f, %f)| = %f\n",
+					n, avgX, avgY, avgZ, avgNorm);
+	return 0;
+}
diff --git a/posts/cmake/v3.cu b/posts/cmake/v3.cu
@@ -0,0 +1,59 @@
+/* Copyright (c) 1993-2015, NVIDIA CORPORATION. All rights reserved.
+ *
+ * Redistribution and use in source and binary forms, with or without
+ * modification, are permitted provided that the following conditions
+ * are met:
+ *  * Redistributions of source code must retain the above copyright
+ *    notice, this list of conditions and the following disclaimer.
+ *  * Redistributions in binary form must reproduce the above copyright
+ *    notice, this list of conditions and the following disclaimer in the
+ *    documentation and/or other materials provided with the distribution.
+ *  * Neither the name of NVIDIA CORPORATION nor the names of its
+ *    contributors may be used to endorse or promote products derived
+ *    from this software without specific prior written permission.
+ *
+ * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS ``AS IS'' AND ANY
+ * EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+ * IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR
+ * PURPOSE ARE DISCLAIMED.  IN NO EVENT SHALL THE COPYRIGHT OWNER OR
+ * CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL,
+ * EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO,
+ * PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR
+ * PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY
+ * OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
+ * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
+ * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+ */
+#include "v3.h"
+#include <math.h>
+
+v3::v3()
+{	randomize(); }
+
+v3::v3(float xIn, float yIn, float zIn) : x(xIn), y(yIn), z(zIn)
+{}
+
+void v3::randomize()
+{
+	x = (float)rand() / (float)RAND_MAX;
+	y = (float)rand() / (float)RAND_MAX;
+	z = (float)rand() / (float)RAND_MAX;
+}
+
+__host__ __device__ void v3::normalize()
+{
+	float t = sqrt(x*x + y*y + z*z);
+	x /= t;
+	y /= t;
+	z /= t;
+}
+
+__host__ __device__ void v3::scramble()
+{
+	float tx = 0.317f*(x + 1.0) + y + z * x * x + y + z;
+	float ty = 0.619f*(y + 1.0) + y * y + x * y * z + y + x;
+	float tz = 0.124f*(z + 1.0) + z * y + x * y * z + y + x;
+	x = tx;
+	y = ty;
+	z = tz;
+}
diff --git a/posts/cmake/v3.h b/posts/cmake/v3.h
@@ -0,0 +1,45 @@
+/* Copyright (c) 1993-2015, NVIDIA CORPORATION. All rights reserved.
+ *
+ * Redistribution and use in source and binary forms, with or without
+ * modification, are permitted provided that the following conditions
+ * are met:
+ *  * Redistributions of source code must retain the above copyright
+ *    notice, this list of conditions and the following disclaimer.
+ *  * Redistributions in binary form must reproduce the above copyright
+ *    notice, this list of conditions and the following disclaimer in the
+ *    documentation and/or other materials provided with the distribution.
+ *  * Neither the name of NVIDIA CORPORATION nor the names of its
+ *    contributors may be used to endorse or promote products derived
+ *    from this software without specific prior written permission.
+ *
+ * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS ``AS IS'' AND ANY
+ * EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+ * IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR
+ * PURPOSE ARE DISCLAIMED.  IN NO EVENT SHALL THE COPYRIGHT OWNER OR
+ * CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL,
+ * EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO,
+ * PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR
+ * PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY
+ * OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
+ * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
+ * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+ */
+#ifndef __v3_h__
+#define __v3_h__
+
+class v3
+{
+public:
+	float x;
+	float y;
+	float z;
+
+	v3();
+	v3(float xIn, float yIn, float zIn);
+	void randomize();
+	__host__ __device__ void normalize();
+	__host__ __device__ void scramble();
+
+};
+
+#endif