diff options
author | Artem Belevich <tra@google.com> | 2018-09-21 17:29:33 +0000 |
---|---|---|
committer | Artem Belevich <tra@google.com> | 2018-09-21 17:29:33 +0000 |
commit | 6e87bc08f54e1d2e90ff9540b83d540cc9ec7ca2 (patch) | |
tree | 012808cb2426a3c65128d6d3eb539cfebbc0c425 /test/CodeGenCUDA | |
parent | 785350882397c919e4149b1347bdfe9639b4e404 (diff) | |
download | clang-6e87bc08f54e1d2e90ff9540b83d540cc9ec7ca2.tar.gz |
[CUDA] Ignore uncallable functions when we check for usual deallocators.
Previously clang considered function variants from both sides of
compilation and that resulted in picking up wrong deallocation function.
Differential Revision: https://reviews.llvm.org/D51808
git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@342749 91177308-0d34-0410-b5e6-96231b3b80d8
Diffstat (limited to 'test/CodeGenCUDA')
-rw-r--r-- | test/CodeGenCUDA/usual-deallocators.cu | 133 |
1 files changed, 133 insertions, 0 deletions
diff --git a/test/CodeGenCUDA/usual-deallocators.cu b/test/CodeGenCUDA/usual-deallocators.cu new file mode 100644 index 0000000000..2d97c8c9f7 --- /dev/null +++ b/test/CodeGenCUDA/usual-deallocators.cu @@ -0,0 +1,133 @@ +// RUN: %clang_cc1 %s --std=c++11 -triple nvptx-unknown-unknown -fcuda-is-device \ +// RUN: -emit-llvm -o - | FileCheck %s --check-prefixes=COMMON,DEVICE +// RUN: %clang_cc1 %s --std=c++11 -triple nvptx-unknown-unknown \ +// RUN: -emit-llvm -o - | FileCheck %s --check-prefixes=COMMON,HOST +// RUN: %clang_cc1 %s --std=c++17 -triple nvptx-unknown-unknown -fcuda-is-device \ +// RUN: -emit-llvm -o - | FileCheck %s --check-prefixes=COMMON,DEVICE +// RUN: %clang_cc1 %s --std=c++17 -triple nvptx-unknown-unknown \ +// RUN: -emit-llvm -o - | FileCheck %s --check-prefixes=COMMON,HOST + +#include "Inputs/cuda.h" +extern "C" __host__ void host_fn(); +extern "C" __device__ void dev_fn(); +extern "C" __host__ __device__ void hd_fn(); + +struct H1D1 { + __host__ void operator delete(void *) { host_fn(); }; + __device__ void operator delete(void *) { dev_fn(); }; +}; + +struct H1D2 { + __host__ void operator delete(void *) { host_fn(); }; + __device__ void operator delete(void *, __SIZE_TYPE__) { dev_fn(); }; +}; + +struct H2D1 { + __host__ void operator delete(void *, __SIZE_TYPE__) { host_fn(); }; + __device__ void operator delete(void *) { dev_fn(); }; +}; + +struct H2D2 { + __host__ void operator delete(void *, __SIZE_TYPE__) { host_fn(); }; + __device__ void operator delete(void *, __SIZE_TYPE__) { dev_fn(); }; +}; + +struct H1D1D2 { + __host__ void operator delete(void *) { host_fn(); }; + __device__ void operator delete(void *) { dev_fn(); }; + __device__ void operator delete(void *, __SIZE_TYPE__) { dev_fn(); }; +}; + +struct H1H2D1 { + __host__ void operator delete(void *) { host_fn(); }; + __host__ void operator delete(void *, __SIZE_TYPE__) { host_fn(); }; + __device__ void operator delete(void *) { dev_fn(); }; +}; + +struct H1H2D2 { + __host__ void operator delete(void *) { host_fn(); }; + __host__ void operator delete(void *, __SIZE_TYPE__) { host_fn(); }; + __device__ void operator delete(void *, __SIZE_TYPE__) { dev_fn(); }; +}; + +struct H1H2D1D2 { + __host__ void operator delete(void *) { host_fn(); }; + __host__ void operator delete(void *, __SIZE_TYPE__) { host_fn(); }; + __device__ void operator delete(void *) { dev_fn(); }; + __device__ void operator delete(void *, __SIZE_TYPE__) { dev_fn(); }; +}; + + +template <typename T> +__host__ __device__ void test_hd(void *p) { + T *t = (T *)p; + delete t; +} + +// Make sure we call the right variant of usual deallocator. +__host__ __device__ void tests_hd(void *t) { + // COMMON-LABEL: define linkonce_odr void @_Z7test_hdI4H1D1EvPv + // COMMON: call void @_ZN4H1D1dlEPv + test_hd<H1D1>(t); + // COMMON-LABEL: define linkonce_odr void @_Z7test_hdI4H1D2EvPv + // DEVICE: call void @_ZN4H1D2dlEPvj(i8* {{.*}}, i32 1) + // HOST: call void @_ZN4H1D2dlEPv(i8* {{.*}}) + test_hd<H1D2>(t); + // COMMON-LABEL: define linkonce_odr void @_Z7test_hdI4H2D1EvPv + // DEVICE: call void @_ZN4H2D1dlEPv(i8* {{.*}}) + // HOST: call void @_ZN4H2D1dlEPvj(i8* %3, i32 1) + test_hd<H2D1>(t); + // COMMON-LABEL: define linkonce_odr void @_Z7test_hdI4H2D2EvPv + // COMMON: call void @_ZN4H2D2dlEPvj(i8* {{.*}}, i32 1) + test_hd<H2D2>(t); + // COMMON-LABEL: define linkonce_odr void @_Z7test_hdI6H1D1D2EvPv + // COMMON: call void @_ZN6H1D1D2dlEPv(i8* %3) + test_hd<H1D1D2>(t); + // COMMON-LABEL: define linkonce_odr void @_Z7test_hdI6H1H2D1EvPv + // COMMON: call void @_ZN6H1H2D1dlEPv(i8* {{.*}}) + test_hd<H1H2D1>(t); + // COMMON-LABEL: define linkonce_odr void @_Z7test_hdI6H1H2D2EvPv + // DEVICE: call void @_ZN6H1H2D2dlEPvj(i8* {{.*}}, i32 1) + // HOST: call void @_ZN6H1H2D2dlEPv(i8* {{.*}}) + test_hd<H1H2D2>(t); + // COMMON-LABEL: define linkonce_odr void @_Z7test_hdI8H1H2D1D2EvPv + // COMMON: call void @_ZN8H1H2D1D2dlEPv(i8* {{.*}}) + test_hd<H1H2D1D2>(t); +} + +// Make sure we've picked deallocator for the correct side of compilation. + +// COMMON-LABEL: define linkonce_odr void @_ZN4H1D1dlEPv(i8*) +// DEVICE: call void @dev_fn() +// HOST: call void @host_fn() + +// DEVICE-LABEL: define linkonce_odr void @_ZN4H1D2dlEPvj(i8*, i32) +// DEVICE: call void @dev_fn() +// HOST-LABEL: define linkonce_odr void @_ZN4H1D2dlEPv(i8*) +// HOST: call void @host_fn() + +// DEVICE-LABEL: define linkonce_odr void @_ZN4H2D1dlEPv(i8*) +// DEVICE: call void @dev_fn() +// HOST-LABEL: define linkonce_odr void @_ZN4H2D1dlEPvj(i8*, i32) +// HOST: call void @host_fn() + +// COMMON-LABEL: define linkonce_odr void @_ZN4H2D2dlEPvj(i8*, i32) +// DEVICE: call void @dev_fn() +// HOST: call void @host_fn() + +// COMMON-LABEL: define linkonce_odr void @_ZN6H1D1D2dlEPv(i8*) +// DEVICE: call void @dev_fn() +// HOST: call void @host_fn() + +// COMMON-LABEL: define linkonce_odr void @_ZN6H1H2D1dlEPv(i8*) +// DEVICE: call void @dev_fn() +// HOST: call void @host_fn() + +// DEVICE-LABEL: define linkonce_odr void @_ZN6H1H2D2dlEPvj(i8*, i32) +// DEVICE: call void @dev_fn() +// HOST-LABEL: define linkonce_odr void @_ZN6H1H2D2dlEPv(i8*) +// HOST: call void @host_fn() + +// COMMON-LABEL: define linkonce_odr void @_ZN8H1H2D1D2dlEPv(i8*) +// DEVICE: call void @dev_fn() +// HOST: call void @host_fn() |