Kyle Roarty has uploaded this change for review. (
https://gem5-review.googlesource.com/c/public/gem5/+/36035 )
Change subject: util: Update gcn Dockefile with better hipDeviceSynchronize
fix
......................................................................
util: Update gcn Dockefile with better hipDeviceSynchronize fix
While the previous patch (that updated HIP) resovled the
hipDeviceSynchronize patch for some cases, we recently came across
other cases where hipDeviceSynchronize failed to work.
This patch manually builds HCC, as we found a couple of places in HCC
that were causing crashes in gem5. One of the crashes was a race
condition that seemed to be creating the hipDeviceSynchronize issue
The other issue had to do with gem5 not properly executing a
shared_future wait, which caused a crash when running with HCC_DB flags
Currently WIP, need to validate against applications
Change-Id: I97dfe240c988bc2c2de84510b612b754eb3f2a96
---
M util/dockerfiles/gcn-gpu/Dockerfile
A util/dockerfiles/gcn-gpu/hcc.patch
2 files changed, 57 insertions(+), 14 deletions(-)
diff --git a/util/dockerfiles/gcn-gpu/Dockerfile
b/util/dockerfiles/gcn-gpu/Dockerfile
index d0fe759..f44119d 100644
--- a/util/dockerfiles/gcn-gpu/Dockerfile
+++ b/util/dockerfiles/gcn-gpu/Dockerfile
@@ -48,13 +48,11 @@
&& dpkg -i h/hsakmt-roct-dev/* \
&& dpkg -i h/hsa-ext-rocr-dev/* \
&& dpkg -i h/hsa-rocr-dev/* \
- && dpkg -i r/rocm-utils/* \
- && dpkg -i h/hcc/* \
- && dpkg -i r/rocm-opencl/* \
- && dpkg -i r/rocm-opencl-dev/*
+ && dpkg -i r/rocm-utils/*
# Get ROCm libraries we need to compile from source (and ROCm-profiler)
-RUN git clone --single-branch https://github.com/ROCm-Developer-Tools/HIP/
&& \
+RUN git clone --single-branch -b roc-1.6.x --recursive
https://github.com/RadeonOpenCompute/hcc/ && \
+ git clone --single-branch https://github.com/ROCm-Developer-Tools/HIP/
&& \
git clone --single-branch
https://github.com/ROCmSoftwarePlatform/hipBLAS/ && \
git clone --single-branch
https://github.com/ROCmSoftwarePlatform/rocBLAS/ && \
git clone --single-branch
https://github.com/ROCmSoftwarePlatform/MIOpenGEMM/ && \
@@ -65,11 +63,14 @@
# Apply patches to various repos
RUN mkdir -p /patch && cd /patch && \
wget ${gem5_dist}/rocm_patches/hipBLAS.patch && \
- wget ${gem5_dist}/rocm_patches/hip.patch_v2 && \
+ wget ${gem5_dist}/rocm_patches/hip.patch && \
wget ${gem5_dist}/rocm_patches/miopen-conv.patch && \
wget ${gem5_dist}/rocm_patches/rocBLAS.patch
-RUN git -C /HIP/ checkout 0e3d824e && git -C /HIP/ apply
/patch/hip.patch_v2 && \
+COPY hcc.patch /patch/hcc.patch
+
+RUN git -C /hcc/ apply /patch/hcc.patch &&\
+ git -C /HIP/ checkout 0e3d824e && git -C /HIP/ apply /patch/hip.patch
&& \
git -C /hipBLAS/ checkout ee57787e && git -C /hipBLAS/ apply
/patch/hipBLAS.patch && \
git -C /rocBLAS/ checkout cbff4b4e && git -C /rocBLAS/ apply
/patch/rocBLAS.patch && \
git -C /MIOpenGEMM/ checkout 9547fb9e && \
@@ -84,7 +85,8 @@
ENV HCC_AMDGPU_TARGET gfx801
# Create build dirs for machine learning ROCm installs
-RUN mkdir -p /HIP/build && \
+RUN mkdir -p /hcc/build && \
+ mkdir -p /HIP/build && \
mkdir -p /rocBLAS/build && \
mkdir -p /hipBLAS/build && \
mkdir -p /rocm-cmake/build && \
@@ -92,15 +94,22 @@
mkdir -p /MIOpen/build
# Do the builds, empty build dir to trim image size
+WORKDIR /hcc/build
+RUN cmake -DCMAKE_INSTALL_PREFIX=/opt/rocm/hcc .. && make -j10 && \
+ make install && rm -rf *
+
+WORKDIR /apt_1.6.2/pool/main
+RUN dpkg -i r/rocm-opencl/* && dpkg -i r/rocm-opencl-dev/*
+
WORKDIR /HIP/build
RUN cmake .. && make -j$(nproc) && make install && rm -rf *
WORKDIR /rocBLAS/build
-RUN CXX=/opt/rocm/bin/hcc cmake
-DCMAKE_CXX_FLAGS="--amdgpu-target=gfx801" .. && \
+RUN CXX=/opt/rocm/hcc/bin/hcc cmake
-DCMAKE_CXX_FLAGS="--amdgpu-target=gfx801" .. && \
make -j$(nproc) && make install && rm -rf *
WORKDIR /hipBLAS/build
-RUN CXX=/opt/rocm/bin/hcc cmake
-DCMAKE_CXX_FLAGS="--amdgpu-target=gfx801" .. && \
+RUN CXX=/opt/rocm/hcc/bin/hcc cmake
-DCMAKE_CXX_FLAGS="--amdgpu-target=gfx801" .. && \
make -j$(nproc) && make install && rm -rf *
WORKDIR /rocm-cmake/build
@@ -114,8 +123,8 @@
# Un-set default c++ version for MIOpen compilation
# As MIOpen 1.7 requires c++14 or higher
-RUN sed
-i 's/INTERFACE_COMPILE_OPTIONS "-std=c++amp;-fPIC;-gline-tables-only"/#&/'
/opt/rocm/hcc-1.0/lib/cmake/hcc/hcc-targets.cmake && \
- sed -i 's/INTERFACE_COMPILE_OPTIONS "-hc"/#&/'
/opt/rocm/hcc-1.0/lib/cmake/hcc/hcc-targets.cmake
+RUN sed
-i 's/INTERFACE_COMPILE_OPTIONS "-std=c++amp;-fPIC;-gline-tables-only"/#&/'
/opt/rocm/hcc/lib/cmake/hcc/hcc-targets.cmake && \
+ sed -i 's/INTERFACE_COMPILE_OPTIONS "-hc"/#&/'
/opt/rocm/hcc/lib/cmake/hcc/hcc-targets.cmake
WORKDIR /MIOpen
# Half is required; This is the version that MIOpen would download
@@ -135,8 +144,8 @@
make -j$(nproc) && make install && rm -rf *
# Re-set defaults
-RUN sed
-i 's/#\(INTERFACE_COMPILE_OPTIONS "-std=c++amp;-fPIC;-gline-tables-only"\)/\1/'
/opt/rocm/hcc-1.0/lib/cmake/hcc/hcc-targets.cmake && \
- sed -i 's/#\(INTERFACE_COMPILE_OPTIONS "-hc"\)/\1/'
/opt/rocm/hcc-1.0/lib/cmake/hcc/hcc-targets.cmake
+RUN sed
-i 's/#\(INTERFACE_COMPILE_OPTIONS "-std=c++amp;-fPIC;-gline-tables-only"\)/\1/'
/opt/rocm/hcc/lib/cmake/hcc/hcc-targets.cmake && \
+ sed -i 's/#\(INTERFACE_COMPILE_OPTIONS "-hc"\)/\1/'
/opt/rocm/hcc/lib/cmake/hcc/hcc-targets.cmake
# Create performance DB for gfx801.
WORKDIR /opt/rocm/miopen/share/miopen/db
diff --git a/util/dockerfiles/gcn-gpu/hcc.patch
b/util/dockerfiles/gcn-gpu/hcc.patch
new file mode 100644
index 0000000..e745468
--- /dev/null
+++ b/util/dockerfiles/gcn-gpu/hcc.patch
@@ -0,0 +1,34 @@
+diff --git a/lib/hsa/mcwamp_hsa.cpp b/lib/hsa/mcwamp_hsa.cpp
+index 6d1a641f..04c2dbc3 100644
+--- a/lib/hsa/mcwamp_hsa.cpp
++++ b/lib/hsa/mcwamp_hsa.cpp
+@@ -1322,11 +1322,14 @@ public:
+ assert(sig.handle != 0);
+ foundFirstValidOp = true;
+ }
+- // wait on valid futures only
+- std::shared_future<void>* future = asyncOp->getFuture();
+- if (future && future->valid()) {
+- future->wait();
+- }
++ // This is commented out becaues calls to future->wait()
are
++ // sometimes broken in gem5. It doesn't break any
functionality
++ // as the call to asyncOps.clear() calls the function that
++ // future waits on (waitComplete())
++ //std::shared_future<void>* future = asyncOp->getFuture();
++ //if (future && future->valid()) {
++ // future->wait();
++ //}
+ }
+ }
+ // clear async operations table
+@@ -1438,7 +1441,8 @@ public:
+ auto&& dependentAsyncOpVector = bufferKernelMap[buffer];
+ for (int i = 0; i < dependentAsyncOpVector.size(); ++i) {
+ auto dependentAsyncOp = dependentAsyncOpVector[i];
+- if (!dependentAsyncOp.expired()) {
++ auto dependentAsyncOpPointer = dependentAsyncOp.lock();
++ if (dependentAsyncOpPointer) {
+ auto dependentAsyncOpPointer = dependentAsyncOp.lock();
+ // wait on valid futures only
+ std::shared_future<void>* future =
dependentAsyncOpPointer->getFuture();
--
To view, visit https://gem5-review.googlesource.com/c/public/gem5/+/36035
To unsubscribe, or for help writing mail filters, visit
https://gem5-review.googlesource.com/settings
Gerrit-Project: public/gem5
Gerrit-Branch: develop
Gerrit-Change-Id: I97dfe240c988bc2c2de84510b612b754eb3f2a96
Gerrit-Change-Number: 36035
Gerrit-PatchSet: 1
Gerrit-Owner: Kyle Roarty <kyleroarty1...@gmail.com>
Gerrit-MessageType: newchange
_______________________________________________
gem5-dev mailing list -- gem5-dev@gem5.org
To unsubscribe send an email to gem5-dev-le...@gem5.org
%(web_page_url)slistinfo%(cgiext)s/%(_internal_name)s
