Poniżej przedstawiony został kod programu realizującego analogiczną
Ope
funkcjonalność, tj. wyznaczającego sumę elementów wektora, w środowisku
nC
OpenCL. Kod zostanie przytoczony w całości, natomiast omówienie ogra-
L
niczy się jedynie do wskazania różnic pomiędzy oboma środowiskami. Na
listingu 3.7 przedstawiona została funkcja kernela realizująca w OpenCL
analogiczne zadanie do funkcji CUDA z listingu 3.5.
Listing 3.7. OpenCL – Algorytm redukcji z sumą – funkcja kernela.
1
# define N_THREADS 256
2
__kernel void reduce ( __global
float * vec ,
3
__global
float * vec_out ,
4
int size )
5
{
6
__local float cache [ N_THREADS ];
7
float sum = 0.0 f;
8
9
int idx = get_global_id (0) ;
10
for ( int i= idx ; i < size ; i += get_global_size (0) )
11
sum += vec [i];
78
3. Model pamięci GPGPU
12
13
cache [ get_local_id (0) ] = sum ;
14
barrier ( CLK_LOCAL_MEM_FENCE );
15
16
for ( int k= get_local_size (0) /2; k; k /=2)
17
{
18
if( get_local_id (0) < k)
19
cache [ get_local_id (0) ] += cache [ get_local_id (0) +k ];
20
barrier ( CLK_LOCAL_MEM_FENCE );
21
}
22
23
if ( get_local_id (0) ==0)
24
vec_out [ get_group_id (0) ] = cache [0];
25
}
W środowisku OpenCL pamięć współdzielona nazywana jest lokalną a
deklaracja zmiennej umieszczonej w tej pamięci musi zawierać specyfikator
__local.
Punktem synchronizacji work-items w obrębie work-group jest funkcja:
void barrier ( cl_mem_fence_flags flags )
Parametr
flags
może
dowolną
kombinację
dwóch
flag:
(1) CLK_LOCAL_MEM_FENCE, która zapewnia spójność zmiennych znajdują-
cych się w pamięci lokalnej oraz (2) CLK_GLOBAL_MEM_FENCE, która zapewnia
spójność pamięci globalnej.
Listing 3.8. OpenCL – Algorytm redukcji z sumą.
1
# include <CL / opencl .h >
2
3
cl_platform_id
platform ;
4
cl_device_id
device ;
5
cl_context
context ;
6
cl_command_queue cmdQueue ;
7
cl_program
hProgram ;
8
cl_kernel
hKernel ;
9
10
const int SIZE
= 67108864;
11
const int N_THREADS = 256;
12
const int N_BLOCKS
= 256;
13
size_t GLOBAL_WS []
= { N_THREADS * N_BLOCKS };
14
size_t LOCAL_WS []
= { N_THREADS };
15
16
int main ( int argv , char * argc [])
17
{
18
float * vec = new float [ SIZE ];
19
float
vec_out [ N_BLOCKS ];
20
cl_mem cl_vec = 0;
3.2. Wykorzystanie pamięci współdzielonej do optymalizacji dostępu do pamięci
urządzenia
79
21
cl_mem cl_vec_out = 0;
22
double time , time2 ;
23
float
sum =0;
24
25
for ( int i =0; i < SIZE ; i ++)
26
vec [i] = 1.0 f -2.0 f* rand () / RAND_MAX ;
27
28
clGetPlatformIDs (1 , & platform , NULL )
29
cl_uint num_dev ;
30
clGetDeviceIDs ( platform , CL_DEVICE_TYPE_GPU , 1, & device ,
31
& num_dev );
32
context = clCreateContext (0 , 1, & device , 0 ,0 ,0) ;
33
cmdQueue = clCreateCommandQueue ( context , device , 0 ,0) ;
34
35
size_t kernelLength ;
36
char * programSource = loadProgSource (" reduce .cl " , "" ,
37
& kernelLength );
38
cmdQueue = clCreateCommandQueue ( context , devices , 0 ,0) ;
39
hProgram = clCreateProgramWithSource ( context , 1,
40
( const char **) & programSource , & kernelLength , 0) ;
41
42
clBuildProgram ( hProgram , 0, 0, 0, 0, 0) ;
43
hKernel = clCreateKernel ( hProgram , " reduce " , 0) ;
44
45
cl_vec = clCreateBuffer ( context , CL_MEM_READ_ONLY |
46
CL_MEM_COPY_HOST_PTR , SIZE * sizeof ( float ), vec , 0) ;
47
cl_vec_out = clCreateBuffer ( context , CL_MEM_WRITE_ONLY ,
48
N_THREADS * sizeof ( float ), 0 ,0) ;
49
50
time = timeStamp () ;
51
clSetKernelArg ( hKernel , 0, sizeof ( cl_mem ), & cl_vec );
52
clSetKernelArg ( hKernel , 1, sizeof ( cl_mem ), & cl_vec_out );
53
clSetKernelArg ( hKernel , 2, sizeof ( int ) , & SIZE );
54
clEnqueueNDRangeKernel ( cmdQueue , hKernel , 1, 0,
55
GLOBAL_WS , LOCAL_WS , 0 ,0 ,0) ;
56
clEnqueueReadBuffer ( cmdQueue , cl_vec_out , CL_TRUE , 0,
57
N_BLOCKS * sizeof ( float ) , vec_out , 0 ,0 ,0) ;
58
clFinish ( cmdQueue );
59
60
for ( int i =0; i < N_BLOCKS ; i ++)
61
sum += vec_out [i];
62
time2 = timeStamp () ;
63
64
cout <<" GPU sum =" << sum <<" , time =" << time2 - time <<"[ ms ]" << endl ; 65
66
clReleaseMemObject ( cl_vec );
67
clReleaseMemObject ( cl_vec_out );
68
delete [] vec ;
69
70
return 0;
71
}
80
3. Model pamięci GPGPU
Po inicjalizacji wektorów oraz samego środowiska OpenCL, w pamięci
globalnej zostały utworzone w liniach 45–48 GPU dwa wektory: wejścio-
wy cu_vec oraz wyjściowy cu_vec_out. Pierwszy wektor został zadeklaro-
wany jako tylko-do-odczytu CL_MEM_READ_ONLY a drugi jako tylko-do-zapisu
CL_MEM_WRITE_ONLY.
W liniach 50–62 przeprowadzony został właściwy test redukujący
wszystkie elementy wektora do 256 sum cząstkowych za pomocą GPU i
następnie obliczający całkowitą sumę już na hoście w liniach 60–61.
Założona ilość wątków w pojedynczym bloku jest standardową, optymal-
ną ilością dla współczesnych kart graficznych, natomiast ilość bloków jest w
zasadzie dowolna a jej konkretna wartość jest podyktowana możliwościami
danego urządzenia, na którym będzie się dany kernel wykonywał. Powinna
to być wartość na tyle duża żeby obsadzić wszystkie jednostki obliczenio-
we urządzenia. Na przykład, karta grafiki, na której testowano powyższe