On Mon, Aug 26, 2013 at 03:59:11PM +0400, Michael V. Zolotukhin wrote:
> As I currently see it, the given code would be expanded to something like
> this:
>
> // Create two versions of V: for host and for target
> int v;
> int v_target __attribute(target);
>
> // The same for TGT function
> int tgt ()
> {
> .. update v ..
> }
> int tgt_target () __attribute(target)
> {
> .. update v_target ..
> }
Actually, not two versions of those during the compilation, you have
just one v and one tgt, both have __attribute__(("omp declare target"))
on them (note, you can't specify that attribute manually).
And just when streaming into .gnu.target_lto_* sections you only stream
everything that has those attributes and types used by it, but nothing else.
>
> float
> bar (int x, int y, int z)
> {
> float b[1024], c[1024], s = 0;
> int i, j;
> baz (b, c, x);
> // #pragma omp target data map(to: b)
> vec<data_descriptor> data_desc;
> data_desc.push ({&b, 1024*sizeof(float), TO});
> GOMP_target_data (&data_desc);
Nope. It would be:
struct data_descriptor data_desc1[1] = { { &b, 1024*sizeof(float), TO } };
GOMP_target_data (-1, data_desc1, 1);
or so. The compiler always knows how many vector elements it needs, there
is no point in making the vector dynamic, and vec<> is a compiler data
structure, while you want to emit runtime code. The -1 in there stands
for missing device(device-id) clause, otherwise it would be the provided
device-id expression. For the if clause, the question is if we want to pass
it down to the runtime library too (as bool, defaulting to true if missing),
or do something else.
> {
> // #pragma omp target map(tofrom: c) map(from:s)
> data_desc.push ({&c, 1024*sizeof(float), TOFROM});
> data_desc.push ({&s, sizeof(float), FROM});
> GOMP_target_data (&data_desc); // Add mapping for S and C variables,
> // mapping for B shouldn't change
Nope, there is only one target data pragma, so you would use here just:
> GOMP_target (foo1, "foo1", &data_desc); // Call either FOO1 or offloaded
> // FOO1_TARGET with arguments
> // from vector DATA_DESC
struct data_descriptor data_desc2[2] = { ... };
GOMP_target (-1, bar.omp_fn.1, "bar.omp_fn.1", data_desc2, 2);
>
> // #pragma omp target update from(b, v)
> vec<data_descriptor> data_desc_update; // target update pragma require a
> // separate vector
> data_desc_update.push ({&b, 1024*sizeof(float), FROM});
> data_desc_update.push ({&v, sizeof(int), FROM});
> GOMP_target_data (&data_desc_update);
Similarly here.
> }
> return s;
> }
> void
> foo1 (vec<data_descriptor> data_desc)
> {
> float b = *data_desc[0].host_address;
> float c = *data_desc[1].host_address;
> float s = 0;
> int i;
> for (i = 0; i < 1024; i++)
> tgt (), s += b[i] * c[i];
> *data_desc[2].host_address = s;
No, I didn't mean you'd do this. omp-lower.c would simply create
a type here that would have the same layout as what would the runtime
library pass to it.
So it would be:
void
bar.omp_fn.1 (struct omp_target_data *.omp_data_in)
{
int i;
*.omp_data_in->s = 0;
for (i = 0; i < 1024; i++)
tgt (), *.omp_data_in->s += .omp_data_in->b[i] * .omp_data_in->c[i];
}
Just look what omplower pass does for normal OpenMP code, say
#pragma omp parallel, task etc.
Jakub
