1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114
115
116
117
118
119
120
121
122
123
124
|
#include <stdio.h>
#include <stdlib.h>
#include <cuda.h>
#include <cuda_runtime.h>
class signal
{
bool _t0,_t1;
public:
__device__
signal()
{
}
__device__
void init(bool valeur)
{
_t0=valeur;
_t1=valeur;
}
__device__
bool t1()
{
return _t1;
}
__device__
bool t0()
{
return _t0;
}
__device__
void w(bool val)
{
_t1=val;
}
__device__
void maj()
{
_t0=_t1;
threadfence();
};
};
__global__
void kernel(int n_cycles,int t_cycle,int* retour)
{
signal s;
s.init(true);
int i;
for(i=0;i<n_cycles;i++)
{
retour[t_cycle*i]=i;
retour[t_cycle*i+1]=s.t0();
retour[t_cycle*i+2]=s.t1();
retour[t_cycle*i+3]=~s.t0();
s.w(~s.t0());
threadfence();
retour[t_cycle*i+4]=s.t1();
s.maj();
threadfence();
}
}
//============================================================host
int main()
{
int t_num_cycle =1;
int t_cycle =4+t_num_cycle;
int n_cycles =100000/t_cycle;
int t_buffer =t_cycle*n_cycles;
int t_bytes_buffer =t_buffer*sizeof(int);
FILE* pFile;
pFile = fopen ("simu2.trace","w");
int retour[t_buffer];
int* retour_device;
cudaFuncSetCacheConfig(kernel, cudaFuncCachePreferL1);
CUevent start, stop;
cudaEventCreate(&start);
cudaEventCreate(&stop);
cudaMalloc ( (void**) &retour_device, t_bytes_buffer);
cudaEventRecord(start, 0);
kernel<<<1,100>>>(n_cycles,t_cycle,retour_device);
cudaThreadSynchronize();
cudaMemcpy(retour,retour_device,t_bytes_buffer,cudaMemcpyDeviceToHost);
cudaEventRecord(stop, 0);
int j;
for(j=0;j<n_cycles;j++)
{
fprintf (pFile, "cycle %5i t0:%i t1:%i ~t0:%i t1=~t0:%i \n",retour[j*t_cycle],retour[j*t_cycle+1],retour[j*t_cycle+2],retour[j*t_cycle+3],retour[j*t_cycle+4]);
}
fclose (pFile);
cudaEventSynchronize(stop);
float elapsedTime;
cudaEventElapsedTime(&elapsedTime, start, stop);
printf("===> %f ms",elapsedTime);
printf("\nnombre de cycles: %i \ntaille d'un cycle: %i \ntaille du buffer en int: %i \ntaille du buffer en octets: %i \n" ,n_cycles,t_cycle,t_buffer,t_bytes_buffer);
return 0;
} |
Exemple simple qui ne marche pas :(
,
Seulement, ce changement n'est tout simplement pas visible.
Répondre avec citation
Envoyé par kiluak
Partager