maxjhandsome
/
llvm-project
forked from OSchip/llvm-project
1
0
Fork 0
Code Issues Pull Requests Packages Releases Wiki Activity
llvm-project/clang/test/OpenMP/target_parallel_debug_codeg...

Ignoring revisions in .git-blame-ignore-revs. Click here to bypass and see the normal blame view.

138 lines
8.7 KiB
C++
Raw Normal View History

[OPENMP, NVPTX] Initial support for L2 parallelism in SPMD mode. Added initial support for L2 parallelism in SPMD mode. Note, though, that the orphaned parallel directives are not currently supported in SPMD mode. llvm-svn: 332016
2018-05-11 02:32:08 +08:00
// RUN: %clang_cc1 -DCK1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -fopenmp-cuda-mode -emit-llvm-bc %s -o %t-ppc-host.bc -fopenmp-version=45
// RUN: %clang_cc1 -DCK1 -verify -fopenmp -x c++ -triple nvptx64-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -fopenmp-cuda-mode -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - -debug-info-kind=limited -fopenmp-version=45 | FileCheck %s
[OPENMP][DEBUG] Set proper address space info if required by target. Arguments, passed to the outlined function, must have correct address space info for proper Debug info support. Patch sets global address space for arguments that are mapped and passed by reference. Also, cuda-gdb does not handle reference types correctly, so reference arguments are represented as pointers. llvm-svn: 310387
2017-08-09 02:04:06 +08:00
// expected-no-diagnostics
[OPENMP]Fix handling of the globals during compilation for the device. Fixed lookup for the target regions in unused virtual functions + fixed processing of the global variables not marked as declare target but emitted during debug info emission. llvm-svn: 346343
2018-11-08 03:11:14 +08:00
template <unsigned *ddd>
struct S {
static int a;
};
extern unsigned aaa;
template<> int S<&aaa>::a;
template struct S<&aaa>;
// CHECK-NOT: @aaa
[OPENMP][DEBUG] Set proper address space info if required by target. Arguments, passed to the outlined function, must have correct address space info for proper Debug info support. Patch sets global address space for arguments that are mapped and passed by reference. Also, cuda-gdb does not handle reference types correctly, so reference arguments are represented as pointers. llvm-svn: 310387
2017-08-09 02:04:06 +08:00
int main() {
/* int(*b)[a]; */
/* int *(**c)[a]; */
[OPENMP] Fix types for the target specific parameters in debug mode. Used incorrect types for target specific parameters in debug mode, should use original pointers rather than the pointee types. llvm-svn: 313186
2017-09-14 04:20:59 +08:00
bool bb;
[OPENMP][DEBUG] Set proper address space info if required by target. Arguments, passed to the outlined function, must have correct address space info for proper Debug info support. Patch sets global address space for arguments that are mapped and passed by reference. Also, cuda-gdb does not handle reference types correctly, so reference arguments are represented as pointers. llvm-svn: 310387
2017-08-09 02:04:06 +08:00
int a;
int b[10][10];
int c[10][10][10];
[OPENMP] Fix types for the target specific parameters in debug mode. Used incorrect types for target specific parameters in debug mode, should use original pointers rather than the pointee types. llvm-svn: 313186
2017-09-14 04:20:59 +08:00
#pragma omp target parallel firstprivate(a, b) map(tofrom \
: c) map(tofrom \
[OPENMP, NVPTX] Initial support for L2 parallelism in SPMD mode. Added initial support for L2 parallelism in SPMD mode. Note, though, that the orphaned parallel directives are not currently supported in SPMD mode. llvm-svn: 332016
2018-05-11 02:32:08 +08:00
: bb) if (target:a)
[OPENMP][DEBUG] Set proper address space info if required by target. Arguments, passed to the outlined function, must have correct address space info for proper Debug info support. Patch sets global address space for arguments that are mapped and passed by reference. Also, cuda-gdb does not handle reference types correctly, so reference arguments are represented as pointers. llvm-svn: 310387
2017-08-09 02:04:06 +08:00
{
int &f = c[1][1][1];
int &g = a;
int &h = b[1][1];
int d = 15;
a = 5;
b[0][a] = 10;
c[0][0][a] = 11;
b[0][a] = c[0][0][a];
[OPENMP] Fix types for the target specific parameters in debug mode. Used incorrect types for target specific parameters in debug mode, should use original pointers rather than the pointee types. llvm-svn: 313186
2017-09-14 04:20:59 +08:00
bb |= b[0][a];
[OPENMP][DEBUG] Set proper address space info if required by target. Arguments, passed to the outlined function, must have correct address space info for proper Debug info support. Patch sets global address space for arguments that are mapped and passed by reference. Also, cuda-gdb does not handle reference types correctly, so reference arguments are represented as pointers. llvm-svn: 310387
2017-08-09 02:04:06 +08:00
}
[OPENMP] Fix types for the target specific parameters in debug mode. Used incorrect types for target specific parameters in debug mode, should use original pointers rather than the pointee types. llvm-svn: 313186
2017-09-14 04:20:59 +08:00
#pragma omp target parallel firstprivate(a) map(tofrom \
: c, b) map(to \
: bb)
[OPENMP][DEBUG] Set proper address space info if required by target. Arguments, passed to the outlined function, must have correct address space info for proper Debug info support. Patch sets global address space for arguments that are mapped and passed by reference. Also, cuda-gdb does not handle reference types correctly, so reference arguments are represented as pointers. llvm-svn: 310387
2017-08-09 02:04:06 +08:00
{
int &f = c[1][1][1];
int &g = a;
int &h = b[1][1];
int d = 15;
a = 5;
b[0][a] = 10;
c[0][0][a] = 11;
b[0][a] = c[0][0][a];
[OPENMP] Fix types for the target specific parameters in debug mode. Used incorrect types for target specific parameters in debug mode, should use original pointers rather than the pointee types. llvm-svn: 313186
2017-09-14 04:20:59 +08:00
d = bb;
[OPENMP][DEBUG] Set proper address space info if required by target. Arguments, passed to the outlined function, must have correct address space info for proper Debug info support. Patch sets global address space for arguments that are mapped and passed by reference. Also, cuda-gdb does not handle reference types correctly, so reference arguments are represented as pointers. llvm-svn: 310387
2017-08-09 02:04:06 +08:00
}
[OPENMP] Fix types for the target specific parameters in debug mode. Used incorrect types for target specific parameters in debug mode, should use original pointers rather than the pointee types. llvm-svn: 313186
2017-09-14 04:20:59 +08:00
#pragma omp target parallel map(tofrom \
: a, c, b) map(from \
: bb)
[OPENMP][DEBUG] Set proper address space info if required by target. Arguments, passed to the outlined function, must have correct address space info for proper Debug info support. Patch sets global address space for arguments that are mapped and passed by reference. Also, cuda-gdb does not handle reference types correctly, so reference arguments are represented as pointers. llvm-svn: 310387
2017-08-09 02:04:06 +08:00
{
int &f = c[1][1][1];
int &g = a;
int &h = b[1][1];
int d = 15;
a = 5;
b[0][a] = 10;
c[0][0][a] = 11;
b[0][a] = c[0][0][a];
[OPENMP] Fix types for the target specific parameters in debug mode. Used incorrect types for target specific parameters in debug mode, should use original pointers rather than the pointee types. llvm-svn: 313186
2017-09-14 04:20:59 +08:00
bb = b[0][a];
[OPENMP][DEBUG] Set proper address space info if required by target. Arguments, passed to the outlined function, must have correct address space info for proper Debug info support. Patch sets global address space for arguments that are mapped and passed by reference. Also, cuda-gdb does not handle reference types correctly, so reference arguments are represented as pointers. llvm-svn: 310387
2017-08-09 02:04:06 +08:00
}
return 0;
}
[OPENMP, NVPTX] Initial support for L2 parallelism in SPMD mode. Added initial support for L2 parallelism in SPMD mode. Note, though, that the orphaned parallel directives are not currently supported in SPMD mode. llvm-svn: 332016
2018-05-11 02:32:08 +08:00
// CHECK: define internal void @__omp_offloading{{[^(]+}}([10 x [10 x [10 x i32]]] addrspace(1)* {{[^,]+}}, i32 {{[^,]+}}, [10 x [10 x i32]]* {{[^,]+}}, i8 addrspace(1)* noalias{{[^,]+}})
[OPENMP][DEBUG] Set proper address space info if required by target. Arguments, passed to the outlined function, must have correct address space info for proper Debug info support. Patch sets global address space for arguments that are mapped and passed by reference. Also, cuda-gdb does not handle reference types correctly, so reference arguments are represented as pointers. llvm-svn: 310387
2017-08-09 02:04:06 +08:00
// CHECK: addrspacecast [10 x [10 x [10 x i32]]] addrspace(1)* %{{.+}} to [10 x [10 x [10 x i32]]]*
[OPENMP] Fix types for the target specific parameters in debug mode. Used incorrect types for target specific parameters in debug mode, should use original pointers rather than the pointee types. llvm-svn: 313186
2017-09-14 04:20:59 +08:00
// CHECK: call void [[NONDEBUG_WRAPPER:.+]](i32* {{[^,]+}}, i32* {{[^,]+}}, [10 x [10 x [10 x i32]]]* {{[^,]+}}, i64 {{[^,]+}}, [10 x [10 x i32]]* {{[^,]+}}, i8* {{[^)]+}})
[OPENMP][DEBUG] Set proper address space info if required by target. Arguments, passed to the outlined function, must have correct address space info for proper Debug info support. Patch sets global address space for arguments that are mapped and passed by reference. Also, cuda-gdb does not handle reference types correctly, so reference arguments are represented as pointers. llvm-svn: 310387
2017-08-09 02:04:06 +08:00
[OPENMP] Fix types for the target specific parameters in debug mode. Used incorrect types for target specific parameters in debug mode, should use original pointers rather than the pointee types. llvm-svn: 313186
2017-09-14 04:20:59 +08:00
// CHECK: define internal void [[DEBUG_PARALLEL:@.+]](i32* {{[^,]+}}, i32* {{[^,]+}}, [10 x [10 x [10 x i32]]] addrspace(1)* noalias{{[^,]+}}, i32 {{[^,]+}}, [10 x [10 x i32]]* noalias{{[^,]+}}, i8 addrspace(1)* noalias{{[^)]+}})
[OPENMP][DEBUG] Set proper address space info if required by target. Arguments, passed to the outlined function, must have correct address space info for proper Debug info support. Patch sets global address space for arguments that are mapped and passed by reference. Also, cuda-gdb does not handle reference types correctly, so reference arguments are represented as pointers. llvm-svn: 310387
2017-08-09 02:04:06 +08:00
// CHECK: addrspacecast [10 x [10 x [10 x i32]]] addrspace(1)* %{{.+}} to [10 x [10 x [10 x i32]]]*
[OPENMP] Fix types for the target specific parameters in debug mode. Used incorrect types for target specific parameters in debug mode, should use original pointers rather than the pointee types. llvm-svn: 313186
2017-09-14 04:20:59 +08:00
// CHECK: define internal void [[NONDEBUG_WRAPPER]](i32* {{[^,]+}}, i32* {{[^,]+}}, [10 x [10 x [10 x i32]]]* dereferenceable{{[^,]+}}, i64 {{[^,]+}}, [10 x [10 x i32]]* dereferenceable{{[^,]+}}, i8* dereferenceable{{[^)]+}})
[OPENMP][DEBUG] Set proper address space info if required by target. Arguments, passed to the outlined function, must have correct address space info for proper Debug info support. Patch sets global address space for arguments that are mapped and passed by reference. Also, cuda-gdb does not handle reference types correctly, so reference arguments are represented as pointers. llvm-svn: 310387
2017-08-09 02:04:06 +08:00
// CHECK: addrspacecast [10 x [10 x [10 x i32]]]* %{{.+}} to [10 x [10 x [10 x i32]]] addrspace(1)*
[OPENMP] Fix types for the target specific parameters in debug mode. Used incorrect types for target specific parameters in debug mode, should use original pointers rather than the pointee types. llvm-svn: 313186
2017-09-14 04:20:59 +08:00
// CHECK: call void [[DEBUG_PARALLEL]](i32* {{[^,]+}}, i32* {{[^,]+}}, [10 x [10 x [10 x i32]]] addrspace(1)* {{[^,]+}}, i32 {{[^,]+}}, [10 x [10 x i32]]* {{[^,]+}}, i8 addrspace(1)* {{[^)]+}})
[OPENMP][DEBUG] Set proper address space info if required by target. Arguments, passed to the outlined function, must have correct address space info for proper Debug info support. Patch sets global address space for arguments that are mapped and passed by reference. Also, cuda-gdb does not handle reference types correctly, so reference arguments are represented as pointers. llvm-svn: 310387
2017-08-09 02:04:06 +08:00
[OPENMP, NVPTX] Fix linkage of the global entries. The linkage of the global entries must be weak to enable support of redefinition of the same target regions in multiple compilation units. llvm-svn: 331768
2018-05-08 22:16:57 +08:00
// CHECK: define weak void @__omp_offloading_{{[^(]+}}([10 x [10 x [10 x i32]]]* dereferenceable{{[^,]+}}, i64 {{[^,]+}}, [10 x [10 x i32]]* dereferenceable{{[^,]+}}, i8* dereferenceable{{[^)]+}})
[OPENMP][DEBUG] Set proper address space info if required by target. Arguments, passed to the outlined function, must have correct address space info for proper Debug info support. Patch sets global address space for arguments that are mapped and passed by reference. Also, cuda-gdb does not handle reference types correctly, so reference arguments are represented as pointers. llvm-svn: 310387
2017-08-09 02:04:06 +08:00
// CHECK: addrspacecast [10 x [10 x [10 x i32]]]* %{{.+}} to [10 x [10 x [10 x i32]]] addrspace(1)*
[OPENMP] Fix types for the target specific parameters in debug mode. Used incorrect types for target specific parameters in debug mode, should use original pointers rather than the pointee types. llvm-svn: 313186
2017-09-14 04:20:59 +08:00
// CHECK: call void @__omp_offloading_{{[^(]+}}([10 x [10 x [10 x i32]]] addrspace(1)* {{[^,]+}}, i32 {{[^,]+}}, [10 x [10 x i32]]* {{[^,]+}}, i8 addrspace(1)* {{[^)]+}})
[OPENMP][DEBUG] Set proper address space info if required by target. Arguments, passed to the outlined function, must have correct address space info for proper Debug info support. Patch sets global address space for arguments that are mapped and passed by reference. Also, cuda-gdb does not handle reference types correctly, so reference arguments are represented as pointers. llvm-svn: 310387
2017-08-09 02:04:06 +08:00
[OPENMP] Fix types for the target specific parameters in debug mode. Used incorrect types for target specific parameters in debug mode, should use original pointers rather than the pointee types. llvm-svn: 313186
2017-09-14 04:20:59 +08:00
// CHECK: define internal void @__omp_offloading_{{[^(]+}}([10 x [10 x [10 x i32]]] addrspace(1)* noalias{{[^,]+}}, i32 {{[^,]+}}, [10 x [10 x i32]] addrspace(1)* noalias{{[^,]+}}, i8 addrspace(1)* noalias{{[^)]+}})
[OPENMP][DEBUG] Set proper address space info if required by target. Arguments, passed to the outlined function, must have correct address space info for proper Debug info support. Patch sets global address space for arguments that are mapped and passed by reference. Also, cuda-gdb does not handle reference types correctly, so reference arguments are represented as pointers. llvm-svn: 310387
2017-08-09 02:04:06 +08:00
// CHECK: addrspacecast [10 x [10 x [10 x i32]]] addrspace(1)* %{{.+}} to [10 x [10 x [10 x i32]]]*
// CHECK: addrspacecast [10 x [10 x i32]] addrspace(1)* %{{.+}} to [10 x [10 x i32]]*
[OPENMP] Fix types for the target specific parameters in debug mode. Used incorrect types for target specific parameters in debug mode, should use original pointers rather than the pointee types. llvm-svn: 313186
2017-09-14 04:20:59 +08:00
// CHECK: call void [[NONDEBUG_WRAPPER:.+]](i32* {{[^,]+}}, i32* {{[^,]+}}, [10 x [10 x [10 x i32]]]* {{[^,]+}}, i64 {{[^,]+}}, [10 x [10 x i32]]* {{[^,]+}}, i8* {{[^)]+}})
[OPENMP][DEBUG] Set proper address space info if required by target. Arguments, passed to the outlined function, must have correct address space info for proper Debug info support. Patch sets global address space for arguments that are mapped and passed by reference. Also, cuda-gdb does not handle reference types correctly, so reference arguments are represented as pointers. llvm-svn: 310387
2017-08-09 02:04:06 +08:00
[OPENMP] Fix types for the target specific parameters in debug mode. Used incorrect types for target specific parameters in debug mode, should use original pointers rather than the pointee types. llvm-svn: 313186
2017-09-14 04:20:59 +08:00
// CHECK: define internal void [[DEBUG_PARALLEL:@.+]](i32* {{[^,]+}}, i32* {{[^,]+}}, [10 x [10 x [10 x i32]]] addrspace(1)* noalias{{[^,]+}}, i32 {{[^,]+}}, [10 x [10 x i32]] addrspace(1)* noalias{{[^,]+}}, i8 addrspace(1)* {{[^)]+}})
[OPENMP][DEBUG] Set proper address space info if required by target. Arguments, passed to the outlined function, must have correct address space info for proper Debug info support. Patch sets global address space for arguments that are mapped and passed by reference. Also, cuda-gdb does not handle reference types correctly, so reference arguments are represented as pointers. llvm-svn: 310387
2017-08-09 02:04:06 +08:00
// CHECK: addrspacecast [10 x [10 x [10 x i32]]] addrspace(1)* %{{.+}} to [10 x [10 x [10 x i32]]]*
// CHECK: addrspacecast [10 x [10 x i32]] addrspace(1)* %{{.+}} to [10 x [10 x i32]]*
[OPENMP] Fix types for the target specific parameters in debug mode. Used incorrect types for target specific parameters in debug mode, should use original pointers rather than the pointee types. llvm-svn: 313186
2017-09-14 04:20:59 +08:00
// CHECK: define internal void [[NONDEBUG_WRAPPER]](i32* {{[^,]+}}, i32* {{[^,]+}}, [10 x [10 x [10 x i32]]]* dereferenceable{{[^,]+}}, i64 {{[^,]+}}, [10 x [10 x i32]]* dereferenceable{{[^,]+}}, i8* dereferenceable{{[^)]+}})
[OPENMP][DEBUG] Set proper address space info if required by target. Arguments, passed to the outlined function, must have correct address space info for proper Debug info support. Patch sets global address space for arguments that are mapped and passed by reference. Also, cuda-gdb does not handle reference types correctly, so reference arguments are represented as pointers. llvm-svn: 310387
2017-08-09 02:04:06 +08:00
// CHECK: addrspacecast [10 x [10 x [10 x i32]]]* %{{.+}} to [10 x [10 x [10 x i32]]] addrspace(1)*
// CHECK: addrspacecast [10 x [10 x i32]]* %{{.+}} to [10 x [10 x i32]] addrspace(1)*
[OPENMP] Fix types for the target specific parameters in debug mode. Used incorrect types for target specific parameters in debug mode, should use original pointers rather than the pointee types. llvm-svn: 313186
2017-09-14 04:20:59 +08:00
// CHECK: call void [[DEBUG_PARALLEL]](i32* {{[^,]+}}, i32* {{[^,]+}}, [10 x [10 x [10 x i32]]] addrspace(1)* {{[^,]+}}, i32 {{[^,]+}}, [10 x [10 x i32]] addrspace(1)* {{[^,]+}}, i8 addrspace(1)* {{[^)]+}})
[OPENMP][DEBUG] Set proper address space info if required by target. Arguments, passed to the outlined function, must have correct address space info for proper Debug info support. Patch sets global address space for arguments that are mapped and passed by reference. Also, cuda-gdb does not handle reference types correctly, so reference arguments are represented as pointers. llvm-svn: 310387
2017-08-09 02:04:06 +08:00
[OPENMP, NVPTX] Fix linkage of the global entries. The linkage of the global entries must be weak to enable support of redefinition of the same target regions in multiple compilation units. llvm-svn: 331768
2018-05-08 22:16:57 +08:00
// CHECK: define weak void @__omp_offloading_{{[^(]+}}([10 x [10 x [10 x i32]]]* dereferenceable{{[^,]+}}, i64 {{[^,]+}}, [10 x [10 x i32]]* dereferenceable{{[^,]+}}, i8* dereferenceable{{[^)]+}})
[OPENMP][DEBUG] Set proper address space info if required by target. Arguments, passed to the outlined function, must have correct address space info for proper Debug info support. Patch sets global address space for arguments that are mapped and passed by reference. Also, cuda-gdb does not handle reference types correctly, so reference arguments are represented as pointers. llvm-svn: 310387
2017-08-09 02:04:06 +08:00
// CHECK: addrspacecast [10 x [10 x [10 x i32]]]* %{{.+}} to [10 x [10 x [10 x i32]]] addrspace(1)*
// CHECK: addrspacecast [10 x [10 x i32]]* %{{.+}} to [10 x [10 x i32]] addrspace(1)*
[OPENMP] Fix types for the target specific parameters in debug mode. Used incorrect types for target specific parameters in debug mode, should use original pointers rather than the pointee types. llvm-svn: 313186
2017-09-14 04:20:59 +08:00
// CHECK: call void @__omp_offloading_{{[^(]+}}([10 x [10 x [10 x i32]]] addrspace(1)* {{[^,]+}}, i32 {{[^,]+}}, [10 x [10 x i32]] addrspace(1)* {{[^,]+}}, i8 addrspace(1)* {{[^)]+}})
[OPENMP][DEBUG] Set proper address space info if required by target. Arguments, passed to the outlined function, must have correct address space info for proper Debug info support. Patch sets global address space for arguments that are mapped and passed by reference. Also, cuda-gdb does not handle reference types correctly, so reference arguments are represented as pointers. llvm-svn: 310387
2017-08-09 02:04:06 +08:00
[OPENMP] Fix types for the target specific parameters in debug mode. Used incorrect types for target specific parameters in debug mode, should use original pointers rather than the pointee types. llvm-svn: 313186
2017-09-14 04:20:59 +08:00
// CHECK: define internal void @__omp_offloading_{{[^(]+}}([10 x [10 x [10 x i32]]] addrspace(1)* noalias{{[^,]+}}, i32 addrspace(1)* noalias{{[^,]+}}, [10 x [10 x i32]] addrspace(1)* noalias{{[^,]+}}, i8 addrspace(1)* noalias{{[^)]+}})
[OPENMP][DEBUG] Set proper address space info if required by target. Arguments, passed to the outlined function, must have correct address space info for proper Debug info support. Patch sets global address space for arguments that are mapped and passed by reference. Also, cuda-gdb does not handle reference types correctly, so reference arguments are represented as pointers. llvm-svn: 310387
2017-08-09 02:04:06 +08:00
// CHECK: addrspacecast [10 x [10 x [10 x i32]]] addrspace(1)* %{{.+}} to [10 x [10 x [10 x i32]]]*
// CHECK: addrspacecast i32 addrspace(1)* %{{.+}} to i32*
// CHECK: addrspacecast [10 x [10 x i32]] addrspace(1)* %{{.+}} to [10 x [10 x i32]]*
[OPENMP] Fix debug info for outlined functions in NVPTX + add more tests. Fixed name of emitted outlined functions in NVPTX target + extra tests for the debug info. llvm-svn: 322022
2018-01-09 04:09:47 +08:00
// CHECK: call void @[[NONDEBUG_WRAPPER:.+]](i32* {{[^,]+}}, i32* {{[^,]+}}, [10 x [10 x [10 x i32]]]* {{[^,]+}}, i32* {{[^,]+}}, [10 x [10 x i32]]* {{[^,]+}}, i8* {{[^)]+}})
[OPENMP][DEBUG] Set proper address space info if required by target. Arguments, passed to the outlined function, must have correct address space info for proper Debug info support. Patch sets global address space for arguments that are mapped and passed by reference. Also, cuda-gdb does not handle reference types correctly, so reference arguments are represented as pointers. llvm-svn: 310387
2017-08-09 02:04:06 +08:00
[OPENMP] Fix debug info for outlined functions in NVPTX + add more tests. Fixed name of emitted outlined functions in NVPTX target + extra tests for the debug info. llvm-svn: 322022
2018-01-09 04:09:47 +08:00
// CHECK: define internal void @[[DEBUG_PARALLEL:.+]](i32* {{[^,]+}}, i32* {{[^,]+}}, [10 x [10 x [10 x i32]]] addrspace(1)* noalias{{[^,]+}}, i32 addrspace(1)* noalias{{[^,]+}}, [10 x [10 x i32]] addrspace(1)* noalias{{[^,]+}}, i8 addrspace(1)* noalias{{[^)]+}})
[OPENMP][DEBUG] Set proper address space info if required by target. Arguments, passed to the outlined function, must have correct address space info for proper Debug info support. Patch sets global address space for arguments that are mapped and passed by reference. Also, cuda-gdb does not handle reference types correctly, so reference arguments are represented as pointers. llvm-svn: 310387
2017-08-09 02:04:06 +08:00
// CHECK: addrspacecast [10 x [10 x [10 x i32]]] addrspace(1)* %{{.+}} to [10 x [10 x [10 x i32]]]*
// CHECK: addrspacecast i32 addrspace(1)* %{{.+}} to i32*
// CHECK: addrspacecast [10 x [10 x i32]] addrspace(1)* %{{.+}} to [10 x [10 x i32]]*
[OPENMP] Fix debug info for outlined functions in NVPTX + add more tests. Fixed name of emitted outlined functions in NVPTX target + extra tests for the debug info. llvm-svn: 322022
2018-01-09 04:09:47 +08:00
// CHECK: define internal void @[[NONDEBUG_WRAPPER]](i32* {{[^,]+}}, i32* {{[^,]+}}, [10 x [10 x [10 x i32]]]* dereferenceable{{[^,]+}}, i32* dereferenceable{{[^,]+}}, [10 x [10 x i32]]* dereferenceable{{[^,]+}}, i8* dereferenceable{{[^)]+}})
[OPENMP] Emit non-debug version of outlined functions with original name. If the host code is compiled with the debug info, while the target without, there is a problem that the compiler is unable to find the debug wrapper. Patch fixes this problem by emitting special name for the debug version of the code. llvm-svn: 310511
2017-08-10 03:38:53 +08:00
// CHECK: addrspacecast [10 x [10 x [10 x i32]]]* %{{.+}} to [10 x [10 x [10 x i32]]] addrspace(1)*
// CHECK: addrspacecast i32* %{{.+}} to i32 addrspace(1)*
// CHECK: addrspacecast [10 x [10 x i32]]* %{{.+}} to [10 x [10 x i32]] addrspace(1)*
[OPENMP] Fix debug info for outlined functions in NVPTX + add more tests. Fixed name of emitted outlined functions in NVPTX target + extra tests for the debug info. llvm-svn: 322022
2018-01-09 04:09:47 +08:00
// CHECK: call void @[[DEBUG_PARALLEL]](i32* {{[^,]+}}, i32* {{[^,]+}}, [10 x [10 x [10 x i32]]] addrspace(1)* {{[^,]+}}, i32 addrspace(1)* {{[^,]+}}, [10 x [10 x i32]] addrspace(1)* {{[^,]+}}, i8 addrspace(1)* {{[^)]+}})
[OPENMP] Emit non-debug version of outlined functions with original name. If the host code is compiled with the debug info, while the target without, there is a problem that the compiler is unable to find the debug wrapper. Patch fixes this problem by emitting special name for the debug version of the code. llvm-svn: 310511
2017-08-10 03:38:53 +08:00
[OPENMP, NVPTX] Fix linkage of the global entries. The linkage of the global entries must be weak to enable support of redefinition of the same target regions in multiple compilation units. llvm-svn: 331768
2018-05-08 22:16:57 +08:00
// CHECK: define weak void @__omp_offloading_{{[^(]+}}([10 x [10 x [10 x i32]]]* dereferenceable{{[^,]+}}, i32* dereferenceable{{[^,]+}}, [10 x [10 x i32]]* dereferenceable{{[^,]+}}, i8* dereferenceable{{[^)]+}})
[OPENMP] Emit non-debug version of outlined functions with original name. If the host code is compiled with the debug info, while the target without, there is a problem that the compiler is unable to find the debug wrapper. Patch fixes this problem by emitting special name for the debug version of the code. llvm-svn: 310511
2017-08-10 03:38:53 +08:00
// CHECK: addrspacecast [10 x [10 x [10 x i32]]]* %{{.+}} to [10 x [10 x [10 x i32]]] addrspace(1)*
// CHECK: addrspacecast i32* %{{.+}} to i32 addrspace(1)*
// CHECK: addrspacecast [10 x [10 x i32]]* %{{.+}} to [10 x [10 x i32]] addrspace(1)*
[OPENMP] Fix types for the target specific parameters in debug mode. Used incorrect types for target specific parameters in debug mode, should use original pointers rather than the pointee types. llvm-svn: 313186
2017-09-14 04:20:59 +08:00
// CHECK: call void @__omp_offloading_{{[^(]+}}([10 x [10 x [10 x i32]]] addrspace(1)* {{[^,]+}}, i32 addrspace(1)* {{[^,]+}}, [10 x [10 x i32]] addrspace(1)* {{[^,]+}}, i8 addrspace(1)* {{[^)]+}})
[OPENMP] Emit non-debug version of outlined functions with original name. If the host code is compiled with the debug info, while the target without, there is a problem that the compiler is unable to find the debug wrapper. Patch fixes this problem by emitting special name for the debug version of the code. llvm-svn: 310511
2017-08-10 03:38:53 +08:00
[OPENMP] Do not mark captured variables as artificial in debug info. Captured variables should not be marked as artificial parameters in outlined functions in debug info. llvm-svn: 318843
2017-11-23 00:02:03 +08:00
// CHECK: !DILocalVariable(name: ".global_tid.",
// CHECK-SAME: DIFlagArtificial
// CHECK: !DILocalVariable(name: ".bound_tid.",
// CHECK-SAME: DIFlagArtificial
// CHECK: !DILocalVariable(name: "c",
[OPENMP]Fix handling of the globals during compilation for the device. Fixed lookup for the target regions in unused virtual functions + fixed processing of the global variables not marked as declare target but emitted during debug info emission. llvm-svn: 346343
2018-11-08 03:11:14 +08:00
// CHECK-SAME: line: 22
[OPENMP] Do not mark captured variables as artificial in debug info. Captured variables should not be marked as artificial parameters in outlined functions in debug info. llvm-svn: 318843
2017-11-23 00:02:03 +08:00
// CHECK: !DILocalVariable(name: "a",
[OPENMP]Fix handling of the globals during compilation for the device. Fixed lookup for the target regions in unused virtual functions + fixed processing of the global variables not marked as declare target but emitted during debug info emission. llvm-svn: 346343
2018-11-08 03:11:14 +08:00
// CHECK-SAME: line: 20
[OPENMP] Do not mark captured variables as artificial in debug info. Captured variables should not be marked as artificial parameters in outlined functions in debug info. llvm-svn: 318843
2017-11-23 00:02:03 +08:00
// CHECK: !DILocalVariable(name: "b",
[OPENMP]Fix handling of the globals during compilation for the device. Fixed lookup for the target regions in unused virtual functions + fixed processing of the global variables not marked as declare target but emitted during debug info emission. llvm-svn: 346343
2018-11-08 03:11:14 +08:00
// CHECK-SAME: line: 21
[OPENMP] Fix debug info for outlined functions in NVPTX + add more tests. Fixed name of emitted outlined functions in NVPTX target + extra tests for the debug info. llvm-svn: 322022
2018-01-09 04:09:47 +08:00
// CHECK-DAG: distinct !DISubprogram(name: "[[NONDEBUG_WRAPPER]]",
// CHECK-DAG: distinct !DISubprogram(name: "[[DEBUG_PARALLEL]]",