Internal compiler error on nvc 26.5 and 26.9

Accelerated Computing HPC Compilers nvc, nvc++ and nvfortran
1

We are getting the error

NVC++-F-0000-Internal compiler error. size of unknown type       0  (nvc_26_5_monitor_nd_ice.c: 16)
NVC++/x86-64 Linux 26.9-0: compilation aborted

When compiling the program below with

nvc -fast -Minfo=accel -acc=gpu -gpu=mem:managed -DOPENACC -c -o o.out nvc_26_5_monitor_nd_ice.c

Interestingly it does work with -O1 instead of fast. The program compiles with 26.1

#include <stdlib.h>

typedef struct {
  double *buffer;
} buffer_state;

#pragma acc routine
void copy_rows(buffer_state *state)
{
  int i;

  for (i = 0; i < 4; i++) {
    #pragma acc atomic write
    state->buffer[i] = 0.0;
  }
}

int main(void)
{
  double *buffer = calloc(4, sizeof(*buffer));
  buffer_state state;

  if (!buffer)
    return 1;

  state.buffer = buffer;

  #pragma acc enter data copyin(buffer[0:4])
  #pragma acc enter data copyin(state)
  #pragma acc enter data attach(state.buffer)
  #pragma acc parallel present(state)
  {
    copy_rows(&state);
  }
  #pragma acc exit data detach(state.buffer)
  #pragma acc exit data copyout(buffer[0:4])
  #pragma acc exit data delete(state)

  for (int i = 0; i < 4; i++) {
    if (buffer[i] != 0.0)
      return 1;
  }

  free(buffer);
  return 0;
}

Any insight would be appreciated, but seems like a bug in the compiler?

2

Hi,

I was able to reproduce your issue and I agree that it looks like a bug! I have reported it to engineering as TPR#39009. I will let you know when I find out more information. I did test a workaround that seemed to work - doing the write through a local pointer instead of the structured one:

#include <stdlib.h>

typedef struct {
  double *buffer;
} buffer_state;

#pragma acc routine
void copy_rows(buffer_state *state)
{
  int i;
  double *p = state->buffer;

  for (i = 0; i < 4; i++) {
    #pragma acc atomic write
    p[i] = 0.0;
  }
}

int main(void)
{
  double *buffer = calloc(4, sizeof(*buffer));
  buffer_state state;

  if (!buffer)
    return 1;

  state.buffer = buffer;

  #pragma acc enter data copyin(buffer[0:4])
  #pragma acc enter data copyin(state)
  #pragma acc enter data attach(state.buffer)
  #pragma acc parallel present(state)
  {
    copy_rows(&state);
  }
  #pragma acc exit data detach(state.buffer)
  #pragma acc exit data copyout(buffer[0:4])
  #pragma acc exit data delete(state)

  for (int i = 0; i < 4; i++) {
    if (buffer[i] != 0.0)
      return 1;
  }

  free(buffer);
  return 0;
}

I hope this helps! Let me know if you have any other problems I can help with! And please report any other issues you find - all feedback makes the compiler better!

Cheers,

Seth.

3

Thank you very much for the swift answer and the workaround, that’s very applicable!

We will do our duty and report any issues we run into.