Discussion:
OpenCL address space and mangling
(too old to reply)
Michele Scandale
2013-07-18 17:11:14 UTC
Permalink
Hello to everybody,

I've noticed a problem in the mangling of function names where address spaces
are used. The problem is related to opencl and cuda address spaces: the target
address space map is used to translate address space, but this conversion can
produce the same mangling from different address space (on X86 all opencl
address spaces are mapped to the address space zero). See attached example.

Commit r174688 introduced the usage of the target translation map. Reverting it
seems be ok as a solution. Do you agree with this solution?

Thanks in advance.

Best regards.

Michele Scandale


-------------- next part --------------
__attribute__((overloadable))
void foo(__private float *a) {
*a = 1;
}

__attribute__((overloadable))
void foo(__global float *a) {
*a = 1;
}

__attribute__((overloadable))
void foo(__local float *a) {
*a = 1;
}
Tanya Lattner
2013-07-19 06:40:41 UTC
Permalink
It doesn't make sense to me to not translate the address space to something meaningful (something other than a random number). Targets should define the address space map if they are supported.

What problem are you trying to solve?

-Tanya
Post by Michele Scandale
Hello to everybody,
I've noticed a problem in the mangling of function names where address spaces
are used. The problem is related to opencl and cuda address spaces: the target
address space map is used to translate address space, but this conversion can
produce the same mangling from different address space (on X86 all opencl
address spaces are mapped to the address space zero). See attached example.
Commit r174688 introduced the usage of the target translation map. Reverting it
seems be ok as a solution. Do you agree with this solution?
Thanks in advance.
Best regards.
Michele Scandale
<test.cl>
_______________________________________________
cfe-commits mailing list
cfe-commits at cs.uiuc.edu
http://lists.cs.uiuc.edu/mailman/listinfo/cfe-commits
Michele Scandale
2013-07-19 07:50:03 UTC
Permalink
Hi Tanya,

thanks for the answer.
Post by Tanya Lattner
It doesn't make sense to me to not translate the address space to something meaningful (something other than a random number). Targets should define the address space map if they are supported.
What problem are you trying to solve?
The fact is that on targets, like X86, the address space translation produces
for opencl (and cuda) address spaces the same number (zero) because as you said
the target itself does not have different physical address spaces.

In the example attached to the previous mail, you can see that if you try to
compile it for X86, clang crashes because:

__attribute__((overloadable)) void foo(__global float *a)

and

__attribute__((overloadable)) void foo(__local float *a)

are mangled in the same way because of address space translation.

The logical address space information is lost with this mangling scheme.

My doubt is about the handling of address space information in the mangling:
- if the information about logical address spaces must be preserved then the
usage of target translation map cannot be used
- otherwise I don't see why I should manage in different way those cases where
the private address space is used from those where the mapping is the private
address space.

From a quick discussion on the IRC channel I found that information must be
preserved because I cannot assume that if two overloaded functions where the
unique difference is an address space qualifier then if the two address spaces
are the same the actions in those function are the same.

I agree with you that the clang internal encoding of opencl/cuda address spaces
is not a wonderful solution, but at least it preserves the information
correctly. Maybe another solution (it's just a claim, I don't know if it's
feasible) can be the mangling specialization for opencl and cuda, adding a
prefix that identifies the programming model and using canonical values for the
address spaces (1 = opencl/cuda global, 2 = opencl_local/cuda_shared, 3 =
opencl/cuda constant).

I hope to have been more clear presenting the problem.

Thanks in advance.

Best Regards,

Michele Scandale
Tanya Lattner
2013-07-22 18:30:48 UTC
Permalink
Michele,
Post by Michele Scandale
Hi Tanya,
thanks for the answer.
Post by Tanya Lattner
It doesn't make sense to me to not translate the address space to something meaningful (something other than a random number). Targets should define the address space map if they are supported.
What problem are you trying to solve?
The fact is that on targets, like X86, the address space translation produces for opencl (and cuda) address spaces the same number (zero) because as you said the target itself does not have different physical address spaces.
__attribute__((overloadable)) void foo(__global float *a)
and
__attribute__((overloadable)) void foo(__local float *a)
are mangled in the same way because of address space translation.
The logical address space information is lost with this mangling scheme.
- if the information about logical address spaces must be preserved then the usage of target translation map cannot be used
- otherwise I don't see why I should manage in different way those cases where the private address space is used from those where the mapping is the private address space.
From a quick discussion on the IRC channel I found that information must be preserved because I cannot assume that if two overloaded functions where the unique difference is an address space qualifier then if the two address spaces are the same the actions in those function are the same.
I understand what you are saying, but I'm arguing that the Target needs to define the address space map in order for this to mean anything at all. If its not defined, then what does global or local mean? What happens when this is lowered to LLVM IR?

Why not define an address space map for X86?

-Tanya
Post by Michele Scandale
I agree with you that the clang internal encoding of opencl/cuda address spaces is not a wonderful solution, but at least it preserves the information correctly. Maybe another solution (it's just a claim, I don't know if it's feasible) can be the mangling specialization for opencl and cuda, adding a prefix that identifies the programming model and using canonical values for the address spaces (1 = opencl/cuda global, 2 = opencl_local/cuda_shared, 3 = opencl/cuda constant).
I hope to have been more clear presenting the problem.
Thanks in advance.
Best Regards,
Michele Scandale
Michele Scandale
2013-07-23 10:21:02 UTC
Permalink
Hi Tanya,
Post by Tanya Lattner
I understand what you are saying, but I'm arguing that the Target needs to define the address space map in order for this to mean anything at all. If its not defined, then what does global or local mean? What happens when this is lowered to LLVM IR?
Why not define an address space map for X86?
The trivial address map is the correct one for X86, since it does not have
separate address spaces.
The problem is that at the language level, address space specifiers carry
additional information over that related to the memory allocation. In
particular, the specifiers cause function prototypes to refer to different
functions even if they differ only for an address space specifier.

Going back to my earlier example, the following function declaration

__attribute__((overloadable)) void foo(__global float *a)

and

__attribute__((overloadable)) void foo(__local float *a)

would be mapped to the same mangled name because of the mapping of local and
global address space to the single physical address space provided by X86, so
the high level information is lost.

Thus, the issue is not that the address spaces are mapped incorrectly, but
rather that the mangling should not be based on the map, but on the original
address space specifiers.

*Ignoring the semantics of the IR addrspace modifier, which is TARGET-specific*,
we could indeed define a custom mapping for X86 which artificially preserved
multiple address spaces, and the above example would work correctly. But the
proposed X86 mapping would actually be a target-independent but language
dependent mapping useful at the higher levels of the compiler (before code
generation).

The fact that using untranslated address spaces is aesthetically unpleasant can
be easily solved with a specialization of the mangling, one for OpenC and one
for CUDA. This solution would be semantically correct and would maintain the
mangling meaningful (see attachment).

Regards,

-Michele
-------------- next part --------------
A non-text attachment was scrubbed...
Name: mangling.patch
Type: text/x-patch
Size: 1357 bytes
Desc: not available
URL: <http://lists.cs.uiuc.edu/pipermail/cfe-commits/attachments/20130723/fc2bb6da/attachment-0001.bin>
Tanya Lattner
2013-07-23 20:14:32 UTC
Permalink
Post by Michele Scandale
Hi Tanya,
Post by Tanya Lattner
I understand what you are saying, but I'm arguing that the Target needs to define the address space map in order for this to mean anything at all. If its not defined, then what does global or local mean? What happens when this is lowered to LLVM IR?
Why not define an address space map for X86?
The trivial address map is the correct one for X86, since it does not have
separate address spaces.
The problem is that at the language level, address space specifiers carry
additional information over that related to the memory allocation. In
particular, the specifiers cause function prototypes to refer to different
functions even if they differ only for an address space specifier.
Going back to my earlier example, the following function declaration
__attribute__((overloadable)) void foo(__global float *a)
and
__attribute__((overloadable)) void foo(__local float *a)
would be mapped to the same mangled name because of the mapping of local and
global address space to the single physical address space provided by X86, so
the high level information is lost.
Thus, the issue is not that the address spaces are mapped incorrectly, but
rather that the mangling should not be based on the map, but on the original
address space specifiers.
*Ignoring the semantics of the IR addrspace modifier, which is TARGET-specific*,
we could indeed define a custom mapping for X86 which artificially preserved
multiple address spaces, and the above example would work correctly. But the
proposed X86 mapping would actually be a target-independent but language
dependent mapping useful at the higher levels of the compiler (before code
generation).
The fact that using untranslated address spaces is aesthetically unpleasant can
be easily solved with a specialization of the mangling, one for OpenC and one
for CUDA. This solution would be semantically correct and would maintain the
mangling meaningful (see attachment).
Ok, this approach will work for me. However, I would prefer to not have CL prefix and keep it as AS#.

Thanks,
Tanya
Post by Michele Scandale
Regards,
-Michele
<mangling.patch>
Michele Scandale
2013-07-23 21:03:21 UTC
Permalink
Post by Tanya Lattner
Ok, this approach will work for me. However, I would prefer to not have CL prefix and keep it as AS#.
This means that you want to revert the commit, removing the usage of the
target-address space map, right?

If so it's still fine :-)

Thanks again.

-Michele
Tanya Lattner
2013-07-23 21:36:55 UTC
Permalink
Post by Tanya Lattner
Ok, this approach will work for me. However, I would prefer to not have CL prefix and keep it as AS#.
This means that you want to revert the commit, removing the usage of the target-address space map, right?
No. I wanted the specialized mangling that you are proposing in your patch but don't use "CL" as the prefix. Use "AS" for CL.

However, I actually should have looked at this closer as it actually doesn't map to what I want it to. I want it to be the following:

1, // opencl_global
3, // opencl_local
2, // opencl_constant

and when there is no address space then it maps to nothing.

So, I don't think your patch is going to work unless the order is changed in the enum. Because this is not clearly defined in the spec and is implementation specific and TARGET specific.. then changing that enum is probably not going to be the right approach either.

So, I'm going back to my original statement to keep it to be Target specific. For your library, are these functions actually implemented differently? Wouldn't they be exactly the same when there is no address space? In our implementation we have an address space map defined for X86 and then the names get mangled "correctly" for all targets. But, all the functionality is the same since the address spaces don't impact codegen for X86.

-Tanya
If so it's still fine :-)
Thanks again.
-Michele
-------------- next part --------------
An HTML attachment was scrubbed...
URL: <http://lists.cs.uiuc.edu/pipermail/cfe-commits/attachments/20130723/2ca7dff7/attachment-0001.html>
Michele Scandale
2013-07-24 14:49:45 UTC
Permalink
Post by Tanya Lattner
However, I actually should have looked at this closer as it actually doesn't map
1, // opencl_global
3, // opencl_local
2, // opencl_constant
and when there is no address space then it maps to nothing.
So, I don't think your patch is going to work unless the order is changed in the
enum. Because this is not clearly defined in the spec and is implementation
specific and TARGET specific.. then changing that enum is probably not going to
be the right approach either.
I see the point but still we need to preserve the source language difference
with the mangling.
Post by Tanya Lattner
So, I'm going back to my original statement to keep it to be Target specific.
For your library, are these functions actually implemented differently? Wouldn't
they be exactly the same when there is no address space? In our implementation
we have an address space map defined for X86 and then the names get mangled
"correctly" for all targets. But, all the functionality is the same since the
address spaces don't impact codegen for X86.
I'd argue that the fact that address spaces do not impact codegen for X86 is
merely incidental: the issue goes beyond X86 and impacts all targets --
including future ones.
By choosing to use a fake address space map (instead of the one fitting the
target description), we introduce in the IR a potential for breaking future
implementations, even though the current X86 target is not affected. Moreover,
by introducing the fake address map we would violate the semantics of the LLVM
IR addrspace modifier.

Even binary compatibility with libraries already distributed can be easily achieved.

This patch (see attachment) aims at preserving the mangling that were generated
by using the target address space map for those targets that override it, while
introducing in the mangling the distinction of opencl/cuda address spaces for
those targets that do not have a non trivial target address space map.

By the way, looking beyond the scope of the current issue of mangling, it would
be IMO interesting to start a public discussion on the mailing list about a way
to represent logical address space information different from target-specific
address space (the case for OpenCL and CUDA) in order to allow the
implementation of custom language specific analysis and/or optimization.
As a temporary solution, if one needed the logical address space information in
the IR too for specific purpose (like OpenCL specific optimization), can still
override the address space map of the target.

Thanks again.

Regards,

-Michele

-------------- next part --------------
A non-text attachment was scrubbed...
Name: mangling-rev4.patch
Type: text/x-patch
Size: 9471 bytes
Desc: not available
URL: <http://lists.cs.uiuc.edu/pipermail/cfe-commits/attachments/20130724/9678d9c6/attachment.bin>
Pekka Jääskeläinen
2013-07-24 15:36:10 UTC
Permalink
Post by Michele Scandale
As a temporary solution, if one needed the logical address space information in
the IR too for specific purpose (like OpenCL specific optimization), can still
override the address space map of the target.
FWIW, we do it this way in pocl now. That is, we force known address space ids
to the OpenCL address spaces via the fake address space map feature so the
OpenCL C kernel compiler passes can refer to non-target-specific address spaces
during the work group generation.

Then, as a final pass, these IDs can be mapped back to the target's preferred
ones. The final pass is not needed if the target can automatically remap them
in the backend, like is the case with single AS CPUs where everything will be
finally mapped to AS0.

BR,
--
--Pekka,
http://pocl.sf.net
Michele Scandale
2013-07-24 16:03:57 UTC
Permalink
Post by Pekka Jääskeläinen
FWIW, we do it this way in pocl now. That is, we force known address space ids
to the OpenCL address spaces via the fake address space map feature so the
OpenCL C kernel compiler passes can refer to non-target-specific address spaces
during the work group generation.
I know there is that option, but it's there only for testing purpose. It can be
a temporary workaround but not a long term solution. I think it's useful for
everybody to find and apply solutions that keep clean both design and code.

As said it would be very interesting to discuss in order to found a clean way to
represent logical address spaces independently from target address spaces so
that in the middle end this information can be exploited for language specific
optimizations.

Maybe a simple extension of the 'addrspace' modifier so that it can represent
either logical or physical address spaces will be fine, but we need to discuss
the implementation details and understand if it's a solution that fits the goals
in term of usability, maintainability and compatibility.

Best Regards,

-Michele
Michele Scandale
2013-08-01 15:33:40 UTC
Permalink
Up.
Post by Michele Scandale
Post by Tanya Lattner
However, I actually should have looked at this closer as it actually doesn't map
1, // opencl_global
3, // opencl_local
2, // opencl_constant
and when there is no address space then it maps to nothing.
So, I don't think your patch is going to work unless the order is changed in the
enum. Because this is not clearly defined in the spec and is implementation
specific and TARGET specific.. then changing that enum is probably not going to
be the right approach either.
I see the point but still we need to preserve the source language difference
with the mangling.
Post by Tanya Lattner
So, I'm going back to my original statement to keep it to be Target specific.
For your library, are these functions actually implemented differently? Wouldn't
they be exactly the same when there is no address space? In our implementation
we have an address space map defined for X86 and then the names get mangled
"correctly" for all targets. But, all the functionality is the same since the
address spaces don't impact codegen for X86.
I'd argue that the fact that address spaces do not impact codegen for X86 is
merely incidental: the issue goes beyond X86 and impacts all targets --
including future ones.
By choosing to use a fake address space map (instead of the one fitting the
target description), we introduce in the IR a potential for breaking future
implementations, even though the current X86 target is not affected. Moreover,
by introducing the fake address map we would violate the semantics of the LLVM
IR addrspace modifier.
Even binary compatibility with libraries already distributed can be easily achieved.
This patch (see attachment) aims at preserving the mangling that were generated
by using the target address space map for those targets that override it, while
introducing in the mangling the distinction of opencl/cuda address spaces for
those targets that do not have a non trivial target address space map.
By the way, looking beyond the scope of the current issue of mangling, it would
be IMO interesting to start a public discussion on the mailing list about a way
to represent logical address space information different from target-specific
address space (the case for OpenCL and CUDA) in order to allow the
implementation of custom language specific analysis and/or optimization.
As a temporary solution, if one needed the logical address space information in
the IR too for specific purpose (like OpenCL specific optimization), can still
override the address space map of the target.
Thanks again.
Regards,
-Michele
Tanya Lattner
2013-08-01 17:45:04 UTC
Permalink
Post by Tanya Lattner
However, I actually should have looked at this closer as it actually doesn't map
1, // opencl_global
3, // opencl_local
2, // opencl_constant
and when there is no address space then it maps to nothing.
So, I don't think your patch is going to work unless the order is changed in the
enum. Because this is not clearly defined in the spec and is implementation
specific and TARGET specific.. then changing that enum is probably not going to
be the right approach either.
I see the point but still we need to preserve the source language difference with the mangling.
Post by Tanya Lattner
So, I'm going back to my original statement to keep it to be Target specific.
For your library, are these functions actually implemented differently? Wouldn't
they be exactly the same when there is no address space? In our implementation
we have an address space map defined for X86 and then the names get mangled
"correctly" for all targets. But, all the functionality is the same since the
address spaces don't impact codegen for X86.
I'd argue that the fact that address spaces do not impact codegen for X86 is merely incidental: the issue goes beyond X86 and impacts all targets -- including future ones.
By choosing to use a fake address space map (instead of the one fitting the target description), we introduce in the IR a potential for breaking future implementations, even though the current X86 target is not affected. Moreover, by introducing the fake address map we would violate the semantics of the LLVM IR addrspace modifier.
Even binary compatibility with libraries already distributed can be easily achieved.
This patch (see attachment) aims at preserving the mangling that were generated by using the target address space map for those targets that override it, while introducing in the mangling the distinction of opencl/cuda address spaces for those targets that do not have a non trivial target address space map.
By the way, looking beyond the scope of the current issue of mangling, it would be IMO interesting to start a public discussion on the mailing list about a way to represent logical address space information different from target-specific address space (the case for OpenCL and CUDA) in order to allow the implementation of custom language specific analysis and/or optimization.
As a temporary solution, if one needed the logical address space information in the IR too for specific purpose (like OpenCL specific optimization), can still override the address space map of the target.
This discussion probably should be moved to the cfe-dev mailing list. I think that its better for Clang to have one consistent way of mangling address spaces regardless of language. I?ve been thinking about this more and another problem I have with not using the Target address space map for mangling is that when you have this:

void __attribute__((__overloadable__)) foo(global int *x)

You get this:
define void @_Z3fooPU10AS16776960b(i8* %x)

The address space on the argument is gone (or zero), but yet you have it in the mangled name. So its not consistent.

I can agree that its right to mangle the names differently from the language perspective, but what you mangle them to is really target specific. If you want to remove this notion from Clang, then maybe all target specific address space maps should go away and a default one for all is used. Then each LLVM backend can interpret it as they wish.

It would be great if some Code owners could weigh in here.

-Tanya

P.S. Please respond to this email address going forward as I will not be using my apple email anymore.


-------------- next part --------------
An HTML attachment was scrubbed...
URL: <http://lists.cs.uiuc.edu/pipermail/cfe-commits/attachments/20130801/7b96bd21/attachment-0001.html>
Michele Scandale
2013-08-01 18:12:51 UTC
Permalink
Post by Tanya Lattner
This discussion probably should be moved to the cfe-dev mailing list. I think
that its better for Clang to have one consistent way of mangling address spaces
regardless of language. I?ve been thinking about this more and another problem
I have with not using the Target address space map for mangling is that when you
void __attribute__((__overloadable__)) foo(global int *x)
The address space on the argument is gone (or zero), but yet you have it in the
mangled name. So its not consistent.
Why it's not consistent? In the IR at this time, the addrspace modifier
represent TARGET/physical address spaces, not LANGUAGE/logical address spaces.
In fact as said it would be useful to represent also logical address spaces in
the IR, but this is another problem not strictly related to the mangler.
Post by Tanya Lattner
I can agree that its right to mangle the names differently from the language
perspective, but what you mangle them to is really target specific. If you want
The last patch I proposed allow to have a translation map for opencl/cuda
address spaces that can be overwritten by targets, so that the mangling can be
tuned in order to have binary compatibility with other/pre-existing libraries.
A default map is used in order to have that all target have a consistent
mangling (no name collisions) in order to fix the problem strictly related to
the mangler.
Post by Tanya Lattner
I can agree that its right to mangle the names differently from the language
perspective, but what you mangle them to is really target specific. If you want
to remove this notion from Clang, then maybe all target specific address space
maps should go away and a default one for all is used. Then each LLVM backend
can interpret it as they wish.
As you said mangling is implementation specific, so this allow to have various
solutions: the fact that the mangling is TARGET+LANGUAGE specific is perfectly
fine. Because of this I proposed the patch that introduce the map to translate
address spaces from the internal representation to the logical representation
(in principle different from the physical representation): with this solution
the mangling reflect its intrinsic dependency from language concepts, like
logical address spaces having the target dependency that allow each target to
change this mapping, e.g. to have binary compatibility.

Again, the fact that logical address spaces are not also inside the IR as
property of pointer types is missing feature IMO. Fixing the mangler would then
allow to start the fixing of this other problem.

Thanks.

Best Regards,

-Michele
Mon Ping Wang
2013-08-02 05:35:50 UTC
Permalink
Post by Tanya Lattner
This discussion probably should be moved to the cfe-dev mailing list. I think
that its better for Clang to have one consistent way of mangling address spaces
regardless of language. I?ve been thinking about this more and another problem
I have with not using the Target address space map for mangling is that when you
void __attribute__((__overloadable__)) foo(global int *x)
The address space on the argument is gone (or zero), but yet you have it in the
mangled name. So its not consistent.
Why it's not consistent? In the IR at this time, the addrspace modifier represent TARGET/physical address spaces, not LANGUAGE/logical address spaces.
In fact as said it would be useful to represent also logical address spaces in the IR, but this is another problem not strictly related to the mangler.
Sorry of being late to this conversation. It doesn?t look consistent me. Address space numbers are not language constructs. The language constructs are global and local. Coming out of clang, I think it is more natural for the AS mangling and the type to match.
In C++, clang will generate different names for structures which can be identical and uses those names consistently to mangle the function, e.g.,
%struct.foo = type { i32, i32 }
define void @_Z4testR3foo(%struct.foo* %foo)

I view the address spaces coming out of clang represent how the target represent memory is a logical. How a particular llvm maps them to physical memory is target dependent. A backend may map them all the address spaces to the same physical memory or to different ones. Due to this, I don?t think it make sense to distinguish between the two in clang for a particular target.

Best regards,
? Mon Ping
Post by Tanya Lattner
I can agree that its right to mangle the names differently from the language
perspective, but what you mangle them to is really target specific. If you want
The last patch I proposed allow to have a translation map for opencl/cuda address spaces that can be overwritten by targets, so that the mangling can be tuned in order to have binary compatibility with other/pre-existing libraries.
A default map is used in order to have that all target have a consistent mangling (no name collisions) in order to fix the problem strictly related to the mangler.
Post by Tanya Lattner
I can agree that its right to mangle the names differently from the language
perspective, but what you mangle them to is really target specific. If you want
to remove this notion from Clang, then maybe all target specific address space
maps should go away and a default one for all is used. Then each LLVM backend
can interpret it as they wish.
As you said mangling is implementation specific, so this allow to have various solutions: the fact that the mangling is TARGET+LANGUAGE specific is perfectly fine. Because of this I proposed the patch that introduce the map to translate address spaces from the internal representation to the logical representation (in principle different from the physical representation): with this solution the mangling reflect its intrinsic dependency from language concepts, like logical address spaces having the target dependency that allow each target to change this mapping, e.g. to have binary compatibility.
Again, the fact that logical address spaces are not also inside the IR as property of pointer types is missing feature IMO. Fixing the mangler would then allow to start the fixing of this other problem.
Thanks.
Best Regards,
-Michele
-------------- next part --------------
An HTML attachment was scrubbed...
URL: <http://lists.cs.uiuc.edu/pipermail/cfe-commits/attachments/20130801/4b3dc7f8/attachment-0001.html>
Michele Scandale
2013-08-02 21:57:09 UTC
Permalink
Hello Mon Ping,

I apologize for the mail length, but I hope to explain as clear as I can the
points I think need to be discussed.
Post by Mon Ping Wang
Sorry of being late to this conversation. It doesn?t look consistent me. Address
space numbers are not language constructs. The language constructs are global
and local. Coming out of clang, I think it is more natural for the AS mangling
and the type to match. In C++, clang will generate different names for
structures which can be identical and uses those names consistently to mangle
the function, e.g.,
%struct.foo = type { i32, i32 }
I view the address spaces coming out of clang represent how the target represent
memory is a logical. How a particular llvm maps them to physical memory is
target dependent. A backend may map them all the address spaces to the same
physical memory or to different ones. Due to this, I don?t think it make sense
to distinguish between the two in clang for a particular target.
I agree, the fact that opencl address spaces are handled like other address
spaces is a technical aspect. To have a common way I don't see a strict
limitation in how address spaces are mangled (they can be numbers decided as
convention in clang, or defined by targets, or whatever), but still mangling
should preserve the differences that are present in the source language.

I want to remark IMO an important aspect:
"Pointer types may have an optional address space attribute defining the
numbered address space where the pointed-to object resides. The default address
space is number zero. The semantics of non-zero address spaces are
target-specific." (http://llvm.org/docs/LangRef.html#pointer-type)

From this description I understand that address spaces in the IR are physical
address spaces. Because of this I consider wrong to use this property as is to
represent inside the IR logical address space. Doing that would imply that each
backend should be aware of language specific mapping: currently this is not the
case and IMO it's a bad idea to have this.

But a derived information from the source language is still useful to perform
optimization, both in the IR and later in the backend: the logical distinction
of address spaces is still useful and IMO shoul be represented in the IR. Have
both logical and physical address spaces information (it's not important to know
is "AS1" means global or local, it's enough to know that 1 is differnt from 2)
would be useful to have a better alias analysis also for those targets that
physically have one unique address space. I consider that this can be solved
independently from the mangling problem.

The answer to both question, I suggested to introduce another map in order to
preserve the distinction between address spaces also for those targets that do
not have physical distinct address spaces, like X86, and through this solve the
problem related to the mangler.

As previously discussed, this is not the only viable solution, the mapping of
logical address spaces to physical address spaces can be delayed till
instruction selection: this would allow the frontend to lower this information
in a target independent manner demanding a late IR pass the mapping task (this
task would be language/target dependent, so basically who builds the pass
pipeline must schedule this language dependent task that requires target
informations). Still here may be useful to preserve the logical information of
address spaces.
This kind of solution is feasible, but simply it does not seem the way chosen in
clang to solve the problem.

My proposal was the one with the minimal impact on the codebase trying to
maintain a desirable flexibility in order to build opencl toolchain compatible
with the past.

Could you explain to me what you are proposing? How the mangler should be fixed?
How address spaces are lowered in the IR? This lowering is target dependent or
not? The mangling is also target dependent?

Thanks in advance.

Best regards,
-Michele
Mon Ping Wang
2013-08-08 19:26:43 UTC
Permalink
HI Michele,
Post by Michele Scandale
Hello Mon Ping,
I apologize for the mail length, but I hope to explain as clear as I can the points I think need to be discussed.
Post by Mon Ping Wang
Sorry of being late to this conversation. It doesn?t look consistent me. Address
space numbers are not language constructs. The language constructs are global
and local. Coming out of clang, I think it is more natural for the AS mangling
and the type to match. In C++, clang will generate different names for
structures which can be identical and uses those names consistently to mangle
the function, e.g.,
%struct.foo = type { i32, i32 }
I view the address spaces coming out of clang represent how the target represent
memory is a logical. How a particular llvm maps them to physical memory is
target dependent. A backend may map them all the address spaces to the same
physical memory or to different ones. Due to this, I don?t think it make sense
to distinguish between the two in clang for a particular target.
I agree, the fact that opencl address spaces are handled like other address spaces is a technical aspect. To have a common way I don't see a strict limitation in how address spaces are mangled (they can be numbers decided as convention in clang, or defined by targets, or whatever), but still mangling should preserve the differences that are present in the source language.
"Pointer types may have an optional address space attribute defining the numbered address space where the pointed-to object resides. The default address space is number zero. The semantics of non-zero address spaces are target-specific." (http://llvm.org/docs/LangRef.html#pointer-type)
From this description I understand that address spaces in the IR are physical address spaces. Because of this I consider wrong to use this property as is to represent inside the IR logical address space. Doing that would imply that each backend should be aware of language specific mapping: currently this is not the case and IMO it's a bad idea to have this.
But a derived information from the source language is still useful to perform optimization, both in the IR and later in the backend: the logical distinction of address spaces is still useful and IMO shoul be represented in the IR. Have both logical and physical address spaces information (it's not important to know is "AS1" means global or local, it's enough to know that 1 is differnt from 2) would be useful to have a better alias analysis also for those targets that physically have one unique address space. I consider that this can be solved independently from the mangling problem.
The answer to both question, I suggested to introduce another map in order to preserve the distinction between address spaces also for those targets that do not have physical distinct address spaces, like X86, and through this solve the problem related to the mangler.
As previously discussed, this is not the only viable solution, the mapping of logical address spaces to physical address spaces can be delayed till instruction selection: this would allow the frontend to lower this information in a target independent manner demanding a late IR pass the mapping task (this task would be language/target dependent, so basically who builds the pass pipeline must schedule this language dependent task that requires target informations). Still here may be useful to preserve the logical information of address spaces.
This kind of solution is feasible, but simply it does not seem the way chosen in clang to solve the problem.
My proposal was the one with the minimal impact on the codebase trying to maintain a desirable flexibility in order to build opencl toolchain compatible with the past.
Could you explain to me what you are proposing? How the mangler should be fixed? How address spaces are lowered in the IR? This lowering is target dependent or not? The mangling is also target dependent?
IMO, the description only indicates that an address space is completely target dependent. For the current x86 target, address spaces > 255 are used for a non-standard address for the stack protector while every other address space overlaps and maps to the same region in memory. A target can defined it differently or make some address spaces illegal but it is up to the target.

When generating code for a particular target, clang need to decide on how to map the global, local, etc.. for a specific target. Currently, for X86, it decides to use different address space to distinguish for overloading knowing that in the target, the address spaces will physically overlap. This keeps the two sides consistent when mangling based on the LLVM IR address space and keeps the overloaded functions to be distinguished for this particular target. This choice, as you noted, is to make the mapping target dependent. If a target wants to map everything to the same address space and wants to overloading of their functions because there is no distinction, it can make that choice at this level.

My objection to the logical map is that by introducing the CL address names to an address space numbering, it looks very target dependent and if the logical address space vs LLVM IR address space doesn?t match, it looks inconsistent. In that case, I think we should do what we are currently doing. Instead of a logical map, if we want to preserve the language constructs in a target independent manner, we should use the language construct names in the overloading as that is language dependent and independent of AS numbers which are LLVM IR concepts; which I believe Eli indicated as well. If we want to preserve compatibility for some target, we can make it target dependent if they want to map use current address space mapping today or use the language mapping. I don?t know how Eli or the other code owners feel about having that compatibly mode which will be useful for people want to preserve the old behavior. Opinions?

Thanks,
? Mon Ping
Post by Michele Scandale
Thanks in advance.
Best regards,
-Michele
_______________________________________________
cfe-commits mailing list
cfe-commits at cs.uiuc.edu
http://lists.cs.uiuc.edu/mailman/listinfo/cfe-commits
Michele Scandale
2013-08-08 20:22:16 UTC
Permalink
Hello Mon Ping,
Post by Mon Ping Wang
IMO, the description only indicates that an address space is completely target dependent. For the current x86 target, address spaces > 255 are used for a non-standard address for the stack protector while every other address space overlaps and maps to the same region in memory. A target can defined it differently or make some address spaces illegal but it is up to the target.
When generating code for a particular target, clang need to decide on how to map the global, local, etc.. for a specific target. Currently, for X86, it decides to use different address space to distinguish for overloading knowing that in the target, the address spaces will physically overlap. This keeps the two sides consistent when mangling based on the LLVM IR address space and keeps the overloaded functions to be distinguished for this particular target. This choice, as you noted, is to make the mapping target dependent. If a target wants to map everything to the same address space and wants to overloading of their functions because there is no distinction, it can make that choice at this level.
So why the addrspace map for X86 is still the trivial one IF the
assumption in the backend is that whatever number I choose less than 255
is the same as 0? Maybe for X86 defining a non trivial map is a correct
fix, but it's not true in general!

What if an hypothetical backend would enforce that there exists ONLY
address space zero? Why I should not be able to produce a correct mangle
for opencl overloaded function that refers to different logical address
spaces?
Post by Mon Ping Wang
My objection to the logical map is that by introducing the CL address names to an address space numbering, it looks very target dependent and if the logical address space vs LLVM IR address space doesn?t match, it looks inconsistent. In that case, I think we should do what we are currently doing. Instead of a logical map, if we want to preserve the language constructs in a target independent manner, we should use the language construct names in the overloading as that is language dependent and independent of AS numbers which are LLVM IR concepts; which I believe Eli indicated as well. If we want to preserve compatibility for some target, we can make it target dependent if they want to map use current address space mapping today or use the language mapping. I don?t know how Eli or the other code owners feel about having that compatibly mode which will be useful for people want to preserve the old behavior. Opinions?
The idea of having something target independent seems considered bad in
the previous messages. IMO the usage of numbers can be unpleasant,
implementation dependent, but I haven't seen a standardized mangling for
OpenCL C.

My point is that *every* target the mangler should produce different
names even if the address space translation map is the trivial one.
How the address space information is propagated in the IR and the
mangling IMO are orthogonal problem: so the inconsistency you underline
conceptually cannot exist by definition.

What I noticed is that the mangler now produces wrong names respect to
its purpose (X86 is only the test case).

Thanks for your reply.

Regards,
-Michele
Michele Scandale
2013-08-20 22:31:02 UTC
Permalink
Up.

Regards,
-Michele
Post by Michele Scandale
The idea of having something target independent seems considered bad in
the previous messages. IMO the usage of numbers can be unpleasant,
implementation dependent, but I haven't seen a standardized mangling for
OpenCL C.
My point is that *every* target the mangler should produce different
names even if the address space translation map is the trivial one.
How the address space information is propagated in the IR and the
mangling IMO are orthogonal problem: so the inconsistency you underline
conceptually cannot exist by definition.
What I noticed is that the mangler now produces wrong names respect to
its purpose (X86 is only the test case).
Mon Ping Wang
2013-08-24 02:21:23 UTC
Permalink
Hi Michele,

Sorry for the delay response.
Post by Michele Scandale
Hello Mon Ping,
Post by Mon Ping Wang
IMO, the description only indicates that an address space is completely target dependent. For the current x86 target, address spaces > 255 are used for a non-standard address for the stack protector while every other address space overlaps and maps to the same region in memory. A target can defined it differently or make some address spaces illegal but it is up to the target.
When generating code for a particular target, clang need to decide on how to map the global, local, etc.. for a specific target. Currently, for X86, it decides to use different address space to distinguish for overloading knowing that in the target, the address spaces will physically overlap. This keeps the two sides consistent when mangling based on the LLVM IR address space and keeps the overloaded functions to be distinguished for this particular target. This choice, as you noted, is to make the mapping target dependent. If a target wants to map everything to the same address space and wants to overloading of their functions because there is no distinction, it can make that choice at this level.
So why the addrspace map for X86 is still the trivial one IF the assumption in the backend is that whatever number I choose less than 255 is the same as 0? Maybe for X86 defining a non trivial map is a correct fix, but it's not true in general!
What if an hypothetical backend would enforce that there exists ONLY address space zero? Why I should not be able to produce a correct mangle for opencl overloaded function that refers to different logical address spaces?
Yes, you are right that in general it is not true. As we both agree, LLVM address spaces are completely target dependent. The question is why would someone want to produce different overloaded functions when the llvm backend address space only supports one. It can?t be for code generation since the code will be the same. A backend may want them to all be mangled to the be the same since they would collapse the number of CL builtin functions they would need to support.

There are cases when clients may want different LLVM IR address spaces and the mangling. One case is if someone uses the address space for alias analysis. Another case is that a platform has a set of devices, some with physical address spaces, and wants to keep the mangled name consistent for the platform. Both of these cases are target dependent on why they want to do so.
Post by Michele Scandale
Post by Mon Ping Wang
My objection to the logical map is that by introducing the CL address names to an address space numbering, it looks very target dependent and if the logical address space vs LLVM IR address space doesn?t match, it looks inconsistent. In that case, I think we should do what we are currently doing. Instead of a logical map, if we want to preserve the language constructs in a target independent manner, we should use the language construct names in the overloading as that is language dependent and independent of AS numbers which are LLVM IR concepts; which I believe Eli indicated as well. If we want to preserve compatibility for some target, we can make it target dependent if they want to map use current address space mapping today or use the language mapping. I don?t know how Eli or the other code owners feel about having that compatibly mode which will be useful for people want to preserve the old behavior. Opinions?
The idea of having something target independent seems considered bad in the previous messages. IMO the usage of numbers can be unpleasant, implementation dependent, but I haven't seen a standardized mangling for OpenCL C.
My point is that *every* target the mangler should produce different names even if the address space translation map is the trivial one.
I?m not convinced on this point. Can you please explain the use case that you want to support again?
Post by Michele Scandale
How the address space information is propagated in the IR and the mangling IMO are orthogonal problem: so the inconsistency you underline conceptually cannot exist by definition.
What I noticed is that the mangler now produces wrong names respect to its purpose (X86 is only the test case).
If we have both a logical map and a llvm address space maps, I think it is confusing that the mangled name address space differs from the physical llvm map. It is like having a type name for managing that has no relationship with the type name in the LLVM IR or the language it is coming from. If we need to support the CL address spaces mangling to be different from the LLVM IR address space, I think it would be better to be target independent and force mangling to be based on the language (global, local, etc..) , which it sounds like you were not opposed of. As noted above, there are cases where we want them to match.


? Mon Ping
Post by Michele Scandale
Thanks for your reply.
Regards,
-Michele
Michele Scandale
2013-08-24 14:02:23 UTC
Permalink
Post by Mon Ping Wang
Hi Michele,
Sorry for the delay response.
Post by Michele Scandale
Hello Mon Ping,
Post by Mon Ping Wang
IMO, the description only indicates that an address space is completely target dependent. For the current x86 target, address spaces > 255 are used for a non-standard address for the stack protector while every other address space overlaps and maps to the same region in memory. A target can defined it differently or make some address spaces illegal but it is up to the target.
When generating code for a particular target, clang need to decide on how to map the global, local, etc.. for a specific target. Currently, for X86, it decides to use different address space to distinguish for overloading knowing that in the target, the address spaces will physically overlap. This keeps the two sides consistent when mangling based on the LLVM IR address space and keeps the overloaded functions to be distinguished for this particular target. This choice, as you noted, is to make the mapping target dependent. If a target wants to map everything to the same address space and wants to overloading of their functions because there is no distinction, it can make that choice at this level.
So why the addrspace map for X86 is still the trivial one IF the assumption in the backend is that whatever number I choose less than 255 is the same as 0? Maybe for X86 defining a non trivial map is a correct fix, but it's not true in general!
What if an hypothetical backend would enforce that there exists ONLY address space zero? Why I should not be able to produce a correct mangle for opencl overloaded function that refers to different logical address spaces?
Yes, you are right that in general it is not true. As we both agree, LLVM address spaces are completely target dependent. The question is why would someone want to produce different overloaded functions when the llvm backend address space only supports one. It can?t be for code generation since the code will be the same. A backend may want them to all be mangled to the be the same since they would collapse the number of CL builtin functions they would need to support.
There are cases when clients may want different LLVM IR address spaces and the mangling. One case is if someone uses the address space for alias analysis. Another case is that a platform has a set of devices, some with physical address spaces, and wants to keep the mangled name consistent for the platform. Both of these cases are target dependent on why they want to do so.
Post by Michele Scandale
Post by Mon Ping Wang
My objection to the logical map is that by introducing the CL address names to an address space numbering, it looks very target dependent and if the logical address space vs LLVM IR address space doesn?t match, it looks inconsistent. In that case, I think we should do what we are currently doing. Instead of a logical map, if we want to preserve the language constructs in a target independent manner, we should use the language construct names in the overloading as that is language dependent and independent of AS numbers which are LLVM IR concepts; which I believe Eli indicated as well. If we want to preserve compatibility for some target, we can make it target dependent if they want to map use current address space mapping today or use the language mapping. I don?t know how Eli or the other code owners feel about having that compatibly mode which will be useful for people want to preserve the old behavior. Opinions?
The idea of having something target independent seems considered bad in the previous messages. IMO the usage of numbers can be unpleasant, implementation dependent, but I haven't seen a standardized mangling for OpenCL C.
My point is that *every* target the mangler should produce different names even if the address space translation map is the trivial one.
I?m not convinced on this point. Can you please explain the use case that you want to support again?
Post by Michele Scandale
How the address space information is propagated in the IR and the mangling IMO are orthogonal problem: so the inconsistency you underline conceptually cannot exist by definition.
What I noticed is that the mangler now produces wrong names respect to its purpose (X86 is only the test case).
If we have both a logical map and a llvm address space maps, I think it is confusing that the mangled name address space differs from the physical llvm map. It is like having a type name for managing that has no relationship with the type name in the LLVM IR or the language it is coming from. If we need to support the CL address spaces mangling to be different from the LLVM IR address space, I think it would be better to be target independent and force mangling to be based on the language (global, local, etc..) , which it sounds like you were not opposed of. As noted above, there are cases where we want them to match.
My use case is OpenCL. In this language the abstraction of address spaces is
explicit, so whatever is the way I implement this abstraction I would need to
have different names for overloaded functions that differs only for address
space qualifiers for pointers, as the mangling is a technique to avoid name
collisions for function with same name but different signatures.

The mangling is just a frontend matter to solve names collision and preserve
source language aspects. For OpenCL I should not care if physical address spaces
exists. Who implements these OpenCL functions (these functions are the OpenCL
builtins) for a target that have already a target library with similar functions
would need only to call them.

The pure solution would be the one proposed by Eli: I don't have any objection
to this solution.
The mangler now has a bug, so it must be fixed. The pure solution implicitly
breaks the binary compatibility. If we do not have problem with this (so we
consider a matter for the users to solve the problem, e.g. with a forced update
of libraries) the right patch is to have a target independent mangling for OpenCL.

Still we would have problems if we consider SPIR: in its specification there is
a fixed mangling scheme (that it's the one produced by the current mangler). In
this case we have two choice: we change the SPIR mangling or we allow targets to
override the target independent mangling for OpenCL with the one based on the
TargetAddrSpaceMap.

*Based on all this would see the mangling proposed by Eli the default except for
targets that explicitly requires a mangling scheme based on the target address
spaces map (e.g. the SPIR target).*


Thanks.

Regards,

-Michele


----- A little digression ------

How OpenCL address spaces are lowered in the IR is another problem orthogonal
problem.

For targets like PTX or R600 we have real distinct address spaces so, also here
is fine to use those.

On X86 we can use target address spaces in the range [0-255] as the backend has
the assumption that they are all equivalent to address space 0.
But this is not the general case for CPUs targets. So on a generic CPU is
correct to map all the OpenCL address spaces to the default target address space
(0).

For alias analysis purpose, I started a discussion (quite huge) in the LLVMDev
mailing list (
http://lists.cs.uiuc.edu/pipermail/llvmdev/2013-August/064620.html ) to find a
way to represent logical address spaces in the IR so to be able to distinguish
between two memory location that are physically in the same address space but
logically in two logical different (maybe also disjoint) address spaces.

A reasonable solution is here (
http://lists.cs.uiuc.edu/pipermail/llvmdev/2013-August/064807.html ): use TBAA
similar metadata to describe relationship between source level address spaces (a
tree structure to represent inclusion relationship and "constant" property) and
add them to load store instructions (as done for the TBAA). From this is
possible to introduce in LLVM a new AliasAnalysis for address spaces that use
these informations to decide aliasing when the physical address space is the
same, and when target address spaces are used the query is done to the target
that knows these address spaces (so it can answer true if two physical address
spaces are disjoint or not.)

In this way in Clang, the target address space map keep the current semantic,
address space metadata should be emitted in the case of OpenCL and attached to
load-store, etc, instructions.
David Tweed
2013-08-27 09:02:51 UTC
Permalink
| The pure solution would be the one proposed by Eli: I don't have any objection
| to this solution.
| The mangler now has a bug, so it must be fixed. The pure solution implicitly
| breaks the binary compatibility. If we do not have problem with this (so we
| consider a matter for the users to solve the problem, e.g. with a forced update
| of libraries) the right patch is to have a target independent mangling for OpenCL.

I think there's another reason for desiring a target independent mangling: a system may contain several OpenCL devices and the actual implementation of address spaces (in particular whether they're "front-end annotations only" or actually denote physically different regions of memory) may depend on the OpenCL device (with its associated backend). (In the conventional OpenCL usage it may not matter since one could postpone the address space resolution to later in the process; once you've got to process already produced SPIR I think it does.)

Cheers,
Dave

-- IMPORTANT NOTICE: The contents of this email and any attachments are confidential and may also be privileged. If you are not the intended recipient, please notify the sender immediately and do not disclose the contents to any other person, use it for any purpose, or store or copy the information in any medium. Thank you.

ARM Limited, Registered office 110 Fulbourn Road, Cambridge CB1 9NJ, Registered in England & Wales, Company No: 2557590
ARM Holdings plc, Registered office 110 Fulbourn Road, Cambridge CB1 9NJ, Registered in England & Wales, Company No: 2548782
Michele Scandale
2013-08-27 23:12:48 UTC
Permalink
Post by Michele Scandale
The pure solution would be the one proposed by Eli: I don't have any objection
to this solution.
The mangler now has a bug, so it must be fixed. The pure solution implicitly
breaks the binary compatibility. If we do not have problem with this (so we
consider a matter for the users to solve the problem, e.g. with a forced update
of libraries) the right patch is to have a target independent mangling for OpenCL.
Still we would have problems if we consider SPIR: in its specification there is
a fixed mangling scheme (that it's the one produced by the current mangler). In
this case we have two choice: we change the SPIR mangling or we allow targets to
override the target independent mangling for OpenCL with the one based on the
TargetAddrSpaceMap.
*Based on all this would see the mangling proposed by Eli the default except for
targets that explicitly requires a mangling scheme based on the target address
spaces map (e.g. the SPIR target).*
In attachment a proposal to implement target independent mangling with the
option for targets to force the use of target address space based mangling.

Regards,
-Michele
-------------- next part --------------
A non-text attachment was scrubbed...
Name: mangling-rev5.patch
Type: text/x-patch
Size: 4792 bytes
Desc: not available
URL: <http://lists.cs.uiuc.edu/pipermail/cfe-commits/attachments/20130828/a3b0c1ed/attachment.bin>
Mon Ping Wang
2013-09-06 10:54:08 UTC
Permalink
This patch looks fine to me.

? Mon Ping
Post by Michele Scandale
Post by Michele Scandale
The pure solution would be the one proposed by Eli: I don't have any objection
to this solution.
The mangler now has a bug, so it must be fixed. The pure solution implicitly
breaks the binary compatibility. If we do not have problem with this (so we
consider a matter for the users to solve the problem, e.g. with a forced update
of libraries) the right patch is to have a target independent mangling for OpenCL.
Still we would have problems if we consider SPIR: in its specification there is
a fixed mangling scheme (that it's the one produced by the current mangler). In
this case we have two choice: we change the SPIR mangling or we allow targets to
override the target independent mangling for OpenCL with the one based on the
TargetAddrSpaceMap.
*Based on all this would see the mangling proposed by Eli the default except for
targets that explicitly requires a mangling scheme based on the target address
spaces map (e.g. the SPIR target).*
In attachment a proposal to implement target independent mangling with the
option for targets to force the use of target address space based mangling.
Regards,
-Michele
<mangling-rev5.patch>_______________________________________________
cfe-commits mailing list
cfe-commits at cs.uiuc.edu
http://lists.cs.uiuc.edu/mailman/listinfo/cfe-commits
Michele Scandale
2013-09-06 11:48:51 UTC
Permalink
Post by Mon Ping Wang
This patch looks fine to me.
If this patch seems generally fine, I would appreciate if someone can commit it
for me, because I haven't commit access.

Thanks in advance.

Best Regards,
-Michele
Post by Mon Ping Wang
? Mon Ping
Post by Michele Scandale
Post by Michele Scandale
The pure solution would be the one proposed by Eli: I don't have any objection
to this solution.
The mangler now has a bug, so it must be fixed. The pure solution implicitly
breaks the binary compatibility. If we do not have problem with this (so we
consider a matter for the users to solve the problem, e.g. with a forced update
of libraries) the right patch is to have a target independent mangling for OpenCL.
Still we would have problems if we consider SPIR: in its specification there is
a fixed mangling scheme (that it's the one produced by the current mangler). In
this case we have two choice: we change the SPIR mangling or we allow targets to
override the target independent mangling for OpenCL with the one based on the
TargetAddrSpaceMap.
*Based on all this would see the mangling proposed by Eli the default except for
targets that explicitly requires a mangling scheme based on the target address
spaces map (e.g. the SPIR target).*
In attachment a proposal to implement target independent mangling with the
option for targets to force the use of target address space based mangling.
Regards,
-Michele
<mangling-rev5.patch>_______________________________________________
cfe-commits mailing list
cfe-commits at cs.uiuc.edu
http://lists.cs.uiuc.edu/mailman/listinfo/cfe-commits
Michele Scandale
2013-09-10 19:23:27 UTC
Permalink
Reup and updated version of the patch!

Thanks in advance.

Regards,
-Michele
Post by Michele Scandale
Post by Mon Ping Wang
This patch looks fine to me.
If this patch seems generally fine, I would appreciate if someone can commit it
for me, because I haven't commit access.
Thanks in advance.
Best Regards,
-Michele
Post by Mon Ping Wang
? Mon Ping
Post by Michele Scandale
Post by Michele Scandale
The pure solution would be the one proposed by Eli: I don't have any objection
to this solution.
The mangler now has a bug, so it must be fixed. The pure solution implicitly
breaks the binary compatibility. If we do not have problem with this (so we
consider a matter for the users to solve the problem, e.g. with a forced update
of libraries) the right patch is to have a target independent mangling for OpenCL.
Still we would have problems if we consider SPIR: in its specification there is
a fixed mangling scheme (that it's the one produced by the current mangler). In
this case we have two choice: we change the SPIR mangling or we allow targets to
override the target independent mangling for OpenCL with the one based on the
TargetAddrSpaceMap.
*Based on all this would see the mangling proposed by Eli the default except for
targets that explicitly requires a mangling scheme based on the target address
spaces map (e.g. the SPIR target).*
In attachment a proposal to implement target independent mangling with the
option for targets to force the use of target address space based mangling.
Regards,
-Michele
<mangling-rev5.patch>_______________________________________________
cfe-commits mailing list
cfe-commits at cs.uiuc.edu
http://lists.cs.uiuc.edu/mailman/listinfo/cfe-commits
-------------- next part --------------
A non-text attachment was scrubbed...
Name: mangling.patch
Type: text/x-patch
Size: 4792 bytes
Desc: not available
URL: <http://lists.cs.uiuc.edu/pipermail/cfe-commits/attachments/20130910/52620f60/attachment.bin>
David Tweed
2013-09-11 09:55:29 UTC
Permalink
Hi,

while I think most people agree with the direction things are going there
look to still
be some fiddly details. As one instance, when I run this on a standard OSS
LLVM build I get
a new test failure in test/CodeGenOpenCL/local.cl. Since behaviour is being
made more sophisticated, it seems
it would be good to have add some tests that verify the new behaviour so we
can detect any
modifications that change it. But the patch looks to be progressing.

Cheers,
Dave

-----Original Message-----
From: Michele Scandale [mailto:michele.scandale at gmail.com]
Sent: 10 September 2013 20:23
To: Mon Ping Wang
Cc: cfe-dev at cs.uiuc.edu; cfe-commits at cs.uiuc.edu; Tanya Lattner; David
Tweed; Eli Friedman
Subject: Re: OpenCL address space and mangling

Reup and updated version of the patch!

Thanks in advance.

Regards,
-Michele
Post by Michele Scandale
Post by Mon Ping Wang
This patch looks fine to me.
If this patch seems generally fine, I would appreciate if someone can commit it
for me, because I haven't commit access.
Thanks in advance.
Best Regards,
-Michele
Post by Mon Ping Wang
- Mon Ping
On Aug 27, 2013, at 4:12 PM, Michele Scandale
Post by Michele Scandale
Post by Michele Scandale
The pure solution would be the one proposed by Eli: I don't have any objection
to this solution.
The mangler now has a bug, so it must be fixed. The pure solution implicitly
breaks the binary compatibility. If we do not have problem with this (so we
consider a matter for the users to solve the problem, e.g. with a forced update
of libraries) the right patch is to have a target independent mangling for OpenCL.
Still we would have problems if we consider SPIR: in its specification there is
a fixed mangling scheme (that it's the one produced by the current mangler). In
this case we have two choice: we change the SPIR mangling or we allow targets to
override the target independent mangling for OpenCL with the one based on the
TargetAddrSpaceMap.
*Based on all this would see the mangling proposed by Eli the default except for
targets that explicitly requires a mangling scheme based on the target address
spaces map (e.g. the SPIR target).*
In attachment a proposal to implement target independent mangling with the
option for targets to force the use of target address space based mangling.
Regards,
-Michele
<mangling-rev5.patch>_______________________________________________
cfe-commits mailing list
cfe-commits at cs.uiuc.edu
http://lists.cs.uiuc.edu/mailman/listinfo/cfe-commits
Michele Scandale
2013-09-12 11:35:06 UTC
Permalink
Post by David Tweed
while I think most people agree with the direction things are going there
look to still
be some fiddly details. As one instance, when I run this on a standard OSS
LLVM build I get
a new test failure in test/CodeGenOpenCL/local.cl. Since behaviour is being
made more sophisticated, it seems
it would be good to have add some tests that verify the new behaviour so we
can detect any
modifications that change it. But the patch looks to be progressing.
Hi David,

I've fixed the test and added another test specific for mangling checking. To
simplify testing I've added a command line option (similar to
-ffake-address-space-map).

In attachment the new version of the patch.

Thanks in advance.

Best Regards,
Michele
-------------- next part --------------
diff --git a/include/clang/AST/ASTContext.h b/include/clang/AST/ASTContext.h
index 377282f..fd351ab 100644
--- a/include/clang/AST/ASTContext.h
+++ b/include/clang/AST/ASTContext.h
@@ -393,6 +393,10 @@ private:
/// \brief The logical -> physical address space map.
const LangAS::Map *AddrSpaceMap;

+ /// \brief Address space map mangling must be used with language specific
+ /// address spaces (e.g. OpenCL/CUDA)
+ bool AddrSpaceMapMangling;
+
friend class ASTDeclReader;
friend class ASTReader;
friend class ASTWriter;
@@ -1920,6 +1924,12 @@ public:
return (*AddrSpaceMap)[AS - LangAS::Offset];
}

+ bool addressSpaceMapManglingFor(unsigned AS) const {
+ return AddrSpaceMapMangling ||
+ AS < LangAS::Offset ||
+ AS >= LangAS::Offset + LangAS::Count;
+ }
+
private:
// Helper for integer ordering
unsigned getIntegerRank(const Type *T) const;
diff --git a/include/clang/Basic/LangOptions.def b/include/clang/Basic/LangOptions.def
index a2e94ff..55db34c 100644
--- a/include/clang/Basic/LangOptions.def
+++ b/include/clang/Basic/LangOptions.def
@@ -142,6 +142,7 @@ LANGOPT(HexagonQdsp6Compat , 1, 0, "hexagon-qdsp6 backward compatibility")
LANGOPT(ObjCAutoRefCount , 1, 0, "Objective-C automated reference counting")
LANGOPT(ObjCARCWeak , 1, 0, "__weak support in the ARC runtime")
LANGOPT(FakeAddressSpaceMap , 1, 0, "OpenCL fake address space map")
+ENUM_LANGOPT(AddressSpaceMapMangling , AddrSpaceMapMangling, 2, ASMM_Target, "OpenCL address space map mangling mode")

LANGOPT(MRTD , 1, 0, "-mrtd calling convention")
BENIGN_LANGOPT(DelayedTemplateParsing , 1, 0, "delayed template parsing")
diff --git a/include/clang/Basic/LangOptions.h b/include/clang/Basic/LangOptions.h
index 21ca7eb..4532054 100644
--- a/include/clang/Basic/LangOptions.h
+++ b/include/clang/Basic/LangOptions.h
@@ -66,6 +66,8 @@ public:
SOB_Trapping // -ftrapv
};

+ enum AddrSpaceMapMangling { ASMM_Target, ASMM_On, ASMM_Off };
+
public:
clang::ObjCRuntime ObjCRuntime;

diff --git a/include/clang/Basic/TargetInfo.h b/include/clang/Basic/TargetInfo.h
index ee3a28d..bda6af3 100644
--- a/include/clang/Basic/TargetInfo.h
+++ b/include/clang/Basic/TargetInfo.h
@@ -202,6 +202,10 @@ protected:
/// zero length bitfield, regardless of the zero length bitfield type.
unsigned ZeroLengthBitfieldBoundary;

+ /// \brief Specify if mangling based on address space map should be used or
+ /// not for language specific address spaces
+ bool UseAddrSpaceMapMangling;
+
public:
IntType getSizeType() const { return SizeType; }
IntType getIntMaxType() const { return IntMaxType; }
@@ -431,6 +435,12 @@ public:
return ComplexLongDoubleUsesFP2Ret;
}

+ /// \brief Specify if mangling based on address space map should be used or
+ /// not for language specific address spaces
+ bool useAddressSpaceMapMangling() const {
+ return UseAddrSpaceMapMangling;
+ }
+
///===---- Other target property query methods --------------------------===//

/// \brief Appends the target-specific \#define values for this
diff --git a/include/clang/Driver/CC1Options.td b/include/clang/Driver/CC1Options.td
index b74f445..d91001f 100644
--- a/include/clang/Driver/CC1Options.td
+++ b/include/clang/Driver/CC1Options.td
@@ -460,6 +460,8 @@ def fno_bitfield_type_align : Flag<["-"], "fno-bitfield-type-align">,
HelpText<"Ignore bit-field types when aligning structures">;
def ffake_address_space_map : Flag<["-"], "ffake-address-space-map">,
HelpText<"Use a fake address space map; OpenCL testing purposes only">;
+def faddress_space_map_mangling_EQ : Joined<["-"], "faddress-space-map-mangling=">, MetaVarName<"<yes|no|target>">,
+ HelpText<"Set the mode for address space map based mangling; OpenCL testing purposes only">;
def funknown_anytype : Flag<["-"], "funknown-anytype">,
HelpText<"Enable parser support for the __unknown_anytype type; for testing purposes only">;
def fdebugger_support : Flag<["-"], "fdebugger-support">,
diff --git a/lib/AST/ASTContext.cpp b/lib/AST/ASTContext.cpp
index 85ac734..6e77f4e 100644
--- a/lib/AST/ASTContext.cpp
+++ b/lib/AST/ASTContext.cpp
@@ -694,6 +694,19 @@ static const LangAS::Map *getAddressSpaceMap(const TargetInfo &T,
}
}

+static bool isAddrSpaceMapManglingEnabled(const TargetInfo &TI,
+ const LangOptions &LangOpts) {
+ switch (LangOpts.getAddressSpaceMapMangling()) {
+ default: return false;
+ case LangOptions::ASMM_Target:
+ return TI.useAddressSpaceMapMangling();
+ case LangOptions::ASMM_On:
+ return true;
+ case LangOptions::ASMM_Off:
+ return false;
+ }
+}
+
ASTContext::ASTContext(LangOptions& LOpts, SourceManager &SM,
const TargetInfo *t,
IdentifierTable &idents, SelectorTable &sels,
@@ -893,6 +906,7 @@ void ASTContext::InitBuiltinTypes(const TargetInfo &Target) {

ABI.reset(createCXXABI(Target));
AddrSpaceMap = getAddressSpaceMap(Target, LangOpts);
+ AddrSpaceMapMangling = isAddrSpaceMapManglingEnabled(Target, LangOpts);

// C99 6.2.5p19.
InitBuiltinType(VoidTy, BuiltinType::Void);
diff --git a/lib/AST/ItaniumMangle.cpp b/lib/AST/ItaniumMangle.cpp
index c3121c0..e135227 100644
--- a/lib/AST/ItaniumMangle.cpp
+++ b/lib/AST/ItaniumMangle.cpp
@@ -1753,15 +1753,33 @@ void CXXNameMangler::mangleQualifiers(Qualifiers Quals) {
Out << 'K';

if (Quals.hasAddressSpace()) {
- // Extension:
+ // Address space extension:
//
- // <type> ::= U <address-space-number>
- //
- // where <address-space-number> is a source name consisting of 'AS'
- // followed by the address space <number>.
+ // <type> ::= U <target-addrspace>
+ // <type> ::= U <OpenCL-addrspace>
+ // <type> ::= U <CUDA-addrspace>
+
SmallString<64> ASString;
- ASString = "AS" + llvm::utostr_32(
- Context.getASTContext().getTargetAddressSpace(Quals.getAddressSpace()));
+ unsigned AS = Quals.getAddressSpace();
+ bool IsLangAS = (LangAS::Offset <= AS) && (AS < LangAS::Last);
+
+ if (Context.getASTContext().addressSpaceMapManglingFor(AS)) {
+ // <target-addrspace> ::= "AS" <address-space-number>
+ unsigned TargetAS = Context.getASTContext().getTargetAddressSpace(AS);
+ ASString = "AS" + llvm::utostr_32(TargetAS);
+ } else {
+ switch (AS) {
+ default: llvm_unreachable("Not a language specific address space");
+ // <OpenCL-addrspace> ::= "CL" [ "global" | "local" | "constant" ]
+ case LangAS::opencl_global: ASString = "CLglobal"; break;
+ case LangAS::opencl_local: ASString = "CLlocal"; break;
+ case LangAS::opencl_constant: ASString = "CLconstant"; break;
+ // <CUDA-addrspace> ::= "CU" [ "device" | "constant" | "shared" ]
+ case LangAS::cuda_device: ASString = "CUdevice"; break;
+ case LangAS::cuda_constant: ASString = "CUconstant"; break;
+ case LangAS::cuda_shared: ASString = "CUshared"; break;
+ }
+ }
Out << 'U' << ASString.size() << ASString;
}

diff --git a/lib/Basic/TargetInfo.cpp b/lib/Basic/TargetInfo.cpp
index 3feaf9e..e993055 100644
--- a/lib/Basic/TargetInfo.cpp
+++ b/lib/Basic/TargetInfo.cpp
@@ -88,6 +88,7 @@ TargetInfo::TargetInfo(const llvm::Triple &T) : TargetOpts(), Triple(T) {

// Default to an empty address space map.
AddrSpaceMap = &DefaultAddrSpaceMap;
+ UseAddrSpaceMapMangling = false;

// Default to an unknown platform name.
PlatformName = "unknown";
diff --git a/lib/Basic/Targets.cpp b/lib/Basic/Targets.cpp
index aa0993d..a6a613d 100644
--- a/lib/Basic/Targets.cpp
+++ b/lib/Basic/Targets.cpp
@@ -1266,6 +1266,7 @@ namespace {
TLSSupported = false;
LongWidth = LongAlign = 64;
AddrSpaceMap = &NVPTXAddrSpaceMap;
+ UseAddrSpaceMapMangling = true;
// Define available target features
// These must be defined in sorted order!
NoAsmVariants = true;
@@ -1424,6 +1425,7 @@ public:
: TargetInfo(Triple), GPU(GK_R600) {
DescriptionString = DescriptionStringR600;
AddrSpaceMap = &R600AddrSpaceMap;
+ UseAddrSpaceMapMangling = true;
}

virtual const char * getClobbers() const {
@@ -4575,6 +4577,7 @@ namespace {
"f32:32:32-f64:32:32-v64:32:32-"
"v128:32:32-a0:0:32-n32";
AddrSpaceMap = &TCEOpenCLAddrSpaceMap;
+ UseAddrSpaceMapMangling = true;
}

virtual void getTargetDefines(const LangOptions &Opts,
@@ -5137,6 +5140,7 @@ namespace {
TLSSupported = false;
LongWidth = LongAlign = 64;
AddrSpaceMap = &SPIRAddrSpaceMap;
+ UseAddrSpaceMapMangling = true;
// Define available target features
// These must be defined in sorted order!
NoAsmVariants = true;
diff --git a/lib/Frontend/CompilerInvocation.cpp b/lib/Frontend/CompilerInvocation.cpp
index ce4fdc3..2947041 100644
--- a/lib/Frontend/CompilerInvocation.cpp
+++ b/lib/Frontend/CompilerInvocation.cpp
@@ -1326,6 +1326,28 @@ static void ParseLangArgs(LangOptions &Opts, ArgList &Args, InputKind IK,
Opts.ApplePragmaPack = Args.hasArg(OPT_fapple_pragma_pack);
Opts.CurrentModule = Args.getLastArgValue(OPT_fmodule_name);

+ if (Arg *A = Args.getLastArg(OPT_faddress_space_map_mangling_EQ)) {
+ switch (llvm::StringSwitch<unsigned>(A->getValue())
+ .Case("target", LangOptions::ASMM_Target)
+ .Case("no", LangOptions::ASMM_Off)
+ .Case("yes", LangOptions::ASMM_On)
+ .Default(255)) {
+ default:
+ Diags.Report(diag::err_drv_invalid_value)
+ << "-faddress-space-map-mangling=" << A->getValue();
+ break;
+ case LangOptions::ASMM_Target:
+ Opts.setAddressSpaceMapMangling(LangOptions::ASMM_Target);
+ break;
+ case LangOptions::ASMM_On:
+ Opts.setAddressSpaceMapMangling(LangOptions::ASMM_On);
+ break;
+ case LangOptions::ASMM_Off:
+ Opts.setAddressSpaceMapMangling(LangOptions::ASMM_Off);
+ break;
+ }
+ }
+
// Check if -fopenmp is specified.
Opts.OpenMP = Args.hasArg(OPT_fopenmp);

diff --git a/test/CodeGenOpenCL/address-spaces-mangling.cl b/test/CodeGenOpenCL/address-spaces-mangling.cl
index e69de29..7ce74d3 100644
--- a/test/CodeGenOpenCL/address-spaces-mangling.cl
+++ b/test/CodeGenOpenCL/address-spaces-mangling.cl
@@ -0,0 +1,22 @@
+// RUN: %clang_cc1 %s -ffake-address-space-map -faddress-space-map-mangling=yes -emit-llvm -o - | FileCheck -check-prefix=ASMANG %s
+// RUN: %clang_cc1 %s -ffake-address-space-map -faddress-space-map-mangling=no -emit-llvm -o - | FileCheck -check-prefix=NOASMANG %s
+
+__attribute__((overloadable))
+void f(private int *arg) { }
+// ASMANG: @_Z1fPi
+// NOASMANG: @_Z1fPi
+
+__attribute__((overloadable))
+void f(global int *arg) { }
+// ASMANG: @_Z1fPU3AS1i
+// NOASMANG: @_Z1fPU8CLglobali
+
+__attribute__((overloadable))
+void f(local int *arg) { }
+// ASMANG: @_Z1fPU3AS2i
+// NOASMANG: @_Z1fPU7CLlocali
+
+__attribute__((overloadable))
+void f(constant int *arg) { }
+// ASMANG: @_Z1fPU3AS3i
+// NOASMANG: @_Z1fPU10CLconstanti
diff --git a/test/CodeGenOpenCL/local.cl b/test/CodeGenOpenCL/local.cl
index 852fa43..b5c67d9 100644
--- a/test/CodeGenOpenCL/local.cl
+++ b/test/CodeGenOpenCL/local.cl
@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 %s -ffake-address-space-map -emit-llvm -o - | FileCheck %s
+// RUN: %clang_cc1 %s -ffake-address-space-map -faddress-space-map-mangling=no -emit-llvm -o - | FileCheck %s

__kernel void foo(void) {
// CHECK: @foo.i = internal addrspace(2)
@@ -6,7 +6,7 @@ __kernel void foo(void) {
++i;
}

-// CHECK-LABEL: define void @_Z3barPU3AS2i
+// CHECK-LABEL: define void @_Z3barPU7CLlocali
__kernel void __attribute__((__overloadable__)) bar(local int *x) {
*x = 5;
}
David Tweed
2013-09-12 12:31:15 UTC
Permalink
Hi Michele,

This patch LGTM. Assuming you'd still like someone to commit it on your behalf, I'll leave it for a day in case
anyone else has any comments or issues, but can commit it for you end of tomorrow if nothing comes up.

Thanks for working on this,

Cheers,
Dave

-----Original Message-----
From: Michele Scandale [mailto:michele.scandale at gmail.com]
Sent: 12 September 2013 12:35
To: David Tweed
Cc: Mon Ping Wang; cfe-dev at cs.uiuc.edu; cfe-commits at cs.uiuc.edu; Tanya Lattner; Eli Friedman
Subject: Re: OpenCL address space and mangling
Post by David Tweed
while I think most people agree with the direction things are going there
look to still
be some fiddly details. As one instance, when I run this on a standard OSS
LLVM build I get
a new test failure in test/CodeGenOpenCL/local.cl. Since behaviour is being
made more sophisticated, it seems
it would be good to have add some tests that verify the new behaviour so we
can detect any
modifications that change it. But the patch looks to be progressing.
Hi David,

I've fixed the test and added another test specific for mangling checking. To
simplify testing I've added a command line option (similar to
-ffake-address-space-map).

In attachment the new version of the patch.

Thanks in advance.

Best Regards,
Michele

-- IMPORTANT NOTICE: The contents of this email and any attachments are confidential and may also be privileged. If you are not the intended recipient, please notify the sender immediately and do not disclose the contents to any other person, use it for any purpose, or store or copy the information in any medium. Thank you.

ARM Limited, Registered office 110 Fulbourn Road, Cambridge CB1 9NJ, Registered in England & Wales, Company No: 2557590
ARM Holdings plc, Registered office 110 Fulbourn Road, Cambridge CB1 9NJ, Registered in England & Wales, Company No: 2548782
Michele Scandale
2013-09-12 13:08:51 UTC
Permalink
Post by Mon Ping Wang
Hi Michele,
This patch LGTM. Assuming you'd still like someone to commit it on your behalf, I'll leave it for a day in case
anyone else has any comments or issues, but can commit it for you end of tomorrow if nothing comes up.
Thanks for working on this,
Yes, if you can commit it I would appreciate :-). I agree with your plan.

Thanks again.

Regards,
-Michele
Post by Mon Ping Wang
Cheers,
Dave
-----Original Message-----
From: Michele Scandale [mailto:michele.scandale at gmail.com]
Sent: 12 September 2013 12:35
To: David Tweed
Cc: Mon Ping Wang; cfe-dev at cs.uiuc.edu; cfe-commits at cs.uiuc.edu; Tanya Lattner; Eli Friedman
Subject: Re: OpenCL address space and mangling
Post by David Tweed
while I think most people agree with the direction things are going there
look to still
be some fiddly details. As one instance, when I run this on a standard OSS
LLVM build I get
a new test failure in test/CodeGenOpenCL/local.cl. Since behaviour is being
made more sophisticated, it seems
it would be good to have add some tests that verify the new behaviour so we
can detect any
modifications that change it. But the patch looks to be progressing.
Hi David,
I've fixed the test and added another test specific for mangling checking. To
simplify testing I've added a command line option (similar to
-ffake-address-space-map).
In attachment the new version of the patch.
Thanks in advance.
Best Regards,
Michele
-- IMPORTANT NOTICE: The contents of this email and any attachments are confidential and may also be privileged. If you are not the intended recipient, please notify the sender immediately and do not disclose the contents to any other person, use it for any purpose, or store or copy the information in any medium. Thank you.
ARM Limited, Registered office 110 Fulbourn Road, Cambridge CB1 9NJ, Registered in England & Wales, Company No: 2557590
ARM Holdings plc, Registered office 110 Fulbourn Road, Cambridge CB1 9NJ, Registered in England & Wales, Company No: 2548782
David Tweed
2013-09-13 12:08:24 UTC
Permalink
Committed r190684.

-----Original Message-----
From: Michele Scandale [mailto:michele.scandale at gmail.com]
Sent: 12 September 2013 14:09
To: David Tweed
Cc: Mon Ping Wang; cfe-dev at cs.uiuc.edu; cfe-commits at cs.uiuc.edu; Tanya Lattner; Eli Friedman
Subject: Re: OpenCL address space and mangling
Post by Mon Ping Wang
Hi Michele,
This patch LGTM. Assuming you'd still like someone to commit it on your behalf, I'll leave it for a day in case
anyone else has any comments or issues, but can commit it for you end of tomorrow if nothing comes up.
Thanks for working on this,
Yes, if you can commit it I would appreciate :-). I agree with your plan.

Thanks again.

Regards,
-Michele
Post by Mon Ping Wang
Cheers,
Dave
-----Original Message-----
From: Michele Scandale [mailto:michele.scandale at gmail.com]
Sent: 12 September 2013 12:35
To: David Tweed
Cc: Mon Ping Wang; cfe-dev at cs.uiuc.edu; cfe-commits at cs.uiuc.edu; Tanya Lattner; Eli Friedman
Subject: Re: OpenCL address space and mangling
Post by David Tweed
while I think most people agree with the direction things are going there
look to still
be some fiddly details. As one instance, when I run this on a standard OSS
LLVM build I get
a new test failure in test/CodeGenOpenCL/local.cl. Since behaviour is being
made more sophisticated, it seems
it would be good to have add some tests that verify the new behaviour so we
can detect any
modifications that change it. But the patch looks to be progressing.
Hi David,
I've fixed the test and added another test specific for mangling checking. To
simplify testing I've added a command line option (similar to
-ffake-address-space-map).
In attachment the new version of the patch.
Thanks in advance.
Best Regards,
Michele
-- IMPORTANT NOTICE: The contents of this email and any attachments are confidential and may also be privileged. If you are not the intended recipient, please notify the sender immediately and do not disclose the contents to any other person, use it for any purpose, or store or copy the information in any medium. Thank you.
ARM Limited, Registered office 110 Fulbourn Road, Cambridge CB1 9NJ, Registered in England & Wales, Company No: 2557590
ARM Holdings plc, Registered office 110 Fulbourn Road, Cambridge CB1 9NJ, Registered in England & Wales, Company No: 2548782
Eli Friedman
2013-08-02 22:57:03 UTC
Permalink
Post by Tanya Lattner
However, I actually should have looked at this closer as it actually doesn't map
1, // opencl_global
3, // opencl_local
2, // opencl_constant
and when there is no address space then it maps to nothing.
So, I don't think your patch is going to work unless the order is changed in the
enum. Because this is not clearly defined in the spec and is implementation
specific and TARGET specific.. then changing that enum is probably not going to
be the right approach either.
I see the point but still we need to preserve the source language difference
with the mangling.
So, I'm going back to my original statement to keep it to be Target specific.
For your library, are these functions actually implemented differently? Wouldn't
they be exactly the same when there is no address space? In our implementation
we have an address space map defined for X86 and then the names get mangled
"correctly" for all targets. But, all the functionality is the same since the
address spaces don't impact codegen for X86.
I'd argue that the fact that address spaces do not impact codegen for X86 is
merely incidental: the issue goes beyond X86 and impacts all targets --
including future ones.
By choosing to use a fake address space map (instead of the one fitting the
target description), we introduce in the IR a potential for breaking future
implementations, even though the current X86 target is not affected.
Moreover, by introducing the fake address map we would violate the semantics
of the LLVM IR addrspace modifier.
Even binary compatibility with libraries already distributed can be easily achieved.
This patch (see attachment) aims at preserving the mangling that were
generated by using the target address space map for those targets that
override it, while introducing in the mangling the distinction of
opencl/cuda address spaces for those targets that do not have a non trivial
target address space map.
By the way, looking beyond the scope of the current issue of mangling, it
would be IMO interesting to start a public discussion on the mailing list
about a way to represent logical address space information different from
target-specific address space (the case for OpenCL and CUDA) in order to
allow the implementation of custom language specific analysis and/or
optimization.
As a temporary solution, if one needed the logical address space information
in the IR too for specific purpose (like OpenCL specific optimization), can
still override the address space map of the target.
This discussion probably should be moved to the cfe-dev mailing list. I
think that its better for Clang to have one consistent way of mangling
address spaces regardless of language. I?ve been thinking about this more
and another problem I have with not using the Target address space map for
void __attribute__((__overloadable__)) foo(global int *x)
The address space on the argument is gone (or zero), but yet you have it in
the mangled name. So its not consistent.
I can agree that its right to mangle the names differently from the language
perspective, but what you mangle them to is really target specific. If you
want to remove this notion from Clang, then maybe all target specific
address space maps should go away and a default one for all is used. Then
each LLVM backend can interpret it as they wish.
It would be great if some Code owners could weigh in here.
The mangling of a symbol shouldn't depend on its machine
representation. For example, in C++, "long*" and "long long*" have a
different mangling, and that mangling is consistent across platforms
regardless of the actual width in bits of "long" and "long long".
This is required to properly represent overloading. The same concept
applies to "__global int *" and "__local int *"

That said, using numbers here is a terrible idea. Strawman
suggestion: the mangling for __global should be "CLglobal".

-Eli
Continue reading on narkive:
Loading...