2 * Copyright 2010-2011 INRIA Saclay
4 * Use of this software is governed by the GNU LGPLv2.1 license
6 * Written by Sven Verdoolaege, INRIA Saclay - Ile-de-France,
7 * Parc Club Orsay Universite, ZAC des vignes, 4 rue Jacques Monod,
14 #include <isl/polynomial.h>
15 #include <isl/union_set.h>
20 #include <isl/schedule.h>
21 #include <isl/options.h>
22 #include <cloog/isl/cloog.h>
25 #include "cuda_common.h"
28 #include "ppcg_options.h"
30 /* The fields stride, shift and shift_map only contain valid information
32 * If so, they express that current index is such that if you add shift,
33 * then the result is always a multiple of stride.
34 * shift_map contains the mapping
36 * i -> (i + shift)/stride
38 struct cuda_array_bound
{
44 isl_basic_map
*shift_map
;
47 struct cuda_array_info
;
49 /* A group of array references in a kernel that should be handled together.
50 * If private_bound is not NULL, then it is mapped to registers.
51 * Otherwise, if shared_bound is not NULL, it is mapped to shared memory.
52 * Otherwise, it is accessed from global memory.
54 struct cuda_array_ref_group
{
55 /* The references in this group access this array. */
56 struct cuda_array_info
*array
;
57 /* Position of this group in the list of reference groups of array. */
60 /* The following fields are use during the construction of the groups.
61 * access is the combined access relation relative to the shared
63 * write is set if any access in the group is a write.
68 /* For each index, size and offset of piece in shared memory. */
69 struct cuda_array_bound
*shared_bound
;
71 /* For each index, size and offset of piece in private memory. */
72 struct cuda_array_bound
*private_bound
;
74 /* References in this group; point to elements of a linked list. */
76 struct cuda_stmt_access
**refs
;
79 struct cuda_array_info
{
83 /* Name of the array. */
85 /* Number of indices. */
87 /* For each index, a bound on the array in that direction. */
89 /* For each index, bound[i] specialized to the current kernel. */
90 isl_pw_aff
**local_bound
;
92 /* All references to this array; point to elements of a linked list. */
94 struct cuda_stmt_access
**refs
;
96 /* The reference groups associated to this array. */
98 struct cuda_array_ref_group
**groups
;
100 /* Last shared memory tile dimension that affects tile of this array. */
102 /* Dimension at which copying to/from shared memory is printed.
103 * if >= 0, then the value is >= last_shared
104 * if -1, then the copying is done at the leaf level.
106 int print_shared_level
;
109 /* Print the name of the local copy of a given group of array references.
111 static void print_array_name(FILE *out
, struct cuda_array_ref_group
*group
)
115 if (group
->private_bound
)
116 fprintf(out
, "private_");
117 else if (group
->shared_bound
)
118 fprintf(out
, "shared_");
121 fprintf(out
, "%s", group
->array
->name
);
122 if (!global
&& group
->array
->n_group
> 1)
123 fprintf(out
, "_%d", group
->nr
);
126 /* Collect all references to the given array and store pointers to them
129 static void collect_references(struct cuda_gen
*gen
,
130 struct cuda_array_info
*array
)
136 for (i
= 0; i
< gen
->n_stmts
; ++i
) {
137 struct cuda_stmt
*stmt
= &gen
->stmts
[i
];
138 struct cuda_stmt_access
*access
;
140 for (access
= stmt
->accesses
; access
; access
= access
->next
) {
142 name
= isl_map_get_tuple_name(access
->access
,
144 if (name
&& !strcmp(array
->name
, name
))
150 array
->refs
= isl_alloc_array(gen
->ctx
, struct cuda_stmt_access
*, n
);
154 for (i
= 0; i
< gen
->n_stmts
; ++i
) {
155 struct cuda_stmt
*stmt
= &gen
->stmts
[i
];
156 struct cuda_stmt_access
*access
;
158 for (access
= stmt
->accesses
; access
; access
= access
->next
) {
160 name
= isl_map_get_tuple_name(access
->access
,
162 if (!name
|| strcmp(array
->name
, name
))
165 array
->refs
[n
++] = access
;
170 static struct cuda_array_bound
*create_bound_list(isl_ctx
*ctx
, int n_index
)
173 struct cuda_array_bound
*bound
;
175 bound
= isl_alloc_array(ctx
, struct cuda_array_bound
, n_index
);
178 for (i
= 0; i
< n_index
; ++i
) {
179 isl_int_init(bound
[i
].size
);
181 isl_int_init(bound
[i
].stride
);
182 bound
[i
].shift
= NULL
;
183 bound
[i
].shift_map
= NULL
;
189 static void free_bound_list(struct cuda_array_bound
*bound
, int n_index
)
196 for (j
= 0; j
< n_index
; ++j
) {
197 isl_int_clear(bound
[j
].size
);
198 isl_int_clear(bound
[j
].stride
);
199 isl_aff_free(bound
[j
].lb
);
200 isl_aff_free(bound
[j
].shift
);
201 isl_basic_map_free(bound
[j
].shift_map
);
206 static struct pet_array
*find_array(struct pet_scop
*scop
,
207 __isl_keep isl_set
*accessed
)
212 id
= isl_set_get_tuple_id(accessed
);
214 for (i
= 0; i
< scop
->n_array
; ++i
) {
217 id_i
= isl_set_get_tuple_id(scop
->arrays
[i
]->extent
);
224 return i
< scop
->n_array
? scop
->arrays
[i
] : NULL
;
227 /* Compute bounds on the host arrays based on the accessed elements
228 * and collect all references to the array.
230 static int extract_array_info(__isl_take isl_set
*array
, void *user
)
233 struct cuda_gen
*gen
= (struct cuda_gen
*)user
;
237 isl_pw_aff
**local_bounds
;
238 struct pet_array
*pa
;
240 n_index
= isl_set_dim(array
, isl_dim_set
);
241 name
= isl_set_get_tuple_name(array
);
242 bounds
= isl_alloc_array(isl_set_get_ctx(array
),
243 isl_pw_aff
*, n_index
);
245 local_bounds
= isl_calloc_array(isl_set_get_ctx(array
),
246 isl_pw_aff
*, n_index
);
247 assert(local_bounds
);
248 gen
->array
[gen
->n_array
].dim
= isl_set_get_space(array
);
249 gen
->array
[gen
->n_array
].name
= strdup(name
);
250 gen
->array
[gen
->n_array
].n_index
= n_index
;
251 gen
->array
[gen
->n_array
].bound
= bounds
;
252 gen
->array
[gen
->n_array
].local_bound
= local_bounds
;
254 pa
= find_array(gen
->scop
, array
);
257 gen
->array
[gen
->n_array
].type
= strdup(pa
->element_type
);
259 for (i
= 0; i
< n_index
; ++i
) {
264 isl_set
*size
= i
== 0 ? array
: pa
->extent
;
266 bound
= isl_set_dim_max(isl_set_copy(size
), i
);
268 dom
= isl_pw_aff_domain(isl_pw_aff_copy(bound
));
269 ls
= isl_local_space_from_space(isl_set_get_space(dom
));
270 one
= isl_aff_zero_on_domain(ls
);
271 one
= isl_aff_add_constant_si(one
, 1);
272 bound
= isl_pw_aff_add(bound
, isl_pw_aff_alloc(dom
, one
));
273 bound
= isl_pw_aff_gist(bound
, isl_set_copy(gen
->context
));
278 collect_references(gen
, &gen
->array
[gen
->n_array
]);
286 void collect_array_info(struct cuda_gen
*gen
)
288 isl_union_set
*arrays
;
290 arrays
= isl_union_map_range(isl_union_map_copy(gen
->read
));
291 arrays
= isl_union_set_union(arrays
,
292 isl_union_map_range(isl_union_map_copy(gen
->write
)));
293 arrays
= isl_union_set_coalesce(arrays
);
295 gen
->n_array
= isl_union_set_n_set(arrays
);
296 gen
->array
= isl_alloc_array(gen
->ctx
,
297 struct cuda_array_info
, gen
->n_array
);
300 isl_union_set_foreach_set(arrays
, &extract_array_info
, gen
);
301 isl_union_set_free(arrays
);
304 static void free_array_info(struct cuda_gen
*gen
)
308 for (i
= 0; i
< gen
->n_array
; ++i
) {
309 int n_index
= gen
->array
[i
].n_index
;
310 free(gen
->array
[i
].type
);
311 free(gen
->array
[i
].name
);
312 for (j
= 0; j
< n_index
; ++j
) {
313 isl_pw_aff_free(gen
->array
[i
].bound
[j
]);
314 isl_pw_aff_free(gen
->array
[i
].local_bound
[j
]);
316 isl_space_free(gen
->array
[i
].dim
);
317 free(gen
->array
[i
].bound
);
318 free(gen
->array
[i
].local_bound
);
319 free(gen
->array
[i
].refs
);
324 static void declare_device_arrays(struct cuda_gen
*gen
)
328 for (i
= 0; i
< gen
->n_array
; ++i
)
329 fprintf(gen
->cuda
.host_c
, "%s *dev_%s;\n",
330 gen
->array
[i
].type
, gen
->array
[i
].name
);
331 fprintf(gen
->cuda
.host_c
, "\n");
334 static void print_array_size(struct cuda_gen
*gen
, FILE *out
,
335 struct cuda_array_info
*array
)
340 prn
= isl_printer_to_file(gen
->ctx
, out
);
341 prn
= isl_printer_set_output_format(prn
, ISL_FORMAT_C
);
342 for (i
= 0; i
< array
->n_index
; ++i
) {
343 prn
= isl_printer_print_str(prn
, "(");
344 prn
= isl_printer_print_pw_aff(prn
, array
->bound
[i
]);
345 prn
= isl_printer_print_str(prn
, ") * ");
347 prn
= isl_printer_print_str(prn
, "sizeof(");
348 prn
= isl_printer_print_str(prn
, array
->type
);
349 prn
= isl_printer_print_str(prn
, ")");
350 isl_printer_free(prn
);
353 static void allocate_device_arrays(struct cuda_gen
*gen
)
357 for (i
= 0; i
< gen
->n_array
; ++i
) {
358 fprintf(gen
->cuda
.host_c
,
359 "cudaCheckReturn(cudaMalloc((void **) &dev_%s, ",
361 print_array_size(gen
, gen
->cuda
.host_c
, &gen
->array
[i
]);
362 fprintf(gen
->cuda
.host_c
, "));\n");
364 fprintf(gen
->cuda
.host_c
, "\n");
367 static void free_device_arrays(struct cuda_gen
*gen
)
371 for (i
= 0; i
< gen
->n_array
; ++i
)
372 fprintf(gen
->cuda
.host_c
, "cudaCheckReturn(cudaFree(dev_%s));\n",
376 /* Check if a cuda array is a scalar. A scalar is a value that is not stored
377 * as an array or through a pointer reference, but as single data element. At
378 * the moment, scalars are represented as zero dimensional arrays.
380 static int cuda_array_is_scalar(struct cuda_array_info
*array
)
382 return (array
->n_index
== 0);
385 static void copy_arrays_to_device(struct cuda_gen
*gen
)
389 for (i
= 0; i
< gen
->n_array
; ++i
) {
394 dim
= isl_space_copy(gen
->array
[i
].dim
);
395 read_i
= isl_union_set_extract_set(gen
->copy_in
, dim
);
396 empty
= isl_set_fast_is_empty(read_i
);
397 isl_set_free(read_i
);
401 fprintf(gen
->cuda
.host_c
, "cudaCheckReturn(cudaMemcpy(dev_%s,",
404 if (cuda_array_is_scalar(&(gen
->array
[i
])))
405 fprintf(gen
->cuda
.host_c
, " &%s, ",
408 fprintf(gen
->cuda
.host_c
, " %s, ", gen
->array
[i
].name
);
410 print_array_size(gen
, gen
->cuda
.host_c
, &gen
->array
[i
]);
411 fprintf(gen
->cuda
.host_c
, ", cudaMemcpyHostToDevice));\n");
413 fprintf(gen
->cuda
.host_c
, "\n");
416 static void copy_arrays_from_device(struct cuda_gen
*gen
)
419 isl_union_set
*write
;
420 write
= isl_union_map_range(isl_union_map_copy(gen
->write
));
422 for (i
= 0; i
< gen
->n_array
; ++i
) {
427 dim
= isl_space_copy(gen
->array
[i
].dim
);
428 write_i
= isl_union_set_extract_set(write
, dim
);
429 empty
= isl_set_fast_is_empty(write_i
);
430 isl_set_free(write_i
);
434 fprintf(gen
->cuda
.host_c
,
435 "cudaCheckReturn(cudaMemcpy(%s, dev_%s, ",
436 gen
->array
[i
].name
, gen
->array
[i
].name
);
437 print_array_size(gen
, gen
->cuda
.host_c
, &gen
->array
[i
]);
438 fprintf(gen
->cuda
.host_c
, ", cudaMemcpyDeviceToHost));\n");
441 isl_union_set_free(write
);
442 fprintf(gen
->cuda
.host_c
, "\n");
445 static void read_sizes_from_file(struct cuda_gen
*gen
, const char *filename
,
451 file
= fopen(filename
, "r");
455 for (i
= 0; i
< len
; ++i
)
456 if (fscanf(file
, "%d", &sizes
[i
]) < 1)
462 static void reverse_list(int *list
, int len
)
467 for (i
= 0; 2 * i
< len
; ++i
) {
469 list
[i
] = list
[len
- 1 - i
];
470 list
[len
- 1 - i
] = t
;
474 /* Read user specified sizes from "tile.sizes", "block.sizes" and "grid.sizes"
475 * after filling in some potentially useful defaults.
477 static void read_sizes(struct cuda_gen
*gen
)
481 gen
->tile_size
= isl_alloc_array(gen
->ctx
, int, gen
->tile_len
);
482 assert(gen
->tile_size
);
483 for (n
= 0; n
< gen
->tile_len
; ++n
)
484 gen
->tile_size
[n
] = gen
->options
->tile_size
;
485 read_sizes_from_file(gen
, "tile.sizes", gen
->tile_size
, gen
->tile_len
);
488 gen
->n_block
= (n
<= 3) ? n
: 3;
489 switch (gen
->n_block
) {
491 gen
->block_dim
[0] = 512;
494 gen
->block_dim
[0] = 32;
495 gen
->block_dim
[1] = 16;
498 gen
->block_dim
[0] = 32;
499 gen
->block_dim
[1] = 4;
500 gen
->block_dim
[2] = 4;
503 read_sizes_from_file(gen
, "block.sizes", gen
->block_dim
, gen
->n_block
);
504 reverse_list(gen
->block_dim
, gen
->n_block
);
506 gen
->n_grid
= (n
<= 2) ? n
: 2;
507 switch (gen
->n_grid
) {
509 gen
->grid_dim
[0] = 32768;
512 gen
->grid_dim
[0] = 256;
513 gen
->grid_dim
[1] = 256;
516 read_sizes_from_file(gen
, "grid.sizes", gen
->grid_dim
, gen
->n_grid
);
517 reverse_list(gen
->grid_dim
, gen
->n_grid
);
520 static void free_stmts(struct cuda_stmt
*stmts
, int n
)
524 for (i
= 0; i
< n
; ++i
) {
525 struct cuda_stmt_access
*access
, *next
;
527 for (access
= stmts
[i
].accesses
; access
; access
= next
) {
529 isl_map_free(access
->access
);
533 isl_set_free(stmts
[i
].domain
);
538 void clear_cuda_gen(struct cuda_gen
*gen
)
540 free_stmts(gen
->stmts
, gen
->n_stmts
);
541 free_array_info(gen
);
542 isl_set_free(gen
->context
);
543 isl_union_set_free(gen
->copy_in
);
544 isl_union_map_free(gen
->sched
);
545 isl_union_map_free(gen
->read
);
546 isl_union_map_free(gen
->write
);
549 static void print_reverse_list(FILE *out
, int len
, int *list
)
553 for (i
= 0; i
< len
; ++i
) {
556 fprintf(out
, "%d", list
[len
- 1 - i
]);
560 static void print_kernel_launch(struct cuda_gen
*gen
,
561 __isl_keep isl_union_set
*arrays
)
568 print_indent(gen
->code
.dst
, gen
->code
.indent
);
569 fprintf(gen
->code
.dst
, "kernel%d <<<k%d_dimGrid, k%d_dimBlock>>> (",
570 gen
->kernel_id
, gen
->kernel_id
, gen
->kernel_id
);
571 fprintf(gen
->cuda
.kernel_c
, "__global__ void kernel%d(",
573 fprintf(gen
->cuda
.kernel_h
, "__global__ void kernel%d(",
576 for (i
= 0; i
< gen
->n_array
; ++i
) {
581 dim
= isl_space_copy(gen
->array
[i
].dim
);
582 arr
= isl_union_set_extract_set(arrays
, dim
);
583 empty
= isl_set_fast_is_empty(arr
);
589 fprintf(gen
->code
.dst
, ", ");
590 fprintf(gen
->cuda
.kernel_c
, ", ");
591 fprintf(gen
->cuda
.kernel_h
, ", ");
594 fprintf(gen
->code
.dst
, "dev_%s", gen
->array
[i
].name
);
595 fprintf(gen
->cuda
.kernel_c
, "%s *%s",
596 gen
->array
[i
].type
, gen
->array
[i
].name
);
597 fprintf(gen
->cuda
.kernel_h
, "%s *%s",
598 gen
->array
[i
].type
, gen
->array
[i
].name
);
603 dim
= isl_union_set_get_space(arrays
);
604 nparam
= isl_space_dim(dim
, isl_dim_param
);
605 for (i
= 0; i
< nparam
; ++i
) {
606 const char *name
= isl_space_get_dim_name(dim
, isl_dim_param
, i
);
608 fprintf(gen
->code
.dst
, ", ");
609 fprintf(gen
->cuda
.kernel_c
, ", ");
610 fprintf(gen
->cuda
.kernel_h
, ", ");
612 fprintf(gen
->code
.dst
, "%s", name
);
613 fprintf(gen
->cuda
.kernel_c
, "int %s", name
);
614 fprintf(gen
->cuda
.kernel_h
, "int %s", name
);
619 for (i
= 0; i
< gen
->tile_first
; ++i
) {
621 fprintf(gen
->code
.dst
, ", ");
622 fprintf(gen
->cuda
.kernel_c
, ", ");
623 fprintf(gen
->cuda
.kernel_h
, ", ");
625 fprintf(gen
->code
.dst
, "h%d", i
);
626 fprintf(gen
->cuda
.kernel_c
, "int h%d", i
);
627 fprintf(gen
->cuda
.kernel_h
, "int h%d", i
);
631 fprintf(gen
->code
.dst
, ");\n");
632 fprintf(gen
->cuda
.kernel_c
, ")\n");
633 fprintf(gen
->cuda
.kernel_h
, ");\n");
635 fprintf(gen
->code
.dst
, "cudaCheckKernel();\n");
638 /* Construct a map from a domain of dimensionality "len"
639 * to a domain of dimensionality "len" + "tile_len" that tiles
640 * the "tile_len" coordinates starting at "first".
641 * In particular, [s_i] -> [s_i / tile_size[i], s_i % tile_size[i]].
642 * "dim" prescribes the parameters.
644 static __isl_give isl_map
*tile(__isl_take isl_space
*dim
, int len
,
645 int first
, int tile_len
, int *tile_size
)
655 dim
= isl_space_add_dims(dim
, isl_dim_in
, len
);
656 dim
= isl_space_add_dims(dim
, isl_dim_out
, len
+ tile_len
);
657 bmap
= isl_basic_map_universe(isl_space_copy(dim
));
658 ls
= isl_local_space_from_space(dim
);
660 for (i
= 0; i
< len
- tile_len
; ++i
) {
661 int j
= i
< first
? i
: i
+ tile_len
;
662 int k
= i
< first
? i
: i
+ 2 * tile_len
;
664 c
= isl_equality_alloc(isl_local_space_copy(ls
));
665 isl_int_set_si(v
, -1);
666 isl_constraint_set_coefficient(c
, isl_dim_in
, j
, v
);
667 isl_int_set_si(v
, 1);
668 isl_constraint_set_coefficient(c
, isl_dim_out
, k
, v
);
669 bmap
= isl_basic_map_add_constraint(bmap
, c
);
672 for (i
= 0; i
< tile_len
; ++i
) {
673 c
= isl_equality_alloc(isl_local_space_copy(ls
));
674 isl_int_set_si(v
, -1);
675 isl_constraint_set_coefficient(c
, isl_dim_in
, first
+ i
, v
);
676 isl_int_set_si(v
, tile_size
[i
]);
677 isl_constraint_set_coefficient(c
, isl_dim_out
, first
+ i
, v
);
678 isl_int_set_si(v
, 1);
679 isl_constraint_set_coefficient(c
, isl_dim_out
,
680 first
+ i
+ tile_len
, v
);
681 bmap
= isl_basic_map_add_constraint(bmap
, c
);
683 c
= isl_inequality_alloc(isl_local_space_copy(ls
));
684 isl_int_set_si(v
, 1);
685 isl_constraint_set_coefficient(c
, isl_dim_out
,
686 first
+ i
+ tile_len
, v
);
687 bmap
= isl_basic_map_add_constraint(bmap
, c
);
689 c
= isl_inequality_alloc(isl_local_space_copy(ls
));
690 isl_int_set_si(v
, -1);
691 isl_constraint_set_coefficient(c
, isl_dim_out
,
692 first
+ i
+ tile_len
, v
);
693 isl_int_set_si(v
, tile_size
[i
] - 1);
694 isl_constraint_set_constant(c
, v
);
695 bmap
= isl_basic_map_add_constraint(bmap
, c
);
698 isl_local_space_free(ls
);
701 return isl_map_from_basic_map(bmap
);
704 /* Construct a map from a domain of dimensionality "len"
705 * to a domain of dimensionality "len" + "wrap_len" that "wraps"
706 * the "wrap_len" coordinates starting at "first" according to "wrap_size".
707 * In particular, [s_i] -> [s_i, s_i % wrap_size[i]].
708 * To do so, we need extra variables corresponding to [s_i / wrap_size[i]],
709 * that are projected out at the end.
710 * "dim" prescribes the parameters.
712 static __isl_give isl_map
*wrap(__isl_take isl_space
*dim
, int len
,
713 int first
, int wrap_len
, int *wrap_size
)
720 dim
= isl_space_add_dims(dim
, isl_dim_in
, len
);
721 dim
= isl_space_add_dims(dim
, isl_dim_out
, len
+ 2 * wrap_len
);
722 bmap
= isl_basic_map_universe(isl_space_copy(dim
));
723 ls
= isl_local_space_from_space(dim
);
725 for (i
= 0; i
< len
; ++i
) {
726 int k
= i
< first
+ wrap_len
? i
: i
+ 2 * wrap_len
;
728 c
= isl_equality_alloc(isl_local_space_copy(ls
));
729 isl_constraint_set_coefficient_si(c
, isl_dim_in
, i
, -1);
730 isl_constraint_set_coefficient_si(c
, isl_dim_out
, k
, 1);
731 bmap
= isl_basic_map_add_constraint(bmap
, c
);
734 for (i
= 0; i
< wrap_len
; ++i
) {
735 c
= isl_equality_alloc(isl_local_space_copy(ls
));
736 isl_constraint_set_coefficient_si(c
, isl_dim_out
,
738 isl_constraint_set_coefficient_si(c
, isl_dim_out
,
739 first
+ wrap_len
+ i
, 1);
740 isl_constraint_set_coefficient_si(c
, isl_dim_out
,
741 first
+ 2 * wrap_len
+ i
, wrap_size
[i
]);
742 bmap
= isl_basic_map_add_constraint(bmap
, c
);
744 c
= isl_inequality_alloc(isl_local_space_copy(ls
));
745 isl_constraint_set_coefficient_si(c
, isl_dim_out
,
746 first
+ wrap_len
+ i
, 1);
747 bmap
= isl_basic_map_add_constraint(bmap
, c
);
749 c
= isl_inequality_alloc(isl_local_space_copy(ls
));
750 isl_constraint_set_coefficient_si(c
, isl_dim_out
,
751 first
+ wrap_len
+ i
, -1);
752 isl_constraint_set_constant_si(c
, wrap_size
[i
] - 1);
753 bmap
= isl_basic_map_add_constraint(bmap
, c
);
756 isl_local_space_free(ls
);
758 bmap
= isl_basic_map_project_out(bmap
, isl_dim_out
,
759 first
+ 2 * wrap_len
, wrap_len
);
761 return isl_map_from_basic_map(bmap
);
764 /* Add "n" parameters named prefix%d.
766 static __isl_give isl_set
*add_params( __isl_take isl_set
*set
,
767 int n
, const char *prefix
)
773 nparam
= isl_set_dim(set
, isl_dim_param
);
774 set
= isl_set_add_dims(set
, isl_dim_param
, n
);
776 for (i
= 0; i
< n
; ++i
) {
777 snprintf(name
, sizeof(name
), "%s%d", prefix
, i
);
778 set
= isl_set_set_dim_name(set
, isl_dim_param
,
785 /* Equate the "n" dimensions of "set" starting at "first" to
786 * freshly created parameters named prefix%d.
788 static __isl_give isl_set
*parametrize(__isl_take isl_set
*set
,
789 int first
, int n
, const char *prefix
)
799 nparam
= isl_set_dim(set
, isl_dim_param
);
801 set
= add_params(set
, n
, prefix
);
803 dim
= isl_set_get_space(set
);
804 bset
= isl_basic_set_universe(isl_space_copy(dim
));
805 ls
= isl_local_space_from_space(dim
);
809 for (i
= 0; i
< n
; ++i
) {
810 c
= isl_equality_alloc(isl_local_space_copy(ls
));
811 isl_int_set_si(v
, -1);
812 isl_constraint_set_coefficient(c
, isl_dim_param
, nparam
+ i
, v
);
813 isl_int_set_si(v
, 1);
814 isl_constraint_set_coefficient(c
, isl_dim_set
, first
+ i
, v
);
815 bset
= isl_basic_set_add_constraint(bset
, c
);
819 isl_local_space_free(ls
);
821 return isl_set_intersect(set
, isl_set_from_basic_set(bset
));
824 static __isl_give isl_set
*parametrization(__isl_take isl_space
*dim
,
825 int len
, int first
, int n
, const char *prefix
)
829 dim
= isl_space_add_dims(dim
, isl_dim_set
, len
);
830 set
= isl_set_universe(dim
);
832 return parametrize(set
, first
, n
, prefix
);
835 /* Tile the B loops over the tile sizes and then tile/wrap
836 * the T1 loops over the blocks.
838 static __isl_give isl_union_map
*tile_schedule(struct cuda_gen
*gen
,
839 __isl_take isl_union_map
*sched
)
842 isl_map
*tiling
, *block_tiling
;
844 dim
= isl_union_map_get_space(sched
);
845 tiling
= tile(isl_space_copy(dim
), gen
->untiled_len
,
846 gen
->tile_first
, gen
->tile_len
, gen
->tile_size
);
848 if (gen
->options
->wrap
)
849 block_tiling
= wrap(dim
, gen
->untiled_len
+ gen
->tile_len
,
850 gen
->tile_first
, gen
->n_grid
, gen
->grid_dim
);
852 block_tiling
= tile(dim
, gen
->untiled_len
+ gen
->tile_len
,
853 gen
->tile_first
, gen
->n_grid
, gen
->grid_dim
);
855 gen
->tiled_len
= gen
->untiled_len
+ gen
->tile_len
+ gen
->n_grid
;
857 tiling
= isl_map_apply_range(tiling
, block_tiling
);
859 sched
= isl_union_map_apply_range(sched
,
860 isl_union_map_from_map(tiling
));
862 gen
->shared_len
= gen
->tile_first
+ gen
->tile_len
+ gen
->n_grid
;
867 static __isl_give isl_union_map
*parametrize_tiled_schedule(
868 struct cuda_gen
*gen
, __isl_take isl_union_map
*sched
)
873 dim
= isl_union_map_get_space(sched
);
874 par
= parametrization(dim
, gen
->tiled_len
, 0, gen
->tile_first
, "h");
875 sched
= isl_union_map_intersect_range(sched
,
876 isl_union_set_from_set(par
));
878 dim
= isl_union_map_get_space(sched
);
879 par
= parametrization(dim
, gen
->tiled_len
,
880 gen
->tile_first
+ gen
->n_grid
, gen
->n_grid
, "b");
881 sched
= isl_union_map_intersect_range(sched
,
882 isl_union_set_from_set(par
));
887 /* Tile/wrap the P1 loops over the threads.
889 static __isl_give isl_union_map
*thread_tile_schedule(struct cuda_gen
*gen
,
890 __isl_take isl_union_map
*sched
)
896 dim
= isl_union_map_get_space(sched
);
898 if (gen
->options
->wrap
)
899 tiling
= wrap(isl_space_copy(dim
), gen
->tiled_len
,
900 gen
->shared_len
, gen
->n_block
, gen
->block_dim
);
902 tiling
= tile(isl_space_copy(dim
), gen
->tiled_len
,
903 gen
->shared_len
, gen
->n_block
, gen
->block_dim
);
904 gen
->thread_tiled_len
= gen
->tiled_len
+ gen
->n_block
;
906 sched
= isl_union_map_apply_range(sched
,
907 isl_union_map_from_map(tiling
));
909 par
= parametrization(dim
, gen
->thread_tiled_len
,
910 gen
->tile_first
+ gen
->tile_len
+ gen
->n_grid
+ gen
->n_block
,
912 sched
= isl_union_map_intersect_range(sched
,
913 isl_union_set_from_set(par
));
915 gen
->shared_len
= gen
->tile_first
+ gen
->tile_len
+ gen
->n_grid
;
920 /* If the user asked for it, scale the shared memory tile loops
921 * (T1P and T2) of "sched" by gen->tile_size[i].
922 * If we are not performing "wrapping", then additionally scale the T1P
923 * loops by gen->grid_dim[i].
925 static __isl_give isl_union_map
*scale_tile_loops(struct cuda_gen
*gen
,
926 __isl_take isl_union_map
*sched
)
930 isl_basic_map
*scale
;
934 if (!gen
->options
->scale_tile_loops
)
937 dim
= isl_union_map_get_space(sched
);
938 dim
= isl_space_add_dims(dim
, isl_dim_in
, gen
->tiled_len
);
939 dim
= isl_space_add_dims(dim
, isl_dim_out
, gen
->tiled_len
);
940 scale
= isl_basic_map_universe(isl_space_copy(dim
));
941 ls
= isl_local_space_from_space(dim
);
943 for (i
= 0; i
< gen
->tiled_len
; ++i
) {
946 if (i
>= gen
->tile_first
&& i
< gen
->tile_first
+ gen
->n_grid
) {
947 f
= gen
->tile_size
[i
- gen
->tile_first
];
948 if (!gen
->options
->wrap
)
949 f
*= gen
->grid_dim
[i
- gen
->tile_first
];
950 } else if (i
>= gen
->tile_first
+ gen
->n_grid
&&
951 i
< gen
->tile_first
+ gen
->n_grid
+ gen
->tile_len
) {
952 f
= gen
->tile_size
[i
- (gen
->tile_first
+ gen
->n_grid
)];
955 c
= isl_equality_alloc(isl_local_space_copy(ls
));
956 isl_constraint_set_coefficient_si(c
, isl_dim_in
, i
, f
);
957 isl_constraint_set_coefficient_si(c
, isl_dim_out
, i
, -1);
958 scale
= isl_basic_map_add_constraint(scale
, c
);
961 isl_local_space_free(ls
);
963 sched
= isl_union_map_apply_range(sched
,
964 isl_union_map_from_map(isl_map_from_basic_map(scale
)));
969 /* If we are not performing "wrapping" and if the user asked for it,
970 * scale the thread tile loops (P1T) of "sched" by gen->block_dim[i].
972 static __isl_give isl_union_map
*scale_thread_tile_loops(struct cuda_gen
*gen
,
973 __isl_take isl_union_map
*sched
)
977 isl_basic_map
*scale
;
981 if (gen
->options
->wrap
)
983 if (!gen
->options
->scale_tile_loops
)
986 dim
= isl_union_map_get_space(sched
);
987 dim
= isl_space_add_dims(dim
, isl_dim_in
, gen
->thread_tiled_len
);
988 dim
= isl_space_add_dims(dim
, isl_dim_out
, gen
->thread_tiled_len
);
989 scale
= isl_basic_map_universe(isl_space_copy(dim
));
990 ls
= isl_local_space_from_space(dim
);
992 for (i
= 0; i
< gen
->thread_tiled_len
; ++i
) {
995 if (i
>= gen
->shared_len
&&
996 i
< gen
->shared_len
+ gen
->n_block
)
997 f
= gen
->block_dim
[i
- gen
->shared_len
];
999 c
= isl_equality_alloc(isl_local_space_copy(ls
));
1000 isl_constraint_set_coefficient_si(c
, isl_dim_in
, i
, f
);
1001 isl_constraint_set_coefficient_si(c
, isl_dim_out
, i
, -1);
1002 scale
= isl_basic_map_add_constraint(scale
, c
);
1005 isl_local_space_free(ls
);
1007 sched
= isl_union_map_apply_range(sched
,
1008 isl_union_map_from_map(isl_map_from_basic_map(scale
)));
1013 /* If we are not performing "wrapping" and if the user asked for it,
1014 * scale the "n_tile" loops starting at "first" of "sched" by gen->block_dim[i].
1016 static __isl_give isl_union_map
*scale_access_tile_loops(struct cuda_gen
*gen
,
1017 __isl_take isl_union_map
*sched
, int len
, int first
, int n_tile
)
1021 isl_basic_map
*scale
;
1023 isl_local_space
*ls
;
1025 if (gen
->options
->wrap
)
1027 if (!gen
->options
->scale_tile_loops
)
1030 dim
= isl_union_map_get_space(sched
);
1031 dim
= isl_space_add_dims(dim
, isl_dim_in
, len
);
1032 dim
= isl_space_add_dims(dim
, isl_dim_out
, len
);
1033 scale
= isl_basic_map_universe(isl_space_copy(dim
));
1034 ls
= isl_local_space_from_space(dim
);
1036 for (i
= 0; i
< len
; ++i
) {
1039 if (i
>= first
&& i
< first
+ n_tile
)
1040 f
= gen
->block_dim
[i
- first
];
1042 c
= isl_equality_alloc(isl_local_space_copy(ls
));
1043 isl_constraint_set_coefficient_si(c
, isl_dim_in
, i
, f
);
1044 isl_constraint_set_coefficient_si(c
, isl_dim_out
, i
, -1);
1045 scale
= isl_basic_map_add_constraint(scale
, c
);
1048 isl_local_space_free(ls
);
1050 sched
= isl_union_map_apply_range(sched
,
1051 isl_union_map_from_map(isl_map_from_basic_map(scale
)));
1056 /* If print_user_stmt is set, we want to print the statements ourselves,
1057 * instead of relying on the C preprocessor. If so, we need to use
1058 * the stop option so that the domains will be saved on the statement
1061 static void print_cloog_shared_body(struct cuda_gen
*gen
,
1062 __isl_keep isl_set
*context
, __isl_keep isl_union_map
*sched
, int len
,
1063 void (*print_user_stmt
)(struct gpucode_info
*info
,
1064 struct clast_user_stmt
*s
),
1068 CloogOptions
*options
;
1069 CloogDomain
*cloog_context
;
1070 CloogUnionDomain
*ud
;
1072 struct clast_stmt
*stmt
;
1075 sched
= isl_union_map_copy(sched
);
1076 sched
= isl_union_map_align_params(sched
, isl_set_get_space(context
));
1078 options
= cloog_options_malloc(gen
->state
);
1079 options
->language
= CLOOG_LANGUAGE_C
;
1080 options
->strides
= 1;
1084 options
->override
= 1;
1085 options
->save_domains
= 1;
1086 options
->noscalars
= 1;
1087 options
->first_unroll
= first_unroll
;
1089 ud
= cloog_union_domain_from_isl_union_map(sched
);
1090 for (i
= 0; i
< len
; ++i
) {
1091 snprintf(name
, sizeof(name
), "c%d", i
);
1092 ud
= cloog_union_domain_set_name(ud
, CLOOG_SCAT
, i
, name
);
1094 cloog_context
= cloog_domain_from_isl_set(isl_set_copy(context
));
1095 input
= cloog_input_alloc(cloog_context
, ud
);
1097 stmt
= cloog_clast_create_from_input(input
, options
);
1099 gen
->stmt_code
.indent
= gen
->kernel_code
.indent
;
1100 gen
->stmt_code
.dst
= gen
->cuda
.kernel_c
;
1101 gen
->stmt_code
.print_user_stmt
= print_user_stmt
;
1102 gen
->stmt_code
.print_user_stmt_list
= NULL
;
1103 gen
->stmt_code
.print_for_head
= NULL
;
1104 gen
->stmt_code
.print_for_foot
= NULL
;
1105 gen
->stmt_code
.user
= gen
;
1106 gpu_print_host_stmt(&gen
->stmt_code
, stmt
);
1108 cloog_clast_free(stmt
);
1109 cloog_options_free(options
);
1112 /* Add "len" parameters p[i] called prefix%d,
1113 * with bounds to 0 <= p[i] < size[i].
1115 __isl_give isl_set
*add_bounded_parameters(__isl_take isl_set
*set
,
1116 int len
, int *size
, const char *prefix
)
1122 isl_basic_set
*bset
;
1124 isl_local_space
*ls
;
1127 nparam
= isl_set_dim(set
, isl_dim_param
);
1128 set
= isl_set_add_dims(set
, isl_dim_param
, len
);
1130 for (i
= 0; i
< len
; ++i
) {
1131 snprintf(name
, sizeof(name
), "%s%d", prefix
, i
);
1132 set
= isl_set_set_dim_name(set
, isl_dim_param
,
1136 dim
= isl_set_get_space(set
);
1137 bset
= isl_basic_set_universe(isl_space_copy(dim
));
1138 ls
= isl_local_space_from_space(dim
);
1142 for (i
= 0; i
< len
; ++i
) {
1143 c
= isl_inequality_alloc(isl_local_space_copy(ls
));
1144 isl_int_set_si(v
, 1);
1145 isl_constraint_set_coefficient(c
, isl_dim_param
, nparam
+ i
, v
);
1146 bset
= isl_basic_set_add_constraint(bset
, c
);
1148 c
= isl_inequality_alloc(isl_local_space_copy(ls
));
1149 isl_int_set_si(v
, -1);
1150 isl_constraint_set_coefficient(c
, isl_dim_param
, nparam
+ i
, v
);
1151 isl_int_set_si(v
, size
[i
] - 1);
1152 isl_constraint_set_constant(c
, v
);
1153 bset
= isl_basic_set_add_constraint(bset
, c
);
1157 isl_local_space_free(ls
);
1159 return isl_set_intersect(set
, isl_set_from_basic_set(bset
));
1162 static void print_shared_body(struct cuda_gen
*gen
,
1163 __isl_keep isl_set
*shared_domain
, __isl_keep isl_union_map
*sched
,
1164 int len
, void (*print_user_stmt
)(struct gpucode_info
*info
,
1165 struct clast_user_stmt
*s
),
1170 context
= isl_set_copy(shared_domain
);
1171 context
= parametrize(context
, 0, gen
->shared_len
, "g");
1172 context
= isl_set_project_out(context
, isl_dim_set
, 0, gen
->shared_len
);
1173 context
= add_bounded_parameters(context
,
1174 gen
->n_block
, gen
->block_dim
, "t");
1176 print_cloog_shared_body(gen
, context
, sched
, len
, print_user_stmt
,
1179 isl_set_free(context
);
1182 /* Given a tile of an array, construct a map that maps each element
1183 * of the tile to a copy of the tile shifted to the origin
1184 * (based on the lower bounds in group->private_bound or group->shared_bound).
1185 * If any of the indices is strided, then {private,shared}_bound[i].shift_map
1186 * is applied to the index first.
1187 * The domain of the resulting map is "access",
1188 * while the range space is anonymous.
1190 static __isl_give isl_map
*shift_access(__isl_take isl_set
*access
,
1191 struct cuda_array_ref_group
*group
)
1195 isl_basic_set
*bset
;
1196 isl_basic_map
*bmap
;
1198 isl_basic_set
*offset
;
1199 isl_basic_map
*shift
;
1200 isl_basic_map
*pre_shift
;
1203 struct cuda_array_bound
*bounds
;
1204 int n_index
= group
->array
->n_index
;
1206 bounds
= group
->private_bound
;
1208 bounds
= group
->shared_bound
;
1210 dim
= isl_set_get_space(access
);
1211 dim
= isl_space_drop_dims(dim
, isl_dim_set
, 0, n_index
);
1212 offset
= isl_basic_set_universe(dim
);
1213 for (i
= 0; i
< n_index
; ++i
) {
1214 lb
= isl_aff_copy(bounds
[i
].lb
);
1215 bmap
= isl_basic_map_from_aff(lb
);
1216 bset
= isl_basic_map_range(bmap
);
1217 offset
= isl_basic_set_flat_product(offset
, bset
);
1219 offset
= isl_basic_set_neg(offset
);
1221 dim
= isl_space_map_from_set(isl_set_get_space(access
));
1222 shift
= isl_basic_map_identity(dim
);
1223 shift
= isl_basic_map_set_tuple_name(shift
, isl_dim_out
, NULL
);
1225 bset
= isl_basic_set_universe(isl_set_get_space(access
));
1226 bmap
= isl_basic_map_from_domain_and_range(bset
, offset
);
1228 shift
= isl_basic_map_sum(shift
, bmap
);
1230 dim
= isl_set_get_space(access
);
1231 dim
= isl_space_drop_dims(dim
, isl_dim_set
, 0, n_index
);
1232 dim
= isl_space_map_from_set(dim
);
1233 pre_shift
= isl_basic_map_universe(isl_space_copy(dim
));
1234 dim
= isl_space_add_dims(dim
, isl_dim_in
, 1);
1235 dim
= isl_space_add_dims(dim
, isl_dim_out
, 1);
1236 for (i
= 0; i
< n_index
; ++i
) {
1237 if (!bounds
[i
].shift_map
)
1238 bmap
= isl_basic_map_identity(isl_space_copy(dim
));
1240 bmap
= isl_basic_map_copy(bounds
[i
].shift_map
);
1241 pre_shift
= isl_basic_map_flat_product(pre_shift
, bmap
);
1243 isl_space_free(dim
);
1244 name
= isl_basic_map_get_tuple_name(shift
, isl_dim_in
);
1245 pre_shift
= isl_basic_map_set_tuple_name(pre_shift
, isl_dim_in
, name
);
1246 pre_shift
= isl_basic_map_set_tuple_name(pre_shift
, isl_dim_out
, name
);
1247 shift
= isl_basic_map_apply_range(pre_shift
, shift
);
1249 sched
= isl_map_from_basic_map(shift
);
1250 sched
= isl_map_intersect_domain(sched
, access
);
1255 /* Construct a schedule for iterating over all elements in the given
1256 * piece of an array. The schedule iterates over a copy of the piece
1257 * that is shifted to the origin.
1258 * We subsequently also perform the tiling/wrapping over the threads.
1260 * In particular, we tile the final iterators so that the final thread
1261 * dimension runs over the final array dimension.
1262 * However, if those final iterators have only a single iteration,
1263 * we try to tile earlier iterators instead.
1265 static __isl_give isl_union_map
*access_schedule(struct cuda_gen
*gen
,
1266 __isl_take isl_set
*access
, struct cuda_array_ref_group
*group
)
1270 isl_union_map
*usched
;
1273 unsigned nvar
= isl_set_dim(access
, isl_dim_set
);
1277 sched
= shift_access(access
, group
);
1279 n_tile
= gen
->n_block
;
1280 if (n_tile
> nvar
) {
1282 sched
= isl_map_insert_dims(sched
,
1283 isl_dim_out
, 0, n_tile
- nvar
);
1284 for (i
= 0; i
< n_tile
- nvar
; ++i
)
1285 sched
= isl_map_fix_si(sched
, isl_dim_out
, i
, 0);
1289 first
= nvar
- n_tile
;
1291 for (; first
> 0; first
--)
1292 if (!isl_map_plain_is_fixed(sched
, isl_dim_out
,
1293 first
+ n_tile
- 1, NULL
))
1296 dim
= isl_map_get_space(sched
);
1297 dim
= isl_space_params(dim
);
1298 if (gen
->options
->wrap
)
1299 tiling
= wrap(isl_space_copy(dim
), nvar
, first
,
1300 n_tile
, gen
->block_dim
);
1302 tiling
= tile(isl_space_copy(dim
), nvar
, first
,
1303 n_tile
, gen
->block_dim
);
1304 sched
= isl_map_apply_range(sched
, tiling
);
1306 par
= parametrization(dim
, nvar
+ n_tile
, first
+ n_tile
, n_tile
, "t");
1307 usched
= isl_union_map_from_map(sched
);
1308 usched
= isl_union_map_intersect_range(usched
,
1309 isl_union_set_from_set(par
));
1311 usched
= scale_access_tile_loops(gen
, usched
, nvar
+ n_tile
,
1317 static void print_shared_access(struct cuda_gen
*gen
,
1318 __isl_keep isl_set
*shared_domain
, __isl_take isl_set
*access
,
1319 const char *type
, struct cuda_array_ref_group
*group
)
1321 const char *array_name
;
1324 isl_union_map
*sched
;
1325 unsigned nvar
= isl_set_dim(access
, isl_dim_set
);
1328 ctx
= isl_set_get_ctx(access
);
1329 array_name
= isl_set_get_tuple_name(access
);
1330 name
= isl_alloc_array(ctx
, char,
1331 strlen(type
) + sizeof("_shared_") + strlen(array_name
) + 20);
1332 if (group
->array
->n_group
> 1)
1333 sprintf(name
, "%s_shared_%s_%d", type
, array_name
, group
->nr
);
1335 sprintf(name
, "%s_shared_%s", type
, array_name
);
1336 access
= isl_set_set_tuple_name(access
, name
);
1339 sched
= access_schedule(gen
, access
, group
);
1341 n_tile
= gen
->n_block
;
1345 print_shared_body(gen
, shared_domain
, sched
, nvar
+ n_tile
, NULL
, -1);
1347 isl_union_map_free(sched
);
1350 /* Return the union of all read (read = 1) and/or write (write = 1)
1351 * access relations in the group.
1353 static __isl_give isl_union_map
*group_access_relation(
1354 struct cuda_array_ref_group
*group
, int read
, int write
)
1357 isl_union_map
*access
;
1359 access
= isl_union_map_empty(isl_map_get_space(group
->access
));
1360 for (i
= 0; i
< group
->n_ref
; ++i
) {
1363 if (!((read
&& group
->refs
[i
]->read
) ||
1364 (write
&& group
->refs
[i
]->write
)))
1366 map_i
= isl_map_copy(group
->refs
[i
]->access
);
1367 access
= isl_union_map_union(access
,
1368 isl_union_map_from_map(map_i
));
1374 /* Check that none of the shared memory tiles involve any strides.
1376 static int no_strides(struct cuda_array_ref_group
*group
)
1379 int n_index
= group
->array
->n_index
;
1381 for (i
= 0; i
< n_index
; ++i
)
1382 if (group
->shared_bound
[i
].shift
)
1388 /* Return a set containing the values of the given index i
1389 * of the elements in the array tile in global memory that corresponds
1390 * to the shared memory copy.
1391 * In particular, if a is the index, we return a set with constraints
1393 * tile_offset <= a <= tile_offset + tile_size - 1
1397 * 0 <= a <= array_size - 1
1400 static __isl_give isl_set
*group_tile_dim(struct cuda_array_ref_group
*group
,
1403 isl_basic_set
*tile
;
1406 isl_local_space
*ls
;
1411 aff
= isl_aff_copy(group
->shared_bound
[i
].lb
);
1412 aff
= isl_aff_add_dims(aff
, isl_dim_in
, 1);
1413 ls
= isl_aff_get_domain_local_space(aff
);
1414 aff
= isl_aff_neg(aff
);
1415 aff
= isl_aff_add_coefficient_si(aff
, isl_dim_in
, 0, 1);
1416 c
= isl_inequality_from_aff(isl_aff_copy(aff
));
1417 tile
= isl_basic_set_from_constraint(c
);
1419 aff
= isl_aff_neg(aff
);
1420 aff
= isl_aff_add_constant(aff
, group
->shared_bound
[i
].size
);
1421 aff
= isl_aff_add_constant_si(aff
, -1);
1422 c
= isl_inequality_from_aff(aff
);
1423 tile
= isl_basic_set_add_constraint(tile
, c
);
1425 aff
= isl_aff_zero_on_domain(ls
);
1426 aff
= isl_aff_add_coefficient_si(aff
, isl_dim_in
, 0, 1);
1427 c
= isl_inequality_from_aff(aff
);
1428 tile
= isl_basic_set_add_constraint(tile
, c
);
1430 bound
= isl_pw_aff_copy(group
->array
->bound
[i
]);
1431 bound
= isl_pw_aff_add_dims(bound
, isl_dim_in
, 1);
1432 ls
= isl_local_space_from_space(isl_pw_aff_get_domain_space(bound
));
1433 aff
= isl_aff_zero_on_domain(ls
);
1434 aff
= isl_aff_add_coefficient_si(aff
, isl_dim_in
, 0, 1);
1435 aff
= isl_aff_add_constant_si(aff
, 1);
1436 dom
= isl_pw_aff_domain(isl_pw_aff_copy(bound
));
1438 tile_set
= isl_pw_aff_ge_set(bound
, isl_pw_aff_alloc(dom
, aff
));
1439 tile_set
= isl_set_align_params(tile_set
, isl_basic_set_get_space(tile
));
1440 tile_set
= isl_set_intersect(tile_set
, isl_set_from_basic_set(tile
));
1445 /* Return a set containing the elements in the array tile in
1446 * global memory that corresponds to the shared memory copy.
1448 static __isl_give isl_set
*group_tile(struct cuda_array_ref_group
*group
)
1451 int n_index
= group
->array
->n_index
;
1454 tile
= group_tile_dim(group
, 0);
1455 for (i
= 1; i
< n_index
; ++i
) {
1458 tile_i
= group_tile_dim(group
, i
);
1459 tile
= isl_set_flat_product(tile
, tile_i
);
1462 tile
= isl_set_set_tuple_name(tile
, group
->array
->name
);
1467 /* Print code for reading into or writing from shared memory
1468 * the given array reference group.
1470 * sched maps the original iteration domains to the shared memory tile loops.
1472 * If we are performing a read from global memory to shared memory,
1473 * if the array involved is not a scalar and if the definition of the
1474 * shared memory tiles does not involve any strides, then we copy
1475 * the entire tile to shared memory. This may result in some extra
1476 * elements getting copied, but it should lead to simpler code
1477 * (which means that fewer registers may be needed) and less divergence.
1479 * Otherwise, we only copy the elements that will be read or have been written
1482 * Note that the absence of stride requirement can easily be lifted.
1483 * We would just need to add constraints of the form
1485 * shift + a = stride * alpha
1487 static int print_group_shared_accesses(struct cuda_gen
*gen
,
1488 struct cuda_array_ref_group
*group
, const char *type
,
1489 __isl_keep isl_set
*shared_domain
, __isl_keep isl_union_map
*sched
)
1492 isl_union_map
*access
;
1493 isl_union_set
*uset
;
1494 isl_set
*access_set
;
1496 if (group
->private_bound
)
1498 if (!group
->shared_bound
)
1501 read
= !strcmp(type
, "read");
1503 access
= group_access_relation(group
, read
, !read
);
1504 access
= isl_union_map_apply_domain(access
, isl_union_map_copy(sched
));
1505 uset
= isl_union_map_range(access
);
1507 if (isl_union_set_is_empty(uset
)) {
1508 isl_union_set_free(uset
);
1512 if (read
&& group
->array
->n_index
> 0 && no_strides(group
)) {
1513 isl_union_set_free(uset
);
1514 access_set
= group_tile(group
);
1515 print_shared_access(gen
, shared_domain
, access_set
,
1520 access_set
= isl_set_from_union_set(uset
);
1521 access_set
= isl_set_coalesce(access_set
);
1523 print_shared_access(gen
, shared_domain
, access_set
, type
, group
);
1528 /* Print code for reading into or writing from shared memory at
1529 * the given level (-1 for innermost).
1531 * If we are not printing at the innermost level, then the dimensionality
1532 * of shared_domain may be smaller than gen->shared_len.
1533 * As the rest of the code assumes that the domain of access has
1534 * gen->shared_len dimensions, we therefore may need to embed this domain
1535 * in a higher dimensional space after intersection with shared_domain.
1537 static void print_shared_accesses(struct cuda_gen
*gen
,
1538 __isl_keep isl_set
*shared_domain
, __isl_keep isl_union_map
*access
,
1539 const char *type
, int level
)
1545 int shared_len
= isl_set_dim(shared_domain
, isl_dim_set
);
1547 isl_union_map
*sched
;
1549 shared_domain
= isl_set_copy(shared_domain
);
1550 sched
= isl_union_map_copy(gen
->tiled_sched
);
1551 dim
= isl_union_map_get_space(sched
);
1552 proj
= projection(dim
, gen
->tiled_len
, shared_len
);
1553 sched
= isl_union_map_apply_range(sched
, isl_union_map_from_map(proj
));
1554 sched
= isl_union_map_intersect_range(sched
,
1555 isl_union_set_from_set(isl_set_copy(shared_domain
)));
1556 if (shared_len
!= gen
->shared_len
) {
1557 dim
= isl_union_map_get_space(sched
);
1558 proj
= projection(dim
, gen
->shared_len
, shared_len
);
1559 proj
= isl_map_reverse(proj
);
1560 shared_domain
= isl_set_apply(shared_domain
,
1561 isl_map_copy(proj
));
1562 sched
= isl_union_map_apply_range(sched
,
1563 isl_union_map_from_map(proj
));
1566 dim
= isl_union_map_get_space(sched
);
1567 par
= parametrization(dim
, gen
->shared_len
, 0, gen
->shared_len
, "g");
1568 sched
= isl_union_map_intersect_range(sched
,
1569 isl_union_set_from_set(par
));
1571 for (i
= 0; i
< gen
->n_array
; ++i
) {
1572 struct cuda_array_info
*array
= &gen
->array
[i
];
1574 if (gen
->array
[i
].print_shared_level
!= level
)
1577 for (j
= 0; j
< array
->n_group
; ++j
) {
1578 if (print_group_shared_accesses(gen
, array
->groups
[j
],
1579 type
, shared_domain
, sched
))
1584 isl_union_map_free(sched
);
1585 isl_set_free(shared_domain
);
1588 print_indent(gen
->cuda
.kernel_c
, gen
->kernel_code
.indent
);
1589 fprintf(gen
->cuda
.kernel_c
, "__syncthreads();\n");
1593 /* Given an index expression into a tile of an array, adjust the expression
1594 * to a shift of the tile to the origin
1595 * (based on the lower bounds in array->shared_bound).
1596 * If the index is strided, then we first add
1597 * bound->shift and divide by bound->stride.
1599 static __isl_give isl_aff
*shift_index(__isl_take isl_aff
*aff
,
1600 struct cuda_array_info
*array
,
1601 struct cuda_array_bound
*bound
, __isl_take isl_set
*domain
)
1607 shift
= bound
->shift
;
1608 shift
= isl_aff_copy(shift
);
1609 shift
= isl_aff_project_domain_on_params(shift
);
1610 shift
= isl_aff_align_params(shift
, isl_aff_get_space(aff
));
1611 aff
= isl_aff_add(aff
, shift
);
1612 aff
= isl_aff_scale_down(aff
, bound
->stride
);
1615 lb
= isl_aff_copy(bound
->lb
);
1616 lb
= isl_aff_project_domain_on_params(lb
);
1618 lb
= isl_aff_align_params(lb
, isl_aff_get_space(aff
));
1620 aff
= isl_aff_sub(aff
, lb
);
1621 aff
= isl_aff_gist(aff
, domain
);
1626 /* This function is called for each access to an array in some statement
1627 * in the original code.
1628 * Replace that access by an access to shared or (linearized) global memory.
1629 * Since the array in shared memory is just
1630 * a shifted copy of part of the original array, we simply need
1631 * to subtract the lower bound, which was computed
1632 * in can_tile_for_shared_memory.
1633 * If any of the indices is strided, then we first add
1634 * shared_bound[i].shift and divide by shared_bound[i].stride.
1636 * If the given array is accessed directly from global memory,
1637 * we don't need to perform any shifting and simply simplify
1638 * expression in the context of the domain instead.
1640 * If the array space (range of access) has no name, then we are
1641 * accessing an iterator in the original program.
1643 static void print_access(struct cuda_gen
*gen
, __isl_take isl_map
*access
,
1649 struct cuda_array_info
*array
= NULL
;
1654 struct cuda_array_bound
*bounds
= NULL
;
1656 access
= isl_map_align_params(access
,
1657 isl_set_get_space(gen
->stmt_domain
));
1659 data_set
= isl_set_apply(isl_set_copy(gen
->stmt_domain
), access
);
1661 name
= isl_set_get_tuple_name(data_set
);
1664 fprintf(gen
->cuda
.kernel_c
, "(");
1666 struct cuda_array_ref_group
*group
;
1668 for (i
= 0; i
< gen
->n_array
; ++i
) {
1669 if (strcmp(name
, gen
->array
[i
].name
))
1671 array
= &gen
->array
[i
];
1674 group
= array
->groups
[group_nr
];
1675 bounds
= group
->private_bound
;
1677 bounds
= group
->shared_bound
;
1679 if (!bounds
&& cuda_array_is_scalar(array
))
1680 fprintf(gen
->cuda
.kernel_c
, "*");
1681 print_array_name(gen
->cuda
.kernel_c
, group
);
1683 if (cuda_array_is_scalar(array
)) {
1684 isl_set_free(data_set
);
1688 fprintf(gen
->cuda
.kernel_c
, "[");
1692 n_index
= isl_set_dim(data_set
, isl_dim_set
);
1693 aff
= isl_set_affine_hull(data_set
);
1695 prn
= isl_printer_to_file(gen
->ctx
, gen
->cuda
.kernel_c
);
1696 prn
= isl_printer_set_output_format(prn
, ISL_FORMAT_C
);
1699 for (i
= 0; i
+ 1 < n_index
; ++i
)
1700 prn
= isl_printer_print_str(prn
, "(");
1702 for (i
= 0; i
< n_index
; ++i
) {
1707 ok
= isl_basic_set_has_defining_equality(aff
,
1708 isl_dim_out
, i
, &c
);
1710 index
= isl_constraint_get_bound(c
, isl_dim_out
, i
);
1711 isl_constraint_free(c
);
1712 index
= isl_aff_project_domain_on_params(index
);
1715 prn
= isl_printer_print_aff(prn
, index
);
1716 isl_aff_free(index
);
1720 domain
= isl_set_copy(gen
->stmt_domain
);
1721 domain
= isl_set_project_out(domain
, isl_dim_set
, 0,
1722 isl_set_dim(domain
, isl_dim_set
));
1724 index
= isl_aff_gist(index
, domain
);
1726 index
= shift_index(index
, array
, &bounds
[i
], domain
);
1730 prn
= isl_printer_print_str(prn
, ") * (");
1731 prn
= isl_printer_print_pw_aff(prn
,
1732 array
->local_bound
[i
]);
1733 prn
= isl_printer_print_str(prn
, ") + ");
1735 prn
= isl_printer_print_str(prn
, "][");
1737 prn
= isl_printer_print_aff(prn
, index
);
1738 isl_aff_free(index
);
1741 prn
= isl_printer_print_str(prn
, ")");
1743 prn
= isl_printer_print_str(prn
, "]");
1744 isl_printer_free(prn
);
1746 isl_basic_set_free(aff
);
1749 static struct cuda_stmt_access
*print_expr(struct cuda_gen
*gen
, FILE *out
,
1750 struct pet_expr
*expr
, struct cuda_stmt_access
*access
, int outer
)
1754 switch (expr
->type
) {
1755 case pet_expr_double
:
1756 fprintf(out
, "%g", expr
->d
);
1758 case pet_expr_access
:
1759 print_access(gen
, isl_map_copy(access
->access
), access
->group
);
1760 access
= access
->next
;
1762 case pet_expr_unary
:
1765 fprintf(out
, " %s ", pet_op_str(expr
->op
));
1766 access
= print_expr(gen
, out
, expr
->args
[pet_un_arg
],
1771 case pet_expr_binary
:
1774 access
= print_expr(gen
, out
, expr
->args
[pet_bin_lhs
],
1776 fprintf(out
, " %s ", pet_op_str(expr
->op
));
1777 access
= print_expr(gen
, out
, expr
->args
[pet_bin_rhs
],
1782 case pet_expr_ternary
:
1785 access
= print_expr(gen
, out
, expr
->args
[pet_ter_cond
],
1787 fprintf(out
, " ? ");
1788 access
= print_expr(gen
, out
, expr
->args
[pet_ter_true
],
1790 fprintf(out
, " : ");
1791 access
= print_expr(gen
, out
, expr
->args
[pet_ter_false
],
1797 fprintf(out
, "%s(", expr
->name
);
1798 for (i
= 0; i
< expr
->n_arg
; ++i
) {
1801 access
= print_expr(gen
, out
, expr
->args
[i
],
1809 static void print_stmt_body(struct cuda_gen
*gen
,
1810 FILE *out
, struct cuda_stmt
*stmt
)
1812 print_expr(gen
, out
, stmt
->body
, stmt
->accesses
, 1);
1813 fprintf(out
, ";\n");
1816 /* This function is called for each leaf in the innermost clast,
1817 * i.e., for each statement.
1818 * We print the statement body, simplifying the accesses based
1821 static void print_statement(struct gpucode_info
*code
,
1822 struct clast_user_stmt
*u
)
1824 struct cuda_gen
*gen
= code
->user
;
1827 isl_set
*stmt_domain
;
1828 isl_union_map
*stmt_sched
;
1829 isl_union_set
*uset
;
1831 struct cuda_stmt
*stmt
;
1833 nr
= atoi(u
->statement
->name
+ 2);
1834 stmt
= &gen
->stmts
[nr
];
1836 stmt_domain
= extract_host_domain(u
);
1838 stmt_sched
= isl_union_map_intersect_range(
1839 isl_union_map_copy(gen
->local_sched
),
1840 isl_union_set_from_set(extend(stmt_domain
,
1841 gen
->thread_tiled_len
)));
1842 dim
= isl_union_map_get_space(stmt_sched
);
1843 par
= parametrization(dim
, gen
->thread_tiled_len
, 0,
1844 gen
->thread_tiled_len
, "c");
1845 stmt_sched
= isl_union_map_intersect_range(stmt_sched
,
1846 isl_union_set_from_set(par
));
1848 uset
= isl_union_map_domain(stmt_sched
);
1849 dim
= isl_union_set_get_space(uset
);
1850 dim
= isl_space_add_dims(dim
, isl_dim_set
,
1851 isl_set_dim(stmt
->domain
, isl_dim_set
));
1852 dim
= isl_space_set_tuple_name(dim
, isl_dim_set
, u
->statement
->name
);
1853 gen
->stmt_domain
= isl_union_set_extract_set(uset
, dim
);
1854 isl_union_set_free(uset
);
1856 print_indent(code
->dst
, code
->indent
);
1857 print_stmt_body(gen
, code
->dst
, stmt
);
1859 isl_set_free(gen
->stmt_domain
);
1862 /* Print an access to the element in the global memory copy of the
1863 * given array that corresponds to element [aff[0]][aff[1]]...
1864 * of the original array.
1865 * The copy in global memory has been linearized, so we need to take
1866 * the array size into account.
1868 static void print_private_global_index(isl_ctx
*ctx
, FILE *out
,
1869 struct cuda_array_info
*array
, __isl_keep isl_aff
**aff
)
1874 fprintf(out
, "%s[", array
->name
);
1875 prn
= isl_printer_to_file(ctx
, out
);
1876 prn
= isl_printer_set_output_format(prn
, ISL_FORMAT_C
);
1877 for (i
= 0; i
+ 1 < array
->n_index
; ++i
)
1878 prn
= isl_printer_print_str(prn
, "(");
1879 for (i
= 0; i
< array
->n_index
; ++i
) {
1881 prn
= isl_printer_print_str(prn
, ") * (");
1882 prn
= isl_printer_print_pw_aff(prn
,
1883 array
->local_bound
[i
]);
1884 prn
= isl_printer_print_str(prn
, ") + ");
1886 prn
= isl_printer_print_aff(prn
, aff
[i
]);
1888 isl_printer_free(prn
);
1892 /* Print an access to the element in the shared memory copy of the
1893 * given array reference group that corresponds to element [affs[0]][affs[1]]...
1894 * of the original array.
1895 * Since the array in shared memory is just a shifted copy of part
1896 * of the original array, we simply need to subtract the lower bound,
1897 * which was computed in can_tile_for_shared_memory.
1898 * If any of the indices is strided, then we first add
1899 * shared_bound[i].shift and divide by shared_bound[i].stride.
1901 static void print_private_local_index(isl_ctx
*ctx
, FILE *out
,
1902 struct cuda_array_ref_group
*group
,
1903 __isl_keep isl_aff
**affs
, __isl_keep isl_set
*domain
)
1907 struct cuda_array_info
*array
= group
->array
;
1908 struct cuda_array_bound
*bounds
= group
->private_bound
;
1910 print_array_name(out
, group
);
1911 for (i
= 0; i
< array
->n_index
; ++i
) {
1912 isl_aff
*aff
= isl_aff_copy(affs
[i
]);
1914 aff
= shift_index(aff
, array
, &bounds
[i
], isl_set_copy(domain
));
1917 prn
= isl_printer_to_file(ctx
, out
);
1918 prn
= isl_printer_set_output_format(prn
, ISL_FORMAT_C
);
1919 prn
= isl_printer_print_aff(prn
, aff
);
1920 isl_printer_free(prn
);
1926 /* This function is called for each leaf in the clast of the code
1927 * for copying to or from private memory.
1928 * The statement name is read_private_<array> or write_private_<array>.
1930 * The schedule iterates over the array elements, so we can use
1931 * the domain of private_sched at the current scheduling position
1932 * as the index of the array.
1934 static void print_private_copy_statement(struct gpucode_info
*code
,
1935 struct clast_user_stmt
*u
)
1937 struct cuda_gen
*gen
= code
->user
;
1940 struct cuda_array_ref_group
*group
= gen
->private_group
;
1952 read
= !strncmp(u
->statement
->name
, "read", 4);
1954 domain
= extract_host_domain(u
);
1957 sched
= isl_map_copy(gen
->private_sched
);
1958 sched
= isl_map_reverse(sched
);
1959 sched
= isl_map_intersect_domain(sched
, domain
);
1960 n_in
= isl_map_dim(sched
, isl_dim_in
);
1961 n_out
= isl_map_dim(sched
, isl_dim_out
);
1962 dim
= isl_map_get_space(sched
);
1963 dim
= isl_space_drop_dims(dim
, isl_dim_in
, 0, n_in
);
1964 dim
= isl_space_drop_dims(dim
, isl_dim_out
, 0, n_out
);
1965 param
= parametrization(dim
, n_in
, 0, n_in
, "c");
1966 sched
= isl_map_align_params(sched
, isl_set_get_space(param
));
1967 sched
= isl_map_intersect_domain(sched
, param
);
1968 index
= isl_map_range(sched
);
1969 domain
= isl_set_copy(index
);
1970 aff
= isl_set_affine_hull(index
);
1971 domain
= isl_set_project_out(domain
, isl_dim_set
, 0, n_out
);
1973 ctx
= isl_basic_set_get_ctx(aff
);
1974 affs
= isl_alloc_array(ctx
, isl_aff
*, n_out
);
1977 for (i
= 0; i
< n_out
; ++i
) {
1981 ok
= isl_basic_set_has_defining_equality(aff
,
1982 isl_dim_set
, i
, &c
);
1984 affs
[i
] = isl_constraint_get_bound(c
, isl_dim_set
, i
);
1985 isl_constraint_free(c
);
1986 affs
[i
] = isl_aff_project_domain_on_params(affs
[i
]);
1989 print_indent(code
->dst
, code
->indent
);
1991 print_private_local_index(ctx
, code
->dst
, group
, affs
, domain
);
1992 fprintf(code
->dst
, " = ");
1993 print_private_global_index(ctx
, code
->dst
, group
->array
, affs
);
1995 print_private_global_index(ctx
, code
->dst
, group
->array
, affs
);
1996 fprintf(code
->dst
, " = ");
1997 print_private_local_index(ctx
, code
->dst
, group
, affs
, domain
);
1999 fprintf(code
->dst
, ";\n");
2001 for (i
= 0; i
< n_out
; ++i
)
2002 isl_aff_free(affs
[i
]);
2005 isl_basic_set_free(aff
);
2006 isl_set_free(domain
);
2009 static void print_private_access(struct cuda_gen
*gen
,
2010 __isl_keep isl_set
*shared_domain
, __isl_take isl_set
*access
,
2011 const char *type
, struct cuda_array_ref_group
*group
)
2013 const char *array_name
;
2016 unsigned nvar
= isl_set_dim(access
, isl_dim_set
);
2017 isl_union_map
*usched
;
2019 if (isl_set_fast_is_empty(access
)) {
2020 isl_set_free(access
);
2024 ctx
= isl_set_get_ctx(access
);
2025 array_name
= isl_set_get_tuple_name(access
);
2026 name
= isl_alloc_array(ctx
, char,
2027 strlen(type
) + sizeof("_private_") + strlen(array_name
) + 20);
2028 if (group
->array
->n_group
> 1)
2029 sprintf(name
, "%s_private_%s_%d", type
, array_name
, group
->nr
);
2031 sprintf(name
, "%s_private_%s", type
, array_name
);
2032 access
= isl_set_set_tuple_name(access
, name
);
2035 gen
->private_sched
= shift_access(access
, group
);
2036 gen
->private_group
= group
;
2038 usched
= isl_union_map_from_map(isl_map_copy(gen
->private_sched
));
2039 print_shared_body(gen
, shared_domain
, usched
, nvar
,
2040 &print_private_copy_statement
, 1);
2041 isl_union_map_free(usched
);
2043 isl_map_free(gen
->private_sched
);
2046 /* Print code for reading into or writing from private memory
2047 * the given array reference group.
2049 * sched maps the original iteration domains to the shared memory tile loops.
2051 static void print_group_private_accesses(struct cuda_gen
*gen
,
2052 struct cuda_array_ref_group
*group
,
2053 const char *type
, __isl_keep isl_set
*shared_domain
,
2054 unsigned first_shared
, int shared_len
, __isl_keep isl_union_map
*sched
)
2057 isl_union_map
*access
;
2058 isl_union_set
*uset
;
2059 isl_set
*access_set
;
2061 if (!group
->private_bound
)
2064 read
= !strcmp(type
, "read");
2066 access
= group_access_relation(group
, read
, !read
);
2067 access
= isl_union_map_apply_domain(access
, isl_union_map_copy(sched
));
2068 access
= isl_union_map_intersect(access
,
2069 isl_union_map_copy(gen
->private_access
));
2070 uset
= isl_union_map_range(access
);
2072 if (isl_union_set_is_empty(uset
)) {
2073 isl_union_set_free(uset
);
2077 access_set
= isl_set_from_union_set(uset
);
2078 access_set
= isl_set_coalesce(access_set
);
2079 access_set
= isl_set_eliminate(access_set
, isl_dim_param
,
2080 first_shared
+ shared_len
,
2081 gen
->shared_len
- shared_len
);
2083 print_private_access(gen
, shared_domain
, access_set
, type
, group
);
2086 /* Print code for reading into or writing from private memory at
2087 * the given level (-1 for innermost).
2089 * If we are not printing at the innermost level, then the dimensionality
2090 * of shared_domain may be smaller than gen->shared_len.
2091 * As the rest of the code assumes that the domain of access has
2092 * gen->shared_len dimensions, we therefore may need to embed this domain
2093 * in a higher dimensional space after intersection with shared_domain.
2095 * This code is very similar to print_shared_accesses.
2096 * The main difference is that we to take into account gen->private_access.
2098 static void print_private_accesses(struct cuda_gen
*gen
,
2099 __isl_keep isl_set
*shared_domain
, __isl_keep isl_union_map
*access
,
2100 const char *type
, int level
)
2105 int shared_len
= isl_set_dim(shared_domain
, isl_dim_set
);
2106 unsigned first_shared
;
2107 isl_union_map
*sched
;
2109 shared_domain
= isl_set_copy(shared_domain
);
2110 sched
= isl_union_map_copy(gen
->tiled_sched
);
2111 dim
= isl_union_map_get_space(sched
);
2112 first_shared
= isl_space_dim(dim
, isl_dim_param
);
2113 proj
= projection(dim
, gen
->tiled_len
, shared_len
);
2114 sched
= isl_union_map_apply_range(sched
, isl_union_map_from_map(proj
));
2115 sched
= isl_union_map_intersect_range(sched
,
2116 isl_union_set_from_set(isl_set_copy(shared_domain
)));
2117 if (shared_len
!= gen
->shared_len
) {
2118 dim
= isl_union_map_get_space(sched
);
2119 proj
= projection(dim
, gen
->shared_len
, shared_len
);
2120 proj
= isl_map_reverse(proj
);
2121 shared_domain
= isl_set_apply(shared_domain
,
2122 isl_map_copy(proj
));
2123 sched
= isl_union_map_apply_range(sched
,
2124 isl_union_map_from_map(proj
));
2127 for (i
= 0; i
< gen
->n_array
; ++i
) {
2128 struct cuda_array_info
*array
= &gen
->array
[i
];
2130 if (gen
->array
[i
].print_shared_level
!= level
)
2133 for (j
= 0; j
< array
->n_group
; ++j
)
2134 print_group_private_accesses(gen
, array
->groups
[j
],
2135 type
, shared_domain
,
2136 first_shared
, shared_len
, sched
);
2139 isl_union_map_free(sched
);
2140 isl_set_free(shared_domain
);
2143 /* Set unroll[j] if the input dimension j is involved in
2144 * the index expression represented by bmap.
2146 static int check_unroll(__isl_take isl_basic_map
*bmap
, void *user
)
2149 int n_in
= isl_basic_map_dim(bmap
, isl_dim_in
);
2150 int n_out
= isl_basic_map_dim(bmap
, isl_dim_out
);
2153 for (i
= 0; i
< n_out
; ++i
) {
2157 ok
= isl_basic_map_has_defining_equality(bmap
,
2158 isl_dim_out
, i
, &c
);
2160 for (j
= 0; j
< n_in
; ++j
)
2161 if (isl_constraint_involves_dims(c
, isl_dim_in
, j
, 1))
2163 isl_constraint_free(c
);
2166 isl_basic_map_free(bmap
);
2170 /* Given an array pos mapping input dimensions to the corresponding
2171 * output dimension, construct the corresponding map.
2173 static __isl_give isl_map
*permutation(__isl_take isl_space
*dim
,
2178 isl_basic_map
*bmap
;
2179 isl_local_space
*ls
;
2181 dim
= isl_space_add_dims(dim
, isl_dim_in
, len
);
2182 dim
= isl_space_add_dims(dim
, isl_dim_out
, len
);
2183 bmap
= isl_basic_map_universe(isl_space_copy(dim
));
2184 ls
= isl_local_space_from_space(dim
);
2186 for (i
= 0; i
< len
; ++i
) {
2187 c
= isl_equality_alloc(isl_local_space_copy(ls
));
2188 isl_constraint_set_coefficient_si(c
, isl_dim_in
, i
, -1);
2189 isl_constraint_set_coefficient_si(c
, isl_dim_out
, pos
[i
], 1);
2190 bmap
= isl_basic_map_add_constraint(bmap
, c
);
2192 isl_local_space_free(ls
);
2194 return isl_map_from_basic_map(bmap
);
2197 /* Find all loops involved in any of the index expressions for any of
2198 * the private accesses, move them innermost and then mark them as
2199 * requiring unrolling by setting gen->first_unroll.
2200 * The loops involved should all be parallel because of the checks
2201 * we performed in check_private_group_access. Moving them innermost
2202 * is therefore a valid transformation.
2204 static __isl_give isl_union_map
*interchange_for_unroll(struct cuda_gen
*gen
,
2205 __isl_take isl_union_map
*sched
)
2208 int unroll
[gen
->thread_tiled_len
];
2209 int perm
[gen
->thread_tiled_len
];
2212 int len
= gen
->shared_len
+ gen
->n_parallel
+ gen
->n_block
;
2214 gen
->first_unroll
= -1;
2216 for (i
= 0; i
< gen
->thread_tiled_len
; ++i
)
2218 for (i
= 0; i
< gen
->n_array
; ++i
) {
2219 struct cuda_array_info
*array
= &gen
->array
[i
];
2221 for (j
= 0; j
< array
->n_group
; ++j
) {
2222 isl_union_map
*access
;
2225 if (!array
->groups
[j
]->private_bound
)
2228 access
= group_access_relation(array
->groups
[j
], 1, 1);
2229 access
= isl_union_map_apply_domain(access
,
2230 isl_union_map_copy(sched
));
2232 acc
= isl_map_from_union_map(access
);
2233 isl_map_foreach_basic_map(acc
, &check_unroll
, unroll
);
2239 for (i
= 0; i
< gen
->shared_len
; ++i
)
2243 for (i
= gen
->shared_len
; i
< len
; ++i
)
2250 for (i
= len
; i
< gen
->thread_tiled_len
; ++i
)
2255 for (i
= 0; i
< gen
->thread_tiled_len
; ++i
)
2258 gen
->first_unroll
= 1 + j
;
2259 for (i
= 0; i
< len
; ++i
)
2263 dim
= isl_union_map_get_space(sched
);
2264 permute
= permutation(dim
, perm
, gen
->thread_tiled_len
);
2265 sched
= isl_union_map_apply_range(sched
,
2266 isl_union_map_from_map(permute
));
2271 /* This function is called for each leaf in the clast of the kernel code.
2272 * We first specialize the schedule to the site of the leaf and
2273 * print code for reading into shared memory, performing the actual
2274 * computations and writing from shared memory, with the required
2277 static void print_kernel_user(struct gpucode_info
*code
,
2278 struct clast_user_stmt
*u
)
2280 struct cuda_gen
*gen
= code
->user
;
2281 isl_set
*shared_domain
;
2283 shared_domain
= extract_entire_host_domain(u
);
2285 print_shared_accesses(gen
, shared_domain
, gen
->read
, "read", -1);
2287 print_private_accesses(gen
, shared_domain
, gen
->read
, "read", -1);
2289 print_shared_body(gen
, shared_domain
, gen
->local_sched
,
2290 gen
->thread_tiled_len
, &print_statement
,
2293 print_private_accesses(gen
, shared_domain
, gen
->write
, "write", -1);
2295 print_indent(gen
->cuda
.kernel_c
, gen
->kernel_code
.indent
);
2296 fprintf(gen
->cuda
.kernel_c
, "__syncthreads();\n");
2298 print_shared_accesses(gen
, shared_domain
, gen
->write
, "write", -1);
2300 isl_set_free(shared_domain
);
2303 /* Check if we need to perform any copying to shared memory at this level
2304 * and if so, print the copying instructions.
2305 * Any array for which we are allowed to print copying instructions at
2306 * this level, but haven't done so already, is printed.
2308 static void print_kernel_for_head(struct gpucode_info
*code
,
2309 struct clast_for
*f
)
2312 struct cuda_gen
*gen
= code
->user
;
2317 domain
= isl_set_from_cloog_domain(cloog_domain_copy(f
->domain
));
2318 level
= isl_set_dim(domain
, isl_dim_set
) - 1;
2320 for (i
= 0; i
< gen
->n_array
; ++i
) {
2321 if (gen
->array
[i
].print_shared_level
>= 0)
2323 if (gen
->array
[i
].last_shared
> level
)
2325 gen
->array
[i
].print_shared_level
= level
;
2330 print_shared_accesses(gen
, domain
, gen
->read
, "read", level
);
2331 print_private_accesses(gen
, domain
, gen
->read
, "read", level
);
2334 isl_set_free(domain
);
2337 /* Print instructions for copying from shared memory for each array
2338 * for which print_kernel_for_head has added copying instructions
2341 static void print_kernel_for_foot(struct gpucode_info
*code
,
2342 struct clast_for
*f
)
2345 struct cuda_gen
*gen
= code
->user
;
2350 domain
= isl_set_from_cloog_domain(cloog_domain_copy(f
->domain
));
2351 level
= isl_set_dim(domain
, isl_dim_set
) - 1;
2353 for (i
= 0; i
< gen
->n_array
; ++i
) {
2354 if (gen
->array
[i
].print_shared_level
!= level
)
2361 print_private_accesses(gen
, domain
, gen
->write
, "write", level
);
2362 print_shared_accesses(gen
, domain
, gen
->write
, "write", level
);
2365 isl_set_free(domain
);
2368 /* Use CLooG to generate code for the outer gen->shared_first loops
2369 * of the local schedule "sched".
2370 * The pretty printing of this code is handled by gpu_print_host_stmt,
2371 * which calls print_kernel_user for each iteration of the shared tile loops.
2373 static void print_cloog_kernel_body(struct cuda_gen
*gen
,
2374 __isl_keep isl_set
*context
, __isl_keep isl_union_map
*sched
)
2377 CloogOptions
*options
;
2378 CloogDomain
*cloog_context
;
2379 CloogUnionDomain
*ud
;
2381 struct clast_stmt
*stmt
;
2384 sched
= isl_union_map_copy(sched
);
2385 sched
= isl_union_map_align_params(sched
, isl_set_get_space(context
));
2387 options
= cloog_options_malloc(gen
->state
);
2388 options
->language
= CLOOG_LANGUAGE_C
;
2389 options
->strides
= 1;
2391 options
->stop
= gen
->shared_len
;
2392 options
->f
= gen
->tiled_len
;
2393 options
->l
= gen
->tiled_len
;
2394 options
->save_domains
= 1;
2395 options
->noscalars
= 1;
2397 ud
= cloog_union_domain_from_isl_union_map(sched
);
2398 for (i
= 0; i
< gen
->shared_len
; ++i
) {
2399 snprintf(name
, sizeof(name
), "g%d", i
);
2400 ud
= cloog_union_domain_set_name(ud
, CLOOG_SCAT
, i
, name
);
2402 cloog_context
= cloog_domain_from_isl_set(isl_set_copy(context
));
2403 input
= cloog_input_alloc(cloog_context
, ud
);
2405 stmt
= cloog_clast_create_from_input(input
, options
);
2407 gen
->kernel_code
.indent
= 4;
2408 gen
->kernel_code
.dst
= gen
->cuda
.kernel_c
;
2409 gen
->kernel_code
.print_user_stmt
= NULL
;
2410 gen
->kernel_code
.print_user_stmt_list
= &print_kernel_user
;
2411 gen
->kernel_code
.print_for_head
= &print_kernel_for_head
;
2412 gen
->kernel_code
.print_for_foot
= &print_kernel_for_foot
;
2413 gen
->kernel_code
.user
= gen
;
2414 gpu_print_host_stmt(&gen
->kernel_code
, stmt
);
2416 cloog_clast_free(stmt
);
2417 cloog_options_free(options
);
2420 static void print_kernel_iterators(struct cuda_gen
*gen
)
2423 const char *block_dims
[] = { "blockIdx.x", "blockIdx.y" };
2424 const char *thread_dims
[] = { "threadIdx.x", "threadIdx.y",
2427 if (gen
->n_grid
> 0) {
2428 print_indent(gen
->cuda
.kernel_c
, 4);
2429 fprintf(gen
->cuda
.kernel_c
, "int ");
2430 for (i
= 0; i
< gen
->n_grid
; ++i
) {
2432 fprintf(gen
->cuda
.kernel_c
, ", ");
2433 fprintf(gen
->cuda
.kernel_c
, "b%d = %s",
2434 i
, block_dims
[gen
->n_grid
- 1 - i
]);
2436 fprintf(gen
->cuda
.kernel_c
, ";\n");
2439 if (gen
->n_block
> 0) {
2440 print_indent(gen
->cuda
.kernel_c
, 4);
2441 fprintf(gen
->cuda
.kernel_c
, "int ");
2442 for (i
= 0; i
< gen
->n_block
; ++i
) {
2444 fprintf(gen
->cuda
.kernel_c
, ", ");
2445 fprintf(gen
->cuda
.kernel_c
, "t%d = %s",
2446 i
, thread_dims
[gen
->n_block
- 1 - i
]);
2448 fprintf(gen
->cuda
.kernel_c
, ";\n");
2452 static void print_group_shared_array(struct cuda_gen
*gen
,
2453 struct cuda_array_ref_group
*group
)
2456 struct cuda_array_bound
*bounds
;
2458 bounds
= group
->private_bound
;
2460 bounds
= group
->shared_bound
;
2464 print_indent(gen
->cuda
.kernel_c
, 4);
2465 fprintf(gen
->cuda
.kernel_c
, "%s%s ",
2466 group
->private_bound
? "" : "__shared__ ", group
->array
->type
);
2467 print_array_name(gen
->cuda
.kernel_c
, group
);
2468 for (j
= 0; j
< group
->array
->n_index
; ++j
) {
2469 fprintf(gen
->cuda
.kernel_c
, "[");
2470 isl_int_print(gen
->cuda
.kernel_c
, bounds
[j
].size
, 0);
2471 fprintf(gen
->cuda
.kernel_c
, "]");
2473 fprintf(gen
->cuda
.kernel_c
, ";\n");
2476 static void print_shared_arrays(struct cuda_gen
*gen
)
2480 for (i
= 0; i
< gen
->n_array
; ++i
) {
2481 struct cuda_array_info
*array
= &gen
->array
[i
];
2483 for (j
= 0; j
< array
->n_group
; ++j
)
2484 print_group_shared_array(gen
, array
->groups
[j
]);
2488 static void print_kernel_body(struct cuda_gen
*gen
,
2489 __isl_keep isl_set
*host_domain
, __isl_keep isl_union_map
*sched
)
2493 context
= isl_set_copy(host_domain
);
2494 context
= parametrize(context
, 0, gen
->tile_first
, "h");
2495 context
= isl_set_project_out(context
, isl_dim_set
, 0, gen
->tile_first
);
2496 context
= add_bounded_parameters(context
,
2497 gen
->n_grid
, gen
->grid_dim
, "b");
2499 print_kernel_iterators(gen
);
2500 print_shared_arrays(gen
);
2502 fprintf(gen
->cuda
.kernel_c
, "\n");
2504 print_cloog_kernel_body(gen
, context
, sched
);
2506 isl_set_free(context
);
2509 /* Given a constraint
2511 * a(p,i) + j = g f(e)
2513 * or -a(p,i) - j = g f(e) if sign < 0,
2514 * store a(p,i) in bound->shift and g (stride) in bound->stride.
2515 * a(p,i) is assumed to be an expression in only the parameters.
2517 static void extract_stride(__isl_keep isl_constraint
*c
,
2518 struct cuda_array_bound
*bound
, isl_int stride
, int sign
)
2526 isl_int_set(bound
->stride
, stride
);
2528 dim
= isl_constraint_get_space(c
);
2529 dim
= isl_space_params(dim
);
2531 nparam
= isl_space_dim(dim
, isl_dim_param
);
2535 isl_constraint_get_constant(c
, &v
);
2538 aff
= isl_aff_zero_on_domain(isl_local_space_from_space(dim
));
2539 aff
= isl_aff_set_constant(aff
, v
);
2541 for (i
= 0; i
< nparam
; ++i
) {
2542 isl_constraint_get_coefficient(c
, isl_dim_param
, i
, &v
);
2543 if (isl_int_is_zero(v
))
2547 aff
= isl_aff_add_coefficient(aff
, isl_dim_param
, i
, v
);
2555 /* Given an equality constraint of a map with a single output dimension j,
2556 * check if the constraint is of the form
2558 * a(p,i) + j = g f(e)
2560 * with a(p,i) an expression in the parameters and input dimensions
2561 * and f(e) an expression in the existentially quantified variables.
2562 * If so, and if g is larger than any such g from a previously considered
2563 * constraint, then call extract_stride. to record the stride information
2566 static int check_stride_constraint(__isl_take isl_constraint
*c
, void *user
)
2571 struct cuda_array_bound
*bound
= user
;
2574 isl_int_init(stride
);
2576 n_div
= isl_constraint_dim(c
, isl_dim_div
);
2577 isl_constraint_get_coefficient(c
, isl_dim_out
, 0, &v
);
2579 if (n_div
&& (isl_int_is_one(v
) || isl_int_is_negone(v
))) {
2580 int s
= isl_int_sgn(v
);
2581 isl_int_set_si(stride
, 0);
2582 for (i
= 0; i
< n_div
; ++i
) {
2583 isl_constraint_get_coefficient(c
, isl_dim_div
, i
, &v
);
2584 isl_int_gcd(stride
, stride
, v
);
2586 if (!isl_int_is_zero(stride
) &&
2587 isl_int_gt(stride
, bound
->stride
))
2588 extract_stride(c
, bound
, stride
, s
);
2591 isl_int_clear(stride
);
2594 isl_constraint_free(c
);
2598 /* Given contraints on an array index i, check if we can find
2599 * a shift a(p) and a stride g such that
2601 * a(p) + i = 0 mod g
2603 * If so, record the information in bound and apply the mapping
2604 * i -> (i + a(p))/g to the array index in bounds and return
2605 * the new constraints.
2606 * If not, simply return the original constraints.
2608 static __isl_give isl_basic_map
*check_stride(struct cuda_gen
*gen
,
2609 struct cuda_array_bound
*bound
, __isl_take isl_basic_map
*bounds
)
2612 isl_basic_map
*shift
;
2615 isl_int_set_si(bound
->stride
, -1);
2617 aff
= isl_basic_map_affine_hull(isl_basic_map_copy(bounds
));
2619 isl_basic_map_foreach_constraint(aff
, &check_stride_constraint
, bound
);
2621 isl_basic_map_free(aff
);
2623 if (isl_int_is_neg(bound
->stride
))
2626 aff_shift
= isl_aff_copy(bound
->shift
);
2627 aff_shift
= isl_aff_add_dims(aff_shift
, isl_dim_in
, 1);
2628 aff_shift
= isl_aff_add_coefficient_si(aff_shift
, isl_dim_in
, 0, 1);
2629 aff_shift
= isl_aff_scale_down(aff_shift
, bound
->stride
);
2630 shift
= isl_basic_map_from_aff(aff_shift
);
2632 bound
->shift_map
= isl_basic_map_copy(shift
);
2633 bounds
= isl_basic_map_apply_range(bounds
, shift
);
2638 struct cuda_size_info
{
2639 isl_basic_set
*bset
;
2640 struct cuda_array_bound
*bound
;
2644 /* Given a constraint from the basic set describing the bounds on
2645 * an array index, check if it is a lower bound, say m i >= b(x), and,
2646 * if so, check whether the expression "i - ceil(b(x)/m) + 1" has a constant
2647 * upper bound. If so, and if this bound is smaller than any bound
2648 * derived from earlier constraints, set the size to this bound on
2649 * the expression and the lower bound to ceil(b(x)/m).
2651 static int compute_size_in_direction(__isl_take isl_constraint
*c
, void *user
)
2653 struct cuda_size_info
*size
= user
;
2658 nparam
= isl_basic_set_dim(size
->bset
, isl_dim_param
);
2659 n_div
= isl_constraint_dim(c
, isl_dim_div
);
2661 if (isl_constraint_involves_dims(c
, isl_dim_div
, 0, n_div
)) {
2662 isl_constraint_free(c
);
2668 isl_constraint_get_coefficient(c
, isl_dim_set
, size
->pos
, &v
);
2670 if (isl_int_is_pos(v
)) {
2673 enum isl_lp_result res
;
2675 aff
= isl_constraint_get_bound(c
, isl_dim_set
, size
->pos
);
2676 aff
= isl_aff_ceil(aff
);
2678 lb
= isl_aff_copy(aff
);
2680 aff
= isl_aff_neg(aff
);
2681 aff
= isl_aff_add_coefficient_si(aff
, isl_dim_in
, size
->pos
, 1);
2683 res
= isl_basic_set_max(size
->bset
, aff
, &v
);
2686 if (res
== isl_lp_ok
) {
2687 isl_int_add_ui(v
, v
, 1);
2688 if (isl_int_is_neg(size
->bound
->size
) ||
2689 isl_int_lt(v
, size
->bound
->size
)) {
2690 isl_int_set(size
->bound
->size
, v
);
2691 lb
= isl_aff_drop_dims(lb
, isl_dim_in
,
2693 isl_aff_free(size
->bound
->lb
);
2694 size
->bound
->lb
= isl_aff_copy(lb
);
2701 isl_constraint_free(c
);
2706 /* Given a basic map "bounds" that maps parameters and input dimensions
2707 * to a single output dimension, look for an expression in the parameters
2708 * and input dimensions such that the range of the output dimension shifted
2709 * by this expression is a constant.
2711 * In particular, we currently only consider lower bounds on the output
2712 * dimension as candidate expressions.
2714 static int compute_array_dim_size(struct cuda_gen
*gen
,
2715 struct cuda_array_bound
*bound
, __isl_take isl_basic_map
*bounds
)
2717 struct cuda_size_info size
;
2719 bounds
= check_stride(gen
, bound
, bounds
);
2721 isl_int_set_si(bound
->size
, -1);
2725 size
.pos
= isl_basic_map_dim(bounds
, isl_dim_in
);
2726 size
.bset
= isl_basic_map_wrap(bounds
);
2727 size
.bset
= isl_basic_set_flatten(size
.bset
);
2728 size
.bset
= isl_set_simple_hull(isl_basic_set_compute_divs(size
.bset
));
2729 isl_basic_set_foreach_constraint(size
.bset
, &compute_size_in_direction
,
2731 isl_basic_set_free(size
.bset
);
2733 return isl_int_is_nonneg(bound
->size
) ? 0 : -1;
2736 /* Check if we can find a shared memory tile for the given array
2737 * based on the given accesses, and if so, put the results
2738 * in array->shared_bound.
2740 * We project the accesses on each index in turn and look for a parametric
2741 * offset such that the size is constant.
2743 static int can_tile_for_shared_memory(struct cuda_gen
*gen
,
2744 struct cuda_array_info
*array
, __isl_keep isl_map
*access
,
2745 struct cuda_array_bound
*bounds
)
2749 for (i
= 0; i
< array
->n_index
; ++i
) {
2751 isl_basic_map
*hull
;
2753 access_i
= isl_map_copy(access
);
2754 access_i
= isl_map_project_out(access_i
, isl_dim_out
, 0, i
);
2755 access_i
= isl_map_project_out(access_i
, isl_dim_out
,
2756 1, array
->n_index
- (i
+ 1));
2757 access_i
= isl_map_compute_divs(access_i
);
2758 hull
= isl_map_simple_hull(access_i
);
2759 if (compute_array_dim_size(gen
, &bounds
[i
], hull
) < 0)
2766 /* Construct a map with input the shared tile loops and the loops that
2767 * will be wrapped around the threads that relates these later loops
2768 * to the thread indices and the projects them out.
2770 static __isl_give isl_map
*compute_privatization(struct cuda_gen
*gen
)
2778 dim
= isl_union_map_get_space(gen
->shared_sched
);
2780 if (gen
->options
->wrap
)
2781 tiling
= wrap(isl_space_copy(dim
), gen
->shared_len
+ gen
->n_block
,
2782 gen
->shared_len
, gen
->n_block
, gen
->block_dim
);
2784 tiling
= tile(isl_space_copy(dim
), gen
->shared_len
+ gen
->n_block
,
2785 gen
->shared_len
, gen
->n_block
, gen
->block_dim
);
2789 par
= parametrization(dim
, gen
->shared_len
+ 2 * gen
->n_block
,
2790 gen
->tile_first
+ gen
->tile_len
+ gen
->n_grid
+ gen
->n_block
,
2793 priv
= isl_map_align_params(priv
, isl_set_get_space(par
));
2794 priv
= isl_map_intersect_range(priv
, par
);
2796 dim
= isl_map_get_space(priv
);
2797 dim
= isl_space_drop_dims(dim
, isl_dim_in
, 0, isl_space_dim(dim
, isl_dim_in
));
2798 dim
= isl_space_drop_dims(dim
, isl_dim_out
, 0, isl_space_dim(dim
, isl_dim_out
));
2799 proj
= projection(dim
, gen
->shared_len
+ 2 * gen
->n_block
,
2802 priv
= isl_map_apply_range(priv
, proj
);
2807 /* Construct a map from domain_dim to domain_dim that increments
2808 * the dimension at position "pos" and leaves all other dimensions
2811 static __isl_give isl_map
*next(__isl_take isl_space
*domain_dim
, int pos
)
2814 int len
= isl_space_dim(domain_dim
, isl_dim_set
);
2816 isl_basic_map
*next
;
2817 isl_local_space
*ls
;
2819 dim
= isl_space_map_from_set(domain_dim
);
2820 next
= isl_basic_map_universe(isl_space_copy(dim
));
2821 ls
= isl_local_space_from_space(dim
);
2823 for (i
= 0; i
< len
; ++i
) {
2826 c
= isl_equality_alloc(isl_local_space_copy(ls
));
2827 isl_constraint_set_coefficient_si(c
, isl_dim_in
, i
, 1);
2828 isl_constraint_set_coefficient_si(c
, isl_dim_out
, i
, -1);
2830 isl_constraint_set_constant_si(c
, 1);
2831 next
= isl_basic_map_add_constraint(next
, c
);
2834 isl_local_space_free(ls
);
2836 return isl_map_from_basic_map(next
);
2839 /* Check if the given access is coalesced.
2840 * That is, check whether incrementing the dimension that will get
2841 * wrapped over the last thread index results in incrementing
2842 * the last array index.
2844 * This function is only called for access relations without reuse.
2846 static int access_is_coalesced(struct cuda_gen
*gen
,
2847 __isl_keep isl_union_map
*access
)
2850 isl_map
*access_map
;
2851 isl_map
*next_thread_x
;
2852 isl_map
*next_element
;
2856 access
= isl_union_map_copy(access
);
2857 access
= isl_union_map_apply_domain(access
,
2858 isl_union_map_copy(gen
->tiled_sched
));
2859 access_map
= isl_map_from_union_map(access
);
2861 dim
= isl_map_get_space(access_map
);
2862 dim
= isl_space_domain(dim
);
2863 next_thread_x
= next(dim
, gen
->shared_len
+ gen
->n_block
- 1);
2865 dim
= isl_map_get_space(access_map
);
2866 dim
= isl_space_range(dim
);
2867 next_element
= next(dim
, isl_space_dim(dim
, isl_dim_set
) - 1);
2869 map
= isl_map_apply_domain(next_thread_x
, isl_map_copy(access_map
));
2870 map
= isl_map_apply_range(map
, access_map
);
2872 coalesced
= isl_map_is_subset(map
, next_element
);
2874 isl_map_free(next_element
);
2880 /* For the given array reference group, check whether the access is private
2881 * to the thread. That is, check that any given array element
2882 * is only accessed by a single thread.
2883 * We compute an access relation that maps the shared tile loop iterators
2884 * and the shared point loop iterators that will be wrapped over the
2885 * threads to the array elements.
2886 * We actually check that those iterators that will be wrapped
2887 * partition the array space. This check is stricter than necessary
2888 * since several iterations may be mapped onto the same thread
2889 * and then they could be allowed to access the same memory elements,
2890 * but our check does not allow this situation.
2892 * We also check that the index expression only depends on parallel
2893 * loops. That way, we can move those loops innermost and unroll them.
2894 * Again, we use a test that is stricter than necessary.
2895 * We actually check whether the index expression only depends
2896 * on the iterators that are wrapped over the threads.
2897 * These are necessarily parallel, but there may be more parallel loops.
2899 * Combining the injectivity of the first test with the single-valuedness
2900 * of the second test, we simply test for bijectivity.
2902 * If it turns out we can use registers, we compute the private memory
2903 * tile size using can_tile_for_shared_memory, after introducing a dependence
2904 * on the thread indices.
2906 * Before performing any of the above computations, we first check
2907 * if there is any reuse on the reference group. If not, we simply
2908 * return. If, moreover, the access is coalesced then we also remove
2909 * the shared memory tiling since we should just use global memory instead.
2911 static void check_private_group_access(struct cuda_gen
*gen
,
2912 struct cuda_array_ref_group
*group
)
2915 isl_union_map
*access
;
2916 int n_index
= group
->array
->n_index
;
2918 access
= group_access_relation(group
, 1, 1);
2919 if (isl_union_map_is_injective(access
)) {
2920 if (group
->shared_bound
&& access_is_coalesced(gen
, access
)) {
2921 free_bound_list(group
->shared_bound
, n_index
);
2922 group
->shared_bound
= NULL
;
2924 isl_union_map_free(access
);
2927 access
= isl_union_map_apply_domain(access
,
2928 isl_union_map_copy(gen
->shared_sched
));
2930 acc
= isl_map_from_union_map(access
);
2932 if (!isl_map_is_bijective(acc
)) {
2937 group
->private_bound
= create_bound_list(gen
->ctx
, n_index
);
2938 acc
= isl_map_align_params(acc
, isl_map_get_space(gen
->privatization
));
2939 acc
= isl_map_apply_domain(acc
, isl_map_copy(gen
->privatization
));
2940 if (!can_tile_for_shared_memory(gen
, group
->array
, acc
,
2941 group
->private_bound
)) {
2942 free_bound_list(group
->private_bound
, n_index
);
2943 group
->private_bound
= NULL
;
2949 /* Look for the last shared tile loop that affects the offset of the
2950 * shared or private tile and store the result in array->last_shared.
2952 static void set_last_shared(struct cuda_gen
*gen
,
2953 struct cuda_array_ref_group
*group
)
2956 struct cuda_array_bound
*bounds
;
2957 unsigned first_shared
= gen
->first_shared
;
2958 int n_index
= group
->array
->n_index
;
2960 bounds
= group
->private_bound
;
2962 bounds
= group
->shared_bound
;
2966 for (j
= gen
->shared_len
- 1; j
>= 0; --j
) {
2967 for (i
= 0; i
< n_index
; ++i
) {
2972 if (isl_aff_involves_dims(lb
, isl_dim_param
,
2973 first_shared
+ j
, 1))
2976 shift
= bounds
[i
].shift
;
2979 if (isl_aff_involves_dims(shift
, isl_dim_param
,
2980 first_shared
+ j
, 1))
2986 group
->array
->last_shared
= j
;
2989 /* Compute the sizes of all private arrays for the current kernel,
2990 * as well as the offsets of the private pieces in the original arrays.
2991 * If we cannot or don't want to privatize a given array group,
2992 * we use the shared memory tile sizes computed in
2993 * compute_group_shared_bound instead.
2995 * If a given Array only has a single reference group and if we have
2996 * been able to find a privated or shared tile,
2997 * we also look for the last shared tile loop that affects the offset
2998 * (and therefore the array tile) and store the result in array->last_shared.
3000 * A privatized copy of all access relations from reference groups that
3001 * are mapped to private memory is stored in gen->privatization.
3003 static void compute_private_size(struct cuda_gen
*gen
)
3006 isl_union_map
*private;
3008 if (!gen
->options
->use_private_memory
)
3011 private = isl_union_map_empty(isl_union_map_get_space(gen
->shared_sched
));
3013 for (i
= 0; i
< gen
->n_array
; ++i
) {
3014 struct cuda_array_info
*array
= &gen
->array
[i
];
3016 for (j
= 0; j
< array
->n_group
; ++j
) {
3017 check_private_group_access(gen
, array
->groups
[j
]);
3019 if (!array
->groups
[j
]->private_bound
)
3022 private = isl_union_map_union(private,
3023 group_access_relation(array
->groups
[j
], 1, 1));
3026 array
->last_shared
= gen
->shared_len
- 1;
3027 array
->print_shared_level
= -1;
3029 if (array
->n_group
!= 1)
3031 set_last_shared(gen
, array
->groups
[0]);
3034 if (isl_union_map_is_empty(private))
3035 isl_union_map_free(private);
3037 isl_union_map
*priv
;
3039 private = isl_union_map_apply_domain(private,
3040 isl_union_map_copy(gen
->shared_sched
));
3041 priv
= isl_union_map_from_map(isl_map_copy(gen
->privatization
));
3042 private = isl_union_map_apply_domain(private, priv
);
3043 gen
->private_access
= private;
3047 /* Fill up the groups array with singleton groups, i.e., one group
3048 * per reference, initializing the array, access, write and refs fields.
3049 * In particular the access field is initialized to the scheduled
3050 * access relation of the array reference.
3052 * Return the number of elements initialized, i.e., the number of
3053 * active references in the current kernel.
3055 static int populate_array_references(struct cuda_gen
*gen
,
3056 struct cuda_array_info
*array
, __isl_keep isl_union_map
*sched
,
3057 struct cuda_array_ref_group
**groups
)
3061 isl_ctx
*ctx
= isl_union_map_get_ctx(sched
);
3064 for (i
= 0; i
< array
->n_ref
; ++i
) {
3065 isl_union_map
*umap
;
3067 struct cuda_array_ref_group
*group
;
3068 struct cuda_stmt_access
*access
= array
->refs
[i
];
3070 map
= isl_map_copy(access
->access
);
3071 umap
= isl_union_map_from_map(map
);
3072 umap
= isl_union_map_apply_domain(umap
,
3073 isl_union_map_copy(sched
));
3075 if (isl_union_map_is_empty(umap
)) {
3076 isl_union_map_free(umap
);
3080 map
= isl_map_from_union_map(umap
);
3082 group
= isl_calloc_type(ctx
, struct cuda_array_ref_group
);
3084 group
->array
= array
;
3085 group
->access
= map
;
3086 group
->write
= access
->write
;
3087 group
->refs
= &array
->refs
[i
];
3089 groups
[n
++] = group
;
3095 static void free_array_ref_group(struct cuda_array_ref_group
*group
,
3100 free_bound_list(group
->shared_bound
, n_index
);
3101 free_bound_list(group
->private_bound
, n_index
);
3102 isl_map_free(group
->access
);
3107 /* If two groups have overlapping access relations and if one of them
3108 * involves a write, then merge the two groups into one.
3110 * We keep track of the grouping in "leader". leader[j] points to
3111 * an earlier group array element that belongs to the same group,
3112 * or the array element j itself if this element is the first in the group.
3114 * Return the number of group leaders.
3116 static int group_overlapping_writes(int n
,
3117 struct cuda_array_ref_group
**groups
, int *leader
)
3122 for (i
= 0; i
< n
; ++i
) {
3124 groups
[l
]->n_ref
= 1;
3125 for (j
= i
- 1; j
>= 0; --j
) {
3131 if (!groups
[l
]->write
&& !groups
[j
]->write
)
3134 map
= isl_map_intersect(isl_map_copy(groups
[l
]->access
),
3135 isl_map_copy(groups
[j
]->access
));
3136 empty
= isl_map_is_empty(map
);
3142 groups
[j
]->access
= isl_map_union(groups
[j
]->access
,
3144 groups
[j
]->write
= 1;
3145 groups
[l
]->access
= NULL
;
3146 groups
[j
]->n_ref
+= groups
[l
]->n_ref
;
3156 /* Compute the size of the shared array corresponding to the given array
3157 * array refrence group, based on the accesses from the current kernel,
3158 * as well as the offset of the shared piece in the original array.
3160 static void compute_group_shared_bound(struct cuda_gen
*gen
,
3161 struct cuda_array_info
*array
, struct cuda_array_ref_group
*group
)
3163 isl_ctx
*ctx
= isl_space_get_ctx(array
->dim
);
3165 if (!gen
->options
->use_shared_memory
)
3168 group
->shared_bound
= create_bound_list(ctx
, array
->n_index
);
3169 if (!can_tile_for_shared_memory(gen
, array
, group
->access
,
3170 group
->shared_bound
)) {
3171 free_bound_list(group
->shared_bound
, array
->n_index
);
3172 group
->shared_bound
= NULL
;
3176 /* Given an initial grouping of array references and shared memory tiles
3177 * for each group that allows for a shared memory tile, merge two groups
3178 * if both have a shared memory tile and if the merged group also has
3179 * a shared memory tile.
3181 * Return the number of group leaders after merging.
3183 static int group_common_shared_memory_tile(struct cuda_gen
*gen
,
3184 struct cuda_array_info
*array
, int n
,
3185 struct cuda_array_ref_group
**groups
, int *leader
, int n_group
)
3188 isl_ctx
*ctx
= isl_space_get_ctx(array
->dim
);
3190 for (i
= 0; n_group
> 1 && i
< n
; ++i
) {
3194 if (!groups
[i
]->shared_bound
)
3196 for (j
= i
- 1; j
>= 0; --j
) {
3199 struct cuda_array_bound
*shared_bound
;
3203 if (!groups
[j
]->shared_bound
)
3206 map
= isl_map_intersect(isl_map_copy(groups
[l
]->access
),
3207 isl_map_copy(groups
[j
]->access
));
3208 empty
= isl_map_is_empty(map
);
3214 map
= isl_map_union(isl_map_copy(groups
[l
]->access
),
3215 isl_map_copy(groups
[j
]->access
));
3216 shared_bound
= create_bound_list(ctx
, array
->n_index
);
3217 if (!can_tile_for_shared_memory(gen
, array
, map
,
3220 free_bound_list(shared_bound
, array
->n_index
);
3224 free_bound_list(groups
[j
]->shared_bound
,
3226 groups
[j
]->shared_bound
= shared_bound
;
3227 isl_map_free(groups
[j
]->access
);
3228 groups
[j
]->access
= map
;
3229 groups
[j
]->n_ref
+= groups
[l
]->n_ref
;
3238 /* Extract an array of array reference groups from the array of references
3239 * and the grouping information in "leader".
3241 * Store the results in array->n_group and array->groups.
3243 static void extract_array_groups(isl_ctx
*ctx
, struct cuda_array_info
*array
,
3244 int n
, struct cuda_array_ref_group
**groups
, int *leader
, int n_group
)
3248 for (i
= 2; i
< n
; ++i
)
3249 leader
[i
] = leader
[leader
[i
]];
3251 array
->n_group
= n_group
;
3252 array
->groups
= isl_alloc_array(ctx
, struct cuda_array_ref_group
*,
3254 assert(array
->groups
);
3257 for (i
= 0; i
< n
; ++i
) {
3259 struct cuda_stmt_access
**refs
;
3261 if (leader
[i
] != i
) {
3262 groups
[i
]->refs
= NULL
;
3263 free_array_ref_group(groups
[i
], array
->n_index
);
3267 refs
= isl_alloc_array(ctx
, struct cuda_stmt_access
*,
3271 for (k
= i
; k
< n
; ++k
)
3272 if (leader
[k
] == i
) {
3273 refs
[l
++] = *groups
[k
]->refs
;
3274 (*groups
[k
]->refs
)->group
= j
;
3277 groups
[i
]->refs
= refs
;
3279 array
->groups
[j
++] = groups
[i
];
3283 /* Group array references that should be considered together when
3284 * deciding whether to access them from private, shared or global memory.
3286 * In particular, if two array references overlap and if one of them
3287 * is a write, then the two references are grouped together.
3288 * Furthermore, if two groups admit a shared memory tile and if the
3289 * combination of the two also admits a shared memory tile, we merge
3292 * During the construction the group->refs field points to a single
3293 * array reference inside the array of array references, while
3294 * group->n_ref contains the number of element in leader that
3295 * (directly or indirectly) point to this group, provided the group
3298 static void group_array_references(struct cuda_gen
*gen
,
3299 struct cuda_array_info
*array
, __isl_keep isl_union_map
*sched
)
3303 isl_ctx
*ctx
= isl_union_map_get_ctx(sched
);
3304 struct cuda_array_ref_group
**groups
;
3307 groups
= isl_calloc_array(ctx
, struct cuda_array_ref_group
*,
3311 n
= populate_array_references(gen
, array
, sched
, groups
);
3313 leader
= isl_alloc_array(ctx
, int, n
);
3316 n_group
= group_overlapping_writes(n
, groups
, leader
);
3318 for (i
= 0; i
< n
; ++i
)
3320 compute_group_shared_bound(gen
, array
, groups
[i
]);
3322 n_group
= group_common_shared_memory_tile(gen
, array
, n
, groups
,
3325 extract_array_groups(ctx
, array
, n
, groups
, leader
, n_group
);
3331 /* Take tiled_sched, project it onto the shared tile loops and
3332 * the loops that will be wrapped over the threads,
3333 * parametrize the shared tile loops and store the result in gen->shared_sched.
3334 * The position of the first of these parameters is stored in gen->first_shared.
3335 * Also compute a projection that projects out the loops that will be
3336 * wrapped over the threads and store this projection in gen->shared_proj.
3338 static void compute_shared_sched(struct cuda_gen
*gen
)
3343 isl_union_map
*sched
;
3345 sched
= isl_union_map_copy(gen
->tiled_sched
);
3347 dim
= isl_union_map_get_space(sched
);
3348 gen
->first_shared
= isl_space_dim(dim
, isl_dim_param
);
3349 proj
= projection(dim
, gen
->tiled_len
, gen
->shared_len
+ gen
->n_block
);
3350 sched
= isl_union_map_apply_range(sched
, isl_union_map_from_map(proj
));
3352 dim
= isl_union_map_get_space(sched
);
3353 par
= parametrization(dim
, gen
->shared_len
+ gen
->n_block
,
3354 0, gen
->shared_len
, "g");
3355 sched
= isl_union_map_intersect_range(sched
,
3356 isl_union_set_from_set(par
));
3358 dim
= isl_union_map_get_space(sched
);
3359 proj
= projection(dim
, gen
->shared_len
+ gen
->n_block
, gen
->shared_len
);
3361 gen
->shared_sched
= sched
;
3362 gen
->shared_proj
= isl_union_map_from_map(proj
);
3365 /* Group references of all arrays in the program.
3367 static void group_references(struct cuda_gen
*gen
)
3370 isl_union_map
*sched
;
3372 sched
= isl_union_map_apply_range(isl_union_map_copy(gen
->shared_sched
),
3373 isl_union_map_copy(gen
->shared_proj
));
3375 for (i
= 0; i
< gen
->n_array
; ++i
)
3376 group_array_references(gen
, &gen
->array
[i
], sched
);
3378 isl_union_map_free(sched
);
3381 /* Free all array information that is local to the current kernel.
3383 static void free_local_array_info(struct cuda_gen
*gen
)
3387 for (i
= 0; i
< gen
->n_array
; ++i
) {
3388 struct cuda_array_info
*array
= &gen
->array
[i
];
3390 for (j
= 0; j
< array
->n_group
; ++j
)
3391 free_array_ref_group(array
->groups
[j
], array
->n_index
);
3392 free(array
->groups
);
3394 if (array
->n_group
== 0)
3396 for (j
= 0; j
< gen
->array
[i
].n_index
; ++j
) {
3397 isl_pw_aff_free(gen
->array
[i
].local_bound
[j
]);
3398 gen
->array
[i
].local_bound
[j
] = NULL
;
3403 static void print_iterator_list(FILE *out
, int len
, const char *prefix
,
3409 for (i
= 0; i
< len
; ++i
) {
3413 fprintf(out
, "(%s%d)", prefix
, i
);
3415 fprintf(out
, "%s%d", prefix
, i
);
3420 /* Print an access to the element in the global memory copy of the
3421 * given array that corresponds to element [a0][a1]... of the original array.
3422 * The copy in global memory has been linearized, so we need to take
3423 * the array size into account.
3425 static void print_global_index(isl_ctx
*ctx
, FILE *out
,
3426 struct cuda_array_info
*array
)
3431 if (cuda_array_is_scalar(array
)) {
3432 fprintf(out
, "*%s", array
->name
);
3436 fprintf(out
, "%s[", array
->name
);
3437 for (i
= 0; i
+ 1 < array
->n_index
; ++i
)
3439 for (i
= 0; i
< array
->n_index
; ++i
) {
3441 prn
= isl_printer_to_file(ctx
, out
);
3442 prn
= isl_printer_set_output_format(prn
, ISL_FORMAT_C
);
3443 prn
= isl_printer_print_str(prn
, ") * (");
3444 prn
= isl_printer_print_pw_aff(prn
,
3445 array
->local_bound
[i
]);
3446 prn
= isl_printer_print_str(prn
, ") + ");
3447 isl_printer_free(prn
);
3449 fprintf(out
, "a%d", i
);
3454 /* Print an access to the element in the shared memory copy of the
3455 * given array that corresponds to element [a0][a1]... of the original array.
3456 * Since the array in shared memory is just a shifted copy of part
3457 * of the original array, we simply need to subtract the lower bound,
3458 * which was computed in can_tile_for_shared_memory.
3459 * If any of the indices is strided, then we first add
3460 * shared_bound[i].shift and divide by shared_bound[i].stride.
3462 static void print_local_index(FILE *out
, struct cuda_array_ref_group
*group
)
3467 struct cuda_array_bound
*bounds
= group
->shared_bound
;
3469 ctx
= isl_space_get_ctx(group
->array
->dim
);
3470 print_array_name(out
, group
);
3471 for (i
= 0; i
< group
->array
->n_index
; ++i
) {
3472 fprintf(out
, "[(a%d", i
);
3473 if (bounds
[i
].shift
) {
3474 fprintf(out
, " + (");
3475 prn
= isl_printer_to_file(ctx
, out
);
3476 prn
= isl_printer_set_output_format(prn
, ISL_FORMAT_C
);
3477 prn
= isl_printer_print_aff(prn
, bounds
[i
].shift
);
3478 prn
= isl_printer_print_str(prn
, "))/");
3479 prn
= isl_printer_print_isl_int(prn
,
3481 isl_printer_free(prn
);
3484 fprintf(out
, " - (");
3485 prn
= isl_printer_to_file(ctx
, out
);
3486 prn
= isl_printer_set_output_format(prn
, ISL_FORMAT_C
);
3487 prn
= isl_printer_print_aff(prn
, bounds
[i
].lb
);
3488 isl_printer_free(prn
);
3493 /* Print '#define's for copying data from global memory to shared
3494 * memory and back for the given array.
3496 static void print_array_copy_defines(struct cuda_gen
*gen
,
3497 struct cuda_array_ref_group
*group
)
3500 const char *type
[] = { "read", "write" };
3501 struct cuda_array_info
*array
= group
->array
;
3502 int n_index
= array
->n_index
;
3504 for (i
= 0; i
< 2; ++i
) {
3505 fprintf(gen
->cuda
.kernel_c
, "#define %s_", type
[i
]);
3506 print_array_name(gen
->cuda
.kernel_c
, group
);
3507 print_iterator_list(gen
->cuda
.kernel_c
, n_index
, "a", 0);
3508 fprintf(gen
->cuda
.kernel_c
, " %s_", type
[i
]);
3509 print_array_name(gen
->cuda
.kernel_c
, group
);
3510 fprintf(gen
->cuda
.kernel_c
, "_");
3511 print_iterator_list(gen
->cuda
.kernel_c
, n_index
, "a", 1);
3512 fprintf(gen
->cuda
.kernel_c
, "\n");
3514 fprintf(gen
->cuda
.kernel_c
, "#define %s_", type
[i
]);
3515 print_array_name(gen
->cuda
.kernel_c
, group
);
3516 fprintf(gen
->cuda
.kernel_c
, "_");
3517 print_iterator_list(gen
->cuda
.kernel_c
, n_index
, "a", 0);
3519 fprintf(gen
->cuda
.kernel_c
, " ");
3520 print_global_index(gen
->ctx
, gen
->cuda
.kernel_c
, array
);
3521 fprintf(gen
->cuda
.kernel_c
, " = ");
3522 print_local_index(gen
->cuda
.kernel_c
, group
);
3524 fprintf(gen
->cuda
.kernel_c
, " ");
3525 print_local_index(gen
->cuda
.kernel_c
, group
);
3526 fprintf(gen
->cuda
.kernel_c
, " = ");
3527 print_global_index(gen
->ctx
, gen
->cuda
.kernel_c
, array
);
3529 fprintf(gen
->cuda
.kernel_c
, "\n");
3533 static void print_copy_defines(struct cuda_gen
*gen
)
3537 for (i
= 0; i
< gen
->n_array
; ++i
) {
3538 struct cuda_array_info
*array
= &gen
->array
[i
];
3540 for (j
= 0; j
< array
->n_group
; ++j
) {
3541 if (array
->groups
[j
]->private_bound
)
3543 if (!array
->groups
[j
]->shared_bound
)
3545 print_array_copy_defines(gen
, array
->groups
[j
]);
3550 /* The sizes of the arrays on the host that have been computed by
3551 * extract_array_info may depend on the parameters. Use the extra
3552 * constraints on the parameters that are valid at "host_domain"
3553 * to simplify these expressions.
3555 static void localize_bounds(struct cuda_gen
*gen
,
3556 __isl_keep isl_set
*host_domain
)
3561 context
= isl_set_copy(host_domain
);
3562 context
= isl_set_params(host_domain
);
3564 for (i
= 0; i
< gen
->n_array
; ++i
) {
3565 struct cuda_array_info
*array
= &gen
->array
[i
];
3567 if (array
->n_group
== 0)
3570 for (j
= 0; j
< array
->n_index
; ++j
) {
3573 pwaff
= isl_pw_aff_copy(array
->bound
[j
]);
3574 pwaff
= isl_pw_aff_gist(pwaff
, isl_set_copy(context
));
3575 array
->local_bound
[j
] = pwaff
;
3578 isl_set_free(context
);
3581 /* Set gen->tile_len and gen->n_parallel to those of the first statement
3582 * in the statement list u.
3583 * Because of the way the schedule is constructed, the other statements
3584 * in the list, if any, should have the same values for these properties.
3586 static void set_tile_len(struct cuda_gen
*gen
, struct clast_user_stmt
*u
)
3589 struct cuda_stmt
*stmt
;
3591 nr
= atoi(u
->statement
->name
+ 2);
3592 stmt
= &gen
->stmts
[nr
];
3594 gen
->tile_len
= stmt
->tile_len
;
3595 gen
->n_parallel
= stmt
->n_parallel
;
3598 /* This function is called for each leaf in the clast of the host code.
3599 * We first specialize the schedule to the site of the leaf, compute
3600 * the size of shared memory and then print the body of host code
3601 * and the associated kernel (through a call to print_kernel_body).
3603 static void print_host_user(struct gpucode_info
*code
,
3604 struct clast_user_stmt
*u
)
3606 struct cuda_gen
*gen
= code
->user
;
3609 isl_set
*host_domain
;
3610 isl_union_map
*access
;
3611 isl_union_map
*local_sched
;
3612 isl_union_set
*arrays
;
3614 set_tile_len(gen
, u
);
3617 host_domain
= extract_entire_host_domain(u
);
3619 local_sched
= isl_union_map_intersect_range(
3620 isl_union_map_copy(gen
->sched
),
3621 isl_union_set_from_set(extend(isl_set_copy(host_domain
),
3622 gen
->untiled_len
)));
3623 access
= isl_union_map_union(isl_union_map_copy(gen
->read
),
3624 isl_union_map_copy(gen
->write
));
3625 access
= isl_union_map_apply_domain(access
,
3626 isl_union_map_copy(local_sched
));
3627 arrays
= isl_union_map_range(access
);
3629 print_indent(code
->dst
, code
->indent
);
3630 fprintf(code
->dst
, "dim3 k%d_dimBlock(", gen
->kernel_id
);
3631 print_reverse_list(code
->dst
, gen
->n_block
, gen
->block_dim
);
3632 fprintf(code
->dst
, ");\n");
3634 print_indent(code
->dst
, code
->indent
);
3635 fprintf(code
->dst
, "dim3 k%d_dimGrid(", gen
->kernel_id
);
3636 print_reverse_list(code
->dst
, gen
->n_grid
, gen
->grid_dim
);
3637 fprintf(code
->dst
, ");\n");
3639 gen
->tiled_sched
= tile_schedule(gen
, local_sched
);
3640 gen
->tiled_sched
= parametrize_tiled_schedule(gen
, gen
->tiled_sched
);
3641 gen
->tiled_sched
= scale_tile_loops(gen
, gen
->tiled_sched
);
3643 gen
->local_sched
= isl_union_map_copy(gen
->tiled_sched
);
3645 dim
= isl_union_map_get_space(gen
->local_sched
);
3646 par
= parametrization(dim
, gen
->tiled_len
, 0, gen
->shared_len
, "g");
3647 gen
->local_sched
= isl_union_map_intersect_range(gen
->local_sched
,
3648 isl_union_set_from_set(par
));
3650 gen
->local_sched
= thread_tile_schedule(gen
, gen
->local_sched
);
3651 gen
->local_sched
= scale_thread_tile_loops(gen
, gen
->local_sched
);
3653 gen
->private_access
= NULL
;
3654 compute_shared_sched(gen
);
3655 gen
->privatization
= compute_privatization(gen
);
3656 group_references(gen
);
3657 compute_private_size(gen
);
3658 localize_bounds(gen
, host_domain
);
3660 gen
->local_sched
= interchange_for_unroll(gen
, gen
->local_sched
);
3662 print_copy_defines(gen
);
3663 print_kernel_launch(gen
, arrays
);
3665 fprintf(gen
->cuda
.kernel_c
, "{\n");
3667 print_kernel_body(gen
, host_domain
, gen
->tiled_sched
);
3669 fprintf(gen
->cuda
.kernel_c
, "}\n");
3671 free_local_array_info(gen
);
3672 isl_map_free(gen
->privatization
);
3673 isl_union_map_free(gen
->private_access
);
3674 isl_union_map_free(gen
->local_sched
);
3675 isl_union_map_free(gen
->tiled_sched
);
3676 isl_union_map_free(gen
->shared_sched
);
3677 isl_union_map_free(gen
->shared_proj
);
3678 isl_union_set_free(arrays
);
3679 isl_set_free(host_domain
);
3681 free(gen
->tile_size
);
3685 /* Use CLooG to generate code for the outer gen->tile_first loops
3686 * of the global schedule in gen->sched.
3687 * The pretty printing of this code is handled by gpu_print_host_stmt,
3688 * which calls print_host_user for each kernel invocation location.
3690 static void print_cloog_host_code(struct cuda_gen
*gen
)
3694 isl_union_map
*sched
;
3695 CloogOptions
*options
;
3696 CloogDomain
*cloog_context
;
3697 CloogUnionDomain
*ud
;
3699 struct clast_stmt
*stmt
;
3702 options
= cloog_options_malloc(gen
->state
);
3703 options
->language
= CLOOG_LANGUAGE_C
;
3705 options
->strides
= 1;
3706 options
->stop
= gen
->tile_first
;
3707 options
->f
= gen
->untiled_len
;
3708 options
->l
= gen
->untiled_len
;
3709 options
->save_domains
= 1;
3710 options
->noscalars
= 1;
3712 sched
= isl_union_map_copy(gen
->sched
);
3713 ud
= cloog_union_domain_from_isl_union_map(sched
);
3714 for (i
= 0; i
< options
->stop
; ++i
) {
3715 snprintf(name
, sizeof(name
), "h%d", i
);
3716 ud
= cloog_union_domain_set_name(ud
, CLOOG_SCAT
, i
, name
);
3718 context
= isl_set_copy(gen
->context
);
3719 cloog_context
= cloog_domain_from_isl_set(context
);
3720 input
= cloog_input_alloc(cloog_context
, ud
);
3722 stmt
= cloog_clast_create_from_input(input
, options
);
3724 gen
->code
.indent
= 0;
3725 gen
->code
.dst
= gen
->cuda
.host_c
;
3726 gen
->code
.print_user_stmt
= NULL
;
3727 gen
->code
.print_user_stmt_list
= &print_host_user
;
3728 gen
->code
.print_for_head
= NULL
;
3729 gen
->code
.print_for_foot
= NULL
;
3730 gen
->code
.user
= gen
;
3731 gpu_print_host_stmt(&gen
->code
, stmt
);
3733 cloog_clast_free(stmt
);
3734 cloog_options_free(options
);
3735 fprintf(gen
->cuda
.host_c
, "\n");
3738 void print_cuda_macros(struct cuda_gen
*gen
)
3740 const char *macros
=
3741 "#define cudaCheckReturn(ret) assert((ret) == cudaSuccess)\n"
3742 "#define cudaCheckKernel()"
3743 " assert(cudaGetLastError() == cudaSuccess)\n\n";
3744 fputs(macros
, gen
->cuda
.host_c
);
3747 void print_host_code(struct cuda_gen
*gen
)
3749 fprintf(gen
->cuda
.host_c
, "{\n");
3750 print_cloog_macros(gen
->cuda
.host_c
);
3751 print_cloog_macros(gen
->cuda
.kernel_c
);
3753 print_cuda_macros(gen
);
3755 declare_device_arrays(gen
);
3757 allocate_device_arrays(gen
);
3758 copy_arrays_to_device(gen
);
3761 print_cloog_host_code(gen
);
3763 copy_arrays_from_device(gen
);
3764 free_device_arrays(gen
);
3766 fprintf(gen
->cuda
.host_c
, "}\n");
3769 __isl_give isl_set
*add_context_from_str(__isl_take isl_set
*set
,
3778 ctx
= isl_set_get_ctx(set
);
3779 context
= isl_set_read_from_str(ctx
, str
);
3780 context
= isl_set_align_params(context
, isl_set_get_space(set
));
3781 set
= isl_set_intersect(set
, context
);
3786 /* Return the union of all iteration domains of the gen->stmts[i].
3788 static __isl_give isl_union_set
*extract_domain(struct cuda_gen
*gen
)
3791 isl_union_set
*domain
;
3793 domain
= isl_union_set_empty(isl_set_get_space(gen
->context
));
3794 for (i
= 0; i
< gen
->n_stmts
; ++i
) {
3797 domain_i
= isl_set_copy(gen
->stmts
[i
].domain
);
3798 domain
= isl_union_set_union(domain
,
3799 isl_union_set_from_set(domain_i
));
3805 /* Information about the outermost tilable bands in the forest of bands.
3807 * tile_len and n_parallel are only sets on band_info structures
3808 * that correspond to outermost bands. For other bands (in particular,
3809 * ancestors of the outermost bands), n_parallal is set to 0.
3811 * prefix is the (padded) schedule leading up to the outermost tilable bands.
3813 * tile_first is the number of schedule dimensions in prefix.
3815 * suffix is the schedule of the outermost tilable bands and their descendants.
3818 struct cuda_gen
*gen
;
3822 isl_union_map
*prefix
;
3823 isl_union_map
*suffix
;
3826 /* Set tile_len and n_parallel of the statement to that of
3827 * their outermost band, recorded in the band_info.
3829 static int set_stmt_tile_len(__isl_take isl_map
*map
, void *user
)
3831 struct band_info
*info
= user
;
3833 struct cuda_stmt
*stmt
;
3835 nr
= atoi(isl_map_get_tuple_name(map
, isl_dim_in
) + 2);
3836 stmt
= &info
->gen
->stmts
[nr
];
3838 stmt
->tile_len
= info
->tile_len
;
3839 stmt
->n_parallel
= info
->n_parallel
;
3846 static void list_select_outer_band(struct cuda_gen
*gen
,
3847 __isl_take isl_band_list
*list
, int pos
, struct band_info
*list_info
);
3849 /* Check if this band has any parallel loops. If so, take it as
3850 * the outermost tilable band. If not, continue looking for the
3851 * outermost tilable band in the children of the current band.
3853 static void band_select_outer_band(struct cuda_gen
*gen
,
3854 __isl_take isl_band
*band
, int pos
, struct band_info
*info
)
3856 int n
= isl_band_n_member(band
);
3859 for (n_parallel
= 0; n_parallel
< n
; ++n_parallel
)
3860 if (!isl_band_member_is_zero_distance(band
, n_parallel
))
3863 info
->n_parallel
= n_parallel
;
3866 info
->tile_first
= pos
;
3868 info
->prefix
= isl_band_get_prefix_schedule(band
);
3869 info
->suffix
= isl_union_map_flat_range_product(
3870 isl_band_get_partial_schedule(band
),
3871 isl_band_get_suffix_schedule(band
));
3872 isl_union_map_foreach_map(info
->prefix
,
3873 &set_stmt_tile_len
, info
);
3875 isl_band_list
*children
;
3876 if (!isl_band_has_children(band
))
3877 isl_die(isl_band_get_ctx(band
), isl_error_unknown
,
3878 "unable to detect any parallelism", abort());
3879 children
= isl_band_get_children(band
);
3880 list_select_outer_band(gen
, children
, pos
+ n
, info
);
3883 isl_band_free(band
);
3886 /* Comparison function that returns a non-zero value for band_infos
3887 * with different tile_len fields or different n_parallel fields.
3889 static int cmp_band(const void *p1
, const void *p2
)
3891 const struct band_info
*info1
= p1
;
3892 const struct band_info
*info2
= p2
;
3894 if (info1
->tile_len
!= info2
->tile_len
)
3895 return info1
->tile_len
- info2
->tile_len
;
3897 return info1
->n_parallel
- info2
->n_parallel
;
3900 /* Extend "umap" with coordinates with fixed value "val"
3901 * to a total length of "dst_len", assuming the original dimension is "src_len".
3903 static __isl_give isl_union_map
*extend_range(__isl_take isl_union_map
*umap
,
3904 int src_len
, int dst_len
, int val
)
3910 dim
= isl_union_map_get_space(umap
);
3911 map
= isl_map_reverse(projection(dim
, dst_len
, src_len
));
3912 for (i
= src_len
; i
< dst_len
; ++i
)
3913 map
= isl_map_fix_si(map
, isl_dim_out
, i
, val
);
3915 umap
= isl_union_map_apply_range(umap
, isl_union_map_from_map(map
));
3920 /* Group bands with the same values for tile_len and n_parallel.
3921 * The prefix schedule is then extended with a fixed coordinate that
3922 * is different for each such group.
3923 * Note that the actual values for this coordinate are not important.
3924 * The bands have already been effectively separated at a higher level
3925 * or they are independent and may be executed in parallel.
3926 * The list of band_info has been sorted before this functions is called.
3928 static void separate_bands(struct band_info
*info
, int n
)
3933 for (i
= 0; i
< n
; ++i
) {
3934 int l
= info
[i
].tile_first
;
3937 (info
[i
].tile_len
!= info
[i
- 1].tile_len
||
3938 info
[i
].n_parallel
!= info
[i
- 1].n_parallel
))
3941 info
[i
].prefix
= extend_range(info
[i
].prefix
,
3943 info
[i
].tile_first
= l
+ 1;
3947 /* Select the outermost bands in the elements of the list, align
3948 * their prefix schedules, separate bands with different values
3949 * for tile_len and/or n_parallel and then combine the resulting
3950 * prefix and suffix schedules into a single pair of prefix and
3951 * suffix schedules for the entire list.
3953 static void list_select_outer_band(struct cuda_gen
*gen
,
3954 __isl_take isl_band_list
*list
, int pos
, struct band_info
*list_info
)
3958 int n
= isl_band_list_n_band(list
);
3959 isl_ctx
*ctx
= isl_band_list_get_ctx(list
);
3960 struct band_info
*info
;
3962 isl_union_map
*prefix
;
3963 isl_union_map
*suffix
;
3966 info
= isl_calloc_array(ctx
, struct band_info
, n
);
3970 for (i
= 0; i
< n
; ++i
) {
3971 band
= isl_band_list_get_band(list
, i
);
3972 band_select_outer_band(gen
, band
, pos
, &info
[i
]);
3973 if (info
[i
].tile_first
> max_tile_first
)
3974 max_tile_first
= info
[i
].tile_first
;
3977 for (i
= 0; i
< n
; ++i
) {
3978 if (info
[i
].tile_first
== max_tile_first
)
3980 info
[i
].prefix
= extend_range(info
[i
].prefix
,
3981 info
[i
].tile_first
, max_tile_first
, 0);
3982 info
[i
].tile_first
= max_tile_first
;
3985 qsort(info
, n
, sizeof(struct band_info
), &cmp_band
);
3987 for (i
= 0; i
< n
- 1; ++i
)
3988 if (info
[i
].tile_len
!= info
[i
+ 1].tile_len
||
3989 info
[i
].n_parallel
!= info
[i
+ 1].n_parallel
)
3993 separate_bands(info
, n
);
3995 prefix
= info
[0].prefix
;
3996 suffix
= info
[0].suffix
;
3998 for (i
= 1; i
< n
; ++i
) {
3999 prefix
= isl_union_map_union(prefix
, info
[i
].prefix
);
4000 suffix
= isl_union_map_union(suffix
, info
[i
].suffix
);
4003 list_info
->tile_first
= info
[0].tile_first
;
4004 list_info
->tile_len
= -1;
4005 list_info
->prefix
= prefix
;
4006 list_info
->suffix
= suffix
;
4008 isl_band_list_free(list
);
4012 /* Set max_out to the maximal number of output dimensions over
4015 static int update_max_out(__isl_take isl_map
*map
, void *user
)
4017 int *max_out
= user
;
4018 int n_out
= isl_map_dim(map
, isl_dim_out
);
4020 if (n_out
> *max_out
)
4027 struct align_range_data
{
4032 /* Extend the dimension of the range of the given map to data->max_out and
4033 * then add the result to data->res.
4035 static int map_align_range(__isl_take isl_map
*map
, void *user
)
4037 struct align_range_data
*data
= user
;
4041 int n_out
= isl_map_dim(map
, isl_dim_out
);
4043 dim
= isl_union_map_get_space(data
->res
);
4044 proj
= isl_map_reverse(projection(dim
, data
->max_out
, n_out
));
4045 for (i
= n_out
; i
< data
->max_out
; ++i
)
4046 proj
= isl_map_fix_si(proj
, isl_dim_out
, i
, 0);
4048 map
= isl_map_apply_range(map
, proj
);
4050 data
->res
= isl_union_map_add_map(data
->res
, map
);
4055 /* Extend the ranges of the maps in the union map such they all have
4056 * the same dimension.
4058 static __isl_give isl_union_map
*align_range(__isl_take isl_union_map
*umap
)
4060 struct align_range_data data
;
4063 isl_union_map_foreach_map(umap
, &update_max_out
, &data
.max_out
);
4065 data
.res
= isl_union_map_empty(isl_union_map_get_space(umap
));
4066 isl_union_map_foreach_map(umap
, &map_align_range
, &data
);
4068 isl_union_map_free(umap
);
4072 /* Select the outermost tilable band that (by construction)
4073 * has at least one parallel loop.
4074 * The starting position of the aligned band is stored in the pair
4076 * The sizes and number of parallel loops may be different in different
4077 * parts of the band forest and are therefore stored in the cuda_stmts.
4079 * Return the complete schedule, with the tilable bands aligned
4080 * at gen->tile_first and padded with zero, if needed.
4082 static __isl_give isl_union_map
*select_outer_tilable_band(struct cuda_gen
*gen
,
4083 __isl_keep isl_schedule
*schedule
)
4085 isl_band_list
*list
;
4086 struct band_info info
;
4088 gen
->n_parallel
= 0;
4091 list
= isl_schedule_get_band_forest(schedule
);
4093 list_select_outer_band(gen
, list
, 0, &info
);
4095 gen
->tile_first
= info
.tile_first
;
4096 info
.suffix
= align_range(info
.suffix
);
4098 return isl_union_map_flat_range_product(info
.prefix
, info
.suffix
);
4101 /* Set gen->untiled_len to the number of scheduling dimensions
4102 * for the schedule of the first domain.
4103 * We assume here that this number is the same for all domains.
4105 static int set_untiled_len(__isl_take isl_map
*map
, void *user
)
4107 unsigned *untiled_len
= user
;
4109 *untiled_len
= isl_map_dim(map
, isl_dim_out
);
4115 /* Compute an appropriate schedule based on the accesses in
4116 * gen->read and gen->write.
4118 * We first compute dependences and then use those to compute
4119 * a schedule that has a parallel loop in each tilable band.
4120 * Finally, we select the outermost tilable band.
4122 static void compute_schedule(struct cuda_gen
*gen
,
4123 __isl_take isl_union_map
*sched
)
4125 isl_ctx
*ctx
= isl_union_map_get_ctx(sched
);
4126 isl_union_set
*domain
;
4127 isl_union_map
*empty
;
4128 isl_union_map
*dep_raw
, *dep2
, *dep3
, *dep
;
4129 isl_union_map
*uninitialized
;
4130 isl_schedule
*schedule
;
4132 empty
= isl_union_map_empty(isl_union_map_get_space(sched
));
4134 isl_union_map_compute_flow(isl_union_map_copy(gen
->read
),
4135 isl_union_map_copy(gen
->write
), empty
,
4136 isl_union_map_copy(sched
),
4137 &dep_raw
, NULL
, &uninitialized
, NULL
);
4138 isl_union_map_compute_flow(isl_union_map_copy(gen
->write
),
4139 isl_union_map_copy(gen
->write
),
4140 isl_union_map_copy(gen
->read
),
4141 isl_union_map_copy(sched
),
4142 &dep2
, &dep3
, NULL
, NULL
);
4143 isl_union_map_free(sched
);
4145 gen
->copy_in
= isl_union_map_range(uninitialized
);
4147 dep
= isl_union_map_union(dep2
, dep3
);
4148 dep
= isl_union_map_union(dep
, dep_raw
);
4149 dep
= isl_union_map_coalesce(dep
);
4151 domain
= extract_domain(gen
);
4152 schedule
= isl_union_set_compute_schedule(isl_union_set_copy(domain
),
4153 isl_union_map_copy(dep
), dep
);
4155 sched
= select_outer_tilable_band(gen
, schedule
);
4157 isl_union_map_foreach_map(sched
, &set_untiled_len
, &gen
->untiled_len
);
4158 sched
= isl_union_map_intersect_domain(sched
, domain
);
4161 isl_schedule_free(schedule
);
4164 static struct cuda_stmt_access
**expr_extract_access(struct pet_expr
*expr
,
4165 struct cuda_stmt_access
**next_access
)
4167 struct cuda_stmt_access
*access
;
4168 isl_ctx
*ctx
= isl_map_get_ctx(expr
->acc
.access
);
4170 access
= isl_alloc_type(ctx
, struct cuda_stmt_access
);
4172 access
->next
= NULL
;
4173 access
->read
= expr
->acc
.read
;
4174 access
->write
= expr
->acc
.write
;
4175 access
->access
= isl_map_copy(expr
->acc
.access
);
4177 *next_access
= access
;
4178 next_access
= &(*next_access
)->next
;
4182 static struct cuda_stmt_access
**expr_extract_accesses(struct pet_expr
*expr
,
4183 struct cuda_stmt_access
**next_access
)
4187 for (i
= 0; i
< expr
->n_arg
; ++i
)
4188 next_access
= expr_extract_accesses(expr
->args
[i
],
4191 if (expr
->type
== pet_expr_access
)
4192 next_access
= expr_extract_access(expr
, next_access
);
4197 static void pet_stmt_extract_accesses(struct cuda_stmt
*stmt
)
4199 struct cuda_stmt_access
**next_access
= &stmt
->accesses
;
4201 stmt
->accesses
= NULL
;
4202 expr_extract_accesses(stmt
->body
, next_access
);
4205 /* Return an array of cuda_stmt representing the statements in "scop".
4207 static struct cuda_stmt
*extract_stmts(isl_ctx
*ctx
, struct pet_scop
*scop
,
4208 __isl_keep isl_set
*context
)
4211 struct cuda_stmt
*stmts
;
4213 stmts
= isl_calloc_array(ctx
, struct cuda_stmt
, scop
->n_stmt
);
4216 for (i
= 0; i
< scop
->n_stmt
; ++i
) {
4217 struct cuda_stmt
*s
= &stmts
[i
];
4219 s
->domain
= isl_set_copy(scop
->stmts
[i
]->domain
);
4220 s
->domain
= isl_set_intersect_params(s
->domain
,
4221 isl_set_copy(context
));
4222 s
->body
= scop
->stmts
[i
]->body
;
4223 pet_stmt_extract_accesses(s
);
4229 /* Replace the scop in the "input" file by equivalent code
4230 * that uses the GPU. "scop" is assumed to correspond to this scop.
4232 * We first compute a schedule that respects the dependences
4233 * of the original program and select the outermost band
4234 * of tilable dimensions that has at least one parallel loop.
4235 * We then have three blocks of dimensions
4239 * The tilable band "B" is first tiled according to "tile.sizes", resulting
4244 * For each iteration of the T loop and for each array, we compute
4245 * the array elements accessed by that iteration, construct a rectangular
4246 * box around it and shift it to the origin. The result is used
4247 * as shared memory for the array.
4249 * We then split off at most 2 parallel loops from the T loops and
4250 * at most 3 parallel loops from the P loops
4254 * The T1/P1 loops are then tiled or "wrapped" over the blocks/threads,
4255 * according to "grid.sizes"/"block.sizes".
4257 * H T1T T1P T2 P1T P1P P2 G
4259 * Finally, the T1P and P1P iterators are equated to the block and
4260 * thread dimensions respectively and so are effectively removed.
4261 * The H loops are run on the host. The T1T, T2, P1T, P2 and G loops
4262 * are run on the GPU.
4264 * Code is generated in three stages. We first generate code for the
4265 * host (the H loops), with iterators h%d. Then, for each leaf node
4266 * of the resulting AST, we generate code for the shared loops (up to
4267 * and including T2), with iterators g%d and after equating the H loops
4268 * to h%d parameters and the T1P loops to the block dimensions.
4269 * Finally, we generate code for the remaining loops in a similar fashion.
4271 int cuda_pet(isl_ctx
*ctx
, struct pet_scop
*scop
, struct ppcg_options
*options
,
4274 isl_union_map
*sched
;
4275 struct cuda_gen gen
;
4280 scop
= pet_scop_align_params(scop
);
4283 gen
.context
= isl_set_copy(scop
->context
);
4284 gen
.context
= add_context_from_str(gen
.context
, options
->ctx
);
4285 gen
.n_stmts
= scop
->n_stmt
;
4286 gen
.stmts
= extract_stmts(ctx
, scop
, gen
.context
);
4287 gen
.read
= pet_scop_collect_reads(scop
);
4288 gen
.write
= pet_scop_collect_writes(scop
);
4289 gen
.options
= options
;
4290 gen
.state
= cloog_isl_state_malloc(gen
.ctx
);
4293 cuda_open_files(&gen
.cuda
, input
);
4295 collect_array_info(&gen
);
4297 sched
= pet_scop_collect_schedule(scop
);
4299 compute_schedule(&gen
, sched
);
4301 print_host_code(&gen
);
4303 cloog_state_free(gen
.state
);
4304 clear_cuda_gen(&gen
);
4306 cuda_close_files(&gen
.cuda
);