-
Notifications
You must be signed in to change notification settings - Fork 10.8k
Commit
This commit does not belong to any branch on this repository, and may belong to a fork outside of the repository.
[CUDA][HIP] Fix overloading resolution of delete operator
Currently clang does not consider host/device preference when resolving delete operator in the file scope, which causes device operator delete selected for class member initialization. Reviewed by: Artem Belevich Differential Revision: https://reviews.llvm.org/D156795
- Loading branch information
Showing
3 changed files
with
113 additions
and
2 deletions.
There are no files selected for viewing
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Original file line number | Diff line number | Diff line change |
---|---|---|
@@ -0,0 +1,73 @@ | ||
// RUN: %clang_cc1 -triple x86_64-pc-linux-gnu -emit-llvm -fexceptions \ | ||
// RUN: -o - -x hip %s | FileCheck %s | ||
|
||
#include "Inputs/cuda.h" | ||
|
||
int* hvar; | ||
__device__ int* dvar; | ||
|
||
// CHECK-LABEL: define {{.*}}@_Znwm | ||
// CHECK: load ptr, ptr @hvar | ||
void* operator new(unsigned long size) { | ||
return hvar; | ||
} | ||
// CHECK-LABEL: define {{.*}}@_ZdlPv | ||
// CHECK: store ptr inttoptr (i64 1 to ptr), ptr @hvar | ||
void operator delete(void *p) { | ||
hvar = (int*)1; | ||
} | ||
|
||
__device__ void* operator new(unsigned long size) { | ||
return dvar; | ||
} | ||
|
||
__device__ void operator delete(void *p) { | ||
dvar = (int*)11; | ||
} | ||
|
||
class A { | ||
int x; | ||
public: | ||
A(){ | ||
x = 123; | ||
} | ||
}; | ||
|
||
template<class T> | ||
class shared_ptr { | ||
int id; | ||
T *ptr; | ||
public: | ||
shared_ptr(T *p) { | ||
id = 2; | ||
ptr = p; | ||
} | ||
}; | ||
|
||
// The constructor of B calls the delete operator to clean up | ||
// the memory allocated by the new operator when exceptions happen. | ||
// Make sure the host delete operator is used on host side. | ||
// | ||
// No need to do similar checks on the device side since it does | ||
// not support exception. | ||
|
||
// CHECK-LABEL: define {{.*}}@main | ||
// CHECK: call void @_ZN1BC1Ev | ||
|
||
// CHECK-LABEL: define {{.*}}@_ZN1BC1Ev | ||
// CHECK: call void @_ZN1BC2Ev | ||
|
||
// CHECK-LABEL: define {{.*}}@_ZN1BC2Ev | ||
// CHECK: call {{.*}}@_Znwm | ||
// CHECK: invoke void @_ZN1AC1Ev | ||
// CHECK: call void @_ZN10shared_ptrI1AEC1EPS0_ | ||
// CHECK: cleanup | ||
// CHECK: call void @_ZdlPv | ||
|
||
struct B{ | ||
shared_ptr<A> pa{new A}; | ||
}; | ||
|
||
int main() { | ||
B b; | ||
} |
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Original file line number | Diff line number | Diff line change |
---|---|---|
@@ -0,0 +1,38 @@ | ||
// RUN: %clang_cc1 -fsyntax-only -verify -fexceptions %s | ||
// expected-no-diagnostics | ||
|
||
#include "Inputs/cuda.h" | ||
|
||
__device__ void operator delete(void *p) {} | ||
|
||
class A { | ||
int x; | ||
public: | ||
A() { | ||
x = 123; | ||
} | ||
}; | ||
|
||
template<class T> | ||
class shared_ptr { | ||
T *ptr; | ||
public: | ||
shared_ptr(T *p) { | ||
ptr = p; | ||
} | ||
}; | ||
|
||
// The constructor of B calls the delete operator to clean up | ||
// the memory allocated by the new operator when exceptions happen. | ||
// Make sure that there are no diagnostics due to the device delete | ||
// operator is used. | ||
// | ||
// No need to do similar checks on the device side since it does | ||
// not support exception. | ||
struct B{ | ||
shared_ptr<A> pa{new A}; | ||
}; | ||
|
||
int main() { | ||
B b; | ||
} |