cuda.c

   1 /*
   2  * Copyright 2012      Ecole Normale Superieure
   3  *
   4  * Use of this software is governed by the GNU LGPLv2.1 license
   5  *
   6  * Written by Sven Verdoolaege,
   7  * Ecole Normale Superieure, 45 rue d’Ulm, 75230 Paris, France
   8  */
   9
  10 #include <isl/aff.h>
  11 #include <isl/ast.h>
  12
  13 #include "cuda_common.h"
  14 #include "cuda.h"
  15 #include "gpu.h"
  16 #include "gpu_print.h"
  17 #include "print.h"
  18
  19 static __isl_give isl_printer *print_cuda_macros(__isl_take isl_printer *p)
  20 {
  21         const char *macros =
  22                 "#define cudaCheckReturn(ret) \\\n"
  23                 "  do { \\\n"
  24                 "    cudaError_t cudaCheckReturn_e = (ret); \\\n"
  25                 "    if (cudaCheckReturn_e != cudaSuccess) { \\\n"
  26                 "      fprintf(stderr, \"CUDA error: %s\\n\", "
  27                 "cudaGetErrorString(cudaCheckReturn_e)); \\\n"
  28                 "      fflush(stderr); \\\n"
  29                 "    } \\\n"
  30                 "    assert(cudaCheckReturn_e == cudaSuccess); \\\n"
  31                 "  } while(0)\n"
  32                 "#define cudaCheckKernel() \\\n"
  33                 "  do { \\\n"
  34                 "    cudaCheckReturn(cudaGetLastError()); \\\n"
  35                 "  } while(0)\n\n";
  36
  37         p = isl_printer_print_str(p, macros);
  38         return p;
  39 }
  40
  41 static __isl_give isl_printer *declare_device_arrays(__isl_take isl_printer *p,
  42         struct gpu_prog *prog)
  43 {
  44         int i;
  45
  46         for (i = 0; i < prog->n_array; ++i) {
  47                 if (gpu_array_is_read_only_scalar(&prog->array[i]))
  48                         continue;
  49                 p = isl_printer_start_line(p);
  50                 p = isl_printer_print_str(p, prog->array[i].type);
  51                 p = isl_printer_print_str(p, " *dev_");
  52                 p = isl_printer_print_str(p, prog->array[i].name);
  53                 p = isl_printer_print_str(p, ";");
  54                 p = isl_printer_end_line(p);
  55         }
  56         p = isl_printer_start_line(p);
  57         p = isl_printer_end_line(p);
  58         return p;
  59 }
  60
  61 static __isl_give isl_printer *allocate_device_arrays(
  62         __isl_take isl_printer *p, struct gpu_prog *prog)
  63 {
  64         int i;
  65
  66         for (i = 0; i < prog->n_array; ++i) {
  67                 if (gpu_array_is_read_only_scalar(&prog->array[i]))
  68                         continue;
  69                 p = isl_printer_start_line(p);
  70                 p = isl_printer_print_str(p,
  71                         "cudaCheckReturn(cudaMalloc((void **) &dev_");
  72                 p = isl_printer_print_str(p, prog->array[i].name);
  73                 p = isl_printer_print_str(p, ", ");
  74                 p = gpu_array_info_print_size(p, &prog->array[i]);
  75                 p = isl_printer_print_str(p, "));");
  76                 p = isl_printer_end_line(p);
  77         }
  78         p = isl_printer_start_line(p);
  79         p = isl_printer_end_line(p);
  80         return p;
  81 }
  82
  83 static __isl_give isl_printer *copy_arrays_to_device(__isl_take isl_printer *p,
  84         struct gpu_prog *prog)
  85 {
  86         int i;
  87
  88         for (i = 0; i < prog->n_array; ++i) {
  89                 isl_space *dim;
  90                 isl_set *read_i;
  91                 int empty;
  92
  93                 if (gpu_array_is_read_only_scalar(&prog->array[i]))
  94                         continue;
  95
  96                 dim = isl_space_copy(prog->array[i].dim);
  97                 read_i = isl_union_set_extract_set(prog->copy_in, dim);
  98                 empty = isl_set_fast_is_empty(read_i);
  99                 isl_set_free(read_i);
 100                 if (empty)
 101                         continue;
 102
 103                 p = isl_printer_print_str(p, "cudaCheckReturn(cudaMemcpy(dev_");
 104                 p = isl_printer_print_str(p, prog->array[i].name);
 105                 p = isl_printer_print_str(p, ", ");
 106
 107                 if (gpu_array_is_scalar(&prog->array[i]))
 108                         p = isl_printer_print_str(p, "&");
 109                 p = isl_printer_print_str(p, prog->array[i].name);
 110                 p = isl_printer_print_str(p, ", ");
 111
 112                 p = gpu_array_info_print_size(p, &prog->array[i]);
 113                 p = isl_printer_print_str(p, ", cudaMemcpyHostToDevice));");
 114                 p = isl_printer_end_line(p);
 115         }
 116         p = isl_printer_start_line(p);
 117         p = isl_printer_end_line(p);
 118         return p;
 119 }
 120
 121 static void print_reverse_list(FILE *out, int len, int *list)
 122 {
 123         int i;
 124
 125         if (len == 0)
 126                 return;
 127
 128         fprintf(out, "(");
 129         for (i = 0; i < len; ++i) {
 130                 if (i)
 131                         fprintf(out, ", ");
 132                 fprintf(out, "%d", list[len - 1 - i]);
 133         }
 134         fprintf(out, ")");
 135 }
 136
 137 /* Print the effective grid size as a list of the sizes in each
 138  * dimension, from innermost to outermost.
 139  */
 140 static __isl_give isl_printer *print_grid_size(__isl_take isl_printer *p,
 141         struct ppcg_kernel *kernel)
 142 {
 143         int i;
 144         int dim;
 145
 146         dim = isl_multi_pw_aff_dim(kernel->grid_size, isl_dim_set);
 147         if (dim == 0)
 148                 return p;
 149
 150         p = isl_printer_print_str(p, "(");
 151         for (i = dim - 1; i >= 0; --i) {
 152                 isl_pw_aff *bound;
 153
 154                 bound = isl_multi_pw_aff_get_pw_aff(kernel->grid_size, i);
 155                 p = isl_printer_print_pw_aff(p, bound);
 156                 isl_pw_aff_free(bound);
 157
 158                 if (i > 0)
 159                         p = isl_printer_print_str(p, ", ");
 160         }
 161
 162         p = isl_printer_print_str(p, ")");
 163
 164         return p;
 165 }
 166
 167 /* Print the grid definition.
 168  */
 169 static __isl_give isl_printer *print_grid(__isl_take isl_printer *p,
 170         struct ppcg_kernel *kernel)
 171 {
 172         p = isl_printer_start_line(p);
 173         p = isl_printer_print_str(p, "dim3 k");
 174         p = isl_printer_print_int(p, kernel->id);
 175         p = isl_printer_print_str(p, "_dimGrid");
 176         p = print_grid_size(p, kernel);
 177         p = isl_printer_print_str(p, ";");
 178         p = isl_printer_end_line(p);
 179
 180         return p;
 181 }
 182
 183 /* Print the arguments to a kernel declaration or call.  If "types" is set,
 184  * then print a declaration (including the types of the arguments).
 185  *
 186  * The arguments are printed in the following order
 187  * - the arrays accessed by the kernel
 188  * - the parameters
 189  * - the host loop iterators
 190  */
 191 static __isl_give isl_printer *print_kernel_arguments(__isl_take isl_printer *p,
 192         struct gpu_prog *prog, struct ppcg_kernel *kernel, int types)
 193 {
 194         int i, n;
 195         int first = 1;
 196         unsigned nparam;
 197         isl_space *space;
 198         const char *type;
 199
 200         for (i = 0; i < prog->n_array; ++i) {
 201                 isl_set *arr;
 202                 int empty;
 203
 204                 space = isl_space_copy(prog->array[i].dim);
 205                 arr = isl_union_set_extract_set(kernel->arrays, space);
 206                 empty = isl_set_fast_is_empty(arr);
 207                 isl_set_free(arr);
 208                 if (empty)
 209                         continue;
 210
 211                 if (!first)
 212                         p = isl_printer_print_str(p, ", ");
 213
 214                 if (types) {
 215                         p = isl_printer_print_str(p, prog->array[i].type);
 216                         p = isl_printer_print_str(p, " ");
 217                 }
 218
 219                 if (gpu_array_is_read_only_scalar(&prog->array[i])) {
 220                         p = isl_printer_print_str(p, prog->array[i].name);
 221                 } else {
 222                         if (types)
 223                                 p = isl_printer_print_str(p, "*");
 224                         else
 225                                 p = isl_printer_print_str(p, "dev_");
 226                         p = isl_printer_print_str(p, prog->array[i].name);
 227                 }
 228
 229                 first = 0;
 230         }
 231
 232         space = isl_union_set_get_space(kernel->arrays);
 233         nparam = isl_space_dim(space, isl_dim_param);
 234         for (i = 0; i < nparam; ++i) {
 235                 const char *name;
 236
 237                 name = isl_space_get_dim_name(space, isl_dim_param, i);
 238
 239                 if (!first)
 240                         p = isl_printer_print_str(p, ", ");
 241                 if (types)
 242                         p = isl_printer_print_str(p, "int ");
 243                 p = isl_printer_print_str(p, name);
 244
 245                 first = 0;
 246         }
 247         isl_space_free(space);
 248
 249         n = isl_space_dim(kernel->space, isl_dim_set);
 250         type = isl_options_get_ast_iterator_type(prog->ctx);
 251         for (i = 0; i < n; ++i) {
 252                 const char *name;
 253                 isl_id *id;
 254
 255                 if (!first)
 256                         p = isl_printer_print_str(p, ", ");
 257                 name = isl_space_get_dim_name(kernel->space, isl_dim_set, i);
 258                 if (types) {
 259                         p = isl_printer_print_str(p, type);
 260                         p = isl_printer_print_str(p, " ");
 261                 }
 262                 p = isl_printer_print_str(p, name);
 263
 264                 first = 0;
 265         }
 266
 267         return p;
 268 }
 269
 270 /* Print the header of the given kernel.
 271  */
 272 static __isl_give isl_printer *print_kernel_header(__isl_take isl_printer *p,
 273         struct gpu_prog *prog, struct ppcg_kernel *kernel)
 274 {
 275         p = isl_printer_start_line(p);
 276         p = isl_printer_print_str(p, "__global__ void kernel");
 277         p = isl_printer_print_int(p, kernel->id);
 278         p = isl_printer_print_str(p, "(");
 279         p = print_kernel_arguments(p, prog, kernel, 1);
 280         p = isl_printer_print_str(p, ")");
 281
 282         return p;
 283 }
 284
 285 /* Print the header of the given kernel to both gen->cuda.kernel_h
 286  * and gen->cuda.kernel_c.
 287  */
 288 static void print_kernel_headers(struct gpu_prog *prog,
 289         struct ppcg_kernel *kernel, struct cuda_info *cuda)
 290 {
 291         isl_printer *p;
 292
 293         p = isl_printer_to_file(prog->ctx, cuda->kernel_h);
 294         p = isl_printer_set_output_format(p, ISL_FORMAT_C);
 295         p = print_kernel_header(p, prog, kernel);
 296         p = isl_printer_print_str(p, ";");
 297         p = isl_printer_end_line(p);
 298         isl_printer_free(p);
 299
 300         p = isl_printer_to_file(prog->ctx, cuda->kernel_c);
 301         p = isl_printer_set_output_format(p, ISL_FORMAT_C);
 302         p = print_kernel_header(p, prog, kernel);
 303         p = isl_printer_end_line(p);
 304         isl_printer_free(p);
 305 }
 306
 307 static void print_indent(FILE *dst, int indent)
 308 {
 309         fprintf(dst, "%*s", indent, "");
 310 }
 311
 312 static void print_kernel_iterators(FILE *out, struct ppcg_kernel *kernel)
 313 {
 314         int i, n_grid;
 315         isl_ctx *ctx = isl_ast_node_get_ctx(kernel->tree);
 316         const char *type;
 317         const char *block_dims[] = { "blockIdx.x", "blockIdx.y" };
 318         const char *thread_dims[] = { "threadIdx.x", "threadIdx.y",
 319                                         "threadIdx.z" };
 320
 321         type = isl_options_get_ast_iterator_type(ctx);
 322
 323         n_grid = isl_multi_pw_aff_dim(kernel->grid_size, isl_dim_set);
 324         if (n_grid > 0) {
 325                 print_indent(out, 4);
 326                 fprintf(out, "%s ", type);
 327                 for (i = 0; i < n_grid; ++i) {
 328                         if (i)
 329                                 fprintf(out, ", ");
 330                         fprintf(out, "b%d = %s",
 331                                 i, block_dims[n_grid - 1 - i]);
 332                 }
 333                 fprintf(out, ";\n");
 334         }
 335
 336         if (kernel->n_block > 0) {
 337                 print_indent(out, 4);
 338                 fprintf(out, "%s ", type);
 339                 for (i = 0; i < kernel->n_block; ++i) {
 340                         if (i)
 341                                 fprintf(out, ", ");
 342                         fprintf(out, "t%d = %s",
 343                                 i, thread_dims[kernel->n_block - 1 - i]);
 344                 }
 345                 fprintf(out, ";\n");
 346         }
 347 }
 348
 349 static __isl_give isl_printer *print_kernel_var(__isl_take isl_printer *p,
 350         struct ppcg_kernel_var *var)
 351 {
 352         int j;
 353
 354         p = isl_printer_start_line(p);
 355         if (var->type == ppcg_access_shared)
 356                 p = isl_printer_print_str(p, "__shared__ ");
 357         p = isl_printer_print_str(p, var->array->type);
 358         p = isl_printer_print_str(p, " ");
 359         p = isl_printer_print_str(p,  var->name);
 360         for (j = 0; j < var->array->n_index; ++j) {
 361                 isl_val *v;
 362
 363                 p = isl_printer_print_str(p, "[");
 364                 v = isl_vec_get_element_val(var->size, j);
 365                 p = isl_printer_print_val(p, v);
 366                 isl_val_free(v);
 367                 p = isl_printer_print_str(p, "]");
 368         }
 369         p = isl_printer_print_str(p, ";");
 370         p = isl_printer_end_line(p);
 371
 372         return p;
 373 }
 374
 375 static __isl_give isl_printer *print_kernel_vars(__isl_take isl_printer *p,
 376         struct ppcg_kernel *kernel)
 377 {
 378         int i;
 379
 380         for (i = 0; i < kernel->n_var; ++i)
 381                 p = print_kernel_var(p, &kernel->var[i]);
 382
 383         return p;
 384 }
 385
 386 /* Print a sync statement.
 387  */
 388 static __isl_give isl_printer *print_sync(__isl_take isl_printer *p,
 389         struct ppcg_kernel_stmt *stmt)
 390 {
 391         p = isl_printer_start_line(p);
 392         p = isl_printer_print_str(p, "__syncthreads();");
 393         p = isl_printer_end_line(p);
 394
 395         return p;
 396 }
 397
 398 /* This function is called for each user statement in the AST,
 399  * i.e., for each kernel body statement, copy statement or sync statement.
 400  */
 401 static __isl_give isl_printer *print_kernel_stmt(__isl_take isl_printer *p,
 402         __isl_take isl_ast_print_options *print_options,
 403         __isl_keep isl_ast_node *node, void *user)
 404 {
 405         isl_id *id;
 406         struct ppcg_kernel_stmt *stmt;
 407
 408         id = isl_ast_node_get_annotation(node);
 409         stmt = isl_id_get_user(id);
 410         isl_id_free(id);
 411
 412         isl_ast_print_options_free(print_options);
 413
 414         switch (stmt->type) {
 415         case ppcg_kernel_copy:
 416                 return ppcg_kernel_print_copy(p, stmt);
 417         case ppcg_kernel_sync:
 418                 return print_sync(p, stmt);
 419         case ppcg_kernel_domain:
 420                 return ppcg_kernel_print_domain(p, stmt);
 421         }
 422
 423         return p;
 424 }
 425
 426 static void print_kernel(struct gpu_prog *prog, struct ppcg_kernel *kernel,
 427         struct cuda_info *cuda)
 428 {
 429         isl_ctx *ctx = isl_ast_node_get_ctx(kernel->tree);
 430         isl_ast_print_options *print_options;
 431         isl_printer *p;
 432
 433         print_kernel_headers(prog, kernel, cuda);
 434         fprintf(cuda->kernel_c, "{\n");
 435         print_kernel_iterators(cuda->kernel_c, kernel);
 436
 437         p = isl_printer_to_file(ctx, cuda->kernel_c);
 438         p = isl_printer_set_output_format(p, ISL_FORMAT_C);
 439         p = isl_printer_indent(p, 4);
 440
 441         p = print_kernel_vars(p, kernel);
 442         p = isl_printer_end_line(p);
 443         p = gpu_print_macros(p, kernel->tree);
 444
 445         print_options = isl_ast_print_options_alloc(ctx);
 446         print_options = isl_ast_print_options_set_print_user(print_options,
 447                                                     &print_kernel_stmt, NULL);
 448         p = isl_ast_node_print(kernel->tree, p, print_options);
 449         isl_printer_free(p);
 450
 451         fprintf(cuda->kernel_c, "}\n");
 452 }
 453
 454 struct print_host_user_data {
 455         struct cuda_info *cuda;
 456         struct gpu_prog *prog;
 457 };
 458
 459 /* Print the user statement of the host code to "p".
 460  *
 461  * In particular, print a block of statements that defines the grid
 462  * and the block and then launches the kernel.
 463  */
 464 static __isl_give isl_printer *print_host_user(__isl_take isl_printer *p,
 465         __isl_take isl_ast_print_options *print_options,
 466         __isl_keep isl_ast_node *node, void *user)
 467 {
 468         isl_id *id;
 469         struct ppcg_kernel *kernel;
 470         struct print_host_user_data *data;
 471
 472         id = isl_ast_node_get_annotation(node);
 473         kernel = isl_id_get_user(id);
 474         isl_id_free(id);
 475
 476         data = (struct print_host_user_data *) user;
 477
 478         p = isl_printer_start_line(p);
 479         p = isl_printer_print_str(p, "{");
 480         p = isl_printer_end_line(p);
 481         p = isl_printer_indent(p, 2);
 482
 483         p = isl_printer_start_line(p);
 484         p = isl_printer_print_str(p, "dim3 k");
 485         p = isl_printer_print_int(p, kernel->id);
 486         p = isl_printer_print_str(p, "_dimBlock");
 487         print_reverse_list(isl_printer_get_file(p),
 488                                 kernel->n_block, kernel->block_dim);
 489         p = isl_printer_print_str(p, ";");
 490         p = isl_printer_end_line(p);
 491
 492         p = print_grid(p, kernel);
 493
 494         p = isl_printer_start_line(p);
 495         p = isl_printer_print_str(p, "kernel");
 496         p = isl_printer_print_int(p, kernel->id);
 497         p = isl_printer_print_str(p, " <<<k");
 498         p = isl_printer_print_int(p, kernel->id);
 499         p = isl_printer_print_str(p, "_dimGrid, k");
 500         p = isl_printer_print_int(p, kernel->id);
 501         p = isl_printer_print_str(p, "_dimBlock>>> (");
 502         p = print_kernel_arguments(p, data->prog, kernel, 0);
 503         p = isl_printer_print_str(p, ");");
 504         p = isl_printer_end_line(p);
 505
 506         p = isl_printer_start_line(p);
 507         p = isl_printer_print_str(p, "cudaCheckKernel();");
 508         p = isl_printer_end_line(p);
 509
 510         p = isl_printer_indent(p, -2);
 511         p = isl_printer_start_line(p);
 512         p = isl_printer_print_str(p, "}");
 513         p = isl_printer_end_line(p);
 514
 515         p = isl_printer_start_line(p);
 516         p = isl_printer_end_line(p);
 517
 518         print_kernel(data->prog, kernel, data->cuda);
 519
 520         isl_ast_print_options_free(print_options);
 521
 522         return p;
 523 }
 524
 525 static __isl_give isl_printer *print_host_code(__isl_take isl_printer *p,
 526         struct gpu_prog *prog, __isl_keep isl_ast_node *tree,
 527         struct cuda_info *cuda)
 528 {
 529         isl_ast_print_options *print_options;
 530         isl_ctx *ctx = isl_ast_node_get_ctx(tree);
 531         struct print_host_user_data data = { cuda, prog };
 532
 533         print_options = isl_ast_print_options_alloc(ctx);
 534         print_options = isl_ast_print_options_set_print_user(print_options,
 535                                                 &print_host_user, &data);
 536
 537         p = gpu_print_macros(p, tree);
 538         p = isl_ast_node_print(tree, p, print_options);
 539
 540         return p;
 541 }
 542
 543 /* For each array that needs to be copied out (based on prog->copy_out),
 544  * copy the contents back from the GPU to the host.
 545  *
 546  * If any element of a given array appears in prog->copy_out, then its
 547  * entire extent is in prog->copy_out.  The bounds on this extent have
 548  * been precomputed in extract_array_info and are used in
 549  * gpu_array_info_print_size.
 550  */
 551 static __isl_give isl_printer *copy_arrays_from_device(
 552         __isl_take isl_printer *p, struct gpu_prog *prog)
 553 {
 554         int i;
 555         isl_union_set *copy_out;
 556         copy_out = isl_union_set_copy(prog->copy_out);
 557
 558         for (i = 0; i < prog->n_array; ++i) {
 559                 isl_space *dim;
 560                 isl_set *copy_out_i;
 561                 int empty;
 562
 563                 dim = isl_space_copy(prog->array[i].dim);
 564                 copy_out_i = isl_union_set_extract_set(copy_out, dim);
 565                 empty = isl_set_fast_is_empty(copy_out_i);
 566                 isl_set_free(copy_out_i);
 567                 if (empty)
 568                         continue;
 569
 570                 p = isl_printer_print_str(p, "cudaCheckReturn(cudaMemcpy(");
 571                 if (gpu_array_is_scalar(&prog->array[i]))
 572                         p = isl_printer_print_str(p, "&");
 573                 p = isl_printer_print_str(p, prog->array[i].name);
 574                 p = isl_printer_print_str(p, ", dev_");
 575                 p = isl_printer_print_str(p, prog->array[i].name);
 576                 p = isl_printer_print_str(p, ", ");
 577                 p = gpu_array_info_print_size(p, &prog->array[i]);
 578                 p = isl_printer_print_str(p, ", cudaMemcpyDeviceToHost));");
 579                 p = isl_printer_end_line(p);
 580         }
 581
 582         isl_union_set_free(copy_out);
 583         p = isl_printer_start_line(p);
 584         p = isl_printer_end_line(p);
 585         return p;
 586 }
 587
 588 static __isl_give isl_printer *free_device_arrays(__isl_take isl_printer *p,
 589         struct gpu_prog *prog)
 590 {
 591         int i;
 592
 593         for (i = 0; i < prog->n_array; ++i) {
 594                 if (gpu_array_is_read_only_scalar(&prog->array[i]))
 595                         continue;
 596                 p = isl_printer_print_str(p, "cudaCheckReturn(cudaFree(dev_");
 597                 p = isl_printer_print_str(p, prog->array[i].name);
 598                 p = isl_printer_print_str(p, "));");
 599                 p = isl_printer_end_line(p);
 600         }
 601
 602         return p;
 603 }
 604
 605 int generate_cuda(isl_ctx *ctx, struct ppcg_scop *scop,
 606         struct ppcg_options *options, const char *input)
 607 {
 608         struct cuda_info cuda;
 609         struct gpu_prog *prog;
 610         isl_ast_node *tree;
 611         isl_printer *p;
 612
 613         if (!scop)
 614                 return -1;
 615
 616         prog = gpu_prog_alloc(ctx, scop);
 617         if (!prog)
 618                 return -1;
 619
 620         tree = generate_gpu(ctx, prog, options);
 621
 622         cuda.start = scop->start;
 623         cuda.end = scop->end;
 624         cuda_open_files(&cuda, input);
 625
 626         p = isl_printer_to_file(ctx, cuda.host_c);
 627         p = isl_printer_set_output_format(p, ISL_FORMAT_C);
 628         p = ppcg_print_exposed_declarations(p, scop);
 629         p = ppcg_start_block(p);
 630
 631         p = print_cuda_macros(p);
 632
 633         p = declare_device_arrays(p, prog);
 634         p = allocate_device_arrays(p, prog);
 635         p = copy_arrays_to_device(p, prog);
 636
 637         p = print_host_code(p, prog, tree, &cuda);
 638         isl_ast_node_free(tree);
 639
 640         p = copy_arrays_from_device(p, prog);
 641         p = free_device_arrays(p, prog);
 642
 643         p = ppcg_end_block(p);
 644         isl_printer_free(p);
 645
 646         cuda_close_files(&cuda);
 647
 648         gpu_prog_free(prog);
 649
 650         return 0;
 651 }