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 /* Given a parameter space "space", create a set of dimension "len"
986 * of which the "n" dimensions starting at "first" are equated to
987 * freshly created parameters named prefix%d.
989 static __isl_give isl_set
*parametrization(__isl_take isl_space
*space
,
990 int len
, int first
, int n
, const char *prefix
)
994 space
= isl_space_set_from_params(space
);
995 space
= isl_space_add_dims(space
, isl_dim_set
, len
);
996 set
= isl_set_universe(space
);
998 return parametrize(set
, first
, n
, prefix
);
1001 /* Tile the B loops over the tile sizes and then tile/wrap
1002 * the T1 loops over the blocks.
1004 static __isl_give isl_union_map
*tile_schedule(struct cuda_gen
*gen
,
1005 __isl_take isl_union_map
*sched
)
1008 isl_map
*tiling
, *block_tiling
;
1010 dim
= isl_union_map_get_space(sched
);
1011 tiling
= tile(isl_space_copy(dim
), gen
->untiled_len
,
1012 gen
->tile_first
, gen
->tile_len
, gen
->tile_size
);
1014 if (gen
->options
->wrap
)
1015 block_tiling
= wrap(dim
, gen
->untiled_len
+ gen
->tile_len
,
1016 gen
->tile_first
, gen
->n_grid
, gen
->grid_dim
);
1018 block_tiling
= tile(dim
, gen
->untiled_len
+ gen
->tile_len
,
1019 gen
->tile_first
, gen
->n_grid
, gen
->grid_dim
);
1021 gen
->tiled_len
= gen
->untiled_len
+ gen
->tile_len
+ gen
->n_grid
;
1023 tiling
= isl_map_apply_range(tiling
, block_tiling
);
1025 sched
= isl_union_map_apply_range(sched
,
1026 isl_union_map_from_map(tiling
));
1028 gen
->shared_len
= gen
->tile_first
+ gen
->tile_len
+ gen
->n_grid
;
1033 /* Equate the "T1P" iterators in the tiled schedule "sched"
1034 * to the block dimensions.
1036 static __isl_give isl_union_map
*parametrize_tiled_schedule(
1037 struct cuda_gen
*gen
, __isl_take isl_union_map
*sched
)
1042 dim
= isl_union_map_get_space(sched
);
1043 par
= parametrization(dim
, gen
->tiled_len
, 0, gen
->tile_first
, "h");
1044 sched
= isl_union_map_intersect_range(sched
,
1045 isl_union_set_from_set(par
));
1047 dim
= isl_union_map_get_space(sched
);
1048 par
= parametrization(dim
, gen
->tiled_len
,
1049 gen
->tile_first
+ gen
->n_grid
, gen
->n_grid
, "b");
1050 sched
= isl_union_map_intersect_range(sched
,
1051 isl_union_set_from_set(par
));
1056 /* Tile/wrap the P1 loops over the threads.
1058 static __isl_give isl_union_map
*thread_tile_schedule(struct cuda_gen
*gen
,
1059 __isl_take isl_union_map
*sched
)
1065 dim
= isl_union_map_get_space(sched
);
1067 if (gen
->options
->wrap
)
1068 tiling
= wrap(isl_space_copy(dim
), gen
->tiled_len
,
1069 gen
->shared_len
, gen
->n_block
, gen
->block_dim
);
1071 tiling
= tile(isl_space_copy(dim
), gen
->tiled_len
,
1072 gen
->shared_len
, gen
->n_block
, gen
->block_dim
);
1073 gen
->thread_tiled_len
= gen
->tiled_len
+ gen
->n_block
;
1075 sched
= isl_union_map_apply_range(sched
,
1076 isl_union_map_from_map(tiling
));
1078 par
= parametrization(dim
, gen
->thread_tiled_len
,
1079 gen
->tile_first
+ gen
->tile_len
+ gen
->n_grid
+ gen
->n_block
,
1081 sched
= isl_union_map_intersect_range(sched
,
1082 isl_union_set_from_set(par
));
1084 gen
->shared_len
= gen
->tile_first
+ gen
->tile_len
+ gen
->n_grid
;
1089 /* If the user asked for it, scale the shared memory tile loops
1090 * (T1T and T2) of "sched" by gen->tile_size[i].
1091 * If we are not performing "wrapping", then additionally scale the T1P
1092 * loops by gen->grid_dim[i].
1094 static __isl_give isl_union_map
*scale_tile_loops(struct cuda_gen
*gen
,
1095 __isl_take isl_union_map
*sched
)
1099 isl_basic_map
*scale
;
1101 isl_local_space
*ls
;
1103 if (!gen
->options
->scale_tile_loops
)
1106 dim
= isl_union_map_get_space(sched
);
1107 dim
= isl_space_add_dims(dim
, isl_dim_in
, gen
->tiled_len
);
1108 dim
= isl_space_add_dims(dim
, isl_dim_out
, gen
->tiled_len
);
1109 scale
= isl_basic_map_universe(isl_space_copy(dim
));
1110 ls
= isl_local_space_from_space(dim
);
1112 for (i
= 0; i
< gen
->tiled_len
; ++i
) {
1115 if (i
>= gen
->tile_first
&& i
< gen
->tile_first
+ gen
->n_grid
) {
1116 f
= gen
->tile_size
[i
- gen
->tile_first
];
1117 if (!gen
->options
->wrap
)
1118 f
*= gen
->grid_dim
[i
- gen
->tile_first
];
1119 } else if (i
>= gen
->tile_first
+ gen
->n_grid
&&
1120 i
< gen
->tile_first
+ gen
->n_grid
+ gen
->tile_len
) {
1121 f
= gen
->tile_size
[i
- (gen
->tile_first
+ gen
->n_grid
)];
1124 c
= isl_equality_alloc(isl_local_space_copy(ls
));
1125 isl_constraint_set_coefficient_si(c
, isl_dim_in
, i
, f
);
1126 isl_constraint_set_coefficient_si(c
, isl_dim_out
, i
, -1);
1127 scale
= isl_basic_map_add_constraint(scale
, c
);
1130 isl_local_space_free(ls
);
1132 sched
= isl_union_map_apply_range(sched
,
1133 isl_union_map_from_map(isl_map_from_basic_map(scale
)));
1138 /* If we are not performing "wrapping" and if the user asked for it,
1139 * scale the thread tile loops (P1T) of "sched" by gen->block_dim[i].
1141 static __isl_give isl_union_map
*scale_thread_tile_loops(struct cuda_gen
*gen
,
1142 __isl_take isl_union_map
*sched
)
1146 isl_basic_map
*scale
;
1148 isl_local_space
*ls
;
1150 if (gen
->options
->wrap
)
1152 if (!gen
->options
->scale_tile_loops
)
1155 dim
= isl_union_map_get_space(sched
);
1156 dim
= isl_space_add_dims(dim
, isl_dim_in
, gen
->thread_tiled_len
);
1157 dim
= isl_space_add_dims(dim
, isl_dim_out
, gen
->thread_tiled_len
);
1158 scale
= isl_basic_map_universe(isl_space_copy(dim
));
1159 ls
= isl_local_space_from_space(dim
);
1161 for (i
= 0; i
< gen
->thread_tiled_len
; ++i
) {
1164 if (i
>= gen
->shared_len
&&
1165 i
< gen
->shared_len
+ gen
->n_block
)
1166 f
= gen
->block_dim
[i
- gen
->shared_len
];
1168 c
= isl_equality_alloc(isl_local_space_copy(ls
));
1169 isl_constraint_set_coefficient_si(c
, isl_dim_in
, i
, f
);
1170 isl_constraint_set_coefficient_si(c
, isl_dim_out
, i
, -1);
1171 scale
= isl_basic_map_add_constraint(scale
, c
);
1174 isl_local_space_free(ls
);
1176 sched
= isl_union_map_apply_range(sched
,
1177 isl_union_map_from_map(isl_map_from_basic_map(scale
)));
1182 /* If we are not performing "wrapping" and if the user asked for it,
1183 * scale the "n_tile" loops starting at "first" of "sched" by gen->block_dim[i].
1185 static __isl_give isl_union_map
*scale_access_tile_loops(struct cuda_gen
*gen
,
1186 __isl_take isl_union_map
*sched
, int len
, int first
, int n_tile
)
1190 isl_basic_map
*scale
;
1192 isl_local_space
*ls
;
1194 if (gen
->options
->wrap
)
1196 if (!gen
->options
->scale_tile_loops
)
1199 dim
= isl_union_map_get_space(sched
);
1200 dim
= isl_space_add_dims(dim
, isl_dim_in
, len
);
1201 dim
= isl_space_add_dims(dim
, isl_dim_out
, len
);
1202 scale
= isl_basic_map_universe(isl_space_copy(dim
));
1203 ls
= isl_local_space_from_space(dim
);
1205 for (i
= 0; i
< len
; ++i
) {
1208 if (i
>= first
&& i
< first
+ n_tile
)
1209 f
= gen
->block_dim
[i
- first
];
1211 c
= isl_equality_alloc(isl_local_space_copy(ls
));
1212 isl_constraint_set_coefficient_si(c
, isl_dim_in
, i
, f
);
1213 isl_constraint_set_coefficient_si(c
, isl_dim_out
, i
, -1);
1214 scale
= isl_basic_map_add_constraint(scale
, c
);
1217 isl_local_space_free(ls
);
1219 sched
= isl_union_map_apply_range(sched
,
1220 isl_union_map_from_map(isl_map_from_basic_map(scale
)));
1225 /* If print_user_stmt is set, we want to print the statements ourselves,
1226 * instead of relying on the C preprocessor. If so, we need to use
1227 * the stop option so that the domains will be saved on the statement
1230 static void print_cloog_shared_body(struct cuda_gen
*gen
,
1231 __isl_keep isl_set
*context
, __isl_keep isl_union_map
*sched
, int len
,
1232 void (*print_user_stmt
)(struct clast_printer_info
*info
,
1233 struct clast_user_stmt
*s
),
1237 CloogOptions
*options
;
1238 CloogDomain
*cloog_context
;
1239 CloogUnionDomain
*ud
;
1241 struct clast_stmt
*stmt
;
1244 sched
= isl_union_map_copy(sched
);
1245 sched
= isl_union_map_align_params(sched
, isl_set_get_space(context
));
1247 options
= cloog_options_malloc(gen
->state
);
1248 options
->language
= CLOOG_LANGUAGE_C
;
1249 options
->strides
= 1;
1253 options
->override
= 1;
1254 options
->save_domains
= 1;
1255 options
->noscalars
= 1;
1256 options
->first_unroll
= first_unroll
;
1258 ud
= cloog_union_domain_from_isl_union_map(sched
);
1259 for (i
= 0; i
< len
; ++i
) {
1260 snprintf(name
, sizeof(name
), "c%d", i
);
1261 ud
= cloog_union_domain_set_name(ud
, CLOOG_SCAT
, i
, name
);
1263 cloog_context
= cloog_domain_from_isl_set(isl_set_copy(context
));
1264 input
= cloog_input_alloc(cloog_context
, ud
);
1266 stmt
= cloog_clast_create_from_input(input
, options
);
1268 gen
->stmt_code
.indent
= gen
->kernel_code
.indent
;
1269 gen
->stmt_code
.dst
= gen
->cuda
.kernel_c
;
1270 gen
->stmt_code
.print_user_stmt
= print_user_stmt
;
1271 gen
->stmt_code
.print_user_stmt_list
= NULL
;
1272 gen
->stmt_code
.print_for_head
= NULL
;
1273 gen
->stmt_code
.print_for_foot
= NULL
;
1274 gen
->stmt_code
.user
= gen
;
1275 print_clast(&gen
->stmt_code
, stmt
);
1277 cloog_clast_free(stmt
);
1278 cloog_options_free(options
);
1281 /* Add "len" parameters p[i] called prefix%d,
1282 * with bounds to 0 <= p[i] < size[i].
1284 __isl_give isl_set
*add_bounded_parameters(__isl_take isl_set
*set
,
1285 int len
, int *size
, const char *prefix
)
1291 isl_basic_set
*bset
;
1293 isl_local_space
*ls
;
1296 nparam
= isl_set_dim(set
, isl_dim_param
);
1297 set
= isl_set_add_dims(set
, isl_dim_param
, len
);
1299 for (i
= 0; i
< len
; ++i
) {
1300 snprintf(name
, sizeof(name
), "%s%d", prefix
, i
);
1301 set
= isl_set_set_dim_name(set
, isl_dim_param
,
1305 dim
= isl_set_get_space(set
);
1306 bset
= isl_basic_set_universe(isl_space_copy(dim
));
1307 ls
= isl_local_space_from_space(dim
);
1311 for (i
= 0; i
< len
; ++i
) {
1312 c
= isl_inequality_alloc(isl_local_space_copy(ls
));
1313 isl_int_set_si(v
, 1);
1314 isl_constraint_set_coefficient(c
, isl_dim_param
, nparam
+ i
, v
);
1315 bset
= isl_basic_set_add_constraint(bset
, c
);
1317 c
= isl_inequality_alloc(isl_local_space_copy(ls
));
1318 isl_int_set_si(v
, -1);
1319 isl_constraint_set_coefficient(c
, isl_dim_param
, nparam
+ i
, v
);
1320 isl_int_set_si(v
, size
[i
] - 1);
1321 isl_constraint_set_constant(c
, v
);
1322 bset
= isl_basic_set_add_constraint(bset
, c
);
1326 isl_local_space_free(ls
);
1328 return isl_set_intersect(set
, isl_set_from_basic_set(bset
));
1331 static void print_shared_body(struct cuda_gen
*gen
,
1332 __isl_keep isl_set
*shared_domain
, __isl_keep isl_union_map
*sched
,
1333 int len
, void (*print_user_stmt
)(struct clast_printer_info
*info
,
1334 struct clast_user_stmt
*s
),
1339 context
= isl_set_copy(shared_domain
);
1340 context
= parametrize(context
, 0, gen
->shared_len
, "g");
1341 context
= isl_set_project_out(context
, isl_dim_set
, 0, gen
->shared_len
);
1342 context
= add_bounded_parameters(context
,
1343 gen
->n_block
, gen
->block_dim
, "t");
1345 print_cloog_shared_body(gen
, context
, sched
, len
, print_user_stmt
,
1348 isl_set_free(context
);
1351 /* Given a tile of an array, construct a map that maps each element
1352 * of the tile to a copy of the tile shifted to the origin
1353 * (based on the lower bounds in group->private_bound or group->shared_bound).
1354 * If any of the indices is strided, then {private,shared}_bound[i].shift_map
1355 * is applied to the index first.
1356 * The domain of the resulting map is "access",
1357 * while the range space is anonymous.
1359 static __isl_give isl_map
*shift_access(__isl_take isl_set
*access
,
1360 struct cuda_array_ref_group
*group
)
1364 isl_basic_set
*bset
;
1365 isl_basic_map
*bmap
;
1367 isl_basic_set
*offset
;
1368 isl_basic_map
*shift
;
1369 isl_basic_map
*pre_shift
;
1372 struct cuda_array_bound
*bounds
;
1373 int n_index
= group
->array
->n_index
;
1375 bounds
= group
->private_bound
;
1377 bounds
= group
->shared_bound
;
1379 dim
= isl_set_get_space(access
);
1380 dim
= isl_space_drop_dims(dim
, isl_dim_set
, 0, n_index
);
1381 offset
= isl_basic_set_universe(dim
);
1382 for (i
= 0; i
< n_index
; ++i
) {
1383 lb
= isl_aff_copy(bounds
[i
].lb
);
1384 bmap
= isl_basic_map_from_aff(lb
);
1385 bset
= isl_basic_map_range(bmap
);
1386 offset
= isl_basic_set_flat_product(offset
, bset
);
1388 offset
= isl_basic_set_neg(offset
);
1390 dim
= isl_space_map_from_set(isl_set_get_space(access
));
1391 shift
= isl_basic_map_identity(dim
);
1392 shift
= isl_basic_map_set_tuple_name(shift
, isl_dim_out
, NULL
);
1394 bset
= isl_basic_set_universe(isl_set_get_space(access
));
1395 bmap
= isl_basic_map_from_domain_and_range(bset
, offset
);
1397 shift
= isl_basic_map_sum(shift
, bmap
);
1399 dim
= isl_set_get_space(access
);
1400 dim
= isl_space_drop_dims(dim
, isl_dim_set
, 0, n_index
);
1401 dim
= isl_space_map_from_set(dim
);
1402 pre_shift
= isl_basic_map_universe(isl_space_copy(dim
));
1403 dim
= isl_space_add_dims(dim
, isl_dim_in
, 1);
1404 dim
= isl_space_add_dims(dim
, isl_dim_out
, 1);
1405 for (i
= 0; i
< n_index
; ++i
) {
1406 if (!bounds
[i
].shift_map
)
1407 bmap
= isl_basic_map_identity(isl_space_copy(dim
));
1409 bmap
= isl_basic_map_copy(bounds
[i
].shift_map
);
1410 pre_shift
= isl_basic_map_flat_product(pre_shift
, bmap
);
1412 isl_space_free(dim
);
1413 name
= isl_basic_map_get_tuple_name(shift
, isl_dim_in
);
1414 pre_shift
= isl_basic_map_set_tuple_name(pre_shift
, isl_dim_in
, name
);
1415 pre_shift
= isl_basic_map_set_tuple_name(pre_shift
, isl_dim_out
, name
);
1416 shift
= isl_basic_map_apply_range(pre_shift
, shift
);
1418 sched
= isl_map_from_basic_map(shift
);
1419 sched
= isl_map_intersect_domain(sched
, access
);
1424 /* Construct a schedule for iterating over all elements in the given
1425 * piece of an array. The schedule iterates over a copy of the piece
1426 * that is shifted to the origin.
1427 * We subsequently also perform the tiling/wrapping over the threads.
1429 * In particular, we tile the final iterators so that the final thread
1430 * dimension runs over the final array dimension.
1431 * However, if those final iterators have only a single iteration,
1432 * we try to tile earlier iterators instead.
1434 static __isl_give isl_union_map
*access_schedule(struct cuda_gen
*gen
,
1435 __isl_take isl_set
*access
, struct cuda_array_ref_group
*group
)
1439 isl_union_map
*usched
;
1442 unsigned nvar
= isl_set_dim(access
, isl_dim_set
);
1446 sched
= shift_access(access
, group
);
1448 n_tile
= gen
->n_block
;
1449 if (n_tile
> nvar
) {
1451 sched
= isl_map_insert_dims(sched
,
1452 isl_dim_out
, 0, n_tile
- nvar
);
1453 for (i
= 0; i
< n_tile
- nvar
; ++i
)
1454 sched
= isl_map_fix_si(sched
, isl_dim_out
, i
, 0);
1458 first
= nvar
- n_tile
;
1460 for (; first
> 0; first
--)
1461 if (!isl_map_plain_is_fixed(sched
, isl_dim_out
,
1462 first
+ n_tile
- 1, NULL
))
1465 dim
= isl_map_get_space(sched
);
1466 dim
= isl_space_params(dim
);
1467 if (gen
->options
->wrap
)
1468 tiling
= wrap(isl_space_copy(dim
), nvar
, first
,
1469 n_tile
, gen
->block_dim
);
1471 tiling
= tile(isl_space_copy(dim
), nvar
, first
,
1472 n_tile
, gen
->block_dim
);
1473 sched
= isl_map_apply_range(sched
, tiling
);
1475 par
= parametrization(dim
, nvar
+ n_tile
, first
+ n_tile
, n_tile
, "t");
1476 sched
= isl_map_intersect_range(sched
, par
);
1478 usched
= isl_union_map_from_map(sched
);
1479 usched
= scale_access_tile_loops(gen
, usched
, nvar
+ n_tile
,
1485 /* Print an access to the element in the global memory copy of the
1486 * given array that corresponds to the element described by "pma".
1487 * of the original array.
1488 * The copy in global memory has been linearized, so we need to take
1489 * the array size into account.
1491 static void print_global_index(FILE *out
,
1492 struct cuda_array_info
*array
, __isl_keep isl_pw_multi_aff
*pma
,
1493 __isl_keep isl_set
*domain
)
1496 isl_ctx
*ctx
= isl_pw_multi_aff_get_ctx(pma
);
1499 if (cuda_array_is_scalar(array
)) {
1500 fprintf(out
, "*%s", array
->name
);
1504 fprintf(out
, "%s[", array
->name
);
1505 prn
= isl_printer_to_file(ctx
, out
);
1506 prn
= isl_printer_set_output_format(prn
, ISL_FORMAT_C
);
1507 for (i
= 0; i
+ 1 < array
->n_index
; ++i
)
1508 prn
= isl_printer_print_str(prn
, "(");
1509 for (i
= 0; i
< array
->n_index
; ++i
) {
1510 isl_pw_aff
*pa
= isl_pw_multi_aff_get_pw_aff(pma
, i
);
1511 pa
= isl_pw_aff_coalesce(pa
);
1512 pa
= isl_pw_aff_gist(pa
, isl_set_copy(domain
));
1514 prn
= isl_printer_print_str(prn
, ") * (");
1515 prn
= isl_printer_print_pw_aff(prn
,
1516 array
->local_bound
[i
]);
1517 prn
= isl_printer_print_str(prn
, ") + ");
1519 prn
= isl_printer_print_pw_aff(prn
, pa
);
1520 isl_pw_aff_free(pa
);
1522 isl_printer_free(prn
);
1526 /* Given an index expression into a tile of an array, adjust the expression
1527 * to a shift of the tile to the origin
1528 * (based on the lower bounds in array->shared_bound).
1529 * If the index is strided, then we first add
1530 * bound->shift and divide by bound->stride.
1532 static __isl_give isl_pw_aff
*shift_index(__isl_take isl_pw_aff
*pa
,
1533 struct cuda_array_info
*array
,
1534 struct cuda_array_bound
*bound
, __isl_take isl_set
*domain
)
1541 shift
= bound
->shift
;
1542 shift
= isl_aff_copy(shift
);
1543 shift
= isl_aff_project_domain_on_params(shift
);
1544 shift
= isl_aff_align_params(shift
, isl_pw_aff_get_space(pa
));
1545 tmp
= isl_pw_aff_alloc(isl_set_copy(domain
), shift
);
1546 pa
= isl_pw_aff_add(pa
, tmp
);
1547 pa
= isl_pw_aff_scale_down(pa
, bound
->stride
);
1550 lb
= isl_aff_copy(bound
->lb
);
1551 lb
= isl_aff_project_domain_on_params(lb
);
1553 lb
= isl_aff_align_params(lb
, isl_pw_aff_get_space(pa
));
1555 tmp
= isl_pw_aff_alloc(isl_set_copy(domain
), lb
);
1556 pa
= isl_pw_aff_sub(pa
, tmp
);
1557 pa
= isl_pw_aff_coalesce(pa
);
1558 pa
= isl_pw_aff_gist(pa
, domain
);
1563 /* Print an access to the element in the private/shared memory copy of the
1564 * given array reference group that corresponds to the element described
1565 * by "pma" of the original array.
1566 * Since the array in private/shared memory is just a shifted copy of part
1567 * of the original array, we simply need to subtract the lower bound,
1568 * which was computed in can_tile_for_shared_memory.
1569 * If any of the indices is strided, then we first add
1570 * bounds[i].shift and divide by bounds[i].stride.
1572 static void print_local_index(FILE *out
,
1573 struct cuda_array_ref_group
*group
, struct cuda_array_bound
*bounds
,
1574 __isl_keep isl_pw_multi_aff
*pma
, __isl_keep isl_set
*domain
)
1577 isl_ctx
*ctx
= isl_pw_multi_aff_get_ctx(pma
);
1579 struct cuda_array_info
*array
= group
->array
;
1581 prn
= isl_printer_to_file(ctx
, out
);
1582 prn
= isl_printer_set_output_format(prn
, ISL_FORMAT_C
);
1584 prn
= print_array_name(prn
, group
);
1585 for (i
= 0; i
< array
->n_index
; ++i
) {
1586 isl_pw_aff
*pa
= isl_pw_multi_aff_get_pw_aff(pma
, i
);
1588 pa
= shift_index(pa
, array
, &bounds
[i
], isl_set_copy(domain
));
1591 prn
= isl_printer_print_pw_aff(prn
, pa
);
1593 isl_pw_aff_free(pa
);
1596 isl_printer_free(prn
);
1599 /* This function is called for each leaf in the clast of the code
1600 * for copying to or from shared/private memory.
1601 * The statement name is {read,write}_{shared,private}_<array>.
1603 * The schedule iterates over the array elements, so we can use
1604 * the domain of copy_sched at the current scheduling position
1605 * as the index of the array.
1607 static void print_copy_statement(struct clast_printer_info
*code
,
1608 struct clast_user_stmt
*u
)
1610 struct cuda_gen
*gen
= code
->user
;
1613 struct cuda_array_ref_group
*group
= gen
->copy_group
;
1614 struct cuda_array_bound
*bounds
= gen
->copy_bound
;
1619 isl_pw_multi_aff
*pma
;
1622 read
= !strncmp(u
->statement
->name
, "read", 4);
1624 domain
= extract_host_domain(u
);
1627 sched
= isl_map_copy(gen
->copy_sched
);
1628 sched
= isl_map_reverse(sched
);
1629 sched
= isl_map_intersect_domain(sched
, domain
);
1630 n_in
= isl_map_dim(sched
, isl_dim_in
);
1631 dim
= isl_map_get_space(sched
);
1632 dim
= isl_space_params(dim
);
1633 param
= parametrization(dim
, n_in
, 0, n_in
, "c");
1634 sched
= isl_map_align_params(sched
, isl_set_get_space(param
));
1635 sched
= isl_map_intersect_domain(sched
, param
);
1636 index
= isl_map_range(sched
);
1637 domain
= isl_set_copy(index
);
1638 pma
= isl_pw_multi_aff_from_set(index
);
1639 pma
= isl_pw_multi_aff_coalesce(pma
);
1640 domain
= isl_set_params(domain
);
1642 print_indent(code
->dst
, code
->indent
);
1644 print_local_index(code
->dst
, group
, bounds
, pma
, domain
);
1645 fprintf(code
->dst
, " = ");
1646 print_global_index(code
->dst
, group
->array
, pma
, domain
);
1648 print_global_index(code
->dst
, group
->array
, pma
, domain
);
1649 fprintf(code
->dst
, " = ");
1650 print_local_index(code
->dst
, group
, bounds
, pma
, domain
);
1652 fprintf(code
->dst
, ";\n");
1654 isl_pw_multi_aff_free(pma
);
1655 isl_set_free(domain
);
1658 static void print_shared_access(struct cuda_gen
*gen
,
1659 __isl_keep isl_set
*shared_domain
, __isl_take isl_set
*access
,
1660 const char *type
, struct cuda_array_ref_group
*group
)
1662 const char *array_name
;
1665 isl_union_map
*sched
;
1666 unsigned nvar
= isl_set_dim(access
, isl_dim_set
);
1669 ctx
= isl_set_get_ctx(access
);
1670 array_name
= isl_set_get_tuple_name(access
);
1671 name
= isl_alloc_array(ctx
, char,
1672 strlen(type
) + sizeof("_shared_") + strlen(array_name
) + 20);
1673 if (group
->array
->n_group
> 1)
1674 sprintf(name
, "%s_shared_%s_%d", type
, array_name
, group
->nr
);
1676 sprintf(name
, "%s_shared_%s", type
, array_name
);
1677 access
= isl_set_set_tuple_name(access
, name
);
1680 sched
= access_schedule(gen
, access
, group
);
1682 n_tile
= gen
->n_block
;
1686 gen
->copy_sched
= isl_map_from_union_map(isl_union_map_copy(sched
));
1687 gen
->copy_group
= group
;
1688 gen
->copy_bound
= group
->shared_bound
;
1690 print_shared_body(gen
, shared_domain
, sched
, nvar
+ n_tile
,
1691 &print_copy_statement
, -1);
1693 isl_union_map_free(sched
);
1694 isl_map_free(gen
->copy_sched
);
1697 /* Return the union of all read (read = 1) and/or write (write = 1)
1698 * access relations in the group.
1700 static __isl_give isl_union_map
*group_access_relation(
1701 struct cuda_array_ref_group
*group
, int read
, int write
)
1704 isl_union_map
*access
;
1706 access
= isl_union_map_empty(isl_map_get_space(group
->access
));
1707 for (i
= 0; i
< group
->n_ref
; ++i
) {
1710 if (!((read
&& group
->refs
[i
]->read
) ||
1711 (write
&& group
->refs
[i
]->write
)))
1713 map_i
= isl_map_copy(group
->refs
[i
]->access
);
1714 access
= isl_union_map_union(access
,
1715 isl_union_map_from_map(map_i
));
1721 /* Check that none of the shared memory tiles involve any strides.
1723 static int no_strides(struct cuda_array_ref_group
*group
)
1726 int n_index
= group
->array
->n_index
;
1728 for (i
= 0; i
< n_index
; ++i
)
1729 if (group
->shared_bound
[i
].shift
)
1735 /* Return a set containing the values of the given index i
1736 * of the elements in the array tile in global memory that corresponds
1737 * to the shared memory copy.
1738 * In particular, if a is the index, we return a set with constraints
1740 * tile_offset <= a <= tile_offset + tile_size - 1
1744 * 0 <= a <= array_size - 1
1747 static __isl_give isl_set
*group_tile_dim(struct cuda_array_ref_group
*group
,
1750 isl_basic_set
*tile
;
1753 isl_local_space
*ls
;
1758 aff
= isl_aff_copy(group
->shared_bound
[i
].lb
);
1759 aff
= isl_aff_add_dims(aff
, isl_dim_in
, 1);
1760 ls
= isl_aff_get_domain_local_space(aff
);
1761 aff
= isl_aff_neg(aff
);
1762 aff
= isl_aff_add_coefficient_si(aff
, isl_dim_in
, 0, 1);
1763 c
= isl_inequality_from_aff(isl_aff_copy(aff
));
1764 tile
= isl_basic_set_from_constraint(c
);
1766 aff
= isl_aff_neg(aff
);
1767 aff
= isl_aff_add_constant(aff
, group
->shared_bound
[i
].size
);
1768 aff
= isl_aff_add_constant_si(aff
, -1);
1769 c
= isl_inequality_from_aff(aff
);
1770 tile
= isl_basic_set_add_constraint(tile
, c
);
1772 aff
= isl_aff_zero_on_domain(ls
);
1773 aff
= isl_aff_add_coefficient_si(aff
, isl_dim_in
, 0, 1);
1774 c
= isl_inequality_from_aff(aff
);
1775 tile
= isl_basic_set_add_constraint(tile
, c
);
1777 bound
= isl_pw_aff_copy(group
->array
->bound
[i
]);
1778 bound
= isl_pw_aff_add_dims(bound
, isl_dim_in
, 1);
1779 ls
= isl_local_space_from_space(isl_pw_aff_get_domain_space(bound
));
1780 aff
= isl_aff_zero_on_domain(ls
);
1781 aff
= isl_aff_add_coefficient_si(aff
, isl_dim_in
, 0, 1);
1782 aff
= isl_aff_add_constant_si(aff
, 1);
1783 dom
= isl_pw_aff_domain(isl_pw_aff_copy(bound
));
1785 tile_set
= isl_pw_aff_ge_set(bound
, isl_pw_aff_alloc(dom
, aff
));
1786 tile_set
= isl_set_align_params(tile_set
, isl_basic_set_get_space(tile
));
1787 tile_set
= isl_set_intersect(tile_set
, isl_set_from_basic_set(tile
));
1792 /* Return a set containing the elements in the array tile in
1793 * global memory that corresponds to the shared memory copy.
1795 static __isl_give isl_set
*group_tile(struct cuda_array_ref_group
*group
)
1798 int n_index
= group
->array
->n_index
;
1801 tile
= group_tile_dim(group
, 0);
1802 for (i
= 1; i
< n_index
; ++i
) {
1805 tile_i
= group_tile_dim(group
, i
);
1806 tile
= isl_set_flat_product(tile
, tile_i
);
1809 tile
= isl_set_set_tuple_name(tile
, group
->array
->name
);
1814 /* Print code for reading into or writing from shared memory
1815 * the given array reference group.
1817 * sched maps the original iteration domains to the shared memory tile loops.
1819 * If we are performing a read from global memory to shared memory,
1820 * if the array involved is not a scalar and if the definition of the
1821 * shared memory tiles does not involve any strides, then we copy
1822 * the entire tile to shared memory. This may result in some extra
1823 * elements getting copied, but it should lead to simpler code
1824 * (which means that fewer registers may be needed) and less divergence.
1826 * Otherwise, we only copy the elements that will be read or have been written
1829 * Note that the absence of stride requirement can easily be lifted.
1830 * We would just need to add constraints of the form
1832 * shift + a = stride * alpha
1834 static int print_group_shared_accesses(struct cuda_gen
*gen
,
1835 struct cuda_array_ref_group
*group
, const char *type
,
1836 __isl_keep isl_set
*shared_domain
, __isl_keep isl_union_map
*sched
)
1839 isl_union_map
*access
;
1840 isl_union_set
*uset
;
1841 isl_set
*access_set
;
1843 if (group
->private_bound
)
1845 if (!group
->shared_bound
)
1848 read
= !strcmp(type
, "read");
1850 access
= group_access_relation(group
, read
, !read
);
1851 access
= isl_union_map_apply_domain(access
, isl_union_map_copy(sched
));
1852 uset
= isl_union_map_range(access
);
1854 if (isl_union_set_is_empty(uset
)) {
1855 isl_union_set_free(uset
);
1859 if (read
&& group
->array
->n_index
> 0 && no_strides(group
)) {
1860 isl_union_set_free(uset
);
1861 access_set
= group_tile(group
);
1862 print_shared_access(gen
, shared_domain
, access_set
,
1867 access_set
= isl_set_from_union_set(uset
);
1868 access_set
= isl_set_coalesce(access_set
);
1870 print_shared_access(gen
, shared_domain
, access_set
, type
, group
);
1875 /* Print code for reading into or writing from shared memory at
1876 * the given level (-1 for innermost).
1878 * If we are not printing at the innermost level, then the dimensionality
1879 * of shared_domain may be smaller than gen->shared_len.
1880 * As the rest of the code assumes that the domain of access has
1881 * gen->shared_len dimensions, we therefore may need to embed this domain
1882 * in a higher dimensional space after intersection with shared_domain.
1884 static void print_shared_accesses(struct cuda_gen
*gen
,
1885 __isl_keep isl_set
*shared_domain
, __isl_keep isl_union_map
*access
,
1886 const char *type
, int level
)
1892 int shared_len
= isl_set_dim(shared_domain
, isl_dim_set
);
1894 isl_union_map
*sched
;
1896 shared_domain
= isl_set_copy(shared_domain
);
1897 sched
= isl_union_map_copy(gen
->tiled_sched
);
1898 dim
= isl_union_map_get_space(sched
);
1899 proj
= projection(dim
, gen
->tiled_len
, shared_len
);
1900 sched
= isl_union_map_apply_range(sched
, isl_union_map_from_map(proj
));
1901 sched
= isl_union_map_intersect_range(sched
,
1902 isl_union_set_from_set(isl_set_copy(shared_domain
)));
1903 if (shared_len
!= gen
->shared_len
) {
1904 dim
= isl_union_map_get_space(sched
);
1905 proj
= projection(dim
, gen
->shared_len
, shared_len
);
1906 proj
= isl_map_reverse(proj
);
1907 shared_domain
= isl_set_apply(shared_domain
,
1908 isl_map_copy(proj
));
1909 sched
= isl_union_map_apply_range(sched
,
1910 isl_union_map_from_map(proj
));
1913 dim
= isl_union_map_get_space(sched
);
1914 par
= parametrization(dim
, gen
->shared_len
, 0, gen
->shared_len
, "g");
1915 sched
= isl_union_map_intersect_range(sched
,
1916 isl_union_set_from_set(par
));
1918 for (i
= 0; i
< gen
->n_array
; ++i
) {
1919 struct cuda_array_info
*array
= &gen
->array
[i
];
1921 for (j
= 0; j
< array
->n_group
; ++j
) {
1922 if (array
->groups
[j
]->print_shared_level
!= level
)
1925 if (print_group_shared_accesses(gen
, array
->groups
[j
],
1926 type
, shared_domain
, sched
))
1931 isl_union_map_free(sched
);
1932 isl_set_free(shared_domain
);
1935 print_indent(gen
->cuda
.kernel_c
, gen
->kernel_code
.indent
);
1936 fprintf(gen
->cuda
.kernel_c
, "__syncthreads();\n");
1940 /* This function is called for each access to an array in some statement
1941 * in the original code.
1942 * Replace that access by an access to shared or (linearized) global memory.
1943 * Since the array in shared memory is just
1944 * a shifted copy of part of the original array, we simply need
1945 * to subtract the lower bound, which was computed
1946 * in can_tile_for_shared_memory.
1947 * If any of the indices is strided, then we first add
1948 * shared_bound[i].shift and divide by shared_bound[i].stride.
1950 * If the given array is accessed directly from global memory,
1951 * we don't need to perform any shifting and simply simplify
1952 * the expression in the context of the domain instead.
1954 * If the array space (range of access) has no name, then we are
1955 * accessing an iterator in the original program.
1957 static __isl_give isl_printer
*print_access(__isl_take isl_printer
*p
,
1958 struct cuda_gen
*gen
, __isl_take isl_map
*access
, int group_nr
)
1963 struct cuda_array_info
*array
= NULL
;
1964 isl_pw_multi_aff
*pma
;
1967 struct cuda_array_bound
*bounds
= NULL
;
1969 access
= isl_map_align_params(access
,
1970 isl_set_get_space(gen
->stmt_domain
));
1972 data_set
= isl_set_apply(isl_set_copy(gen
->stmt_domain
), access
);
1974 name
= isl_set_get_tuple_name(data_set
);
1977 fprintf(gen
->cuda
.kernel_c
, "(");
1979 struct cuda_array_ref_group
*group
;
1981 for (i
= 0; i
< gen
->n_array
; ++i
) {
1982 if (strcmp(name
, gen
->array
[i
].name
))
1984 array
= &gen
->array
[i
];
1987 group
= array
->groups
[group_nr
];
1988 bounds
= group
->private_bound
;
1990 bounds
= group
->shared_bound
;
1992 if (!bounds
&& cuda_array_is_scalar(array
) && !array
->read_only
)
1993 fprintf(gen
->cuda
.kernel_c
, "*");
1994 p
= print_array_name(p
, group
);
1996 if (cuda_array_is_scalar(array
)) {
1997 isl_set_free(data_set
);
2001 fprintf(gen
->cuda
.kernel_c
, "[");
2005 n_index
= isl_set_dim(data_set
, isl_dim_set
);
2006 pma
= isl_pw_multi_aff_from_set(data_set
);
2007 pma
= isl_pw_multi_aff_coalesce(pma
);
2010 for (i
= 0; i
+ 1 < n_index
; ++i
)
2011 p
= isl_printer_print_str(p
, "(");
2013 for (i
= 0; i
< n_index
; ++i
) {
2016 index
= isl_pw_multi_aff_get_pw_aff(pma
, i
);
2019 p
= isl_printer_print_pw_aff(p
, index
);
2020 isl_pw_aff_free(index
);
2024 domain
= isl_set_copy(gen
->stmt_domain
);
2025 domain
= isl_set_params(domain
);
2027 index
= isl_pw_aff_coalesce(index
);
2028 index
= isl_pw_aff_gist(index
, domain
);
2030 index
= shift_index(index
, array
, &bounds
[i
], domain
);
2034 p
= isl_printer_print_str(p
, ") * (");
2035 p
= isl_printer_print_pw_aff(p
,
2036 array
->local_bound
[i
]);
2037 p
= isl_printer_print_str(p
, ") + ");
2039 p
= isl_printer_print_str(p
, "][");
2041 p
= isl_printer_print_pw_aff(p
, index
);
2042 isl_pw_aff_free(index
);
2045 p
= isl_printer_print_str(p
, ")");
2047 p
= isl_printer_print_str(p
, "]");
2049 isl_pw_multi_aff_free(pma
);
2054 struct cuda_access_print_info
{
2055 struct cuda_gen
*gen
;
2056 struct cuda_stmt_access
*access
;
2059 /* To print the cuda accesses we walk the list of cuda accesses simultaneously
2060 * with the pet printer. This means that whenever the pet printer prints a
2061 * pet access expression we have the corresponding cuda access available and can
2062 * print the modified access.
2064 static __isl_give isl_printer
*print_cuda_access(__isl_take isl_printer
*p
,
2065 struct pet_expr
*expr
, void *usr
)
2067 struct cuda_access_print_info
*info
=
2068 (struct cuda_access_print_info
*) usr
;
2070 p
= print_access(p
, info
->gen
, isl_map_copy(info
->access
->access
),
2071 info
->access
->group
);
2072 info
->access
= info
->access
->next
;
2077 static void print_stmt_body(struct cuda_gen
*gen
,
2078 FILE *out
, struct cuda_stmt
*stmt
)
2080 struct cuda_access_print_info info
;
2083 p
= isl_printer_to_file(gen
->ctx
, out
);
2084 p
= isl_printer_set_output_format(p
, ISL_FORMAT_C
);
2087 info
.access
= stmt
->accesses
;
2089 p
= print_pet_expr(p
, stmt
->body
, &print_cuda_access
, &info
);
2090 fprintf(out
, ";\n");
2092 isl_printer_free(p
);
2095 /* This function is called for each leaf in the innermost clast,
2096 * i.e., for each statement.
2097 * We print the statement body, simplifying the accesses based
2100 static void print_statement(struct clast_printer_info
*code
,
2101 struct clast_user_stmt
*u
)
2103 struct cuda_gen
*gen
= code
->user
;
2106 isl_set
*stmt_domain
;
2107 isl_union_map
*stmt_sched
;
2108 isl_union_set
*uset
;
2110 struct cuda_stmt
*stmt
;
2112 nr
= atoi(u
->statement
->name
+ 2);
2113 stmt
= &gen
->stmts
[nr
];
2115 stmt_domain
= extract_host_domain(u
);
2117 stmt_sched
= isl_union_map_intersect_range(
2118 isl_union_map_copy(gen
->local_sched
),
2119 isl_union_set_from_set(extend(stmt_domain
,
2120 gen
->thread_tiled_len
)));
2121 dim
= isl_union_map_get_space(stmt_sched
);
2122 par
= parametrization(dim
, gen
->thread_tiled_len
, 0,
2123 gen
->thread_tiled_len
, "c");
2124 stmt_sched
= isl_union_map_intersect_range(stmt_sched
,
2125 isl_union_set_from_set(par
));
2127 uset
= isl_union_map_domain(stmt_sched
);
2128 dim
= isl_union_set_get_space(uset
);
2129 dim
= isl_space_add_dims(dim
, isl_dim_set
,
2130 isl_set_dim(stmt
->domain
, isl_dim_set
));
2131 dim
= isl_space_set_tuple_name(dim
, isl_dim_set
, u
->statement
->name
);
2132 gen
->stmt_domain
= isl_union_set_extract_set(uset
, dim
);
2133 isl_union_set_free(uset
);
2135 print_indent(code
->dst
, code
->indent
);
2136 print_stmt_body(gen
, code
->dst
, stmt
);
2138 isl_set_free(gen
->stmt_domain
);
2141 static void print_private_access(struct cuda_gen
*gen
,
2142 __isl_keep isl_set
*shared_domain
, __isl_take isl_set
*access
,
2143 const char *type
, struct cuda_array_ref_group
*group
)
2145 const char *array_name
;
2148 unsigned nvar
= isl_set_dim(access
, isl_dim_set
);
2149 isl_union_map
*usched
;
2151 if (isl_set_fast_is_empty(access
)) {
2152 isl_set_free(access
);
2156 ctx
= isl_set_get_ctx(access
);
2157 array_name
= isl_set_get_tuple_name(access
);
2158 name
= isl_alloc_array(ctx
, char,
2159 strlen(type
) + sizeof("_private_") + strlen(array_name
) + 20);
2160 if (group
->array
->n_group
> 1)
2161 sprintf(name
, "%s_private_%s_%d", type
, array_name
, group
->nr
);
2163 sprintf(name
, "%s_private_%s", type
, array_name
);
2164 access
= isl_set_set_tuple_name(access
, name
);
2167 gen
->copy_sched
= shift_access(access
, group
);
2168 gen
->copy_group
= group
;
2169 gen
->copy_bound
= group
->private_bound
;
2171 usched
= isl_union_map_from_map(isl_map_copy(gen
->copy_sched
));
2172 print_shared_body(gen
, shared_domain
, usched
, nvar
,
2173 &print_copy_statement
, 1);
2174 isl_union_map_free(usched
);
2176 isl_map_free(gen
->copy_sched
);
2179 /* Print code for reading into or writing from private memory
2180 * the given array reference group.
2182 * sched maps the original iteration domains to the shared memory tile loops.
2184 static void print_group_private_accesses(struct cuda_gen
*gen
,
2185 struct cuda_array_ref_group
*group
,
2186 const char *type
, __isl_keep isl_set
*shared_domain
,
2187 unsigned first_shared
, int shared_len
, __isl_keep isl_union_map
*sched
)
2190 isl_union_map
*access
;
2191 isl_union_set
*uset
;
2192 isl_set
*access_set
;
2194 if (!group
->private_bound
)
2197 read
= !strcmp(type
, "read");
2199 access
= group_access_relation(group
, read
, !read
);
2200 access
= isl_union_map_apply_domain(access
, isl_union_map_copy(sched
));
2201 access
= isl_union_map_intersect(access
,
2202 isl_union_map_copy(gen
->private_access
));
2203 uset
= isl_union_map_range(access
);
2205 if (isl_union_set_is_empty(uset
)) {
2206 isl_union_set_free(uset
);
2210 access_set
= isl_set_from_union_set(uset
);
2211 access_set
= isl_set_coalesce(access_set
);
2212 access_set
= isl_set_eliminate(access_set
, isl_dim_param
,
2213 first_shared
+ shared_len
,
2214 gen
->shared_len
- shared_len
);
2216 print_private_access(gen
, shared_domain
, access_set
, type
, group
);
2219 /* Print code for reading into or writing from private memory at
2220 * the given level (-1 for innermost).
2222 * If we are not printing at the innermost level, then the dimensionality
2223 * of shared_domain may be smaller than gen->shared_len.
2224 * As the rest of the code assumes that the domain of access has
2225 * gen->shared_len dimensions, we therefore may need to embed this domain
2226 * in a higher dimensional space after intersection with shared_domain.
2228 * This code is very similar to print_shared_accesses.
2229 * The main difference is that we to take into account gen->private_access.
2231 static void print_private_accesses(struct cuda_gen
*gen
,
2232 __isl_keep isl_set
*shared_domain
, __isl_keep isl_union_map
*access
,
2233 const char *type
, int level
)
2238 int shared_len
= isl_set_dim(shared_domain
, isl_dim_set
);
2239 unsigned first_shared
;
2240 isl_union_map
*sched
;
2242 shared_domain
= isl_set_copy(shared_domain
);
2243 sched
= isl_union_map_copy(gen
->tiled_sched
);
2244 dim
= isl_union_map_get_space(sched
);
2245 first_shared
= isl_space_dim(dim
, isl_dim_param
);
2246 proj
= projection(dim
, gen
->tiled_len
, shared_len
);
2247 sched
= isl_union_map_apply_range(sched
, isl_union_map_from_map(proj
));
2248 sched
= isl_union_map_intersect_range(sched
,
2249 isl_union_set_from_set(isl_set_copy(shared_domain
)));
2250 if (shared_len
!= gen
->shared_len
) {
2251 dim
= isl_union_map_get_space(sched
);
2252 proj
= projection(dim
, gen
->shared_len
, shared_len
);
2253 proj
= isl_map_reverse(proj
);
2254 shared_domain
= isl_set_apply(shared_domain
,
2255 isl_map_copy(proj
));
2256 sched
= isl_union_map_apply_range(sched
,
2257 isl_union_map_from_map(proj
));
2260 for (i
= 0; i
< gen
->n_array
; ++i
) {
2261 struct cuda_array_info
*array
= &gen
->array
[i
];
2263 for (j
= 0; j
< array
->n_group
; ++j
) {
2264 if (array
->groups
[j
]->print_shared_level
!= level
)
2267 print_group_private_accesses(gen
, array
->groups
[j
],
2268 type
, shared_domain
,
2269 first_shared
, shared_len
, sched
);
2273 isl_union_map_free(sched
);
2274 isl_set_free(shared_domain
);
2277 /* Set unroll[j] if the input dimension j is involved in
2278 * the index expression represented by bmap.
2280 static int check_unroll(__isl_take isl_basic_map
*bmap
, void *user
)
2283 int n_in
= isl_basic_map_dim(bmap
, isl_dim_in
);
2284 int n_out
= isl_basic_map_dim(bmap
, isl_dim_out
);
2287 for (i
= 0; i
< n_out
; ++i
) {
2291 ok
= isl_basic_map_has_defining_equality(bmap
,
2292 isl_dim_out
, i
, &c
);
2294 for (j
= 0; j
< n_in
; ++j
)
2295 if (isl_constraint_involves_dims(c
, isl_dim_in
, j
, 1))
2297 isl_constraint_free(c
);
2300 isl_basic_map_free(bmap
);
2304 /* Given an array pos mapping input dimensions to the corresponding
2305 * output dimension, construct the corresponding map.
2307 static __isl_give isl_map
*permutation(__isl_take isl_space
*dim
,
2312 isl_basic_map
*bmap
;
2313 isl_local_space
*ls
;
2315 dim
= isl_space_add_dims(dim
, isl_dim_in
, len
);
2316 dim
= isl_space_add_dims(dim
, isl_dim_out
, len
);
2317 bmap
= isl_basic_map_universe(isl_space_copy(dim
));
2318 ls
= isl_local_space_from_space(dim
);
2320 for (i
= 0; i
< len
; ++i
) {
2321 c
= isl_equality_alloc(isl_local_space_copy(ls
));
2322 isl_constraint_set_coefficient_si(c
, isl_dim_in
, i
, -1);
2323 isl_constraint_set_coefficient_si(c
, isl_dim_out
, pos
[i
], 1);
2324 bmap
= isl_basic_map_add_constraint(bmap
, c
);
2326 isl_local_space_free(ls
);
2328 return isl_map_from_basic_map(bmap
);
2331 /* Find all loops involved in any of the index expressions for any of
2332 * the private accesses, move them innermost and then mark them as
2333 * requiring unrolling by setting gen->first_unroll.
2334 * The loops involved should all be parallel because of the checks
2335 * we performed in check_private_group_access. Moving them innermost
2336 * is therefore a valid transformation.
2338 * Loops up to gen->shared_len are generated before the mapping to
2339 * threads is applied. They should therefore be ignored.
2341 static __isl_give isl_union_map
*interchange_for_unroll(struct cuda_gen
*gen
,
2342 __isl_take isl_union_map
*sched
)
2345 int unroll
[gen
->thread_tiled_len
];
2346 int perm
[gen
->thread_tiled_len
];
2349 int len
= gen
->shared_len
+ gen
->n_parallel
+ gen
->n_block
;
2351 gen
->first_unroll
= -1;
2353 for (i
= 0; i
< gen
->thread_tiled_len
; ++i
)
2355 for (i
= 0; i
< gen
->n_array
; ++i
) {
2356 struct cuda_array_info
*array
= &gen
->array
[i
];
2358 for (j
= 0; j
< array
->n_group
; ++j
) {
2359 isl_union_map
*access
;
2362 if (!array
->groups
[j
]->private_bound
)
2365 access
= group_access_relation(array
->groups
[j
], 1, 1);
2366 access
= isl_union_map_apply_domain(access
,
2367 isl_union_map_copy(sched
));
2369 acc
= isl_map_from_union_map(access
);
2370 isl_map_foreach_basic_map(acc
, &check_unroll
, unroll
);
2376 for (i
= gen
->shared_len
; i
< len
; ++i
)
2383 for (i
= len
; i
< gen
->thread_tiled_len
; ++i
)
2388 for (i
= 0; i
< gen
->shared_len
; ++i
)
2390 for (i
= gen
->shared_len
; i
< gen
->thread_tiled_len
; ++i
)
2393 gen
->first_unroll
= 1 + j
;
2394 for (i
= gen
->shared_len
; i
< len
; ++i
)
2398 dim
= isl_union_map_get_space(sched
);
2399 permute
= permutation(dim
, perm
, gen
->thread_tiled_len
);
2400 sched
= isl_union_map_apply_range(sched
,
2401 isl_union_map_from_map(permute
));
2406 /* This function is called for each leaf in the clast of the kernel code.
2407 * We first specialize the schedule to the site of the leaf and
2408 * print code for reading into shared memory, performing the actual
2409 * computations and writing from shared memory, with the required
2412 static void print_kernel_user(struct clast_printer_info
*code
,
2413 struct clast_user_stmt
*u
)
2415 struct cuda_gen
*gen
= code
->user
;
2416 isl_set
*shared_domain
;
2418 shared_domain
= extract_entire_host_domain(&u
->stmt
);
2420 print_shared_accesses(gen
, shared_domain
, gen
->read
, "read", -1);
2422 print_private_accesses(gen
, shared_domain
, gen
->read
, "read", -1);
2424 print_shared_body(gen
, shared_domain
, gen
->local_sched
,
2425 gen
->thread_tiled_len
, &print_statement
,
2428 print_private_accesses(gen
, shared_domain
, gen
->write
, "write", -1);
2430 print_indent(gen
->cuda
.kernel_c
, gen
->kernel_code
.indent
);
2431 fprintf(gen
->cuda
.kernel_c
, "__syncthreads();\n");
2433 print_shared_accesses(gen
, shared_domain
, gen
->write
, "write", -1);
2435 isl_set_free(shared_domain
);
2438 /* Check if we need to perform any copying to shared memory at this level
2439 * and if so, print the copying instructions.
2440 * Any array for which we are allowed to print copying instructions at
2441 * this level, but haven't done so already, is printed.
2443 static void copy_to_local(struct cuda_gen
*gen
, __isl_keep isl_set
*domain
)
2449 level
= isl_set_dim(domain
, isl_dim_set
);
2451 for (i
= 0; i
< gen
->n_array
; ++i
) {
2452 struct cuda_array_info
*array
= &gen
->array
[i
];
2454 for (j
= 0; j
< array
->n_group
; ++j
) {
2455 if (array
->groups
[j
]->print_shared_level
>= 0)
2457 if (array
->groups
[j
]->last_shared
>= level
)
2459 array
->groups
[j
]->print_shared_level
= level
;
2465 print_shared_accesses(gen
, domain
, gen
->read
, "read", level
);
2466 print_private_accesses(gen
, domain
, gen
->read
, "read", level
);
2471 /* This function is called for each for loop in the clast,
2472 * right after the opening brace has been printed.
2474 * Print copying instructions to shared or private memory if needed.
2476 static void print_kernel_for_head(struct clast_printer_info
*code
,
2477 struct clast_for
*f
)
2479 struct cuda_gen
*gen
= code
->user
;
2482 domain
= isl_set_from_cloog_domain(cloog_domain_copy(f
->domain
));
2483 copy_to_local(gen
, domain
);
2485 isl_set_free(domain
);
2488 /* Print instructions for copying from shared memory for each array
2489 * for which print_kernel_for_head has added copying instructions
2492 static void copy_from_local(struct cuda_gen
*gen
, __isl_keep isl_set
*domain
)
2498 level
= isl_set_dim(domain
, isl_dim_set
);
2500 for (i
= 0; i
< gen
->n_array
; ++i
) {
2501 struct cuda_array_info
*array
= &gen
->array
[i
];
2503 for (j
= 0; j
< array
->n_group
; ++j
) {
2504 if (array
->groups
[j
]->print_shared_level
!= level
)
2514 print_private_accesses(gen
, domain
, gen
->write
, "write", level
);
2515 print_shared_accesses(gen
, domain
, gen
->write
, "write", level
);
2519 /* This function is called for each for loop in the clast,
2520 * right before the closing brace is printed.
2522 * Print copying instructions from shared or private memory if needed.
2524 static void print_kernel_for_foot(struct clast_printer_info
*code
,
2525 struct clast_for
*f
)
2527 struct cuda_gen
*gen
= code
->user
;
2530 domain
= isl_set_from_cloog_domain(cloog_domain_copy(f
->domain
));
2531 copy_from_local(gen
, domain
);
2533 isl_set_free(domain
);
2536 /* Use CLooG to generate code for the outer gen->shared_first loops
2537 * of the local schedule "sched".
2538 * The pretty printing of this code is handled by print_clast,
2539 * which calls print_kernel_user for each iteration of the shared tile loops.
2541 static void print_cloog_kernel_body(struct cuda_gen
*gen
,
2542 __isl_keep isl_set
*context
, __isl_keep isl_union_map
*sched
)
2545 CloogOptions
*options
;
2546 CloogDomain
*cloog_context
;
2547 CloogUnionDomain
*ud
;
2549 struct clast_stmt
*stmt
;
2552 sched
= isl_union_map_copy(sched
);
2553 sched
= isl_union_map_align_params(sched
, isl_set_get_space(context
));
2555 options
= cloog_options_malloc(gen
->state
);
2556 options
->language
= CLOOG_LANGUAGE_C
;
2557 options
->strides
= 1;
2559 options
->stop
= gen
->shared_len
;
2560 options
->f
= gen
->tiled_len
;
2561 options
->l
= gen
->tiled_len
;
2562 options
->save_domains
= 1;
2563 options
->noscalars
= 1;
2565 ud
= cloog_union_domain_from_isl_union_map(sched
);
2566 for (i
= 0; i
< gen
->shared_len
; ++i
) {
2567 snprintf(name
, sizeof(name
), "g%d", i
);
2568 ud
= cloog_union_domain_set_name(ud
, CLOOG_SCAT
, i
, name
);
2570 cloog_context
= cloog_domain_from_isl_set(isl_set_copy(context
));
2571 input
= cloog_input_alloc(cloog_context
, ud
);
2573 stmt
= cloog_clast_create_from_input(input
, options
);
2575 gen
->kernel_code
.indent
= 4;
2576 gen
->kernel_code
.dst
= gen
->cuda
.kernel_c
;
2577 gen
->kernel_code
.print_user_stmt
= NULL
;
2578 gen
->kernel_code
.print_user_stmt_list
= &print_kernel_user
;
2579 gen
->kernel_code
.print_for_head
= &print_kernel_for_head
;
2580 gen
->kernel_code
.print_for_foot
= &print_kernel_for_foot
;
2581 gen
->kernel_code
.user
= gen
;
2582 copy_to_local(gen
, context
);
2583 print_clast(&gen
->kernel_code
, stmt
);
2584 copy_from_local(gen
, context
);
2586 cloog_clast_free(stmt
);
2587 cloog_options_free(options
);
2590 static void print_kernel_iterators(struct cuda_gen
*gen
)
2593 const char *block_dims
[] = { "blockIdx.x", "blockIdx.y" };
2594 const char *thread_dims
[] = { "threadIdx.x", "threadIdx.y",
2597 if (gen
->n_grid
> 0) {
2598 print_indent(gen
->cuda
.kernel_c
, 4);
2599 fprintf(gen
->cuda
.kernel_c
, "int ");
2600 for (i
= 0; i
< gen
->n_grid
; ++i
) {
2602 fprintf(gen
->cuda
.kernel_c
, ", ");
2603 fprintf(gen
->cuda
.kernel_c
, "b%d = %s",
2604 i
, block_dims
[gen
->n_grid
- 1 - i
]);
2606 fprintf(gen
->cuda
.kernel_c
, ";\n");
2609 if (gen
->n_block
> 0) {
2610 print_indent(gen
->cuda
.kernel_c
, 4);
2611 fprintf(gen
->cuda
.kernel_c
, "int ");
2612 for (i
= 0; i
< gen
->n_block
; ++i
) {
2614 fprintf(gen
->cuda
.kernel_c
, ", ");
2615 fprintf(gen
->cuda
.kernel_c
, "t%d = %s",
2616 i
, thread_dims
[gen
->n_block
- 1 - i
]);
2618 fprintf(gen
->cuda
.kernel_c
, ";\n");
2622 static void print_group_shared_array(struct cuda_gen
*gen
,
2623 struct cuda_array_ref_group
*group
)
2626 struct cuda_array_bound
*bounds
;
2629 bounds
= group
->private_bound
;
2631 bounds
= group
->shared_bound
;
2635 print_indent(gen
->cuda
.kernel_c
, 4);
2636 fprintf(gen
->cuda
.kernel_c
, "%s%s ",
2637 group
->private_bound
? "" : "__shared__ ", group
->array
->type
);
2638 p
= isl_printer_to_file(gen
->ctx
, gen
->cuda
.kernel_c
);
2639 p
= print_array_name(p
, group
);
2640 isl_printer_free(p
);
2641 for (j
= 0; j
< group
->array
->n_index
; ++j
) {
2642 fprintf(gen
->cuda
.kernel_c
, "[");
2643 isl_int_print(gen
->cuda
.kernel_c
, bounds
[j
].size
, 0);
2644 fprintf(gen
->cuda
.kernel_c
, "]");
2646 fprintf(gen
->cuda
.kernel_c
, ";\n");
2649 static void print_shared_arrays(struct cuda_gen
*gen
)
2653 for (i
= 0; i
< gen
->n_array
; ++i
) {
2654 struct cuda_array_info
*array
= &gen
->array
[i
];
2656 for (j
= 0; j
< array
->n_group
; ++j
)
2657 print_group_shared_array(gen
, array
->groups
[j
]);
2661 static void print_kernel_body(struct cuda_gen
*gen
,
2662 __isl_keep isl_set
*host_domain
, __isl_keep isl_union_map
*sched
)
2666 context
= isl_set_copy(host_domain
);
2667 context
= parametrize(context
, 0, gen
->tile_first
, "h");
2668 context
= isl_set_project_out(context
, isl_dim_set
, 0, gen
->tile_first
);
2669 context
= add_bounded_parameters(context
,
2670 gen
->n_grid
, gen
->grid_dim
, "b");
2672 print_kernel_iterators(gen
);
2673 print_shared_arrays(gen
);
2675 fprintf(gen
->cuda
.kernel_c
, "\n");
2677 print_cloog_kernel_body(gen
, context
, sched
);
2679 isl_set_free(context
);
2682 /* Given a constraint
2684 * a(p,i) + j = g f(e)
2686 * or -a(p,i) - j = g f(e) if sign < 0,
2687 * store a(p,i) in bound->shift and g (stride) in bound->stride.
2688 * a(p,i) is assumed to be an expression in only the parameters.
2690 static void extract_stride(__isl_keep isl_constraint
*c
,
2691 struct cuda_array_bound
*bound
, isl_int stride
, int sign
)
2699 isl_int_set(bound
->stride
, stride
);
2701 dim
= isl_constraint_get_space(c
);
2702 dim
= isl_space_params(dim
);
2704 nparam
= isl_space_dim(dim
, isl_dim_param
);
2708 isl_constraint_get_constant(c
, &v
);
2711 aff
= isl_aff_zero_on_domain(isl_local_space_from_space(dim
));
2712 aff
= isl_aff_set_constant(aff
, v
);
2714 for (i
= 0; i
< nparam
; ++i
) {
2715 isl_constraint_get_coefficient(c
, isl_dim_param
, i
, &v
);
2716 if (isl_int_is_zero(v
))
2720 aff
= isl_aff_add_coefficient(aff
, isl_dim_param
, i
, v
);
2728 /* Given an equality constraint of a map with a single output dimension j,
2729 * check if the constraint is of the form
2731 * a(p,i) + j = g f(e)
2733 * with a(p,i) an expression in the parameters and input dimensions
2734 * and f(e) an expression in the existentially quantified variables.
2735 * If so, and if g is larger than any such g from a previously considered
2736 * constraint, then call extract_stride to record the stride information
2739 static int check_stride_constraint(__isl_take isl_constraint
*c
, void *user
)
2744 struct cuda_array_bound
*bound
= user
;
2747 isl_int_init(stride
);
2749 n_div
= isl_constraint_dim(c
, isl_dim_div
);
2750 isl_constraint_get_coefficient(c
, isl_dim_out
, 0, &v
);
2752 if (n_div
&& (isl_int_is_one(v
) || isl_int_is_negone(v
))) {
2753 int s
= isl_int_sgn(v
);
2754 isl_int_set_si(stride
, 0);
2755 for (i
= 0; i
< n_div
; ++i
) {
2756 isl_constraint_get_coefficient(c
, isl_dim_div
, i
, &v
);
2757 isl_int_gcd(stride
, stride
, v
);
2759 if (!isl_int_is_zero(stride
) &&
2760 isl_int_gt(stride
, bound
->stride
))
2761 extract_stride(c
, bound
, stride
, s
);
2764 isl_int_clear(stride
);
2767 isl_constraint_free(c
);
2771 /* Given contraints on an array index i, check if we can find
2772 * a shift a(p) and a stride g such that
2774 * a(p) + i = 0 mod g
2776 * If so, record the information in bound and apply the mapping
2777 * i -> (i + a(p))/g to the array index in bounds and return
2778 * the new constraints.
2779 * If not, simply return the original constraints.
2781 static __isl_give isl_basic_map
*check_stride(struct cuda_array_bound
*bound
,
2782 __isl_take isl_basic_map
*bounds
)
2785 isl_basic_map
*shift
;
2788 isl_int_set_si(bound
->stride
, -1);
2790 aff
= isl_basic_map_affine_hull(isl_basic_map_copy(bounds
));
2792 isl_basic_map_foreach_constraint(aff
, &check_stride_constraint
, bound
);
2794 isl_basic_map_free(aff
);
2796 if (isl_int_is_neg(bound
->stride
))
2799 aff_shift
= isl_aff_copy(bound
->shift
);
2800 aff_shift
= isl_aff_add_dims(aff_shift
, isl_dim_in
, 1);
2801 aff_shift
= isl_aff_add_coefficient_si(aff_shift
, isl_dim_in
, 0, 1);
2802 aff_shift
= isl_aff_scale_down(aff_shift
, bound
->stride
);
2803 shift
= isl_basic_map_from_aff(aff_shift
);
2805 bound
->shift_map
= isl_basic_map_copy(shift
);
2806 bounds
= isl_basic_map_apply_range(bounds
, shift
);
2811 struct cuda_size_info
{
2812 isl_basic_set
*bset
;
2813 struct cuda_array_bound
*bound
;
2817 /* Given a constraint from the basic set describing the bounds on
2818 * an array index, check if it is a lower bound, say m i >= b(x), and,
2819 * if so, check whether the expression "i - ceil(b(x)/m) + 1" has a constant
2820 * upper bound. If so, and if this bound is smaller than any bound
2821 * derived from earlier constraints, set the size to this bound on
2822 * the expression and the lower bound to ceil(b(x)/m).
2824 static int compute_size_in_direction(__isl_take isl_constraint
*c
, void *user
)
2826 struct cuda_size_info
*size
= user
;
2831 nparam
= isl_basic_set_dim(size
->bset
, isl_dim_param
);
2832 n_div
= isl_constraint_dim(c
, isl_dim_div
);
2834 if (isl_constraint_involves_dims(c
, isl_dim_div
, 0, n_div
)) {
2835 isl_constraint_free(c
);
2841 isl_constraint_get_coefficient(c
, isl_dim_set
, size
->pos
, &v
);
2843 if (isl_int_is_pos(v
)) {
2846 enum isl_lp_result res
;
2848 aff
= isl_constraint_get_bound(c
, isl_dim_set
, size
->pos
);
2849 aff
= isl_aff_ceil(aff
);
2851 lb
= isl_aff_copy(aff
);
2853 aff
= isl_aff_neg(aff
);
2854 aff
= isl_aff_add_coefficient_si(aff
, isl_dim_in
, size
->pos
, 1);
2856 res
= isl_basic_set_max(size
->bset
, aff
, &v
);
2859 if (res
== isl_lp_ok
) {
2860 isl_int_add_ui(v
, v
, 1);
2861 if (isl_int_is_neg(size
->bound
->size
) ||
2862 isl_int_lt(v
, size
->bound
->size
)) {
2863 isl_int_set(size
->bound
->size
, v
);
2864 lb
= isl_aff_drop_dims(lb
, isl_dim_in
,
2866 isl_aff_free(size
->bound
->lb
);
2867 size
->bound
->lb
= isl_aff_copy(lb
);
2874 isl_constraint_free(c
);
2879 /* Given a basic map "bounds" that maps parameters and input dimensions
2880 * to a single output dimension, look for an expression in the parameters
2881 * and input dimensions such that the range of the output dimension shifted
2882 * by this expression is a constant.
2884 * In particular, we currently only consider lower bounds on the output
2885 * dimension as candidate expressions.
2887 static int compute_array_dim_size(struct cuda_array_bound
*bound
,
2888 __isl_take isl_basic_map
*bounds
)
2890 struct cuda_size_info size
;
2892 bounds
= isl_basic_map_detect_equalities(bounds
);
2893 bounds
= check_stride(bound
, bounds
);
2895 isl_int_set_si(bound
->size
, -1);
2899 size
.pos
= isl_basic_map_dim(bounds
, isl_dim_in
);
2900 size
.bset
= isl_basic_map_wrap(bounds
);
2901 size
.bset
= isl_basic_set_flatten(size
.bset
);
2902 size
.bset
= isl_set_simple_hull(isl_basic_set_compute_divs(size
.bset
));
2903 isl_basic_set_foreach_constraint(size
.bset
, &compute_size_in_direction
,
2905 isl_basic_set_free(size
.bset
);
2907 return isl_int_is_nonneg(bound
->size
) ? 0 : -1;
2910 /* Check if we can find a shared memory tile for the given array
2911 * based on the given accesses, and if so, put the results
2912 * in array->shared_bound.
2914 * We project the accesses on each index in turn and look for a parametric
2915 * offset such that the size is constant.
2917 static int can_tile_for_shared_memory(struct cuda_array_info
*array
,
2918 __isl_keep isl_map
*access
, struct cuda_array_bound
*bounds
)
2922 for (i
= 0; i
< array
->n_index
; ++i
) {
2924 isl_basic_map
*hull
;
2926 access_i
= isl_map_copy(access
);
2927 access_i
= isl_map_project_out(access_i
, isl_dim_out
, 0, i
);
2928 access_i
= isl_map_project_out(access_i
, isl_dim_out
,
2929 1, array
->n_index
- (i
+ 1));
2930 access_i
= isl_map_compute_divs(access_i
);
2931 hull
= isl_map_simple_hull(access_i
);
2932 if (compute_array_dim_size(&bounds
[i
], hull
) < 0)
2939 /* Construct a map with input the shared tile loops and the loops that
2940 * will be wrapped around the threads that relates these later loops
2941 * to the thread indices and then projects them out.
2943 static __isl_give isl_map
*compute_privatization(struct cuda_gen
*gen
)
2951 dim
= isl_union_map_get_space(gen
->shared_sched
);
2953 if (gen
->options
->wrap
)
2954 tiling
= wrap(isl_space_copy(dim
), gen
->shared_len
+ gen
->n_block
,
2955 gen
->shared_len
, gen
->n_block
, gen
->block_dim
);
2957 tiling
= tile(isl_space_copy(dim
), gen
->shared_len
+ gen
->n_block
,
2958 gen
->shared_len
, gen
->n_block
, gen
->block_dim
);
2962 par
= parametrization(dim
, gen
->shared_len
+ 2 * gen
->n_block
,
2963 gen
->tile_first
+ gen
->tile_len
+ gen
->n_grid
+ gen
->n_block
,
2966 priv
= isl_map_align_params(priv
, isl_set_get_space(par
));
2967 priv
= isl_map_intersect_range(priv
, par
);
2969 dim
= isl_map_get_space(priv
);
2970 dim
= isl_space_drop_dims(dim
, isl_dim_in
, 0, isl_space_dim(dim
, isl_dim_in
));
2971 dim
= isl_space_drop_dims(dim
, isl_dim_out
, 0, isl_space_dim(dim
, isl_dim_out
));
2972 proj
= projection(dim
, gen
->shared_len
+ 2 * gen
->n_block
,
2975 priv
= isl_map_apply_range(priv
, proj
);
2980 /* Construct a map from domain_dim to domain_dim that increments
2981 * the dimension at position "pos" and leaves all other dimensions
2984 static __isl_give isl_map
*next(__isl_take isl_space
*domain_dim
, int pos
)
2987 int len
= isl_space_dim(domain_dim
, isl_dim_set
);
2989 isl_basic_map
*next
;
2990 isl_local_space
*ls
;
2992 dim
= isl_space_map_from_set(domain_dim
);
2993 next
= isl_basic_map_universe(isl_space_copy(dim
));
2994 ls
= isl_local_space_from_space(dim
);
2996 for (i
= 0; i
< len
; ++i
) {
2999 c
= isl_equality_alloc(isl_local_space_copy(ls
));
3000 isl_constraint_set_coefficient_si(c
, isl_dim_in
, i
, 1);
3001 isl_constraint_set_coefficient_si(c
, isl_dim_out
, i
, -1);
3003 isl_constraint_set_constant_si(c
, 1);
3004 next
= isl_basic_map_add_constraint(next
, c
);
3007 isl_local_space_free(ls
);
3009 return isl_map_from_basic_map(next
);
3012 /* Check if the given access is coalesced.
3013 * That is, check whether incrementing the dimension that will get
3014 * wrapped over the last thread index results in incrementing
3015 * the last array index.
3017 * This function is only called for access relations without reuse.
3019 static int access_is_coalesced(struct cuda_gen
*gen
,
3020 __isl_keep isl_union_map
*access
)
3023 isl_map
*access_map
;
3024 isl_map
*next_thread_x
;
3025 isl_map
*next_element
;
3029 access
= isl_union_map_copy(access
);
3030 access
= isl_union_map_apply_domain(access
,
3031 isl_union_map_copy(gen
->tiled_sched
));
3032 access_map
= isl_map_from_union_map(access
);
3034 dim
= isl_map_get_space(access_map
);
3035 dim
= isl_space_domain(dim
);
3036 next_thread_x
= next(dim
, gen
->shared_len
+ gen
->n_block
- 1);
3038 dim
= isl_map_get_space(access_map
);
3039 dim
= isl_space_range(dim
);
3040 next_element
= next(dim
, isl_space_dim(dim
, isl_dim_set
) - 1);
3042 map
= isl_map_apply_domain(next_thread_x
, isl_map_copy(access_map
));
3043 map
= isl_map_apply_range(map
, access_map
);
3045 coalesced
= isl_map_is_subset(map
, next_element
);
3047 isl_map_free(next_element
);
3053 /* For the given array reference group, check whether the access is private
3054 * to the thread. That is, check that any given array element
3055 * is only accessed by a single thread.
3056 * We compute an access relation that maps the shared tile loop iterators
3057 * and the shared point loop iterators that will be wrapped over the
3058 * threads to the array elements.
3059 * We actually check that those iterators that will be wrapped
3060 * partition the array space. This check is stricter than necessary
3061 * since several iterations may be mapped onto the same thread
3062 * and then they could be allowed to access the same memory elements,
3063 * but our check does not allow this situation.
3065 * We also check that the index expression only depends on parallel
3066 * loops. That way, we can move those loops innermost and unroll them.
3067 * Again, we use a test that is stricter than necessary.
3068 * We actually check whether the index expression only depends
3069 * on the iterators that are wrapped over the threads.
3070 * These are necessarily parallel, but there may be more parallel loops.
3072 * Combining the injectivity of the first test with the single-valuedness
3073 * of the second test, we simply test for bijectivity.
3075 * If it turns out we can use registers, we compute the private memory
3076 * tile size using can_tile_for_shared_memory, after introducing a dependence
3077 * on the thread indices.
3079 * Before performing any of the above computations, we first check
3080 * if there is any reuse on the reference group. If not, we simply
3081 * return. If, moreover, the access is coalesced then we also remove
3082 * the shared memory tiling since we should just use global memory instead.
3084 static void check_private_group_access(struct cuda_gen
*gen
,
3085 struct cuda_array_ref_group
*group
)
3088 isl_union_map
*access
;
3089 int n_index
= group
->array
->n_index
;
3091 access
= group_access_relation(group
, 1, 1);
3092 if (isl_union_map_is_injective(access
)) {
3093 if (group
->shared_bound
&& access_is_coalesced(gen
, access
)) {
3094 free_bound_list(group
->shared_bound
, n_index
);
3095 group
->shared_bound
= NULL
;
3097 isl_union_map_free(access
);
3100 access
= isl_union_map_apply_domain(access
,
3101 isl_union_map_copy(gen
->shared_sched
));
3103 acc
= isl_map_from_union_map(access
);
3105 if (!isl_map_is_bijective(acc
)) {
3110 group
->private_bound
= create_bound_list(gen
->ctx
, n_index
);
3111 acc
= isl_map_align_params(acc
, isl_map_get_space(gen
->privatization
));
3112 acc
= isl_map_apply_domain(acc
, isl_map_copy(gen
->privatization
));
3113 if (!can_tile_for_shared_memory(group
->array
, acc
,
3114 group
->private_bound
)) {
3115 free_bound_list(group
->private_bound
, n_index
);
3116 group
->private_bound
= NULL
;
3122 /* Look for the last shared tile loop that affects the offset of the
3123 * shared or private tile and store the result in array->last_shared.
3125 static void set_last_shared(struct cuda_gen
*gen
,
3126 struct cuda_array_ref_group
*group
)
3129 struct cuda_array_bound
*bounds
;
3130 unsigned first_shared
= gen
->first_shared
;
3131 int n_index
= group
->array
->n_index
;
3133 bounds
= group
->private_bound
;
3135 bounds
= group
->shared_bound
;
3139 for (j
= gen
->shared_len
- 1; j
>= 0; --j
) {
3140 for (i
= 0; i
< n_index
; ++i
) {
3145 if (isl_aff_involves_dims(lb
, isl_dim_param
,
3146 first_shared
+ j
, 1))
3149 shift
= bounds
[i
].shift
;
3152 if (isl_aff_involves_dims(shift
, isl_dim_param
,
3153 first_shared
+ j
, 1))
3159 group
->last_shared
= j
;
3162 /* Compute the sizes of all private arrays for the current kernel,
3163 * as well as the offsets of the private pieces in the original arrays.
3164 * If we cannot or don't want to privatize a given array group,
3165 * we use the shared memory tile sizes computed in
3166 * compute_group_shared_bound instead.
3168 * If we have been able to find a private or shared tile,
3169 * we also look for the last shared tile loop that affects the offset
3170 * (and therefore the group tile) and store the result in group->last_shared.
3172 * A privatized copy of all access relations from reference groups that
3173 * are mapped to private memory is stored in gen->privatization.
3175 static void compute_private_size(struct cuda_gen
*gen
)
3178 isl_union_map
*private;
3180 if (!gen
->options
->use_private_memory
)
3183 private = isl_union_map_empty(isl_union_map_get_space(gen
->shared_sched
));
3185 for (i
= 0; i
< gen
->n_array
; ++i
) {
3186 struct cuda_array_info
*array
= &gen
->array
[i
];
3188 if (cuda_array_is_read_only_scalar(array
))
3191 for (j
= 0; j
< array
->n_group
; ++j
) {
3192 check_private_group_access(gen
, array
->groups
[j
]);
3194 if (!array
->groups
[j
]->private_bound
)
3197 private = isl_union_map_union(private,
3198 group_access_relation(array
->groups
[j
], 1, 1));
3201 for (j
= 0; j
< array
->n_group
; ++j
) {
3202 array
->groups
[j
]->last_shared
= gen
->shared_len
- 1;
3203 array
->groups
[j
]->print_shared_level
= -1;
3204 set_last_shared(gen
, array
->groups
[j
]);
3208 if (isl_union_map_is_empty(private))
3209 isl_union_map_free(private);
3211 isl_union_map
*priv
;
3213 private = isl_union_map_apply_domain(private,
3214 isl_union_map_copy(gen
->shared_sched
));
3215 priv
= isl_union_map_from_map(isl_map_copy(gen
->privatization
));
3216 private = isl_union_map_apply_domain(private, priv
);
3217 gen
->private_access
= private;
3221 /* Compute the size of the tile specified by the list "bound" of n_index
3222 * cuda_array_bounds in number of elements and put the result in *size.
3224 static void tile_size(unsigned n_index
, struct cuda_array_bound
*bound
,
3229 isl_int_set_si(*size
, 1);
3231 for (i
= 0; i
< n_index
; ++i
)
3232 isl_int_mul(*size
, *size
, bound
[i
].size
);
3235 /* If max_shared_memory is not set to infinity (-1), then make
3236 * sure that the total amount of shared memory required by the
3237 * array reference groups mapped to shared memory is no larger
3238 * than this maximum.
3240 * We apply a greedy approach and discard (keep in global memory)
3241 * those groups that would result in a total memory size that
3242 * is larger than the maximum.
3244 static void check_shared_memory_bound(struct cuda_gen
*gen
)
3249 if (gen
->options
->max_shared_memory
< 0)
3254 isl_int_set_si(left
, gen
->options
->max_shared_memory
);
3256 for (i
= 0; i
< gen
->n_array
; ++i
) {
3257 struct cuda_array_info
*array
= &gen
->array
[i
];
3259 for (j
= 0; j
< array
->n_group
; ++j
) {
3260 struct cuda_array_ref_group
*group
;
3262 group
= array
->groups
[j
];
3263 if (!group
->shared_bound
)
3266 tile_size(array
->n_index
, group
->shared_bound
, &size
);
3267 isl_int_mul_ui(size
, size
, array
->size
);
3269 if (isl_int_le(size
, left
)) {
3270 isl_int_sub(left
, left
, size
);
3274 free_bound_list(group
->shared_bound
, array
->n_index
);
3275 group
->shared_bound
= NULL
;
3279 isl_int_clear(size
);
3280 isl_int_clear(left
);
3283 /* Fill up the groups array with singleton groups, i.e., one group
3284 * per reference, initializing the array, access, write and refs fields.
3285 * In particular the access field is initialized to the scheduled
3286 * access relation of the array reference.
3288 * Return the number of elements initialized, i.e., the number of
3289 * active references in the current kernel.
3291 static int populate_array_references(struct cuda_array_info
*array
,
3292 __isl_keep isl_union_map
*sched
, struct cuda_array_ref_group
**groups
)
3296 isl_ctx
*ctx
= isl_union_map_get_ctx(sched
);
3299 for (i
= 0; i
< array
->n_ref
; ++i
) {
3300 isl_union_map
*umap
;
3302 struct cuda_array_ref_group
*group
;
3303 struct cuda_stmt_access
*access
= array
->refs
[i
];
3305 map
= isl_map_copy(access
->access
);
3306 umap
= isl_union_map_from_map(map
);
3307 umap
= isl_union_map_apply_domain(umap
,
3308 isl_union_map_copy(sched
));
3310 if (isl_union_map_is_empty(umap
)) {
3311 isl_union_map_free(umap
);
3315 map
= isl_map_from_union_map(umap
);
3316 map
= isl_map_detect_equalities(map
);
3318 group
= isl_calloc_type(ctx
, struct cuda_array_ref_group
);
3320 group
->array
= array
;
3321 group
->access
= map
;
3322 group
->write
= access
->write
;
3323 group
->refs
= &array
->refs
[i
];
3325 groups
[n
++] = group
;
3331 static void free_array_ref_group(struct cuda_array_ref_group
*group
,
3336 free_bound_list(group
->shared_bound
, n_index
);
3337 free_bound_list(group
->private_bound
, n_index
);
3338 isl_map_free(group
->access
);
3343 /* Given a set where the parameters gen->first_shared up to
3344 * gen->first_shared + gen->shared_len represent the tile loops,
3345 * eliminate the innermost of those that have a fixed value
3346 * until we reach one that does not (obviously) have a fixed value.
3348 static __isl_give isl_set
*eliminate_fixed_inner_loops(struct cuda_gen
*gen
,
3349 __isl_take isl_set
*access
)
3353 for (i
= gen
->shared_len
- 1; i
>= 0; --i
) {
3354 int pos
= gen
->first_shared
+ i
;
3355 if (!isl_set_plain_is_fixed(access
, isl_dim_param
, pos
, NULL
))
3357 access
= isl_set_eliminate(access
, isl_dim_param
, pos
, 1);
3362 /* Check if the accessed set of group1 and group2 overlap within
3363 * the innermost loop. In particular, ignore any inner dimension
3364 * with a fixed value.
3365 * The copying to and from shared memory will be performed within
3366 * the innermost actual loop so we are only allowed to consider
3367 * the dimensions up to that innermost loop while checking whether
3368 * two access sets overlap.
3370 static int accesses_overlap(struct cuda_gen
*gen
,
3371 struct cuda_array_ref_group
*group1
,
3372 struct cuda_array_ref_group
*group2
)
3375 isl_set
*access1
, *access2
;
3377 access1
= isl_map_range(isl_map_copy(group1
->access
));
3378 access1
= eliminate_fixed_inner_loops(gen
, access1
);
3379 access2
= isl_map_range(isl_map_copy(group2
->access
));
3380 access2
= eliminate_fixed_inner_loops(gen
, access2
);
3381 access1
= isl_set_intersect(access1
, access2
);
3382 empty
= isl_set_is_empty(access1
);
3383 isl_set_free(access1
);
3388 /* If two groups have overlapping access relations (within the innermost
3389 * loop) and if one of them involves a write, then merge the two groups
3392 * We keep track of the grouping in "leader". leader[j] points to
3393 * an earlier group array element that belongs to the same group,
3394 * or the array element j itself if this element is the first in the group.
3396 * Return the number of group leaders.
3398 static int group_overlapping_writes(struct cuda_gen
*gen
, int n
,
3399 struct cuda_array_ref_group
**groups
, int *leader
)
3404 for (i
= 0; i
< n
; ++i
) {
3406 groups
[l
]->n_ref
= 1;
3407 for (j
= i
- 1; j
>= 0; --j
) {
3410 if (!groups
[l
]->write
&& !groups
[j
]->write
)
3413 if (!accesses_overlap(gen
, groups
[l
], groups
[j
]))
3416 groups
[j
]->access
= isl_map_union(groups
[j
]->access
,
3418 groups
[j
]->write
= 1;
3419 groups
[l
]->access
= NULL
;
3420 groups
[j
]->n_ref
+= groups
[l
]->n_ref
;
3430 /* Compute the size of the shared array corresponding to the given
3431 * array reference group, based on the accesses from the current kernel,
3432 * as well as the offset of the shared piece in the original array.
3434 static void compute_group_shared_bound(struct cuda_gen
*gen
,
3435 struct cuda_array_info
*array
, struct cuda_array_ref_group
*group
)
3437 isl_ctx
*ctx
= isl_space_get_ctx(array
->dim
);
3439 if (!gen
->options
->use_shared_memory
)
3441 if (cuda_array_is_read_only_scalar(array
))
3444 group
->shared_bound
= create_bound_list(ctx
, array
->n_index
);
3445 if (!can_tile_for_shared_memory(array
, group
->access
,
3446 group
->shared_bound
)) {
3447 free_bound_list(group
->shared_bound
, array
->n_index
);
3448 group
->shared_bound
= NULL
;
3452 /* Is the size of the tile specified by "bound" smaller than the sum of
3453 * the sizes of the tiles specified by "bound1" and "bound2"?
3455 static int smaller_tile(unsigned n_index
, struct cuda_array_bound
*bound
,
3456 struct cuda_array_bound
*bound1
, struct cuda_array_bound
*bound2
)
3459 isl_int size
, size1
, size2
;
3462 isl_int_init(size1
);
3463 isl_int_init(size2
);
3465 tile_size(n_index
, bound
, &size
);
3466 tile_size(n_index
, bound1
, &size1
);
3467 tile_size(n_index
, bound2
, &size2
);
3469 isl_int_sub(size
, size
, size1
);
3470 isl_int_sub(size
, size
, size2
);
3471 smaller
= isl_int_is_neg(size
);
3473 isl_int_clear(size2
);
3474 isl_int_clear(size1
);
3475 isl_int_clear(size
);
3480 /* Given an initial grouping of array references and shared memory tiles
3481 * for each group that allows for a shared memory tile, merge two groups
3482 * if both have a shared memory tile, the merged group also has
3483 * a shared memory tile and the size of the tile for the merge group
3484 * is smaller than the sum of the tile sizes of the individual groups.
3486 * Return the number of group leaders after merging.
3488 static int group_common_shared_memory_tile(struct cuda_gen
*gen
,
3489 struct cuda_array_info
*array
, int n
,
3490 struct cuda_array_ref_group
**groups
, int *leader
, int n_group
)
3493 isl_ctx
*ctx
= isl_space_get_ctx(array
->dim
);
3495 for (i
= 0; n_group
> 1 && i
< n
; ++i
) {
3499 if (!groups
[i
]->shared_bound
)
3501 for (j
= i
- 1; j
>= 0; --j
) {
3504 struct cuda_array_bound
*shared_bound
;
3508 if (!groups
[j
]->shared_bound
)
3511 map
= isl_map_intersect(isl_map_copy(groups
[l
]->access
),
3512 isl_map_copy(groups
[j
]->access
));
3513 empty
= isl_map_is_empty(map
);
3519 map
= isl_map_union(isl_map_copy(groups
[l
]->access
),
3520 isl_map_copy(groups
[j
]->access
));
3521 shared_bound
= create_bound_list(ctx
, array
->n_index
);
3522 if (!can_tile_for_shared_memory(array
, map
,
3524 !smaller_tile(array
->n_index
, shared_bound
,
3525 groups
[l
]->shared_bound
,
3526 groups
[j
]->shared_bound
)) {
3528 free_bound_list(shared_bound
, array
->n_index
);
3532 free_bound_list(groups
[j
]->shared_bound
,
3534 groups
[j
]->shared_bound
= shared_bound
;
3535 isl_map_free(groups
[j
]->access
);
3536 groups
[j
]->access
= map
;
3537 groups
[j
]->n_ref
+= groups
[l
]->n_ref
;
3546 /* Extract an array of array reference groups from the array of references
3547 * and the grouping information in "leader".
3549 * Store the results in array->n_group and array->groups.
3551 static void extract_array_groups(isl_ctx
*ctx
, struct cuda_array_info
*array
,
3552 int n
, struct cuda_array_ref_group
**groups
, int *leader
, int n_group
)
3556 for (i
= 2; i
< n
; ++i
)
3557 leader
[i
] = leader
[leader
[i
]];
3559 array
->n_group
= n_group
;
3560 array
->groups
= isl_alloc_array(ctx
, struct cuda_array_ref_group
*,
3562 assert(array
->groups
);
3565 for (i
= 0; i
< n
; ++i
) {
3567 struct cuda_stmt_access
**refs
;
3569 if (leader
[i
] != i
) {
3570 groups
[i
]->refs
= NULL
;
3571 free_array_ref_group(groups
[i
], array
->n_index
);
3575 refs
= isl_alloc_array(ctx
, struct cuda_stmt_access
*,
3579 for (k
= i
; k
< n
; ++k
)
3580 if (leader
[k
] == i
) {
3581 refs
[l
++] = *groups
[k
]->refs
;
3582 (*groups
[k
]->refs
)->group
= j
;
3585 groups
[i
]->refs
= refs
;
3587 array
->groups
[j
++] = groups
[i
];
3591 /* Group array references that should be considered together when
3592 * deciding whether to access them from private, shared or global memory.
3594 * In particular, if two array references overlap and if one of them
3595 * is a write, then the two references are grouped together.
3596 * Furthermore, if two groups admit a shared memory tile and if the
3597 * combination of the two also admits a shared memory tile, we merge
3600 * During the construction the group->refs field points to a single
3601 * array reference inside the array of array references, while
3602 * group->n_ref contains the number of element in leader that
3603 * (directly or indirectly) point to this group, provided the group
3606 static void group_array_references(struct cuda_gen
*gen
,
3607 struct cuda_array_info
*array
, __isl_keep isl_union_map
*sched
)
3611 isl_ctx
*ctx
= isl_union_map_get_ctx(sched
);
3612 struct cuda_array_ref_group
**groups
;
3615 groups
= isl_calloc_array(ctx
, struct cuda_array_ref_group
*,
3619 n
= populate_array_references(array
, sched
, groups
);
3621 leader
= isl_alloc_array(ctx
, int, n
);
3624 n_group
= group_overlapping_writes(gen
, n
, groups
, leader
);
3626 for (i
= 0; i
< n
; ++i
)
3628 compute_group_shared_bound(gen
, array
, groups
[i
]);
3630 n_group
= group_common_shared_memory_tile(gen
, array
, n
, groups
,
3633 extract_array_groups(ctx
, array
, n
, groups
, leader
, n_group
);
3639 /* Take tiled_sched, project it onto the shared tile loops and
3640 * the loops that will be wrapped over the threads,
3641 * parametrize the shared tile loops and store the result in gen->shared_sched.
3642 * The position of the first of these parameters is stored in gen->first_shared.
3643 * Also compute a projection that projects out the loops that will be
3644 * wrapped over the threads and store this projection in gen->shared_proj.
3646 static void compute_shared_sched(struct cuda_gen
*gen
)
3651 isl_union_map
*sched
;
3653 sched
= isl_union_map_copy(gen
->tiled_sched
);
3655 dim
= isl_union_map_get_space(sched
);
3656 gen
->first_shared
= isl_space_dim(dim
, isl_dim_param
);
3657 proj
= projection(dim
, gen
->tiled_len
, gen
->shared_len
+ gen
->n_block
);
3658 sched
= isl_union_map_apply_range(sched
, isl_union_map_from_map(proj
));
3660 dim
= isl_union_map_get_space(sched
);
3661 par
= parametrization(dim
, gen
->shared_len
+ gen
->n_block
,
3662 0, gen
->shared_len
, "g");
3663 sched
= isl_union_map_intersect_range(sched
,
3664 isl_union_set_from_set(par
));
3666 dim
= isl_union_map_get_space(sched
);
3667 proj
= projection(dim
, gen
->shared_len
+ gen
->n_block
, gen
->shared_len
);
3669 gen
->shared_sched
= sched
;
3670 gen
->shared_proj
= isl_union_map_from_map(proj
);
3673 /* Group references of all arrays in the program.
3675 static void group_references(struct cuda_gen
*gen
)
3678 isl_union_map
*sched
;
3680 sched
= isl_union_map_apply_range(isl_union_map_copy(gen
->shared_sched
),
3681 isl_union_map_copy(gen
->shared_proj
));
3683 for (i
= 0; i
< gen
->n_array
; ++i
)
3684 group_array_references(gen
, &gen
->array
[i
], sched
);
3686 isl_union_map_free(sched
);
3689 /* Free all array information that is local to the current kernel.
3691 static void free_local_array_info(struct cuda_gen
*gen
)
3695 for (i
= 0; i
< gen
->n_array
; ++i
) {
3696 struct cuda_array_info
*array
= &gen
->array
[i
];
3698 for (j
= 0; j
< array
->n_group
; ++j
)
3699 free_array_ref_group(array
->groups
[j
], array
->n_index
);
3700 free(array
->groups
);
3702 if (array
->n_group
== 0)
3704 for (j
= 0; j
< gen
->array
[i
].n_index
; ++j
) {
3705 isl_pw_aff_free(gen
->array
[i
].local_bound
[j
]);
3706 gen
->array
[i
].local_bound
[j
] = NULL
;
3711 /* The sizes of the arrays on the host that have been computed by
3712 * extract_array_info may depend on the parameters. Use the extra
3713 * constraints on the parameters that are valid at "host_domain"
3714 * to simplify these expressions.
3716 static void localize_bounds(struct cuda_gen
*gen
,
3717 __isl_keep isl_set
*host_domain
)
3722 context
= isl_set_copy(host_domain
);
3723 context
= isl_set_params(context
);
3725 for (i
= 0; i
< gen
->n_array
; ++i
) {
3726 struct cuda_array_info
*array
= &gen
->array
[i
];
3728 if (array
->n_group
== 0)
3731 for (j
= 0; j
< array
->n_index
; ++j
) {
3734 pwaff
= isl_pw_aff_copy(array
->bound
[j
]);
3735 pwaff
= isl_pw_aff_gist(pwaff
, isl_set_copy(context
));
3736 array
->local_bound
[j
] = pwaff
;
3739 isl_set_free(context
);
3742 /* Set gen->tile_len and gen->n_parallel to those of the first statement
3743 * in the statement list u.
3744 * Because of the way the schedule is constructed, the other statements
3745 * in the list, if any, should have the same values for these properties.
3747 static void set_tile_len(struct cuda_gen
*gen
, struct clast_user_stmt
*u
)
3750 struct cuda_stmt
*stmt
;
3752 nr
= atoi(u
->statement
->name
+ 2);
3753 stmt
= &gen
->stmts
[nr
];
3755 gen
->tile_len
= stmt
->tile_len
;
3756 gen
->n_parallel
= stmt
->n_parallel
;
3759 /* Extract a description of the grid, i.e., the possible values
3760 * of the block ids, from gen->tiled_sched.
3761 * The block ids are parameters in gen->tiled_sched.
3762 * We simply need to change them into set dimensions.
3764 static __isl_give isl_set
*extract_grid(struct cuda_gen
*gen
)
3769 grid
= isl_union_map_params(isl_union_map_copy(gen
->tiled_sched
));
3770 grid
= isl_set_from_params(grid
);
3771 grid
= isl_set_add_dims(grid
, isl_dim_set
, gen
->n_grid
);
3772 for (i
= 0; i
< gen
->n_grid
; ++i
) {
3776 snprintf(name
, sizeof(name
), "b%d", i
);
3777 pos
= isl_set_find_dim_by_name(grid
, isl_dim_param
, name
);
3779 grid
= isl_set_equate(grid
, isl_dim_param
, pos
, isl_dim_set
, i
);
3780 grid
= isl_set_project_out(grid
, isl_dim_param
, pos
, 1);
3786 /* Print the effective grid size as a list of the sizes in each
3787 * dimension, from innermost to outermost.
3789 * The grid size specified by the user or set by default
3790 * in read_grid_sizes() and applied in tile_schedule(),
3791 * may be too large for the given code in the sense that
3792 * it may contain blocks that don't need to execute anything.
3793 * We therefore don't print this grid size, but instead the
3794 * smallest grid size that ensures that all blocks that actually
3795 * execute code are included in the grid.
3797 * For each block dimension, we compute the maximal value of the block id
3800 static void print_grid_size(struct cuda_gen
*gen
, __isl_take isl_set
*context
)
3806 if (gen
->n_grid
== 0) {
3807 isl_set_free(context
);
3811 grid
= extract_grid(gen
);
3813 prn
= isl_printer_to_file(gen
->ctx
, gen
->cuda
.host_c
);
3814 prn
= isl_printer_set_output_format(prn
, ISL_FORMAT_C
);
3816 prn
= isl_printer_print_str(prn
, "(");
3817 for (i
= gen
->n_grid
- 1; i
>= 0; --i
) {
3820 isl_pw_aff
*bound
= isl_set_dim_max(isl_set_copy(grid
), i
);
3822 bound
= isl_pw_aff_coalesce(bound
);
3823 bound
= isl_pw_aff_gist(bound
, isl_set_copy(context
));
3825 space
= isl_pw_aff_get_domain_space(bound
);
3826 one
= isl_aff_zero_on_domain(isl_local_space_from_space(space
));
3827 one
= isl_aff_add_constant_si(one
, 1);
3828 bound
= isl_pw_aff_add(bound
, isl_pw_aff_from_aff(one
));
3829 prn
= isl_printer_print_pw_aff(prn
, bound
);
3830 isl_pw_aff_free(bound
);
3833 prn
= isl_printer_print_str(prn
, ", ");
3835 prn
= isl_printer_print_str(prn
, ")");
3837 isl_printer_free(prn
);
3839 isl_set_free(context
);
3842 /* This function is called for each leaf in the clast of the host code.
3843 * We first specialize the schedule to the site of the leaf, compute
3844 * the size of shared memory and then print the body of host code
3845 * and the associated kernel (through a call to print_kernel_body).
3847 static void print_host_user(struct clast_printer_info
*code
,
3848 struct clast_user_stmt
*u
)
3850 struct cuda_gen
*gen
= code
->user
;
3853 isl_set
*host_domain
;
3854 isl_union_map
*access
;
3855 isl_union_map
*local_sched
;
3856 isl_union_set
*arrays
;
3858 set_tile_len(gen
, u
);
3861 host_domain
= extract_entire_host_domain(&u
->stmt
);
3863 local_sched
= isl_union_map_intersect_range(
3864 isl_union_map_copy(gen
->sched
),
3865 isl_union_set_from_set(extend(isl_set_copy(host_domain
),
3866 gen
->untiled_len
)));
3867 access
= isl_union_map_union(isl_union_map_copy(gen
->read
),
3868 isl_union_map_copy(gen
->write
));
3869 access
= isl_union_map_apply_domain(access
,
3870 isl_union_map_copy(local_sched
));
3871 arrays
= isl_union_map_range(access
);
3873 print_indent(code
->dst
, code
->indent
);
3874 fprintf(code
->dst
, "dim3 k%d_dimBlock", gen
->kernel_id
);
3875 print_reverse_list(code
->dst
, gen
->n_block
, gen
->block_dim
);
3876 fprintf(code
->dst
, ";\n");
3878 gen
->tiled_sched
= tile_schedule(gen
, local_sched
);
3879 gen
->tiled_sched
= parametrize_tiled_schedule(gen
, gen
->tiled_sched
);
3880 gen
->tiled_sched
= scale_tile_loops(gen
, gen
->tiled_sched
);
3882 print_indent(code
->dst
, code
->indent
);
3883 fprintf(code
->dst
, "dim3 k%d_dimGrid", gen
->kernel_id
);
3884 print_grid_size(gen
, isl_set_params(isl_set_copy(host_domain
)));
3885 fprintf(code
->dst
, ";\n");
3887 gen
->local_sched
= isl_union_map_copy(gen
->tiled_sched
);
3889 dim
= isl_union_map_get_space(gen
->local_sched
);
3890 par
= parametrization(dim
, gen
->tiled_len
, 0, gen
->shared_len
, "g");
3891 gen
->local_sched
= isl_union_map_intersect_range(gen
->local_sched
,
3892 isl_union_set_from_set(par
));
3894 gen
->local_sched
= thread_tile_schedule(gen
, gen
->local_sched
);
3895 gen
->local_sched
= scale_thread_tile_loops(gen
, gen
->local_sched
);
3897 gen
->private_access
= NULL
;
3898 compute_shared_sched(gen
);
3899 gen
->privatization
= compute_privatization(gen
);
3900 group_references(gen
);
3901 compute_private_size(gen
);
3902 check_shared_memory_bound(gen
);
3903 localize_bounds(gen
, host_domain
);
3905 gen
->local_sched
= interchange_for_unroll(gen
, gen
->local_sched
);
3907 print_kernel_launch(gen
, arrays
);
3909 fprintf(gen
->cuda
.kernel_c
, "{\n");
3911 print_kernel_body(gen
, host_domain
, gen
->tiled_sched
);
3913 fprintf(gen
->cuda
.kernel_c
, "}\n");
3915 free_local_array_info(gen
);
3916 isl_map_free(gen
->privatization
);
3917 isl_union_map_free(gen
->private_access
);
3918 isl_union_map_free(gen
->local_sched
);
3919 isl_union_map_free(gen
->tiled_sched
);
3920 isl_union_map_free(gen
->shared_sched
);
3921 isl_union_map_free(gen
->shared_proj
);
3922 isl_union_set_free(arrays
);
3923 isl_set_free(host_domain
);
3925 free(gen
->tile_size
);
3929 /* Use CLooG to generate code for the outer gen->tile_first loops
3930 * of the global schedule in gen->sched.
3931 * The pretty printing of this code is handled by print_clast,
3932 * which calls print_host_user for each kernel invocation location.
3934 static void print_cloog_host_code(struct cuda_gen
*gen
)
3938 isl_union_map
*sched
;
3939 CloogOptions
*options
;
3940 CloogDomain
*cloog_context
;
3941 CloogUnionDomain
*ud
;
3943 struct clast_stmt
*stmt
;
3946 options
= cloog_options_malloc(gen
->state
);
3947 options
->language
= CLOOG_LANGUAGE_C
;
3949 options
->strides
= 1;
3950 options
->stop
= gen
->tile_first
;
3951 options
->f
= gen
->untiled_len
;
3952 options
->l
= gen
->untiled_len
;
3953 options
->save_domains
= 1;
3954 options
->noscalars
= 1;
3956 sched
= isl_union_map_copy(gen
->sched
);
3957 ud
= cloog_union_domain_from_isl_union_map(sched
);
3958 for (i
= 0; i
< options
->stop
; ++i
) {
3959 snprintf(name
, sizeof(name
), "h%d", i
);
3960 ud
= cloog_union_domain_set_name(ud
, CLOOG_SCAT
, i
, name
);
3962 context
= isl_set_copy(gen
->context
);
3963 cloog_context
= cloog_domain_from_isl_set(context
);
3964 input
= cloog_input_alloc(cloog_context
, ud
);
3966 stmt
= cloog_clast_create_from_input(input
, options
);
3968 gen
->code
.indent
= 0;
3969 gen
->code
.dst
= gen
->cuda
.host_c
;
3970 gen
->code
.print_user_stmt
= NULL
;
3971 gen
->code
.print_user_stmt_list
= &print_host_user
;
3972 gen
->code
.print_for_head
= NULL
;
3973 gen
->code
.print_for_foot
= NULL
;
3974 gen
->code
.user
= gen
;
3975 print_clast(&gen
->code
, stmt
);
3977 cloog_clast_free(stmt
);
3978 cloog_options_free(options
);
3979 fprintf(gen
->cuda
.host_c
, "\n");
3982 void print_cuda_macros(struct cuda_gen
*gen
)
3984 const char *macros
=
3985 "#define cudaCheckReturn(ret) assert((ret) == cudaSuccess)\n"
3986 "#define cudaCheckKernel()"
3987 " assert(cudaGetLastError() == cudaSuccess)\n\n";
3988 fputs(macros
, gen
->cuda
.host_c
);
3991 void print_host_code(struct cuda_gen
*gen
)
3993 fprintf(gen
->cuda
.host_c
, "{\n");
3994 print_cloog_macros(gen
->cuda
.host_c
);
3995 print_cloog_macros(gen
->cuda
.kernel_c
);
3997 print_cuda_macros(gen
);
3999 declare_device_arrays(gen
);
4001 allocate_device_arrays(gen
);
4002 copy_arrays_to_device(gen
);
4005 print_cloog_host_code(gen
);
4007 copy_arrays_from_device(gen
);
4008 free_device_arrays(gen
);
4010 fprintf(gen
->cuda
.host_c
, "}\n");
4013 __isl_give isl_set
*add_context_from_str(__isl_take isl_set
*set
,
4022 ctx
= isl_set_get_ctx(set
);
4023 context
= isl_set_read_from_str(ctx
, str
);
4024 context
= isl_set_align_params(context
, isl_set_get_space(set
));
4025 set
= isl_set_intersect(set
, context
);
4030 __isl_give isl_union_map
*extract_sizes_from_str(isl_ctx
*ctx
, const char *str
)
4034 return isl_union_map_read_from_str(ctx
, str
);
4037 /* Return the union of all iteration domains of the gen->stmts[i].
4039 static __isl_give isl_union_set
*extract_domain(struct cuda_gen
*gen
)
4042 isl_union_set
*domain
;
4044 domain
= isl_union_set_empty(isl_set_get_space(gen
->context
));
4045 for (i
= 0; i
< gen
->n_stmts
; ++i
) {
4048 domain_i
= isl_set_copy(gen
->stmts
[i
].domain
);
4049 domain
= isl_union_set_union(domain
,
4050 isl_union_set_from_set(domain_i
));
4056 /* Information about the outermost tilable bands in the forest of bands.
4058 * tile_len and n_parallel are only sets on band_info structures
4059 * that correspond to outermost bands. For other bands (in particular,
4060 * ancestors of the outermost bands), n_parallal is set to 0.
4062 * prefix is the (padded) schedule leading up to the outermost tilable bands.
4064 * tile_first is the number of schedule dimensions in prefix.
4066 * suffix is the schedule of the outermost tilable bands and their descendants.
4069 struct cuda_gen
*gen
;
4073 isl_union_map
*prefix
;
4074 isl_union_map
*suffix
;
4077 /* Set tile_len and n_parallel of the statement to that of
4078 * their outermost band, recorded in the band_info.
4080 static int set_stmt_tile_len(__isl_take isl_map
*map
, void *user
)
4082 struct band_info
*info
= user
;
4084 struct cuda_stmt
*stmt
;
4086 nr
= atoi(isl_map_get_tuple_name(map
, isl_dim_in
) + 2);
4087 stmt
= &info
->gen
->stmts
[nr
];
4089 stmt
->tile_len
= info
->tile_len
;
4090 stmt
->n_parallel
= info
->n_parallel
;
4097 static void list_select_outer_band(struct cuda_gen
*gen
,
4098 __isl_take isl_band_list
*list
, int pos
, struct band_info
*list_info
);
4100 /* Check if this band has any parallel loops. If so, take it as
4101 * the outermost tilable band. If not, continue looking for the
4102 * outermost tilable band in the children of the current band.
4104 static void band_select_outer_band(struct cuda_gen
*gen
,
4105 __isl_take isl_band
*band
, int pos
, struct band_info
*info
)
4107 int n
= isl_band_n_member(band
);
4110 for (n_parallel
= 0; n_parallel
< n
; ++n_parallel
)
4111 if (!isl_band_member_is_zero_distance(band
, n_parallel
))
4114 info
->n_parallel
= n_parallel
;
4117 info
->tile_first
= pos
;
4119 info
->prefix
= isl_band_get_prefix_schedule(band
);
4120 info
->suffix
= isl_union_map_flat_range_product(
4121 isl_band_get_partial_schedule(band
),
4122 isl_band_get_suffix_schedule(band
));
4123 isl_union_map_foreach_map(info
->prefix
,
4124 &set_stmt_tile_len
, info
);
4125 } else if (isl_band_has_children(band
)) {
4126 isl_band_list
*children
;
4127 children
= isl_band_get_children(band
);
4128 list_select_outer_band(gen
, children
, pos
+ n
, info
);
4131 info
->tile_first
= pos
+ n
;
4133 info
->prefix
= isl_union_map_flat_range_product(
4134 isl_band_get_prefix_schedule(band
),
4135 isl_band_get_partial_schedule(band
));
4136 info
->suffix
= isl_band_get_suffix_schedule(band
);
4137 isl_union_map_foreach_map(info
->prefix
,
4138 &set_stmt_tile_len
, info
);
4141 isl_band_free(band
);
4144 /* Comparison function that returns a non-zero value for band_infos
4145 * with different tile_len fields or different n_parallel fields.
4147 static int cmp_band(const void *p1
, const void *p2
)
4149 const struct band_info
*info1
= p1
;
4150 const struct band_info
*info2
= p2
;
4152 if (info1
->tile_len
!= info2
->tile_len
)
4153 return info1
->tile_len
- info2
->tile_len
;
4155 return info1
->n_parallel
- info2
->n_parallel
;
4158 /* Extend "umap" with coordinates with fixed value "val"
4159 * to a total length of "dst_len", assuming the original dimension is "src_len".
4161 static __isl_give isl_union_map
*extend_range(__isl_take isl_union_map
*umap
,
4162 int src_len
, int dst_len
, int val
)
4168 dim
= isl_union_map_get_space(umap
);
4169 map
= isl_map_reverse(projection(dim
, dst_len
, src_len
));
4170 for (i
= src_len
; i
< dst_len
; ++i
)
4171 map
= isl_map_fix_si(map
, isl_dim_out
, i
, val
);
4173 umap
= isl_union_map_apply_range(umap
, isl_union_map_from_map(map
));
4178 /* Group bands with the same values for tile_len and n_parallel.
4179 * The prefix schedule is then extended with a fixed coordinate that
4180 * is different for each such group.
4181 * Note that the actual values for this coordinate are not important.
4182 * The bands have already been effectively separated at a higher level
4183 * or they are independent and may be executed in parallel.
4184 * The list of band_info has been sorted before this functions is called.
4186 static void separate_bands(struct band_info
*info
, int n
)
4191 for (i
= 0; i
< n
; ++i
) {
4192 int l
= info
[i
].tile_first
;
4195 (info
[i
].tile_len
!= info
[i
- 1].tile_len
||
4196 info
[i
].n_parallel
!= info
[i
- 1].n_parallel
))
4199 info
[i
].prefix
= extend_range(info
[i
].prefix
,
4201 info
[i
].tile_first
= l
+ 1;
4205 /* Select the outermost bands in the elements of the list, align
4206 * their prefix schedules, separate bands with different values
4207 * for tile_len and/or n_parallel and then combine the resulting
4208 * prefix and suffix schedules into a single pair of prefix and
4209 * suffix schedules for the entire list.
4211 static void list_select_outer_band(struct cuda_gen
*gen
,
4212 __isl_take isl_band_list
*list
, int pos
, struct band_info
*list_info
)
4216 int n
= isl_band_list_n_band(list
);
4217 isl_ctx
*ctx
= isl_band_list_get_ctx(list
);
4218 struct band_info
*info
;
4220 isl_union_map
*prefix
;
4221 isl_union_map
*suffix
;
4224 info
= isl_calloc_array(ctx
, struct band_info
, n
);
4228 for (i
= 0; i
< n
; ++i
) {
4229 band
= isl_band_list_get_band(list
, i
);
4230 band_select_outer_band(gen
, band
, pos
, &info
[i
]);
4231 if (info
[i
].tile_first
> max_tile_first
)
4232 max_tile_first
= info
[i
].tile_first
;
4235 for (i
= 0; i
< n
; ++i
) {
4236 if (info
[i
].tile_first
== max_tile_first
)
4238 info
[i
].prefix
= extend_range(info
[i
].prefix
,
4239 info
[i
].tile_first
, max_tile_first
, 0);
4240 info
[i
].tile_first
= max_tile_first
;
4243 qsort(info
, n
, sizeof(struct band_info
), &cmp_band
);
4245 for (i
= 0; i
< n
- 1; ++i
)
4246 if (info
[i
].tile_len
!= info
[i
+ 1].tile_len
||
4247 info
[i
].n_parallel
!= info
[i
+ 1].n_parallel
)
4251 separate_bands(info
, n
);
4253 prefix
= info
[0].prefix
;
4254 suffix
= info
[0].suffix
;
4256 for (i
= 1; i
< n
; ++i
) {
4257 prefix
= isl_union_map_union(prefix
, info
[i
].prefix
);
4258 suffix
= isl_union_map_union(suffix
, info
[i
].suffix
);
4261 list_info
->tile_first
= info
[0].tile_first
;
4262 list_info
->tile_len
= -1;
4263 list_info
->prefix
= prefix
;
4264 list_info
->suffix
= suffix
;
4266 isl_band_list_free(list
);
4270 /* Select the outermost tilable band that (by construction)
4271 * has at least one parallel loop.
4272 * The starting position of the aligned band is stored in the pair
4274 * The sizes and number of parallel loops may be different in different
4275 * parts of the band forest and are therefore stored in the cuda_stmts.
4277 * Return the complete schedule, with the tilable bands aligned
4278 * at gen->tile_first and padded with zero, if needed.
4280 static __isl_give isl_union_map
*select_outer_tilable_band(struct cuda_gen
*gen
,
4281 __isl_keep isl_schedule
*schedule
)
4283 isl_band_list
*list
;
4284 struct band_info info
;
4286 gen
->n_parallel
= 0;
4289 list
= isl_schedule_get_band_forest(schedule
);
4291 list_select_outer_band(gen
, list
, 0, &info
);
4293 gen
->tile_first
= info
.tile_first
;
4294 info
.suffix
= align_range(info
.suffix
);
4296 return isl_union_map_flat_range_product(info
.prefix
, info
.suffix
);
4299 /* Set gen->untiled_len to the number of scheduling dimensions
4300 * for the schedule of the first domain.
4301 * We assume here that this number is the same for all domains.
4303 static int set_untiled_len(__isl_take isl_map
*map
, void *user
)
4305 unsigned *untiled_len
= user
;
4307 *untiled_len
= isl_map_dim(map
, isl_dim_out
);
4313 /* Compute an appropriate schedule based on the accesses in
4314 * gen->read and gen->write.
4316 * We first compute dependences and then use those to compute
4317 * a schedule that has a parallel loop in each tilable band.
4318 * Finally, we select the outermost tilable band.
4320 static void compute_schedule(struct cuda_gen
*gen
,
4321 __isl_take isl_union_map
*sched
)
4323 isl_union_set
*domain
;
4324 isl_union_map
*empty
;
4325 isl_union_map
*dep_raw
, *dep2
, *dep3
, *dep
;
4326 isl_union_map
*uninitialized
;
4327 isl_schedule
*schedule
;
4329 empty
= isl_union_map_empty(isl_union_map_get_space(sched
));
4331 isl_union_map_compute_flow(isl_union_map_copy(gen
->read
),
4332 isl_union_map_copy(gen
->write
), empty
,
4333 isl_union_map_copy(sched
),
4334 &dep_raw
, NULL
, &uninitialized
, NULL
);
4335 isl_union_map_compute_flow(isl_union_map_copy(gen
->write
),
4336 isl_union_map_copy(gen
->write
),
4337 isl_union_map_copy(gen
->read
),
4338 isl_union_map_copy(sched
),
4339 &dep2
, &dep3
, NULL
, NULL
);
4340 isl_union_map_free(sched
);
4342 gen
->copy_in
= isl_union_map_range(uninitialized
);
4344 dep
= isl_union_map_union(dep2
, dep3
);
4345 dep
= isl_union_map_union(dep
, dep_raw
);
4346 dep
= isl_union_map_coalesce(dep
);
4348 domain
= extract_domain(gen
);
4349 schedule
= isl_union_set_compute_schedule(isl_union_set_copy(domain
),
4350 isl_union_map_copy(dep
), dep
);
4352 sched
= select_outer_tilable_band(gen
, schedule
);
4354 isl_union_map_foreach_map(sched
, &set_untiled_len
, &gen
->untiled_len
);
4355 sched
= isl_union_map_intersect_domain(sched
, domain
);
4358 isl_schedule_free(schedule
);
4361 static struct cuda_stmt_access
**expr_extract_access(struct pet_expr
*expr
,
4362 struct cuda_stmt_access
**next_access
)
4364 struct cuda_stmt_access
*access
;
4365 isl_ctx
*ctx
= isl_map_get_ctx(expr
->acc
.access
);
4367 access
= isl_alloc_type(ctx
, struct cuda_stmt_access
);
4369 access
->next
= NULL
;
4370 access
->read
= expr
->acc
.read
;
4371 access
->write
= expr
->acc
.write
;
4372 access
->access
= isl_map_copy(expr
->acc
.access
);
4374 *next_access
= access
;
4375 next_access
= &(*next_access
)->next
;
4379 static struct cuda_stmt_access
**expr_extract_accesses(struct pet_expr
*expr
,
4380 struct cuda_stmt_access
**next_access
)
4384 for (i
= 0; i
< expr
->n_arg
; ++i
)
4385 next_access
= expr_extract_accesses(expr
->args
[i
],
4388 if (expr
->type
== pet_expr_access
)
4389 next_access
= expr_extract_access(expr
, next_access
);
4394 static void pet_stmt_extract_accesses(struct cuda_stmt
*stmt
)
4396 struct cuda_stmt_access
**next_access
= &stmt
->accesses
;
4398 stmt
->accesses
= NULL
;
4399 expr_extract_accesses(stmt
->body
, next_access
);
4402 /* Return an array of cuda_stmt representing the statements in "scop".
4404 static struct cuda_stmt
*extract_stmts(isl_ctx
*ctx
, struct pet_scop
*scop
,
4405 __isl_keep isl_set
*context
)
4408 struct cuda_stmt
*stmts
;
4410 stmts
= isl_calloc_array(ctx
, struct cuda_stmt
, scop
->n_stmt
);
4413 for (i
= 0; i
< scop
->n_stmt
; ++i
) {
4414 struct cuda_stmt
*s
= &stmts
[i
];
4416 s
->domain
= isl_set_copy(scop
->stmts
[i
]->domain
);
4417 s
->domain
= isl_set_intersect_params(s
->domain
,
4418 isl_set_copy(context
));
4419 s
->body
= scop
->stmts
[i
]->body
;
4420 pet_stmt_extract_accesses(s
);
4426 /* Replace the scop in the "input" file by equivalent code
4427 * that uses the GPU. "scop" is assumed to correspond to this scop.
4429 * We first compute a schedule that respects the dependences
4430 * of the original program and select the outermost band
4431 * of tilable dimensions that has at least one parallel loop.
4432 * We then have three blocks of dimensions
4436 * The tilable band "B" is first tiled according to "tile" sizes, resulting
4441 * For each iteration of the T loop and for each array, we compute
4442 * the array elements accessed by that iteration, construct a rectangular
4443 * box around it and shift it to the origin. The result is used
4444 * as shared memory for the array.
4446 * We then split off at most 2 parallel loops from the T loops and
4447 * at most 3 parallel loops from the P loops
4451 * The T1/P1 loops are then tiled or "wrapped" over the blocks/threads,
4452 * according to "grid"/"block" sizes.
4454 * H T1T T1P T2 P1T P1P P2 G
4456 * Finally, the T1P and P1P iterators are equated to the block and
4457 * thread dimensions respectively and so are effectively removed.
4458 * The H loops are run on the host. The T1T, T2, P1T, P2 and G loops
4459 * are run on the GPU.
4461 * Code is generated in three stages. We first generate code for the
4462 * host (the H loops), with iterators h%d. Then, for each leaf node
4463 * of the resulting AST, we generate code for the shared loops (up to
4464 * and including T2), with iterators g%d and after equating the H loops
4465 * to h%d parameters and the T1P loops to the block dimensions.
4466 * Finally, we generate code for the remaining loops in a similar fashion.
4468 int generate_cuda(isl_ctx
*ctx
, struct pet_scop
*scop
,
4469 struct ppcg_options
*options
, const char *input
)
4471 isl_union_map
*sched
;
4472 struct cuda_gen gen
;
4477 scop
= pet_scop_align_params(scop
);
4480 gen
.context
= isl_set_copy(scop
->context
);
4481 gen
.context
= add_context_from_str(gen
.context
, options
->ctx
);
4482 gen
.sizes
= extract_sizes_from_str(ctx
, options
->sizes
);
4483 gen
.n_stmts
= scop
->n_stmt
;
4484 gen
.stmts
= extract_stmts(ctx
, scop
, gen
.context
);
4485 gen
.read
= pet_scop_collect_reads(scop
);
4486 gen
.write
= pet_scop_collect_writes(scop
);
4487 gen
.options
= options
;
4488 gen
.state
= cloog_isl_state_malloc(gen
.ctx
);
4491 cuda_open_files(&gen
.cuda
, input
);
4493 collect_array_info(&gen
);
4495 sched
= pet_scop_collect_schedule(scop
);
4497 compute_schedule(&gen
, sched
);
4499 print_host_code(&gen
);
4501 cloog_state_free(gen
.state
);
4502 clear_cuda_gen(&gen
);
4504 cuda_close_files(&gen
.cuda
);