
When offloading was enabled, top-level 'asm' were added to the offloading section, confusing assemblers which did not support the syntax. Additionally, with offloading and -flto, the top-level assembler code did not end up in the host files. As r14-321-g9a41d2cdbcd added top-level 'asm' to one libstdc++ header file, the issue became more apparent, causing fails with nvptx for some C++ testcases. PR libstdc++/109816 gcc/ChangeLog: * lto-cgraph.cc (output_symtab): Guard lto_output_toplevel_asms by '!lto_stream_offload_p'. libgomp/ChangeLog: * testsuite/libgomp.c++/target-map-class-1.C: New test. * testsuite/libgomp.c++/target-map-class-2.C: New test.
98 lines
1.9 KiB
C
98 lines
1.9 KiB
C
/* PR middle-end/109816 */
|
|
|
|
/* This variant: without -flto, see target-map-class-2.C for -flto. */
|
|
|
|
/* iostream.h adds 'globl _ZSt21ios_base_library_initv' with _GLIBCXX_SYMVER_GNU,
|
|
but it shouldn't end up in the offload assembly but only in the host assembly. */
|
|
|
|
/* Example based on sollve_vv's test_target_data_map_classes.cpp; however,
|
|
relevant is only the 'include' and not the actual executable code. */
|
|
|
|
#include <iostream>
|
|
#include <omp.h>
|
|
|
|
using namespace std;
|
|
|
|
#define N 1000
|
|
|
|
struct A
|
|
{
|
|
int *h_array;
|
|
int size, sum;
|
|
|
|
A (int *array, const int s) : h_array(array), size(s), sum(0) { }
|
|
~A() { h_array = NULL; }
|
|
};
|
|
|
|
void
|
|
test_map_tofrom_class_heap ()
|
|
{
|
|
int *array = new int[N];
|
|
A *obj = new A (array, N);
|
|
|
|
#pragma omp target map(from: array[:N]) map(tofrom: obj[:1])
|
|
{
|
|
int *tmp_h_array = obj->h_array;
|
|
obj->h_array = array;
|
|
int tmp = 0;
|
|
for (int i = 0; i < N; ++i)
|
|
{
|
|
obj->h_array[i] = 4*i;
|
|
tmp += 3;
|
|
}
|
|
obj->h_array = tmp_h_array;
|
|
obj->sum = tmp;
|
|
}
|
|
|
|
for (int i = 0; i < N; ++i)
|
|
if (obj->h_array[i] != 4*i)
|
|
__builtin_abort ();
|
|
|
|
if (3*N != obj->sum)
|
|
{
|
|
std::cout << "sum: " << obj->sum << std::endl;
|
|
__builtin_abort ();
|
|
}
|
|
|
|
delete obj;
|
|
delete[] array;
|
|
}
|
|
|
|
void
|
|
test_map_tofrom_class_stack ()
|
|
{
|
|
int array[N];
|
|
A obj(array, N);
|
|
|
|
#pragma omp target map(from: array[:N]) map(tofrom: obj)
|
|
{
|
|
int *tmp_h_array = obj.h_array;
|
|
obj.h_array = array;
|
|
int tmp = 0;
|
|
for (int i = 0; i < N; ++i)
|
|
{
|
|
obj.h_array[i] = 7*i;
|
|
tmp += 5;
|
|
}
|
|
obj.h_array = tmp_h_array;
|
|
obj.sum = tmp;
|
|
}
|
|
|
|
for (int i = 0; i < N; ++i)
|
|
if (obj.h_array[i] != 7*i)
|
|
__builtin_abort ();
|
|
|
|
if (5*N != obj.sum)
|
|
{
|
|
std::cout << "sum: " << obj.sum << std::endl;
|
|
__builtin_abort ();
|
|
}
|
|
}
|
|
|
|
int
|
|
main()
|
|
{
|
|
test_map_tofrom_class_heap();
|
|
test_map_tofrom_class_stack();
|
|
return 0;
|
|
}
|