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"
26 #include "clast_printer.h"
28 #include "pet_printer.h"
29 #include "ppcg_options.h"
31 /* The fields stride, shift and shift_map only contain valid information
33 * If so, they express that current index is such that if you add shift,
34 * then the result is always a multiple of stride.
35 * shift_map contains the mapping
37 * i -> (i + shift)/stride
39 struct cuda_array_bound
{
45 isl_basic_map
*shift_map
;
48 struct cuda_array_info
;
50 /* A group of array references in a kernel that should be handled together.
51 * If private_bound is not NULL, then it is mapped to registers.
52 * Otherwise, if shared_bound is not NULL, it is mapped to shared memory.
53 * Otherwise, it is accessed from global memory.
55 struct cuda_array_ref_group
{
56 /* The references in this group access this array. */
57 struct cuda_array_info
*array
;
58 /* Position of this group in the list of reference groups of array. */
61 /* The following fields are use during the construction of the groups.
62 * access is the combined access relation relative to the shared
64 * write is set if any access in the group is a write.
69 /* For each index, size and offset of piece in shared memory. */
70 struct cuda_array_bound
*shared_bound
;
72 /* For each index, size and offset of piece in private memory. */
73 struct cuda_array_bound
*private_bound
;
75 /* References in this group; point to elements of a linked list. */
77 struct cuda_stmt_access
**refs
;
79 /* Last shared memory tile dimension that affects tile of this group. */
81 /* Dimension at which copying to/from shared memory is printed.
82 * if >= 0, then the value is >= last_shared
83 * if -1, then the copying is done at the leaf level.
85 int print_shared_level
;
88 struct cuda_array_info
{
94 /* Name of the array. */
96 /* Number of indices. */
98 /* For each index, a bound on the array in that direction. */
100 /* For each index, bound[i] specialized to the current kernel. */
101 isl_pw_aff
**local_bound
;
103 /* All references to this array; point to elements of a linked list. */
105 struct cuda_stmt_access
**refs
;
107 /* The reference groups associated to this array. */
109 struct cuda_array_ref_group
**groups
;
111 /* For scalars, is this scalar read-only within the entire program? */
115 /* Print the name of the local copy of a given group of array references.
117 static __isl_give isl_printer
*print_array_name(__isl_take isl_printer
*p
,
118 struct cuda_array_ref_group
*group
)
122 if (group
->private_bound
)
123 p
= isl_printer_print_str(p
, "private_");
124 else if (group
->shared_bound
)
125 p
= isl_printer_print_str(p
, "shared_");
128 p
= isl_printer_print_str(p
, group
->array
->name
);
129 if (!global
&& group
->array
->n_group
> 1) {
130 p
= isl_printer_print_str(p
, "_");
131 p
= isl_printer_print_int(p
, group
->nr
);
137 /* Collect all references to the given array and store pointers to them
140 static void collect_references(struct cuda_gen
*gen
,
141 struct cuda_array_info
*array
)
147 for (i
= 0; i
< gen
->n_stmts
; ++i
) {
148 struct cuda_stmt
*stmt
= &gen
->stmts
[i
];
149 struct cuda_stmt_access
*access
;
151 for (access
= stmt
->accesses
; access
; access
= access
->next
) {
153 name
= isl_map_get_tuple_name(access
->access
,
155 if (name
&& !strcmp(array
->name
, name
))
161 array
->refs
= isl_alloc_array(gen
->ctx
, struct cuda_stmt_access
*, n
);
165 for (i
= 0; i
< gen
->n_stmts
; ++i
) {
166 struct cuda_stmt
*stmt
= &gen
->stmts
[i
];
167 struct cuda_stmt_access
*access
;
169 for (access
= stmt
->accesses
; access
; access
= access
->next
) {
171 name
= isl_map_get_tuple_name(access
->access
,
173 if (!name
|| strcmp(array
->name
, name
))
176 array
->refs
[n
++] = access
;
181 static struct cuda_array_bound
*create_bound_list(isl_ctx
*ctx
, int n_index
)
184 struct cuda_array_bound
*bound
;
186 bound
= isl_alloc_array(ctx
, struct cuda_array_bound
, n_index
);
189 for (i
= 0; i
< n_index
; ++i
) {
190 isl_int_init(bound
[i
].size
);
192 isl_int_init(bound
[i
].stride
);
193 bound
[i
].shift
= NULL
;
194 bound
[i
].shift_map
= NULL
;
200 static void free_bound_list(struct cuda_array_bound
*bound
, int n_index
)
207 for (j
= 0; j
< n_index
; ++j
) {
208 isl_int_clear(bound
[j
].size
);
209 isl_int_clear(bound
[j
].stride
);
210 isl_aff_free(bound
[j
].lb
);
211 isl_aff_free(bound
[j
].shift
);
212 isl_basic_map_free(bound
[j
].shift_map
);
217 static struct pet_array
*find_array(struct pet_scop
*scop
,
218 __isl_keep isl_set
*accessed
)
223 id
= isl_set_get_tuple_id(accessed
);
225 for (i
= 0; i
< scop
->n_array
; ++i
) {
228 id_i
= isl_set_get_tuple_id(scop
->arrays
[i
]->extent
);
235 return i
< scop
->n_array
? scop
->arrays
[i
] : NULL
;
238 /* Compute bounds on the host arrays based on the accessed elements
239 * and collect all references to the array.
241 * If the array is zero-dimensional, i.e., a scalar, we check
242 * whether it is read-only.
244 static int extract_array_info(__isl_take isl_set
*array
, void *user
)
247 struct cuda_gen
*gen
= (struct cuda_gen
*)user
;
251 isl_pw_aff
**local_bounds
;
252 struct pet_array
*pa
;
254 n_index
= isl_set_dim(array
, isl_dim_set
);
255 name
= isl_set_get_tuple_name(array
);
256 bounds
= isl_alloc_array(isl_set_get_ctx(array
),
257 isl_pw_aff
*, n_index
);
259 local_bounds
= isl_calloc_array(isl_set_get_ctx(array
),
260 isl_pw_aff
*, n_index
);
261 assert(local_bounds
);
262 gen
->array
[gen
->n_array
].dim
= isl_set_get_space(array
);
263 gen
->array
[gen
->n_array
].name
= strdup(name
);
264 gen
->array
[gen
->n_array
].n_index
= n_index
;
265 gen
->array
[gen
->n_array
].bound
= bounds
;
266 gen
->array
[gen
->n_array
].local_bound
= local_bounds
;
268 pa
= find_array(gen
->scop
, array
);
271 gen
->array
[gen
->n_array
].type
= strdup(pa
->element_type
);
272 gen
->array
[gen
->n_array
].size
= pa
->element_size
;
276 isl_union_map
*write
;
279 write
= isl_union_map_copy(gen
->write
);
280 space
= isl_set_universe(isl_set_get_space(array
));
281 write
= isl_union_map_intersect_range(write
,
282 isl_union_set_from_set(space
));
283 empty
= isl_union_map_is_empty(write
);
284 isl_union_map_free(write
);
286 gen
->array
[gen
->n_array
].read_only
= empty
;
289 for (i
= 0; i
< n_index
; ++i
) {
294 isl_set
*size
= i
== 0 ? array
: pa
->extent
;
296 bound
= isl_set_dim_max(isl_set_copy(size
), i
);
298 dom
= isl_pw_aff_domain(isl_pw_aff_copy(bound
));
299 ls
= isl_local_space_from_space(isl_set_get_space(dom
));
300 one
= isl_aff_zero_on_domain(ls
);
301 one
= isl_aff_add_constant_si(one
, 1);
302 bound
= isl_pw_aff_add(bound
, isl_pw_aff_alloc(dom
, one
));
303 bound
= isl_pw_aff_gist(bound
, isl_set_copy(gen
->context
));
308 collect_references(gen
, &gen
->array
[gen
->n_array
]);
316 void collect_array_info(struct cuda_gen
*gen
)
318 isl_union_set
*arrays
;
320 arrays
= isl_union_map_range(isl_union_map_copy(gen
->read
));
321 arrays
= isl_union_set_union(arrays
,
322 isl_union_map_range(isl_union_map_copy(gen
->write
)));
323 arrays
= isl_union_set_coalesce(arrays
);
325 gen
->n_array
= isl_union_set_n_set(arrays
);
326 gen
->array
= isl_alloc_array(gen
->ctx
,
327 struct cuda_array_info
, gen
->n_array
);
330 isl_union_set_foreach_set(arrays
, &extract_array_info
, gen
);
331 isl_union_set_free(arrays
);
334 static void free_array_info(struct cuda_gen
*gen
)
338 for (i
= 0; i
< gen
->n_array
; ++i
) {
339 int n_index
= gen
->array
[i
].n_index
;
340 free(gen
->array
[i
].type
);
341 free(gen
->array
[i
].name
);
342 for (j
= 0; j
< n_index
; ++j
) {
343 isl_pw_aff_free(gen
->array
[i
].bound
[j
]);
344 isl_pw_aff_free(gen
->array
[i
].local_bound
[j
]);
346 isl_space_free(gen
->array
[i
].dim
);
347 free(gen
->array
[i
].bound
);
348 free(gen
->array
[i
].local_bound
);
349 free(gen
->array
[i
].refs
);
354 /* Check if a cuda array is a scalar. A scalar is a value that is not stored
355 * as an array or through a pointer reference, but as single data element. At
356 * the moment, scalars are represented as zero dimensional arrays.
358 static int cuda_array_is_scalar(struct cuda_array_info
*array
)
360 return (array
->n_index
== 0);
363 /* Is "array" a read-only scalar?
365 static int cuda_array_is_read_only_scalar(struct cuda_array_info
*array
)
367 return cuda_array_is_scalar(array
) && array
->read_only
;
370 static void declare_device_arrays(struct cuda_gen
*gen
)
374 for (i
= 0; i
< gen
->n_array
; ++i
) {
375 if (cuda_array_is_read_only_scalar(&gen
->array
[i
]))
377 fprintf(gen
->cuda
.host_c
, "%s *dev_%s;\n",
378 gen
->array
[i
].type
, gen
->array
[i
].name
);
380 fprintf(gen
->cuda
.host_c
, "\n");
383 static void print_array_size(struct cuda_gen
*gen
, FILE *out
,
384 struct cuda_array_info
*array
)
389 prn
= isl_printer_to_file(gen
->ctx
, out
);
390 prn
= isl_printer_set_output_format(prn
, ISL_FORMAT_C
);
391 for (i
= 0; i
< array
->n_index
; ++i
) {
392 prn
= isl_printer_print_str(prn
, "(");
393 prn
= isl_printer_print_pw_aff(prn
, array
->bound
[i
]);
394 prn
= isl_printer_print_str(prn
, ") * ");
396 prn
= isl_printer_print_str(prn
, "sizeof(");
397 prn
= isl_printer_print_str(prn
, array
->type
);
398 prn
= isl_printer_print_str(prn
, ")");
399 isl_printer_free(prn
);
402 static void allocate_device_arrays(struct cuda_gen
*gen
)
406 for (i
= 0; i
< gen
->n_array
; ++i
) {
407 if (cuda_array_is_read_only_scalar(&gen
->array
[i
]))
409 fprintf(gen
->cuda
.host_c
,
410 "cudaCheckReturn(cudaMalloc((void **) &dev_%s, ",
412 print_array_size(gen
, gen
->cuda
.host_c
, &gen
->array
[i
]);
413 fprintf(gen
->cuda
.host_c
, "));\n");
415 fprintf(gen
->cuda
.host_c
, "\n");
418 static void free_device_arrays(struct cuda_gen
*gen
)
422 for (i
= 0; i
< gen
->n_array
; ++i
) {
423 if (cuda_array_is_read_only_scalar(&gen
->array
[i
]))
425 fprintf(gen
->cuda
.host_c
, "cudaCheckReturn(cudaFree(dev_%s));\n",
430 static void copy_arrays_to_device(struct cuda_gen
*gen
)
434 for (i
= 0; i
< gen
->n_array
; ++i
) {
439 if (cuda_array_is_read_only_scalar(&gen
->array
[i
]))
442 dim
= isl_space_copy(gen
->array
[i
].dim
);
443 read_i
= isl_union_set_extract_set(gen
->copy_in
, dim
);
444 empty
= isl_set_fast_is_empty(read_i
);
445 isl_set_free(read_i
);
449 fprintf(gen
->cuda
.host_c
, "cudaCheckReturn(cudaMemcpy(dev_%s,",
452 if (cuda_array_is_scalar(&(gen
->array
[i
])))
453 fprintf(gen
->cuda
.host_c
, " &%s, ",
456 fprintf(gen
->cuda
.host_c
, " %s, ", gen
->array
[i
].name
);
458 print_array_size(gen
, gen
->cuda
.host_c
, &gen
->array
[i
]);
459 fprintf(gen
->cuda
.host_c
, ", cudaMemcpyHostToDevice));\n");
461 fprintf(gen
->cuda
.host_c
, "\n");
464 static void copy_arrays_from_device(struct cuda_gen
*gen
)
467 isl_union_set
*write
;
468 write
= isl_union_map_range(isl_union_map_copy(gen
->write
));
470 for (i
= 0; i
< gen
->n_array
; ++i
) {
475 dim
= isl_space_copy(gen
->array
[i
].dim
);
476 write_i
= isl_union_set_extract_set(write
, dim
);
477 empty
= isl_set_fast_is_empty(write_i
);
478 isl_set_free(write_i
);
482 fprintf(gen
->cuda
.host_c
, "cudaCheckReturn(cudaMemcpy(");
483 if (cuda_array_is_scalar(&gen
->array
[i
]))
484 fprintf(gen
->cuda
.host_c
, "&%s, ", gen
->array
[i
].name
);
486 fprintf(gen
->cuda
.host_c
, "%s, ", gen
->array
[i
].name
);
487 fprintf(gen
->cuda
.host_c
, "dev_%s, ", gen
->array
[i
].name
);
488 print_array_size(gen
, gen
->cuda
.host_c
, &gen
->array
[i
]);
489 fprintf(gen
->cuda
.host_c
, ", cudaMemcpyDeviceToHost));\n");
492 isl_union_set_free(write
);
493 fprintf(gen
->cuda
.host_c
, "\n");
496 /* Internal data structure for extract_size_of_type.
497 * "type" specifies the name of the space that we want to extract.
498 * "res" is used to store the subset of that space.
500 struct ppcg_extract_size_data
{
505 /* This function is called for each set in a union_set.
506 * If the name of the set matches data->type, we store the
509 static int extract_size_of_type(__isl_take isl_set
*size
, void *user
)
511 struct ppcg_extract_size_data
*data
= user
;
514 name
= isl_set_get_tuple_name(size
);
515 if (name
&& !strcmp(name
, data
->type
)) {
524 /* Given a union map { kernel[i] -> *[...] },
525 * return the range in the space called "type" for the kernel with
526 * sequence number "id".
528 static __isl_give isl_set
*extract_sizes(__isl_keep isl_union_map
*sizes
,
529 const char *type
, int id
)
533 isl_union_set
*local_sizes
;
534 struct ppcg_extract_size_data data
= { type
, NULL
};
539 space
= isl_union_map_get_space(sizes
);
540 space
= isl_space_set_from_params(space
);
541 space
= isl_space_add_dims(space
, isl_dim_set
, 1);
542 space
= isl_space_set_tuple_name(space
, isl_dim_set
, "kernel");
543 dom
= isl_set_universe(space
);
544 dom
= isl_set_fix_si(dom
, isl_dim_set
, 0, id
);
546 local_sizes
= isl_union_set_apply(isl_union_set_from_set(dom
),
547 isl_union_map_copy(sizes
));
548 isl_union_set_foreach_set(local_sizes
, &extract_size_of_type
, &data
);
549 isl_union_set_free(local_sizes
);
553 /* Given a singleton set, extract the first (at most *len) elements
554 * of the single integer tuple into *sizes and update *len if needed.
556 static void read_sizes_from_set(__isl_take isl_set
*set
, int *sizes
, int *len
)
565 dim
= isl_set_dim(set
, isl_dim_set
);
571 for (i
= 0; i
< *len
; ++i
) {
574 ok
= isl_set_plain_is_fixed(set
, isl_dim_set
, i
, &v
);
577 sizes
[i
] = isl_int_get_si(v
);
585 /* Extract user specified "tile" sizes from the "sizes" command line option,
586 * defaulting to option->tile_size in each dimension.
588 static void read_tile_sizes(struct cuda_gen
*gen
)
593 gen
->tile_size
= isl_alloc_array(gen
->ctx
, int, gen
->tile_len
);
594 assert(gen
->tile_size
);
595 for (n
= 0; n
< gen
->tile_len
; ++n
)
596 gen
->tile_size
[n
] = gen
->options
->tile_size
;
598 size
= extract_sizes(gen
->sizes
, "tile", gen
->kernel_id
);
599 read_sizes_from_set(size
, gen
->tile_size
, &gen
->tile_len
);
601 if (gen
->n_parallel
> gen
->tile_len
)
602 gen
->n_parallel
= gen
->tile_len
;
605 /* Extract user specified "block" sizes from the "sizes" command line option,
606 * after filling in some potentially useful defaults.
608 static void read_block_sizes(struct cuda_gen
*gen
)
614 gen
->n_block
= (n
<= 3) ? n
: 3;
615 switch (gen
->n_block
) {
617 gen
->block_dim
[0] = 512;
620 gen
->block_dim
[0] = 32;
621 gen
->block_dim
[1] = 16;
624 gen
->block_dim
[0] = 32;
625 gen
->block_dim
[1] = 4;
626 gen
->block_dim
[2] = 4;
630 size
= extract_sizes(gen
->sizes
, "block", gen
->kernel_id
);
631 read_sizes_from_set(size
, gen
->block_dim
, &gen
->n_block
);
634 /* Extract user specified "grid" sizes from the "sizes" command line option,
635 * after filling in some potentially useful defaults.
637 static void read_grid_sizes(struct cuda_gen
*gen
)
639 int n
= gen
->n_parallel
;
642 gen
->n_grid
= (n
<= 2) ? n
: 2;
643 switch (gen
->n_grid
) {
645 gen
->grid_dim
[0] = 32768;
648 gen
->grid_dim
[0] = 256;
649 gen
->grid_dim
[1] = 256;
653 size
= extract_sizes(gen
->sizes
, "grid", gen
->kernel_id
);
654 read_sizes_from_set(size
, gen
->grid_dim
, &gen
->n_grid
);
657 /* Extract user specified sizes from the "sizes" command line option
658 * after filling in some potentially useful defaults.
660 static void read_sizes(struct cuda_gen
*gen
)
662 read_tile_sizes(gen
);
663 read_block_sizes(gen
);
664 read_grid_sizes(gen
);
667 static void free_stmts(struct cuda_stmt
*stmts
, int n
)
671 for (i
= 0; i
< n
; ++i
) {
672 struct cuda_stmt_access
*access
, *next
;
674 for (access
= stmts
[i
].accesses
; access
; access
= next
) {
676 isl_map_free(access
->access
);
680 isl_set_free(stmts
[i
].domain
);
685 void clear_cuda_gen(struct cuda_gen
*gen
)
687 free_stmts(gen
->stmts
, gen
->n_stmts
);
688 free_array_info(gen
);
689 isl_union_map_free(gen
->sizes
);
690 isl_set_free(gen
->context
);
691 isl_union_set_free(gen
->copy_in
);
692 isl_union_map_free(gen
->sched
);
693 isl_union_map_free(gen
->read
);
694 isl_union_map_free(gen
->write
);
697 static void print_reverse_list(FILE *out
, int len
, int *list
)
705 for (i
= 0; i
< len
; ++i
) {
708 fprintf(out
, "%d", list
[len
- 1 - i
]);
713 static void print_kernel_launch(struct cuda_gen
*gen
,
714 __isl_keep isl_union_set
*arrays
)
721 print_indent(gen
->code
.dst
, gen
->code
.indent
);
722 fprintf(gen
->code
.dst
, "kernel%d <<<k%d_dimGrid, k%d_dimBlock>>> (",
723 gen
->kernel_id
, gen
->kernel_id
, gen
->kernel_id
);
724 fprintf(gen
->cuda
.kernel_c
, "__global__ void kernel%d(",
726 fprintf(gen
->cuda
.kernel_h
, "__global__ void kernel%d(",
729 for (i
= 0; i
< gen
->n_array
; ++i
) {
734 dim
= isl_space_copy(gen
->array
[i
].dim
);
735 arr
= isl_union_set_extract_set(arrays
, dim
);
736 empty
= isl_set_fast_is_empty(arr
);
742 fprintf(gen
->code
.dst
, ", ");
743 fprintf(gen
->cuda
.kernel_c
, ", ");
744 fprintf(gen
->cuda
.kernel_h
, ", ");
747 if (cuda_array_is_read_only_scalar(&gen
->array
[i
])) {
748 fprintf(gen
->code
.dst
, "%s", gen
->array
[i
].name
);
749 fprintf(gen
->cuda
.kernel_c
, "%s %s",
750 gen
->array
[i
].type
, gen
->array
[i
].name
);
751 fprintf(gen
->cuda
.kernel_h
, "%s %s",
752 gen
->array
[i
].type
, gen
->array
[i
].name
);
754 fprintf(gen
->code
.dst
, "dev_%s", gen
->array
[i
].name
);
755 fprintf(gen
->cuda
.kernel_c
, "%s *%s",
756 gen
->array
[i
].type
, gen
->array
[i
].name
);
757 fprintf(gen
->cuda
.kernel_h
, "%s *%s",
758 gen
->array
[i
].type
, gen
->array
[i
].name
);
764 dim
= isl_union_set_get_space(arrays
);
765 nparam
= isl_space_dim(dim
, isl_dim_param
);
766 for (i
= 0; i
< nparam
; ++i
) {
767 const char *name
= isl_space_get_dim_name(dim
, isl_dim_param
, i
);
769 fprintf(gen
->code
.dst
, ", ");
770 fprintf(gen
->cuda
.kernel_c
, ", ");
771 fprintf(gen
->cuda
.kernel_h
, ", ");
773 fprintf(gen
->code
.dst
, "%s", name
);
774 fprintf(gen
->cuda
.kernel_c
, "int %s", name
);
775 fprintf(gen
->cuda
.kernel_h
, "int %s", name
);
780 for (i
= 0; i
< gen
->tile_first
; ++i
) {
782 fprintf(gen
->code
.dst
, ", ");
783 fprintf(gen
->cuda
.kernel_c
, ", ");
784 fprintf(gen
->cuda
.kernel_h
, ", ");
786 fprintf(gen
->code
.dst
, "h%d", i
);
787 fprintf(gen
->cuda
.kernel_c
, "int h%d", i
);
788 fprintf(gen
->cuda
.kernel_h
, "int h%d", i
);
792 fprintf(gen
->code
.dst
, ");\n");
793 fprintf(gen
->cuda
.kernel_c
, ")\n");
794 fprintf(gen
->cuda
.kernel_h
, ");\n");
796 fprintf(gen
->code
.dst
, "cudaCheckKernel();\n");
799 /* Construct a map from a domain of dimensionality "len"
800 * to a domain of dimensionality "len" + "tile_len" that tiles
801 * the "tile_len" coordinates starting at "first".
802 * In particular, [s_i] -> [s_i / tile_size[i], s_i % tile_size[i]].
803 * "dim" prescribes the parameters.
805 static __isl_give isl_map
*tile(__isl_take isl_space
*dim
, int len
,
806 int first
, int tile_len
, int *tile_size
)
816 dim
= isl_space_add_dims(dim
, isl_dim_in
, len
);
817 dim
= isl_space_add_dims(dim
, isl_dim_out
, len
+ tile_len
);
818 bmap
= isl_basic_map_universe(isl_space_copy(dim
));
819 ls
= isl_local_space_from_space(dim
);
821 for (i
= 0; i
< len
- tile_len
; ++i
) {
822 int j
= i
< first
? i
: i
+ tile_len
;
823 int k
= i
< first
? i
: i
+ 2 * tile_len
;
825 c
= isl_equality_alloc(isl_local_space_copy(ls
));
826 isl_int_set_si(v
, -1);
827 isl_constraint_set_coefficient(c
, isl_dim_in
, j
, v
);
828 isl_int_set_si(v
, 1);
829 isl_constraint_set_coefficient(c
, isl_dim_out
, k
, v
);
830 bmap
= isl_basic_map_add_constraint(bmap
, c
);
833 for (i
= 0; i
< tile_len
; ++i
) {
834 c
= isl_equality_alloc(isl_local_space_copy(ls
));
835 isl_int_set_si(v
, -1);
836 isl_constraint_set_coefficient(c
, isl_dim_in
, first
+ i
, v
);
837 isl_int_set_si(v
, tile_size
[i
]);
838 isl_constraint_set_coefficient(c
, isl_dim_out
, first
+ i
, v
);
839 isl_int_set_si(v
, 1);
840 isl_constraint_set_coefficient(c
, isl_dim_out
,
841 first
+ i
+ tile_len
, v
);
842 bmap
= isl_basic_map_add_constraint(bmap
, c
);
844 c
= isl_inequality_alloc(isl_local_space_copy(ls
));
845 isl_int_set_si(v
, 1);
846 isl_constraint_set_coefficient(c
, isl_dim_out
,
847 first
+ i
+ tile_len
, v
);
848 bmap
= isl_basic_map_add_constraint(bmap
, c
);
850 c
= isl_inequality_alloc(isl_local_space_copy(ls
));
851 isl_int_set_si(v
, -1);
852 isl_constraint_set_coefficient(c
, isl_dim_out
,
853 first
+ i
+ tile_len
, v
);
854 isl_int_set_si(v
, tile_size
[i
] - 1);
855 isl_constraint_set_constant(c
, v
);
856 bmap
= isl_basic_map_add_constraint(bmap
, c
);
859 isl_local_space_free(ls
);
862 return isl_map_from_basic_map(bmap
);
865 /* Construct a map from a domain of dimensionality "len"
866 * to a domain of dimensionality "len" + "wrap_len" that "wraps"
867 * the "wrap_len" coordinates starting at "first" according to "wrap_size".
868 * In particular, [s_i] -> [s_i, s_i % wrap_size[i]].
869 * To do so, we need extra variables corresponding to [s_i / wrap_size[i]],
870 * that are projected out at the end.
871 * "dim" prescribes the parameters.
873 static __isl_give isl_map
*wrap(__isl_take isl_space
*dim
, int len
,
874 int first
, int wrap_len
, int *wrap_size
)
881 dim
= isl_space_add_dims(dim
, isl_dim_in
, len
);
882 dim
= isl_space_add_dims(dim
, isl_dim_out
, len
+ 2 * wrap_len
);
883 bmap
= isl_basic_map_universe(isl_space_copy(dim
));
884 ls
= isl_local_space_from_space(dim
);
886 for (i
= 0; i
< len
; ++i
) {
887 int k
= i
< first
+ wrap_len
? i
: i
+ 2 * wrap_len
;
889 c
= isl_equality_alloc(isl_local_space_copy(ls
));
890 isl_constraint_set_coefficient_si(c
, isl_dim_in
, i
, -1);
891 isl_constraint_set_coefficient_si(c
, isl_dim_out
, k
, 1);
892 bmap
= isl_basic_map_add_constraint(bmap
, c
);
895 for (i
= 0; i
< wrap_len
; ++i
) {
896 c
= isl_equality_alloc(isl_local_space_copy(ls
));
897 isl_constraint_set_coefficient_si(c
, isl_dim_out
,
899 isl_constraint_set_coefficient_si(c
, isl_dim_out
,
900 first
+ wrap_len
+ i
, 1);
901 isl_constraint_set_coefficient_si(c
, isl_dim_out
,
902 first
+ 2 * wrap_len
+ i
, wrap_size
[i
]);
903 bmap
= isl_basic_map_add_constraint(bmap
, c
);
905 c
= isl_inequality_alloc(isl_local_space_copy(ls
));
906 isl_constraint_set_coefficient_si(c
, isl_dim_out
,
907 first
+ wrap_len
+ i
, 1);
908 bmap
= isl_basic_map_add_constraint(bmap
, c
);
910 c
= isl_inequality_alloc(isl_local_space_copy(ls
));
911 isl_constraint_set_coefficient_si(c
, isl_dim_out
,
912 first
+ wrap_len
+ i
, -1);
913 isl_constraint_set_constant_si(c
, wrap_size
[i
] - 1);
914 bmap
= isl_basic_map_add_constraint(bmap
, c
);
917 isl_local_space_free(ls
);
919 bmap
= isl_basic_map_project_out(bmap
, isl_dim_out
,
920 first
+ 2 * wrap_len
, wrap_len
);
922 return isl_map_from_basic_map(bmap
);
925 /* Add "n" parameters named prefix%d.
927 static __isl_give isl_set
*add_params( __isl_take isl_set
*set
,
928 int n
, const char *prefix
)
934 nparam
= isl_set_dim(set
, isl_dim_param
);
935 set
= isl_set_add_dims(set
, isl_dim_param
, n
);
937 for (i
= 0; i
< n
; ++i
) {
938 snprintf(name
, sizeof(name
), "%s%d", prefix
, i
);
939 set
= isl_set_set_dim_name(set
, isl_dim_param
,
946 /* Equate the "n" dimensions of "set" starting at "first" to
947 * freshly created parameters named prefix%d.
949 static __isl_give isl_set
*parametrize(__isl_take isl_set
*set
,
950 int first
, int n
, const char *prefix
)
960 nparam
= isl_set_dim(set
, isl_dim_param
);
962 set
= add_params(set
, n
, prefix
);
964 dim
= isl_set_get_space(set
);
965 bset
= isl_basic_set_universe(isl_space_copy(dim
));
966 ls
= isl_local_space_from_space(dim
);
970 for (i
= 0; i
< n
; ++i
) {
971 c
= isl_equality_alloc(isl_local_space_copy(ls
));
972 isl_int_set_si(v
, -1);
973 isl_constraint_set_coefficient(c
, isl_dim_param
, nparam
+ i
, v
);
974 isl_int_set_si(v
, 1);
975 isl_constraint_set_coefficient(c
, isl_dim_set
, first
+ i
, v
);
976 bset
= isl_basic_set_add_constraint(bset
, c
);
980 isl_local_space_free(ls
);
982 return isl_set_intersect(set
, isl_set_from_basic_set(bset
));
985 static __isl_give isl_set
*parametrization(__isl_take isl_space
*dim
,
986 int len
, int first
, int n
, const char *prefix
)
990 dim
= isl_space_add_dims(dim
, isl_dim_set
, len
);
991 set
= isl_set_universe(dim
);
993 return parametrize(set
, first
, n
, prefix
);
996 /* Tile the B loops over the tile sizes and then tile/wrap
997 * the T1 loops over the blocks.
999 static __isl_give isl_union_map
*tile_schedule(struct cuda_gen
*gen
,
1000 __isl_take isl_union_map
*sched
)
1003 isl_map
*tiling
, *block_tiling
;
1005 dim
= isl_union_map_get_space(sched
);
1006 tiling
= tile(isl_space_copy(dim
), gen
->untiled_len
,
1007 gen
->tile_first
, gen
->tile_len
, gen
->tile_size
);
1009 if (gen
->options
->wrap
)
1010 block_tiling
= wrap(dim
, gen
->untiled_len
+ gen
->tile_len
,
1011 gen
->tile_first
, gen
->n_grid
, gen
->grid_dim
);
1013 block_tiling
= tile(dim
, gen
->untiled_len
+ gen
->tile_len
,
1014 gen
->tile_first
, gen
->n_grid
, gen
->grid_dim
);
1016 gen
->tiled_len
= gen
->untiled_len
+ gen
->tile_len
+ gen
->n_grid
;
1018 tiling
= isl_map_apply_range(tiling
, block_tiling
);
1020 sched
= isl_union_map_apply_range(sched
,
1021 isl_union_map_from_map(tiling
));
1023 gen
->shared_len
= gen
->tile_first
+ gen
->tile_len
+ gen
->n_grid
;
1028 static __isl_give isl_union_map
*parametrize_tiled_schedule(
1029 struct cuda_gen
*gen
, __isl_take isl_union_map
*sched
)
1034 dim
= isl_union_map_get_space(sched
);
1035 par
= parametrization(dim
, gen
->tiled_len
, 0, gen
->tile_first
, "h");
1036 sched
= isl_union_map_intersect_range(sched
,
1037 isl_union_set_from_set(par
));
1039 dim
= isl_union_map_get_space(sched
);
1040 par
= parametrization(dim
, gen
->tiled_len
,
1041 gen
->tile_first
+ gen
->n_grid
, gen
->n_grid
, "b");
1042 sched
= isl_union_map_intersect_range(sched
,
1043 isl_union_set_from_set(par
));
1048 /* Tile/wrap the P1 loops over the threads.
1050 static __isl_give isl_union_map
*thread_tile_schedule(struct cuda_gen
*gen
,
1051 __isl_take isl_union_map
*sched
)
1057 dim
= isl_union_map_get_space(sched
);
1059 if (gen
->options
->wrap
)
1060 tiling
= wrap(isl_space_copy(dim
), gen
->tiled_len
,
1061 gen
->shared_len
, gen
->n_block
, gen
->block_dim
);
1063 tiling
= tile(isl_space_copy(dim
), gen
->tiled_len
,
1064 gen
->shared_len
, gen
->n_block
, gen
->block_dim
);
1065 gen
->thread_tiled_len
= gen
->tiled_len
+ gen
->n_block
;
1067 sched
= isl_union_map_apply_range(sched
,
1068 isl_union_map_from_map(tiling
));
1070 par
= parametrization(dim
, gen
->thread_tiled_len
,
1071 gen
->tile_first
+ gen
->tile_len
+ gen
->n_grid
+ gen
->n_block
,
1073 sched
= isl_union_map_intersect_range(sched
,
1074 isl_union_set_from_set(par
));
1076 gen
->shared_len
= gen
->tile_first
+ gen
->tile_len
+ gen
->n_grid
;
1081 /* If the user asked for it, scale the shared memory tile loops
1082 * (T1T and T2) of "sched" by gen->tile_size[i].
1083 * If we are not performing "wrapping", then additionally scale the T1P
1084 * loops by gen->grid_dim[i].
1086 static __isl_give isl_union_map
*scale_tile_loops(struct cuda_gen
*gen
,
1087 __isl_take isl_union_map
*sched
)
1091 isl_basic_map
*scale
;
1093 isl_local_space
*ls
;
1095 if (!gen
->options
->scale_tile_loops
)
1098 dim
= isl_union_map_get_space(sched
);
1099 dim
= isl_space_add_dims(dim
, isl_dim_in
, gen
->tiled_len
);
1100 dim
= isl_space_add_dims(dim
, isl_dim_out
, gen
->tiled_len
);
1101 scale
= isl_basic_map_universe(isl_space_copy(dim
));
1102 ls
= isl_local_space_from_space(dim
);
1104 for (i
= 0; i
< gen
->tiled_len
; ++i
) {
1107 if (i
>= gen
->tile_first
&& i
< gen
->tile_first
+ gen
->n_grid
) {
1108 f
= gen
->tile_size
[i
- gen
->tile_first
];
1109 if (!gen
->options
->wrap
)
1110 f
*= gen
->grid_dim
[i
- gen
->tile_first
];
1111 } else if (i
>= gen
->tile_first
+ gen
->n_grid
&&
1112 i
< gen
->tile_first
+ gen
->n_grid
+ gen
->tile_len
) {
1113 f
= gen
->tile_size
[i
- (gen
->tile_first
+ gen
->n_grid
)];
1116 c
= isl_equality_alloc(isl_local_space_copy(ls
));
1117 isl_constraint_set_coefficient_si(c
, isl_dim_in
, i
, f
);
1118 isl_constraint_set_coefficient_si(c
, isl_dim_out
, i
, -1);
1119 scale
= isl_basic_map_add_constraint(scale
, c
);
1122 isl_local_space_free(ls
);
1124 sched
= isl_union_map_apply_range(sched
,
1125 isl_union_map_from_map(isl_map_from_basic_map(scale
)));
1130 /* If we are not performing "wrapping" and if the user asked for it,
1131 * scale the thread tile loops (P1T) of "sched" by gen->block_dim[i].
1133 static __isl_give isl_union_map
*scale_thread_tile_loops(struct cuda_gen
*gen
,
1134 __isl_take isl_union_map
*sched
)
1138 isl_basic_map
*scale
;
1140 isl_local_space
*ls
;
1142 if (gen
->options
->wrap
)
1144 if (!gen
->options
->scale_tile_loops
)
1147 dim
= isl_union_map_get_space(sched
);
1148 dim
= isl_space_add_dims(dim
, isl_dim_in
, gen
->thread_tiled_len
);
1149 dim
= isl_space_add_dims(dim
, isl_dim_out
, gen
->thread_tiled_len
);
1150 scale
= isl_basic_map_universe(isl_space_copy(dim
));
1151 ls
= isl_local_space_from_space(dim
);
1153 for (i
= 0; i
< gen
->thread_tiled_len
; ++i
) {
1156 if (i
>= gen
->shared_len
&&
1157 i
< gen
->shared_len
+ gen
->n_block
)
1158 f
= gen
->block_dim
[i
- gen
->shared_len
];
1160 c
= isl_equality_alloc(isl_local_space_copy(ls
));
1161 isl_constraint_set_coefficient_si(c
, isl_dim_in
, i
, f
);
1162 isl_constraint_set_coefficient_si(c
, isl_dim_out
, i
, -1);
1163 scale
= isl_basic_map_add_constraint(scale
, c
);
1166 isl_local_space_free(ls
);
1168 sched
= isl_union_map_apply_range(sched
,
1169 isl_union_map_from_map(isl_map_from_basic_map(scale
)));
1174 /* If we are not performing "wrapping" and if the user asked for it,
1175 * scale the "n_tile" loops starting at "first" of "sched" by gen->block_dim[i].
1177 static __isl_give isl_union_map
*scale_access_tile_loops(struct cuda_gen
*gen
,
1178 __isl_take isl_union_map
*sched
, int len
, int first
, int n_tile
)
1182 isl_basic_map
*scale
;
1184 isl_local_space
*ls
;
1186 if (gen
->options
->wrap
)
1188 if (!gen
->options
->scale_tile_loops
)
1191 dim
= isl_union_map_get_space(sched
);
1192 dim
= isl_space_add_dims(dim
, isl_dim_in
, len
);
1193 dim
= isl_space_add_dims(dim
, isl_dim_out
, len
);
1194 scale
= isl_basic_map_universe(isl_space_copy(dim
));
1195 ls
= isl_local_space_from_space(dim
);
1197 for (i
= 0; i
< len
; ++i
) {
1200 if (i
>= first
&& i
< first
+ n_tile
)
1201 f
= gen
->block_dim
[i
- first
];
1203 c
= isl_equality_alloc(isl_local_space_copy(ls
));
1204 isl_constraint_set_coefficient_si(c
, isl_dim_in
, i
, f
);
1205 isl_constraint_set_coefficient_si(c
, isl_dim_out
, i
, -1);
1206 scale
= isl_basic_map_add_constraint(scale
, c
);
1209 isl_local_space_free(ls
);
1211 sched
= isl_union_map_apply_range(sched
,
1212 isl_union_map_from_map(isl_map_from_basic_map(scale
)));
1217 /* If print_user_stmt is set, we want to print the statements ourselves,
1218 * instead of relying on the C preprocessor. If so, we need to use
1219 * the stop option so that the domains will be saved on the statement
1222 static void print_cloog_shared_body(struct cuda_gen
*gen
,
1223 __isl_keep isl_set
*context
, __isl_keep isl_union_map
*sched
, int len
,
1224 void (*print_user_stmt
)(struct clast_printer_info
*info
,
1225 struct clast_user_stmt
*s
),
1229 CloogOptions
*options
;
1230 CloogDomain
*cloog_context
;
1231 CloogUnionDomain
*ud
;
1233 struct clast_stmt
*stmt
;
1236 sched
= isl_union_map_copy(sched
);
1237 sched
= isl_union_map_align_params(sched
, isl_set_get_space(context
));
1239 options
= cloog_options_malloc(gen
->state
);
1240 options
->language
= CLOOG_LANGUAGE_C
;
1241 options
->strides
= 1;
1245 options
->override
= 1;
1246 options
->save_domains
= 1;
1247 options
->noscalars
= 1;
1248 options
->first_unroll
= first_unroll
;
1250 ud
= cloog_union_domain_from_isl_union_map(sched
);
1251 for (i
= 0; i
< len
; ++i
) {
1252 snprintf(name
, sizeof(name
), "c%d", i
);
1253 ud
= cloog_union_domain_set_name(ud
, CLOOG_SCAT
, i
, name
);
1255 cloog_context
= cloog_domain_from_isl_set(isl_set_copy(context
));
1256 input
= cloog_input_alloc(cloog_context
, ud
);
1258 stmt
= cloog_clast_create_from_input(input
, options
);
1260 gen
->stmt_code
.indent
= gen
->kernel_code
.indent
;
1261 gen
->stmt_code
.dst
= gen
->cuda
.kernel_c
;
1262 gen
->stmt_code
.print_user_stmt
= print_user_stmt
;
1263 gen
->stmt_code
.print_user_stmt_list
= NULL
;
1264 gen
->stmt_code
.print_for_head
= NULL
;
1265 gen
->stmt_code
.print_for_foot
= NULL
;
1266 gen
->stmt_code
.user
= gen
;
1267 print_clast(&gen
->stmt_code
, stmt
);
1269 cloog_clast_free(stmt
);
1270 cloog_options_free(options
);
1273 /* Add "len" parameters p[i] called prefix%d,
1274 * with bounds to 0 <= p[i] < size[i].
1276 __isl_give isl_set
*add_bounded_parameters(__isl_take isl_set
*set
,
1277 int len
, int *size
, const char *prefix
)
1283 isl_basic_set
*bset
;
1285 isl_local_space
*ls
;
1288 nparam
= isl_set_dim(set
, isl_dim_param
);
1289 set
= isl_set_add_dims(set
, isl_dim_param
, len
);
1291 for (i
= 0; i
< len
; ++i
) {
1292 snprintf(name
, sizeof(name
), "%s%d", prefix
, i
);
1293 set
= isl_set_set_dim_name(set
, isl_dim_param
,
1297 dim
= isl_set_get_space(set
);
1298 bset
= isl_basic_set_universe(isl_space_copy(dim
));
1299 ls
= isl_local_space_from_space(dim
);
1303 for (i
= 0; i
< len
; ++i
) {
1304 c
= isl_inequality_alloc(isl_local_space_copy(ls
));
1305 isl_int_set_si(v
, 1);
1306 isl_constraint_set_coefficient(c
, isl_dim_param
, nparam
+ i
, v
);
1307 bset
= isl_basic_set_add_constraint(bset
, c
);
1309 c
= isl_inequality_alloc(isl_local_space_copy(ls
));
1310 isl_int_set_si(v
, -1);
1311 isl_constraint_set_coefficient(c
, isl_dim_param
, nparam
+ i
, v
);
1312 isl_int_set_si(v
, size
[i
] - 1);
1313 isl_constraint_set_constant(c
, v
);
1314 bset
= isl_basic_set_add_constraint(bset
, c
);
1318 isl_local_space_free(ls
);
1320 return isl_set_intersect(set
, isl_set_from_basic_set(bset
));
1323 static void print_shared_body(struct cuda_gen
*gen
,
1324 __isl_keep isl_set
*shared_domain
, __isl_keep isl_union_map
*sched
,
1325 int len
, void (*print_user_stmt
)(struct clast_printer_info
*info
,
1326 struct clast_user_stmt
*s
),
1331 context
= isl_set_copy(shared_domain
);
1332 context
= parametrize(context
, 0, gen
->shared_len
, "g");
1333 context
= isl_set_project_out(context
, isl_dim_set
, 0, gen
->shared_len
);
1334 context
= add_bounded_parameters(context
,
1335 gen
->n_block
, gen
->block_dim
, "t");
1337 print_cloog_shared_body(gen
, context
, sched
, len
, print_user_stmt
,
1340 isl_set_free(context
);
1343 /* Given a tile of an array, construct a map that maps each element
1344 * of the tile to a copy of the tile shifted to the origin
1345 * (based on the lower bounds in group->private_bound or group->shared_bound).
1346 * If any of the indices is strided, then {private,shared}_bound[i].shift_map
1347 * is applied to the index first.
1348 * The domain of the resulting map is "access",
1349 * while the range space is anonymous.
1351 static __isl_give isl_map
*shift_access(__isl_take isl_set
*access
,
1352 struct cuda_array_ref_group
*group
)
1356 isl_basic_set
*bset
;
1357 isl_basic_map
*bmap
;
1359 isl_basic_set
*offset
;
1360 isl_basic_map
*shift
;
1361 isl_basic_map
*pre_shift
;
1364 struct cuda_array_bound
*bounds
;
1365 int n_index
= group
->array
->n_index
;
1367 bounds
= group
->private_bound
;
1369 bounds
= group
->shared_bound
;
1371 dim
= isl_set_get_space(access
);
1372 dim
= isl_space_drop_dims(dim
, isl_dim_set
, 0, n_index
);
1373 offset
= isl_basic_set_universe(dim
);
1374 for (i
= 0; i
< n_index
; ++i
) {
1375 lb
= isl_aff_copy(bounds
[i
].lb
);
1376 bmap
= isl_basic_map_from_aff(lb
);
1377 bset
= isl_basic_map_range(bmap
);
1378 offset
= isl_basic_set_flat_product(offset
, bset
);
1380 offset
= isl_basic_set_neg(offset
);
1382 dim
= isl_space_map_from_set(isl_set_get_space(access
));
1383 shift
= isl_basic_map_identity(dim
);
1384 shift
= isl_basic_map_set_tuple_name(shift
, isl_dim_out
, NULL
);
1386 bset
= isl_basic_set_universe(isl_set_get_space(access
));
1387 bmap
= isl_basic_map_from_domain_and_range(bset
, offset
);
1389 shift
= isl_basic_map_sum(shift
, bmap
);
1391 dim
= isl_set_get_space(access
);
1392 dim
= isl_space_drop_dims(dim
, isl_dim_set
, 0, n_index
);
1393 dim
= isl_space_map_from_set(dim
);
1394 pre_shift
= isl_basic_map_universe(isl_space_copy(dim
));
1395 dim
= isl_space_add_dims(dim
, isl_dim_in
, 1);
1396 dim
= isl_space_add_dims(dim
, isl_dim_out
, 1);
1397 for (i
= 0; i
< n_index
; ++i
) {
1398 if (!bounds
[i
].shift_map
)
1399 bmap
= isl_basic_map_identity(isl_space_copy(dim
));
1401 bmap
= isl_basic_map_copy(bounds
[i
].shift_map
);
1402 pre_shift
= isl_basic_map_flat_product(pre_shift
, bmap
);
1404 isl_space_free(dim
);
1405 name
= isl_basic_map_get_tuple_name(shift
, isl_dim_in
);
1406 pre_shift
= isl_basic_map_set_tuple_name(pre_shift
, isl_dim_in
, name
);
1407 pre_shift
= isl_basic_map_set_tuple_name(pre_shift
, isl_dim_out
, name
);
1408 shift
= isl_basic_map_apply_range(pre_shift
, shift
);
1410 sched
= isl_map_from_basic_map(shift
);
1411 sched
= isl_map_intersect_domain(sched
, access
);
1416 /* Construct a schedule for iterating over all elements in the given
1417 * piece of an array. The schedule iterates over a copy of the piece
1418 * that is shifted to the origin.
1419 * We subsequently also perform the tiling/wrapping over the threads.
1421 * In particular, we tile the final iterators so that the final thread
1422 * dimension runs over the final array dimension.
1423 * However, if those final iterators have only a single iteration,
1424 * we try to tile earlier iterators instead.
1426 static __isl_give isl_union_map
*access_schedule(struct cuda_gen
*gen
,
1427 __isl_take isl_set
*access
, struct cuda_array_ref_group
*group
)
1431 isl_union_map
*usched
;
1434 unsigned nvar
= isl_set_dim(access
, isl_dim_set
);
1438 sched
= shift_access(access
, group
);
1440 n_tile
= gen
->n_block
;
1441 if (n_tile
> nvar
) {
1443 sched
= isl_map_insert_dims(sched
,
1444 isl_dim_out
, 0, n_tile
- nvar
);
1445 for (i
= 0; i
< n_tile
- nvar
; ++i
)
1446 sched
= isl_map_fix_si(sched
, isl_dim_out
, i
, 0);
1450 first
= nvar
- n_tile
;
1452 for (; first
> 0; first
--)
1453 if (!isl_map_plain_is_fixed(sched
, isl_dim_out
,
1454 first
+ n_tile
- 1, NULL
))
1457 dim
= isl_map_get_space(sched
);
1458 dim
= isl_space_params(dim
);
1459 if (gen
->options
->wrap
)
1460 tiling
= wrap(isl_space_copy(dim
), nvar
, first
,
1461 n_tile
, gen
->block_dim
);
1463 tiling
= tile(isl_space_copy(dim
), nvar
, first
,
1464 n_tile
, gen
->block_dim
);
1465 sched
= isl_map_apply_range(sched
, tiling
);
1467 par
= parametrization(dim
, nvar
+ n_tile
, first
+ n_tile
, n_tile
, "t");
1468 usched
= isl_union_map_from_map(sched
);
1469 usched
= isl_union_map_intersect_range(usched
,
1470 isl_union_set_from_set(par
));
1472 usched
= scale_access_tile_loops(gen
, usched
, nvar
+ n_tile
,
1478 /* Print an access to the element in the global memory copy of the
1479 * given array that corresponds to the element described by "pma".
1480 * of the original array.
1481 * The copy in global memory has been linearized, so we need to take
1482 * the array size into account.
1484 static void print_global_index(FILE *out
,
1485 struct cuda_array_info
*array
, __isl_keep isl_pw_multi_aff
*pma
,
1486 __isl_keep isl_set
*domain
)
1489 isl_ctx
*ctx
= isl_pw_multi_aff_get_ctx(pma
);
1492 if (cuda_array_is_scalar(array
)) {
1493 fprintf(out
, "*%s", array
->name
);
1497 fprintf(out
, "%s[", array
->name
);
1498 prn
= isl_printer_to_file(ctx
, out
);
1499 prn
= isl_printer_set_output_format(prn
, ISL_FORMAT_C
);
1500 for (i
= 0; i
+ 1 < array
->n_index
; ++i
)
1501 prn
= isl_printer_print_str(prn
, "(");
1502 for (i
= 0; i
< array
->n_index
; ++i
) {
1503 isl_pw_aff
*pa
= isl_pw_multi_aff_get_pw_aff(pma
, i
);
1504 pa
= isl_pw_aff_coalesce(pa
);
1505 pa
= isl_pw_aff_gist(pa
, isl_set_copy(domain
));
1507 prn
= isl_printer_print_str(prn
, ") * (");
1508 prn
= isl_printer_print_pw_aff(prn
,
1509 array
->local_bound
[i
]);
1510 prn
= isl_printer_print_str(prn
, ") + ");
1512 prn
= isl_printer_print_pw_aff(prn
, pa
);
1513 isl_pw_aff_free(pa
);
1515 isl_printer_free(prn
);
1519 /* Given an index expression into a tile of an array, adjust the expression
1520 * to a shift of the tile to the origin
1521 * (based on the lower bounds in array->shared_bound).
1522 * If the index is strided, then we first add
1523 * bound->shift and divide by bound->stride.
1525 static __isl_give isl_pw_aff
*shift_index(__isl_take isl_pw_aff
*pa
,
1526 struct cuda_array_info
*array
,
1527 struct cuda_array_bound
*bound
, __isl_take isl_set
*domain
)
1534 shift
= bound
->shift
;
1535 shift
= isl_aff_copy(shift
);
1536 shift
= isl_aff_project_domain_on_params(shift
);
1537 shift
= isl_aff_align_params(shift
, isl_pw_aff_get_space(pa
));
1538 tmp
= isl_pw_aff_alloc(isl_set_copy(domain
), shift
);
1539 pa
= isl_pw_aff_add(pa
, tmp
);
1540 pa
= isl_pw_aff_scale_down(pa
, bound
->stride
);
1543 lb
= isl_aff_copy(bound
->lb
);
1544 lb
= isl_aff_project_domain_on_params(lb
);
1546 lb
= isl_aff_align_params(lb
, isl_pw_aff_get_space(pa
));
1548 tmp
= isl_pw_aff_alloc(isl_set_copy(domain
), lb
);
1549 pa
= isl_pw_aff_sub(pa
, tmp
);
1550 pa
= isl_pw_aff_coalesce(pa
);
1551 pa
= isl_pw_aff_gist(pa
, domain
);
1556 /* Print an access to the element in the private/shared memory copy of the
1557 * given array reference group that corresponds to the element described
1558 * by "pma" of the original array.
1559 * Since the array in private/shared memory is just a shifted copy of part
1560 * of the original array, we simply need to subtract the lower bound,
1561 * which was computed in can_tile_for_shared_memory.
1562 * If any of the indices is strided, then we first add
1563 * bounds[i].shift and divide by bounds[i].stride.
1565 static void print_local_index(FILE *out
,
1566 struct cuda_array_ref_group
*group
, struct cuda_array_bound
*bounds
,
1567 __isl_keep isl_pw_multi_aff
*pma
, __isl_keep isl_set
*domain
)
1570 isl_ctx
*ctx
= isl_pw_multi_aff_get_ctx(pma
);
1572 struct cuda_array_info
*array
= group
->array
;
1574 prn
= isl_printer_to_file(ctx
, out
);
1575 prn
= isl_printer_set_output_format(prn
, ISL_FORMAT_C
);
1577 prn
= print_array_name(prn
, group
);
1578 for (i
= 0; i
< array
->n_index
; ++i
) {
1579 isl_pw_aff
*pa
= isl_pw_multi_aff_get_pw_aff(pma
, i
);
1581 pa
= shift_index(pa
, array
, &bounds
[i
], isl_set_copy(domain
));
1584 prn
= isl_printer_print_pw_aff(prn
, pa
);
1586 isl_pw_aff_free(pa
);
1589 isl_printer_free(prn
);
1592 /* This function is called for each leaf in the clast of the code
1593 * for copying to or from shared/private memory.
1594 * The statement name is {read,write}_{shared,private}_<array>.
1596 * The schedule iterates over the array elements, so we can use
1597 * the domain of copy_sched at the current scheduling position
1598 * as the index of the array.
1600 static void print_copy_statement(struct clast_printer_info
*code
,
1601 struct clast_user_stmt
*u
)
1603 struct cuda_gen
*gen
= code
->user
;
1606 struct cuda_array_ref_group
*group
= gen
->copy_group
;
1607 struct cuda_array_bound
*bounds
= gen
->copy_bound
;
1613 isl_pw_multi_aff
*pma
;
1616 read
= !strncmp(u
->statement
->name
, "read", 4);
1618 domain
= extract_host_domain(u
);
1621 sched
= isl_map_copy(gen
->copy_sched
);
1622 sched
= isl_map_reverse(sched
);
1623 sched
= isl_map_intersect_domain(sched
, domain
);
1624 n_in
= isl_map_dim(sched
, isl_dim_in
);
1625 n_out
= isl_map_dim(sched
, isl_dim_out
);
1626 dim
= isl_map_get_space(sched
);
1627 dim
= isl_space_drop_dims(dim
, isl_dim_in
, 0, n_in
);
1628 dim
= isl_space_drop_dims(dim
, isl_dim_out
, 0, n_out
);
1629 param
= parametrization(dim
, n_in
, 0, n_in
, "c");
1630 sched
= isl_map_align_params(sched
, isl_set_get_space(param
));
1631 sched
= isl_map_intersect_domain(sched
, param
);
1632 index
= isl_map_range(sched
);
1633 domain
= isl_set_copy(index
);
1634 pma
= isl_pw_multi_aff_from_set(index
);
1635 pma
= isl_pw_multi_aff_coalesce(pma
);
1636 domain
= isl_set_params(domain
);
1638 print_indent(code
->dst
, code
->indent
);
1640 print_local_index(code
->dst
, group
, bounds
, pma
, domain
);
1641 fprintf(code
->dst
, " = ");
1642 print_global_index(code
->dst
, group
->array
, pma
, domain
);
1644 print_global_index(code
->dst
, group
->array
, pma
, domain
);
1645 fprintf(code
->dst
, " = ");
1646 print_local_index(code
->dst
, group
, bounds
, pma
, domain
);
1648 fprintf(code
->dst
, ";\n");
1650 isl_pw_multi_aff_free(pma
);
1651 isl_set_free(domain
);
1654 static void print_shared_access(struct cuda_gen
*gen
,
1655 __isl_keep isl_set
*shared_domain
, __isl_take isl_set
*access
,
1656 const char *type
, struct cuda_array_ref_group
*group
)
1658 const char *array_name
;
1661 isl_union_map
*sched
;
1662 unsigned nvar
= isl_set_dim(access
, isl_dim_set
);
1665 ctx
= isl_set_get_ctx(access
);
1666 array_name
= isl_set_get_tuple_name(access
);
1667 name
= isl_alloc_array(ctx
, char,
1668 strlen(type
) + sizeof("_shared_") + strlen(array_name
) + 20);
1669 if (group
->array
->n_group
> 1)
1670 sprintf(name
, "%s_shared_%s_%d", type
, array_name
, group
->nr
);
1672 sprintf(name
, "%s_shared_%s", type
, array_name
);
1673 access
= isl_set_set_tuple_name(access
, name
);
1676 sched
= access_schedule(gen
, access
, group
);
1678 n_tile
= gen
->n_block
;
1682 gen
->copy_sched
= isl_map_from_union_map(isl_union_map_copy(sched
));
1683 gen
->copy_group
= group
;
1684 gen
->copy_bound
= group
->shared_bound
;
1686 print_shared_body(gen
, shared_domain
, sched
, nvar
+ n_tile
,
1687 &print_copy_statement
, -1);
1689 isl_union_map_free(sched
);
1690 isl_map_free(gen
->copy_sched
);
1693 /* Return the union of all read (read = 1) and/or write (write = 1)
1694 * access relations in the group.
1696 static __isl_give isl_union_map
*group_access_relation(
1697 struct cuda_array_ref_group
*group
, int read
, int write
)
1700 isl_union_map
*access
;
1702 access
= isl_union_map_empty(isl_map_get_space(group
->access
));
1703 for (i
= 0; i
< group
->n_ref
; ++i
) {
1706 if (!((read
&& group
->refs
[i
]->read
) ||
1707 (write
&& group
->refs
[i
]->write
)))
1709 map_i
= isl_map_copy(group
->refs
[i
]->access
);
1710 access
= isl_union_map_union(access
,
1711 isl_union_map_from_map(map_i
));
1717 /* Check that none of the shared memory tiles involve any strides.
1719 static int no_strides(struct cuda_array_ref_group
*group
)
1722 int n_index
= group
->array
->n_index
;
1724 for (i
= 0; i
< n_index
; ++i
)
1725 if (group
->shared_bound
[i
].shift
)
1731 /* Return a set containing the values of the given index i
1732 * of the elements in the array tile in global memory that corresponds
1733 * to the shared memory copy.
1734 * In particular, if a is the index, we return a set with constraints
1736 * tile_offset <= a <= tile_offset + tile_size - 1
1740 * 0 <= a <= array_size - 1
1743 static __isl_give isl_set
*group_tile_dim(struct cuda_array_ref_group
*group
,
1746 isl_basic_set
*tile
;
1749 isl_local_space
*ls
;
1754 aff
= isl_aff_copy(group
->shared_bound
[i
].lb
);
1755 aff
= isl_aff_add_dims(aff
, isl_dim_in
, 1);
1756 ls
= isl_aff_get_domain_local_space(aff
);
1757 aff
= isl_aff_neg(aff
);
1758 aff
= isl_aff_add_coefficient_si(aff
, isl_dim_in
, 0, 1);
1759 c
= isl_inequality_from_aff(isl_aff_copy(aff
));
1760 tile
= isl_basic_set_from_constraint(c
);
1762 aff
= isl_aff_neg(aff
);
1763 aff
= isl_aff_add_constant(aff
, group
->shared_bound
[i
].size
);
1764 aff
= isl_aff_add_constant_si(aff
, -1);
1765 c
= isl_inequality_from_aff(aff
);
1766 tile
= isl_basic_set_add_constraint(tile
, c
);
1768 aff
= isl_aff_zero_on_domain(ls
);
1769 aff
= isl_aff_add_coefficient_si(aff
, isl_dim_in
, 0, 1);
1770 c
= isl_inequality_from_aff(aff
);
1771 tile
= isl_basic_set_add_constraint(tile
, c
);
1773 bound
= isl_pw_aff_copy(group
->array
->bound
[i
]);
1774 bound
= isl_pw_aff_add_dims(bound
, isl_dim_in
, 1);
1775 ls
= isl_local_space_from_space(isl_pw_aff_get_domain_space(bound
));
1776 aff
= isl_aff_zero_on_domain(ls
);
1777 aff
= isl_aff_add_coefficient_si(aff
, isl_dim_in
, 0, 1);
1778 aff
= isl_aff_add_constant_si(aff
, 1);
1779 dom
= isl_pw_aff_domain(isl_pw_aff_copy(bound
));
1781 tile_set
= isl_pw_aff_ge_set(bound
, isl_pw_aff_alloc(dom
, aff
));
1782 tile_set
= isl_set_align_params(tile_set
, isl_basic_set_get_space(tile
));
1783 tile_set
= isl_set_intersect(tile_set
, isl_set_from_basic_set(tile
));
1788 /* Return a set containing the elements in the array tile in
1789 * global memory that corresponds to the shared memory copy.
1791 static __isl_give isl_set
*group_tile(struct cuda_array_ref_group
*group
)
1794 int n_index
= group
->array
->n_index
;
1797 tile
= group_tile_dim(group
, 0);
1798 for (i
= 1; i
< n_index
; ++i
) {
1801 tile_i
= group_tile_dim(group
, i
);
1802 tile
= isl_set_flat_product(tile
, tile_i
);
1805 tile
= isl_set_set_tuple_name(tile
, group
->array
->name
);
1810 /* Print code for reading into or writing from shared memory
1811 * the given array reference group.
1813 * sched maps the original iteration domains to the shared memory tile loops.
1815 * If we are performing a read from global memory to shared memory,
1816 * if the array involved is not a scalar and if the definition of the
1817 * shared memory tiles does not involve any strides, then we copy
1818 * the entire tile to shared memory. This may result in some extra
1819 * elements getting copied, but it should lead to simpler code
1820 * (which means that fewer registers may be needed) and less divergence.
1822 * Otherwise, we only copy the elements that will be read or have been written
1825 * Note that the absence of stride requirement can easily be lifted.
1826 * We would just need to add constraints of the form
1828 * shift + a = stride * alpha
1830 static int print_group_shared_accesses(struct cuda_gen
*gen
,
1831 struct cuda_array_ref_group
*group
, const char *type
,
1832 __isl_keep isl_set
*shared_domain
, __isl_keep isl_union_map
*sched
)
1835 isl_union_map
*access
;
1836 isl_union_set
*uset
;
1837 isl_set
*access_set
;
1839 if (group
->private_bound
)
1841 if (!group
->shared_bound
)
1844 read
= !strcmp(type
, "read");
1846 access
= group_access_relation(group
, read
, !read
);
1847 access
= isl_union_map_apply_domain(access
, isl_union_map_copy(sched
));
1848 uset
= isl_union_map_range(access
);
1850 if (isl_union_set_is_empty(uset
)) {
1851 isl_union_set_free(uset
);
1855 if (read
&& group
->array
->n_index
> 0 && no_strides(group
)) {
1856 isl_union_set_free(uset
);
1857 access_set
= group_tile(group
);
1858 print_shared_access(gen
, shared_domain
, access_set
,
1863 access_set
= isl_set_from_union_set(uset
);
1864 access_set
= isl_set_coalesce(access_set
);
1866 print_shared_access(gen
, shared_domain
, access_set
, type
, group
);
1871 /* Print code for reading into or writing from shared memory at
1872 * the given level (-1 for innermost).
1874 * If we are not printing at the innermost level, then the dimensionality
1875 * of shared_domain may be smaller than gen->shared_len.
1876 * As the rest of the code assumes that the domain of access has
1877 * gen->shared_len dimensions, we therefore may need to embed this domain
1878 * in a higher dimensional space after intersection with shared_domain.
1880 static void print_shared_accesses(struct cuda_gen
*gen
,
1881 __isl_keep isl_set
*shared_domain
, __isl_keep isl_union_map
*access
,
1882 const char *type
, int level
)
1888 int shared_len
= isl_set_dim(shared_domain
, isl_dim_set
);
1890 isl_union_map
*sched
;
1892 shared_domain
= isl_set_copy(shared_domain
);
1893 sched
= isl_union_map_copy(gen
->tiled_sched
);
1894 dim
= isl_union_map_get_space(sched
);
1895 proj
= projection(dim
, gen
->tiled_len
, shared_len
);
1896 sched
= isl_union_map_apply_range(sched
, isl_union_map_from_map(proj
));
1897 sched
= isl_union_map_intersect_range(sched
,
1898 isl_union_set_from_set(isl_set_copy(shared_domain
)));
1899 if (shared_len
!= gen
->shared_len
) {
1900 dim
= isl_union_map_get_space(sched
);
1901 proj
= projection(dim
, gen
->shared_len
, shared_len
);
1902 proj
= isl_map_reverse(proj
);
1903 shared_domain
= isl_set_apply(shared_domain
,
1904 isl_map_copy(proj
));
1905 sched
= isl_union_map_apply_range(sched
,
1906 isl_union_map_from_map(proj
));
1909 dim
= isl_union_map_get_space(sched
);
1910 par
= parametrization(dim
, gen
->shared_len
, 0, gen
->shared_len
, "g");
1911 sched
= isl_union_map_intersect_range(sched
,
1912 isl_union_set_from_set(par
));
1914 for (i
= 0; i
< gen
->n_array
; ++i
) {
1915 struct cuda_array_info
*array
= &gen
->array
[i
];
1917 for (j
= 0; j
< array
->n_group
; ++j
) {
1918 if (array
->groups
[j
]->print_shared_level
!= level
)
1921 if (print_group_shared_accesses(gen
, array
->groups
[j
],
1922 type
, shared_domain
, sched
))
1927 isl_union_map_free(sched
);
1928 isl_set_free(shared_domain
);
1931 print_indent(gen
->cuda
.kernel_c
, gen
->kernel_code
.indent
);
1932 fprintf(gen
->cuda
.kernel_c
, "__syncthreads();\n");
1936 /* This function is called for each access to an array in some statement
1937 * in the original code.
1938 * Replace that access by an access to shared or (linearized) global memory.
1939 * Since the array in shared memory is just
1940 * a shifted copy of part of the original array, we simply need
1941 * to subtract the lower bound, which was computed
1942 * in can_tile_for_shared_memory.
1943 * If any of the indices is strided, then we first add
1944 * shared_bound[i].shift and divide by shared_bound[i].stride.
1946 * If the given array is accessed directly from global memory,
1947 * we don't need to perform any shifting and simply simplify
1948 * expression in the context of the domain instead.
1950 * If the array space (range of access) has no name, then we are
1951 * accessing an iterator in the original program.
1953 static void print_access(struct cuda_gen
*gen
, __isl_take isl_map
*access
,
1959 struct cuda_array_info
*array
= NULL
;
1961 isl_pw_multi_aff
*pma
;
1964 struct cuda_array_bound
*bounds
= NULL
;
1966 access
= isl_map_align_params(access
,
1967 isl_set_get_space(gen
->stmt_domain
));
1969 data_set
= isl_set_apply(isl_set_copy(gen
->stmt_domain
), access
);
1971 name
= isl_set_get_tuple_name(data_set
);
1974 fprintf(gen
->cuda
.kernel_c
, "(");
1976 struct cuda_array_ref_group
*group
;
1978 for (i
= 0; i
< gen
->n_array
; ++i
) {
1979 if (strcmp(name
, gen
->array
[i
].name
))
1981 array
= &gen
->array
[i
];
1984 group
= array
->groups
[group_nr
];
1985 bounds
= group
->private_bound
;
1987 bounds
= group
->shared_bound
;
1989 if (!bounds
&& cuda_array_is_scalar(array
) && !array
->read_only
)
1990 fprintf(gen
->cuda
.kernel_c
, "*");
1991 prn
= isl_printer_to_file(gen
->ctx
, gen
->cuda
.kernel_c
);
1992 prn
= print_array_name(prn
, group
);
1993 isl_printer_free(prn
);
1995 if (cuda_array_is_scalar(array
)) {
1996 isl_set_free(data_set
);
2000 fprintf(gen
->cuda
.kernel_c
, "[");
2004 n_index
= isl_set_dim(data_set
, isl_dim_set
);
2005 pma
= isl_pw_multi_aff_from_set(data_set
);
2006 pma
= isl_pw_multi_aff_coalesce(pma
);
2008 prn
= isl_printer_to_file(gen
->ctx
, gen
->cuda
.kernel_c
);
2009 prn
= isl_printer_set_output_format(prn
, ISL_FORMAT_C
);
2012 for (i
= 0; i
+ 1 < n_index
; ++i
)
2013 prn
= isl_printer_print_str(prn
, "(");
2015 for (i
= 0; i
< n_index
; ++i
) {
2018 index
= isl_pw_multi_aff_get_pw_aff(pma
, i
);
2021 prn
= isl_printer_print_pw_aff(prn
, index
);
2022 isl_pw_aff_free(index
);
2026 domain
= isl_set_copy(gen
->stmt_domain
);
2027 domain
= isl_set_params(domain
);
2029 index
= isl_pw_aff_coalesce(index
);
2030 index
= isl_pw_aff_gist(index
, domain
);
2032 index
= shift_index(index
, array
, &bounds
[i
], domain
);
2036 prn
= isl_printer_print_str(prn
, ") * (");
2037 prn
= isl_printer_print_pw_aff(prn
,
2038 array
->local_bound
[i
]);
2039 prn
= isl_printer_print_str(prn
, ") + ");
2041 prn
= isl_printer_print_str(prn
, "][");
2043 prn
= isl_printer_print_pw_aff(prn
, index
);
2044 isl_pw_aff_free(index
);
2047 prn
= isl_printer_print_str(prn
, ")");
2049 prn
= isl_printer_print_str(prn
, "]");
2050 isl_printer_free(prn
);
2052 isl_pw_multi_aff_free(pma
);
2055 struct cuda_access_print_info
{
2056 struct cuda_gen
*gen
;
2057 struct cuda_stmt_access
*access
;
2060 /* To print the cuda accesses we walk the list of cuda accesses simultaneously
2061 * with the pet printer. This means that whenever the pet printer prints a
2062 * pet access expression we have the corresponding cuda access available and can
2063 * print the modified access.
2065 static void print_cuda_access(struct pet_expr
*expr
, void *usr
)
2067 struct cuda_access_print_info
*info
=
2068 (struct cuda_access_print_info
*) usr
;
2069 print_access(info
->gen
, isl_map_copy(info
->access
->access
),
2070 info
->access
->group
);
2071 info
->access
= info
->access
->next
;
2074 static void print_stmt_body(struct cuda_gen
*gen
,
2075 FILE *out
, struct cuda_stmt
*stmt
)
2077 struct cuda_access_print_info info
;
2080 info
.access
= stmt
->accesses
;
2082 print_pet_expr(out
, stmt
->body
, print_cuda_access
, &info
);
2083 fprintf(out
, ";\n");
2086 /* This function is called for each leaf in the innermost clast,
2087 * i.e., for each statement.
2088 * We print the statement body, simplifying the accesses based
2091 static void print_statement(struct clast_printer_info
*code
,
2092 struct clast_user_stmt
*u
)
2094 struct cuda_gen
*gen
= code
->user
;
2097 isl_set
*stmt_domain
;
2098 isl_union_map
*stmt_sched
;
2099 isl_union_set
*uset
;
2101 struct cuda_stmt
*stmt
;
2103 nr
= atoi(u
->statement
->name
+ 2);
2104 stmt
= &gen
->stmts
[nr
];
2106 stmt_domain
= extract_host_domain(u
);
2108 stmt_sched
= isl_union_map_intersect_range(
2109 isl_union_map_copy(gen
->local_sched
),
2110 isl_union_set_from_set(extend(stmt_domain
,
2111 gen
->thread_tiled_len
)));
2112 dim
= isl_union_map_get_space(stmt_sched
);
2113 par
= parametrization(dim
, gen
->thread_tiled_len
, 0,
2114 gen
->thread_tiled_len
, "c");
2115 stmt_sched
= isl_union_map_intersect_range(stmt_sched
,
2116 isl_union_set_from_set(par
));
2118 uset
= isl_union_map_domain(stmt_sched
);
2119 dim
= isl_union_set_get_space(uset
);
2120 dim
= isl_space_add_dims(dim
, isl_dim_set
,
2121 isl_set_dim(stmt
->domain
, isl_dim_set
));
2122 dim
= isl_space_set_tuple_name(dim
, isl_dim_set
, u
->statement
->name
);
2123 gen
->stmt_domain
= isl_union_set_extract_set(uset
, dim
);
2124 isl_union_set_free(uset
);
2126 print_indent(code
->dst
, code
->indent
);
2127 print_stmt_body(gen
, code
->dst
, stmt
);
2129 isl_set_free(gen
->stmt_domain
);
2132 static void print_private_access(struct cuda_gen
*gen
,
2133 __isl_keep isl_set
*shared_domain
, __isl_take isl_set
*access
,
2134 const char *type
, struct cuda_array_ref_group
*group
)
2136 const char *array_name
;
2139 unsigned nvar
= isl_set_dim(access
, isl_dim_set
);
2140 isl_union_map
*usched
;
2142 if (isl_set_fast_is_empty(access
)) {
2143 isl_set_free(access
);
2147 ctx
= isl_set_get_ctx(access
);
2148 array_name
= isl_set_get_tuple_name(access
);
2149 name
= isl_alloc_array(ctx
, char,
2150 strlen(type
) + sizeof("_private_") + strlen(array_name
) + 20);
2151 if (group
->array
->n_group
> 1)
2152 sprintf(name
, "%s_private_%s_%d", type
, array_name
, group
->nr
);
2154 sprintf(name
, "%s_private_%s", type
, array_name
);
2155 access
= isl_set_set_tuple_name(access
, name
);
2158 gen
->copy_sched
= shift_access(access
, group
);
2159 gen
->copy_group
= group
;
2160 gen
->copy_bound
= group
->private_bound
;
2162 usched
= isl_union_map_from_map(isl_map_copy(gen
->copy_sched
));
2163 print_shared_body(gen
, shared_domain
, usched
, nvar
,
2164 &print_copy_statement
, 1);
2165 isl_union_map_free(usched
);
2167 isl_map_free(gen
->copy_sched
);
2170 /* Print code for reading into or writing from private memory
2171 * the given array reference group.
2173 * sched maps the original iteration domains to the shared memory tile loops.
2175 static void print_group_private_accesses(struct cuda_gen
*gen
,
2176 struct cuda_array_ref_group
*group
,
2177 const char *type
, __isl_keep isl_set
*shared_domain
,
2178 unsigned first_shared
, int shared_len
, __isl_keep isl_union_map
*sched
)
2181 isl_union_map
*access
;
2182 isl_union_set
*uset
;
2183 isl_set
*access_set
;
2185 if (!group
->private_bound
)
2188 read
= !strcmp(type
, "read");
2190 access
= group_access_relation(group
, read
, !read
);
2191 access
= isl_union_map_apply_domain(access
, isl_union_map_copy(sched
));
2192 access
= isl_union_map_intersect(access
,
2193 isl_union_map_copy(gen
->private_access
));
2194 uset
= isl_union_map_range(access
);
2196 if (isl_union_set_is_empty(uset
)) {
2197 isl_union_set_free(uset
);
2201 access_set
= isl_set_from_union_set(uset
);
2202 access_set
= isl_set_coalesce(access_set
);
2203 access_set
= isl_set_eliminate(access_set
, isl_dim_param
,
2204 first_shared
+ shared_len
,
2205 gen
->shared_len
- shared_len
);
2207 print_private_access(gen
, shared_domain
, access_set
, type
, group
);
2210 /* Print code for reading into or writing from private memory at
2211 * the given level (-1 for innermost).
2213 * If we are not printing at the innermost level, then the dimensionality
2214 * of shared_domain may be smaller than gen->shared_len.
2215 * As the rest of the code assumes that the domain of access has
2216 * gen->shared_len dimensions, we therefore may need to embed this domain
2217 * in a higher dimensional space after intersection with shared_domain.
2219 * This code is very similar to print_shared_accesses.
2220 * The main difference is that we to take into account gen->private_access.
2222 static void print_private_accesses(struct cuda_gen
*gen
,
2223 __isl_keep isl_set
*shared_domain
, __isl_keep isl_union_map
*access
,
2224 const char *type
, int level
)
2229 int shared_len
= isl_set_dim(shared_domain
, isl_dim_set
);
2230 unsigned first_shared
;
2231 isl_union_map
*sched
;
2233 shared_domain
= isl_set_copy(shared_domain
);
2234 sched
= isl_union_map_copy(gen
->tiled_sched
);
2235 dim
= isl_union_map_get_space(sched
);
2236 first_shared
= isl_space_dim(dim
, isl_dim_param
);
2237 proj
= projection(dim
, gen
->tiled_len
, shared_len
);
2238 sched
= isl_union_map_apply_range(sched
, isl_union_map_from_map(proj
));
2239 sched
= isl_union_map_intersect_range(sched
,
2240 isl_union_set_from_set(isl_set_copy(shared_domain
)));
2241 if (shared_len
!= gen
->shared_len
) {
2242 dim
= isl_union_map_get_space(sched
);
2243 proj
= projection(dim
, gen
->shared_len
, shared_len
);
2244 proj
= isl_map_reverse(proj
);
2245 shared_domain
= isl_set_apply(shared_domain
,
2246 isl_map_copy(proj
));
2247 sched
= isl_union_map_apply_range(sched
,
2248 isl_union_map_from_map(proj
));
2251 for (i
= 0; i
< gen
->n_array
; ++i
) {
2252 struct cuda_array_info
*array
= &gen
->array
[i
];
2254 for (j
= 0; j
< array
->n_group
; ++j
) {
2255 if (array
->groups
[j
]->print_shared_level
!= level
)
2258 print_group_private_accesses(gen
, array
->groups
[j
],
2259 type
, shared_domain
,
2260 first_shared
, shared_len
, sched
);
2264 isl_union_map_free(sched
);
2265 isl_set_free(shared_domain
);
2268 /* Set unroll[j] if the input dimension j is involved in
2269 * the index expression represented by bmap.
2271 static int check_unroll(__isl_take isl_basic_map
*bmap
, void *user
)
2274 int n_in
= isl_basic_map_dim(bmap
, isl_dim_in
);
2275 int n_out
= isl_basic_map_dim(bmap
, isl_dim_out
);
2278 for (i
= 0; i
< n_out
; ++i
) {
2282 ok
= isl_basic_map_has_defining_equality(bmap
,
2283 isl_dim_out
, i
, &c
);
2285 for (j
= 0; j
< n_in
; ++j
)
2286 if (isl_constraint_involves_dims(c
, isl_dim_in
, j
, 1))
2288 isl_constraint_free(c
);
2291 isl_basic_map_free(bmap
);
2295 /* Given an array pos mapping input dimensions to the corresponding
2296 * output dimension, construct the corresponding map.
2298 static __isl_give isl_map
*permutation(__isl_take isl_space
*dim
,
2303 isl_basic_map
*bmap
;
2304 isl_local_space
*ls
;
2306 dim
= isl_space_add_dims(dim
, isl_dim_in
, len
);
2307 dim
= isl_space_add_dims(dim
, isl_dim_out
, len
);
2308 bmap
= isl_basic_map_universe(isl_space_copy(dim
));
2309 ls
= isl_local_space_from_space(dim
);
2311 for (i
= 0; i
< len
; ++i
) {
2312 c
= isl_equality_alloc(isl_local_space_copy(ls
));
2313 isl_constraint_set_coefficient_si(c
, isl_dim_in
, i
, -1);
2314 isl_constraint_set_coefficient_si(c
, isl_dim_out
, pos
[i
], 1);
2315 bmap
= isl_basic_map_add_constraint(bmap
, c
);
2317 isl_local_space_free(ls
);
2319 return isl_map_from_basic_map(bmap
);
2322 /* Find all loops involved in any of the index expressions for any of
2323 * the private accesses, move them innermost and then mark them as
2324 * requiring unrolling by setting gen->first_unroll.
2325 * The loops involved should all be parallel because of the checks
2326 * we performed in check_private_group_access. Moving them innermost
2327 * is therefore a valid transformation.
2329 static __isl_give isl_union_map
*interchange_for_unroll(struct cuda_gen
*gen
,
2330 __isl_take isl_union_map
*sched
)
2333 int unroll
[gen
->thread_tiled_len
];
2334 int perm
[gen
->thread_tiled_len
];
2337 int len
= gen
->shared_len
+ gen
->n_parallel
+ gen
->n_block
;
2339 gen
->first_unroll
= -1;
2341 for (i
= 0; i
< gen
->thread_tiled_len
; ++i
)
2343 for (i
= 0; i
< gen
->n_array
; ++i
) {
2344 struct cuda_array_info
*array
= &gen
->array
[i
];
2346 for (j
= 0; j
< array
->n_group
; ++j
) {
2347 isl_union_map
*access
;
2350 if (!array
->groups
[j
]->private_bound
)
2353 access
= group_access_relation(array
->groups
[j
], 1, 1);
2354 access
= isl_union_map_apply_domain(access
,
2355 isl_union_map_copy(sched
));
2357 acc
= isl_map_from_union_map(access
);
2358 isl_map_foreach_basic_map(acc
, &check_unroll
, unroll
);
2364 for (i
= 0; i
< gen
->shared_len
; ++i
)
2368 for (i
= gen
->shared_len
; i
< len
; ++i
)
2375 for (i
= len
; i
< gen
->thread_tiled_len
; ++i
)
2380 for (i
= 0; i
< gen
->thread_tiled_len
; ++i
)
2383 gen
->first_unroll
= 1 + j
;
2384 for (i
= 0; i
< len
; ++i
)
2388 dim
= isl_union_map_get_space(sched
);
2389 permute
= permutation(dim
, perm
, gen
->thread_tiled_len
);
2390 sched
= isl_union_map_apply_range(sched
,
2391 isl_union_map_from_map(permute
));
2396 /* This function is called for each leaf in the clast of the kernel code.
2397 * We first specialize the schedule to the site of the leaf and
2398 * print code for reading into shared memory, performing the actual
2399 * computations and writing from shared memory, with the required
2402 static void print_kernel_user(struct clast_printer_info
*code
,
2403 struct clast_user_stmt
*u
)
2405 struct cuda_gen
*gen
= code
->user
;
2406 isl_set
*shared_domain
;
2408 shared_domain
= extract_entire_host_domain(&u
->stmt
);
2410 print_shared_accesses(gen
, shared_domain
, gen
->read
, "read", -1);
2412 print_private_accesses(gen
, shared_domain
, gen
->read
, "read", -1);
2414 print_shared_body(gen
, shared_domain
, gen
->local_sched
,
2415 gen
->thread_tiled_len
, &print_statement
,
2418 print_private_accesses(gen
, shared_domain
, gen
->write
, "write", -1);
2420 print_indent(gen
->cuda
.kernel_c
, gen
->kernel_code
.indent
);
2421 fprintf(gen
->cuda
.kernel_c
, "__syncthreads();\n");
2423 print_shared_accesses(gen
, shared_domain
, gen
->write
, "write", -1);
2425 isl_set_free(shared_domain
);
2428 /* Check if we need to perform any copying to shared memory at this level
2429 * and if so, print the copying instructions.
2430 * Any array for which we are allowed to print copying instructions at
2431 * this level, but haven't done so already, is printed.
2433 static void copy_to_local(struct cuda_gen
*gen
, __isl_keep isl_set
*domain
)
2439 level
= isl_set_dim(domain
, isl_dim_set
);
2441 for (i
= 0; i
< gen
->n_array
; ++i
) {
2442 struct cuda_array_info
*array
= &gen
->array
[i
];
2444 for (j
= 0; j
< array
->n_group
; ++j
) {
2445 if (array
->groups
[j
]->print_shared_level
>= 0)
2447 if (array
->groups
[j
]->last_shared
>= level
)
2449 array
->groups
[j
]->print_shared_level
= level
;
2455 print_shared_accesses(gen
, domain
, gen
->read
, "read", level
);
2456 print_private_accesses(gen
, domain
, gen
->read
, "read", level
);
2461 /* This function is called for each for loop in the clast,
2462 * right after the opening brace has been printed.
2464 * Print copying instructions to shared or private memory if needed.
2466 static void print_kernel_for_head(struct clast_printer_info
*code
,
2467 struct clast_for
*f
)
2469 struct cuda_gen
*gen
= code
->user
;
2472 domain
= isl_set_from_cloog_domain(cloog_domain_copy(f
->domain
));
2473 copy_to_local(gen
, domain
);
2475 isl_set_free(domain
);
2478 /* Print instructions for copying from shared memory for each array
2479 * for which print_kernel_for_head has added copying instructions
2482 static void copy_from_local(struct cuda_gen
*gen
, __isl_keep isl_set
*domain
)
2488 level
= isl_set_dim(domain
, isl_dim_set
);
2490 for (i
= 0; i
< gen
->n_array
; ++i
) {
2491 struct cuda_array_info
*array
= &gen
->array
[i
];
2493 for (j
= 0; j
< array
->n_group
; ++j
) {
2494 if (array
->groups
[j
]->print_shared_level
!= level
)
2504 print_private_accesses(gen
, domain
, gen
->write
, "write", level
);
2505 print_shared_accesses(gen
, domain
, gen
->write
, "write", level
);
2509 /* This function is called for each for loop in the clast,
2510 * right before the closing brace is printed.
2512 * Print copying instructions from shared or private memory if needed.
2514 static void print_kernel_for_foot(struct clast_printer_info
*code
,
2515 struct clast_for
*f
)
2517 struct cuda_gen
*gen
= code
->user
;
2520 domain
= isl_set_from_cloog_domain(cloog_domain_copy(f
->domain
));
2521 copy_from_local(gen
, domain
);
2523 isl_set_free(domain
);
2526 /* Use CLooG to generate code for the outer gen->shared_first loops
2527 * of the local schedule "sched".
2528 * The pretty printing of this code is handled by print_clast,
2529 * which calls print_kernel_user for each iteration of the shared tile loops.
2531 static void print_cloog_kernel_body(struct cuda_gen
*gen
,
2532 __isl_keep isl_set
*context
, __isl_keep isl_union_map
*sched
)
2535 CloogOptions
*options
;
2536 CloogDomain
*cloog_context
;
2537 CloogUnionDomain
*ud
;
2539 struct clast_stmt
*stmt
;
2542 sched
= isl_union_map_copy(sched
);
2543 sched
= isl_union_map_align_params(sched
, isl_set_get_space(context
));
2545 options
= cloog_options_malloc(gen
->state
);
2546 options
->language
= CLOOG_LANGUAGE_C
;
2547 options
->strides
= 1;
2549 options
->stop
= gen
->shared_len
;
2550 options
->f
= gen
->tiled_len
;
2551 options
->l
= gen
->tiled_len
;
2552 options
->save_domains
= 1;
2553 options
->noscalars
= 1;
2555 ud
= cloog_union_domain_from_isl_union_map(sched
);
2556 for (i
= 0; i
< gen
->shared_len
; ++i
) {
2557 snprintf(name
, sizeof(name
), "g%d", i
);
2558 ud
= cloog_union_domain_set_name(ud
, CLOOG_SCAT
, i
, name
);
2560 cloog_context
= cloog_domain_from_isl_set(isl_set_copy(context
));
2561 input
= cloog_input_alloc(cloog_context
, ud
);
2563 stmt
= cloog_clast_create_from_input(input
, options
);
2565 gen
->kernel_code
.indent
= 4;
2566 gen
->kernel_code
.dst
= gen
->cuda
.kernel_c
;
2567 gen
->kernel_code
.print_user_stmt
= NULL
;
2568 gen
->kernel_code
.print_user_stmt_list
= &print_kernel_user
;
2569 gen
->kernel_code
.print_for_head
= &print_kernel_for_head
;
2570 gen
->kernel_code
.print_for_foot
= &print_kernel_for_foot
;
2571 gen
->kernel_code
.user
= gen
;
2572 copy_to_local(gen
, context
);
2573 print_clast(&gen
->kernel_code
, stmt
);
2574 copy_from_local(gen
, context
);
2576 cloog_clast_free(stmt
);
2577 cloog_options_free(options
);
2580 static void print_kernel_iterators(struct cuda_gen
*gen
)
2583 const char *block_dims
[] = { "blockIdx.x", "blockIdx.y" };
2584 const char *thread_dims
[] = { "threadIdx.x", "threadIdx.y",
2587 if (gen
->n_grid
> 0) {
2588 print_indent(gen
->cuda
.kernel_c
, 4);
2589 fprintf(gen
->cuda
.kernel_c
, "int ");
2590 for (i
= 0; i
< gen
->n_grid
; ++i
) {
2592 fprintf(gen
->cuda
.kernel_c
, ", ");
2593 fprintf(gen
->cuda
.kernel_c
, "b%d = %s",
2594 i
, block_dims
[gen
->n_grid
- 1 - i
]);
2596 fprintf(gen
->cuda
.kernel_c
, ";\n");
2599 if (gen
->n_block
> 0) {
2600 print_indent(gen
->cuda
.kernel_c
, 4);
2601 fprintf(gen
->cuda
.kernel_c
, "int ");
2602 for (i
= 0; i
< gen
->n_block
; ++i
) {
2604 fprintf(gen
->cuda
.kernel_c
, ", ");
2605 fprintf(gen
->cuda
.kernel_c
, "t%d = %s",
2606 i
, thread_dims
[gen
->n_block
- 1 - i
]);
2608 fprintf(gen
->cuda
.kernel_c
, ";\n");
2612 static void print_group_shared_array(struct cuda_gen
*gen
,
2613 struct cuda_array_ref_group
*group
)
2616 struct cuda_array_bound
*bounds
;
2619 bounds
= group
->private_bound
;
2621 bounds
= group
->shared_bound
;
2625 print_indent(gen
->cuda
.kernel_c
, 4);
2626 fprintf(gen
->cuda
.kernel_c
, "%s%s ",
2627 group
->private_bound
? "" : "__shared__ ", group
->array
->type
);
2628 p
= isl_printer_to_file(gen
->ctx
, gen
->cuda
.kernel_c
);
2629 p
= print_array_name(p
, group
);
2630 isl_printer_free(p
);
2631 for (j
= 0; j
< group
->array
->n_index
; ++j
) {
2632 fprintf(gen
->cuda
.kernel_c
, "[");
2633 isl_int_print(gen
->cuda
.kernel_c
, bounds
[j
].size
, 0);
2634 fprintf(gen
->cuda
.kernel_c
, "]");
2636 fprintf(gen
->cuda
.kernel_c
, ";\n");
2639 static void print_shared_arrays(struct cuda_gen
*gen
)
2643 for (i
= 0; i
< gen
->n_array
; ++i
) {
2644 struct cuda_array_info
*array
= &gen
->array
[i
];
2646 for (j
= 0; j
< array
->n_group
; ++j
)
2647 print_group_shared_array(gen
, array
->groups
[j
]);
2651 static void print_kernel_body(struct cuda_gen
*gen
,
2652 __isl_keep isl_set
*host_domain
, __isl_keep isl_union_map
*sched
)
2656 context
= isl_set_copy(host_domain
);
2657 context
= parametrize(context
, 0, gen
->tile_first
, "h");
2658 context
= isl_set_project_out(context
, isl_dim_set
, 0, gen
->tile_first
);
2659 context
= add_bounded_parameters(context
,
2660 gen
->n_grid
, gen
->grid_dim
, "b");
2662 print_kernel_iterators(gen
);
2663 print_shared_arrays(gen
);
2665 fprintf(gen
->cuda
.kernel_c
, "\n");
2667 print_cloog_kernel_body(gen
, context
, sched
);
2669 isl_set_free(context
);
2672 /* Given a constraint
2674 * a(p,i) + j = g f(e)
2676 * or -a(p,i) - j = g f(e) if sign < 0,
2677 * store a(p,i) in bound->shift and g (stride) in bound->stride.
2678 * a(p,i) is assumed to be an expression in only the parameters.
2680 static void extract_stride(__isl_keep isl_constraint
*c
,
2681 struct cuda_array_bound
*bound
, isl_int stride
, int sign
)
2689 isl_int_set(bound
->stride
, stride
);
2691 dim
= isl_constraint_get_space(c
);
2692 dim
= isl_space_params(dim
);
2694 nparam
= isl_space_dim(dim
, isl_dim_param
);
2698 isl_constraint_get_constant(c
, &v
);
2701 aff
= isl_aff_zero_on_domain(isl_local_space_from_space(dim
));
2702 aff
= isl_aff_set_constant(aff
, v
);
2704 for (i
= 0; i
< nparam
; ++i
) {
2705 isl_constraint_get_coefficient(c
, isl_dim_param
, i
, &v
);
2706 if (isl_int_is_zero(v
))
2710 aff
= isl_aff_add_coefficient(aff
, isl_dim_param
, i
, v
);
2718 /* Given an equality constraint of a map with a single output dimension j,
2719 * check if the constraint is of the form
2721 * a(p,i) + j = g f(e)
2723 * with a(p,i) an expression in the parameters and input dimensions
2724 * and f(e) an expression in the existentially quantified variables.
2725 * If so, and if g is larger than any such g from a previously considered
2726 * constraint, then call extract_stride. to record the stride information
2729 static int check_stride_constraint(__isl_take isl_constraint
*c
, void *user
)
2734 struct cuda_array_bound
*bound
= user
;
2737 isl_int_init(stride
);
2739 n_div
= isl_constraint_dim(c
, isl_dim_div
);
2740 isl_constraint_get_coefficient(c
, isl_dim_out
, 0, &v
);
2742 if (n_div
&& (isl_int_is_one(v
) || isl_int_is_negone(v
))) {
2743 int s
= isl_int_sgn(v
);
2744 isl_int_set_si(stride
, 0);
2745 for (i
= 0; i
< n_div
; ++i
) {
2746 isl_constraint_get_coefficient(c
, isl_dim_div
, i
, &v
);
2747 isl_int_gcd(stride
, stride
, v
);
2749 if (!isl_int_is_zero(stride
) &&
2750 isl_int_gt(stride
, bound
->stride
))
2751 extract_stride(c
, bound
, stride
, s
);
2754 isl_int_clear(stride
);
2757 isl_constraint_free(c
);
2761 /* Given contraints on an array index i, check if we can find
2762 * a shift a(p) and a stride g such that
2764 * a(p) + i = 0 mod g
2766 * If so, record the information in bound and apply the mapping
2767 * i -> (i + a(p))/g to the array index in bounds and return
2768 * the new constraints.
2769 * If not, simply return the original constraints.
2771 static __isl_give isl_basic_map
*check_stride(struct cuda_gen
*gen
,
2772 struct cuda_array_bound
*bound
, __isl_take isl_basic_map
*bounds
)
2775 isl_basic_map
*shift
;
2778 isl_int_set_si(bound
->stride
, -1);
2780 aff
= isl_basic_map_affine_hull(isl_basic_map_copy(bounds
));
2782 isl_basic_map_foreach_constraint(aff
, &check_stride_constraint
, bound
);
2784 isl_basic_map_free(aff
);
2786 if (isl_int_is_neg(bound
->stride
))
2789 aff_shift
= isl_aff_copy(bound
->shift
);
2790 aff_shift
= isl_aff_add_dims(aff_shift
, isl_dim_in
, 1);
2791 aff_shift
= isl_aff_add_coefficient_si(aff_shift
, isl_dim_in
, 0, 1);
2792 aff_shift
= isl_aff_scale_down(aff_shift
, bound
->stride
);
2793 shift
= isl_basic_map_from_aff(aff_shift
);
2795 bound
->shift_map
= isl_basic_map_copy(shift
);
2796 bounds
= isl_basic_map_apply_range(bounds
, shift
);
2801 struct cuda_size_info
{
2802 isl_basic_set
*bset
;
2803 struct cuda_array_bound
*bound
;
2807 /* Given a constraint from the basic set describing the bounds on
2808 * an array index, check if it is a lower bound, say m i >= b(x), and,
2809 * if so, check whether the expression "i - ceil(b(x)/m) + 1" has a constant
2810 * upper bound. If so, and if this bound is smaller than any bound
2811 * derived from earlier constraints, set the size to this bound on
2812 * the expression and the lower bound to ceil(b(x)/m).
2814 static int compute_size_in_direction(__isl_take isl_constraint
*c
, void *user
)
2816 struct cuda_size_info
*size
= user
;
2821 nparam
= isl_basic_set_dim(size
->bset
, isl_dim_param
);
2822 n_div
= isl_constraint_dim(c
, isl_dim_div
);
2824 if (isl_constraint_involves_dims(c
, isl_dim_div
, 0, n_div
)) {
2825 isl_constraint_free(c
);
2831 isl_constraint_get_coefficient(c
, isl_dim_set
, size
->pos
, &v
);
2833 if (isl_int_is_pos(v
)) {
2836 enum isl_lp_result res
;
2838 aff
= isl_constraint_get_bound(c
, isl_dim_set
, size
->pos
);
2839 aff
= isl_aff_ceil(aff
);
2841 lb
= isl_aff_copy(aff
);
2843 aff
= isl_aff_neg(aff
);
2844 aff
= isl_aff_add_coefficient_si(aff
, isl_dim_in
, size
->pos
, 1);
2846 res
= isl_basic_set_max(size
->bset
, aff
, &v
);
2849 if (res
== isl_lp_ok
) {
2850 isl_int_add_ui(v
, v
, 1);
2851 if (isl_int_is_neg(size
->bound
->size
) ||
2852 isl_int_lt(v
, size
->bound
->size
)) {
2853 isl_int_set(size
->bound
->size
, v
);
2854 lb
= isl_aff_drop_dims(lb
, isl_dim_in
,
2856 isl_aff_free(size
->bound
->lb
);
2857 size
->bound
->lb
= isl_aff_copy(lb
);
2864 isl_constraint_free(c
);
2869 /* Given a basic map "bounds" that maps parameters and input dimensions
2870 * to a single output dimension, look for an expression in the parameters
2871 * and input dimensions such that the range of the output dimension shifted
2872 * by this expression is a constant.
2874 * In particular, we currently only consider lower bounds on the output
2875 * dimension as candidate expressions.
2877 static int compute_array_dim_size(struct cuda_gen
*gen
,
2878 struct cuda_array_bound
*bound
, __isl_take isl_basic_map
*bounds
)
2880 struct cuda_size_info size
;
2882 bounds
= isl_basic_map_detect_equalities(bounds
);
2883 bounds
= check_stride(gen
, bound
, bounds
);
2885 isl_int_set_si(bound
->size
, -1);
2889 size
.pos
= isl_basic_map_dim(bounds
, isl_dim_in
);
2890 size
.bset
= isl_basic_map_wrap(bounds
);
2891 size
.bset
= isl_basic_set_flatten(size
.bset
);
2892 size
.bset
= isl_set_simple_hull(isl_basic_set_compute_divs(size
.bset
));
2893 isl_basic_set_foreach_constraint(size
.bset
, &compute_size_in_direction
,
2895 isl_basic_set_free(size
.bset
);
2897 return isl_int_is_nonneg(bound
->size
) ? 0 : -1;
2900 /* Check if we can find a shared memory tile for the given array
2901 * based on the given accesses, and if so, put the results
2902 * in array->shared_bound.
2904 * We project the accesses on each index in turn and look for a parametric
2905 * offset such that the size is constant.
2907 static int can_tile_for_shared_memory(struct cuda_gen
*gen
,
2908 struct cuda_array_info
*array
, __isl_keep isl_map
*access
,
2909 struct cuda_array_bound
*bounds
)
2913 for (i
= 0; i
< array
->n_index
; ++i
) {
2915 isl_basic_map
*hull
;
2917 access_i
= isl_map_copy(access
);
2918 access_i
= isl_map_project_out(access_i
, isl_dim_out
, 0, i
);
2919 access_i
= isl_map_project_out(access_i
, isl_dim_out
,
2920 1, array
->n_index
- (i
+ 1));
2921 access_i
= isl_map_compute_divs(access_i
);
2922 hull
= isl_map_simple_hull(access_i
);
2923 if (compute_array_dim_size(gen
, &bounds
[i
], hull
) < 0)
2930 /* Construct a map with input the shared tile loops and the loops that
2931 * will be wrapped around the threads that relates these later loops
2932 * to the thread indices and then projects them out.
2934 static __isl_give isl_map
*compute_privatization(struct cuda_gen
*gen
)
2942 dim
= isl_union_map_get_space(gen
->shared_sched
);
2944 if (gen
->options
->wrap
)
2945 tiling
= wrap(isl_space_copy(dim
), gen
->shared_len
+ gen
->n_block
,
2946 gen
->shared_len
, gen
->n_block
, gen
->block_dim
);
2948 tiling
= tile(isl_space_copy(dim
), gen
->shared_len
+ gen
->n_block
,
2949 gen
->shared_len
, gen
->n_block
, gen
->block_dim
);
2953 par
= parametrization(dim
, gen
->shared_len
+ 2 * gen
->n_block
,
2954 gen
->tile_first
+ gen
->tile_len
+ gen
->n_grid
+ gen
->n_block
,
2957 priv
= isl_map_align_params(priv
, isl_set_get_space(par
));
2958 priv
= isl_map_intersect_range(priv
, par
);
2960 dim
= isl_map_get_space(priv
);
2961 dim
= isl_space_drop_dims(dim
, isl_dim_in
, 0, isl_space_dim(dim
, isl_dim_in
));
2962 dim
= isl_space_drop_dims(dim
, isl_dim_out
, 0, isl_space_dim(dim
, isl_dim_out
));
2963 proj
= projection(dim
, gen
->shared_len
+ 2 * gen
->n_block
,
2966 priv
= isl_map_apply_range(priv
, proj
);
2971 /* Construct a map from domain_dim to domain_dim that increments
2972 * the dimension at position "pos" and leaves all other dimensions
2975 static __isl_give isl_map
*next(__isl_take isl_space
*domain_dim
, int pos
)
2978 int len
= isl_space_dim(domain_dim
, isl_dim_set
);
2980 isl_basic_map
*next
;
2981 isl_local_space
*ls
;
2983 dim
= isl_space_map_from_set(domain_dim
);
2984 next
= isl_basic_map_universe(isl_space_copy(dim
));
2985 ls
= isl_local_space_from_space(dim
);
2987 for (i
= 0; i
< len
; ++i
) {
2990 c
= isl_equality_alloc(isl_local_space_copy(ls
));
2991 isl_constraint_set_coefficient_si(c
, isl_dim_in
, i
, 1);
2992 isl_constraint_set_coefficient_si(c
, isl_dim_out
, i
, -1);
2994 isl_constraint_set_constant_si(c
, 1);
2995 next
= isl_basic_map_add_constraint(next
, c
);
2998 isl_local_space_free(ls
);
3000 return isl_map_from_basic_map(next
);
3003 /* Check if the given access is coalesced.
3004 * That is, check whether incrementing the dimension that will get
3005 * wrapped over the last thread index results in incrementing
3006 * the last array index.
3008 * This function is only called for access relations without reuse.
3010 static int access_is_coalesced(struct cuda_gen
*gen
,
3011 __isl_keep isl_union_map
*access
)
3014 isl_map
*access_map
;
3015 isl_map
*next_thread_x
;
3016 isl_map
*next_element
;
3020 access
= isl_union_map_copy(access
);
3021 access
= isl_union_map_apply_domain(access
,
3022 isl_union_map_copy(gen
->tiled_sched
));
3023 access_map
= isl_map_from_union_map(access
);
3025 dim
= isl_map_get_space(access_map
);
3026 dim
= isl_space_domain(dim
);
3027 next_thread_x
= next(dim
, gen
->shared_len
+ gen
->n_block
- 1);
3029 dim
= isl_map_get_space(access_map
);
3030 dim
= isl_space_range(dim
);
3031 next_element
= next(dim
, isl_space_dim(dim
, isl_dim_set
) - 1);
3033 map
= isl_map_apply_domain(next_thread_x
, isl_map_copy(access_map
));
3034 map
= isl_map_apply_range(map
, access_map
);
3036 coalesced
= isl_map_is_subset(map
, next_element
);
3038 isl_map_free(next_element
);
3044 /* For the given array reference group, check whether the access is private
3045 * to the thread. That is, check that any given array element
3046 * is only accessed by a single thread.
3047 * We compute an access relation that maps the shared tile loop iterators
3048 * and the shared point loop iterators that will be wrapped over the
3049 * threads to the array elements.
3050 * We actually check that those iterators that will be wrapped
3051 * partition the array space. This check is stricter than necessary
3052 * since several iterations may be mapped onto the same thread
3053 * and then they could be allowed to access the same memory elements,
3054 * but our check does not allow this situation.
3056 * We also check that the index expression only depends on parallel
3057 * loops. That way, we can move those loops innermost and unroll them.
3058 * Again, we use a test that is stricter than necessary.
3059 * We actually check whether the index expression only depends
3060 * on the iterators that are wrapped over the threads.
3061 * These are necessarily parallel, but there may be more parallel loops.
3063 * Combining the injectivity of the first test with the single-valuedness
3064 * of the second test, we simply test for bijectivity.
3066 * If it turns out we can use registers, we compute the private memory
3067 * tile size using can_tile_for_shared_memory, after introducing a dependence
3068 * on the thread indices.
3070 * Before performing any of the above computations, we first check
3071 * if there is any reuse on the reference group. If not, we simply
3072 * return. If, moreover, the access is coalesced then we also remove
3073 * the shared memory tiling since we should just use global memory instead.
3075 static void check_private_group_access(struct cuda_gen
*gen
,
3076 struct cuda_array_ref_group
*group
)
3079 isl_union_map
*access
;
3080 int n_index
= group
->array
->n_index
;
3082 access
= group_access_relation(group
, 1, 1);
3083 if (isl_union_map_is_injective(access
)) {
3084 if (group
->shared_bound
&& access_is_coalesced(gen
, access
)) {
3085 free_bound_list(group
->shared_bound
, n_index
);
3086 group
->shared_bound
= NULL
;
3088 isl_union_map_free(access
);
3091 access
= isl_union_map_apply_domain(access
,
3092 isl_union_map_copy(gen
->shared_sched
));
3094 acc
= isl_map_from_union_map(access
);
3096 if (!isl_map_is_bijective(acc
)) {
3101 group
->private_bound
= create_bound_list(gen
->ctx
, n_index
);
3102 acc
= isl_map_align_params(acc
, isl_map_get_space(gen
->privatization
));
3103 acc
= isl_map_apply_domain(acc
, isl_map_copy(gen
->privatization
));
3104 if (!can_tile_for_shared_memory(gen
, group
->array
, acc
,
3105 group
->private_bound
)) {
3106 free_bound_list(group
->private_bound
, n_index
);
3107 group
->private_bound
= NULL
;
3113 /* Look for the last shared tile loop that affects the offset of the
3114 * shared or private tile and store the result in array->last_shared.
3116 static void set_last_shared(struct cuda_gen
*gen
,
3117 struct cuda_array_ref_group
*group
)
3120 struct cuda_array_bound
*bounds
;
3121 unsigned first_shared
= gen
->first_shared
;
3122 int n_index
= group
->array
->n_index
;
3124 bounds
= group
->private_bound
;
3126 bounds
= group
->shared_bound
;
3130 for (j
= gen
->shared_len
- 1; j
>= 0; --j
) {
3131 for (i
= 0; i
< n_index
; ++i
) {
3136 if (isl_aff_involves_dims(lb
, isl_dim_param
,
3137 first_shared
+ j
, 1))
3140 shift
= bounds
[i
].shift
;
3143 if (isl_aff_involves_dims(shift
, isl_dim_param
,
3144 first_shared
+ j
, 1))
3150 group
->last_shared
= j
;
3153 /* Compute the sizes of all private arrays for the current kernel,
3154 * as well as the offsets of the private pieces in the original arrays.
3155 * If we cannot or don't want to privatize a given array group,
3156 * we use the shared memory tile sizes computed in
3157 * compute_group_shared_bound instead.
3159 * If we have been able to find a private or shared tile,
3160 * we also look for the last shared tile loop that affects the offset
3161 * (and therefore the group tile) and store the result in group->last_shared.
3163 * A privatized copy of all access relations from reference groups that
3164 * are mapped to private memory is stored in gen->privatization.
3166 static void compute_private_size(struct cuda_gen
*gen
)
3169 isl_union_map
*private;
3171 if (!gen
->options
->use_private_memory
)
3174 private = isl_union_map_empty(isl_union_map_get_space(gen
->shared_sched
));
3176 for (i
= 0; i
< gen
->n_array
; ++i
) {
3177 struct cuda_array_info
*array
= &gen
->array
[i
];
3179 for (j
= 0; j
< array
->n_group
; ++j
) {
3180 check_private_group_access(gen
, array
->groups
[j
]);
3182 if (!array
->groups
[j
]->private_bound
)
3185 private = isl_union_map_union(private,
3186 group_access_relation(array
->groups
[j
], 1, 1));
3189 for (j
= 0; j
< array
->n_group
; ++j
) {
3190 array
->groups
[j
]->last_shared
= gen
->shared_len
- 1;
3191 array
->groups
[j
]->print_shared_level
= -1;
3192 set_last_shared(gen
, array
->groups
[j
]);
3196 if (isl_union_map_is_empty(private))
3197 isl_union_map_free(private);
3199 isl_union_map
*priv
;
3201 private = isl_union_map_apply_domain(private,
3202 isl_union_map_copy(gen
->shared_sched
));
3203 priv
= isl_union_map_from_map(isl_map_copy(gen
->privatization
));
3204 private = isl_union_map_apply_domain(private, priv
);
3205 gen
->private_access
= private;
3209 /* Compute the size of the tile specified by the list "bound" of n_index
3210 * cuda_array_bounds in number of elements and put the result in *size.
3212 static void tile_size(unsigned n_index
, struct cuda_array_bound
*bound
,
3217 isl_int_set_si(*size
, 1);
3219 for (i
= 0; i
< n_index
; ++i
)
3220 isl_int_mul(*size
, *size
, bound
[i
].size
);
3223 /* If max_shared_memory is not set to infinity (-1), then make
3224 * sure that the total amount of shared memory required by the
3225 * array reference groups mapped to shared memory is no larger
3226 * than this maximum.
3228 * We apply a greedy approach and discard (keep in global memory)
3229 * those groups that would result in a total memory size that
3230 * is larger than the maximum.
3232 static void check_shared_memory_bound(struct cuda_gen
*gen
)
3237 if (gen
->options
->max_shared_memory
< 0)
3242 isl_int_set_si(left
, gen
->options
->max_shared_memory
);
3244 for (i
= 0; i
< gen
->n_array
; ++i
) {
3245 struct cuda_array_info
*array
= &gen
->array
[i
];
3247 for (j
= 0; j
< array
->n_group
; ++j
) {
3248 struct cuda_array_ref_group
*group
;
3250 group
= array
->groups
[j
];
3251 if (!group
->shared_bound
)
3254 tile_size(array
->n_index
, group
->shared_bound
, &size
);
3255 isl_int_mul_ui(size
, size
, array
->size
);
3257 if (isl_int_le(size
, left
)) {
3258 isl_int_sub(left
, left
, size
);
3262 free_bound_list(group
->shared_bound
, array
->n_index
);
3263 group
->shared_bound
= NULL
;
3267 isl_int_clear(size
);
3268 isl_int_clear(left
);
3271 /* Fill up the groups array with singleton groups, i.e., one group
3272 * per reference, initializing the array, access, write and refs fields.
3273 * In particular the access field is initialized to the scheduled
3274 * access relation of the array reference.
3276 * Return the number of elements initialized, i.e., the number of
3277 * active references in the current kernel.
3279 static int populate_array_references(struct cuda_gen
*gen
,
3280 struct cuda_array_info
*array
, __isl_keep isl_union_map
*sched
,
3281 struct cuda_array_ref_group
**groups
)
3285 isl_ctx
*ctx
= isl_union_map_get_ctx(sched
);
3288 for (i
= 0; i
< array
->n_ref
; ++i
) {
3289 isl_union_map
*umap
;
3291 struct cuda_array_ref_group
*group
;
3292 struct cuda_stmt_access
*access
= array
->refs
[i
];
3294 map
= isl_map_copy(access
->access
);
3295 umap
= isl_union_map_from_map(map
);
3296 umap
= isl_union_map_apply_domain(umap
,
3297 isl_union_map_copy(sched
));
3299 if (isl_union_map_is_empty(umap
)) {
3300 isl_union_map_free(umap
);
3304 map
= isl_map_from_union_map(umap
);
3305 map
= isl_map_detect_equalities(map
);
3307 group
= isl_calloc_type(ctx
, struct cuda_array_ref_group
);
3309 group
->array
= array
;
3310 group
->access
= map
;
3311 group
->write
= access
->write
;
3312 group
->refs
= &array
->refs
[i
];
3314 groups
[n
++] = group
;
3320 static void free_array_ref_group(struct cuda_array_ref_group
*group
,
3325 free_bound_list(group
->shared_bound
, n_index
);
3326 free_bound_list(group
->private_bound
, n_index
);
3327 isl_map_free(group
->access
);
3332 /* Given a set where the parameters gen->first_shared up to
3333 * gen->first_shared + gen->shared_len represent the tile loops,
3334 * eliminate the innermost of those that have a fixed value
3335 * until we reach one that does not (obviously) have a fixed value.
3337 static __isl_give isl_set
*eliminate_fixed_inner_loops(struct cuda_gen
*gen
,
3338 __isl_take isl_set
*access
)
3342 for (i
= gen
->shared_len
- 1; i
>= 0; --i
) {
3343 int pos
= gen
->first_shared
+ i
;
3344 if (!isl_set_plain_is_fixed(access
, isl_dim_param
, pos
, NULL
))
3346 access
= isl_set_eliminate(access
, isl_dim_param
, pos
, 1);
3351 /* Check if the accessed set of group1 and group2 overlap within
3352 * the innermost loop. In particular, ignore any inner dimension
3353 * with a fixed value.
3354 * The copying to and from shared memory will be performed within
3355 * the innermost actual loop so we are only allowed to consider
3356 * the dimensions up to that innermost loop while checking whether
3357 * two access sets overlap.
3359 static int accesses_overlap(struct cuda_gen
*gen
,
3360 struct cuda_array_ref_group
*group1
,
3361 struct cuda_array_ref_group
*group2
)
3364 isl_set
*access1
, *access2
;
3366 access1
= isl_map_range(isl_map_copy(group1
->access
));
3367 access1
= eliminate_fixed_inner_loops(gen
, access1
);
3368 access2
= isl_map_range(isl_map_copy(group2
->access
));
3369 access2
= eliminate_fixed_inner_loops(gen
, access2
);
3370 access1
= isl_set_intersect(access1
, access2
);
3371 empty
= isl_set_is_empty(access1
);
3372 isl_set_free(access1
);
3377 /* If two groups have overlapping access relations (within the innermost
3378 * loop) and if one of them involves a write, then merge the two groups
3381 * We keep track of the grouping in "leader". leader[j] points to
3382 * an earlier group array element that belongs to the same group,
3383 * or the array element j itself if this element is the first in the group.
3385 * Return the number of group leaders.
3387 static int group_overlapping_writes(struct cuda_gen
*gen
, int n
,
3388 struct cuda_array_ref_group
**groups
, int *leader
)
3393 for (i
= 0; i
< n
; ++i
) {
3395 groups
[l
]->n_ref
= 1;
3396 for (j
= i
- 1; j
>= 0; --j
) {
3399 if (!groups
[l
]->write
&& !groups
[j
]->write
)
3402 if (!accesses_overlap(gen
, groups
[l
], groups
[j
]))
3405 groups
[j
]->access
= isl_map_union(groups
[j
]->access
,
3407 groups
[j
]->write
= 1;
3408 groups
[l
]->access
= NULL
;
3409 groups
[j
]->n_ref
+= groups
[l
]->n_ref
;
3419 /* Compute the size of the shared array corresponding to the given array
3420 * array refrence group, based on the accesses from the current kernel,
3421 * as well as the offset of the shared piece in the original array.
3423 static void compute_group_shared_bound(struct cuda_gen
*gen
,
3424 struct cuda_array_info
*array
, struct cuda_array_ref_group
*group
)
3426 isl_ctx
*ctx
= isl_space_get_ctx(array
->dim
);
3428 if (!gen
->options
->use_shared_memory
)
3430 if (cuda_array_is_read_only_scalar(array
))
3433 group
->shared_bound
= create_bound_list(ctx
, array
->n_index
);
3434 if (!can_tile_for_shared_memory(gen
, array
, group
->access
,
3435 group
->shared_bound
)) {
3436 free_bound_list(group
->shared_bound
, array
->n_index
);
3437 group
->shared_bound
= NULL
;
3441 /* Is the size of the tile specified by "bound" smaller than the sum of
3442 * the sizes of the tiles specified by "bound1" and "bound2"?
3444 static int smaller_tile(unsigned n_index
, struct cuda_array_bound
*bound
,
3445 struct cuda_array_bound
*bound1
, struct cuda_array_bound
*bound2
)
3448 isl_int size
, size1
, size2
;
3451 isl_int_init(size1
);
3452 isl_int_init(size2
);
3454 tile_size(n_index
, bound
, &size
);
3455 tile_size(n_index
, bound1
, &size1
);
3456 tile_size(n_index
, bound2
, &size2
);
3458 isl_int_sub(size
, size
, size1
);
3459 isl_int_sub(size
, size
, size2
);
3460 smaller
= isl_int_is_neg(size
);
3462 isl_int_clear(size2
);
3463 isl_int_clear(size1
);
3464 isl_int_clear(size
);
3469 /* Given an initial grouping of array references and shared memory tiles
3470 * for each group that allows for a shared memory tile, merge two groups
3471 * if both have a shared memory tile, the merged group also has
3472 * a shared memory tile and the size of the tile for the merge group
3473 * is smaller than the sum of the tile sizes of the individual groups.
3475 * Return the number of group leaders after merging.
3477 static int group_common_shared_memory_tile(struct cuda_gen
*gen
,
3478 struct cuda_array_info
*array
, int n
,
3479 struct cuda_array_ref_group
**groups
, int *leader
, int n_group
)
3482 isl_ctx
*ctx
= isl_space_get_ctx(array
->dim
);
3484 for (i
= 0; n_group
> 1 && i
< n
; ++i
) {
3488 if (!groups
[i
]->shared_bound
)
3490 for (j
= i
- 1; j
>= 0; --j
) {
3493 struct cuda_array_bound
*shared_bound
;
3497 if (!groups
[j
]->shared_bound
)
3500 map
= isl_map_intersect(isl_map_copy(groups
[l
]->access
),
3501 isl_map_copy(groups
[j
]->access
));
3502 empty
= isl_map_is_empty(map
);
3508 map
= isl_map_union(isl_map_copy(groups
[l
]->access
),
3509 isl_map_copy(groups
[j
]->access
));
3510 shared_bound
= create_bound_list(ctx
, array
->n_index
);
3511 if (!can_tile_for_shared_memory(gen
, array
, map
,
3513 !smaller_tile(array
->n_index
, shared_bound
,
3514 groups
[l
]->shared_bound
,
3515 groups
[j
]->shared_bound
)) {
3517 free_bound_list(shared_bound
, array
->n_index
);
3521 free_bound_list(groups
[j
]->shared_bound
,
3523 groups
[j
]->shared_bound
= shared_bound
;
3524 isl_map_free(groups
[j
]->access
);
3525 groups
[j
]->access
= map
;
3526 groups
[j
]->n_ref
+= groups
[l
]->n_ref
;
3535 /* Extract an array of array reference groups from the array of references
3536 * and the grouping information in "leader".
3538 * Store the results in array->n_group and array->groups.
3540 static void extract_array_groups(isl_ctx
*ctx
, struct cuda_array_info
*array
,
3541 int n
, struct cuda_array_ref_group
**groups
, int *leader
, int n_group
)
3545 for (i
= 2; i
< n
; ++i
)
3546 leader
[i
] = leader
[leader
[i
]];
3548 array
->n_group
= n_group
;
3549 array
->groups
= isl_alloc_array(ctx
, struct cuda_array_ref_group
*,
3551 assert(array
->groups
);
3554 for (i
= 0; i
< n
; ++i
) {
3556 struct cuda_stmt_access
**refs
;
3558 if (leader
[i
] != i
) {
3559 groups
[i
]->refs
= NULL
;
3560 free_array_ref_group(groups
[i
], array
->n_index
);
3564 refs
= isl_alloc_array(ctx
, struct cuda_stmt_access
*,
3568 for (k
= i
; k
< n
; ++k
)
3569 if (leader
[k
] == i
) {
3570 refs
[l
++] = *groups
[k
]->refs
;
3571 (*groups
[k
]->refs
)->group
= j
;
3574 groups
[i
]->refs
= refs
;
3576 array
->groups
[j
++] = groups
[i
];
3580 /* Group array references that should be considered together when
3581 * deciding whether to access them from private, shared or global memory.
3583 * In particular, if two array references overlap and if one of them
3584 * is a write, then the two references are grouped together.
3585 * Furthermore, if two groups admit a shared memory tile and if the
3586 * combination of the two also admits a shared memory tile, we merge
3589 * During the construction the group->refs field points to a single
3590 * array reference inside the array of array references, while
3591 * group->n_ref contains the number of element in leader that
3592 * (directly or indirectly) point to this group, provided the group
3595 static void group_array_references(struct cuda_gen
*gen
,
3596 struct cuda_array_info
*array
, __isl_keep isl_union_map
*sched
)
3600 isl_ctx
*ctx
= isl_union_map_get_ctx(sched
);
3601 struct cuda_array_ref_group
**groups
;
3604 groups
= isl_calloc_array(ctx
, struct cuda_array_ref_group
*,
3608 n
= populate_array_references(gen
, array
, sched
, groups
);
3610 leader
= isl_alloc_array(ctx
, int, n
);
3613 n_group
= group_overlapping_writes(gen
, n
, groups
, leader
);
3615 for (i
= 0; i
< n
; ++i
)
3617 compute_group_shared_bound(gen
, array
, groups
[i
]);
3619 n_group
= group_common_shared_memory_tile(gen
, array
, n
, groups
,
3622 extract_array_groups(ctx
, array
, n
, groups
, leader
, n_group
);
3628 /* Take tiled_sched, project it onto the shared tile loops and
3629 * the loops that will be wrapped over the threads,
3630 * parametrize the shared tile loops and store the result in gen->shared_sched.
3631 * The position of the first of these parameters is stored in gen->first_shared.
3632 * Also compute a projection that projects out the loops that will be
3633 * wrapped over the threads and store this projection in gen->shared_proj.
3635 static void compute_shared_sched(struct cuda_gen
*gen
)
3640 isl_union_map
*sched
;
3642 sched
= isl_union_map_copy(gen
->tiled_sched
);
3644 dim
= isl_union_map_get_space(sched
);
3645 gen
->first_shared
= isl_space_dim(dim
, isl_dim_param
);
3646 proj
= projection(dim
, gen
->tiled_len
, gen
->shared_len
+ gen
->n_block
);
3647 sched
= isl_union_map_apply_range(sched
, isl_union_map_from_map(proj
));
3649 dim
= isl_union_map_get_space(sched
);
3650 par
= parametrization(dim
, gen
->shared_len
+ gen
->n_block
,
3651 0, gen
->shared_len
, "g");
3652 sched
= isl_union_map_intersect_range(sched
,
3653 isl_union_set_from_set(par
));
3655 dim
= isl_union_map_get_space(sched
);
3656 proj
= projection(dim
, gen
->shared_len
+ gen
->n_block
, gen
->shared_len
);
3658 gen
->shared_sched
= sched
;
3659 gen
->shared_proj
= isl_union_map_from_map(proj
);
3662 /* Group references of all arrays in the program.
3664 static void group_references(struct cuda_gen
*gen
)
3667 isl_union_map
*sched
;
3669 sched
= isl_union_map_apply_range(isl_union_map_copy(gen
->shared_sched
),
3670 isl_union_map_copy(gen
->shared_proj
));
3672 for (i
= 0; i
< gen
->n_array
; ++i
)
3673 group_array_references(gen
, &gen
->array
[i
], sched
);
3675 isl_union_map_free(sched
);
3678 /* Free all array information that is local to the current kernel.
3680 static void free_local_array_info(struct cuda_gen
*gen
)
3684 for (i
= 0; i
< gen
->n_array
; ++i
) {
3685 struct cuda_array_info
*array
= &gen
->array
[i
];
3687 for (j
= 0; j
< array
->n_group
; ++j
)
3688 free_array_ref_group(array
->groups
[j
], array
->n_index
);
3689 free(array
->groups
);
3691 if (array
->n_group
== 0)
3693 for (j
= 0; j
< gen
->array
[i
].n_index
; ++j
) {
3694 isl_pw_aff_free(gen
->array
[i
].local_bound
[j
]);
3695 gen
->array
[i
].local_bound
[j
] = NULL
;
3700 /* The sizes of the arrays on the host that have been computed by
3701 * extract_array_info may depend on the parameters. Use the extra
3702 * constraints on the parameters that are valid at "host_domain"
3703 * to simplify these expressions.
3705 static void localize_bounds(struct cuda_gen
*gen
,
3706 __isl_keep isl_set
*host_domain
)
3711 context
= isl_set_copy(host_domain
);
3712 context
= isl_set_params(host_domain
);
3714 for (i
= 0; i
< gen
->n_array
; ++i
) {
3715 struct cuda_array_info
*array
= &gen
->array
[i
];
3717 if (array
->n_group
== 0)
3720 for (j
= 0; j
< array
->n_index
; ++j
) {
3723 pwaff
= isl_pw_aff_copy(array
->bound
[j
]);
3724 pwaff
= isl_pw_aff_gist(pwaff
, isl_set_copy(context
));
3725 array
->local_bound
[j
] = pwaff
;
3728 isl_set_free(context
);
3731 /* Set gen->tile_len and gen->n_parallel to those of the first statement
3732 * in the statement list u.
3733 * Because of the way the schedule is constructed, the other statements
3734 * in the list, if any, should have the same values for these properties.
3736 static void set_tile_len(struct cuda_gen
*gen
, struct clast_user_stmt
*u
)
3739 struct cuda_stmt
*stmt
;
3741 nr
= atoi(u
->statement
->name
+ 2);
3742 stmt
= &gen
->stmts
[nr
];
3744 gen
->tile_len
= stmt
->tile_len
;
3745 gen
->n_parallel
= stmt
->n_parallel
;
3748 /* Extract a description of the grid, i.e., the possible values
3749 * of the block ids, from gen->tiled_sched.
3750 * The block ids are parameters in gen->tiled_sched.
3751 * We simply need to change them into set dimensions.
3753 static __isl_give isl_set
*extract_grid(struct cuda_gen
*gen
)
3758 grid
= isl_union_map_params(isl_union_map_copy(gen
->tiled_sched
));
3759 grid
= isl_set_from_params(grid
);
3760 grid
= isl_set_add_dims(grid
, isl_dim_set
, gen
->n_grid
);
3761 for (i
= 0; i
< gen
->n_grid
; ++i
) {
3765 snprintf(name
, sizeof(name
), "b%d", i
);
3766 pos
= isl_set_find_dim_by_name(grid
, isl_dim_param
, name
);
3768 grid
= isl_set_equate(grid
, isl_dim_param
, pos
, isl_dim_set
, i
);
3769 grid
= isl_set_project_out(grid
, isl_dim_param
, pos
, 1);
3775 /* Print the effective grid size as a list of the sizes in each
3776 * dimension, from innermost to outermost.
3778 * The grid size specified by the user or set by default
3779 * in read_grid_sizes() and applied in tile_schedule(),
3780 * may be too large for the given code in the sense that
3781 * it may contain blocks that don't need to execute anything.
3782 * We therefore don't print this grid size, but instead the
3783 * smallest grid size that ensures that all blocks that actually
3784 * execute code are included in the grid.
3786 * For each block dimension, we compute the maximal value of the block id
3789 static void print_grid_size(struct cuda_gen
*gen
, __isl_take isl_set
*context
)
3795 if (gen
->n_grid
== 0) {
3796 isl_set_free(context
);
3800 grid
= extract_grid(gen
);
3802 prn
= isl_printer_to_file(gen
->ctx
, gen
->cuda
.host_c
);
3803 prn
= isl_printer_set_output_format(prn
, ISL_FORMAT_C
);
3805 prn
= isl_printer_print_str(prn
, "(");
3806 for (i
= gen
->n_grid
- 1; i
>= 0; --i
) {
3809 isl_pw_aff
*bound
= isl_set_dim_max(isl_set_copy(grid
), i
);
3811 bound
= isl_pw_aff_coalesce(bound
);
3812 bound
= isl_pw_aff_gist(bound
, isl_set_copy(context
));
3814 space
= isl_pw_aff_get_domain_space(bound
);
3815 one
= isl_aff_zero_on_domain(isl_local_space_from_space(space
));
3816 one
= isl_aff_add_constant_si(one
, 1);
3817 bound
= isl_pw_aff_add(bound
, isl_pw_aff_from_aff(one
));
3818 prn
= isl_printer_print_pw_aff(prn
, bound
);
3819 isl_pw_aff_free(bound
);
3822 prn
= isl_printer_print_str(prn
, ", ");
3824 prn
= isl_printer_print_str(prn
, ")");
3826 isl_printer_free(prn
);
3828 isl_set_free(context
);
3831 /* This function is called for each leaf in the clast of the host code.
3832 * We first specialize the schedule to the site of the leaf, compute
3833 * the size of shared memory and then print the body of host code
3834 * and the associated kernel (through a call to print_kernel_body).
3836 static void print_host_user(struct clast_printer_info
*code
,
3837 struct clast_user_stmt
*u
)
3839 struct cuda_gen
*gen
= code
->user
;
3842 isl_set
*host_domain
;
3843 isl_union_map
*access
;
3844 isl_union_map
*local_sched
;
3845 isl_union_set
*arrays
;
3847 set_tile_len(gen
, u
);
3850 host_domain
= extract_entire_host_domain(&u
->stmt
);
3852 local_sched
= isl_union_map_intersect_range(
3853 isl_union_map_copy(gen
->sched
),
3854 isl_union_set_from_set(extend(isl_set_copy(host_domain
),
3855 gen
->untiled_len
)));
3856 access
= isl_union_map_union(isl_union_map_copy(gen
->read
),
3857 isl_union_map_copy(gen
->write
));
3858 access
= isl_union_map_apply_domain(access
,
3859 isl_union_map_copy(local_sched
));
3860 arrays
= isl_union_map_range(access
);
3862 print_indent(code
->dst
, code
->indent
);
3863 fprintf(code
->dst
, "dim3 k%d_dimBlock", gen
->kernel_id
);
3864 print_reverse_list(code
->dst
, gen
->n_block
, gen
->block_dim
);
3865 fprintf(code
->dst
, ";\n");
3867 gen
->tiled_sched
= tile_schedule(gen
, local_sched
);
3868 gen
->tiled_sched
= parametrize_tiled_schedule(gen
, gen
->tiled_sched
);
3869 gen
->tiled_sched
= scale_tile_loops(gen
, gen
->tiled_sched
);
3871 print_indent(code
->dst
, code
->indent
);
3872 fprintf(code
->dst
, "dim3 k%d_dimGrid", gen
->kernel_id
);
3873 print_grid_size(gen
, isl_set_params(isl_set_copy(host_domain
)));
3874 fprintf(code
->dst
, ";\n");
3876 gen
->local_sched
= isl_union_map_copy(gen
->tiled_sched
);
3878 dim
= isl_union_map_get_space(gen
->local_sched
);
3879 par
= parametrization(dim
, gen
->tiled_len
, 0, gen
->shared_len
, "g");
3880 gen
->local_sched
= isl_union_map_intersect_range(gen
->local_sched
,
3881 isl_union_set_from_set(par
));
3883 gen
->local_sched
= thread_tile_schedule(gen
, gen
->local_sched
);
3884 gen
->local_sched
= scale_thread_tile_loops(gen
, gen
->local_sched
);
3886 gen
->private_access
= NULL
;
3887 compute_shared_sched(gen
);
3888 gen
->privatization
= compute_privatization(gen
);
3889 group_references(gen
);
3890 compute_private_size(gen
);
3891 check_shared_memory_bound(gen
);
3892 localize_bounds(gen
, host_domain
);
3894 gen
->local_sched
= interchange_for_unroll(gen
, gen
->local_sched
);
3896 print_kernel_launch(gen
, arrays
);
3898 fprintf(gen
->cuda
.kernel_c
, "{\n");
3900 print_kernel_body(gen
, host_domain
, gen
->tiled_sched
);
3902 fprintf(gen
->cuda
.kernel_c
, "}\n");
3904 free_local_array_info(gen
);
3905 isl_map_free(gen
->privatization
);
3906 isl_union_map_free(gen
->private_access
);
3907 isl_union_map_free(gen
->local_sched
);
3908 isl_union_map_free(gen
->tiled_sched
);
3909 isl_union_map_free(gen
->shared_sched
);
3910 isl_union_map_free(gen
->shared_proj
);
3911 isl_union_set_free(arrays
);
3912 isl_set_free(host_domain
);
3914 free(gen
->tile_size
);
3918 /* Use CLooG to generate code for the outer gen->tile_first loops
3919 * of the global schedule in gen->sched.
3920 * The pretty printing of this code is handled by print_clast,
3921 * which calls print_host_user for each kernel invocation location.
3923 static void print_cloog_host_code(struct cuda_gen
*gen
)
3927 isl_union_map
*sched
;
3928 CloogOptions
*options
;
3929 CloogDomain
*cloog_context
;
3930 CloogUnionDomain
*ud
;
3932 struct clast_stmt
*stmt
;
3935 options
= cloog_options_malloc(gen
->state
);
3936 options
->language
= CLOOG_LANGUAGE_C
;
3938 options
->strides
= 1;
3939 options
->stop
= gen
->tile_first
;
3940 options
->f
= gen
->untiled_len
;
3941 options
->l
= gen
->untiled_len
;
3942 options
->save_domains
= 1;
3943 options
->noscalars
= 1;
3945 sched
= isl_union_map_copy(gen
->sched
);
3946 ud
= cloog_union_domain_from_isl_union_map(sched
);
3947 for (i
= 0; i
< options
->stop
; ++i
) {
3948 snprintf(name
, sizeof(name
), "h%d", i
);
3949 ud
= cloog_union_domain_set_name(ud
, CLOOG_SCAT
, i
, name
);
3951 context
= isl_set_copy(gen
->context
);
3952 cloog_context
= cloog_domain_from_isl_set(context
);
3953 input
= cloog_input_alloc(cloog_context
, ud
);
3955 stmt
= cloog_clast_create_from_input(input
, options
);
3957 gen
->code
.indent
= 0;
3958 gen
->code
.dst
= gen
->cuda
.host_c
;
3959 gen
->code
.print_user_stmt
= NULL
;
3960 gen
->code
.print_user_stmt_list
= &print_host_user
;
3961 gen
->code
.print_for_head
= NULL
;
3962 gen
->code
.print_for_foot
= NULL
;
3963 gen
->code
.user
= gen
;
3964 print_clast(&gen
->code
, stmt
);
3966 cloog_clast_free(stmt
);
3967 cloog_options_free(options
);
3968 fprintf(gen
->cuda
.host_c
, "\n");
3971 void print_cuda_macros(struct cuda_gen
*gen
)
3973 const char *macros
=
3974 "#define cudaCheckReturn(ret) assert((ret) == cudaSuccess)\n"
3975 "#define cudaCheckKernel()"
3976 " assert(cudaGetLastError() == cudaSuccess)\n\n";
3977 fputs(macros
, gen
->cuda
.host_c
);
3980 void print_host_code(struct cuda_gen
*gen
)
3982 fprintf(gen
->cuda
.host_c
, "{\n");
3983 print_cloog_macros(gen
->cuda
.host_c
);
3984 print_cloog_macros(gen
->cuda
.kernel_c
);
3986 print_cuda_macros(gen
);
3988 declare_device_arrays(gen
);
3990 allocate_device_arrays(gen
);
3991 copy_arrays_to_device(gen
);
3994 print_cloog_host_code(gen
);
3996 copy_arrays_from_device(gen
);
3997 free_device_arrays(gen
);
3999 fprintf(gen
->cuda
.host_c
, "}\n");
4002 __isl_give isl_set
*add_context_from_str(__isl_take isl_set
*set
,
4011 ctx
= isl_set_get_ctx(set
);
4012 context
= isl_set_read_from_str(ctx
, str
);
4013 context
= isl_set_align_params(context
, isl_set_get_space(set
));
4014 set
= isl_set_intersect(set
, context
);
4019 __isl_give isl_union_map
*extract_sizes_from_str(isl_ctx
*ctx
, const char *str
)
4023 return isl_union_map_read_from_str(ctx
, str
);
4026 /* Return the union of all iteration domains of the gen->stmts[i].
4028 static __isl_give isl_union_set
*extract_domain(struct cuda_gen
*gen
)
4031 isl_union_set
*domain
;
4033 domain
= isl_union_set_empty(isl_set_get_space(gen
->context
));
4034 for (i
= 0; i
< gen
->n_stmts
; ++i
) {
4037 domain_i
= isl_set_copy(gen
->stmts
[i
].domain
);
4038 domain
= isl_union_set_union(domain
,
4039 isl_union_set_from_set(domain_i
));
4045 /* Information about the outermost tilable bands in the forest of bands.
4047 * tile_len and n_parallel are only sets on band_info structures
4048 * that correspond to outermost bands. For other bands (in particular,
4049 * ancestors of the outermost bands), n_parallal is set to 0.
4051 * prefix is the (padded) schedule leading up to the outermost tilable bands.
4053 * tile_first is the number of schedule dimensions in prefix.
4055 * suffix is the schedule of the outermost tilable bands and their descendants.
4058 struct cuda_gen
*gen
;
4062 isl_union_map
*prefix
;
4063 isl_union_map
*suffix
;
4066 /* Set tile_len and n_parallel of the statement to that of
4067 * their outermost band, recorded in the band_info.
4069 static int set_stmt_tile_len(__isl_take isl_map
*map
, void *user
)
4071 struct band_info
*info
= user
;
4073 struct cuda_stmt
*stmt
;
4075 nr
= atoi(isl_map_get_tuple_name(map
, isl_dim_in
) + 2);
4076 stmt
= &info
->gen
->stmts
[nr
];
4078 stmt
->tile_len
= info
->tile_len
;
4079 stmt
->n_parallel
= info
->n_parallel
;
4086 static void list_select_outer_band(struct cuda_gen
*gen
,
4087 __isl_take isl_band_list
*list
, int pos
, struct band_info
*list_info
);
4089 /* Check if this band has any parallel loops. If so, take it as
4090 * the outermost tilable band. If not, continue looking for the
4091 * outermost tilable band in the children of the current band.
4093 static void band_select_outer_band(struct cuda_gen
*gen
,
4094 __isl_take isl_band
*band
, int pos
, struct band_info
*info
)
4096 int n
= isl_band_n_member(band
);
4099 for (n_parallel
= 0; n_parallel
< n
; ++n_parallel
)
4100 if (!isl_band_member_is_zero_distance(band
, n_parallel
))
4103 info
->n_parallel
= n_parallel
;
4106 info
->tile_first
= pos
;
4108 info
->prefix
= isl_band_get_prefix_schedule(band
);
4109 info
->suffix
= isl_union_map_flat_range_product(
4110 isl_band_get_partial_schedule(band
),
4111 isl_band_get_suffix_schedule(band
));
4112 isl_union_map_foreach_map(info
->prefix
,
4113 &set_stmt_tile_len
, info
);
4114 } else if (isl_band_has_children(band
)) {
4115 isl_band_list
*children
;
4116 children
= isl_band_get_children(band
);
4117 list_select_outer_band(gen
, children
, pos
+ n
, info
);
4120 info
->tile_first
= pos
+ n
;
4122 info
->prefix
= isl_union_map_flat_range_product(
4123 isl_band_get_prefix_schedule(band
),
4124 isl_band_get_partial_schedule(band
));
4125 info
->suffix
= isl_band_get_suffix_schedule(band
);
4126 isl_union_map_foreach_map(info
->prefix
,
4127 &set_stmt_tile_len
, info
);
4130 isl_band_free(band
);
4133 /* Comparison function that returns a non-zero value for band_infos
4134 * with different tile_len fields or different n_parallel fields.
4136 static int cmp_band(const void *p1
, const void *p2
)
4138 const struct band_info
*info1
= p1
;
4139 const struct band_info
*info2
= p2
;
4141 if (info1
->tile_len
!= info2
->tile_len
)
4142 return info1
->tile_len
- info2
->tile_len
;
4144 return info1
->n_parallel
- info2
->n_parallel
;
4147 /* Extend "umap" with coordinates with fixed value "val"
4148 * to a total length of "dst_len", assuming the original dimension is "src_len".
4150 static __isl_give isl_union_map
*extend_range(__isl_take isl_union_map
*umap
,
4151 int src_len
, int dst_len
, int val
)
4157 dim
= isl_union_map_get_space(umap
);
4158 map
= isl_map_reverse(projection(dim
, dst_len
, src_len
));
4159 for (i
= src_len
; i
< dst_len
; ++i
)
4160 map
= isl_map_fix_si(map
, isl_dim_out
, i
, val
);
4162 umap
= isl_union_map_apply_range(umap
, isl_union_map_from_map(map
));
4167 /* Group bands with the same values for tile_len and n_parallel.
4168 * The prefix schedule is then extended with a fixed coordinate that
4169 * is different for each such group.
4170 * Note that the actual values for this coordinate are not important.
4171 * The bands have already been effectively separated at a higher level
4172 * or they are independent and may be executed in parallel.
4173 * The list of band_info has been sorted before this functions is called.
4175 static void separate_bands(struct band_info
*info
, int n
)
4180 for (i
= 0; i
< n
; ++i
) {
4181 int l
= info
[i
].tile_first
;
4184 (info
[i
].tile_len
!= info
[i
- 1].tile_len
||
4185 info
[i
].n_parallel
!= info
[i
- 1].n_parallel
))
4188 info
[i
].prefix
= extend_range(info
[i
].prefix
,
4190 info
[i
].tile_first
= l
+ 1;
4194 /* Select the outermost bands in the elements of the list, align
4195 * their prefix schedules, separate bands with different values
4196 * for tile_len and/or n_parallel and then combine the resulting
4197 * prefix and suffix schedules into a single pair of prefix and
4198 * suffix schedules for the entire list.
4200 static void list_select_outer_band(struct cuda_gen
*gen
,
4201 __isl_take isl_band_list
*list
, int pos
, struct band_info
*list_info
)
4205 int n
= isl_band_list_n_band(list
);
4206 isl_ctx
*ctx
= isl_band_list_get_ctx(list
);
4207 struct band_info
*info
;
4209 isl_union_map
*prefix
;
4210 isl_union_map
*suffix
;
4213 info
= isl_calloc_array(ctx
, struct band_info
, n
);
4217 for (i
= 0; i
< n
; ++i
) {
4218 band
= isl_band_list_get_band(list
, i
);
4219 band_select_outer_band(gen
, band
, pos
, &info
[i
]);
4220 if (info
[i
].tile_first
> max_tile_first
)
4221 max_tile_first
= info
[i
].tile_first
;
4224 for (i
= 0; i
< n
; ++i
) {
4225 if (info
[i
].tile_first
== max_tile_first
)
4227 info
[i
].prefix
= extend_range(info
[i
].prefix
,
4228 info
[i
].tile_first
, max_tile_first
, 0);
4229 info
[i
].tile_first
= max_tile_first
;
4232 qsort(info
, n
, sizeof(struct band_info
), &cmp_band
);
4234 for (i
= 0; i
< n
- 1; ++i
)
4235 if (info
[i
].tile_len
!= info
[i
+ 1].tile_len
||
4236 info
[i
].n_parallel
!= info
[i
+ 1].n_parallel
)
4240 separate_bands(info
, n
);
4242 prefix
= info
[0].prefix
;
4243 suffix
= info
[0].suffix
;
4245 for (i
= 1; i
< n
; ++i
) {
4246 prefix
= isl_union_map_union(prefix
, info
[i
].prefix
);
4247 suffix
= isl_union_map_union(suffix
, info
[i
].suffix
);
4250 list_info
->tile_first
= info
[0].tile_first
;
4251 list_info
->tile_len
= -1;
4252 list_info
->prefix
= prefix
;
4253 list_info
->suffix
= suffix
;
4255 isl_band_list_free(list
);
4259 /* Set max_out to the maximal number of output dimensions over
4262 static int update_max_out(__isl_take isl_map
*map
, void *user
)
4264 int *max_out
= user
;
4265 int n_out
= isl_map_dim(map
, isl_dim_out
);
4267 if (n_out
> *max_out
)
4274 struct align_range_data
{
4279 /* Extend the dimension of the range of the given map to data->max_out and
4280 * then add the result to data->res.
4282 static int map_align_range(__isl_take isl_map
*map
, void *user
)
4284 struct align_range_data
*data
= user
;
4288 int n_out
= isl_map_dim(map
, isl_dim_out
);
4290 dim
= isl_union_map_get_space(data
->res
);
4291 proj
= isl_map_reverse(projection(dim
, data
->max_out
, n_out
));
4292 for (i
= n_out
; i
< data
->max_out
; ++i
)
4293 proj
= isl_map_fix_si(proj
, isl_dim_out
, i
, 0);
4295 map
= isl_map_apply_range(map
, proj
);
4297 data
->res
= isl_union_map_add_map(data
->res
, map
);
4302 /* Extend the ranges of the maps in the union map such they all have
4303 * the same dimension.
4305 static __isl_give isl_union_map
*align_range(__isl_take isl_union_map
*umap
)
4307 struct align_range_data data
;
4310 isl_union_map_foreach_map(umap
, &update_max_out
, &data
.max_out
);
4312 data
.res
= isl_union_map_empty(isl_union_map_get_space(umap
));
4313 isl_union_map_foreach_map(umap
, &map_align_range
, &data
);
4315 isl_union_map_free(umap
);
4319 /* Select the outermost tilable band that (by construction)
4320 * has at least one parallel loop.
4321 * The starting position of the aligned band is stored in the pair
4323 * The sizes and number of parallel loops may be different in different
4324 * parts of the band forest and are therefore stored in the cuda_stmts.
4326 * Return the complete schedule, with the tilable bands aligned
4327 * at gen->tile_first and padded with zero, if needed.
4329 static __isl_give isl_union_map
*select_outer_tilable_band(struct cuda_gen
*gen
,
4330 __isl_keep isl_schedule
*schedule
)
4332 isl_band_list
*list
;
4333 struct band_info info
;
4335 gen
->n_parallel
= 0;
4338 list
= isl_schedule_get_band_forest(schedule
);
4340 list_select_outer_band(gen
, list
, 0, &info
);
4342 gen
->tile_first
= info
.tile_first
;
4343 info
.suffix
= align_range(info
.suffix
);
4345 return isl_union_map_flat_range_product(info
.prefix
, info
.suffix
);
4348 /* Set gen->untiled_len to the number of scheduling dimensions
4349 * for the schedule of the first domain.
4350 * We assume here that this number is the same for all domains.
4352 static int set_untiled_len(__isl_take isl_map
*map
, void *user
)
4354 unsigned *untiled_len
= user
;
4356 *untiled_len
= isl_map_dim(map
, isl_dim_out
);
4362 /* Compute an appropriate schedule based on the accesses in
4363 * gen->read and gen->write.
4365 * We first compute dependences and then use those to compute
4366 * a schedule that has a parallel loop in each tilable band.
4367 * Finally, we select the outermost tilable band.
4369 static void compute_schedule(struct cuda_gen
*gen
,
4370 __isl_take isl_union_map
*sched
)
4372 isl_union_set
*domain
;
4373 isl_union_map
*empty
;
4374 isl_union_map
*dep_raw
, *dep2
, *dep3
, *dep
;
4375 isl_union_map
*uninitialized
;
4376 isl_schedule
*schedule
;
4378 empty
= isl_union_map_empty(isl_union_map_get_space(sched
));
4380 isl_union_map_compute_flow(isl_union_map_copy(gen
->read
),
4381 isl_union_map_copy(gen
->write
), empty
,
4382 isl_union_map_copy(sched
),
4383 &dep_raw
, NULL
, &uninitialized
, NULL
);
4384 isl_union_map_compute_flow(isl_union_map_copy(gen
->write
),
4385 isl_union_map_copy(gen
->write
),
4386 isl_union_map_copy(gen
->read
),
4387 isl_union_map_copy(sched
),
4388 &dep2
, &dep3
, NULL
, NULL
);
4389 isl_union_map_free(sched
);
4391 gen
->copy_in
= isl_union_map_range(uninitialized
);
4393 dep
= isl_union_map_union(dep2
, dep3
);
4394 dep
= isl_union_map_union(dep
, dep_raw
);
4395 dep
= isl_union_map_coalesce(dep
);
4397 domain
= extract_domain(gen
);
4398 schedule
= isl_union_set_compute_schedule(isl_union_set_copy(domain
),
4399 isl_union_map_copy(dep
), dep
);
4401 sched
= select_outer_tilable_band(gen
, schedule
);
4403 isl_union_map_foreach_map(sched
, &set_untiled_len
, &gen
->untiled_len
);
4404 sched
= isl_union_map_intersect_domain(sched
, domain
);
4407 isl_schedule_free(schedule
);
4410 static struct cuda_stmt_access
**expr_extract_access(struct pet_expr
*expr
,
4411 struct cuda_stmt_access
**next_access
)
4413 struct cuda_stmt_access
*access
;
4414 isl_ctx
*ctx
= isl_map_get_ctx(expr
->acc
.access
);
4416 access
= isl_alloc_type(ctx
, struct cuda_stmt_access
);
4418 access
->next
= NULL
;
4419 access
->read
= expr
->acc
.read
;
4420 access
->write
= expr
->acc
.write
;
4421 access
->access
= isl_map_copy(expr
->acc
.access
);
4423 *next_access
= access
;
4424 next_access
= &(*next_access
)->next
;
4428 static struct cuda_stmt_access
**expr_extract_accesses(struct pet_expr
*expr
,
4429 struct cuda_stmt_access
**next_access
)
4433 for (i
= 0; i
< expr
->n_arg
; ++i
)
4434 next_access
= expr_extract_accesses(expr
->args
[i
],
4437 if (expr
->type
== pet_expr_access
)
4438 next_access
= expr_extract_access(expr
, next_access
);
4443 static void pet_stmt_extract_accesses(struct cuda_stmt
*stmt
)
4445 struct cuda_stmt_access
**next_access
= &stmt
->accesses
;
4447 stmt
->accesses
= NULL
;
4448 expr_extract_accesses(stmt
->body
, next_access
);
4451 /* Return an array of cuda_stmt representing the statements in "scop".
4453 static struct cuda_stmt
*extract_stmts(isl_ctx
*ctx
, struct pet_scop
*scop
,
4454 __isl_keep isl_set
*context
)
4457 struct cuda_stmt
*stmts
;
4459 stmts
= isl_calloc_array(ctx
, struct cuda_stmt
, scop
->n_stmt
);
4462 for (i
= 0; i
< scop
->n_stmt
; ++i
) {
4463 struct cuda_stmt
*s
= &stmts
[i
];
4465 s
->domain
= isl_set_copy(scop
->stmts
[i
]->domain
);
4466 s
->domain
= isl_set_intersect_params(s
->domain
,
4467 isl_set_copy(context
));
4468 s
->body
= scop
->stmts
[i
]->body
;
4469 pet_stmt_extract_accesses(s
);
4475 /* Replace the scop in the "input" file by equivalent code
4476 * that uses the GPU. "scop" is assumed to correspond to this scop.
4478 * We first compute a schedule that respects the dependences
4479 * of the original program and select the outermost band
4480 * of tilable dimensions that has at least one parallel loop.
4481 * We then have three blocks of dimensions
4485 * The tilable band "B" is first tiled according to "tile" sizes, resulting
4490 * For each iteration of the T loop and for each array, we compute
4491 * the array elements accessed by that iteration, construct a rectangular
4492 * box around it and shift it to the origin. The result is used
4493 * as shared memory for the array.
4495 * We then split off at most 2 parallel loops from the T loops and
4496 * at most 3 parallel loops from the P loops
4500 * The T1/P1 loops are then tiled or "wrapped" over the blocks/threads,
4501 * according to "grid"/"block" sizes.
4503 * H T1T T1P T2 P1T P1P P2 G
4505 * Finally, the T1P and P1P iterators are equated to the block and
4506 * thread dimensions respectively and so are effectively removed.
4507 * The H loops are run on the host. The T1T, T2, P1T, P2 and G loops
4508 * are run on the GPU.
4510 * Code is generated in three stages. We first generate code for the
4511 * host (the H loops), with iterators h%d. Then, for each leaf node
4512 * of the resulting AST, we generate code for the shared loops (up to
4513 * and including T2), with iterators g%d and after equating the H loops
4514 * to h%d parameters and the T1P loops to the block dimensions.
4515 * Finally, we generate code for the remaining loops in a similar fashion.
4517 int generate_cuda(isl_ctx
*ctx
, struct pet_scop
*scop
,
4518 struct ppcg_options
*options
, const char *input
)
4520 isl_union_map
*sched
;
4521 struct cuda_gen gen
;
4526 scop
= pet_scop_align_params(scop
);
4529 gen
.context
= isl_set_copy(scop
->context
);
4530 gen
.context
= add_context_from_str(gen
.context
, options
->ctx
);
4531 gen
.sizes
= extract_sizes_from_str(ctx
, options
->sizes
);
4532 gen
.n_stmts
= scop
->n_stmt
;
4533 gen
.stmts
= extract_stmts(ctx
, scop
, gen
.context
);
4534 gen
.read
= pet_scop_collect_reads(scop
);
4535 gen
.write
= pet_scop_collect_writes(scop
);
4536 gen
.options
= options
;
4537 gen
.state
= cloog_isl_state_malloc(gen
.ctx
);
4540 cuda_open_files(&gen
.cuda
, input
);
4542 collect_array_info(&gen
);
4544 sched
= pet_scop_collect_schedule(scop
);
4546 compute_schedule(&gen
, sched
);
4548 print_host_code(&gen
);
4550 cloog_state_free(gen
.state
);
4551 clear_cuda_gen(&gen
);
4553 cuda_close_files(&gen
.cuda
);