rjmccall (John McCall)
User

Projects

User does not belong to any projects.

User Details

User Since
Jan 18 2013, 11:30 AM (274 w, 3 d)

Recent Activity

Yesterday

rjmccall added inline comments to D45476: [C++2a] Implement operator<=> CodeGen and ExprConstant.
Mon, Apr 23, 9:05 PM
rjmccall added a comment to D45766: [Sema] Add -Wno-self-assign-overloaded.

Is there anything else in the "-w" namespace other than the literal "-w" so
far?

Mon, Apr 23, 3:29 PM
rjmccall added a comment to D45766: [Sema] Add -Wno-self-assign-overloaded.

FWIW I don't fundamentalyl object to also having something like -wtest.
Probably needs a better name though (unfortunately the double-negative gets
confusing... - like you want to describe the set of diagnostics that should
not be used in test code, so that as a group might be "-Wnon-test" but then
"-Wno-non-test" is pretty awkward) - probably worth chatting to Richard
Smith about that, I reckon.

Mon, Apr 23, 3:22 PM
rjmccall added inline comments to D45382: [CodeGen] Avoid destructing a struct type that has already been destructed by a delegated constructor.
Mon, Apr 23, 3:17 PM
rjmccall added inline comments to D45382: [CodeGen] Avoid destructing a struct type that has already been destructed by a delegated constructor.
Mon, Apr 23, 1:41 PM
rjmccall added inline comments to D45451: [ItaniumMangle] Undeduced auto type doesn't belong in the substitution table.
Mon, Apr 23, 11:06 AM

Sat, Apr 21

rjmccall added inline comments to D45451: [ItaniumMangle] Undeduced auto type doesn't belong in the substitution table.
Sat, Apr 21, 5:56 AM
rjmccall added inline comments to D45451: [ItaniumMangle] Undeduced auto type doesn't belong in the substitution table.
Sat, Apr 21, 5:56 AM
Herald added a reviewer for D45382: [CodeGen] Avoid destructing a struct type that has already been destructed by a delegated constructor: javed.absar.
Sat, Apr 21, 5:49 AM
rjmccall added a comment to D44435: CUDA ctor/dtor Module-Unique Symbol Name.

This does not address my review. My review is suggesting that we avoid this issue completely by fixing IRGen to use an external linkage for internal declarations in your emission mode. That would allow you to just emit the module ctors as truly internal in the first place, removing any need to mangle them.

Sat, Apr 21, 5:38 AM
rjmccall added a reviewer for D45898: [SemaCXX] Mark destructor as referenced : doug.gregor.
Sat, Apr 21, 5:33 AM
rjmccall added a comment to D45900: CodeGen: Fix invalid bitcast for lifetime.start/end.

These functions must predate the addition of CreateLifetimeStart and CreateLifetimeEnd methods on IRBuilder; we should just be calling those.

Sat, Apr 21, 4:24 AM

Wed, Apr 18

rjmccall added a comment to D45223: [CUDA] Set LLVM calling convention for CUDA kernel.

Yes, I'm sorry, I think you're right. I had misunderstood the language problem when I suggested going down this road.

Wed, Apr 18, 1:13 PM
rjmccall added a comment to D45685: [Sema] Add -wtest global flag that silences -Wself-assign for overloaded operators..

Let me try to summarize where I think we are.

Wed, Apr 18, 5:56 AM
rjmccall added a comment to D45685: [Sema] Add -wtest global flag that silences -Wself-assign for overloaded operators..

The goal of having a unified option is that you can uniformly suppress warnings that don't always make sense in unit tests. It's future-proofing against the addition of other warnings (probably C++ warnings) that might not make sense for unit tests, like extending the x &= x warning (which is not in -Wself-assign) to user-defined operators. You don't think you would be able to take advantage of that? Because -Wno-self-assign-overloaded-nonfield is rather, uh, pretty precisely targeted for exactly that use case; I can't imagine why someone would use -Wself-assign-overloaded-nofield *positively*, or even *negatively* for any purpose besides suppressing a unit-test problem.

Wed, Apr 18, 4:54 AM
rjmccall added a comment to D45685: [Sema] Add -wtest global flag that silences -Wself-assign for overloaded operators..

I'm sorry, that warning *is* in self-assign, although I'm not sure it should be.

Wed, Apr 18, 4:54 AM

Mon, Apr 16

rjmccall accepted D45310: Warn about memcpy'ing non-trivial C structs.

Thanks; LGTM.

Mon, Apr 16, 3:17 PM
rjmccall added inline comments to D45310: Warn about memcpy'ing non-trivial C structs.
Mon, Apr 16, 3:00 PM
rjmccall accepted D45578: Add a command line option `fregister_global_dtors_with_atexit` to register destructor functions annotated with __attribute__((destructor)) using __cxa_atexit or atexit..

Thank you. Minor request and then LGTM.

Mon, Apr 16, 2:39 AM

Sat, Apr 14

rjmccall accepted D45441: [HIP] Add predefined macros __HIPCC__ and __HIP_DEVICE_COMPILE__.

If so, LGTM.

Sat, Apr 14, 6:24 AM
rjmccall accepted D45489: [HIP] Add input type for HIP.

LGTM.

Sat, Apr 14, 6:20 AM
rjmccall added inline comments to D44984: [HIP] Add hip input kind and codegen for kernel launching.
Sat, Apr 14, 2:59 AM

Fri, Apr 13

rjmccall added inline comments to D45578: Add a command line option `fregister_global_dtors_with_atexit` to register destructor functions annotated with __attribute__((destructor)) using __cxa_atexit or atexit..
Fri, Apr 13, 12:29 AM
rjmccall accepted D45212: [HIP] Let CUDA toolchain support HIP language mode and amdgpu.

I'd still prefer if someone with more driver-design expertise weighed in, but we might not have any specialists there.

Fri, Apr 13, 12:16 AM
rjmccall added inline comments to D45451: [ItaniumMangle] Undeduced auto type doesn't belong in the substitution table.
Fri, Apr 13, 12:12 AM
rjmccall added inline comments to D44984: [HIP] Add hip input kind and codegen for kernel launching.
Fri, Apr 13, 12:07 AM

Sat, Apr 7

rjmccall accepted D45410: [Sema] Remove dead code in BuildAnonymousStructUnionMemberReference. NFCI.

I think there was a point when we weren't always creating CXXThisExprs eagerly for these accesses. Now that we are, yeah, this should be dead.

Sat, Apr 7, 9:22 PM
rjmccall accepted D45412: [Sema] Fix PR22637 - IndirectFieldDecl's discard qualifiers during template instantiation..

Okay, LGTM.

Sat, Apr 7, 9:19 PM
rjmccall accepted D45411: [Sema] Fix PR35832 - Ambiguity accessing anonymous struct/union with multiple bases..

Well, that is a really silly bug. Fix LGTM.

Sat, Apr 7, 9:17 PM
rjmccall committed rC329513: Generalize the swiftcall API since being passed indirectly isn't.
Generalize the swiftcall API since being passed indirectly isn't
Sat, Apr 7, 1:21 PM
rjmccall committed rL329513: Generalize the swiftcall API since being passed indirectly isn't.
Generalize the swiftcall API since being passed indirectly isn't
Sat, Apr 7, 1:21 PM
rjmccall committed rC329508: Allow equality comparisons between block pointers and.
Allow equality comparisons between block pointers and
Sat, Apr 7, 10:45 AM
rjmccall committed rL329508: Allow equality comparisons between block pointers and.
Allow equality comparisons between block pointers and
Sat, Apr 7, 10:45 AM
rjmccall closed D44580: Sema: allow comparison between blocks & block-compatible objc types.

Committed as r329508.

Sat, Apr 7, 10:45 AM · Restricted Project
rjmccall accepted D45305: ObjCGNU: Fix empty v3 protocols being emitted two fields short.

Hmm. Alright, I guess.

Sat, Apr 7, 10:40 AM · Restricted Project

Fri, Apr 6

rjmccall added a comment to D45382: [CodeGen] Avoid destructing a struct type that has already been destructed by a delegated constructor.

If that's the problem, then I think the right design is for CallArg to include an optional cleanup reference for a cleanup that can be deactivated at the instant of the call (we should assert that this exists for parameters that are destroyed in the callee), and then for forwarding it's just a matter of getting that cleanup from parameter emission to forwarding-argument emission.

Fri, Apr 6, 9:17 PM
rjmccall added a comment to D45384: [ObjC++] Never pass structs with `__weak` fields in registers .

Just a couple minor requests; if you accept them, feel free to commit.

Fri, Apr 6, 8:33 PM
rjmccall added a comment to D45212: [HIP] Let CUDA toolchain support HIP language mode and amdgpu.

This looks okay to me, but I think it would better if someone with more expertise in the design of the driver and frontend code could review this.

Fri, Apr 6, 8:23 PM
rjmccall added a comment to D45384: [ObjC++] Never pass structs with `__weak` fields in registers .

Yes. I intended it as a property that propagates to classes that contain or derive from the type.

Would it make it less confusing if I merged CXXRecordDecl::CanPassInRegisters and RecordDecl::CannotPassInRegisters into a single enum? For example, the enum could have three enumerators, "CanPass", "CannotPass", "CanNeverPass", or something. Both "CannotPass" and "CanNeverPass" would force the type to be passed indirectly, and the only difference is that "CanNeverPass" propagates its property outwards and "CannotPass" doesn't. C structs are either "CanPass" or "CanNeverPass" while C++ structs can take any of the three values.

Fri, Apr 6, 4:25 PM
rjmccall added a comment to D45384: [ObjC++] Never pass structs with `__weak` fields in registers .

Well, but I think CanPassInRegisters==false in the base class does always mean CanPassInRegisters==false in the subclass.

I think there are cases that is not true. If I compile the following code, S0 (base class) is passed indirectly and S1 (derived class) is passed directly.

Fri, Apr 6, 2:31 PM
rjmccall accepted D44616: Update existed CodeGen TBAA tests.

Yes, this is fine.

Fri, Apr 6, 2:08 PM
rjmccall added a comment to D45384: [ObjC++] Never pass structs with `__weak` fields in registers .

Well, but I think CanPassInRegisters==false in the base class does always mean CanPassInRegisters==false in the subclass.

Fri, Apr 6, 12:43 PM
rjmccall added a comment to D45382: [CodeGen] Avoid destructing a struct type that has already been destructed by a delegated constructor.

Hmm. I'm not actually sure *why* it's not okay to forward callee-cleanup arguments. Do we just not have the necessary logic to disable the cleanup in the caller?

Fri, Apr 6, 12:41 PM
rjmccall added inline comments to D45310: Warn about memcpy'ing non-trivial C structs.
Fri, Apr 6, 10:29 AM

Thu, Apr 5

rjmccall added a comment to D42092: implement C++ dr388 for the Itanium C++ ABI: proper handling of catching exceptions by reference.

Issue #3 is tricky; it's arguably not okay to force a landing at a landing pad if we're not actually catching anything.

I think the only way this is observable (at least within the bounds of well-defined C++ code) is whether -- and how much -- the stack is unwound before a call to terminate() in the case where an exception is not caught.

[except.handle]p9 says: "If no matching handler is found, the function std::terminate() is called; whether or not the stack is unwound before this call to std::terminate() is implementation-defined (18.5.1)."
[except.terminate]p2 says: " In the situation where no matching handler is found, it is implementation-defined whether or not the stack is unwound before std::terminate() is called. In the situation where the search for a handler (18.3) encounters the outermost block of a function with a non-throwing exception specification (18.4), it is implementation-defined whether the stack is unwound, unwound partially, or not unwound at all before std::terminate() is called. In all other situations, the stack shall not be unwound before std::terminate() is called."

So I think you're right; this solution to issue #3 is non-conforming, because it will result in the stack being partially unwound on the path to a call to std::terminate() for an unhandled exception. Even if we changed the rule for terminate to allow partial unwinding when no handler is found, this patch would still give a QoI regression for that case (because the call to the terminate_handler would happen after some unwinding had been performed).

I'm not particularly interested in pushing this patch further if it can't actually make catch-by-reference work properly. Let's drop this for now and encourage interested parties to work on an alternative personality function.

Thu, Apr 5, 1:23 PM
rjmccall accepted D45306: PR36992 don't overwrite virtual bases in tail padding.

Okay, LGTM.

Thu, Apr 5, 1:20 PM
rjmccall added inline comments to D45310: Warn about memcpy'ing non-trivial C structs.
Thu, Apr 5, 12:23 PM
rjmccall added a comment to D45306: PR36992 don't overwrite virtual bases in tail padding.

Changing the requirements on the return-value slot would be an ABI break, no? And it would force to us to treat all complete-object constructions as nvsize-limited unless they're constructions of known sizeof-sized allocations.

Thu, Apr 5, 12:12 PM
rjmccall added a comment to D42092: implement C++ dr388 for the Itanium C++ ABI: proper handling of catching exceptions by reference.

This is definitely something that the personality function should handle. If we want to do heroic things in the absence of personality function support, we can, but the code should at least be written to be conditional on personality support.

I suppose the question is, do we want to do such heroic things?

Also, we don't have a personality function that can get this right, nor even a concrete design proposed for one. It doesn't seem too hard -- we'd need to extend the EH table to indicate catch-by-reference -- but there's nothing for this to be made conditional on right now. I can easily add a CodeGenOption that turns this all off, if you'd like.

Thu, Apr 5, 11:53 AM
rjmccall added a comment to D45305: ObjCGNU: Fix empty v3 protocols being emitted two fields short.

The change LGTM, but please add a test.

Thu, Apr 5, 11:18 AM · Restricted Project

Wed, Apr 4

rjmccall added inline comments to D45289: Disable -fmerge-all-constants as default..
Wed, Apr 4, 3:40 PM
rjmccall added a comment to D45289: Disable -fmerge-all-constants as default..

It's not unreasonable to just add -fmerge-all-constants to the command line invocations for the individual tests, yeah. We can take those off as necessary later.

Wed, Apr 4, 3:25 PM
rjmccall accepted D45277: [CUDA] Add amdgpu sub archs.

Thanks for splitting the patch up. LGTM.

Wed, Apr 4, 11:49 AM
rjmccall added inline comments to D45212: [HIP] Let CUDA toolchain support HIP language mode and amdgpu.
Wed, Apr 4, 9:11 AM

Tue, Apr 3

rjmccall added a comment to D45212: [HIP] Let CUDA toolchain support HIP language mode and amdgpu.

The other changes I see here seem reasonable, but please do split the patch.

Tue, Apr 3, 9:32 PM
rjmccall accepted D44580: Sema: allow comparison between blocks & block-compatible objc types.

Okay, LGTM with the reduced set of changes to the functionality.

Tue, Apr 3, 9:25 PM · Restricted Project
rjmccall accepted D45176: implement recent "standard-layout" changes.

Okay. LGTM. Thank you for putting the effort into maintaining both rules simultaneously.

Tue, Apr 3, 9:22 PM
rjmccall added a comment to D45179: [libc++] Add _LIBCPP_FORCE_NODISCARD define to force-enable nodiscard in pre-C++17.

Yes, if we think that the committee is likely to include questionable functions on the [[nodiscard]] list which we don't want to warn about pre-C++17, then it makes sense to have two internal macros, one for functions like std::move that should be unconditionally warned about and one for the iffy-bu-standard-directed cases.

Tue, Apr 3, 3:11 PM
rjmccall added inline comments to D45176: implement recent "standard-layout" changes.
Tue, Apr 3, 2:43 PM
rjmccall added a comment to D45223: [CUDA] Set LLVM calling convention for CUDA kernel.

I think the appropriate place to do this is in IsStandardConversion, immediately after the call to ResolveAddressOfOverloadedFunction. You might want to add a general utility for getting the type-of-reference of a function decl.

Tue, Apr 3, 2:33 PM
rjmccall added a comment to D45179: [libc++] Add _LIBCPP_FORCE_NODISCARD define to force-enable nodiscard in pre-C++17.

Is Marshall arguing that the standard doesn't allow compilers to warn about failing to use these function results prior to C++17? Because I don't think that's true; warnings are thoroughly non-normative.

Tue, Apr 3, 2:31 PM
rjmccall added a comment to D45163: [Sema] -Wunused-value: diagnose unused std::move() call results..

Trust me, I understand that this is an important function, but a lot of functions are important, and we're not going to hardcode knowledge about all of them in the compiler.

Tue, Apr 3, 12:50 PM
rjmccall added a comment to D45163: [Sema] -Wunused-value: diagnose unused std::move() call results..

Well, we should feel welcome to submit patches to the C++ library implementations, I think. I can understand why Marshall is waiting to just do this until he gets a comprehensive committee paper, but he might consider taking a patch in the meantime. But I'm not sure what the rush is if it doesn't change what goes into any concrete release of the compiler + library.

Tue, Apr 3, 10:44 AM

Mon, Apr 2

rjmccall added a comment to D44536: Avoid segfault when destructor is not yet known.

It seems that we have two options: either we valiantly try to support this:

  • we keep a list of types for which we've tried to form a delete-expression, but found that the type was incomplete
  • when such a type is completed, we mark the destructor as used, triggering its synthesis or instantiation if necessary, and likewise for the operator delete
  • we change our CXXDeleteExpr representation so that we don't store any information that might be changed by completing the class on the expression, such as the selected OperatorDelete function and whether we're performing a sized delete
Mon, Apr 2, 4:59 PM
rjmccall added a comment to D44883: [Sema] Extend -Wself-assign and -Wself-assign-field to warn on overloaded self-assignment (classes).

No, the analysis is intentionally syntactic; it should apply even on dependently-typed arguments (but not get re-checked on instantiation). But I agree that the warning ought to be disabled in unevaluated contexts.

Mon, Apr 2, 4:37 PM
rjmccall added a comment to D45163: [Sema] -Wunused-value: diagnose unused std::move() call results..

That seems reasonable. And this summer would still be in time for the next release.

Mon, Apr 2, 11:42 AM
rjmccall added a comment to D45163: [Sema] -Wunused-value: diagnose unused std::move() call results..

Yeah, actually, I'm second-guessing myself. Maybe this should just be a libc++ / libstdc++ bug.

Mon, Apr 2, 10:41 AM
rjmccall added a comment to D45163: [Sema] -Wunused-value: diagnose unused std::move() call results..

LGTM. I think it wouldn't be unreasonable to ask standard library maintainers to add [[nodiscard]] to std::move, but it's also not unreasonable for us to special-case some functions.

Mon, Apr 2, 10:14 AM
rjmccall accepted D44985: [CUDA] Let device-side shared variables be initialized with undef.

LGTM, thanks.

Mon, Apr 2, 10:12 AM
rjmccall added a comment to D44883: [Sema] Extend -Wself-assign and -Wself-assign-field to warn on overloaded self-assignment (classes).

Right. I think it's fair to acknowledge that many data structure unit tests will contain a legitimate use of a user-defined self-assignment without feeling that that disqualifies the warning.

Mon, Apr 2, 9:13 AM

Sun, Apr 1

rjmccall committed rC328942: Fix a major swiftcall ABI bug with trivial C++ class types..
Fix a major swiftcall ABI bug with trivial C++ class types.
Sun, Apr 1, 2:08 PM
rjmccall committed rL328942: Fix a major swiftcall ABI bug with trivial C++ class types..
Fix a major swiftcall ABI bug with trivial C++ class types.
Sun, Apr 1, 2:07 PM

Sat, Mar 31

rjmccall accepted D44883: [Sema] Extend -Wself-assign and -Wself-assign-field to warn on overloaded self-assignment (classes).

Thanks. LGTM whenever you've cleared up the self-hosting problems.

Sat, Mar 31, 4:25 PM
rjmccall accepted D45112: [MS] Emit vftable thunks for functions with incomplete prototypes.

LGTM.

Sat, Mar 31, 10:04 AM
rjmccall added inline comments to D44883: [Sema] Extend -Wself-assign and -Wself-assign-field to warn on overloaded self-assignment (classes).
Sat, Mar 31, 9:59 AM
rjmccall added a comment to D45082: [RFC][unittests] ADT: silence -Wself-assign diagnostics.
  • Rebase
  • Drop comments
  • Silence the warning via "*&" contraption, not by casting to the reference.

    The main question still stands, do we want to recommend everyone to use that in their codebases, or double the amount of diag flags and provide a way to silence this specific diagnostic?
Sat, Mar 31, 9:53 AM

Fri, Mar 30

rjmccall added a comment to D44238: [CFG] Fix automatic destructors when a member is bound to a reference..

Seems fine to me, but you might want someone with analyzer experience to review.

Fri, Mar 30, 8:28 PM
rjmccall accepted D45101: [ObjC] Use the name specified by objc_runtime_name instead of the class identifier.

Note that CGObjCNonFragileABIMac::EmitClassRef also passes the class identifier to CGObjCNonFragileABIMac::EmitClassRefFromId, but it doesn't cause a problem. CGObjCNonFragileABIMac::EmitClassRefFromId uses the identifier only when the ObjCInterfaceDecl passed to it is null and that happens only when it is called from CGObjCNonFragileABIMac::EmitNSAutoreleasePoolClassRef.

Fri, Mar 30, 8:19 PM

Thu, Mar 29

rjmccall added inline comments to D44985: [CUDA] Let device-side shared variables be initialized with undef.
Thu, Mar 29, 2:46 PM
rjmccall added a comment to D44559: [Sema] Wrong width of result of mul operation.

What would the design for that warning be? C promotes all arithmetic on sub-int types to int, and if that's assigned back to a variable of the original type, there's a conversion. But you seem to only want to warn about truncating the result of multiplication and not, say, addition or negation. Is there a principle to this? Just the likelihood of escaping the range of the original type?

Thu, Mar 29, 11:43 AM · Restricted Project
rjmccall added a comment to D44985: [CUDA] Let device-side shared variables be initialized with undef.

What exactly are you trying to express here? Are you just trying to make these external declarations when compiling for the device because __shared__ variables are actually defined on the host? That should be handled by the frontend by setting up the AST so that these declarations are not definitions.

No. These variables are not like external symbols defined on the host. They behave like global variables in the kernel code but never initialized. Currently no targets are able to initialize them and it is users' responsibility to initialize them explicitly.

Giving them an initial value will cause error in some backends since they cannot handle them, therefore put undef as initializer.

So undef is being used as a special marker to the backends that it's okay not to try to initialize these variables?

I think undef as the initializer tells the llvm passes and backend that this global variable contains undefined value. I am not sure if this is better than without an initializer. I saw code in CodeGenModule::getOrCreateStaticVarDecl

// Local address space cannot have an initializer.
llvm::Constant *Init = nullptr;
if (Ty.getAddressSpace() != LangAS::opencl_local)
  Init = EmitNullConstant(Ty);
else
  Init = llvm::UndefValue::get(LTy);

which means OpenCL static variable in local address space (equivalent to CUDA shared address space) gets an undef initializer.

For CUDA shared variable, in CodeGenFunction::EmitStaticVarDecl, it first goes through call of CodeGenModule::getOrCreateStaticVarDecl and gets a zeroinitializer, then it reaches line 400

// Whatever initializer such variable may have when it gets here is
  // a no-op and should not be emitted.
  bool isCudaSharedVar = getLangOpts().CUDA && getLangOpts().CUDAIsDevice &&
                         D.hasAttr<CUDASharedAttr>();
  // If this value has an initializer, emit it.
  if (D.getInit() && !isCudaSharedVar)
    var = AddInitializerToStaticVarDecl(D, var);

Although this disables adding initializer from D, var already has a zeroinitializer from CodeGenModule::getOrCreateStaticVarDecl, therefore its initializer needs to be overwritten by undef.

Probably a better solution would be do it in CodeGenModule::getOrCreateStaticVarDecl, side by side by the OpenCL code.

Thu, Mar 29, 9:25 AM
rjmccall added a comment to D44883: [Sema] Extend -Wself-assign and -Wself-assign-field to warn on overloaded self-assignment (classes).

Ok, tried llvm stage2 build, and it failed as expected, in code that intentionally does self-assignment:

Thu, Mar 29, 9:19 AM

Wed, Mar 28

rjmccall added a comment to D44536: Avoid segfault when destructor is not yet known.

I see, so Sema::CheckCompletedCXXClass probably isn't the right place to call DeclareImplicitDestructor as that could significantly increase the size of the AST.

Wed, Mar 28, 9:41 PM
rjmccall added a comment to D44536: Avoid segfault when destructor is not yet known.

I think it's part of an effort to avoid creating implicit declarations for all the special members of every struct we parse from system headers.

Wed, Mar 28, 6:50 PM
rjmccall accepted D44968: Generalize NRVO to cover C structs.

LGTM!

Wed, Mar 28, 3:34 PM
rjmccall added inline comments to D44883: [Sema] Extend -Wself-assign and -Wself-assign-field to warn on overloaded self-assignment (classes).
Wed, Mar 28, 3:34 PM
rjmccall added inline comments to D44883: [Sema] Extend -Wself-assign and -Wself-assign-field to warn on overloaded self-assignment (classes).
Wed, Mar 28, 3:04 PM
rjmccall added a comment to D44968: Generalize NRVO to cover C structs.

Wow, the IR improvements here are amazing.

Wed, Mar 28, 2:52 PM
rjmccall accepted D44989: [ast] Do not auto-initialize Objective-C for-loop variables in Objective-C++ in templatized code under ARC.

LGTM. I guess it just didn't matter for Objective-C because the variable always has to be an Objective-C pointer and the extra store of nil never hurt anything.

Wed, Mar 28, 2:42 PM
rjmccall accepted D44747: Set calling convention for CUDA kernel.

LGTM.

Wed, Mar 28, 2:40 PM
rjmccall accepted D44987: Disable emitting static extern C aliases for amdgcn target for CUDA.

LGTM.

Wed, Mar 28, 2:39 PM
rjmccall added a comment to D44985: [CUDA] Let device-side shared variables be initialized with undef.

What exactly are you trying to express here? Are you just trying to make these external declarations when compiling for the device because __shared__ variables are actually defined on the host? That should be handled by the frontend by setting up the AST so that these declarations are not definitions.

No. These variables are not like external symbols defined on the host. They behave like global variables in the kernel code but never initialized. Currently no targets are able to initialize them and it is users' responsibility to initialize them explicitly.

Giving them an initial value will cause error in some backends since they cannot handle them, therefore put undef as initializer.

Wed, Mar 28, 2:38 PM
rjmccall added a comment to D44984: [HIP] Add hip input kind and codegen for kernel launching.

You should send an RFC to cfe-dev about adding this new language mode. I understand that it's very similar to an existing language mode that we already support, and that's definitely we'll consider, but we shouldn't just agree to add new language modes in patch review.

Wed, Mar 28, 12:04 PM
rjmccall added inline comments to D44987: Disable emitting static extern C aliases for amdgcn target for CUDA.
Wed, Mar 28, 12:00 PM
rjmccall added a comment to D44985: [CUDA] Let device-side shared variables be initialized with undef.

What exactly are you trying to express here? Are you just trying to make these external declarations when compiling for the device because __shared__ variables are actually defined on the host? That should be handled by the frontend by setting up the AST so that these declarations are not definitions.

Wed, Mar 28, 11:55 AM
rjmccall added a comment to D44747: Set calling convention for CUDA kernel.

If __global__ is supported in C++ structures, you might also need to make sure that member function constants (&A::kernel_function) drop the CC. And it might be a good idea to make sure that decltype(kernel_function) doesn't have a problem with it, either, since that does do some special-case work.

Wed, Mar 28, 11:50 AM
rjmccall added inline comments to D44883: [Sema] Extend -Wself-assign and -Wself-assign-field to warn on overloaded self-assignment (classes).
Wed, Mar 28, 11:45 AM

Tue, Mar 27

rjmccall added a comment to D44968: Generalize NRVO to cover C structs.

Is it possible to just do this for all structs? I don't think it hurts anything to do it for structs that are trivial and returned indirectly, and it would certainly be nice to construct C return values in-place even if they're guilty of nothing more than being, y'know, really really big.

Tue, Mar 27, 9:05 PM
rjmccall added inline comments to D44747: Set calling convention for CUDA kernel.
Tue, Mar 27, 1:22 PM
rjmccall added inline comments to D44883: [Sema] Extend -Wself-assign and -Wself-assign-field to warn on overloaded self-assignment (classes).
Tue, Mar 27, 11:05 AM