feat(warmup): add solutions

Student_Cuda_Warmup/src/core/01_addArray/00_procedurale/addArrayProcedurale.cu

@@ -9,6 +9,7 @@
 #include "cudas.h"
 #include "GM.h"
 #include "Kernel.h"
+#include "Hardware.h"
 
 using std::cout;
 using std::endl;
@@ -32,9 +33,34 @@ static __global__ void kaddArray(float* ptrGMV1 , float* ptrGMV2 , float* ptrGMW
  * ptrW receptionne le resultat
  * n nombre de case
  */
-__host__ bool addArray_procedurale(float* ptrV1 , float* ptrV2 , float* ptrW , int n) // __host__ facultatif
+__host__ void addArray_procedurale(float* ptrV1 , float* ptrV2 , float* ptrW , int n) // __host__ facultatif
     {
-    // TODO addArray
+    size_t sizeVector = sizeof(float) * n;
+
+    float* ptrGMV1;
+    float* ptrGMV2;
+    float* ptrGMW;
+
+    GM::malloc(&ptrGMV1, sizeVector);
+    GM::malloc(&ptrGMV2, sizeVector);
+    GM::malloc(&ptrGMW, sizeVector);
+
+    GM::memcpyHToD(ptrGMV1, ptrV1, sizeVector);
+    GM::memcpyHToD(ptrGMV2, ptrV2, sizeVector);
+
+    const int MP = Hardware::getMPCount();
+    const int CORE_MP = Hardware::getCoreCountMP();
+
+    dim3 dg(MP, 3, 1);
+    dim3 db(CORE_MP, 2, 1); // <=1024
+
+    kaddArray<<<dg,db>>>( ptrGMV1 , ptrGMV2 , ptrGMW , n);
+
+    GM::memcpyDToH(ptrW, ptrGMW, sizeVector);
+
+    GM::free(ptrGMV1);
+    GM::free(ptrGMV2);
+    GM::free(ptrGMW);
     }
 
 /*--------------------------------------*\
@@ -50,8 +76,12 @@ __global__ void kaddArray(float* ptrGMV1 , float* ptrGMV2 , float* ptrGMW , int
     const int TID = Thread2D::tid();
 
     // pattern entrelacement
-
+    int s = TID;
-    // TODO addArray
+    while (s < n)
+	{
+	ptrGMW[s] = ptrGMV1[s] + ptrGMV2[s];
+	s += NB_THREAD;
+	}
     }
 
 /*----------------------------------------------------------------------*\

Student_Cuda_Warmup/src/core/01_addArray/00_procedurale/useAddArrayProcedurale.cpp

@@ -37,16 +37,17 @@ bool exemple_addArray_procedurale()
 	ArrayTools::print(ptrV2, n);
 	cout << "--------------------------------------------------------------------" << endl;
 	ArrayTools::print(ptrW, n);
+	cout<<endl;
 	}
 
     // Delete
 	{
 	delete ptrV1;
-	delete ptrV1;
+	delete ptrV2;
 	delete ptrW;
 	}
 
-    return isOk;
+    return true;
     }
 
 /*----------------------------------------------------------------------*\

Student_Cuda_Warmup/src/core/01_addArray/01_poo/01_pattern_entrelacement/device/addArray_device.cu

@@ -11,23 +11,18 @@
 /**
  * output : void required, because kernel is asynchrone
  */
-__global__ void addArray(float* ptrGMV1 , float* ptrGMV2 , float* ptrGMW , int n)
+__global__ void addArray(float* ptrGMV1 , float* ptrGMV2 , float* ptrGMW , int n) {
-    {
     const int NB_THREAD = Thread2D::nbThread();
     const int TID = Thread2D::tid();
 
-    // Debug, facultatif
-//    if (TID == 0)
-//	{
-//	printf("Coucou from device tid = %d", TID);
-//	}
-
     // pattern entrelacement
-
+    int s = TID;
-    // TODO addArray
+    while (s < n) {
+    	ptrGMW[s] = ptrGMV1[s] + ptrGMV2[s];
+    	s += NB_THREAD;
+	}
 }
 
 /*----------------------------------------------------------------------*\
  |*			End	 					*|
  \*---------------------------------------------------------------------*/
-

Student_Cuda_Warmup/src/core/01_addArray/01_poo/01_pattern_entrelacement/host/AddArray.cu

@@ -12,7 +12,6 @@ using std::endl;
 using std::to_string;
 using std::string;
 
-
 /*--------------------------------------*\
  |*		Imported	 	*|
  \*-------------------------------------*/
@@ -33,25 +32,23 @@ AddArray::AddArray(const Grid& grid , float* ptrV1 , float* ptrV2 , float* ptrW
 	ptrW(ptrW), //
 	n(n), //
 	dg(grid.dg), //
-	db(grid.db)
+	db(grid.db) {
-    {
+        this->sizeVector = sizeof(float) * n;
-    this->sizeVector = -1; // TODO addArray // octet
 
         // MM (malloc Device)
     	{
            	GM::malloc(&ptrGMV1, sizeVector);
-	// TODO addArray
+           	GM::malloc(&ptrGMV2, sizeVector);
+           	GM::malloc(&ptrGMW, sizeVector);
     	}
 
     }
 
-AddArray::~AddArray()
+AddArray::~AddArray() {
-    {
     //MM (device free)
-	{
 	GM::free(ptrGMV1);
-	// TODO addArray
+	GM::free(ptrGMV2);
-	}
+	GM::free(ptrGMW);
 }
 
 /*--------------------------------------*\
@@ -61,21 +58,21 @@ AddArray::~AddArray()
 /**
  * override
  */
-void AddArray::run()
+void AddArray::run() {
-    {
     // MM (copy Host->Device)
     {
     	GM::memcpyHToD(ptrGMV1, ptrV1, sizeVector);
-	// TODO addArray
+    	GM::memcpyHToD(ptrGMV2, ptrV2, sizeVector);
 	}
 
-    // TODO addArray // call kernel // assynchrone
+
+    addArray<<<dg,db>>>( ptrGMV1 , ptrGMV2 , ptrGMW , n);
 
     //Kernel::synchronize();// inutile
 
     // MM (Device -> Host)
 	{
-	// TODO addArray // MM barier de synchronisation implicite
+	    GM::memcpyDToH(ptrW, ptrGMW, sizeVector);
 	}
 }
 

Student_Cuda_Warmup/src/core/01_addArray/01_poo/01_pattern_entrelacement/host/AddArray.h

@@ -9,8 +9,7 @@
  \*---------------------------------------------------------------------*/
 
 
-class AddArray
+class AddArray {
-    {
 	/*--------------------------------------*\
 	|*		Constructor		*|
 	 \*-------------------------------------*/

Student_Cuda_Warmup/src/core/01_addArray/01_poo/01_pattern_entrelacement/useAddArrayObject.cpp

@@ -45,12 +45,13 @@ bool exemple_addArray_object()
 	ArrayTools::print(ptrV2, n);
 	cout << "--------------------------------------------------------------------" << endl;
 	ArrayTools::print(ptrW, n);
+	cout<<endl;
 	}
 
     // Delete
 	{
 	delete ptrV1;
-	delete ptrV1;
+	delete ptrV2;
 	delete ptrW;
 	}
 
@@ -69,16 +70,10 @@ static Grid createGrid()
     const int MP = Hardware::getMPCount();
     const int CORE_MP = Hardware::getCoreCountMP();
 
-    dim3 dg(1, 1, 1); // TODO addArray
+    dim3 dg(MP, 2, 1);
-    dim3 db(1, 1, 1); // TODO addArray  // produit <=1024
+    dim3 db(CORE_MP, 3, 1);   // produit <=1024
     Grid grid(dg, db);
 
-    // to remove once coded
-	{
-	Couts::redln("aie aie aie, your best grid won t build itself");
-	assert(false);
-	}
-
     return grid;
     }
 

Student_Cuda_Warmup/src/core/01_addArray/01_poo/02_pattern_11/useAddArray11.cpp

@@ -45,12 +45,13 @@ bool exemple_addArray_11()
 	ArrayTools::print(ptrV2, n);
 	cout << "--------------------------------------------------------------------" << endl;
 	ArrayTools::print(ptrW, n);
+	cout<<endl;
 	}
 
     // Delete
 	{
 	delete ptrV1;
-	delete ptrV1;
+	delete ptrV2;
 	delete ptrW;
 	}
 

Student_Cuda_Warmup/src/main/main.cpp

@@ -2,6 +2,7 @@
 #include <stdlib.h>
 
 #include "Couts.h"
+#include "Hardware.h"
 
 using std::cerr;
 using std::cout;
@@ -29,18 +30,30 @@ static void scalar(bool& isOk);
 
 int main(int argc , char** argv)
     {
+    // Facultatif
+	{
+	Hardware::loadCudaDriverAll();
+	Hardware::printAllSimple();
+
+	const int DEVICE_ID = 0;
+	Hardware::setDevice(DEVICE_ID);
+	}
+
     const int IS_VERBOSE = true;
 
+    // run
+	{
 	bool isOk = true;
 
 	// Commenter ce dont vous n'avez pas besoin ci-dessous
-    scalar(isOk);	 // commenter dans la methode ci-dessous ce que vous ne voulez pas lancer
+	//scalar(isOk);	 // commenter dans la methode ci-dessous ce que vous ne voulez pas lancer
 	array(isOk);	// commenter dans la methode ci-dessous ce que vous ne voulez pas lancer
 
 	Couts::statusln(isOk);
 
 	return isOk ? EXIT_SUCCESS : EXIT_FAILURE;
 	}
+    }
 
 /*-------------------------------------*\
  |*		Private 		*|
@@ -56,7 +69,7 @@ static void array(bool& isOk)
     {
     isOk &= exemple_addArray_procedurale();
     isOk &= exemple_addArray_object();
-    isOk &= exemple_addArray_11();
+    // isOk &= exemple_addArray_11();
     }
 
 /*----------------------------------------------------------------------*\