Anastasia added inline comments.
================
Comment at: include/clang/Basic/BuiltinsAMDGPU.def:49
+
+// FIXME: Need to disallow constant address space.
BUILTIN(__builtin_amdgcn_div_scale, "dddbb*", "n")
----------------
arsenm wrote:
> Anastasia wrote:
> > Do you plan to provide the support for it later? Or if else perhaps we
> > should elaborate more what's to be done.
> I'm not sure. I don't know how to best enforce this
The only way I guess if we list overloads with each address space explicitly
(apart from `constant`)? Or may be with the use of `generic` AS, although that
will only work for CL2.0.
================
Comment at: lib/AST/ASTContext.cpp:9093
unsigned AddrSpace = strtoul(Str, &End, 10);
- if (End != Str && AddrSpace != 0) {
- Type = Context.getAddrSpaceQualType(Type,
- getLangASFromTargetAS(AddrSpace));
+ if (End != Str) {
+ // Note AddrSpace == 0 is not the same as an unspecified address space.
----------------
arsenm wrote:
> Anastasia wrote:
> > Could we check against LangAS::Default instead of removing this completely.
> I don't think that really make sense, since that would be leaving this the
> same. I don't really need it for this patch, but I fundamentally think
> specifying address space 0 is different from an unspecified address space.
> According to the description for builtins, if no address space is specified
> than any address space will be accepted. This is different from a builtin
> requiring address space 0
I thought `Default` AS was meant to be for the case no AS is specified but I
guess it doesn't work the same in Builtins specification syntax.
================
Comment at: lib/CodeGen/CGBuiltin.cpp:3500
+ if (auto *PtrTy = dyn_cast<llvm::PointerType>(PTy)) {
+ if (PtrTy->getAddressSpace() !=
+ ArgValue->getType()->getPointerAddressSpace()) {
----------------
arsenm wrote:
> Anastasia wrote:
> > Would this be correct for OpenCL? Should we use `isAddressSpaceSupersetOf`
> > helper instead? Would it also sort the issue with constant AS (at least for
> > OpenCL)?
> The issue I mentioned for the other builtin is that it modifies the memory,
> and doesn't have to do with the casting.
>
> At this point the AddrSpaceCast has to be emitted. The checking if the cast
> is legal I guess would be in the SemaExpr part. I know at one point I was
> trying to use isAddressSpaceSupersetOf in rewriteBuiltinFunctionDecl, but
> there was some problem with that. I think it didn't make sense with the magic
> where the builtin without an address space is supposed to accept any address
> space or something along those lines.
Yes, I think Sema has to check it before indeed. I am not sure it works right
with OpenCL rules though for the Builtin functions. Would it make sense to add
a negative test for this then?
https://reviews.llvm.org/D47154
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
