0votos

Número de caminos en una cuadrícula en C

por josejuan usando OpenCL hace 6 años

Usando OpenCL (ver notas sobre rendimiento). Pongo ambos archivos juntos (código en C y código OpenCL) por comodidad.

Dada una cuadrícula de NxN nodos, ¿cuantos caminos diferentes hay (sin pasar dos veces por el mismo nodos) al ir de una esquina a su opuesta diagonal?

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
125
126
127
128
129
130
131
132
133
134
135
136
137
138
139
140
141
142
143
144
145
146
147
148
149
150
151
152
153
154
155
156
157
158
159
160
161
162
163
164
165
166
167
168
169
170
171
172
173
174
175
176
177
178
179
180
181
182
183
184
185
186
187
188
189
190
191
192
193
194
195
196
197
198
199
200
201
202
203
204
205
206
207
208
209
210
211
212
213
214
215
216
217
218
219
220
221
222
223
224
225
226
227
228
229
230
231
232
233
234
235
236
237
238
239
240
241
242
243
244
245
246
247
248
249
250
251
252
253
254
255
256
257
258
259
260
261
262
263
264
265
266
267
268
269
270
271
272
273
274
275
276
277
278
279
280
281
282
283
284
285
286
287
288
289
290
291
292
293
294
295
296
297
298
299
300
301
302
303
304
305
306
307
308
309
310
311
312
313
314
315
316
317
318
319
320
321
322
323
324
325
326
327
328
329
330
331
332
333
334
335
336
337
338
339
340
341
342
343
344
345
346
347
348
349
350
351
352
353
354
355
356
357
358
359
360
361
362
363
364
365
366
367
368
369
370
371
372
/* 
 
En mi máquina, tengo: 
 
* un AMD con 6 cores a 2,7Ghz. 
* una Radeon HD 6570 con 6 unidades de procesamiento con un total de 480 cores a 650MHz. 
 
A priori, usando CPU tenemos 6 * 2,7 = 16,2 GHz de procesamiento. 
     Usando GPU tenemos 480 * 0,65 = 312 GHz de procesamiento. 
 
Esto tendría que hacer un factor de unas 19 veces más rápido usando GPU que CPU (o 10, la cuestión es que debería ser significativamente más rápido). 
 
Sin embargo, el rendimiento es similar (ajustando finamente la GPU, que es difícil). 
 
El motivo es que dentro de cada unidad de procesamiento de la GPU (en total yo tengo 6) todas las unidades de procesamiento (stream processors) ejecutan las instrucciones a la vez. Dada la máquina de estado que resulta ser "runGPUCell" (con muchos saltos y bifurcaciones), tenemos una ejecución bastante pobre que fácilmente explica esa bajada de rendimiento en 19 veces para quedarse en las mismas (un poco peor, la verdad). 
 
Desenrollar la máquina de estados de "runGPUCell" no parece nada fácil, por lo que evitar las bifurcaciones no va a ser posible. 
 
Así, para sacar provecho de la GPU, ésta tiene que tener tantas unidades de procesamiento o más que CPU tengamos. Claro que en mi caso hay 80 stream processors por core, y podría ser que tener más ¡fuera peor! al haber más solapamientos y bloqueos en el flujo de instrucciones. 
 
En cualquier caso, me ha gustado esta Kata. XD XD XD 
 
*/ 
 
//=== archivo "paths.c" ============================================ 
#include <stdio.h> 
#include <string.h> 
#include <stdlib.h> 
#include <CL/cl.h> 
 
// Tamaño del problema como constante ¡debe ser igual que en "runGPUCell.cl"! 
#define SIZE 7 
#define SIZE2 (SIZE + 2) 
 
#define UP          '^' 
#define RIGHT       '>' 
#define DOWN        'v' 
#define LEFT        '<' 
#define OBJETIVE    'O' 
#define EMPTY       ' ' 
 
#define SOLID       '*' 
#define EXPLORED    '+' 
#define NULLSTATE   '#' 
 
#define COORD(x, y) (y * SIZE2 + x) 
 
#define MAX_SOURCE_SIZE (10 * 1024 * 1024) 
 
typedef struct __GPUStack__ { 
    unsigned short X; // x coord 
    unsigned short Y; // y coord 
    char S; // status 
} GPUStack; 
 
typedef struct __GPUCell__ { 
    char G[SIZE2 * SIZE2];  // [U, D, L, R] - direction, E - empty, O - objetive 
    int pstack;             // stack pointer 
    long counter;           // path counter 
    char halt;              // halted process 
    GPUStack stack[4 * SIZE2 * SIZE2];  // stack 
} GPUCell; 
 
// similar to "runGPUCell.cl", is used here to expand paths before split 
void runGPUCell(GPUCell *cell, int MAXCOUNT) { 
    if(!cell->halt) { 
        int ps = cell->pstack; 
        GPUStack s; 
        while(ps >= 0 && (s = cell->stack[ps]).S != SOLID) { 
            switch(s.S) { 
            case NULLSTATE: 
                ps--; 
                break; 
            case EXPLORED: 
                cell->G[COORD(s.X, s.Y)] = EMPTY; 
                ps--; 
                break; 
            default: { 
                char *cc = cell->G + COORD(s.X, s.Y); 
                if(*cc == OBJETIVE) { 
                    cell->counter++; 
                    ps--; 
                    if(cell->counter >= MAXCOUNT) { 
                        cell->pstack = ps; 
                        return; 
                } else 
                    if(*cc == EMPTY) { 
                        cell->stack[ps].S = EXPLORED; 
                        *cc = s.S; 
                        switch(s.S) { 
                        case UP: 
                            ps++; cell->stack[ps].S =    UP; cell->stack[ps].X = s.X; cell->stack[ps].Y = s.Y - 1; 
                            ps++; cell->stack[ps].S =  LEFT; cell->stack[ps].X = s.X - 1; cell->stack[ps].Y = s.Y; 
                            ps++; cell->stack[ps].S = RIGHT; cell->stack[ps].X = s.X + 1; cell->stack[ps].Y = s.Y; 
                            break; 
                        case DOWN: 
                            ps++; cell->stack[ps].S =  DOWN; cell->stack[ps].X = s.X; cell->stack[ps].Y = s.Y + 1; 
                            ps++; cell->stack[ps].S =  LEFT; cell->stack[ps].X = s.X - 1; cell->stack[ps].Y = s.Y; 
                            ps++; cell->stack[ps].S = RIGHT; cell->stack[ps].X = s.X + 1; cell->stack[ps].Y = s.Y; 
                            break; 
                        case LEFT: 
                            ps++; cell->stack[ps].S =   UP; cell->stack[ps].X = s.X; cell->stack[ps].Y = s.Y - 1; 
                            ps++; cell->stack[ps].S = LEFT; cell->stack[ps].X = s.X - 1; cell->stack[ps].Y = s.Y; 
                            ps++; cell->stack[ps].S = DOWN; cell->stack[ps].X = s.X; cell->stack[ps].Y = s.Y + 1; 
                            break; 
                        case RIGHT: 
                            ps++; cell->stack[ps].S =  DOWN; cell->stack[ps].X = s.X; cell->stack[ps].Y = s.Y + 1; 
                            ps++; cell->stack[ps].S =    UP; cell->stack[ps].X = s.X; cell->stack[ps].Y = s.Y - 1; 
                            ps++; cell->stack[ps].S = RIGHT; cell->stack[ps].X = s.X + 1; cell->stack[ps].Y = s.Y; 
                            break; 
                    } else { 
                        // no route 
                        ps--; 
                break; 
        cell->pstack = ps; 
        cell->halt = 1; 
 
// search halted cells and split running cells 
int populateHalted(GPUCell *cell, int n, long *count) { 
    int hc = -1, pc = -1, ps; 
 
    // search one halted cell 
    while(++hc < n) if(cell[hc].halt) break; if(hc >= n) return 0; 
 
    // search pending cell and pending state 
    do { 
        while(++pc < n) if(!cell[pc].halt) break; if(pc >= n) return 0; 
 
        // run state to full expanded path 
        runGPUCell(&cell[pc], 1); 
 
        // search pending state 
        for(ps = 0; ps <= cell[pc].pstack; ps++) { 
            char s = cell[pc].stack[ps].S; 
            if(s == UP || s == DOWN || s == LEFT || s == RIGHT) 
                break; 
    } while(ps > cell[pc].pstack); 
 
    // save halted count 
    (*count) += cell[hc].counter; 
 
    // clone state 
    memcpy(&cell[hc], &cell[pc], sizeof(GPUCell)); 
 
    // reset count 
    cell[hc].counter = 0; 
    cell[hc].pstack = ps; // start on pending 
 
    { // remove steps 
        int c; 
        for(c = ps + 1; c <= cell[pc].pstack; c++) 
            if(cell[hc].stack[c].S == EXPLORED) 
                cell[hc].G[COORD(cell[hc].stack[c].X, cell[hc].stack[c].Y)] = EMPTY; 
 
    // avoid pending on source cell 
    cell[pc].stack[ps].S = NULLSTATE; 
 
    // can be splited 
    return 1; 
 
// generate start path, all others halt cells 
GPUCell *generateStartPaths(int n, long *count) { 
    GPUCell *cell = (GPUCell *) malloc(sizeof(GPUCell) * n); 
 
    { // initial problem (pending down from (1, 1) cell to (1, 2) cell) 
        int x, y; 
        for(y = 0; y < SIZE2; y++) 
            for(x = 0; x < SIZE2; x++) 
                cell[0].G[COORD(x, y)] = 
                    y == 0    || y == (SIZE + 1) ? LEFT     : 
                    x == 0    || x == (SIZE + 1) ? UP       : 
                    x == 1    && y == 1          ? UP       : 
                    x == SIZE && y == SIZE       ? OBJETIVE : EMPTY; 
        cell[0].pstack = 0; 
        cell[0].counter = 0; 
        cell[0].stack[0].X = 1; 
        cell[0].stack[0].Y = 2; 
        cell[0].stack[0].S = DOWN; 
        cell[0].halt = 0; 
 
    { // all rest halted 
        int c; 
        for(c = 1; c < n; c++) { 
            cell[c].halt = 1; 
            cell[c].counter = 0; 
 
    // populate halted cells 
    populateHalted(cell, n, count); 
 
    return cell; 
 
int main(int argc, char **argv) { 
 
    int n = 256; // cells to run as parallel 
    long count = 0; 
 
    GPUCell *cell = generateStartPaths(n, &count); 
 
    char *source_str; 
    size_t source_size; 
    { // read kernel 
        FILE *fp = fopen("runGPUCell.cl", "r"); 
        if (!fp) {fprintf(stderr, "Failed to load kernel.\n"); exit(1);} 
        source_str = (char*) malloc(MAX_SOURCE_SIZE); 
        source_size = fread(source_str, 1, MAX_SOURCE_SIZE, fp); 
        fclose(fp); 
 
    // Get platform and device information 
    cl_platform_id platform_id = NULL; 
    cl_uint ret_num_devices, ret_num_platforms, ret = clGetPlatformIDs(1, &platform_id, &ret_num_platforms); 
    cl_device_id device_id = NULL; 
    ret = clGetDeviceIDs( platform_id, 
            CL_DEVICE_TYPE_GPU, // CL_DEVICE_TYPE_GPU or CL_DEVICE_TYPE_CPU 
            1, &device_id, &ret_num_devices); 
 
    cl_context context = clCreateContext( NULL, 1, &device_id, NULL, NULL, &ret); 
    cl_command_queue command_queue = clCreateCommandQueue(context, device_id, 0, &ret); 
    cl_mem cell_mem_obj = clCreateBuffer(context, CL_MEM_READ_WRITE, n * sizeof(GPUCell), NULL, &ret); 
    cl_program program = clCreateProgramWithSource(context, 1, (const char **) &source_str, (const size_t *) &source_size, &ret); 
    ret = clBuildProgram(program, 1, &device_id, NULL, NULL, NULL); 
    cl_kernel kernel = clCreateKernel(program, "runGPUCell", &ret); 
    ret = clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *)&cell_mem_obj); 
    size_t global_item_size = n; 
 
    while(populateHalted(cell, n, &count)); 
 
    for(;;) { 
        // actual halted and running cores 
        int h = 0, c; 
        long s = count; 
        for(c = 0; c < n; c++) {s += cell[c].counter; if(cell[c].halt) h++;} 
        if(h == n) break; 
 
        // populate halted cells 
        while(populateHalted(cell, n, &count)); 
 
        // run step 
        ret = clEnqueueWriteBuffer(command_queue, cell_mem_obj, CL_TRUE, 0, n * sizeof(GPUCell), cell, 0, NULL, NULL); 
        ret = clEnqueueNDRangeKernel(command_queue, kernel, 1, NULL, &global_item_size, NULL, 0, NULL, NULL); 
        ret = clEnqueueReadBuffer(command_queue, cell_mem_obj, CL_TRUE, 0, n * sizeof(GPUCell), cell, 0, NULL, NULL); 
 
 
    ret = clFlush(command_queue); 
    ret = clFinish(command_queue); 
    ret = clReleaseKernel(kernel); 
    ret = clReleaseProgram(program); 
    ret = clReleaseMemObject(cell_mem_obj); 
    ret = clReleaseCommandQueue(command_queue); 
    ret = clReleaseContext(context); 
 
    return 0; 
 
 
//=== archivo "runGPUCell.cl" ======================================== 
// Tamaño del problema como constante  
#define SIZE 7 
#define MAXCOUNT 1000 // ¡ajustar para rendimiento! 
#define SIZE2 (SIZE + 2) 
 
#define UP          '^' 
#define RIGHT       '>' 
#define DOWN        'v' 
#define LEFT        '<' 
#define OBJETIVE    'O' 
#define EMPTY       ' ' 
 
#define SOLID       '*' 
#define EXPLORED    '+' 
#define NULLSTATE   '#' 
 
#define COORD(x, y) (y * SIZE2 + x) 
 
typedef struct __GPUStack__ { 
    unsigned short X; // x coord 
    unsigned short Y; // y coord 
    char S; // status 
} GPUStack; 
 
typedef struct __GPUCell__ { 
    char G[SIZE2 * SIZE2];  // [U, D, L, R] - direction, E - empty, O - objetive 
    int pstack;             // stack pointer 
    long counter;           // path counter 
    char halt;              // halted process 
    GPUStack stack[4 * SIZE2 * SIZE2];  // stack 
} GPUCell; 
 
__kernel void runGPUCell(__global GPUCell *cell) { 
 
    // Get the index of the current element to be processed 
    int i = get_global_id(0); 
 
    if(!cell[i].halt) { 
        int ps = cell[i].pstack; 
        GPUStack s; 
        int maxCount = MAXCOUNT; 
        while(ps >= 0 && (s = cell[i].stack[ps]).S != SOLID) { 
 
            // como máximo 'n' estados 
            if(--maxCount <= 0) { 
                cell[i].pstack = ps; 
                return; 
 
            switch(s.S) { 
            case NULLSTATE: 
                ps--; 
                break; 
            case EXPLORED: 
                cell[i].G[COORD(s.X, s.Y)] = EMPTY; 
                ps--; 
                break; 
            default: { 
                char cc = cell[i].G[COORD(s.X, s.Y)]; 
                if(cc == OBJETIVE) { 
                    cell[i].counter++; 
                    ps--; 
                } else 
                    if(cc == EMPTY) { 
                        cell[i].stack[ps].S = EXPLORED; 
                        cell[i].G[COORD(s.X, s.Y)] = s.S; 
                        switch(s.S) { 
                        case UP: 
                            ps++; cell[i].stack[ps].S =    UP; cell[i].stack[ps].X = s.X; cell[i].stack[ps].Y = s.Y - 1; 
                            ps++; cell[i].stack[ps].S =  LEFT; cell[i].stack[ps].X = s.X - 1; cell[i].stack[ps].Y = s.Y; 
                            ps++; cell[i].stack[ps].S = RIGHT; cell[i].stack[ps].X = s.X + 1; cell[i].stack[ps].Y = s.Y; 
                            break; 
                        case DOWN: 
                            ps++; cell[i].stack[ps].S =  DOWN; cell[i].stack[ps].X = s.X; cell[i].stack[ps].Y = s.Y + 1; 
                            ps++; cell[i].stack[ps].S =  LEFT; cell[i].stack[ps].X = s.X - 1; cell[i].stack[ps].Y = s.Y; 
                            ps++; cell[i].stack[ps].S = RIGHT; cell[i].stack[ps].X = s.X + 1; cell[i].stack[ps].Y = s.Y; 
                            break; 
                        case LEFT: 
                            ps++; cell[i].stack[ps].S =   UP; cell[i].stack[ps].X = s.X; cell[i].stack[ps].Y = s.Y - 1; 
                            ps++; cell[i].stack[ps].S = LEFT; cell[i].stack[ps].X = s.X - 1; cell[i].stack[ps].Y = s.Y; 
                            ps++; cell[i].stack[ps].S = DOWN; cell[i].stack[ps].X = s.X; cell[i].stack[ps].Y = s.Y + 1; 
                            break; 
                        case RIGHT: 
                            ps++; cell[i].stack[ps].S =  DOWN; cell[i].stack[ps].X = s.X; cell[i].stack[ps].Y = s.Y + 1; 
                            ps++; cell[i].stack[ps].S =    UP; cell[i].stack[ps].X = s.X; cell[i].stack[ps].Y = s.Y - 1; 
                            ps++; cell[i].stack[ps].S = RIGHT; cell[i].stack[ps].X = s.X + 1; cell[i].stack[ps].Y = s.Y; 
                            break; 
                    } else { 
                        // no route 
                        ps--; 
                break; 
        cell[i].pstack = ps; 
        cell[i].halt = 1; 

Comenta la solución

Tienes que identificarte para poder publicar tu comentario.