arsenm 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")
----------------
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
================
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.
----------------
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
================
Comment at: lib/CodeGen/CGBuiltin.cpp:3500
+ if (auto *PtrTy = dyn_cast<llvm::PointerType>(PTy)) {
+ if (PtrTy->getAddressSpace() !=
+ ArgValue->getType()->getPointerAddressSpace()) {
----------------
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.
================
Comment at: test/CodeGenOpenCL/numbered-address-space.cl:36
+#if 0
+// XXX: Should this compile?
+void
test_generic_as_to_builtin_parameter_explicit_cast_numeric(__attribute__((address_space(3)))
int *local_ptr, float src) {
----------------
Anastasia wrote:
> `__attribute__((address_space(N)))` is not an OpenCL feature and I think it's
> not specified in C either? But I think generally non matching address spaces
> don't compile in Clang. So it might be useful to disallow this?
I'm pretty sure it's a C extension. The way things seem to work now is address
spaces are accepted anywhere and everywhere.
https://reviews.llvm.org/D47154
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
