포인터가 참조하는 구조체의 동적으로 할당 된 구조체 멤버 배열에 값을 할당 할 때 openacc 오류

Hendrik van Eerten

동적으로 할당 된 멤버를 포함하는 구조체에 대한 포인터와 openacc를 결합하여 머리를 감싸려고합니다. 아래 코드는 실패합니다.

cuStreamSynchronize에 대한 Thread : 1 호출에 실패하면 오류 700이 반환되었습니다. 커널 실행 중 잘못된 주소

nvc ( "x86-64 Linux의 nvc 20.9-0 LLVM 64 비트 대상 -tp haswell")를 사용하여 컴파일 할 때. 내가 말할 수있는 한 나는 OpenACC '시작하기'가이드에서 제안 된 접근 방식을 따르고있다. 그러나 어떻게 든 아마도 포인터가 장치에 붙어 있지 않습니다 (?). 여기에서 무엇이 잘못되었는지 아는 사람이 있습니까?

#include <stdlib.h>
#include <stdio.h>

typedef struct grid
{
  int N;
  double *X;
} grid;

void allocate(grid* g, int N)
{
  g->N = N;
  g->X = (double*) malloc(sizeof(double) * g->N);

  #pragma acc enter data create(g[0:1])
  #pragma acc enter data create(g->X[0:N])
}

void release(grid* g)
{
  #pragma acc exit data delete(g->X[0:g->N])
  #pragma acc exit data delete(g[0:1])

  free(g->X);
}

void fill(grid * g)
{
  int i;

  #pragma acc parallel loop
  for (i = 0; i < g->N; i++)
  {
    g->X[i] = 42; // the cuprit, commenting this removes the error too
  }
}

int main()
{
  grid g;

  allocate(&g, 10);

  fill(&g);

  release(&g);

  return 0;
}```
Mat Colgrove

컴파일러 피드백 메시지에서 다음과 같은 내용을 볼 수 있습니다.

     fill:
          32, Accelerator restriction: size of the GPU copy of g is unknown
              Generating Tesla code
              32, #pragma acc loop gang, vector(128) /* blockIdx.x threadIdx.x */
          32, Generating implicit copyin(g) [if not already present]
          37, Generating update self(g->X[:g->N])

문제는 컴파일러가 동적 데이터 멤버를 사용하여 집계 유형을 암시 적으로 복사 할 수 없으므로 g가 이미 장치임을 나타 내기 위해 "present (g)"를 추가해야한다는 것입니다.

또한 장치에서 N 값을 가져오고 종료 데이터 삭제 지시문에 배열 모양을 포함 할 필요가 없도록 복사하고 싶을 것입니다. 예를 들면 :

% cat test.c
#include <stdlib.h>
#include <stdio.h>

typedef struct grid
{
  int N;
  double *X;
} grid;

void allocate(grid* g, int N)
{
  g->N = N;
  g->X = (double*) malloc(sizeof(double) * g->N);

  #pragma acc enter data copyin(g[0:1])
  #pragma acc enter data create(g->X[0:N])
}

void release(grid* g)
{
  #pragma acc exit data delete(g->X)
  #pragma acc exit data delete(g)

  free(g->X);
}

void fill(grid * g)
{
  int i;

  #pragma acc parallel loop present(g)
  for (i = 0; i < g->N; i++)
  {
    g->X[i] = 42; // the cuprit, commenting this removes the error too
  }
  #pragma acc update self(g->X[:g->N])
  for (i = 0; i < 4; i++)
  {
    printf("%d : %f \n",i,g->X[i]);
  }
}

int main()
{
  grid g;

  allocate(&g, 10);

  fill(&g);

  release(&g);

  return 0;
}

% nvc -acc test.c -Minfo=accel -V20.9 ; a.out
allocate:
     17, Generating enter data copyin(g[:1])
         Generating enter data create(g->X[:N])
release:
     24, Generating exit data delete(g[:1],g->X[:1])
fill:
     32, Generating present(g[:1])
         Generating Tesla code
         32, #pragma acc loop gang, vector(128) /* blockIdx.x threadIdx.x */
     37, Generating update self(g->X[:g->N])
0 : 42.000000
1 : 42.000000
2 : 42.000000
3 : 42.000000

이 기사는 인터넷에서 수집됩니다. 재 인쇄 할 때 출처를 알려주십시오.

침해가 발생한 경우 연락 주시기 바랍니다[email protected] 삭제

에서 수정
0

몇 마디 만하겠습니다

0리뷰
로그인참여 후 검토

관련 기사

Related 관련 기사

뜨겁다태그

보관