Add support for AMD GPUs

[jakurzak]

Overview

This pull request adds support for AMD Instinct GPUs.

Changes

CMakeLists.txt

• The newly added section of the CMake script introduces logic to detect the presence of the hipcc compiler, which is part of AMD's ROCm software platform. If the hipcc compiler is detected, CMake will configure the build system to compile the HIP source files using hipcc.

Makefile

• Add HIPCC=hipcc and HIPCCFLAGS = -O3 variables, similarly to the existing NVCC=nvcc and HIPCCFLAGS = -O3 variables.
• Add .PHONY: qsim-hip, .PHONY: hip-tests, and .PHONY: run-hip-tests sections, similarly to the existing .PHONY: qsim-cuda, .PHONY: cuda-tests, and .PHONY: run-cuda-tests sections.

apps/Makefile

• Generate the list of HIP_TARGETS wih the .hip.x suffix from all files with the .cuda.cu suffix, similarly to how the list of CUDA_TARGETS is generated.
• Add the .PHONY: qsim-hip and the %hip.x: %cuda.cu section, similarly to the existing .PHONY: qsim-cuda and %cuda.x: %cuda.cu sections.

apps/make.sh

• The added section compiles CUDA source files with the hipcc compiler instead of nvcc when hipcc is available.

docs/_book.yaml

• Add "AMD GPU support" article to the "Other tutorials" section.

docs/tutorials/amd_gpu.md

• Add basic documentation for the AMD GPU support.

lib/cuda2hip.h

• Define a series of macros that map CUDA function calls, types, constants, and utilities to their corresponding HIP counterparts.
• Provide a template function __shfl_down_sync(), as a wrapper around HIP's __shfl_down() function, since the former is not currently available in HIP.

lib/fuser_mqubit.h

• Remove unused variable to eliminate Clang warning.

lib/simulator_cuda_kernels.h

• Add a section for including HIP headers if the HIP compiler is detected.

lib/statespace_cuda.h

• Add a section for including HIP headers if the HIP compiler is detected.
• Add a missing error check for cudaMemset().

lib/statespace_cuda_kernels.h

• Add a section for including HIP headers if the HIP compiler is detected.

lib/util_cuda.h

• Add a section for including HIP headers if the HIP compiler is detected.

lib/vectorspace_cuda.h

• Add a section for including HIP headers if the HIP compiler is detected.
• Add missing error checks for cudaFree(), cudaMemcpy(), and cudaDeviceSynchronize().

pybind_interface/Makefile

• Add QSIMLIB_HIP, similar to the existing QSIMLIB_CUDA.
• Rename pybind-gpu and decide-gpu to pybind-cuda and decide-cuda.
• Add pybind-hip and decide-hip.
• Add removal of HIP files in the clean section.

pybind_interface/decide/CMakeLists.txt

• To add HIP support to the CMake file, checks were introduced to detect whether hipcc is available. If hipcc is found, the project is configured to use HIP by setting HIP as one of the project languages. The CMake file then adjusts the build process to use HIP-specific commands (hip_add_library) and properties for compiling the relevant source files. It also includes the necessary directories and sets the module properties for creating a Python extension module with HIP support, allowing the same build system to work across both CUDA and HIP platforms depending on the available compiler.

pybind_interface/decide/decide.cpp

• Add HIP = 2 to enum GPUCapabilities.
• Set GPUCapabilities gpu = HIP if HIP compiler is detected.

pybind_interface/hip/CMakeLists.txt

• Set up HIP compilation for a Python extension module written in HIP and C++. Configure compiler optimizations and OpenMP settings. Integrate Pybind11 for C++/Python interfacing, and locate necessary Python libraries.

pybind_interface/hip/pybind_main_hip.cpp

• Integrate with Pybind11 for creating Python bindings by creating a factory pattern for object creation and setting up the necessary environment.

pybind_interface/hip/pybind_main_hip.h

• Introduce Python bindings for a module called qsim_hip using Pybind11.

pybind_interface/pybind_main.cpp

• Remove unused lambda captures: circuit, ncircuit, and state_space.

qsimcirq/__init__.py

• Add detection of HIP. If qsim_decide.detect_gpu() returns 2, import the HIP-compatible module qsimcirq.qsim_hip.

setup.py

• Add import shutil.
• If hipcc is found, append arguments to the cmake_args list to specify hipcc as the C and C++ compiler.
• Specify that the qsim_hip directory contains a CMake project to be built as a Python extension with support for HIP.

tests/Makefile

• Add rules for building HIP tests, similar to the rules for building CUDA tests.

tests/make.sh

• Add rules for building HIP tests, similar to the rules for building CUDA tests.

Building

make -j qsim      # to build CPU qsim
make -j qsim-hip  # to build HIP qsim
make -j pybind    # to build Python bindings
make -j cxx-tests # to build CPU tests
make -j hip-tests # to build HIP tests
pip install .

Testing

Simulator

make run-cxx-tests # to run CPU tests
make run-hip-tests # to run HIP tests

or

cd tests
for file in *.x; do ./"$file"; done          # to run all tests
for file in *_hip_test.x; do ./"$file"; done # to run HIP tests only

Python bindings

make run-py-tests

or

cd qsimcirq_tests
python3 -m pytest -v qsimcirq_test.py

[google-cla]

Thanks for your pull request! It looks like this may be your first contribution to a Google open source project. Before we can look at your pull request, you'll need to sign a Contributor License Agreement (CLA).

View this failed invocation of the CLA check for more information.

For the most up to date status, view the checks section at the bottom of the pull request.

[tanujkhattar]

@jakurzak Can you please add a PR description ?

@sergeisakov Can you take a look at this PR?

[sergeisakov]

@sergeisakov Can you take a look at this PR?

Yes, I'll take a look at this PR.

[sergeisakov]

Thank you for adding support for AMD GPUs. I can't really test if it works or not. Nevertheless the PR looks good to me. Could you fix a few formatting issues? Could you also add some doc?

[sergeisakov]

There should be a four space indent if the line breaks after (.

[jakurzak]

Fixed

[jakurzak]

Thank you for taking the time to look at the PR!
I am out of office this week and traveling overseas.
I'll do my best to address the comments ASAP when I'm back in office.

[jakurzak]

I fixed the style issues and commented on the use of HIPCC for CPU code when compiling Python bindings.
I'll be soliciting help with updating of documentation.

[jakurzak]

I added basic documentation for AMD GPU support.
I added detailed description to the PR.
Please let me know if anything else needs to be done before the PR can be accepted.

[sergeisakov]

I added basic documentation for AMD GPU support. I added detailed description to the PR. Please let me know if anything else needs to be done before the PR can be accepted.

Thank you. Could you replace Qsim with qsim in amd_gpu.md? Could you also add something about how to install the build tools? Is it ROCm?

[jakurzak]

Thank you. Could you replace Qsim with qsim in amd_gpu.md? Could you also add something about how to install the build tools? Is it ROCm?

done

[sergeisakov]

LGTM

[jakurzak]

Looks like it's getting hung up on the kokoro test.
Any clue why?

[95-martin-orion]

The cuQuantum test failed during bringup:

Logs

#21 ERROR: process "/bin/sh -c ${PIP_CMD} install --upgrade --force-reinstall ${tensorflow_pip_spec} ${extra_pip_specs}" did not complete successfully: exit code: 2
------
 > [17/26] RUN python3.9 -m pip install --upgrade --force-reinstall tf-nightly :
78.73     for chunk in response.raw.stream(
78.73   File "/usr/local/lib/python3.9/dist-packages/pip/_vendor/urllib3/response.py", line 622, in stream
78.73     data = self.read(amt=amt, decode_content=decode_content)
78.73   File "/usr/local/lib/python3.9/dist-packages/pip/_vendor/urllib3/response.py", line 587, in read
78.73     raise IncompleteRead(self._fp_bytes_read, self.length_remaining)
78.73   File "/usr/lib/python3.9/contextlib.py", line 137, in __exit__
78.73     self.gen.throw(typ, value, traceback)
78.73   File "/usr/local/lib/python3.9/dist-packages/pip/_vendor/urllib3/response.py", line 443, in _error_catcher
78.73     raise ReadTimeoutError(self._pool, None, "Read timed out.")
78.73 pip._vendor.urllib3.exceptions.ReadTimeoutError: HTTPSConnectionPool(host='files.pythonhosted.org', port=443): Read timed out.
------
Dockerfile_ubuntu_2004_tf_cuda_11:91
--------------------
  89 |     RUN ${PIP_CMD} install wheel
  90 |     RUN ${PIP_CMD} install absl-py
  91 | >>> RUN ${PIP_CMD} install --upgrade --force-reinstall ${tensorflow_pip_spec} ${extra_pip_specs}
  92 |     
  93 |     RUN ${PIP_CMD} install tfds-nightly
--------------------
ERROR: failed to solve: process "/bin/sh -c ${PIP_CMD} install --upgrade --force-reinstall ${tensorflow_pip_spec} ${extra_pip_specs}" did not complete successfully: exit code: 2

From the logs, it's possible this was a network error. I'll re-trigger the test, but in the meantime check the logs to see if this PR could have caused the issue.

[95-martin-orion]

Hmm. Doesn't look like I can re-trigger the kokoro test specifically. Could you push a commit so Github picks up on it?

[jakurzak]

No cuQuantum support in the AMD port, BTW, as there's no AMD equivalent.
The HIP port should not affect anything in the CUDA implementation, though.
Pushed a commit to re-trigger.

[jakurzak]

@sergeisakov, @95-martin-orion, can you help move this forward?
Thanks in advance!