Page MenuHomePhabricator

[Sema] Produce diagnostics when unavailable aligned allocation/deallocation functions are called
ClosedPublic

Authored by ahatanak on Jun 4 2018, 7:25 PM.

Details

Summary

Deallocation functions are called from deleting destructors that are declared virtual. Issue an error if an aligned deallocation function is selected but is not available.

Diff Detail

Repository
rC Clang

Event Timeline

ahatanak created this revision.Jun 4 2018, 7:25 PM

I think we should sink the DiagnoseUnavailableAlignedAllocation into DiagnoseUseOfDecl, and then adds call's in the few places they're needed. Most paths we care about already pass through there.

include/clang/Sema/Sema.h
5169

DiagnoseUnavailableAlignedAllocation.

Also, I think we can drop the IsDelete parameter and instead deduce it using FD.getDeclName().getCXXOverloadedOperator().

ahatanak updated this revision to Diff 150211.Jun 6 2018, 4:04 PM
ahatanak marked an inline comment as done.
ahatanak retitled this revision from [Sema] Diagnose unavailable aligned deallocation functions called from deleting destructors. to [Sema] Produce diagnostics when unavailable aligned allocation/deallocation functions are called.

Sink diagnoseUnavailableAlignedAllocation into DiagnoseUseOfDecl and add calls to it in a few other places so that diagnostics are produced for calls to aligned operator and builtin operator new/delete in addition to aligned deallocation functions called from deleting destructors.

I had to make changes to two test cases (dr2xx.cpp and call-host-fn-from-device.cu) that have nothing to do with aligned allocation/deallocation functions. The warning in dr2xx.cpp seems correct to me judging from the comment left few lines above it ("We're also missing the -Wused-but-marked-unused "). I'm not sure about the diagnostic in call-host-fn-from-device.cu. The original comment says the sized delete function annotated with __device__ is called, but it seems that the non-sized __host__ function is being called, in which case I think the diagnostic is correct.

ahatanak added a subscriber: jlebar.

ping.

@jlebar, is the change I made to call-host-fn-from-device.cu correct?

@jlebar, is the change I made to call-host-fn-from-device.cu correct?

I don't think so -- that's a change in overloading behavior afaict.

The original comment said "call sized device delete even though host has preferable non-sized version", but it seems that the non-sized host version 'T::operator delete(void*)' is being called in the IR, not the sized device version of delete. Is that a bug in overload resolution?

This is the command I used:

"clang -cc1 -internal-isystem --std=c++11 -triple nvptx-unknown-unknown -fcuda-is-device -emit-llvm -o - test.cu"

$ cat test.cu

#define __constant__ __attribute__((constant))
#define __device__ __attribute__((device))
#define __host__ __attribute__((host))

typedef __SIZE_TYPE__ size_t;

struct T {
  __host__ void operator delete(void*);
  __device__ void operator delete(void*, size_t);
};

__host__ __device__ void class_specific_delete(T *t) {
  delete t;
}

I mean ToT clang (without my patch applied) seems to select the non-sized host version 'T::operator delete(void*)'.

I mean ToT clang (without my patch applied) seems to select the non-sized host version 'T::operator delete(void*)'.

OK, if this is just making an error out of something which previously silently didn't work (and should result in a compile error further down the line when we try to and can't resolve that function), then this is totally fine.

Yes, that is the case. It doesn't change overloading resolution, it is just producing a diagnostic.

jlebar edited reviewers, added: tra; removed: jlebar.Aug 9 2018, 1:46 PM

+tra in the hopes that perhaps he's comfortable reviewing this (sorry that I'm not).

tra added inline comments.Aug 9 2018, 2:08 PM
test/SemaCUDA/call-host-fn-from-device.cu
88 ↗(On Diff #150211)

The C++ magic is way above my paygrade, but as far as CUDA goes this is a regression, compared to what nvcc does. This code in NVCC produced a warning and clang should not error out at this point in time either as
it's not an error to call a host function from HD unless we use HD in a host function, and we would not know how it's used until later. I think the error should be postponed until codegen.

rsmith added inline comments.Aug 9 2018, 3:17 PM
lib/Sema/SemaExprCXX.cpp
3465

Are we also missing a MarkFunctionReferenced call here? (I don't think it matters much for the predefined new/delete, since they can't be inline or templated, but it's still wrong to not mark the function the builtin will call as referenced.)

test/CXX/drs/dr2xx.cpp
721–722 ↗(On Diff #150211)

These are -> This is

test/SemaCUDA/call-host-fn-from-device.cu
88 ↗(On Diff #150211)

We're in -fcuda-is-device mode, so IIUC it's correct to reject a call to a host function here (because __host__ __device__ is treated as basically meaning __device__ in that mode for the purpose of checking whether a call is valid), right?

However, the comment suggests that the intent was that this would instead call the device version. Did that actually previously happen (in which case this patch is somehow affecting overload resolution and should be fixed), or is the comment prior to this patch wrong and we were silently calling a host function from a device function (in which case this patch is fine, but we should add a FIXME here to select the device delete function if we think that's appropriate)?

rsmith added inline comments.Aug 9 2018, 3:20 PM
test/SemaCUDA/call-host-fn-from-device.cu
88 ↗(On Diff #150211)

OK, I see from prior review comments (that phab is helpfully hiding from view) that this is just adding a diagnostic and the overload resolution behavior is unchanged. So I think this change is correct. @tra, can you confirm? My testing shows that

__host__ void f(); __host__ __device__ void g() { f(); }

is accepted by default but rejected in -fcuda-is-device mode, which is consistent with the behavior after this patch is applied.

ahatanak updated this revision to Diff 160046.Aug 9 2018, 7:02 PM
ahatanak marked 2 inline comments as done.

Call MarkFunctionReferenced to mark the function __builtin_operator_new or __builtin_operator_delete will call as referenced.

tra added inline comments.Aug 13 2018, 11:21 AM
test/SemaCUDA/call-host-fn-from-device.cu
88 ↗(On Diff #150211)

"-fcuda-is-device" does not necessarily mean that the host device function will be used.

In the example above the error is indeed correct, as g() is considered to be externally visible and we will attempt to generate code for it, and we can't call f() on device.

However, if you make it static, there should be no error:

__host__ void f(); 
static __host__ __device__ void g() { f(); }

CUDA is somewhat weird when it comes to what's considered available and what is not.
If you want some of the gory details, see D12453 and https://goo.gl/EXnymm

@jlebar has details on how we handle the errors in cases like that:
https://github.com/llvm-mirror/clang/blob/master/lib/Sema/SemaCUDA.cpp#L590

@tra and @rsmith: Can we move forward and fix the incorrect cuda diagnostics in a separate patch?

tra added a comment.Aug 17 2018, 12:02 PM

@tra and @rsmith: Can we move forward and fix the incorrect cuda diagnostics in a separate patch?

Doing that in a separate patch is OK, provided that that patch will be committed along with this one.

It's a regression. There's a decent chance it breaks someone and this patch, if committed by itself, will end up being rolled back.

In D47757#1204561, @tra wrote:

It's a regression. There's a decent chance it breaks someone and this patch, if committed by itself, will end up being rolled back.

Is the regression you are referring to about the static function case? I don't see a difference between ToT clang and my patch in the diagnostics they produce when I compile the following code:

__host__ void f();
static __host__ __device__ void g() { f(); }
__host__ __device__ void g2() { g(); }

Both error out when -fcuda-is-device is provided. If I comment out the definition of g2, it compiles fine.

tra added a comment.Aug 17 2018, 1:07 PM
In D47757#1204561, @tra wrote:

It's a regression. There's a decent chance it breaks someone and this patch, if committed by itself, will end up being rolled back.

Is the regression you are referring to about the static function case? I don't see a difference between ToT clang and my patch in the diagnostics they produce when I compile the following code:

__host__ void f();
static __host__ __device__ void g() { f(); }
__host__ __device__ void g2() { g(); }

Both error out when -fcuda-is-device is provided. If I comment out the definition of g2, it compiles fine.

The example above *is* expected to produce the error on device side, bacause g2() is externally visible, uses g(), which in turn uses host-only f().

I'm talking about a case where g() {f()} is present in the source code, but will not be codegen'ed on device side.

The code below is expected to compile. Note that g2() is host-only.

__host__ void f(); 
static __host__ __device__ void g() { f(); }
__host__ void g2() { g(); }

The code you showed does compile with or without -fcuda-is-device after applying my patch.

tra added a comment.Aug 17 2018, 3:27 PM

Talked to @ahatanak over IRC. It appears that this patch may have exposed a preexisting bug.
Apparently delete t; in test/SemaCUDA/call-host-fn-from-device.cu does actually end up calling __host__ operator delete. It should've picked __device__ operator delete, but it does not, so reporting an error here appears to be correct.

It's visible in AST and the IR.

@rsmith -- the original change was done a while back in rL283830. I assume it worked at that time and wonder if it's a (possibly not-so-)recent regression.

tra added a comment.Aug 23 2018, 11:06 AM

I've confirmed that the patch does not break anything in our CUDA code, so it's good to go as far as CUDA is concerned.

I'll fix the exposed CUDA issue in a separate patch.

In D47757#1211276, @tra wrote:

I've confirmed that the patch does not break anything in our CUDA code, so it's good to go as far as CUDA is concerned.

Thanks. @rsmith, do you have any other comments about the patch?

ahatanak updated this revision to Diff 176214.Nov 30 2018, 3:19 PM

Rebase & ping.

I've reverted the changes I made to test/SemaCUDA/call-host-fn-from-device.cu since r342749 fixed the overload resolution bug.

rsmith accepted this revision.Dec 20 2018, 1:19 PM
rsmith added inline comments.
lib/Sema/SemaExpr.cpp
54 ↗(On Diff #176214)

Does this also need to be updated?

This revision is now accepted and ready to land.Dec 20 2018, 1:19 PM
ahatanak updated this revision to Diff 179195.Dec 20 2018, 4:11 PM
ahatanak marked 2 inline comments as done.

Check whether the declaration passed to Sema::CanUseDecl is an aligned allocation/deallocation function that is unavailable.

vsapsai accepted this revision.Dec 20 2018, 5:08 PM

The change itself looks correct. Cannot really tell if you need to make changes in other places. For that I rely on Richard's opinion.

This revision was automatically updated to reflect the committed changes.