Klagarge/MSE-GPU-Cuda
1
0
Fork 0
Code Issues Pull Requests Releases Activity

Compare commits

base: Klagarge:2daed2f9eb59d0b9f9637800ab16954abaae0f00
Branches Tags
Klagarge:main
...
compare: Klagarge:main
Branches Tags
Klagarge:main

10 Commits
2daed2f9eb ... main

Author SHA1 Message Date
Klagarge
bda49a20c9 chore: delivery preparation 2025-11-27 18:49:46 +01:00
Klagarge
09ee32383f feat(Mandelbrot): add TP mandelbrot 2025-11-27 18:41:30 +01:00
Klagarge
a892100308 doc(sliceSM): add output for sliceSM lab 2025-11-24 23:22:33 +01:00
Klagarge
bf9a0a2941 feat(sliceSM): add sliceSM 2025-11-24 23:21:59 +01:00
Klagarge
cb6856ccde feat(reduction): add reduction generic PII long 2025-11-17 19:13:08 +01:00
Klagarge
1ec4ef198e feat(reduction): add reduction generic PII int 2025-11-17 18:46:08 +01:00
Klagarge
78fb87e978 feat/reduction): add reduction generic PI int 2025-11-17 18:34:58 +01:00
Klagarge
135f425fd7 feat(reduction): add reduction PII add 2025-11-17 18:00:54 +01:00
Klagarge
dbaa9440b4 chore: add task and clangd 2025-11-15 18:03:49 +01:00
Klagarge
1543a0f121 feat(reduction): add reduction PI add 2025-11-15 18:03:11 +01:00
76 changed files with 2886 additions and 1690 deletions
Expand all files Collapse all files

15
.zed/tasks.json Normal file
View File

@@ -0,0 +1,15 @@
[
{
"label": "Run CUDA - Student_Cuda",
"command": "cbicc cuda clean jall run",
"cwd": "Student_Cuda",
"use_new_terminal": false,
"allow_concurrent_runs": false,
"reveal": "always",
"reveal_target": "dock",
"hide": "never",
"shell": "system",
"show_summary": true,
"show_command": true
}
]

10
Student_Cuda/.clangd Normal file
View File

@@ -0,0 +1,10 @@
CompileFlags:
Add:
- "-I/home/mse15/CUDA/toStudent/code/WCudaStudent/Student_Cuda/INC_SYMLINK/EXT"
- "-I/home/mse15/CUDA/toStudent/code/WCudaStudent/Student_Cuda/INC_SYMLINK/PROJECT"
- "-std=c++17"
- "-x"
- "cuda"
---
Diagnostics:
Suppress: "*"

14
Student_Cuda/.zed/tasks.json Normal file
View File

@@ -0,0 +1,14 @@
[
{
"label": "Run CUDA",
"command": "cbicc cuda clean jall run",
"use_new_terminal": false,
"allow_concurrent_runs": false,
"reveal": "always",
"reveal_target": "dock",
"hide": "never",
"shell": "system",
"show_summary": true,
"show_command": true
}
]

6
Student_Cuda/out/SliceGMHOST_justesse.html
View File

@@ -87,7 +87,7 @@ Designed by <a href="https://github.com/cpptest/cpptest">CppTest</a>
<td style="width:30%" class="tablecell_success">0</td> <td style="width:30%" class="tablecell_success">0</td>
<td style="width:30%" class="tablecell_success">0</td> <td style="width:30%" class="tablecell_success">0</td>
<td style="width:30%" class="tablecell_success">100%</td> <td style="width:30%" class="tablecell_success">100%</td>
<td style="width:10%" class="tablecell_success">4.000000</td> <td style="width:10%" class="tablecell_success">3.000000</td>
</tr> </tr>
</table> </table>
<hr /> <hr />
@@ -106,7 +106,7 @@ Designed by <a href="https://github.com/cpptest/cpptest">CppTest</a>
<td style="width:10%" class="tablecell_success">14</td> <td style="width:10%" class="tablecell_success">14</td>
<td style="width:10%" class="tablecell_success">0</td> <td style="width:10%" class="tablecell_success">0</td>
<td style="width:10%" class="tablecell_success">100%</td> <td style="width:10%" class="tablecell_success">100%</td>
<td style="width:10%" class="tablecell_success">4.000000</td> <td style="width:10%" class="tablecell_success">3.000000</td>
</tr> </tr>
</table> </table>
<hr /> <hr />
@@ -189,7 +189,7 @@ Designed by <a href="https://github.com/cpptest/cpptest">CppTest</a>
<td class="tablecell_success">testMonoBlock</td> <td class="tablecell_success">testMonoBlock</td>
<td class="tablecell_success">0</td> <td class="tablecell_success">0</td>
<td class="tablecell_success">true</td> <td class="tablecell_success">true</td>
<td class="tablecell_success">3.000000</td> <td class="tablecell_success">2.000000</td>
</tr> </tr>
<tr> <tr>
<td class="tablecell_success">testMonoThread</td> <td class="tablecell_success">testMonoThread</td>

6
Student_Cuda/out/SliceGMHOST_performance.html
View File

@@ -87,7 +87,7 @@ Designed by <a href="https://github.com/cpptest/cpptest">CppTest</a>
<td style="width:30%" class="tablecell_success">1</td> <td style="width:30%" class="tablecell_success">1</td>
<td style="width:30%" class="tablecell_success">0</td> <td style="width:30%" class="tablecell_success">0</td>
<td style="width:30%" class="tablecell_success">100%</td> <td style="width:30%" class="tablecell_success">100%</td>
<td style="width:10%" class="tablecell_success">11.000000</td> <td style="width:10%" class="tablecell_success">10.000000</td>
</tr> </tr>
</table> </table>
<hr /> <hr />
@@ -106,7 +106,7 @@ Designed by <a href="https://github.com/cpptest/cpptest">CppTest</a>
<td style="width:10%" class="tablecell_success">1</td> <td style="width:10%" class="tablecell_success">1</td>
<td style="width:10%" class="tablecell_success">0</td> <td style="width:10%" class="tablecell_success">0</td>
<td style="width:10%" class="tablecell_success">100%</td> <td style="width:10%" class="tablecell_success">100%</td>
<td style="width:10%" class="tablecell_success">11.000000</td> <td style="width:10%" class="tablecell_success">10.000000</td>
</tr> </tr>
</table> </table>
<hr /> <hr />
@@ -123,7 +123,7 @@ Designed by <a href="https://github.com/cpptest/cpptest">CppTest</a>
<td class="tablecell_success">performanceOnly</td> <td class="tablecell_success">performanceOnly</td>
<td class="tablecell_success">0</td> <td class="tablecell_success">0</td>
<td class="tablecell_success">true</td> <td class="tablecell_success">true</td>
<td class="tablecell_success">11.000000</td> <td class="tablecell_success">10.000000</td>
</tr> </tr>
</table> </table>
<p class="spaced"><a href="#top">Back to top</a> <p class="spaced"><a href="#top">Back to top</a>

4
Student_Cuda/out/SliceGM_justesse.html
View File

@@ -153,7 +153,7 @@ Designed by <a href="https://github.com/cpptest/cpptest">CppTest</a>
<td class="tablecell_success">testDB64</td> <td class="tablecell_success">testDB64</td>
<td class="tablecell_success">0</td> <td class="tablecell_success">0</td>
<td class="tablecell_success">true</td> <td class="tablecell_success">true</td>
<td class="tablecell_success">0.000000</td> <td class="tablecell_success">1.000000</td>
</tr> </tr>
<tr> <tr>
<td class="tablecell_success">testDB128</td> <td class="tablecell_success">testDB128</td>
@@ -183,7 +183,7 @@ Designed by <a href="https://github.com/cpptest/cpptest">CppTest</a>
<td class="tablecell_success">testGrid</td> <td class="tablecell_success">testGrid</td>
<td class="tablecell_success">0</td> <td class="tablecell_success">0</td>
<td class="tablecell_success">true</td> <td class="tablecell_success">true</td>
<td class="tablecell_success">1.000000</td> <td class="tablecell_success">0.000000</td>
</tr> </tr>
<tr> <tr>
<td class="tablecell_success">testMonoBlock</td> <td class="tablecell_success">testMonoBlock</td>

6
Student_Cuda/out/SliceGM_performance.html
View File

@@ -87,7 +87,7 @@ Designed by <a href="https://github.com/cpptest/cpptest">CppTest</a>
<td style="width:30%" class="tablecell_success">1</td> <td style="width:30%" class="tablecell_success">1</td>
<td style="width:30%" class="tablecell_success">0</td> <td style="width:30%" class="tablecell_success">0</td>
<td style="width:30%" class="tablecell_success">100%</td> <td style="width:30%" class="tablecell_success">100%</td>
<td style="width:10%" class="tablecell_success">10.000000</td> <td style="width:10%" class="tablecell_success">11.000000</td>
</tr> </tr>
</table> </table>
<hr /> <hr />
@@ -106,7 +106,7 @@ Designed by <a href="https://github.com/cpptest/cpptest">CppTest</a>
<td style="width:10%" class="tablecell_success">1</td> <td style="width:10%" class="tablecell_success">1</td>
<td style="width:10%" class="tablecell_success">0</td> <td style="width:10%" class="tablecell_success">0</td>
<td style="width:10%" class="tablecell_success">100%</td> <td style="width:10%" class="tablecell_success">100%</td>
<td style="width:10%" class="tablecell_success">10.000000</td> <td style="width:10%" class="tablecell_success">11.000000</td>
</tr> </tr>
</table> </table>
<hr /> <hr />
@@ -123,7 +123,7 @@ Designed by <a href="https://github.com/cpptest/cpptest">CppTest</a>
<td class="tablecell_success">performanceOnly</td> <td class="tablecell_success">performanceOnly</td>
<td class="tablecell_success">0</td> <td class="tablecell_success">0</td>
<td class="tablecell_success">true</td> <td class="tablecell_success">true</td>
<td class="tablecell_success">10.000000</td> <td class="tablecell_success">11.000000</td>
</tr> </tr>
</table> </table>
<p class="spaced"><a href="#top">Back to top</a> <p class="spaced"><a href="#top">Back to top</a>

212
Student_Cuda/out/SliceSM_justesse.html Normal file
View File

@@ -0,0 +1,212 @@
<!DOCTYPE html PUBLIC "-//W3C//DTD XHTML 1.0 Strict//EN" "http://www.w3.org/TR/xhtml1/DTD/xhtml1-strict.dtd">
<html xmlns="http://www.w3.org/1999/xhtml" xml:lang="en" lang="en">
<head>
<meta http-equiv="content-type" content="text/html; charset=utf-8" />
<meta name="generator" content="CppTest - https://github.com/cpptest/cpptest" />
<title>SliceSM_justesse </title>
<style type="text/css" media="screen">
<!--
hr {
width: 100%;
border-width: 0px;
height: 1px;
color: #cccccc;
background-color: #cccccc;
padding: 0px;
}
table {
width:100%;
border-collapse:separate;
border-spacing: 2px;
border:0px;
}
tr {
margin:0px;
padding:0px;
}
td {
margin:0px;
padding:1px;
}
.table_summary {
}
.table_suites {
}
.table_suite {
}
.table_result {
margin: 0px 0px 1em 0px;
}
.tablecell_title {
background-color: #a5cef7;
font-weight: bold;
}
.tablecell_success {
background-color: #efefe7;
}
.tablecell_error {
color: #ff0808;
background-color: #efefe7;
font-weight: bold;
}
p.spaced {
margin: 0px;
padding: 1em 0px 2em 0px;
}
p.unspaced {
margin: 0px;
padding: 0px 0px 2em 0px;
}
-->
</style>
</head>
<body>
<h1><a name="top"></a>SliceSM_justesse </h1>
<div style="text-align:right">
Designed by <a href="https://github.com/cpptest/cpptest">CppTest</a>
</div>
<hr />
<h2>Summary</h2>
<table summary="Summary of test results" class="table_summary">
<tr>
<td style="width:30%" class="tablecell_title">Tests</td>
<td style="width:30%" class="tablecell_title">Errors</td>
<td style="width:30%" class="tablecell_title">Success</td>
<td style="width:10%" class="tablecell_title">Time (s)</td>
</tr>
<tr>
<td style="width:30%" class="tablecell_success">0</td>
<td style="width:30%" class="tablecell_success">0</td>
<td style="width:30%" class="tablecell_success">100%</td>
<td style="width:10%" class="tablecell_success">0.000000</td>
</tr>
</table>
<hr />
<h2>Test suites</h2>
<table summary="Test Suites" class="table_suites">
<tr>
<td class="tablecell_title">Name</td>
<td style="width:10%" class="tablecell_title">Tests</td>
<td style="width:10%" class="tablecell_title">Errors</td>
<td style="width:10%" class="tablecell_title">Success</td>
<td style="width:10%" class="tablecell_title">Time (s)</td>
</tr>
<tr>
<td class="tablecell_success"><a href="#TestSliceSM">TestSliceSM</a></td>
<td style="width:10%" class="tablecell_success">13</td>
<td style="width:10%" class="tablecell_success">0</td>
<td style="width:10%" class="tablecell_success">100%</td>
<td style="width:10%" class="tablecell_success">0.000000</td>
</tr>
</table>
<hr />
<h3><a name="TestSliceSM"></a>Suite: TestSliceSM</h3>
<table summary="Details for suite TestSliceSM" class="table_suite">
<tr>
<td class="tablecell_title">Name</td>
<td style="width:10%" class="tablecell_title">Errors</td>
<td style="width:10%" class="tablecell_title">Success</td>
<td style="width:10%" class="tablecell_title">Time (s)</td>
</tr>
<tr>
<td class="tablecell_success">testDB2</td>
<td class="tablecell_success">0</td>
<td class="tablecell_success">true</td>
<td class="tablecell_success">0.000000</td>
</tr>
<tr>
<td class="tablecell_success">testDB4</td>
<td class="tablecell_success">0</td>
<td class="tablecell_success">true</td>
<td class="tablecell_success">0.000000</td>
</tr>
<tr>
<td class="tablecell_success">testDB8</td>
<td class="tablecell_success">0</td>
<td class="tablecell_success">true</td>
<td class="tablecell_success">0.000000</td>
</tr>
<tr>
<td class="tablecell_success">testDB16</td>
<td class="tablecell_success">0</td>
<td class="tablecell_success">true</td>
<td class="tablecell_success">0.000000</td>
</tr>
<tr>
<td class="tablecell_success">testDB32</td>
<td class="tablecell_success">0</td>
<td class="tablecell_success">true</td>
<td class="tablecell_success">0.000000</td>
</tr>
<tr>
<td class="tablecell_success">testDB64</td>
<td class="tablecell_success">0</td>
<td class="tablecell_success">true</td>
<td class="tablecell_success">0.000000</td>
</tr>
<tr>
<td class="tablecell_success">testDB128</td>
<td class="tablecell_success">0</td>
<td class="tablecell_success">true</td>
<td class="tablecell_success">0.000000</td>
</tr>
<tr>
<td class="tablecell_success">testDB256</td>
<td class="tablecell_success">0</td>
<td class="tablecell_success">true</td>
<td class="tablecell_success">0.000000</td>
</tr>
<tr>
<td class="tablecell_success">testDB512</td>
<td class="tablecell_success">0</td>
<td class="tablecell_success">true</td>
<td class="tablecell_success">0.000000</td>
</tr>
<tr>
<td class="tablecell_success">testDB1024</td>
<td class="tablecell_success">0</td>
<td class="tablecell_success">true</td>
<td class="tablecell_success">0.000000</td>
</tr>
<tr>
<td class="tablecell_success">testGrid</td>
<td class="tablecell_success">0</td>
<td class="tablecell_success">true</td>
<td class="tablecell_success">0.000000</td>
</tr>
<tr>
<td class="tablecell_success">testMonoBlock</td>
<td class="tablecell_success">0</td>
<td class="tablecell_success">true</td>
<td class="tablecell_success">0.000000</td>
</tr>
<tr>
<td class="tablecell_success">testBestGrid</td>
<td class="tablecell_success">0</td>
<td class="tablecell_success">true</td>
<td class="tablecell_success">0.000000</td>
</tr>
</table>
<p class="spaced"><a href="#top">Back to top</a>
</p>
<hr />
<p>
<a href="http://validator.w3.org/#validate-by-upload">
Valid XHTML 1.0 Strict
</a>
</p>
</body>
</html>

140
Student_Cuda/out/SliceSM_performance.html Normal file
View File

@@ -0,0 +1,140 @@
<!DOCTYPE html PUBLIC "-//W3C//DTD XHTML 1.0 Strict//EN" "http://www.w3.org/TR/xhtml1/DTD/xhtml1-strict.dtd">
<html xmlns="http://www.w3.org/1999/xhtml" xml:lang="en" lang="en">
<head>
<meta http-equiv="content-type" content="text/html; charset=utf-8" />
<meta name="generator" content="CppTest - https://github.com/cpptest/cpptest" />
<title>SliceSM_performance </title>
<style type="text/css" media="screen">
<!--
hr {
width: 100%;
border-width: 0px;
height: 1px;
color: #cccccc;
background-color: #cccccc;
padding: 0px;
}
table {
width:100%;
border-collapse:separate;
border-spacing: 2px;
border:0px;
}
tr {
margin:0px;
padding:0px;
}
td {
margin:0px;
padding:1px;
}
.table_summary {
}
.table_suites {
}
.table_suite {
}
.table_result {
margin: 0px 0px 1em 0px;
}
.tablecell_title {
background-color: #a5cef7;
font-weight: bold;
}
.tablecell_success {
background-color: #efefe7;
}
.tablecell_error {
color: #ff0808;
background-color: #efefe7;
font-weight: bold;
}
p.spaced {
margin: 0px;
padding: 1em 0px 2em 0px;
}
p.unspaced {
margin: 0px;
padding: 0px 0px 2em 0px;
}
-->
</style>
</head>
<body>
<h1><a name="top"></a>SliceSM_performance </h1>
<div style="text-align:right">
Designed by <a href="https://github.com/cpptest/cpptest">CppTest</a>
</div>
<hr />
<h2>Summary</h2>
<table summary="Summary of test results" class="table_summary">
<tr>
<td style="width:30%" class="tablecell_title">Tests</td>
<td style="width:30%" class="tablecell_title">Errors</td>
<td style="width:30%" class="tablecell_title">Success</td>
<td style="width:10%" class="tablecell_title">Time (s)</td>
</tr>
<tr>
<td style="width:30%" class="tablecell_success">1</td>
<td style="width:30%" class="tablecell_success">0</td>
<td style="width:30%" class="tablecell_success">100%</td>
<td style="width:10%" class="tablecell_success">11.000000</td>
</tr>
</table>
<hr />
<h2>Test suites</h2>
<table summary="Test Suites" class="table_suites">
<tr>
<td class="tablecell_title">Name</td>
<td style="width:10%" class="tablecell_title">Tests</td>
<td style="width:10%" class="tablecell_title">Errors</td>
<td style="width:10%" class="tablecell_title">Success</td>
<td style="width:10%" class="tablecell_title">Time (s)</td>
</tr>
<tr>
<td class="tablecell_success"><a href="#TestPerformance_RunnableGPU_A">TestPerformance_RunnableGPU_A</a></td>
<td style="width:10%" class="tablecell_success">1</td>
<td style="width:10%" class="tablecell_success">0</td>
<td style="width:10%" class="tablecell_success">100%</td>
<td style="width:10%" class="tablecell_success">11.000000</td>
</tr>
</table>
<hr />
<h3><a name="TestPerformance_RunnableGPU_A"></a>Suite: TestPerformance_RunnableGPU_A</h3>
<table summary="Details for suite TestPerformance_RunnableGPU_A" class="table_suite">
<tr>
<td class="tablecell_title">Name</td>
<td style="width:10%" class="tablecell_title">Errors</td>
<td style="width:10%" class="tablecell_title">Success</td>
<td style="width:10%" class="tablecell_title">Time (s)</td>
</tr>
<tr>
<td class="tablecell_success">performanceOnly</td>
<td class="tablecell_success">0</td>
<td class="tablecell_success">true</td>
<td class="tablecell_success">11.000000</td>
</tr>
</table>
<p class="spaced"><a href="#top">Back to top</a>
</p>
<hr />
<p>
<a href="http://validator.w3.org/#validate-by-upload">
Valid XHTML 1.0 Strict
</a>
</p>
</body>
</html>

10
Student_Cuda/out/bruteforce/SliceSM_21474836_db.csv Normal file
View File

@@ -0,0 +1,10 @@
64,128,256,512,1024
64,128,256,512,1024
64,128,256,512,1024
64,128,256,512,1024
64,128,256,512,1024
64,128,256,512,1024
64,128,256,512,1024
64,128,256,512,1024
64,128,256,512,1024
64,128,256,512,1024
1 64 128 256 512 1024
2 64 128 256 512 1024
3 64 128 256 512 1024
4 64 128 256 512 1024
5 64 128 256 512 1024
6 64 128 256 512 1024
7 64 128 256 512 1024
8 64 128 256 512 1024
9 64 128 256 512 1024
10 64 128 256 512 1024

10
Student_Cuda/out/bruteforce/SliceSM_21474836_dg.csv Normal file
View File

@@ -0,0 +1,10 @@
68,68,68,68,68
136,136,136,136,136
204,204,204,204,204
272,272,272,272,272
340,340,340,340,340
408,408,408,408,408
476,476,476,476,476
544,544,544,544,544
612,612,612,612,612
680,680,680,680,680
1 68 68 68 68 68
2 136 136 136 136 136
3 204 204 204 204 204
4 272 272 272 272 272
5 340 340 340 340 340
6 408 408 408 408 408
7 476 476 476 476 476
8 544 544 544 544 544
9 612 612 612 612 612
10 680 680 680 680 680

10
Student_Cuda/out/bruteforce/SliceSM_21474836_fps.csv Normal file
View File

@@ -0,0 +1,10 @@
14872,23559,29077,29179,29245
23644,29250,29265,29249,29235
23620,29205,29208,29217,23873
29202,29184,29182,29183,23568
23586,29205,29282,29279,23595
29261,29312,29262,29297,23584
29235,29140,29239,29173,19773
29160,29123,29176,29187,19713
29179,29213,29187,29235,19736
26261,29069,29066,29051,19642
1 14872 23559 29077 29179 29245
2 23644 29250 29265 29249 29235
3 23620 29205 29208 29217 23873
4 29202 29184 29182 29183 23568
5 23586 29205 29282 29279 23595
6 29261 29312 29262 29297 23584
7 29235 29140 29239 29173 19773
8 29160 29123 29176 29187 19713
9 29179 29213 29187 29235 19736
10 26261 29069 29066 29051 19642

39
Student_Cuda/src/core/01_student/03_Slice/03_Slice_SM/device/sliceSM_device.cu
View File

@@ -17,16 +17,21 @@ static __device__ float f(float x);
|* Implementation *| |* Implementation *|
\*---------------------------------------------------------------------*/ \*---------------------------------------------------------------------*/
__global__ void sliceSM(int nbSlice , float* ptrPiHatGM) __global__
{ void sliceSM(int nbSlice , float* ptrPiHatGM) {
// TODO SliceSM // TODO SliceSM
// Reception tabSM // Reception tabSM
extern __shared__ float tabSM[];
// ReductionIntraThread // ReductionIntraThread
reductionIntraThread(tabSM, nbSlice);
__syncthreads();
// Reduction de tabSM (use tools ReductionAdd) // Reduction de tabSM (use tools ReductionAdd)
ReductionAdd::reduce(tabSM, ptrPiHatGM);
// __syncthreads(); necessaire? ou? pas a la fin en tout cas // __syncthreads(); necessaire? ou? pas a la fin en tout cas
} }
/*--------------------------------------*\ /*--------------------------------------*\
|* Private *| |* Private *|
@@ -35,19 +40,33 @@ __global__ void sliceSM(int nbSlice , float* ptrPiHatGM)
/** /**
* remplit la sm * remplit la sm
*/ */
void reductionIntraThread(float* tabSM , int nbSlice) static
{ __device__
void reductionIntraThread(float* tabSM , int nbSlice) {
// TODO SliceSM // TODO SliceSM
// Warning: Il faut employer TID et TID_LOCAL // Warning: Il faut employer TID et TID_LOCAL
} const int TID = Thread2D::tid();
const int localTID = Thread2D::tidLocal();
const int NB_THREAD = Thread2D::nbThread();
__device__ float f(float x) const float delta_x = 1.f / (float)nbSlice;
{
return 4.f / (1.f + x * x); int s = TID;
tabSM[localTID] = 0.f;
while (s < nbSlice) {
float xi = s * delta_x;
tabSM[localTID] += f(xi);
s += NB_THREAD;
} }
}
__device__
float f(float x) {
return 4.f / (1.f + x * x);
}
/*----------------------------------------------------------------------*\ /*----------------------------------------------------------------------*\
|* End *| |* End *|
\*---------------------------------------------------------------------*/ \*---------------------------------------------------------------------*/

30
Student_Cuda/src/core/01_student/03_Slice/03_Slice_SM/host/SliceSM.cu
View File

@@ -25,34 +25,31 @@ extern __global__ void sliceSM(int nbSlice,float* ptrPiHatGM);
\*-------------------------------------*/ \*-------------------------------------*/
SliceSM::SliceSM(const Grid& grid , int nbSlice , double* ptrPiHat , bool isVerbose) : SliceSM::SliceSM(const Grid& grid , int nbSlice , double* ptrPiHat , bool isVerbose) :
RunnableGPU(grid, "SliceSM_" + to_string(nbSlice), isVerbose), // classe parente RunnableGPU(grid, "SliceSM_" + to_string(nbSlice), isVerbose),
// ptrPiHat(ptrPiHat),
ptrPiHat(ptrPiHat), // nbSlice(nbSlice) {
nbSlice(nbSlice) // this->sizeSM = grid.threadByBlock() * sizeof(float); //TODO SliceSM
{
this->sizeSM = -1; //TODO SliceSM
// MM // MM
{ {
// TODO SliceSM (pas oublier de mettre a zero, avec mallocfloat0 par exemple) // TODO SliceSM (pas oublier de mettre a zero, avec mallocfloat0 par exemple)
GM::mallocFloat0(&ptrPiHatGM);
// Tip: Il y a une methode dedier pour malloquer un float cote device et l'initialiser a zero // Tip: Il y a une methode dedier pour malloquer un float cote device et l'initialiser a zero
// //
// GM::mallocfloat0(&ptrPiHatGM); // GM::mallocfloat0(&ptrPiHatGM);
} }
} }
SliceSM::~SliceSM(void) SliceSM::~SliceSM(void) {
{
//TODO SliceSM //TODO SliceSM
} GM::free(ptrPiHatGM);
}
/*--------------------------------------*\ /*--------------------------------------*\
|* Methode *| |* Methode *|
\*-------------------------------------*/ \*-------------------------------------*/
void SliceSM::run() void SliceSM::run() {
{
// Etape 1 : lancer le kernel // Etape 1 : lancer le kernel
// Etape 2 : recuperer le resultat coter host (par exemple avec memcpyDToH_float) // Etape 2 : recuperer le resultat coter host (par exemple avec memcpyDToH_float)
// Etape 3 : finaliser le calcul de PI // Etape 3 : finaliser le calcul de PI
@@ -60,7 +57,12 @@ void SliceSM::run()
// Solution : double result; // et ramener dans result, transferer et finaliser ensuite dans ptrPiHat // Solution : double result; // et ramener dans result, transferer et finaliser ensuite dans ptrPiHat
// TODO SliceSM // TODO SliceSM
} sliceSM<<<dg,db,this->sizeSM>>>(this->nbSlice, this->ptrPiHatGM);
float result;
GM::memcpyDToH_float(&result, this->ptrPiHatGM);
const double delta_x = 1.0 / (double) this->nbSlice;
*this->ptrPiHat = (double) result * delta_x;
}
///////////////////////// /////////////////////////
// Rappel: // Rappel:

5
Student_Cuda/src/core/01_student/03_Slice/03_Slice_SM/host/SliceSM.h
View File

@@ -8,8 +8,7 @@
|* Declaration *| |* Declaration *|
\*---------------------------------------------------------------------*/ \*---------------------------------------------------------------------*/
class SliceSM: public RunnableGPU class SliceSM: public RunnableGPU {
{
/*--------------------------------------*\ /*--------------------------------------*\
|* Constructor *| |* Constructor *|
\*-------------------------------------*/ \*-------------------------------------*/
@@ -51,7 +50,7 @@ class SliceSM: public RunnableGPU
size_t sizeSM; // [octet] size_t sizeSM; // [octet]
float* ptrPiHatGM; float* ptrPiHatGM;
}; };
/*----------------------------------------------------------------------*\ /*----------------------------------------------------------------------*\
|* End *| |* End *|

25
Student_Cuda/src/core/01_student/03_Slice/03_Slice_SM/host/SliceSM_BestGrid.h
View File

@@ -11,29 +11,26 @@
|* Impelmentation *| |* Impelmentation *|
\*---------------------------------------------------------------------*/ \*---------------------------------------------------------------------*/
namespace sliceSM namespace sliceSM {
{
class BestGrid class BestGrid {
{
public: public:
static Grid get() static Grid get() {
{ const int MP = Hardware::getMPCount();
const int MP = Hardware::getMPCount(); const int CORE_MP = Hardware::getCoreCountMP();
// TODO SliceGMHOST grid // TODO SliceGMHOST grid
dim3 dg(MP, 6, 1);
dim3 db(CORE_MP, 2, 1);
Grid grid(dg, db);
// to remove once coded return grid;
{
Couts::redln("aie aie aie, your best grid won t build itself");
assert(false);
}
} }
}; };
} }
/*----------------------------------------------------------------------*\ /*----------------------------------------------------------------------*\
|* End *| |* End *|

2
Student_Cuda/src/main/main.cpp
View File

@@ -29,7 +29,7 @@ int main(int argc , char** argv)
// public // public
{ {
cudaContext.deviceId = 1; // in [0,2] width Server Cuda3 cudaContext.deviceId = 1; // in [0,2] width Server Cuda3
cudaContext.launchMode = LaunchModeMOO::USE; // USE TEST BENCHMARK FORCEBRUT cudaContext.launchMode = LaunchModeMOO::TEST; // USE TEST BENCHMARK FORCEBRUT
cudaContext.deviceDriver = DeviceDriver::LOAD_ALL; // LOAD_CURRENT LOAD_ALL cudaContext.deviceDriver = DeviceDriver::LOAD_ALL; // LOAD_CURRENT LOAD_ALL
cudaContext.deviceInfo = DeviceInfo::ALL_SIMPLE; // NONE ALL ALL_SIMPLE CURRENT cudaContext.deviceInfo = DeviceInfo::ALL_SIMPLE; // NONE ALL ALL_SIMPLE CURRENT

5
Student_Cuda/src/main/mainBenchmark.cpp
View File

@@ -34,6 +34,7 @@ using std::endl;
static void sliceGMHOST(); static void sliceGMHOST();
static void sliceGM(); static void sliceGM();
static void sliceSM();
static void montecarloMono(); static void montecarloMono();
@@ -59,8 +60,8 @@ int mainBenchmark()
// Slice // Slice
{ {
// sliceGMHOST(); // sliceGMHOST();
sliceGM(); // sliceGM();
//sliceSM(); sliceSM();
//sliceMulti(); //sliceMulti();
} }

4
Student_Cuda/src/main/mainBruteforce.cpp
View File

@@ -75,8 +75,8 @@ int mainBrutforce()
// Slice // Slice
{ {
// sliceGMHOST(&matlab); // sliceGMHOST(&matlab);
sliceGM(&matlab); // sliceGM(&matlab);
// sliceSM(&matlab); sliceSM(&matlab);
} }
// Montecarlo // Montecarlo

15
Student_Cuda/src/main/mainTest.cpp
View File

@@ -52,17 +52,16 @@ int mainTest()
/** /**
* activer ci-dessous la version souhaiter * activer ci-dessous la version souhaiter
*/ */
void slice() void slice() {
{ VTSliceGMHOST test1;
// VTSliceGMHOST test1;
VTSliceGM test2; VTSliceGM test2;
// VTSliceSM test3; VTSliceSM test3;
// test1.run(); test1.run();
test2.run(); test2.run();
// test3.run(); test3.run();
} }
/** /**
* activer ci-dessous la version souhaiter * activer ci-dessous la version souhaiter

10
Student_Cuda/src/main/mainUse.cpp
View File

@@ -66,13 +66,13 @@ int mainUse()
*/ */
void slice(bool& isOk) void slice(bool& isOk)
{ {
// SliceGmHostUse sliceGmHostUse(IS_VERBOSE); SliceGmHostUse sliceGmHostUse(IS_VERBOSE);
SliceGmUse sliceGmUse(IS_VERBOSE); SliceGmUse sliceGmUse(IS_VERBOSE);
// SliceSmUse sliceSmUse(IS_VERBOSE); SliceSmUse sliceSmUse(IS_VERBOSE);
// isOk &= sliceGmHostUse.isOk(IS_VERBOSE); isOk &= sliceGmHostUse.isOk(IS_VERBOSE);
isOk &= sliceGmUse.isOk(IS_VERBOSE); isOk &= sliceGmUse.isOk(IS_VERBOSE);
// isOk &= sliceSmUse.isOk(IS_VERBOSE); isOk &= sliceSmUse.isOk(IS_VERBOSE);
} }
/** /**

140
Student_Cuda_Image/out/Mandelbrot_fp16_justesse.html Normal file
View File

@@ -0,0 +1,140 @@
<!DOCTYPE html PUBLIC "-//W3C//DTD XHTML 1.0 Strict//EN" "http://www.w3.org/TR/xhtml1/DTD/xhtml1-strict.dtd">
<html xmlns="http://www.w3.org/1999/xhtml" xml:lang="en" lang="en">
<head>
<meta http-equiv="content-type" content="text/html; charset=utf-8" />
<meta name="generator" content="CppTest - https://github.com/cpptest/cpptest" />
<title>Mandelbrot_fp16_justesse </title>
<style type="text/css" media="screen">
<!--
hr {
width: 100%;
border-width: 0px;
height: 1px;
color: #cccccc;
background-color: #cccccc;
padding: 0px;
}
table {
width:100%;
border-collapse:separate;
border-spacing: 2px;
border:0px;
}
tr {
margin:0px;
padding:0px;
}
td {
margin:0px;
padding:1px;
}
.table_summary {
}
.table_suites {
}
.table_suite {
}
.table_result {
margin: 0px 0px 1em 0px;
}
.tablecell_title {
background-color: #a5cef7;
font-weight: bold;
}
.tablecell_success {
background-color: #efefe7;
}
.tablecell_error {
color: #ff0808;
background-color: #efefe7;
font-weight: bold;
}
p.spaced {
margin: 0px;
padding: 1em 0px 2em 0px;
}
p.unspaced {
margin: 0px;
padding: 0px 0px 2em 0px;
}
-->
</style>
</head>
<body>
<h1><a name="top"></a>Mandelbrot_fp16_justesse </h1>
<div style="text-align:right">
Designed by <a href="https://github.com/cpptest/cpptest">CppTest</a>
</div>
<hr />
<h2>Summary</h2>
<table summary="Summary of test results" class="table_summary">
<tr>
<td style="width:30%" class="tablecell_title">Tests</td>
<td style="width:30%" class="tablecell_title">Errors</td>
<td style="width:30%" class="tablecell_title">Success</td>
<td style="width:10%" class="tablecell_title">Time (s)</td>
</tr>
<tr>
<td style="width:30%" class="tablecell_success">1</td>
<td style="width:30%" class="tablecell_success">0</td>
<td style="width:30%" class="tablecell_success">100%</td>
<td style="width:10%" class="tablecell_success">2.000000</td>
</tr>
</table>
<hr />
<h2>Test suites</h2>
<table summary="Test Suites" class="table_suites">
<tr>
<td class="tablecell_title">Name</td>
<td style="width:10%" class="tablecell_title">Tests</td>
<td style="width:10%" class="tablecell_title">Errors</td>
<td style="width:10%" class="tablecell_title">Success</td>
<td style="width:10%" class="tablecell_title">Time (s)</td>
</tr>
<tr>
<td class="tablecell_success"><a href="#TestImageCuda">TestImageCuda</a></td>
<td style="width:10%" class="tablecell_success">1</td>
<td style="width:10%" class="tablecell_success">0</td>
<td style="width:10%" class="tablecell_success">100%</td>
<td style="width:10%" class="tablecell_success">2.000000</td>
</tr>
</table>
<hr />
<h3><a name="TestImageCuda"></a>Suite: TestImageCuda</h3>
<table summary="Details for suite TestImageCuda" class="table_suite">
<tr>
<td class="tablecell_title">Name</td>
<td style="width:10%" class="tablecell_title">Errors</td>
<td style="width:10%" class="tablecell_title">Success</td>
<td style="width:10%" class="tablecell_title">Time (s)</td>
</tr>
<tr>
<td class="tablecell_success">allTests</td>
<td class="tablecell_success">0</td>
<td class="tablecell_success">true</td>
<td class="tablecell_success">2.000000</td>
</tr>
</table>
<p class="spaced"><a href="#top">Back to top</a>
</p>
<hr />
<p>
<a href="http://validator.w3.org/#validate-by-upload">
Valid XHTML 1.0 Strict
</a>
</p>
</body>
</html>

140
Student_Cuda_Image/out/Mandelbrot_fp16_performance.html Normal file
View File

@@ -0,0 +1,140 @@
<!DOCTYPE html PUBLIC "-//W3C//DTD XHTML 1.0 Strict//EN" "http://www.w3.org/TR/xhtml1/DTD/xhtml1-strict.dtd">
<html xmlns="http://www.w3.org/1999/xhtml" xml:lang="en" lang="en">
<head>
<meta http-equiv="content-type" content="text/html; charset=utf-8" />
<meta name="generator" content="CppTest - https://github.com/cpptest/cpptest" />
<title>Mandelbrot_fp16_performance </title>
<style type="text/css" media="screen">
<!--
hr {
width: 100%;
border-width: 0px;
height: 1px;
color: #cccccc;
background-color: #cccccc;
padding: 0px;
}
table {
width:100%;
border-collapse:separate;
border-spacing: 2px;
border:0px;
}
tr {
margin:0px;
padding:0px;
}
td {
margin:0px;
padding:1px;
}
.table_summary {
}
.table_suites {
}
.table_suite {
}
.table_result {
margin: 0px 0px 1em 0px;
}
.tablecell_title {
background-color: #a5cef7;
font-weight: bold;
}
.tablecell_success {
background-color: #efefe7;
}
.tablecell_error {
color: #ff0808;
background-color: #efefe7;
font-weight: bold;
}
p.spaced {
margin: 0px;
padding: 1em 0px 2em 0px;
}
p.unspaced {
margin: 0px;
padding: 0px 0px 2em 0px;
}
-->
</style>
</head>
<body>
<h1><a name="top"></a>Mandelbrot_fp16_performance </h1>
<div style="text-align:right">
Designed by <a href="https://github.com/cpptest/cpptest">CppTest</a>
</div>
<hr />
<h2>Summary</h2>
<table summary="Summary of test results" class="table_summary">
<tr>
<td style="width:30%" class="tablecell_title">Tests</td>
<td style="width:30%" class="tablecell_title">Errors</td>
<td style="width:30%" class="tablecell_title">Success</td>
<td style="width:10%" class="tablecell_title">Time (s)</td>
</tr>
<tr>
<td style="width:30%" class="tablecell_success">1</td>
<td style="width:30%" class="tablecell_success">0</td>
<td style="width:30%" class="tablecell_success">100%</td>
<td style="width:10%" class="tablecell_success">12.000000</td>
</tr>
</table>
<hr />
<h2>Test suites</h2>
<table summary="Test Suites" class="table_suites">
<tr>
<td class="tablecell_title">Name</td>
<td style="width:10%" class="tablecell_title">Tests</td>
<td style="width:10%" class="tablecell_title">Errors</td>
<td style="width:10%" class="tablecell_title">Success</td>
<td style="width:10%" class="tablecell_title">Time (s)</td>
</tr>
<tr>
<td class="tablecell_success"><a href="#TestPerformance_RunnableGPU_A">TestPerformance_RunnableGPU_A</a></td>
<td style="width:10%" class="tablecell_success">1</td>
<td style="width:10%" class="tablecell_success">0</td>
<td style="width:10%" class="tablecell_success">100%</td>
<td style="width:10%" class="tablecell_success">12.000000</td>
</tr>
</table>
<hr />
<h3><a name="TestPerformance_RunnableGPU_A"></a>Suite: TestPerformance_RunnableGPU_A</h3>
<table summary="Details for suite TestPerformance_RunnableGPU_A" class="table_suite">
<tr>
<td class="tablecell_title">Name</td>
<td style="width:10%" class="tablecell_title">Errors</td>
<td style="width:10%" class="tablecell_title">Success</td>
<td style="width:10%" class="tablecell_title">Time (s)</td>
</tr>
<tr>
<td class="tablecell_success">performanceOnly</td>
<td class="tablecell_success">0</td>
<td class="tablecell_success">true</td>
<td class="tablecell_success">12.000000</td>
</tr>
</table>
<p class="spaced"><a href="#top">Back to top</a>
</p>
<hr />
<p>
<a href="http://validator.w3.org/#validate-by-upload">
Valid XHTML 1.0 Strict
</a>
</p>
</body>
</html>

6
Student_Cuda_Image/out/Rippling_performance.html
View File

@@ -87,7 +87,7 @@ Designed by <a href="https://github.com/cpptest/cpptest">CppTest</a>
<td style="width:30%" class="tablecell_success">1</td> <td style="width:30%" class="tablecell_success">1</td>
<td style="width:30%" class="tablecell_success">0</td> <td style="width:30%" class="tablecell_success">0</td>
<td style="width:30%" class="tablecell_success">100%</td> <td style="width:30%" class="tablecell_success">100%</td>
<td style="width:10%" class="tablecell_success">8.000000</td> <td style="width:10%" class="tablecell_success">16.000000</td>
</tr> </tr>
</table> </table>
<hr /> <hr />
@@ -106,7 +106,7 @@ Designed by <a href="https://github.com/cpptest/cpptest">CppTest</a>
<td style="width:10%" class="tablecell_success">1</td> <td style="width:10%" class="tablecell_success">1</td>
<td style="width:10%" class="tablecell_success">0</td> <td style="width:10%" class="tablecell_success">0</td>
<td style="width:10%" class="tablecell_success">100%</td> <td style="width:10%" class="tablecell_success">100%</td>
<td style="width:10%" class="tablecell_success">8.000000</td> <td style="width:10%" class="tablecell_success">16.000000</td>
</tr> </tr>
</table> </table>
<hr /> <hr />
@@ -123,7 +123,7 @@ Designed by <a href="https://github.com/cpptest/cpptest">CppTest</a>
<td class="tablecell_success">performanceOnly</td> <td class="tablecell_success">performanceOnly</td>
<td class="tablecell_success">0</td> <td class="tablecell_success">0</td>
<td class="tablecell_success">true</td> <td class="tablecell_success">true</td>
<td class="tablecell_success">8.000000</td> <td class="tablecell_success">16.000000</td>
</tr> </tr>
</table> </table>
<p class="spaced"><a href="#top">Back to top</a> <p class="spaced"><a href="#top">Back to top</a>

10
Student_Cuda_Image/out/bruteforce/Mandelbrot-Cuda-uchar4-fp16-nMin20-nMax120_db.csv Normal file
View File

@@ -0,0 +1,10 @@
64,128,192,256,320,384,448,512,576,640,704,768,832,896,960,1024
64,128,192,256,320,384,448,512,576,640,704,768,832,896,960,1024
64,128,192,256,320,384,448,512,576,640,704,768,832,896,960,1024
64,128,192,256,320,384,448,512,576,640,704,768,832,896,960,1024
64,128,192,256,320,384,448,512,576,640,704,768,832,896,960,1024
64,128,192,256,320,384,448,512,576,640,704,768,832,896,960,1024
64,128,192,256,320,384,448,512,576,640,704,768,832,896,960,1024
64,128,192,256,320,384,448,512,576,640,704,768,832,896,960,1024
64,128,192,256,320,384,448,512,576,640,704,768,832,896,960,1024
64,128,192,256,320,384,448,512,576,640,704,768,832,896,960,1024
1 64 128 192 256 320 384 448 512 576 640 704 768 832 896 960 1024
2 64 128 192 256 320 384 448 512 576 640 704 768 832 896 960 1024
3 64 128 192 256 320 384 448 512 576 640 704 768 832 896 960 1024
4 64 128 192 256 320 384 448 512 576 640 704 768 832 896 960 1024
5 64 128 192 256 320 384 448 512 576 640 704 768 832 896 960 1024
6 64 128 192 256 320 384 448 512 576 640 704 768 832 896 960 1024
7 64 128 192 256 320 384 448 512 576 640 704 768 832 896 960 1024
8 64 128 192 256 320 384 448 512 576 640 704 768 832 896 960 1024
9 64 128 192 256 320 384 448 512 576 640 704 768 832 896 960 1024
10 64 128 192 256 320 384 448 512 576 640 704 768 832 896 960 1024

10
Student_Cuda_Image/out/bruteforce/Mandelbrot-Cuda-uchar4-fp16-nMin20-nMax120_dg.csv Normal file
View File

@@ -0,0 +1,10 @@
68,68,68,68,68,68,68,68,68,68,68,68,68,68,68,68
136,136,136,136,136,136,136,136,136,136,136,136,136,136,136,136
204,204,204,204,204,204,204,204,204,204,204,204,204,204,204,204
272,272,272,272,272,272,272,272,272,272,272,272,272,272,272,272
340,340,340,340,340,340,340,340,340,340,340,340,340,340,340,340
408,408,408,408,408,408,408,408,408,408,408,408,408,408,408,408
476,476,476,476,476,476,476,476,476,476,476,476,476,476,476,476
544,544,544,544,544,544,544,544,544,544,544,544,544,544,544,544
612,612,612,612,612,612,612,612,612,612,612,612,612,612,612,612
680,680,680,680,680,680,680,680,680,680,680,680,680,680,680,680
1 68 68 68 68 68 68 68 68 68 68 68 68 68 68 68 68
2 136 136 136 136 136 136 136 136 136 136 136 136 136 136 136 136
3 204 204 204 204 204 204 204 204 204 204 204 204 204 204 204 204
4 272 272 272 272 272 272 272 272 272 272 272 272 272 272 272 272
5 340 340 340 340 340 340 340 340 340 340 340 340 340 340 340 340
6 408 408 408 408 408 408 408 408 408 408 408 408 408 408 408 408
7 476 476 476 476 476 476 476 476 476 476 476 476 476 476 476 476
8 544 544 544 544 544 544 544 544 544 544 544 544 544 544 544 544
9 612 612 612 612 612 612 612 612 612 612 612 612 612 612 612 612
10 680 680 680 680 680 680 680 680 680 680 680 680 680 680 680 680

10
Student_Cuda_Image/out/bruteforce/Mandelbrot-Cuda-uchar4-fp16-nMin20-nMax120_fps.csv Normal file
View File

@@ -0,0 +1,10 @@
11261,21977,27704,28654,28277,28547,28629,28404,28310,28099,28086,28083,27716,27519,27791,28011
19038,28186,28672,29112,28426,28273,28146,28144,27978,27624,27759,27741,27325,27066,27341,27544
27172,28378,27808,28260,28006,28038,27917,27990,27496,27172,27343,27313,26925,26740,26918,27046
27677,28301,27408,28080,27452,27837,27162,27780,27046,26611,26934,26825,26609,26218,26593,26529
27585,27504,27616,27705,27863,27826,27102,27073,26633,26380,26509,26507,26361,25949,26339,26248
26661,27570,27783,28103,27632,27613,27589,27606,26069,25936,26088,26036,26011,25575,25894,25830
27632,27463,27391,27489,27242,27606,27584,26880,25818,25521,25840,25693,25640,25316,25482,25309
27657,27392,27302,27627,27038,27244,26594,27163,25196,24997,25035,25023,24712,24685,24668,24540
25842,24673,26280,26057,26492,25605,26130,25341,24431,23305,24874,23801,24877,23471,24577,23293
23284,25709,25062,26164,25511,27316,25176,24768,23921,22530,23529,22545,23692,22471,24333,22815
1 11261 21977 27704 28654 28277 28547 28629 28404 28310 28099 28086 28083 27716 27519 27791 28011
2 19038 28186 28672 29112 28426 28273 28146 28144 27978 27624 27759 27741 27325 27066 27341 27544
3 27172 28378 27808 28260 28006 28038 27917 27990 27496 27172 27343 27313 26925 26740 26918 27046
4 27677 28301 27408 28080 27452 27837 27162 27780 27046 26611 26934 26825 26609 26218 26593 26529
5 27585 27504 27616 27705 27863 27826 27102 27073 26633 26380 26509 26507 26361 25949 26339 26248
6 26661 27570 27783 28103 27632 27613 27589 27606 26069 25936 26088 26036 26011 25575 25894 25830
7 27632 27463 27391 27489 27242 27606 27584 26880 25818 25521 25840 25693 25640 25316 25482 25309
8 27657 27392 27302 27627 27038 27244 26594 27163 25196 24997 25035 25023 24712 24685 24668 24540
9 25842 24673 26280 26057 26492 25605 26130 25341 24431 23305 24874 23801 24877 23471 24577 23293
10 23284 25709 25062 26164 25511 27316 25176 24768 23921 22530 23529 22545 23692 22471 24333 22815

14
Student_Cuda_Image/src/core/01_student/01_Rippling/host/Rippling_BestGrid.h
View File

@@ -18,15 +18,15 @@ namespace rippling {
public: public:
static Grid get() { static Grid get() {
const int MP = Hardware::getMPCount(); const int MP = Hardware::getMPCount();
const int CORE_MP = Hardware::getCoreCountMP(); const int CORE_MP = Hardware::getCoreCountMP();
// TODO Rippling // TODO Rippling
dim3 dg(MP, 1, 1); dim3 dg(MP, 1, 1);
dim3 db(CORE_MP, 4, 1); dim3 db(CORE_MP, 4, 1);
Grid grid(dg, db); Grid grid(dg, db);
return grid; return grid;
} }

40
Student_Cuda_Image/src/core/01_student/02_Mandelbrot/device/mandelbrotDevice.cu
View File

@@ -20,19 +20,36 @@ static __device__ void color(uchar4* ptrColor,int k, int n);
|* Implementation *| |* Implementation *|
\*---------------------------------------------------------------------*/ \*---------------------------------------------------------------------*/
__global__ void mandelbrot(uchar4* tabPixelsGM , uint w , uint h , DomaineMath domaineMath , int n) __global__
{ void mandelbrot(uchar4* tabPixelsGM , uint w , uint h , DomaineMath domaineMath , int n) {
// TODO Mandelbrot : // TODO Mandelbrot :
// //
// entrelacement // entrelacement
// s -> (i,j) -> (x,y) // s -> (i,j) -> (x,y)
// appeler colorXY // appeler colorXY
double x; double x = 0;
double y; double y = 0;
const int TID = Thread2D::tid();
const int NB_THREAD = Thread2D::nbThread();
const int WH = w * h;
MandelbrotMath mm(n);
int s = TID;
while ( s < WH) {
int i = s / w;
int j = s % w;
domaineMath.toXY(i, j, &x, &y);
mm.colorXY(&tabPixelsGM[s], (real)x, (real)y);
s += NB_THREAD;
}
//domaineMath.toXY(i, j, &x, &y); // x et y doivent etre en double! Caster ensuite en real lors du passage à colorXY //domaineMath.toXY(i, j, &x, &y); // x et y doivent etre en double! Caster ensuite en real lors du passage à colorXY
} }
/*----------------------------------------------------------------------*\ /*----------------------------------------------------------------------*\
|* private *| |* private *|
@@ -42,8 +59,8 @@ __global__ void mandelbrot(uchar4* tabPixelsGM , uint w , uint h , DomaineMath d
* optimisation lookup table color (facultatif) * optimisation lookup table color (facultatif)
*/ */
__inline__ __inline__
__device__ void fill(uchar4* tabSM,int n) __device__
{ void fill(uchar4* tabSM,int n) {
// Indications: // Indications:
// (I1) tabSM a n cases // (I1) tabSM a n cases
// (I2) La case k contient les couleurs en RVBA lorsque la suite s est arreter a k // (I2) La case k contient les couleurs en RVBA lorsque la suite s est arreter a k
@@ -54,18 +71,17 @@ __device__ void fill(uchar4* tabSM,int n)
// Warning // Warning
// (W1) Commencer d'abord sans cette piste d'optimisation // (W1) Commencer d'abord sans cette piste d'optimisation
} }
/** /**
* optimisation lookup table color (facultatif) * optimisation lookup table color (facultatif)
*/ */
__inline__ __inline__
__device__ void color(uchar4* ptrColor,int k, int n) __device__
{ void color(uchar4* ptrColor,int k, int n) {
} }
/*----------------------------------------------------------------------*\ /*----------------------------------------------------------------------*\
|* End *| |* End *|
\*---------------------------------------------------------------------*/ \*---------------------------------------------------------------------*/

91
Student_Cuda_Image/src/core/01_student/02_Mandelbrot/device/math/MandelbrotMath.cu.h
View File

@@ -33,25 +33,24 @@ using mandelbrotReal::real;
#define ZERO 0.0 #define ZERO 0.0
#endif #endif
class MandelbrotMath class MandelbrotMath {
{
/*--------------------------------------*\ /*--------------------------------------*\
|* Constructeur *| |* Constructeur *|
\*-------------------------------------*/ \*-------------------------------------*/
public: public:
__device__ MandelbrotMath(int n) : // __device__
n(n) MandelbrotMath(int n) : //
{ n(n) {
// rien // rien
} }
__device__ __device__
virtual ~MandelbrotMath() virtual
{ ~MandelbrotMath() {
// rien // rien
} }
/*--------------------------------------*\ /*--------------------------------------*\
|* Methode *| |* Methode *|
@@ -60,8 +59,7 @@ class MandelbrotMath
public: public:
__device__ __device__
void colorXY(uchar4* ptrColorIJ , real x , real y) void colorXY(uchar4* ptrColorIJ , real x , real y) {
{
// TODO Mandelbrot // TODO Mandelbrot
// Calculer la suite en (x,y) et recuperer l'indice d'arret de la suite // Calculer la suite en (x,y) et recuperer l'indice d'arret de la suite
@@ -75,12 +73,30 @@ class MandelbrotMath
// Etape 2: Puis une fois que l'image grise est valider, attaquer montecarlo // Etape 2: Puis une fois que l'image grise est valider, attaquer montecarlo
// debug temp // debug temp
// { // {
// ptrColorIJ->x = 128; // ptrColorIJ->x = 128;
// ptrColorIJ->y = 128; // ptrColorIJ->y = 128;
// ptrColorIJ->z = 128; // ptrColorIJ->z = 128;
// ptrColorIJ->w = 255; // opacity facultatif // ptrColorIJ->w = 255; // opacity facultatif
} // }
int itr = suite(x, y);
if (itr >= n) {
ptrColorIJ->x = 0;
ptrColorIJ->y = 0;
ptrColorIJ->z = 0;
ptrColorIJ->w = 255; // opacity facultatif
} else {
Colors::HSB_TO_RVB(
( (float)itr ) / ( (float)(this->n-1) ),
1.f,
1.f,
ptrColorIJ
);
}
}
private: private:
@@ -90,20 +106,35 @@ class MandelbrotMath
*/ */
__inline__ __inline__
__device__ __device__
int suite(real x , real y) int suite(real x , real y) { // TODO Mandelbrot
{ // TODO Mandelbrot
// Utiliser dans vos formules les variable : // Utiliser dans vos formules les variable :
// //
// DEUX // DEUX
// QUATRE // QUATRE
// ZERO // ZERO
// //
// definit au debut de ce fichier. Est-utile pour passer facilement d'une version fp64 (double) fp32(float) fp16(half) // definit au debut de ce fichier. Est-utile pour passer facilement d'une version fp64 (double) fp32(float) fp16(half)
// Calculer la suite en (x,y) jusqu'à n, à moins que critere arret soit atteint avant // Calculer la suite en (x,y) jusqu'à n, à moins que critere arret soit atteint avant
// return le nombre d'element de la suite calculer, ie un entier // return le nombre d'element de la suite calculer, ie un entier
}
real a = ZERO;
real b = ZERO;
real tmp = ZERO;
int itr = -1;
while ( (a*a + b*b <= QUATRE) && (itr < n) ) {
tmp = a;
a = x + (a*a - b*b);
b = DEUX * tmp * b + y;
itr++;
}
return itr;
}
///////////////////// /////////////////////
// Warning: // Warning:
@@ -150,7 +181,7 @@ class MandelbrotMath
// Inputs // Inputs
int n; int n;
}; };
/*----------------------------------------------------------------------*\ /*----------------------------------------------------------------------*\
|* End *| |* End *|

23
Student_Cuda_Image/src/core/01_student/02_Mandelbrot/host/Mandelbrot.cu
View File

@@ -40,10 +40,9 @@ Mandelbrot::Mandelbrot(const Grid& grid , uint w , uint h , const DomaineMath& d
// Tools // Tools
this->t = nMin; this->t = nMin;
} }
Mandelbrot::~Mandelbrot() Mandelbrot::~Mandelbrot() {
{
// rien // rien
} }
@@ -55,28 +54,24 @@ Mandelbrot::~Mandelbrot()
* Override * Override
* Call periodicly by the API * Call periodicly by the API
*/ */
void Mandelbrot::process(uchar4* tabPixelsGM , uint w , uint h , const DomaineMath& domaineMath) void Mandelbrot::process(uchar4* tabPixelsGM , uint w , uint h , const DomaineMath& domaineMath) {
{
assert(false); // to be removed once implemented
// TODO Mandelbrot // TODO Mandelbrot
// lauch kernel (you find at line 18) // lauch kernel (you find at line 18)
} mandelbrot<<<dg,db>>>(tabPixelsGM,w,h,domaineMath,t);
}
/** /**
* Override * Override
* Call periodicly by the API * Call periodicly by the API
*/ */
void Mandelbrot::animationStep() void Mandelbrot::animationStep() {
{
this->t = variateurT.varierAndGet(); this->t = variateurT.varierAndGet();
} }
string titre(int nMin , int nMax) string titre(int nMin , int nMax) {
{
return "Mandelbrot-Cuda-uchar4-" + mandelbrotReal::realToString() + "-nMin" + std::to_string(nMin) + "-nMax" + std::to_string(nMax); return "Mandelbrot-Cuda-uchar4-" + mandelbrotReal::realToString() + "-nMin" + std::to_string(nMin) + "-nMax" + std::to_string(nMax);
} }
/*----------------------------------------------------------------------*\ /*----------------------------------------------------------------------*\
|* End *| |* End *|
\*---------------------------------------------------------------------*/ \*---------------------------------------------------------------------*/

5
Student_Cuda_Image/src/core/01_student/02_Mandelbrot/host/Mandelbrot.h
View File

@@ -10,8 +10,7 @@
|* Declaration *| |* Declaration *|
\*---------------------------------------------------------------------*/ \*---------------------------------------------------------------------*/
class Mandelbrot: public Animable_I<uchar4> class Mandelbrot: public Animable_I<uchar4> {
{
/*--------------------------------------*\ /*--------------------------------------*\
|* Constructor *| |* Constructor *|
\*-------------------------------------*/ \*-------------------------------------*/
@@ -53,7 +52,7 @@ class Mandelbrot: public Animable_I<uchar4>
// Tools // Tools
Variateur<int> variateurT; Variateur<int> variateurT;
}; };
/*----------------------------------------------------------------------*\ /*----------------------------------------------------------------------*\
|* End *| |* End *|

77
Student_Cuda_Image/src/core/01_student/02_Mandelbrot/host/Mandelbrot_BestGrid.h
View File

@@ -12,49 +12,56 @@
|* Impelmentation *| |* Impelmentation *|
\*---------------------------------------------------------------------*/ \*---------------------------------------------------------------------*/
namespace mandelbrot namespace mandelbrot {
{
class BestGrid class BestGrid {
{
public: public:
static Grid get() static Grid get() {
{ // const int MP = Hardware::getMPCount();
const int MP = Hardware::getMPCount(); // const int CORE_MP = Hardware::getCoreCountMP();
const int CORE_MP = Hardware::getCoreCountMP();
// fp64 (float 64 bits) // fp64 (float 64 bits)
#ifdef MANDELBROT_DOUBLE #ifdef MANDELBROT_DOUBLE
//TODO Mandelbrot const int MP = Hardware::getMPCount();
// to remove once coded const int CORE_MP = Hardware::getCoreCountMP();
{
Couts::redln("aie aie aie, your best grid won t build itself");
assert(false);
}
#endif
// fp32 (float 32 bits) dim3 dg(MP, 1, 1);
#ifdef MANDELBROT_FLOAT dim3 db(CORE_MP, 1, 1);
//TODO Mandelbrot Grid grid(dg, db);
// to remove once coded
{
Couts::redln("aie aie aie, your best grid won t build itself");
assert(false);
}
#endif
// fp16 (float 32 bits) return grid;
#ifdef MANDELBROT_HALF #endif
//TODO Mandelbrot
// to remove once coded // fp32 (float 32 bits)
{ #ifdef MANDELBROT_FLOAT
Couts::redln("aie aie aie, your best grid won t build itself"); const int MP = Hardware::getMPCount();
assert(false); const int CORE_MP = Hardware::getCoreCountMP();
}
#endif dim3 dg(MP, 1, 1);
dim3 db(CORE_MP, 1, 1);
Grid grid(dg, db);
return grid;
#endif
// fp16 (float 32 bits)
#ifdef MANDELBROT_HALF
const int MP = Hardware::getMPCount();
const int CORE_MP = Hardware::getCoreCountMP();
dim3 dg(MP, 2, 1);
dim3 db(CORE_MP, 4, 1);
Grid grid(dg, db);
return grid;
#endif
// default
Couts::redln("Oups, something went wrong on BestGrid settings");
assert(false); // unknown type
return Grid();
} }
}; };

168
Student_Cuda_Image/src/core/01_student/02_Mandelbrot/real_mandelbrot.h
View File

@@ -14,107 +14,101 @@
\*---------------------*/ \*---------------------*/
// Choose one of the two (either/either): // Choose one of the two (either/either):
//#define MANDELBROT_DOUBLE // #define MANDELBROT_DOUBLE // fp64
//#define MANDELBROT_FLOAT // #define MANDELBROT_FLOAT // fp32
#define MANDELBROT_HALF #define MANDELBROT_HALF // fp16
/*----------------------*\ /*----------------------*\
|* private *| |* private *|
\*---------------------*/ \*---------------------*/
namespace mandelbrotReal namespace mandelbrotReal {
{
// fp64 (float 64 bits) // fp64 (float 64 bits)
#ifdef MANDELBROT_DOUBLE #ifdef MANDELBROT_DOUBLE
//#define real double //#define real double
using real=double; using real=double;
#endif #endif
// fp32 (float 32 bits) // fp32 (float 32 bits)
#ifdef MANDELBROT_FLOAT #ifdef MANDELBROT_FLOAT
//#define real float //#define real float
using real=float; using real=float;
#endif #endif
// fp16 (float 16 bits) // fp16 (float 16 bits)
#ifdef MANDELBROT_HALF #ifdef MANDELBROT_HALF
//#define real half //#define real half
using real=half; using real=half;
#endif #endif
static std::string realToString() static std::string realToString() {
{
#ifdef MANDELBROT_DOUBLE #ifdef MANDELBROT_DOUBLE
return "fp64"; return "fp64";
#endif #endif
#ifdef MANDELBROT_FLOAT #ifdef MANDELBROT_FLOAT
return "fp32"; return "fp32";
#endif #endif
#ifdef MANDELBROT_HALF #ifdef MANDELBROT_HALF
return "fp16"; return "fp16";
#endif #endif
}
static bool isFp16()
{
#ifdef MANDELBROT_DOUBLE
return false;
#endif
#ifdef MANDELBROT_FLOAT
return false;
#endif
#ifdef MANDELBROT_HALF
return true;
#endif
assert(false);
}
static bool isFp32()
{
#ifdef MANDELBROT_DOUBLE
return false;
#endif
#ifdef MANDELBROT_FLOAT
return true;
#endif
#ifdef MANDELBROT_HALF
return false;
#endif
assert(false);
}
static bool isFp64()
{
#ifdef MANDELBROT_DOUBLE
return true;
#endif
#ifdef MANDELBROT_FLOAT
return false;
#endif
#ifdef MANDELBROT_HALF
return false;
#endif
assert(false);
}
} }
static bool isFp16() {
#ifdef MANDELBROT_DOUBLE
return false;
#endif
#ifdef MANDELBROT_FLOAT
return false;
#endif
#ifdef MANDELBROT_HALF
return true;
#endif
assert(false);
}
static bool isFp32() {
#ifdef MANDELBROT_DOUBLE
return false;
#endif
#ifdef MANDELBROT_FLOAT
return true;
#endif
#ifdef MANDELBROT_HALF
return false;
#endif
assert(false);
}
static bool isFp64() {
#ifdef MANDELBROT_DOUBLE
return true;
#endif
#ifdef MANDELBROT_FLOAT
return false;
#endif
#ifdef MANDELBROT_HALF
return false;
#endif
assert(false);
}
}
/*----------------------------------------------------------------------*\ /*----------------------------------------------------------------------*\
|* End *| |* End *|
\*---------------------------------------------------------------------*/ \*---------------------------------------------------------------------*/

4
Student_Cuda_Image/src/main/mainBenchmark.cpp
View File

@@ -34,8 +34,8 @@ static void raytracingCM2SM();
int mainBenchmark() { int mainBenchmark() {
// Please, un a la fois! // Please, un a la fois!
rippling(); // rippling();
// mandelbrot(); // Conseil : use nFixe (by example nMin=nMax=80) mandelbrot(); // Conseil : use nFixe (by example nMin=nMax=80)
// //
// raytracingGM(); // raytracingGM();
// raytracingSM(); // raytracingSM();

4
Student_Cuda_Image/src/main/mainBruteforce.cpp
View File

@@ -44,8 +44,8 @@ int mainBrutforce() {
Matlab matlab; Matlab matlab;
// Please, un a la fois! // Please, un a la fois!
rippling(&matlab); // rippling(&matlab);
// mandelbrot(&matlab); // Conseil : use nFixe (by example nMin=nMax=100) mandelbrot(&matlab); // Conseil : use nFixe (by example nMin=nMax=100)
// raytracingGM(&matlab); // raytracingGM(&matlab);
// raytracingCM(&matlab); // raytracingCM(&matlab);
// raytracingSM(&matlab); // raytracingSM(&matlab);

4
Student_Cuda_Image/src/main/mainImage.cpp
View File

@@ -31,8 +31,8 @@ int mainImage(const Args& args) {
ImageOption zoomable(true, true, true, true); ImageOption zoomable(true, true, true, true);
ImageOption nozoomable(false, true, false, false); ImageOption nozoomable(false, true, false, false);
Viewer<RipplingProvider> rippling(nozoomable, 0, 0); // imageOption px py // Viewer<RipplingProvider> rippling(nozoomable, 0, 0); // imageOption px py
// Viewer<MandelbrotProvider> mandelbrot(zoomable, 0, 0); Viewer<MandelbrotProvider> mandelbrot(zoomable, 0, 0);
// Viewer<RaytracingProviderGM> raytracingGM(nozoomable, 0, 0); // Viewer<RaytracingProviderGM> raytracingGM(nozoomable, 0, 0);
// Viewer<RaytracingProviderCM> raytracingCM(nozoomable, 0, 0); // Viewer<RaytracingProviderCM> raytracingCM(nozoomable, 0, 0);
// Viewer<RaytracingProviderSM> raytracingSM(nozoomable, 0, 0); // Viewer<RaytracingProviderSM> raytracingSM(nozoomable, 0, 0);

6
Student_Cuda_Image/src/main/mainTest.cpp
View File

@@ -34,9 +34,9 @@ static void raytracing();
int mainTest() { int mainTest() {
// activer ci-dessous seulement le TP voulu (pas tous) // activer ci-dessous seulement le TP voulu (pas tous)
rippling(); rippling();
//mandelbrot(); // fp16 only mandelbrot(); // fp16 only
//raytracing(); // voir code ci-dessous pour activer la version voulue // raytracing(); // voir code ci-dessous pour activer la version voulue
return EXIT_SUCCESS; return EXIT_SUCCESS;
} }

10
Student_Cuda_Tools_Reduction/.clangd Normal file
View File

@@ -0,0 +1,10 @@
CompileFlags:
Add:
- "-I/home/mse15/CUDA/toStudent/code/WCudaStudent/Student_Cuda_Tools_Reduction/INC_SYMLINK/EXT"
- "-I/home/mse15/CUDA/toStudent/code/WCudaStudent/Student_Cuda_Tools_Reduction/INC_SYMLINK/PROJECT"
- "-std=c++17"
- "-x"
- "cuda"
---
Diagnostics:
Suppress: "*"

14
Student_Cuda_Tools_Reduction/.zed/tasks.json Normal file
View File

@@ -0,0 +1,14 @@
[
{
"label": "Run CUDA",
"command": "cbicc cuda clean jall run",
"use_new_terminal": false,
"allow_concurrent_runs": false,
"reveal": "always",
"reveal_target": "dock",
"hide": "never",
"shell": "system",
"show_summary": true,
"show_command": true
}
]

2
Student_Cuda_Tools_Reduction/INC_SYMLINK/PROJECT/Reduction.cu.h
View File

@@ -1 +1 @@
/home/bilat/CBI/Dropbox/02_CBI_LINUX/CoursCuda/toStudent/code/WCudaStudent/Student_Cuda_Tools_Reduction/src/core/01_algorithme/generic/Reduction.cu.h /home/mse15/CUDA/toStudent/code/WCudaStudent/Student_Cuda_Tools_Reduction/src/core/01_algorithme/generic/Reduction.cu.h

2
Student_Cuda_Tools_Reduction/INC_SYMLINK/PROJECT/ReductionAdd.cu.h
View File

@@ -1 +1 @@
/home/bilat/CBI/Dropbox/02_CBI_LINUX/CoursCuda/toStudent/code/WCudaStudent/Student_Cuda_Tools_Reduction/src/core/01_algorithme/add/ReductionAdd.cu.h /home/mse15/CUDA/toStudent/code/WCudaStudent/Student_Cuda_Tools_Reduction/src/core/01_algorithme/add/ReductionAdd.cu.h

2
Student_Cuda_Tools_Reduction/INC_SYMLINK/PROJECT/ReductionAddIntI.h
View File

@@ -1 +1 @@
/home/bilat/CBI/Dropbox/02_CBI_LINUX/CoursCuda/toStudent/code/WCudaStudent/Student_Cuda_Tools_Reduction/src/core/02_use_protocole/add/int/PI/host/ReductionAddIntI.h /home/mse15/CUDA/toStudent/code/WCudaStudent/Student_Cuda_Tools_Reduction/src/core/02_use_protocole/add/int/PI/host/ReductionAddIntI.h

2
Student_Cuda_Tools_Reduction/INC_SYMLINK/PROJECT/ReductionAddIntII.h
View File

@@ -1 +1 @@
/home/bilat/CBI/Dropbox/02_CBI_LINUX/CoursCuda/toStudent/code/WCudaStudent/Student_Cuda_Tools_Reduction/src/core/02_use_protocole/add/int/PII/host/ReductionAddIntII.h /home/mse15/CUDA/toStudent/code/WCudaStudent/Student_Cuda_Tools_Reduction/src/core/02_use_protocole/add/int/PII/host/ReductionAddIntII.h

2
Student_Cuda_Tools_Reduction/INC_SYMLINK/PROJECT/ReductionIntI.h
View File

@@ -1 +1 @@
/home/bilat/CBI/Dropbox/02_CBI_LINUX/CoursCuda/toStudent/code/WCudaStudent/Student_Cuda_Tools_Reduction/src/core/02_use_protocole/generic/int/PI/host/ReductionIntI.h /home/mse15/CUDA/toStudent/code/WCudaStudent/Student_Cuda_Tools_Reduction/src/core/02_use_protocole/generic/int/PI/host/ReductionIntI.h

2
Student_Cuda_Tools_Reduction/INC_SYMLINK/PROJECT/ReductionIntII.h
View File

@@ -1 +1 @@
/home/bilat/CBI/Dropbox/02_CBI_LINUX/CoursCuda/toStudent/code/WCudaStudent/Student_Cuda_Tools_Reduction/src/core/02_use_protocole/generic/int/PII/host/ReductionIntII.h /home/mse15/CUDA/toStudent/code/WCudaStudent/Student_Cuda_Tools_Reduction/src/core/02_use_protocole/generic/int/PII/host/ReductionIntII.h

2
Student_Cuda_Tools_Reduction/INC_SYMLINK/PROJECT/ReductionLongII.h
View File

@@ -1 +1 @@
/home/bilat/CBI/Dropbox/02_CBI_LINUX/CoursCuda/toStudent/code/WCudaStudent/Student_Cuda_Tools_Reduction/src/core/02_use_protocole/generic/long/PII/host/ReductionLongII.h /home/mse15/CUDA/toStudent/code/WCudaStudent/Student_Cuda_Tools_Reduction/src/core/02_use_protocole/generic/long/PII/host/ReductionLongII.h

194
Student_Cuda_Tools_Reduction/out/Reduce_Add_IntII_justesse.html Normal file
View File

@@ -0,0 +1,194 @@
<!DOCTYPE html PUBLIC "-//W3C//DTD XHTML 1.0 Strict//EN" "http://www.w3.org/TR/xhtml1/DTD/xhtml1-strict.dtd">
<html xmlns="http://www.w3.org/1999/xhtml" xml:lang="en" lang="en">
<head>
<meta http-equiv="content-type" content="text/html; charset=utf-8" />
<meta name="generator" content="CppTest - https://github.com/cpptest/cpptest" />
<title>Reduce_Add_IntII_justesse </title>
<style type="text/css" media="screen">
<!--
hr {
width: 100%;
border-width: 0px;
height: 1px;
color: #cccccc;
background-color: #cccccc;
padding: 0px;
}
table {
width:100%;
border-collapse:separate;
border-spacing: 2px;
border:0px;
}
tr {
margin:0px;
padding:0px;
}
td {
margin:0px;
padding:1px;
}
.table_summary {
}
.table_suites {
}
.table_suite {
}
.table_result {
margin: 0px 0px 1em 0px;
}
.tablecell_title {
background-color: #a5cef7;
font-weight: bold;
}
.tablecell_success {
background-color: #efefe7;
}
.tablecell_error {
color: #ff0808;
background-color: #efefe7;
font-weight: bold;
}
p.spaced {
margin: 0px;
padding: 1em 0px 2em 0px;
}
p.unspaced {
margin: 0px;
padding: 0px 0px 2em 0px;
}
-->
</style>
</head>
<body>
<h1><a name="top"></a>Reduce_Add_IntII_justesse </h1>
<div style="text-align:right">
Designed by <a href="https://github.com/cpptest/cpptest">CppTest</a>
</div>
<hr />
<h2>Summary</h2>
<table summary="Summary of test results" class="table_summary">
<tr>
<td style="width:30%" class="tablecell_title">Tests</td>
<td style="width:30%" class="tablecell_title">Errors</td>
<td style="width:30%" class="tablecell_title">Success</td>
<td style="width:10%" class="tablecell_title">Time (s)</td>
</tr>
<tr>
<td style="width:30%" class="tablecell_success">0</td>
<td style="width:30%" class="tablecell_success">0</td>
<td style="width:30%" class="tablecell_success">100%</td>
<td style="width:10%" class="tablecell_success">0.000000</td>
</tr>
</table>
<hr />
<h2>Test suites</h2>
<table summary="Test Suites" class="table_suites">
<tr>
<td class="tablecell_title">Name</td>
<td style="width:10%" class="tablecell_title">Tests</td>
<td style="width:10%" class="tablecell_title">Errors</td>
<td style="width:10%" class="tablecell_title">Success</td>
<td style="width:10%" class="tablecell_title">Time (s)</td>
</tr>
<tr>
<td class="tablecell_success"><a href="#TestReductionAddII">TestReductionAddII</a></td>
<td style="width:10%" class="tablecell_success">10</td>
<td style="width:10%" class="tablecell_success">0</td>
<td style="width:10%" class="tablecell_success">100%</td>
<td style="width:10%" class="tablecell_success">0.000000</td>
</tr>
</table>
<hr />
<h3><a name="TestReductionAddII"></a>Suite: TestReductionAddII</h3>
<table summary="Details for suite TestReductionAddII" class="table_suite">
<tr>
<td class="tablecell_title">Name</td>
<td style="width:10%" class="tablecell_title">Errors</td>
<td style="width:10%" class="tablecell_title">Success</td>
<td style="width:10%" class="tablecell_title">Time (s)</td>
</tr>
<tr>
<td class="tablecell_success">testDB2</td>
<td class="tablecell_success">0</td>
<td class="tablecell_success">true</td>
<td class="tablecell_success">0.000000</td>
</tr>
<tr>
<td class="tablecell_success">testDB4</td>
<td class="tablecell_success">0</td>
<td class="tablecell_success">true</td>
<td class="tablecell_success">0.000000</td>
</tr>
<tr>
<td class="tablecell_success">testDB8</td>
<td class="tablecell_success">0</td>
<td class="tablecell_success">true</td>
<td class="tablecell_success">0.000000</td>
</tr>
<tr>
<td class="tablecell_success">testDB16</td>
<td class="tablecell_success">0</td>
<td class="tablecell_success">true</td>
<td class="tablecell_success">0.000000</td>
</tr>
<tr>
<td class="tablecell_success">testDB32</td>
<td class="tablecell_success">0</td>
<td class="tablecell_success">true</td>
<td class="tablecell_success">0.000000</td>
</tr>
<tr>
<td class="tablecell_success">testDB64</td>
<td class="tablecell_success">0</td>
<td class="tablecell_success">true</td>
<td class="tablecell_success">0.000000</td>
</tr>
<tr>
<td class="tablecell_success">testDB128</td>
<td class="tablecell_success">0</td>
<td class="tablecell_success">true</td>
<td class="tablecell_success">0.000000</td>
</tr>
<tr>
<td class="tablecell_success">testGrid</td>
<td class="tablecell_success">0</td>
<td class="tablecell_success">true</td>
<td class="tablecell_success">0.000000</td>
</tr>
<tr>
<td class="tablecell_success">testMonoBlock</td>
<td class="tablecell_success">0</td>
<td class="tablecell_success">true</td>
<td class="tablecell_success">0.000000</td>
</tr>
<tr>
<td class="tablecell_success">testSpecialeMax</td>
<td class="tablecell_success">0</td>
<td class="tablecell_success">true</td>
<td class="tablecell_success">0.000000</td>
</tr>
</table>
<p class="spaced"><a href="#top">Back to top</a>
</p>
<hr />
<p>
<a href="http://validator.w3.org/#validate-by-upload">
Valid XHTML 1.0 Strict
</a>
</p>
</body>
</html>

140
Student_Cuda_Tools_Reduction/out/Reduce_Add_IntII_performance.html Normal file
View File

@@ -0,0 +1,140 @@
<!DOCTYPE html PUBLIC "-//W3C//DTD XHTML 1.0 Strict//EN" "http://www.w3.org/TR/xhtml1/DTD/xhtml1-strict.dtd">
<html xmlns="http://www.w3.org/1999/xhtml" xml:lang="en" lang="en">
<head>
<meta http-equiv="content-type" content="text/html; charset=utf-8" />
<meta name="generator" content="CppTest - https://github.com/cpptest/cpptest" />
<title>Reduce_Add_IntII_performance </title>
<style type="text/css" media="screen">
<!--
hr {
width: 100%;
border-width: 0px;
height: 1px;
color: #cccccc;
background-color: #cccccc;
padding: 0px;
}
table {
width:100%;
border-collapse:separate;
border-spacing: 2px;
border:0px;
}
tr {
margin:0px;
padding:0px;
}
td {
margin:0px;
padding:1px;
}
.table_summary {
}
.table_suites {
}
.table_suite {
}
.table_result {
margin: 0px 0px 1em 0px;
}
.tablecell_title {
background-color: #a5cef7;
font-weight: bold;
}
.tablecell_success {
background-color: #efefe7;
}
.tablecell_error {
color: #ff0808;
background-color: #efefe7;
font-weight: bold;
}
p.spaced {
margin: 0px;
padding: 1em 0px 2em 0px;
}
p.unspaced {
margin: 0px;
padding: 0px 0px 2em 0px;
}
-->
</style>
</head>
<body>
<h1><a name="top"></a>Reduce_Add_IntII_performance </h1>
<div style="text-align:right">
Designed by <a href="https://github.com/cpptest/cpptest">CppTest</a>
</div>
<hr />
<h2>Summary</h2>
<table summary="Summary of test results" class="table_summary">
<tr>
<td style="width:30%" class="tablecell_title">Tests</td>
<td style="width:30%" class="tablecell_title">Errors</td>
<td style="width:30%" class="tablecell_title">Success</td>
<td style="width:10%" class="tablecell_title">Time (s)</td>
</tr>
<tr>
<td style="width:30%" class="tablecell_success">1</td>
<td style="width:30%" class="tablecell_success">0</td>
<td style="width:30%" class="tablecell_success">100%</td>
<td style="width:10%" class="tablecell_success">15.000000</td>
</tr>
</table>
<hr />
<h2>Test suites</h2>
<table summary="Test Suites" class="table_suites">
<tr>
<td class="tablecell_title">Name</td>
<td style="width:10%" class="tablecell_title">Tests</td>
<td style="width:10%" class="tablecell_title">Errors</td>
<td style="width:10%" class="tablecell_title">Success</td>
<td style="width:10%" class="tablecell_title">Time (s)</td>
</tr>
<tr>
<td class="tablecell_success"><a href="#TestPerformance_RunnableGPU_A">TestPerformance_RunnableGPU_A</a></td>
<td style="width:10%" class="tablecell_success">1</td>
<td style="width:10%" class="tablecell_success">0</td>
<td style="width:10%" class="tablecell_success">100%</td>
<td style="width:10%" class="tablecell_success">15.000000</td>
</tr>
</table>
<hr />
<h3><a name="TestPerformance_RunnableGPU_A"></a>Suite: TestPerformance_RunnableGPU_A</h3>
<table summary="Details for suite TestPerformance_RunnableGPU_A" class="table_suite">
<tr>
<td class="tablecell_title">Name</td>
<td style="width:10%" class="tablecell_title">Errors</td>
<td style="width:10%" class="tablecell_title">Success</td>
<td style="width:10%" class="tablecell_title">Time (s)</td>
</tr>
<tr>
<td class="tablecell_success">performanceOnly</td>
<td class="tablecell_success">0</td>
<td class="tablecell_success">true</td>
<td class="tablecell_success">15.000000</td>
</tr>
</table>
<p class="spaced"><a href="#top">Back to top</a>
</p>
<hr />
<p>
<a href="http://validator.w3.org/#validate-by-upload">
Valid XHTML 1.0 Strict
</a>
</p>
</body>
</html>

1228
Student_Cuda_Tools_Reduction/out/Reduce_Add_IntI_justesse.html
View File

File diff suppressed because it is too large Load Diff

194
Student_Cuda_Tools_Reduction/out/Reduce_Generic_IntII_justesse.html Normal file
View File

@@ -0,0 +1,194 @@
<!DOCTYPE html PUBLIC "-//W3C//DTD XHTML 1.0 Strict//EN" "http://www.w3.org/TR/xhtml1/DTD/xhtml1-strict.dtd">
<html xmlns="http://www.w3.org/1999/xhtml" xml:lang="en" lang="en">
<head>
<meta http-equiv="content-type" content="text/html; charset=utf-8" />
<meta name="generator" content="CppTest - https://github.com/cpptest/cpptest" />
<title>Reduce_Generic_IntII_justesse </title>
<style type="text/css" media="screen">
<!--
hr {
width: 100%;
border-width: 0px;
height: 1px;
color: #cccccc;
background-color: #cccccc;
padding: 0px;
}
table {
width:100%;
border-collapse:separate;
border-spacing: 2px;
border:0px;
}
tr {
margin:0px;
padding:0px;
}
td {
margin:0px;
padding:1px;
}
.table_summary {
}
.table_suites {
}
.table_suite {
}
.table_result {
margin: 0px 0px 1em 0px;
}
.tablecell_title {
background-color: #a5cef7;
font-weight: bold;
}
.tablecell_success {
background-color: #efefe7;
}
.tablecell_error {
color: #ff0808;
background-color: #efefe7;
font-weight: bold;
}
p.spaced {
margin: 0px;
padding: 1em 0px 2em 0px;
}
p.unspaced {
margin: 0px;
padding: 0px 0px 2em 0px;
}
-->
</style>
</head>
<body>
<h1><a name="top"></a>Reduce_Generic_IntII_justesse </h1>
<div style="text-align:right">
Designed by <a href="https://github.com/cpptest/cpptest">CppTest</a>
</div>
<hr />
<h2>Summary</h2>
<table summary="Summary of test results" class="table_summary">
<tr>
<td style="width:30%" class="tablecell_title">Tests</td>
<td style="width:30%" class="tablecell_title">Errors</td>
<td style="width:30%" class="tablecell_title">Success</td>
<td style="width:10%" class="tablecell_title">Time (s)</td>
</tr>
<tr>
<td style="width:30%" class="tablecell_success">0</td>
<td style="width:30%" class="tablecell_success">0</td>
<td style="width:30%" class="tablecell_success">100%</td>
<td style="width:10%" class="tablecell_success">1.000000</td>
</tr>
</table>
<hr />
<h2>Test suites</h2>
<table summary="Test Suites" class="table_suites">
<tr>
<td class="tablecell_title">Name</td>
<td style="width:10%" class="tablecell_title">Tests</td>
<td style="width:10%" class="tablecell_title">Errors</td>
<td style="width:10%" class="tablecell_title">Success</td>
<td style="width:10%" class="tablecell_title">Time (s)</td>
</tr>
<tr>
<td class="tablecell_success"><a href="#TestReductionGenericII">TestReductionGenericII</a></td>
<td style="width:10%" class="tablecell_success">10</td>
<td style="width:10%" class="tablecell_success">0</td>
<td style="width:10%" class="tablecell_success">100%</td>
<td style="width:10%" class="tablecell_success">1.000000</td>
</tr>
</table>
<hr />
<h3><a name="TestReductionGenericII"></a>Suite: TestReductionGenericII</h3>
<table summary="Details for suite TestReductionGenericII" class="table_suite">
<tr>
<td class="tablecell_title">Name</td>
<td style="width:10%" class="tablecell_title">Errors</td>
<td style="width:10%" class="tablecell_title">Success</td>
<td style="width:10%" class="tablecell_title">Time (s)</td>
</tr>
<tr>
<td class="tablecell_success">testDB2</td>
<td class="tablecell_success">0</td>
<td class="tablecell_success">true</td>
<td class="tablecell_success">0.000000</td>
</tr>
<tr>
<td class="tablecell_success">testDB4</td>
<td class="tablecell_success">0</td>
<td class="tablecell_success">true</td>
<td class="tablecell_success">0.000000</td>
</tr>
<tr>
<td class="tablecell_success">testDB8</td>
<td class="tablecell_success">0</td>
<td class="tablecell_success">true</td>
<td class="tablecell_success">0.000000</td>
</tr>
<tr>
<td class="tablecell_success">testDB16</td>
<td class="tablecell_success">0</td>
<td class="tablecell_success">true</td>
<td class="tablecell_success">0.000000</td>
</tr>
<tr>
<td class="tablecell_success">testDB32</td>
<td class="tablecell_success">0</td>
<td class="tablecell_success">true</td>
<td class="tablecell_success">0.000000</td>
</tr>
<tr>
<td class="tablecell_success">testDB64</td>
<td class="tablecell_success">0</td>
<td class="tablecell_success">true</td>
<td class="tablecell_success">0.000000</td>
</tr>
<tr>
<td class="tablecell_success">testDB128</td>
<td class="tablecell_success">0</td>
<td class="tablecell_success">true</td>
<td class="tablecell_success">0.000000</td>
</tr>
<tr>
<td class="tablecell_success">testGrid</td>
<td class="tablecell_success">0</td>
<td class="tablecell_success">true</td>
<td class="tablecell_success">0.000000</td>
</tr>
<tr>
<td class="tablecell_success">testMonoBlock</td>
<td class="tablecell_success">0</td>
<td class="tablecell_success">true</td>
<td class="tablecell_success">1.000000</td>
</tr>
<tr>
<td class="tablecell_success">testSpecialeMax</td>
<td class="tablecell_success">0</td>
<td class="tablecell_success">true</td>
<td class="tablecell_success">0.000000</td>
</tr>
</table>
<p class="spaced"><a href="#top">Back to top</a>
</p>
<hr />
<p>
<a href="http://validator.w3.org/#validate-by-upload">
Valid XHTML 1.0 Strict
</a>
</p>
</body>
</html>

140
Student_Cuda_Tools_Reduction/out/Reduce_Generic_IntII_performance.html Normal file
View File

@@ -0,0 +1,140 @@
<!DOCTYPE html PUBLIC "-//W3C//DTD XHTML 1.0 Strict//EN" "http://www.w3.org/TR/xhtml1/DTD/xhtml1-strict.dtd">
<html xmlns="http://www.w3.org/1999/xhtml" xml:lang="en" lang="en">
<head>
<meta http-equiv="content-type" content="text/html; charset=utf-8" />
<meta name="generator" content="CppTest - https://github.com/cpptest/cpptest" />
<title>Reduce_Generic_IntII_performance </title>
<style type="text/css" media="screen">
<!--
hr {
width: 100%;
border-width: 0px;
height: 1px;
color: #cccccc;
background-color: #cccccc;
padding: 0px;
}
table {
width:100%;
border-collapse:separate;
border-spacing: 2px;
border:0px;
}
tr {
margin:0px;
padding:0px;
}
td {
margin:0px;
padding:1px;
}
.table_summary {
}
.table_suites {
}
.table_suite {
}
.table_result {
margin: 0px 0px 1em 0px;
}
.tablecell_title {
background-color: #a5cef7;
font-weight: bold;
}
.tablecell_success {
background-color: #efefe7;
}
.tablecell_error {
color: #ff0808;
background-color: #efefe7;
font-weight: bold;
}
p.spaced {
margin: 0px;
padding: 1em 0px 2em 0px;
}
p.unspaced {
margin: 0px;
padding: 0px 0px 2em 0px;
}
-->
</style>
</head>
<body>
<h1><a name="top"></a>Reduce_Generic_IntII_performance </h1>
<div style="text-align:right">
Designed by <a href="https://github.com/cpptest/cpptest">CppTest</a>
</div>
<hr />
<h2>Summary</h2>
<table summary="Summary of test results" class="table_summary">
<tr>
<td style="width:30%" class="tablecell_title">Tests</td>
<td style="width:30%" class="tablecell_title">Errors</td>
<td style="width:30%" class="tablecell_title">Success</td>
<td style="width:10%" class="tablecell_title">Time (s)</td>
</tr>
<tr>
<td style="width:30%" class="tablecell_success">1</td>
<td style="width:30%" class="tablecell_success">0</td>
<td style="width:30%" class="tablecell_success">100%</td>
<td style="width:10%" class="tablecell_success">16.000000</td>
</tr>
</table>
<hr />
<h2>Test suites</h2>
<table summary="Test Suites" class="table_suites">
<tr>
<td class="tablecell_title">Name</td>
<td style="width:10%" class="tablecell_title">Tests</td>
<td style="width:10%" class="tablecell_title">Errors</td>
<td style="width:10%" class="tablecell_title">Success</td>
<td style="width:10%" class="tablecell_title">Time (s)</td>
</tr>
<tr>
<td class="tablecell_success"><a href="#TestPerformance_RunnableGPU_A">TestPerformance_RunnableGPU_A</a></td>
<td style="width:10%" class="tablecell_success">1</td>
<td style="width:10%" class="tablecell_success">0</td>
<td style="width:10%" class="tablecell_success">100%</td>
<td style="width:10%" class="tablecell_success">16.000000</td>
</tr>
</table>
<hr />
<h3><a name="TestPerformance_RunnableGPU_A"></a>Suite: TestPerformance_RunnableGPU_A</h3>
<table summary="Details for suite TestPerformance_RunnableGPU_A" class="table_suite">
<tr>
<td class="tablecell_title">Name</td>
<td style="width:10%" class="tablecell_title">Errors</td>
<td style="width:10%" class="tablecell_title">Success</td>
<td style="width:10%" class="tablecell_title">Time (s)</td>
</tr>
<tr>
<td class="tablecell_success">performanceOnly</td>
<td class="tablecell_success">0</td>
<td class="tablecell_success">true</td>
<td class="tablecell_success">16.000000</td>
</tr>
</table>
<p class="spaced"><a href="#top">Back to top</a>
</p>
<hr />
<p>
<a href="http://validator.w3.org/#validate-by-upload">
Valid XHTML 1.0 Strict
</a>
</p>
</body>
</html>

212
Student_Cuda_Tools_Reduction/out/Reduce_Generic_IntI_justesse.html Normal file
View File

@@ -0,0 +1,212 @@
<!DOCTYPE html PUBLIC "-//W3C//DTD XHTML 1.0 Strict//EN" "http://www.w3.org/TR/xhtml1/DTD/xhtml1-strict.dtd">
<html xmlns="http://www.w3.org/1999/xhtml" xml:lang="en" lang="en">
<head>
<meta http-equiv="content-type" content="text/html; charset=utf-8" />
<meta name="generator" content="CppTest - https://github.com/cpptest/cpptest" />
<title>Reduce_Generic_IntI_justesse </title>
<style type="text/css" media="screen">
<!--
hr {
width: 100%;
border-width: 0px;
height: 1px;
color: #cccccc;
background-color: #cccccc;
padding: 0px;
}
table {
width:100%;
border-collapse:separate;
border-spacing: 2px;
border:0px;
}
tr {
margin:0px;
padding:0px;
}
td {
margin:0px;
padding:1px;
}
.table_summary {
}
.table_suites {
}
.table_suite {
}
.table_result {
margin: 0px 0px 1em 0px;
}
.tablecell_title {
background-color: #a5cef7;
font-weight: bold;
}
.tablecell_success {
background-color: #efefe7;
}
.tablecell_error {
color: #ff0808;
background-color: #efefe7;
font-weight: bold;
}
p.spaced {
margin: 0px;
padding: 1em 0px 2em 0px;
}
p.unspaced {
margin: 0px;
padding: 0px 0px 2em 0px;
}
-->
</style>
</head>
<body>
<h1><a name="top"></a>Reduce_Generic_IntI_justesse </h1>
<div style="text-align:right">
Designed by <a href="https://github.com/cpptest/cpptest">CppTest</a>
</div>
<hr />
<h2>Summary</h2>
<table summary="Summary of test results" class="table_summary">
<tr>
<td style="width:30%" class="tablecell_title">Tests</td>
<td style="width:30%" class="tablecell_title">Errors</td>
<td style="width:30%" class="tablecell_title">Success</td>
<td style="width:10%" class="tablecell_title">Time (s)</td>
</tr>
<tr>
<td style="width:30%" class="tablecell_success">0</td>
<td style="width:30%" class="tablecell_success">0</td>
<td style="width:30%" class="tablecell_success">100%</td>
<td style="width:10%" class="tablecell_success">1.000000</td>
</tr>
</table>
<hr />
<h2>Test suites</h2>
<table summary="Test Suites" class="table_suites">
<tr>
<td class="tablecell_title">Name</td>
<td style="width:10%" class="tablecell_title">Tests</td>
<td style="width:10%" class="tablecell_title">Errors</td>
<td style="width:10%" class="tablecell_title">Success</td>
<td style="width:10%" class="tablecell_title">Time (s)</td>
</tr>
<tr>
<td class="tablecell_success"><a href="#TestReductionGenericI">TestReductionGenericI</a></td>
<td style="width:10%" class="tablecell_success">13</td>
<td style="width:10%" class="tablecell_success">0</td>
<td style="width:10%" class="tablecell_success">100%</td>
<td style="width:10%" class="tablecell_success">1.000000</td>
</tr>
</table>
<hr />
<h3><a name="TestReductionGenericI"></a>Suite: TestReductionGenericI</h3>
<table summary="Details for suite TestReductionGenericI" class="table_suite">
<tr>
<td class="tablecell_title">Name</td>
<td style="width:10%" class="tablecell_title">Errors</td>
<td style="width:10%" class="tablecell_title">Success</td>
<td style="width:10%" class="tablecell_title">Time (s)</td>
</tr>
<tr>
<td class="tablecell_success">testDB2</td>
<td class="tablecell_success">0</td>
<td class="tablecell_success">true</td>
<td class="tablecell_success">0.000000</td>
</tr>
<tr>
<td class="tablecell_success">testDB4</td>
<td class="tablecell_success">0</td>
<td class="tablecell_success">true</td>
<td class="tablecell_success">0.000000</td>
</tr>
<tr>
<td class="tablecell_success">testDB8</td>
<td class="tablecell_success">0</td>
<td class="tablecell_success">true</td>
<td class="tablecell_success">0.000000</td>
</tr>
<tr>
<td class="tablecell_success">testDB16</td>
<td class="tablecell_success">0</td>
<td class="tablecell_success">true</td>
<td class="tablecell_success">0.000000</td>
</tr>
<tr>
<td class="tablecell_success">testDB32</td>
<td class="tablecell_success">0</td>
<td class="tablecell_success">true</td>
<td class="tablecell_success">0.000000</td>
</tr>
<tr>
<td class="tablecell_success">testDB64</td>
<td class="tablecell_success">0</td>
<td class="tablecell_success">true</td>
<td class="tablecell_success">0.000000</td>
</tr>
<tr>
<td class="tablecell_success">testDB128</td>
<td class="tablecell_success">0</td>
<td class="tablecell_success">true</td>
<td class="tablecell_success">0.000000</td>
</tr>
<tr>
<td class="tablecell_success">testDB256</td>
<td class="tablecell_success">0</td>
<td class="tablecell_success">true</td>
<td class="tablecell_success">0.000000</td>
</tr>
<tr>
<td class="tablecell_success">testDB512</td>
<td class="tablecell_success">0</td>
<td class="tablecell_success">true</td>
<td class="tablecell_success">0.000000</td>
</tr>
<tr>
<td class="tablecell_success">testDB1024</td>
<td class="tablecell_success">0</td>
<td class="tablecell_success">true</td>
<td class="tablecell_success">0.000000</td>
</tr>
<tr>
<td class="tablecell_success">testGrid</td>
<td class="tablecell_success">0</td>
<td class="tablecell_success">true</td>
<td class="tablecell_success">1.000000</td>
</tr>
<tr>
<td class="tablecell_success">testMonoBlock</td>
<td class="tablecell_success">0</td>
<td class="tablecell_success">true</td>
<td class="tablecell_success">0.000000</td>
</tr>
<tr>
<td class="tablecell_success">testSpecialeMax</td>
<td class="tablecell_success">0</td>
<td class="tablecell_success">true</td>
<td class="tablecell_success">0.000000</td>
</tr>
</table>
<p class="spaced"><a href="#top">Back to top</a>
</p>
<hr />
<p>
<a href="http://validator.w3.org/#validate-by-upload">
Valid XHTML 1.0 Strict
</a>
</p>
</body>
</html>

140
Student_Cuda_Tools_Reduction/out/Reduce_Generic_IntI_performance.html Normal file
View File

@@ -0,0 +1,140 @@
<!DOCTYPE html PUBLIC "-//W3C//DTD XHTML 1.0 Strict//EN" "http://www.w3.org/TR/xhtml1/DTD/xhtml1-strict.dtd">
<html xmlns="http://www.w3.org/1999/xhtml" xml:lang="en" lang="en">
<head>
<meta http-equiv="content-type" content="text/html; charset=utf-8" />
<meta name="generator" content="CppTest - https://github.com/cpptest/cpptest" />
<title>Reduce_Generic_IntI_performance </title>
<style type="text/css" media="screen">
<!--
hr {
width: 100%;
border-width: 0px;
height: 1px;
color: #cccccc;
background-color: #cccccc;
padding: 0px;
}
table {
width:100%;
border-collapse:separate;
border-spacing: 2px;
border:0px;
}
tr {
margin:0px;
padding:0px;
}
td {
margin:0px;
padding:1px;
}
.table_summary {
}
.table_suites {
}
.table_suite {
}
.table_result {
margin: 0px 0px 1em 0px;
}
.tablecell_title {
background-color: #a5cef7;
font-weight: bold;
}
.tablecell_success {
background-color: #efefe7;
}
.tablecell_error {
color: #ff0808;
background-color: #efefe7;
font-weight: bold;
}
p.spaced {
margin: 0px;
padding: 1em 0px 2em 0px;
}
p.unspaced {
margin: 0px;
padding: 0px 0px 2em 0px;
}
-->
</style>
</head>
<body>
<h1><a name="top"></a>Reduce_Generic_IntI_performance </h1>
<div style="text-align:right">
Designed by <a href="https://github.com/cpptest/cpptest">CppTest</a>
</div>
<hr />
<h2>Summary</h2>
<table summary="Summary of test results" class="table_summary">
<tr>
<td style="width:30%" class="tablecell_title">Tests</td>
<td style="width:30%" class="tablecell_title">Errors</td>
<td style="width:30%" class="tablecell_title">Success</td>
<td style="width:10%" class="tablecell_title">Time (s)</td>
</tr>
<tr>
<td style="width:30%" class="tablecell_success">1</td>
<td style="width:30%" class="tablecell_success">0</td>
<td style="width:30%" class="tablecell_success">100%</td>
<td style="width:10%" class="tablecell_success">11.000000</td>
</tr>
</table>
<hr />
<h2>Test suites</h2>
<table summary="Test Suites" class="table_suites">
<tr>
<td class="tablecell_title">Name</td>
<td style="width:10%" class="tablecell_title">Tests</td>
<td style="width:10%" class="tablecell_title">Errors</td>
<td style="width:10%" class="tablecell_title">Success</td>
<td style="width:10%" class="tablecell_title">Time (s)</td>
</tr>
<tr>
<td class="tablecell_success"><a href="#TestPerformance_RunnableGPU_A">TestPerformance_RunnableGPU_A</a></td>
<td style="width:10%" class="tablecell_success">1</td>
<td style="width:10%" class="tablecell_success">0</td>
<td style="width:10%" class="tablecell_success">100%</td>
<td style="width:10%" class="tablecell_success">11.000000</td>
</tr>
</table>
<hr />
<h3><a name="TestPerformance_RunnableGPU_A"></a>Suite: TestPerformance_RunnableGPU_A</h3>
<table summary="Details for suite TestPerformance_RunnableGPU_A" class="table_suite">
<tr>
<td class="tablecell_title">Name</td>
<td style="width:10%" class="tablecell_title">Errors</td>
<td style="width:10%" class="tablecell_title">Success</td>
<td style="width:10%" class="tablecell_title">Time (s)</td>
</tr>
<tr>
<td class="tablecell_success">performanceOnly</td>
<td class="tablecell_success">0</td>
<td class="tablecell_success">true</td>
<td class="tablecell_success">11.000000</td>
</tr>
</table>
<p class="spaced"><a href="#top">Back to top</a>
</p>
<hr />
<p>
<a href="http://validator.w3.org/#validate-by-upload">
Valid XHTML 1.0 Strict
</a>
</p>
</body>
</html>

224
Student_Cuda_Tools_Reduction/out/Reduce_Generic_LongII_justesse.html Normal file
View File

@@ -0,0 +1,224 @@
<!DOCTYPE html PUBLIC "-//W3C//DTD XHTML 1.0 Strict//EN" "http://www.w3.org/TR/xhtml1/DTD/xhtml1-strict.dtd">
<html xmlns="http://www.w3.org/1999/xhtml" xml:lang="en" lang="en">
<head>
<meta http-equiv="content-type" content="text/html; charset=utf-8" />
<meta name="generator" content="CppTest - https://github.com/cpptest/cpptest" />
<title>Reduce_Generic_LongII_justesse </title>
<style type="text/css" media="screen">
<!--
hr {
width: 100%;
border-width: 0px;
height: 1px;
color: #cccccc;
background-color: #cccccc;
padding: 0px;
}
table {
width:100%;
border-collapse:separate;
border-spacing: 2px;
border:0px;
}
tr {
margin:0px;
padding:0px;
}
td {
margin:0px;
padding:1px;
}
.table_summary {
}
.table_suites {
}
.table_suite {
}
.table_result {
margin: 0px 0px 1em 0px;
}
.tablecell_title {
background-color: #a5cef7;
font-weight: bold;
}
.tablecell_success {
background-color: #efefe7;
}
.tablecell_error {
color: #ff0808;
background-color: #efefe7;
font-weight: bold;
}
p.spaced {
margin: 0px;
padding: 1em 0px 2em 0px;
}
p.unspaced {
margin: 0px;
padding: 0px 0px 2em 0px;
}
-->
</style>
</head>
<body>
<h1><a name="top"></a>Reduce_Generic_LongII_justesse </h1>
<div style="text-align:right">
Designed by <a href="https://github.com/cpptest/cpptest">CppTest</a>
</div>
<hr />
<h2>Summary</h2>
<table summary="Summary of test results" class="table_summary">
<tr>
<td style="width:30%" class="tablecell_title">Tests</td>
<td style="width:30%" class="tablecell_title">Errors</td>
<td style="width:30%" class="tablecell_title">Success</td>
<td style="width:10%" class="tablecell_title">Time (s)</td>
</tr>
<tr>
<td style="width:30%" class="tablecell_success">0</td>
<td style="width:30%" class="tablecell_success">0</td>
<td style="width:30%" class="tablecell_success">100%</td>
<td style="width:10%" class="tablecell_success">87.000000</td>
</tr>
</table>
<hr />
<h2>Test suites</h2>
<table summary="Test Suites" class="table_suites">
<tr>
<td class="tablecell_title">Name</td>
<td style="width:10%" class="tablecell_title">Tests</td>
<td style="width:10%" class="tablecell_title">Errors</td>
<td style="width:10%" class="tablecell_title">Success</td>
<td style="width:10%" class="tablecell_title">Time (s)</td>
</tr>
<tr>
<td class="tablecell_success"><a href="#TestReductionGenericLongII">TestReductionGenericLongII</a></td>
<td style="width:10%" class="tablecell_success">15</td>
<td style="width:10%" class="tablecell_success">0</td>
<td style="width:10%" class="tablecell_success">100%</td>
<td style="width:10%" class="tablecell_success">87.000000</td>
</tr>
</table>
<hr />
<h3><a name="TestReductionGenericLongII"></a>Suite: TestReductionGenericLongII</h3>
<table summary="Details for suite TestReductionGenericLongII" class="table_suite">
<tr>
<td class="tablecell_title">Name</td>
<td style="width:10%" class="tablecell_title">Errors</td>
<td style="width:10%" class="tablecell_title">Success</td>
<td style="width:10%" class="tablecell_title">Time (s)</td>
</tr>
<tr>
<td class="tablecell_success">testDB2</td>
<td class="tablecell_success">0</td>
<td class="tablecell_success">true</td>
<td class="tablecell_success">0.000000</td>
</tr>
<tr>
<td class="tablecell_success">testDB4</td>
<td class="tablecell_success">0</td>
<td class="tablecell_success">true</td>
<td class="tablecell_success">0.000000</td>
</tr>
<tr>
<td class="tablecell_success">testDB8</td>
<td class="tablecell_success">0</td>
<td class="tablecell_success">true</td>
<td class="tablecell_success">0.000000</td>
</tr>
<tr>
<td class="tablecell_success">testDB16</td>
<td class="tablecell_success">0</td>
<td class="tablecell_success">true</td>
<td class="tablecell_success">0.000000</td>
</tr>
<tr>
<td class="tablecell_success">testDB32</td>
<td class="tablecell_success">0</td>
<td class="tablecell_success">true</td>
<td class="tablecell_success">0.000000</td>
</tr>
<tr>
<td class="tablecell_success">testDB64</td>
<td class="tablecell_success">0</td>
<td class="tablecell_success">true</td>
<td class="tablecell_success">0.000000</td>
</tr>
<tr>
<td class="tablecell_success">testDB128</td>
<td class="tablecell_success">0</td>
<td class="tablecell_success">true</td>
<td class="tablecell_success">0.000000</td>
</tr>
<tr>
<td class="tablecell_success">testDB256</td>
<td class="tablecell_success">0</td>
<td class="tablecell_success">true</td>
<td class="tablecell_success">0.000000</td>
</tr>
<tr>
<td class="tablecell_success">testDB512</td>
<td class="tablecell_success">0</td>
<td class="tablecell_success">true</td>
<td class="tablecell_success">0.000000</td>
</tr>
<tr>
<td class="tablecell_success">testDB1024</td>
<td class="tablecell_success">0</td>
<td class="tablecell_success">true</td>
<td class="tablecell_success">0.000000</td>
</tr>
<tr>
<td class="tablecell_success">testGrid</td>
<td class="tablecell_success">0</td>
<td class="tablecell_success">true</td>
<td class="tablecell_success">1.000000</td>
</tr>
<tr>
<td class="tablecell_success">testMonoBlock</td>
<td class="tablecell_success">0</td>
<td class="tablecell_success">true</td>
<td class="tablecell_success">0.000000</td>
</tr>
<tr>
<td class="tablecell_success">testspecialGridDGXMAX</td>
<td class="tablecell_success">0</td>
<td class="tablecell_success">true</td>
<td class="tablecell_success">42.000000</td>
</tr>
<tr>
<td class="tablecell_success">testSpecialGrid2</td>
<td class="tablecell_success">0</td>
<td class="tablecell_success">true</td>
<td class="tablecell_success">43.000000</td>
</tr>
<tr>
<td class="tablecell_success">testSpecialeMax</td>
<td class="tablecell_success">0</td>
<td class="tablecell_success">true</td>
<td class="tablecell_success">1.000000</td>
</tr>
</table>
<p class="spaced"><a href="#top">Back to top</a>
</p>
<hr />
<p>
<a href="http://validator.w3.org/#validate-by-upload">
Valid XHTML 1.0 Strict
</a>
</p>
</body>
</html>

140
Student_Cuda_Tools_Reduction/out/Reduce_Generic_LongII_performance.html Normal file
View File

@@ -0,0 +1,140 @@
<!DOCTYPE html PUBLIC "-//W3C//DTD XHTML 1.0 Strict//EN" "http://www.w3.org/TR/xhtml1/DTD/xhtml1-strict.dtd">
<html xmlns="http://www.w3.org/1999/xhtml" xml:lang="en" lang="en">
<head>
<meta http-equiv="content-type" content="text/html; charset=utf-8" />
<meta name="generator" content="CppTest - https://github.com/cpptest/cpptest" />
<title>Reduce_Generic_LongII_performance </title>
<style type="text/css" media="screen">
<!--
hr {
width: 100%;
border-width: 0px;
height: 1px;
color: #cccccc;
background-color: #cccccc;
padding: 0px;
}
table {
width:100%;
border-collapse:separate;
border-spacing: 2px;
border:0px;
}
tr {
margin:0px;
padding:0px;
}
td {
margin:0px;
padding:1px;
}
.table_summary {
}
.table_suites {
}
.table_suite {
}
.table_result {
margin: 0px 0px 1em 0px;
}
.tablecell_title {
background-color: #a5cef7;
font-weight: bold;
}
.tablecell_success {
background-color: #efefe7;
}
.tablecell_error {
color: #ff0808;
background-color: #efefe7;
font-weight: bold;
}
p.spaced {
margin: 0px;
padding: 1em 0px 2em 0px;
}
p.unspaced {
margin: 0px;
padding: 0px 0px 2em 0px;
}
-->
</style>
</head>
<body>
<h1><a name="top"></a>Reduce_Generic_LongII_performance </h1>
<div style="text-align:right">
Designed by <a href="https://github.com/cpptest/cpptest">CppTest</a>
</div>
<hr />
<h2>Summary</h2>
<table summary="Summary of test results" class="table_summary">
<tr>
<td style="width:30%" class="tablecell_title">Tests</td>
<td style="width:30%" class="tablecell_title">Errors</td>
<td style="width:30%" class="tablecell_title">Success</td>
<td style="width:10%" class="tablecell_title">Time (s)</td>
</tr>
<tr>
<td style="width:30%" class="tablecell_success">1</td>
<td style="width:30%" class="tablecell_success">0</td>
<td style="width:30%" class="tablecell_success">100%</td>
<td style="width:10%" class="tablecell_success">10.000000</td>
</tr>
</table>
<hr />
<h2>Test suites</h2>
<table summary="Test Suites" class="table_suites">
<tr>
<td class="tablecell_title">Name</td>
<td style="width:10%" class="tablecell_title">Tests</td>
<td style="width:10%" class="tablecell_title">Errors</td>
<td style="width:10%" class="tablecell_title">Success</td>
<td style="width:10%" class="tablecell_title">Time (s)</td>
</tr>
<tr>
<td class="tablecell_success"><a href="#TestPerformance_RunnableGPU_A">TestPerformance_RunnableGPU_A</a></td>
<td style="width:10%" class="tablecell_success">1</td>
<td style="width:10%" class="tablecell_success">0</td>
<td style="width:10%" class="tablecell_success">100%</td>
<td style="width:10%" class="tablecell_success">10.000000</td>
</tr>
</table>
<hr />
<h3><a name="TestPerformance_RunnableGPU_A"></a>Suite: TestPerformance_RunnableGPU_A</h3>
<table summary="Details for suite TestPerformance_RunnableGPU_A" class="table_suite">
<tr>
<td class="tablecell_title">Name</td>
<td style="width:10%" class="tablecell_title">Errors</td>
<td style="width:10%" class="tablecell_title">Success</td>
<td style="width:10%" class="tablecell_title">Time (s)</td>
</tr>
<tr>
<td class="tablecell_success">performanceOnly</td>
<td class="tablecell_success">0</td>
<td class="tablecell_success">true</td>
<td class="tablecell_success">10.000000</td>
</tr>
</table>
<p class="spaced"><a href="#top">Back to top</a>
</p>
<hr />
<p>
<a href="http://validator.w3.org/#validate-by-upload">
Valid XHTML 1.0 Strict
</a>
</p>
</body>
</html>

57
Student_Cuda_Tools_Reduction/src/core/01_algorithme/add/ReductionAdd.cu.h
View File

@@ -1,13 +1,12 @@
#pragma once #pragma once
#include "Thread1D.cu.h" #include "Thread2D.cu.h"
/*----------------------------------------------------------------------*\ /*----------------------------------------------------------------------*\
|* Implementation *| |* Implementation *|
\*---------------------------------------------------------------------*/ \*---------------------------------------------------------------------*/
class ReductionAdd class ReductionAdd {
{
public: public:
/** /**
@@ -46,8 +45,9 @@ class ReductionAdd
* *
*/ */
template <typename T> template <typename T>
static __device__ void reduce(T* tabSM, T* ptrResultGM) static
{ __device__
void reduce(T* tabSM, T* ptrResultGM) {
// Rappel : // Rappel :
// |ThreadByBlock|=|tabSM| . // |ThreadByBlock|=|tabSM| .
// Il y autant de case en SM que de thread par block. // Il y autant de case en SM que de thread par block.
@@ -55,11 +55,11 @@ class ReductionAdd
// 1 thread <---> 1 armoire // 1 thread <---> 1 armoire
// TODO ReductionAdd // TODO ReductionAdd
// reductionIntraBlock reductionIntraBlock(tabSM);
// reductionInterblock reductionInterBlock(tabSM, ptrResultGM);
// __syncthreads();// pour touts les threads d'un meme block, necessaires? ou? pas a le fin en tous les cas // __syncthreads();// pour touts les threads d'un meme block, necessaires? ou? pas a le fin en tous les cas
} }
private: private:
@@ -71,8 +71,9 @@ class ReductionAdd
* used by reductionIntraBlock * used by reductionIntraBlock
*/ */
template <typename T> template <typename T>
static __device__ void ecrasement(T* tabSM, int middle) static
{ __device__
void ecrasement(T* tabSM, int middle) {
// Indications : // Indications :
// (I1) je suis un thread, je dois faire quoi ? // (I1) je suis un thread, je dois faire quoi ?
// (I2) Tous les threads doivent-ils faire quelquechose? // (I2) Tous les threads doivent-ils faire quelquechose?
@@ -80,39 +81,59 @@ class ReductionAdd
// TODO ReductionAdd // TODO ReductionAdd
const int localTID = Thread2D::tidLocal();
if(localTID < middle) {
tabSM[localTID] = tabSM[localTID] + tabSM[localTID + middle];
}
// __syncthreads();// pour touts les threads d'un meme block, necessaires? ou? // __syncthreads();// pour touts les threads d'un meme block, necessaires? ou?
} }
/** /**
* Sur place, le resultat est dans tabSM[0] * Sur place, le resultat est dans tabSM[0]
*/ */
template <typename T> template <typename T>
static __device__ void reductionIntraBlock(T* tabSM) static
{ __device__
void reductionIntraBlock(T* tabSM) { // Reduce tab SM (all in [0])
// Ecrasement sucessifs dans une boucle (utiliser la methode ecrasement ci-dessus) // Ecrasement sucessifs dans une boucle (utiliser la methode ecrasement ci-dessus)
// TODO ReductionAdd // TODO ReductionAdd
const int NB_THREAD_LOCAL = Thread2D::nbThreadLocal();
int middle = NB_THREAD_LOCAL>>1;
while (middle > 0) {
ecrasement(tabSM, middle);
__syncthreads();
middle = middle >> 1;
}
// __syncthreads();// pour touts les threads d'un meme block, necessaires? ou? // __syncthreads();// pour touts les threads d'un meme block, necessaires? ou?
} }
/*--------------------------------------*\ /*--------------------------------------*\
|* reductionInterblock *| |* reductionInterblock *|
\*-------------------------------------*/ \*-------------------------------------*/
template <typename T> template <typename T>
static __device__ void reductionInterBlock(T* tabSM, T* ptrResultGM) static
{ __device__
void reductionInterBlock(T* tabSM, T* ptrResultGM) { // SM -> GM
// Indication: // Indication:
// (I1) Utiliser atomicAdd(pointeurDestination, valeurSource); // (I1) Utiliser atomicAdd(pointeurDestination, valeurSource);
// (i2) Travailler sous l hypothese d'une grid2d,avec Thread2D // (i2) Travailler sous l hypothese d'une grid2d,avec Thread2D
// TODO ReductionAdd // TODO ReductionAdd
if(Thread2D::tidLocal() == 0) {
atomicAdd(ptrResultGM, tabSM[0]);
}
// __syncthreads();// pour touts les threads d'un meme block, necessaires? ou? // __syncthreads();// pour touts les threads d'un meme block, necessaires? ou?
} }
}; };
/*----------------------------------------------------------------------*\ /*----------------------------------------------------------------------*\
|* End *| |* End *|

56
Student_Cuda_Tools_Reduction/src/core/01_algorithme/generic/Reduction.cu.h
View File

@@ -1,7 +1,7 @@
#pragma once #pragma once
#include "Lock.cu.h" #include "Lock.cu.h"
#include "Thread1D.cu.h" #include "Thread2D.cu.h"
/*----------------------------------------------------------------------*\ /*----------------------------------------------------------------------*\
|* prt fonction / reduction *| |* prt fonction / reduction *|
@@ -14,8 +14,7 @@
|* Implementation *| |* Implementation *|
\*---------------------------------------------------------------------*/ \*---------------------------------------------------------------------*/
class Reduction class Reduction {
{
public: public:
/** /**
@@ -50,14 +49,18 @@ class Reduction
* ReductionGeneric::reduce(add,addAtomic,tabSm,ptrResultGM); * ReductionGeneric::reduce(add,addAtomic,tabSm,ptrResultGM);
*/ */
template <typename T> template <typename T>
static __device__ void reduce(BinaryOperator(OP) ,AtomicOp(ATOMIC_OP), T* tabSM, T* ptrResultGM) static
__device__
void reduce(BinaryOperator(OP) ,AtomicOp(ATOMIC_OP), T* tabSM, T* ptrResultGM)
//static __device__ void reduce(T (*OP)(T, T) ,void (*ATOMIC_OP)(T*, T), T* tabSM, T* ptrResultGM) // idem ci-dessus mais sans define //static __device__ void reduce(T (*OP)(T, T) ,void (*ATOMIC_OP)(T*, T), T* tabSM, T* ptrResultGM) // idem ci-dessus mais sans define
{ {
// Meme principe que ReductionAdd // Meme principe que ReductionAdd
reductionIntraBlock<T>(OP,tabSM);
// TODO ReductionGeneric // TODO ReductionGeneric
// Meme principe que ReductionAdd // Meme principe que ReductionAdd
} reductionInterBlock<T>(ATOMIC_OP,tabSM, ptrResultGM);
}
private: private:
@@ -69,37 +72,58 @@ class Reduction
* used by reductionIntraBlock * used by reductionIntraBlock
*/ */
template <typename T> template <typename T>
static __device__ void ecrasement(BinaryOperator(OP),T* tabSM, int middle) static
{ __device__
void ecrasement(BinaryOperator(OP),T* tabSM, int middle) {
// TODO ReductionGeneric // TODO ReductionGeneric
// Meme principe que ReductionAdd // Meme principe que ReductionAdd
// OP est la variable representant l'operateur binaire // OP est la variable representant l'operateur binaire
}
const int tidLocal = Thread2D::tidLocal();
if (tidLocal < middle) {
tabSM[tidLocal] = OP(tabSM[tidLocal], tabSM[tidLocal + middle]);
}
}
/** /**
* Sur place, le resultat est dans tabSM[0] * Sur place, le resultat est dans tabSM[0]
*/ */
template <typename T> template <typename T>
static __device__ void reductionIntraBlock(BinaryOperator(OP),T* tabSM) static
{ __device__
void reductionIntraBlock(BinaryOperator(OP),T* tabSM) {
// TODO ReductionGeneric // TODO ReductionGeneric
// Meme principe que ReductionAdd // Meme principe que ReductionAdd
// OP est la variable representant l'operateur binaire // OP est la variable representant l'operateur binaire
}
const int NB_THEAD_LOCAL = Thread2D::nbThreadLocal();
int middle = NB_THEAD_LOCAL >> 1;
while (middle > 0) {
ecrasement<T>(OP,tabSM, middle);
__syncthreads();
middle = middle >> 1;
}
}
/*--------------------------------------*\ /*--------------------------------------*\
|* reductionInterblock *| |* reductionInterblock *|
\*-------------------------------------*/ \*-------------------------------------*/
template <typename T> template <typename T>
static __device__ void reductionInterBlock(AtomicOp(ATOMIC_OP), T* tabSM, T* ptrResultGM) static
{ __device__
void reductionInterBlock(AtomicOp(ATOMIC_OP), T* tabSM, T* ptrResultGM) {
// TODO ReductionGeneric // TODO ReductionGeneric
// Meme principe que ReductionAdd // Meme principe que ReductionAdd
// ATOMIC_OP est la variable representant l'operateur binaire atomic // ATOMIC_OP est la variable representant l'operateur binaire atomic
}
}; if (Thread2D::tidLocal() == 0) {
ATOMIC_OP(ptrResultGM, tabSM[0]);
}
}
};
/*----------------------------------------------------------------------*\ /*----------------------------------------------------------------------*\
|* End *| |* End *|

27
Student_Cuda_Tools_Reduction/src/core/02_use_protocole/add/int/PI/device/reductionAddIntI_device.cu
View File

@@ -1,4 +1,4 @@
#include "Thread1D.cu.h" #include "Thread2D.cu.h"
#include "cudas.h" #include "cudas.h"
#include "ReductionAdd.cu.h" #include "ReductionAdd.cu.h"
@@ -9,7 +9,9 @@
|* Private *| |* Private *|
\*---------------------------------------------------------------------*/ \*---------------------------------------------------------------------*/
static __device__ void reductionIntraThread(int* tabSM); static
__device__
void reductionIntraThread(int* tabSM);
/*----------------------------------------------------------------------*\ /*----------------------------------------------------------------------*\
|* Implementation *| |* Implementation *|
@@ -18,13 +20,19 @@ static __device__ void reductionIntraThread(int* tabSM);
/** /**
* 1 partout en tabSM * 1 partout en tabSM
*/ */
__global__ void KAddIntProtocoleI(int* ptrSumGM) __global__
{ void KAddIntProtocoleI(int* ptrSumGM) {
// TODO ReductionAddIntI // TODO ReductionAddIntI
// Reception tabSM // Reception tabSM
extern __shared__ int tabSM[];
// ReductionIntraThread // ReductionIntraThread
reductionIntraThread(tabSM);
__syncthreads();
// ReductionAdd // ReductionAdd
ReductionAdd::reduce(tabSM, ptrSumGM);
// __syncthreads(); // des threads de meme block!// Question : utile? ou? // __syncthreads(); // des threads de meme block!// Question : utile? ou?
} }
@@ -36,12 +44,15 @@ __global__ void KAddIntProtocoleI(int* ptrSumGM)
/** /**
* 1 partout en tabSM * 1 partout en tabSM
*/ */
__device__ void reductionIntraThread(int* tabSM) __device__
{ void reductionIntraThread(int* tabSM) {
// TODO ReductionAddIntI // TODO ReductionAddIntI
} const int localTID = Thread2D::tidLocal();
tabSM[localTID] = 1;
}
/*----------------------------------------------------------------------*\ /*----------------------------------------------------------------------*\
|* End *| |* End *|
\*---------------------------------------------------------------------*/ \*---------------------------------------------------------------------*/

21
Student_Cuda_Tools_Reduction/src/core/02_use_protocole/add/int/PI/host/ReductionAddIntI.cu
View File

@@ -27,32 +27,33 @@ extern __global__ void KAddIntProtocoleI(int* ptrSumGM);
ReductionAddIntI::ReductionAddIntI(const Grid& grid , int* ptrSum , bool isVerbose) : ReductionAddIntI::ReductionAddIntI(const Grid& grid , int* ptrSum , bool isVerbose) :
//RunnableGPU(grid, "Reduce_Add_IntI_" + to_string(grid.threadCounts()),isVerbose), // classe parente //RunnableGPU(grid, "Reduce_Add_IntI_" + to_string(grid.threadCounts()),isVerbose), // classe parente
RunnableGPU(grid, "Reduce_Add_IntI", isVerbose), // classe parente RunnableGPU(grid, "Reduce_Add_IntI", isVerbose), // classe parente
ptrSum(ptrSum) ptrSum(ptrSum) {
{
// TODO ReductionAddIntI // TODO ReductionAddIntI
// MM pour ptrSumGM (oubliez pas initialisation) // MM pour ptrSumGM (oubliez pas initialisation)
this->sizeSM = -1; this->sizeSM = grid.threadByBlock() * sizeof(int);
// Tip: Il y a une methode dedier pour malloquer un int cote device et l'initialiser a zero // Tip: Il y a une methode dedier pour malloquer un int cote device et l'initialiser a zero
// //
// GM::mallocInt0(&ptrSumGM); GM::mallocInt0(&ptrSumGM);
} }
ReductionAddIntI::~ReductionAddIntI() ReductionAddIntI::~ReductionAddIntI() {
{
// TODO ReductionAddIntI // TODO ReductionAddIntI
} GM::free(ptrSumGM);
}
/*--------------------------------------*\ /*--------------------------------------*\
|* Methode *| |* Methode *|
\*-------------------------------------*/ \*-------------------------------------*/
void ReductionAddIntI::run() void ReductionAddIntI::run() {
{
// TODO ReductionAddIntI // TODO ReductionAddIntI
// appeler le kernel // appeler le kernel
// recuperer le resulat coter host // recuperer le resulat coter host
KAddIntProtocoleI<<<dg,db,this->sizeSM>>>(ptrSumGM);
GM::memcpyDToH_int(ptrSum, ptrSumGM);
// Tip: Il y a une methode dedier ramener coter host un int // Tip: Il y a une methode dedier ramener coter host un int
// //
// GM::memcpyDtoH_int(ptrDestination, ptrSourceGM);); // GM::memcpyDtoH_int(ptrDestination, ptrSourceGM););

3
Student_Cuda_Tools_Reduction/src/core/02_use_protocole/add/int/PI/host/ReductionAddIntI.h
View File

@@ -8,8 +8,7 @@
|* Declaration *| |* Declaration *|
\*---------------------------------------------------------------------*/ \*---------------------------------------------------------------------*/
class ReductionAddIntI: public RunnableGPU class ReductionAddIntI: public RunnableGPU {
{
/*--------------------------------------*\ /*--------------------------------------*\
|* Constructor *| |* Constructor *|
\*-------------------------------------*/ \*-------------------------------------*/

30
Student_Cuda_Tools_Reduction/src/core/02_use_protocole/add/int/PII/device/reductionAddIntII_device.cu
View File

@@ -1,4 +1,4 @@
#include "Thread1D.cu.h" #include "Thread2D.cu.h"
#include "cudas.h" #include "cudas.h"
#include "ReductionAdd.cu.h" #include "ReductionAdd.cu.h"
@@ -9,7 +9,9 @@
|* Private *| |* Private *|
\*---------------------------------------------------------------------*/ \*---------------------------------------------------------------------*/
static __device__ void reductionIntraThread(int* tabSM); static
__device__
void reductionIntraThread(int* tabSM);
/*----------------------------------------------------------------------*\ /*----------------------------------------------------------------------*\
|* Implementation *| |* Implementation *|
@@ -18,10 +20,16 @@ static __device__ void reductionIntraThread(int* tabSM);
/** /**
* TID partout en tabSM * TID partout en tabSM
*/ */
__global__ void KAddIntProtocoleII(int* ptrSumGM) __global__
{ void KAddIntProtocoleII(int* ptrSumGM) {
// TODO ReductionAddIntII // TODO ReductionAddIntII
}
extern __shared__ int tabSM[];
reductionIntraThread(tabSM);
__syncthreads();
ReductionAdd::reduce(tabSM, ptrSumGM);
}
/*--------------------------------------*\ /*--------------------------------------*\
|* Private *| |* Private *|
@@ -30,12 +38,16 @@ __global__ void KAddIntProtocoleII(int* ptrSumGM)
/** /**
* TID partout en tabSM * TID partout en tabSM
*/ */
__device__ void reductionIntraThread(int* tabSM) __device__
{ void reductionIntraThread(int* tabSM) {
// TODO ReductionAddIntII // TODO ReductionAddIntII
} const int TID = Thread2D::tid();
const int localTID = Thread2D::tidLocal();
tabSM[localTID] = TID;
}
/*----------------------------------------------------------------------*\ /*----------------------------------------------------------------------*\
|* End *| |* End *|
\*---------------------------------------------------------------------*/ \*---------------------------------------------------------------------*/

27
Student_Cuda_Tools_Reduction/src/core/02_use_protocole/add/int/PII/host/ReductionAddIntII.cu
View File

@@ -14,7 +14,9 @@ using std::to_string;
|* Imported *| |* Imported *|
\*---------------------------------------------------------------------*/ \*---------------------------------------------------------------------*/
extern __global__ void KAddIntProtocoleII(int* ptrSumGM); extern
__global__
void KAddIntProtocoleII(int* ptrSumGM);
/*----------------------------------------------------------------------*\ /*----------------------------------------------------------------------*\
|* Implementation *| |* Implementation *|
@@ -27,25 +29,28 @@ extern __global__ void KAddIntProtocoleII(int* ptrSumGM);
ReductionAddIntII::ReductionAddIntII(const Grid& grid , int* ptrSum,bool isVerbose) : ReductionAddIntII::ReductionAddIntII(const Grid& grid , int* ptrSum,bool isVerbose) :
//RunnableGPU(grid, "Reduce_AddInt_II_" + to_string(grid.threadCounts()),isVerbose), // classe parente //RunnableGPU(grid, "Reduce_AddInt_II_" + to_string(grid.threadCounts()),isVerbose), // classe parente
RunnableGPU(grid, "Reduce_AddInt_II",isVerbose), // classe parente RunnableGPU(grid, "Reduce_AddInt_II",isVerbose), // classe parente
ptrSum(ptrSum) ptrSum(ptrSum) {
{
// TODO ReductionAddIntII // TODO ReductionAddIntII
this->sizeSM = -1; this->sizeSM = grid.threadByBlock() * sizeof(int);
}
ReductionAddIntII::~ReductionAddIntII() GM::mallocInt0(&ptrSumGM);
{ }
ReductionAddIntII::~ReductionAddIntII() {
// TODO ReductionAddIntII // TODO ReductionAddIntII
} GM::free(ptrSumGM);
}
/*--------------------------------------*\ /*--------------------------------------*\
|* Methode *| |* Methode *|
\*-------------------------------------*/ \*-------------------------------------*/
void ReductionAddIntII::run() void ReductionAddIntII::run() {
{
// TODO ReductionAddIntII // TODO ReductionAddIntII
}
KAddIntProtocoleII<<<dg,db, this->sizeSM>>>(ptrSumGM);
GM::memcpyDToH_int(ptrSum, ptrSumGM);
}
/*----------------------------------------------------------------------*\ /*----------------------------------------------------------------------*\
|* End *| |* End *|

5
Student_Cuda_Tools_Reduction/src/core/02_use_protocole/add/int/PII/host/ReductionAddIntII.h
View File

@@ -8,8 +8,7 @@
|* Declaration *| |* Declaration *|
\*---------------------------------------------------------------------*/ \*---------------------------------------------------------------------*/
class ReductionAddIntII: public RunnableGPU class ReductionAddIntII: public RunnableGPU {
{
/*--------------------------------------*\ /*--------------------------------------*\
|* Constructor *| |* Constructor *|
\*-------------------------------------*/ \*-------------------------------------*/
@@ -44,7 +43,7 @@ class ReductionAddIntII: public RunnableGPU
int* ptrSumGM; int* ptrSumGM;
size_t sizeSM; size_t sizeSM;
}; };
/*----------------------------------------------------------------------*\ /*----------------------------------------------------------------------*\
|* End *| |* End *|

44
Student_Cuda_Tools_Reduction/src/core/02_use_protocole/generic/int/PI/device/reductionIntI_device.cu
View File

@@ -1,4 +1,4 @@
#include "Thread1D.cu.h" #include "Thread2D.cu.h"
#include "cudas.h" #include "cudas.h"
#include "Reduction.cu.h" #include "Reduction.cu.h"
@@ -23,10 +23,16 @@ static __device__ void addAtomicV2(int* ptrX , int y);
/** /**
* 1 partout en tabSM * 1 partout en tabSM
*/ */
__global__ void KIntProtocoleI(int* ptrSumGM) __global__
{ void KIntProtocoleI(int* ptrSumGM) {
// TODO ReductionIntI // TODO ReductionIntI
}
extern __shared__ int tabSM[];
reductionIntraThread(tabSM);
__syncthreads();
Reduction::reduce(add, addAtomicV1, tabSM, ptrSumGM);
}
/*--------------------------------------*\ /*--------------------------------------*\
|* Private *| |* Private *|
@@ -35,29 +41,33 @@ __global__ void KIntProtocoleI(int* ptrSumGM)
/** /**
* 1 partout en tabSM * 1 partout en tabSM
*/ */
__device__ void reductionIntraThread(int* tabSM) __device__
{ void reductionIntraThread(int* tabSM) {
// TODO ReductionIntI // TODO ReductionIntI
} const int tidLocal = Thread2D::tidLocal();
tabSM[tidLocal] = 1;
}
/*----------------------------*\ /*----------------------------*\
|* Operateur reduction *| |* Operateur reduction *|
\*---------------------------*/ \*---------------------------*/
__inline__ __inline__
__device__ int add(int x , int y) __device__
{ int add(int x , int y) {
// TODO ReductionIntI // TODO ReductionIntI
} return x + y;
}
/** /**
* Utiliser la methode system : atomicAdd(pointeurDestination, valeurSource); * Utiliser la methode system : atomicAdd(pointeurDestination, valeurSource);
*/ */
__inline__ __inline__
__device__ void addAtomicV1(int* ptrX , int y) __device__
{ void addAtomicV1(int* ptrX , int y) {
// TODO ReductionIntI // TODO ReductionIntI
} atomicAdd(ptrX, y);
}
/** /**
* 10x plus lent,mais plus flexible! * 10x plus lent,mais plus flexible!
@@ -65,17 +75,17 @@ __device__ void addAtomicV1(int* ptrX , int y)
* Necessaire aussi pour des objets par exemple * Necessaire aussi pour des objets par exemple
*/ */
__device__ int volatile mutex = 0; //variable global __device__ int volatile mutex = 0; //variable global
__device__ void addAtomicV2(int* ptrX , int y) __device__
{ void addAtomicV2(int* ptrX , int y) {
Lock locker(&mutex); Lock locker(&mutex);
locker.lock(); locker.lock();
// TODO ReductionIntI // TODO ReductionIntI
*ptrX = *ptrX + y;
locker.unlock(); locker.unlock();
} }
/*----------------------------------------------------------------------*\ /*----------------------------------------------------------------------*\
|* End *| |* End *|
\*---------------------------------------------------------------------*/ \*---------------------------------------------------------------------*/

27
Student_Cuda_Tools_Reduction/src/core/02_use_protocole/generic/int/PI/host/ReductionIntI.cu
View File

@@ -14,7 +14,9 @@ using std::to_string;
|* Imported *| |* Imported *|
\*---------------------------------------------------------------------*/ \*---------------------------------------------------------------------*/
extern __global__ void KIntProtocoleI(int* ptrSumGM); extern
__global__
void KIntProtocoleI(int* ptrSumGM);
/*----------------------------------------------------------------------*\ /*----------------------------------------------------------------------*\
|* Implementation *| |* Implementation *|
@@ -27,25 +29,28 @@ extern __global__ void KIntProtocoleI(int* ptrSumGM);
ReductionIntI::ReductionIntI(const Grid& grid , int* ptrSum,bool isVerbose) : ReductionIntI::ReductionIntI(const Grid& grid , int* ptrSum,bool isVerbose) :
//RunnableGPU(grid, "Redude_Generic_IntI_" + to_string(grid.threadCounts()),isVerbose), // classe parente //RunnableGPU(grid, "Redude_Generic_IntI_" + to_string(grid.threadCounts()),isVerbose), // classe parente
RunnableGPU(grid, "Reduce_Generic_IntI",isVerbose), // classe parente RunnableGPU(grid, "Reduce_Generic_IntI",isVerbose), // classe parente
ptrSum(ptrSum) ptrSum(ptrSum) {
{
// TODO ReductionIntI // TODO ReductionIntI
this->sizeSM = -1; this->sizeSM = grid.threadByBlock() * sizeof(int);
}
ReductionIntI::~ReductionIntI() GM::mallocInt0(&ptrSumGM);
{ }
ReductionIntI::~ReductionIntI() {
// TODO ReductionIntI // TODO ReductionIntI
} GM::free(ptrSumGM);
}
/*--------------------------------------*\ /*--------------------------------------*\
|* Methode *| |* Methode *|
\*-------------------------------------*/ \*-------------------------------------*/
void ReductionIntI::run() void ReductionIntI::run() {
{
// TODO ReductionIntI // TODO ReductionIntI
}
KIntProtocoleI<<<dg,db,this->sizeSM>>>(ptrSumGM);
GM::memcpyDToH_int(ptrSum, ptrSumGM);
}
/*----------------------------------------------------------------------*\ /*----------------------------------------------------------------------*\
|* End *| |* End *|

5
Student_Cuda_Tools_Reduction/src/core/02_use_protocole/generic/int/PI/host/ReductionIntI.h
View File

@@ -8,8 +8,7 @@
|* Declaration *| |* Declaration *|
\*---------------------------------------------------------------------*/ \*---------------------------------------------------------------------*/
class ReductionIntI: public RunnableGPU class ReductionIntI: public RunnableGPU {
{
/*--------------------------------------*\ /*--------------------------------------*\
|* Constructor *| |* Constructor *|
\*-------------------------------------*/ \*-------------------------------------*/
@@ -44,7 +43,7 @@ class ReductionIntI: public RunnableGPU
int* ptrSumGM; int* ptrSumGM;
size_t sizeSM; size_t sizeSM;
}; };
/*----------------------------------------------------------------------*\ /*----------------------------------------------------------------------*\
|* End *| |* End *|

46
Student_Cuda_Tools_Reduction/src/core/02_use_protocole/generic/int/PII/device/reductionIntII_device.cu
View File

@@ -1,4 +1,4 @@
#include "Thread1D.cu.h" #include "Thread2D.cu.h"
#include "cudas.h" #include "cudas.h"
#include "Reduction.cu.h" #include "Reduction.cu.h"
@@ -24,10 +24,16 @@ static __device__ void addAtomicV2(int* ptrX , int y);
/** /**
* TID partout en tabSM * TID partout en tabSM
*/ */
__global__ void KIntProtocoleII(int* ptrSumGM) __global__
{ void KIntProtocoleII(int* ptrSumGM) {
// TODO ReductionIntII // TODO ReductionIntII
}
extern __shared__ int tabSM[];
reductionIntraThread(tabSM);
__syncthreads();
Reduction::reduce(add, addAtomicV1, tabSM, ptrSumGM);
}
/*--------------------------------------*\ /*--------------------------------------*\
|* Private *| |* Private *|
@@ -36,46 +42,52 @@ __global__ void KIntProtocoleII(int* ptrSumGM)
/** /**
* TID partout en tabSM * TID partout en tabSM
*/ */
__device__ void reductionIntraThread(int* tabSM) __device__
{ void reductionIntraThread(int* tabSM) {
// TODO ReductionIntII // TODO ReductionIntII
} const int TID = Thread2D::tid();
const int tidLocal = Thread2D::tidLocal();
tabSM[tidLocal] = TID;
}
/*----------------------------*\ /*----------------------------*\
|* Operateur reduction *| |* Operateur reduction *|
\*---------------------------*/ \*---------------------------*/
__inline__ __inline__
__device__ int add(int x , int y) __device__
{ int add(int x , int y) {
// TODO ReductionIntII // TODO ReductionIntII
} return x + y;
}
/** /**
* Utiliser la methode system : atomicAdd(pointeurDestination, valeurSource); * Utiliser la methode system : atomicAdd(pointeurDestination, valeurSource);
*/ */
__inline__ __inline__
__device__ void addAtomicV1(int* ptrX , int y) __device__
{ void addAtomicV1(int* ptrX , int y) {
// TODO ReductionIntII // TODO ReductionIntII
} atomicAdd(ptrX, y);
}
/** /**
* Une alternative, moins performante, mais generalisable serait d'employer un lock * Une alternative, moins performante, mais generalisable serait d'employer un lock
* Tip : le Lock est implementer avec deux methodes atomic * Tip : le Lock est implementer avec deux methodes atomic
*/ */
__device__ int volatile mutex = 0; //variable global __device__ int volatile mutex = 0; //variable global
__device__ void addAtomicV2(int* ptrX , int y) __device__
{ void addAtomicV2(int* ptrX , int y) {
Lock locker(&mutex); Lock locker(&mutex);
locker.lock(); locker.lock();
// TODO ReductionIntII // TODO ReductionIntII
*ptrX = *ptrX + y;
locker.unlock(); locker.unlock();
} }
/*----------------------------------------------------------------------*\ /*----------------------------------------------------------------------*\
|* End *| |* End *|
\*---------------------------------------------------------------------*/ \*---------------------------------------------------------------------*/

24
Student_Cuda_Tools_Reduction/src/core/02_use_protocole/generic/int/PII/host/ReductionIntII.cu Executable file → Normal file
View File

@@ -27,25 +27,29 @@ extern __global__ void KIntProtocoleII(int* ptrSumGM);
ReductionIntII::ReductionIntII(const Grid& grid , int* ptrSum,bool isVerbose) : ReductionIntII::ReductionIntII(const Grid& grid , int* ptrSum,bool isVerbose) :
//RunnableGPU(grid, "Reduce_Generic_IntII_" + to_string(grid.threadCounts()),isVerbose), // classe parente //RunnableGPU(grid, "Reduce_Generic_IntII_" + to_string(grid.threadCounts()),isVerbose), // classe parente
RunnableGPU(grid, "Reduce_Generic_IntII" ,isVerbose), // classe parente RunnableGPU(grid, "Reduce_Generic_IntII" ,isVerbose), // classe parente
ptrSum(ptrSum) ptrSum(ptrSum) {
{
// TODO ReductionIntII // TODO ReductionIntII
this->sizeSM = -1; this->sizeSM = grid.threadByBlock() * sizeof(int);
}
ReductionIntII::~ReductionIntII() GM::mallocInt0(&ptrSumGM);
{ }
ReductionIntII::~ReductionIntII() {
// TODO ReductionIntII // TODO ReductionIntII
}
GM::free(ptrSumGM);
}
/*--------------------------------------*\ /*--------------------------------------*\
|* Methode *| |* Methode *|
\*-------------------------------------*/ \*-------------------------------------*/
void ReductionIntII::run() void ReductionIntII::run() {
{
// TODO ReductionIntII // TODO ReductionIntII
}
KIntProtocoleII<<<dg, db, this->sizeSM>>>(ptrSumGM);
GM::memcpyDToH_int(ptrSum, ptrSumGM);
}
/*----------------------------------------------------------------------*\ /*----------------------------------------------------------------------*\
|* End *| |* End *|

5
Student_Cuda_Tools_Reduction/src/core/02_use_protocole/generic/int/PII/host/ReductionIntII.h
View File

@@ -8,8 +8,7 @@
|* Declaration *| |* Declaration *|
\*---------------------------------------------------------------------*/ \*---------------------------------------------------------------------*/
class ReductionIntII: public RunnableGPU class ReductionIntII: public RunnableGPU {
{
/*--------------------------------------*\ /*--------------------------------------*\
|* Constructor *| |* Constructor *|
\*-------------------------------------*/ \*-------------------------------------*/
@@ -44,7 +43,7 @@ class ReductionIntII: public RunnableGPU
int* ptrSumGM; int* ptrSumGM;
size_t sizeSM; size_t sizeSM;
}; };
/*----------------------------------------------------------------------*\ /*----------------------------------------------------------------------*\
|* End *| |* End *|

43
Student_Cuda_Tools_Reduction/src/core/02_use_protocole/generic/long/PII/device/reductionLongII_device.cu
View File

@@ -1,4 +1,4 @@
#include "Thread1D_long.cu.h" #include "Thread2D_long.cu.h"
#include "cudas.h" #include "cudas.h"
#include "Reduction.cu.h" #include "Reduction.cu.h"
@@ -23,10 +23,16 @@ static __device__ void addAtomic(long* ptrX , long y);
/** /**
* TID partout en tabSM * TID partout en tabSM
*/ */
__global__ void KLongProtocoleII(long* ptrSumGM) __global__
{ void KLongProtocoleII(long* ptrSumGM) {
// TODO ReductionLongII // TODO ReductionLongII
} //
extern __shared__ long tabSM[];
reductionIntraThread(tabSM);
__syncthreads();
Reduction::reduce(add, addAtomic, tabSM, ptrSumGM);
}
/*--------------------------------------*\ /*--------------------------------------*\
|* Private *| |* Private *|
@@ -35,8 +41,9 @@ __global__ void KLongProtocoleII(long* ptrSumGM)
/** /**
* TID partout en tabSM * TID partout en tabSM
*/ */
__device__ void reductionIntraThread(long* tabSM) __device__
{ // Rappel : Dans le protocoleII on cherche a calculer void reductionIntraThread(long* tabSM) {
// Rappel : Dans le protocoleII on cherche a calculer
// //
// x=x+i avec i in [0,N] // x=x+i avec i in [0,N]
// //
@@ -70,7 +77,12 @@ __device__ void reductionIntraThread(long* tabSM)
// pour TID utiliser const long TID=Thread2D_long::tid(); // (nouvelle methode) // pour TID utiliser const long TID=Thread2D_long::tid(); // (nouvelle methode)
// pour TID_LOCAL utiliser const int TID_LOCAL=Thread2D::tidLocal(); // (methode habituelle) // pour TID_LOCAL utiliser const int TID_LOCAL=Thread2D::tidLocal(); // (methode habituelle)
} //
const long TID=Thread2D_long::tid();
const int TID_LOCAL=Thread2D::tidLocal();
tabSM[TID_LOCAL] = TID;
}
/*----------------------------*\ /*----------------------------*\
@@ -78,10 +90,11 @@ __device__ void reductionIntraThread(long* tabSM)
\*---------------------------*/ \*---------------------------*/
__inline__ __inline__
__device__ long add(long x , long y) __device__
{ long add(long x , long y) {
// TODO ReductionLongII // TODO ReductionLongII
} return x + y;
}
/** /**
* Utiliser la methode system, si elle existe * Utiliser la methode system, si elle existe
@@ -90,20 +103,20 @@ __device__ long add(long x , long y)
* *
* ou la technique du lock vu precedement! * ou la technique du lock vu precedement!
* *
* Question : atomicAdd pour les long existe? * Question : atomicAdd pour les long existe? non ;(
*/ */
__device__ int volatile mutex = 0; //variable global __device__ int volatile mutex = 0; //variable global
__device__ void addAtomic(long* ptrX , long y) __device__
{ void addAtomic(long* ptrX , long y) {
Lock locker(&mutex); Lock locker(&mutex);
locker.lock(); locker.lock();
// TODO ReductionLongII // TODO ReductionLongII
*ptrX = *ptrX + y;
locker.unlock(); locker.unlock();
} }
/*----------------------------------------------------------------------*\ /*----------------------------------------------------------------------*\
|* End *| |* End *|
\*---------------------------------------------------------------------*/ \*---------------------------------------------------------------------*/

22
Student_Cuda_Tools_Reduction/src/core/02_use_protocole/generic/long/PII/host/ReductionLongII.cu
View File

@@ -27,25 +27,27 @@ extern __global__ void KLongProtocoleII(long* ptrSumGM);
ReductionLongII::ReductionLongII(const Grid& grid , long* ptrSum,bool isVerbose) : ReductionLongII::ReductionLongII(const Grid& grid , long* ptrSum,bool isVerbose) :
//RunnableGPU(grid, "Reduce_Generic_LongII_" + to_string(grid.threadCounts()),isVerbose), // classe parente //RunnableGPU(grid, "Reduce_Generic_LongII_" + to_string(grid.threadCounts()),isVerbose), // classe parente
RunnableGPU(grid, "Reduce_Generic_LongII",isVerbose), // classe parente RunnableGPU(grid, "Reduce_Generic_LongII",isVerbose), // classe parente
ptrSum(ptrSum) ptrSum(ptrSum) {
{
// TODO ReductionLongII // TODO ReductionLongII
this->sizeSM = -1; this->sizeSM = grid.threadByBlock() * sizeof(long);
}
ReductionLongII::~ReductionLongII() GM::mallocLong0(&ptrSumGM);
{ }
ReductionLongII::~ReductionLongII() {
// TODO ReductionLongII // TODO ReductionLongII
} GM::free(ptrSumGM);
}
/*--------------------------------------*\ /*--------------------------------------*\
|* Methode *| |* Methode *|
\*-------------------------------------*/ \*-------------------------------------*/
void ReductionLongII::run() void ReductionLongII::run() {
{
// TODO ReductionLongII // TODO ReductionLongII
} KLongProtocoleII<<<dg, db, this->sizeSM>>>(ptrSumGM);
GM::memcpyDToH_long(ptrSum, ptrSumGM);
}
/*----------------------------------------------------------------------*\ /*----------------------------------------------------------------------*\
|* End *| |* End *|

5
Student_Cuda_Tools_Reduction/src/core/02_use_protocole/generic/long/PII/host/ReductionLongII.h
View File

@@ -8,8 +8,7 @@
|* Declaration *| |* Declaration *|
\*---------------------------------------------------------------------*/ \*---------------------------------------------------------------------*/
class ReductionLongII: public RunnableGPU class ReductionLongII: public RunnableGPU {
{
/*--------------------------------------*\ /*--------------------------------------*\
|* Constructor *| |* Constructor *|
\*-------------------------------------*/ \*-------------------------------------*/
@@ -44,7 +43,7 @@ class ReductionLongII: public RunnableGPU
long* ptrSumGM; long* ptrSumGM;
size_t sizeSM; size_t sizeSM;
}; };
/*----------------------------------------------------------------------*\ /*----------------------------------------------------------------------*\
|* End *| |* End *|

10
Student_Cuda_Tools_Reduction/src/main/main.cpp
View File

@@ -18,16 +18,15 @@ extern int mainTest();
|* Implementation *| |* Implementation *|
\*---------------------------------------------------------------------*/ \*---------------------------------------------------------------------*/
int main(int argc , char** argv) int main(int argc , char** argv) {
{
// Limits::show(); // Limits::show();
CudaContext cudaContext; CudaContext cudaContext;
// public // public
{ {
cudaContext.deviceId = 0; // in [0,2] width Server Cuda3 cudaContext.deviceId = 1; // in [0,2] width Server Cuda3
cudaContext.launchMode = LaunchModeMOO::USE; // USE TEST (only) cudaContext.launchMode = LaunchModeMOO::TEST; // USE TEST (only)
cudaContext.deviceDriver = DeviceDriver::LOAD_ALL; // LOAD_CURRENT LOAD_ALL cudaContext.deviceDriver = DeviceDriver::LOAD_ALL; // LOAD_CURRENT LOAD_ALL
cudaContext.deviceInfo = DeviceInfo::ALL_SIMPLE; // NONE ALL ALL_SIMPLE CURRENT cudaContext.deviceInfo = DeviceInfo::ALL_SIMPLE; // NONE ALL ALL_SIMPLE CURRENT
@@ -40,9 +39,8 @@ int main(int argc , char** argv)
} }
return cudaContext.process(); return cudaContext.process();
} }
/*----------------------------------------------------------------------*\ /*----------------------------------------------------------------------*\
|* End *| |* End *|
\*---------------------------------------------------------------------*/ \*---------------------------------------------------------------------*/

24
Student_Cuda_Tools_Reduction/src/main/mainTest.cpp
View File

@@ -26,15 +26,14 @@ static void generic();
|* Implementation *| |* Implementation *|
\*---------------------------------------------------------------------*/ \*---------------------------------------------------------------------*/
int mainTest() int mainTest() {
{
// activer ci-dessous seulement le TP voulu (pas tous) // activer ci-dessous seulement le TP voulu (pas tous)
add(); add();
//generic(); generic();
return EXIT_SUCCESS; return EXIT_SUCCESS;
} }
/*--------------------------------------*\ /*--------------------------------------*\
|* private *| |* private *|
@@ -43,30 +42,27 @@ int mainTest()
/** /**
* activer ci-dessous la version souhaiter * activer ci-dessous la version souhaiter
*/ */
void add() void add() {
{
VTReductionAddIntI test1; VTReductionAddIntI test1;
VTReductionAddIntII test2; VTReductionAddIntII test2;
test1.run(); test1.run();
//test2.run(); test2.run();
} }
/** /**
* activer ci-dessous la version souhaiter * activer ci-dessous la version souhaiter
*/ */
void generic() void generic() {
{
VTReductionGenericI test1; VTReductionGenericI test1;
VTReductionGenericII test2; VTReductionGenericII test2;
VTReductionGenericLongII test3; VTReductionGenericLongII test3;
test1.run(); test1.run();
// test2.run(); test2.run();
// test3.run(); test3.run();
} }
/*----------------------------------------------------------------------*\ /*----------------------------------------------------------------------*\
|* End *| |* End *|
\*---------------------------------------------------------------------*/ \*---------------------------------------------------------------------*/

31
Student_Cuda_Tools_Reduction/src/main/mainUse.cpp
View File

@@ -31,19 +31,18 @@ static void print(bool isSuccess);
static const int IS_VERBOSE = true; static const int IS_VERBOSE = true;
int mainUse() int mainUse() {
{
// activer ci-dessous seulement le TP voulu (pas tous) // activer ci-dessous seulement le TP voulu (pas tous)
bool isOk = true; bool isOk = true;
reduction_add(isOk); // voir code ci-dessous pour activer la version voulue reduction_add(isOk); // voir code ci-dessous pour activer la version voulue
//reduction_generic(isOk); // voir code ci-dessous pour activer la version voulue reduction_generic(isOk); // voir code ci-dessous pour activer la version voulue
print(isOk); print(isOk);
return isOk ? EXIT_SUCCESS : EXIT_FAILURE; return isOk ? EXIT_SUCCESS : EXIT_FAILURE;
} }
/*----------------------------------------------------------------------*\ /*----------------------------------------------------------------------*\
|* TP *| |* TP *|
@@ -52,8 +51,7 @@ int mainUse()
/** /**
* activer ci-dessous la version souhaiter * activer ci-dessous la version souhaiter
*/ */
void reduction_add(bool& isOk) void reduction_add(bool& isOk) {
{
// InbI // InbI
{ {
UseReductionAddIntI algo(IS_VERBOSE); UseReductionAddIntI algo(IS_VERBOSE);
@@ -61,17 +59,16 @@ void reduction_add(bool& isOk)
} }
// IntII // IntII
// { {
// UseReductionAddIntII algo(IS_VERBOSE); UseReductionAddIntII algo(IS_VERBOSE);
// isOk &= algo.isOk(IS_VERBOSE); isOk &= algo.isOk(IS_VERBOSE);
// } }
} }
/** /**
* activer ci-dessous la version souhaiter * activer ci-dessous la version souhaiter
*/ */
void reduction_generic(bool& isOk) void reduction_generic(bool& isOk) {
{
// InbI // InbI
{ {
UseReductionIntI algo(IS_VERBOSE); UseReductionIntI algo(IS_VERBOSE);
@@ -89,22 +86,20 @@ void reduction_generic(bool& isOk)
UseReductionLongII algo(IS_VERBOSE); UseReductionLongII algo(IS_VERBOSE);
isOk &= algo.isOk(IS_VERBOSE); isOk &= algo.isOk(IS_VERBOSE);
} }
} }
/*----------------------------------------------------------------------*\ /*----------------------------------------------------------------------*\
|* Tools *| |* Tools *|
\*---------------------------------------------------------------------*/ \*---------------------------------------------------------------------*/
void print(bool isSuccess) void print(bool isSuccess) {
{
cout << endl << Couts::REVERSE; cout << endl << Couts::REVERSE;
Couts::status(isSuccess, "Success, Congratulations !", "Failed, sorry!"); Couts::status(isSuccess, "Success, Congratulations !", "Failed, sorry!");
cout << endl << Couts::RESET; cout << endl << Couts::RESET;
} }
/*----------------------------------------------------------------------*\ /*----------------------------------------------------------------------*\
|* End *| |* End *|
\*---------------------------------------------------------------------*/ \*---------------------------------------------------------------------*/