Accelerator restriction: unsupported call to support routine 'memcmp'

Accelerated Computing HPC Compilers Legacy PGI Compilers
1

When I compile pgc++, I encountered following errors. I am sure such error has no relationship with my code (line 259). No matter what code I comment out, I got the same error. I just simply add #pragma acc routine in front of my subroutine

operator new (unsigned long, void *):
38, include “iostream”
39, include “ostream”
38, include “ios”
42, include “ios_base.h”
41, include “locale_classes.h”
40, include “string”
41, include “allocator.h”
46, include “c++allocator.h”
33, include “new_allocator.h”
33, include “new”
130, Generating implicit acc routine seq
Generating acc routine seq
Generating Tesla code
operator delete (void *, void *):
38, include “iostream”
39, include “ostream”
38, include “ios”
42, include “ios_base.h”
41, include “locale_classes.h”
40, include “string”
41, include “allocator.h”
46, include “c++allocator.h”
33, include “new_allocator.h”
33, include “new”
135, Generating implicit acc routine seq
Generating acc routine seq
Generating Tesla code
PGCC-S-0000-Internal compiler error. Call in OpenACC region to support routine - memcmp (/xxx.cpp: 259)
PGCC-W-0155-Compiler failed to translate accelerator region (see -Minfo messages): Missing branch target block (/xxx.cpp: 1)
std::char_traits::compare(const char *, const char *, unsigned long):
38, include “iostream”
39, include “ostream”
38, include “ios”
40, include “char_traits.h”
259, Generating implicit acc routine seq
Generating acc routine seq
Generating Tesla code
262, Accelerator restriction: unsupported call to support routine ‘memcmp’
42, include “ios_base.h”
PGCC-F-0704-Compilation aborted due to previous errors. (/xxx.cpp)
PGCC/x86 Linux 18.4-0: compilation aborted

2

Hi Po Chun LAI,

All subroutine calls made from within an OpenACC compute region must have a device version available. This means that the OpenACC “routine” directive must be added to all of these subroutines so the compiler knows to generate the device version.

In C++, adding “routine” to all subroutines would be very laborious. Also, there may be cases where the user doesn’t have access to the subroutine’s definition so can’t explicitly add “routine”. Hence, pgc++ will implicitly add “routine” to all visible subroutines called from within a compute region (as well as all routines that they call).

In this case, it looks like you are declaring a class object within the device code so the compiler is having to implicitly create the “new” and “delete” operators for this class. Either in a later routine or within “delete”, you have a call to “std::char_traits::compare” which the compiler is attempting to implicitly create a device routine. However, this routine is calling the system OS routine “memcmp” for which there’s no device equivalent.

I/O and strings are not supported within device code. My best guess is that in your class’ delete operator, you have a “cout” or some other type of I/O operation. If this is the case, you might try using conditional compilation so the I/O is not used when compiling with OpenACC. For example:

#ifndef _OPENACC
  std::cout << "something" << std::endl;
#endif

Longer term, you may want to track down why the class object is getting created. If your class doesn’t do any allocation in “new”, then it’s not an issue. But if it does, device side allocation is very slow and the default device heap is very small (8MB) so should be avoided.

-Mat