Skip to content

[Issue]: [OpenMP] amdclang runtime error with schedule(dynamic) at -O0 #267

Open
@jrodgers-github

Description

@jrodgers-github

Problem Description

Executing the following reproducer script:

#!/bin/bash

module load rocm/6.4.0

TEST_NAME=amdclang-target-offload-schedule-failure

cat > ${TEST_NAME}.c << EOF
#include <omp.h>

int main()
{
  int N = 100;
  double a[N];
  double sum1 = 0;
  for (int i=0; i<N; i++)
    a[i]=i;
#pragma omp target teams distribute parallel for map(tofrom:sum1) reduction(+:sum1) schedule(dynamic,10)
  for (int j = 0; j < N; j=j+1)
    sum1 += a[j];
  return 0;
}
EOF

[[ -f ${TEST_NAME}.exe ]] && rm ${TEST_NAME}.exe
[[ -f ${TEST_NAME}.o ]] && rm ${TEST_NAME}.o
OPTIONS=" -O0 -fopenmp -fopenmp-targets=amdgcn-amd-amdhsa -Xopenmp-target=amdgcn-amd-amdhsa -march=gfx908 "

echo "AMDClang version"
which amdclang
echo "Compile"
amdclang -c ${TEST_NAME}.c -o ${TEST_NAME}.o ${OPTIONS}
echo "Link"
amdclang -o ${TEST_NAME}.exe ${TEST_NAME}.o ${OPTIONS}
echo "Run"
./${TEST_NAME}.exe

I'm finding the following behavior:

AMDClang version
/opt/rocm-6.4.0/bin/amdclang
Compile
Link
Run
AMDGPU fatal error 1: Memory access fault by GPU 2 (agent 0x1280c060) at virtual address 0x7f0f8521b000. Reasons: Page not present or supervisor privilege, Write access to a read-only page
./amdclang-reprod-debug-O0.sh: line 35: 67155 Aborted                 (core dumped) ./${TEST_NAME}.exe

Executing the same reproducer with any other -O optimization level (-O{1,2,3}), the fatal error is avoided.

Operating System

SLES15-SP6

CPU

AMD EPYC 7742 64-Core Processor

GPU

AMD Instinct MI100

ROCm Version

ROCm 6.4.0, ROCm 6.3.0, ROCm 6.2.0

ROCm Component

No response

Steps to Reproduce

See description for bash reproducer.

(Optional for Linux users) Output of /opt/rocm/bin/rocminfo --support

No response

Additional Information

No response

Metadata

Metadata

Assignees

No one assigned

    Type

    No type

    Projects

    No projects

    Milestone

    No milestone

    Relationships

    None yet

    Development

    No branches or pull requests

    Issue actions