diff --git a/.github/workflows/pretest.yml b/.github/workflows/pretest.yml index 76d7fea8c97..86043e6610f 100644 --- a/.github/workflows/pretest.yml +++ b/.github/workflows/pretest.yml @@ -30,6 +30,7 @@ jobs: - name: Setup pre-commit run: | + pip install --upgrade autopep8 pip install pre-commit - name: Check @@ -106,4 +107,4 @@ jobs: - name: Build & Test run: | - docker run --rm -v "${PWD}:/src" -w /src "rocm/dev-ubuntu-22.04:6.0" bash .github/workflows/scripts/pretest-rocm-test.sh + docker run --rm -v "${PWD}:/src" -w /src "rocm/dev-ubuntu-22.04:6.4" bash .github/workflows/scripts/pretest-rocm-test.sh diff --git a/cupy/_core/core.pyx b/cupy/_core/core.pyx index 3fdaac15c9c..cf6e7cc468f 100644 --- a/cupy/_core/core.pyx +++ b/cupy/_core/core.pyx @@ -261,9 +261,6 @@ cdef class _ndarray_base: @property def __cuda_array_interface__(self): - if runtime._is_hip_environment: - raise AttributeError( - 'HIP/ROCm does not support cuda array interface') cdef dict desc = { 'shape': self.shape, 'typestr': self.dtype.str, @@ -1800,8 +1797,6 @@ cdef class _ndarray_base: # `_kernel._preprocess_args`. check = (hasattr(x, '__cuda_array_interface__') or hasattr(x, '__cupy_get_ndarray__')) - if runtime._is_hip_environment and isinstance(x, ndarray): - check = True if (not check and not type(x) in _scalar.scalar_type_set and not isinstance(x, numpy.ndarray)): @@ -3000,9 +2995,6 @@ cpdef _ndarray_base asfortranarray(_ndarray_base a, dtype=None): cpdef _ndarray_base _convert_object_with_cuda_array_interface(a): - if runtime._is_hip_environment: - raise RuntimeError( - 'HIP/ROCm does not support cuda array interface') cdef Py_ssize_t sh, st cdef dict desc = a.__cuda_array_interface__ diff --git a/cupy/cuda/cupy_cufft.h b/cupy/cuda/cupy_cufft.h index f52fe367534..168b34287df 100644 --- a/cupy/cuda/cupy_cufft.h +++ b/cupy/cuda/cupy_cufft.h @@ -12,7 +12,12 @@ #include #elif defined(CUPY_USE_HIP) +#include //for HIP_VERSION +#if HIP_VERSION >= 50530600 +#include +#else #include +#endif extern "C" { diff --git a/cupy/cuda/cupy_thrust.cu b/cupy/cuda/cupy_thrust.cu index 069af242a0b..65b54ab93c3 100644 --- a/cupy/cuda/cupy_thrust.cu +++ b/cupy/cuda/cupy_thrust.cu @@ -476,7 +476,7 @@ struct _argsort { #else thrust::make_counting_iterator(0), thrust::make_counting_iterator(size), - thrust::make_constant_iterator(shape[ndim-1]), + thrust::make_constant_iterator(shape[ndim-1]), #endif dp_keys_first, thrust::divides()); diff --git a/cupy/cuda/cupy_thrust.h b/cupy/cuda/cupy_thrust.h index 76cfe9ff7f8..838142b981e 100644 --- a/cupy/cuda/cupy_thrust.h +++ b/cupy/cuda/cupy_thrust.h @@ -4,7 +4,14 @@ #ifndef CUPY_NO_CUDA #include #include + +#ifndef CUPY_USE_HIP #include // for THRUST_VERSION +#else +// WAR #9098: +// rocThrust 3.3.0 (ROCm 6.4.0) cannot be compiled by host compiler +#define THRUST_VERSION 0 +#endif void thrust_sort(int, void *, size_t *, const std::vector&, intptr_t, void *); void thrust_lexsort(int, size_t *, void *, size_t, size_t, intptr_t, void *); diff --git a/cupy/cuda/memory.pyx b/cupy/cuda/memory.pyx index 3a272fdf968..b9aea1c7af0 100644 --- a/cupy/cuda/memory.pyx +++ b/cupy/cuda/memory.pyx @@ -506,10 +506,7 @@ cdef class MemoryPointer: """ stream_ptr = stream_module.get_current_stream_ptr() - if ( - not runtime._is_hip_environment - and runtime.streamIsCapturing(stream_ptr) - ): + if runtime.streamIsCapturing(stream_ptr): raise RuntimeError( 'the current stream is capturing, so synchronous API calls ' 'are disallowed') @@ -553,10 +550,7 @@ cdef class MemoryPointer: """ stream_ptr = stream_module.get_current_stream_ptr() - if ( - not runtime._is_hip_environment - and runtime.streamIsCapturing(stream_ptr) - ): + if runtime.streamIsCapturing(stream_ptr): raise RuntimeError( 'the current stream is capturing, so synchronous API calls ' 'are disallowed') @@ -580,10 +574,7 @@ cdef class MemoryPointer: stream_ptr = stream_module.get_current_stream_ptr() else: stream_ptr = stream.ptr - if ( - not runtime._is_hip_environment - and runtime.streamIsCapturing(stream_ptr) - ): + if runtime.streamIsCapturing(stream_ptr): raise RuntimeError( 'the current stream is capturing, so H2D transfers ' 'are disallowed') @@ -651,10 +642,7 @@ cdef class MemoryPointer: """ stream_ptr = stream_module.get_current_stream_ptr() - if ( - not runtime._is_hip_environment - and runtime.streamIsCapturing(stream_ptr) - ): + if runtime.streamIsCapturing(stream_ptr): raise RuntimeError( 'the current stream is capturing, so synchronous API calls ' 'are disallowed') @@ -678,10 +666,7 @@ cdef class MemoryPointer: stream_ptr = stream_module.get_current_stream_ptr() else: stream_ptr = stream.ptr - if ( - not runtime._is_hip_environment - and runtime.streamIsCapturing(stream_ptr) - ): + if runtime.streamIsCapturing(stream_ptr): raise RuntimeError( 'the current stream is capturing, so D2H transfers ' 'are disallowed') @@ -705,10 +690,7 @@ cdef class MemoryPointer: """ stream_ptr = stream_module.get_current_stream_ptr() - if ( - not runtime._is_hip_environment - and runtime.streamIsCapturing(stream_ptr) - ): + if runtime.streamIsCapturing(stream_ptr): raise RuntimeError( 'the current stream is capturing, so synchronous API calls ' 'are disallowed') diff --git a/cupy/cuda/stream.pyx b/cupy/cuda/stream.pyx index eb8ed8d1a17..4aaafd61c0c 100644 --- a/cupy/cuda/stream.pyx +++ b/cupy/cuda/stream.pyx @@ -369,8 +369,6 @@ class _BaseStream: https://docs.nvidia.com/cuda/cuda-runtime-api/group__CUDART__STREAM.html#group__CUDART__STREAM_1g793d7d4e474388ddfda531603dc34aa3 """ - if runtime._is_hip_environment: - raise RuntimeError('This function is not supported on HIP') if self.ptr == 0 or self.ptr == 1: raise RuntimeError('cannot capture on the default (legacy) stream') if mode is None: @@ -399,8 +397,6 @@ class _BaseStream: https://docs.nvidia.com/cuda/cuda-runtime-api/group__CUDART__STREAM.html#group__CUDART__STREAM_1gf5a0efebc818054ceecd1e3e5e76d93e """ - if runtime._is_hip_environment: - raise RuntimeError('This function is not supported on HIP') cdef intptr_t g = runtime.streamEndCapture(self.ptr) return graph.Graph.from_stream(g) @@ -415,9 +411,6 @@ class _BaseStream: Programming Guide for detail. """ - # TODO(leofang): is it better to be a property? - if runtime._is_hip_environment: - raise RuntimeError('This function is not supported on HIP') try: return runtime.streamIsCapturing(self.ptr) except RuntimeError: # can be RuntimeError or CUDARuntimeError @@ -483,9 +476,6 @@ class Stream(_BaseStream): ptr = 0 device_id = -1 elif ptds: - if runtime._is_hip_environment: - raise ValueError('HIP does not support per-thread ' - 'default stream (ptds)') ptr = runtime.streamPerThread device_id = -1 else: @@ -544,5 +534,4 @@ class ExternalStream(_BaseStream): Stream.null = Stream(null=True) -if not runtime._is_hip_environment: - Stream.ptds = Stream(ptds=True) +Stream.ptds = Stream(ptds=True) diff --git a/cupy/random/cupy_distributions.cuh b/cupy/random/cupy_distributions.cuh index f010bff6cac..87678790620 100644 --- a/cupy/random/cupy_distributions.cuh +++ b/cupy/random/cupy_distributions.cuh @@ -34,7 +34,9 @@ struct rk_binomial_state { // When compiling cython extensions with hip 4.0 // gcc will be used, but the hiprand_kernel can only be compiled with llvm // so we need to explicitly declare stubs for the functions -#if HIP_VERSION > 400 +#if HIP_VERSION >= 50530600 +#include +#elif HIP_VERSION > 400 #include #else #include diff --git a/cupy_backends/cuda/api/runtime.pyx b/cupy_backends/cuda/api/runtime.pyx index d65b734def2..652664ace5a 100644 --- a/cupy_backends/cuda/api/runtime.pyx +++ b/cupy_backends/cuda/api/runtime.pyx @@ -530,8 +530,8 @@ cpdef intptr_t mallocArray(intptr_t descPtr, size_t width, size_t height, cpdef intptr_t mallocAsync(size_t size, intptr_t stream) except? 0: cdef void* ptr - if _is_hip_environment: - raise RuntimeError('HIP does not support mallocAsync') + if _is_hip_environment and 0 < CUPY_HIP_VERSION < 60200000: + raise RuntimeError('mallocAsync is supported since ROCm 6.2') with nogil: status = cudaMallocAsync(&ptr, size, stream) check_status(status) @@ -581,8 +581,8 @@ cpdef freeArray(intptr_t ptr): check_status(status) cpdef freeAsync(intptr_t ptr, intptr_t stream): - if _is_hip_environment: - raise RuntimeError('HIP does not support freeAsync') + if _is_hip_environment and 0 < CUPY_HIP_VERSION < 60200000: + raise RuntimeError('freeAsync is supported since ROCm 6.2') with nogil: status = cudaFreeAsync(ptr, stream) check_status(status) @@ -715,7 +715,7 @@ cpdef PointerAttributes pointerGetAttributes(intptr_t ptr): cdef _PointerAttributes attrs status = cudaPointerGetAttributes(&attrs, ptr) check_status(status) - IF CUPY_CUDA_VERSION > 0: + IF CUPY_CUDA_VERSION > 0 or CUPY_HIP_VERSION > 60000000: return PointerAttributes( attrs.device, attrs.devicePointer, @@ -896,9 +896,6 @@ cdef _HostFnFunc(void* func_arg) with gil: cpdef streamAddCallback(intptr_t stream, callback, intptr_t arg, unsigned int flags=0): - if _is_hip_environment and stream == 0: - raise RuntimeError('HIP does not allow adding callbacks to the ' - 'default (null) stream') func_arg = (callback, arg) cpython.Py_INCREF(func_arg) with nogil: @@ -909,9 +906,6 @@ cpdef streamAddCallback(intptr_t stream, callback, intptr_t arg, cpdef launchHostFunc(intptr_t stream, callback, intptr_t arg): - if _is_hip_environment: - raise RuntimeError('This feature is not supported on HIP') - func_arg = (callback, arg) cpython.Py_INCREF(func_arg) with nogil: @@ -933,8 +927,6 @@ cpdef streamWaitEvent(intptr_t stream, intptr_t event, unsigned int flags=0): cpdef streamBeginCapture(intptr_t stream, int mode=streamCaptureModeRelaxed): - if _is_hip_environment: - raise RuntimeError('streamBeginCapture is not supported in ROCm') # TODO(leofang): check and raise if stream == 0? with nogil: status = cudaStreamBeginCapture(stream, @@ -945,8 +937,6 @@ cpdef streamBeginCapture(intptr_t stream, int mode=streamCaptureModeRelaxed): cpdef intptr_t streamEndCapture(intptr_t stream) except? 0: # TODO(leofang): check and raise if stream == 0? cdef Graph g - if _is_hip_environment: - raise RuntimeError('streamEndCapture is not supported in ROCm') with nogil: status = cudaStreamEndCapture(stream, &g) check_status(status) @@ -955,8 +945,6 @@ cpdef intptr_t streamEndCapture(intptr_t stream) except? 0: cpdef bint streamIsCapturing(intptr_t stream) except*: cdef StreamCaptureStatus s - if _is_hip_environment: - raise RuntimeError('streamIsCapturing is not supported in ROCm') with nogil: status = cudaStreamIsCapturing(stream, &s) check_status(status) # cudaErrorStreamCaptureImplicit could be raised here diff --git a/cupy_backends/cuda/libs/cublas.pyx b/cupy_backends/cuda/libs/cublas.pyx index 2e94f9c9c19..89d86d41c22 100644 --- a/cupy_backends/cuda/libs/cublas.pyx +++ b/cupy_backends/cuda/libs/cublas.pyx @@ -508,7 +508,7 @@ cpdef setStream(intptr_t handle, size_t stream): # https://docs.nvidia.com/cuda/cublas/index.html#CUDA-graphs # Before we come up with a robust strategy to test the support conditions, # we disable this functionality. - if not runtime._is_hip_environment and runtime.streamIsCapturing(stream): + if runtime.streamIsCapturing(stream): raise NotImplementedError( 'calling cuBLAS API during stream capture is currently ' 'unsupported') diff --git a/cupy_backends/cuda/libs/curand.pyx b/cupy_backends/cuda/libs/curand.pyx index d8e8df97036..3812f2e0ec5 100644 --- a/cupy_backends/cuda/libs/curand.pyx +++ b/cupy_backends/cuda/libs/curand.pyx @@ -115,7 +115,7 @@ cpdef int getVersion() except? -1: cpdef setStream(size_t generator, size_t stream): # TODO(leofang): The support of stream capture is not mentioned at all in # the cuRAND docs (as of CUDA 11.5), so we disable this functionality. - if not runtime._is_hip_environment and runtime.streamIsCapturing(stream): + if runtime.streamIsCapturing(stream): raise NotImplementedError( 'calling cuRAND API during stream capture is currently ' 'unsupported') diff --git a/cupy_backends/cuda/libs/cusolver.pyx b/cupy_backends/cuda/libs/cusolver.pyx index c994d4d6646..dcc5d6c25be 100644 --- a/cupy_backends/cuda/libs/cusolver.pyx +++ b/cupy_backends/cuda/libs/cusolver.pyx @@ -1136,7 +1136,7 @@ cpdef spDestroy(intptr_t handle): cpdef setStream(intptr_t handle, size_t stream): # TODO(leofang): The support of stream capture is not mentioned at all in # the cuSOLVER docs (as of CUDA 11.5), so we disable this functionality. - if not runtime._is_hip_environment and runtime.streamIsCapturing(stream): + if runtime.streamIsCapturing(stream): raise NotImplementedError( 'calling cuSOLVER API during stream capture is currently ' 'unsupported') diff --git a/cupy_backends/cuda/libs/cusparse.pyx b/cupy_backends/cuda/libs/cusparse.pyx index 0b7a99c7018..5994a53197f 100644 --- a/cupy_backends/cuda/libs/cusparse.pyx +++ b/cupy_backends/cuda/libs/cusparse.pyx @@ -1630,7 +1630,7 @@ cpdef void setStream(intptr_t handle, size_t stream) except *: # https://docs.nvidia.com/cuda/cusparse/index.html#optimization-notes # Before we come up with a robust strategy to test the support conditions, # we disable this functionality. - if not runtime._is_hip_environment and runtime.streamIsCapturing(stream): + if runtime.streamIsCapturing(stream): raise NotImplementedError( 'calling cuSPARSE API during stream capture is currently ' 'unsupported') diff --git a/cupy_backends/hip/cupy_hip_common.h b/cupy_backends/hip/cupy_hip_common.h index 8f428f4903e..48337125851 100644 --- a/cupy_backends/hip/cupy_hip_common.h +++ b/cupy_backends/hip/cupy_hip_common.h @@ -2,8 +2,13 @@ #define INCLUDE_GUARD_HIP_CUPY_COMMON_H #include +#if HIP_VERSION >= 50530600 +#include +#include +#else #include #include +#endif #define CUDA_VERSION 0 diff --git a/cupy_backends/hip/cupy_hip_runtime.h b/cupy_backends/hip/cupy_hip_runtime.h index 9a514d32afb..459f8653409 100644 --- a/cupy_backends/hip/cupy_hip_runtime.h +++ b/cupy_backends/hip/cupy_hip_runtime.h @@ -80,8 +80,12 @@ cudaError_t cudaDeviceGetLimit(size_t* pValue, cudaLimit limit) { } cudaError_t cudaDeviceSetLimit(cudaLimit limit, size_t value) { +#if HIP_VERSION >= 50300000 + return hipDeviceSetLimit(limit, value); +#else // see https://github.com/ROCm-Developer-Tools/HIP/issues/1632 return hipErrorUnknown; +#endif } // IPC operations @@ -90,11 +94,13 @@ cudaError_t cudaIpcCloseMemHandle(void* devPtr) { } cudaError_t cudaIpcGetEventHandle(cudaIpcEventHandle_t* handle, cudaEvent_t event) { - return hipErrorUnknown; - +#if HIP_VERSION >= 40300000 + return hipIpcGetEventHandle(handle, event); +#else // TODO(leofang): this is supported after ROCm-Developer-Tools/HIP#1996 is released; // as of ROCm 3.5.0 it is still not supported - //return hipIpcGetEventHandle(handle, event); + return hipErrorUnknown; +#endif } cudaError_t cudaIpcGetMemHandle(cudaIpcMemHandle_t* handle, void* devPtr) { @@ -102,11 +108,13 @@ cudaError_t cudaIpcGetMemHandle(cudaIpcMemHandle_t* handle, void* devPtr) { } cudaError_t cudaIpcOpenEventHandle(cudaEvent_t* event, cudaIpcEventHandle_t handle) { - return hipErrorUnknown; - +#if HIP_VERSION >= 40300000 + return hipIpcOpenEventHandle(event, handle); +#else // TODO(leofang): this is supported after ROCm-Developer-Tools/HIP#1996 is released; // as of ROCm 3.5.0 it is still not supported - //return hipIpcOpenEventHandle(event, handle); + return hipErrorUnknown; +#endif } cudaError_t cudaIpcOpenMemHandle(void** devPtr, cudaIpcMemHandle_t handle, unsigned int flags) { @@ -141,8 +149,12 @@ cudaError_t cudaMallocArray(...) { return hipErrorUnknown; } -cudaError_t cudaMallocAsync(...) { +cudaError_t cudaMallocAsync(void** ptr, size_t size, cudaStream_t stream) { +#if HIP_VERSION >= 60200000 + return hipMallocAsync(ptr, size, stream); +#else return hipErrorUnknown; +#endif } cudaError_t cudaHostAlloc(void** ptr, size_t size, unsigned int flags) { @@ -390,7 +402,11 @@ cudaError_t cudaStreamAddCallback(cudaStream_t stream, } cudaError_t cudaLaunchHostFunc(cudaStream_t stream, cudaHostFn_t fn, void* userData) { +#if HIP_VERSION >= 50300000 + return hipLaunchHostFunc(stream, fn, userData); +#else return hipErrorUnknown; +#endif } cudaError_t cudaStreamQuery(cudaStream_t stream) { @@ -542,8 +558,12 @@ cudaError_t cudaGraphLaunch(cudaGraphExec_t graphExec, cudaStream_t stream) { #endif } -cudaError_t cudaGraphUpload(...) { +cudaError_t cudaGraphUpload(cudaGraphExec_t graphExec, cudaStream_t stream) { +#if HIP_VERSION >= 60000000 + return hipGraphUpload(graphExec, stream); +#else return hipErrorUnknown; +#endif } cudaError_t cudaGraphDebugDotPrint(cudaGraph_t graph, const char* path, unsigned int flags) { diff --git a/cupy_backends/hip/cupy_hipblas.h b/cupy_backends/hip/cupy_hipblas.h index 0dc3d7891d1..83b36772894 100644 --- a/cupy_backends/hip/cupy_hipblas.h +++ b/cupy_backends/hip/cupy_hipblas.h @@ -2,7 +2,11 @@ #define INCLUDE_GUARD_HIP_CUPY_HIPBLAS_H #include "cupy_hip_common.h" +#if HIP_VERSION >= 50530600 +#include +#else #include +#endif #include // for HIP_VERSION #include // for gcc 10 diff --git a/cupy_backends/hip/cupy_hipsparse.h b/cupy_backends/hip/cupy_hipsparse.h index 5b2e9388ef2..db008ebcc0d 100644 --- a/cupy_backends/hip/cupy_hipsparse.h +++ b/cupy_backends/hip/cupy_hipsparse.h @@ -2,7 +2,11 @@ #ifndef INCLUDE_GUARD_HIP_CUPY_HIPSPARSE_H #define INCLUDE_GUARD_HIP_CUPY_HIPSPARSE_H +#if HIP_VERSION >= 50530600 +#include +#else #include +#endif #include // for HIP_VERSION #include // for hipDataType #include // for gcc 10.0 diff --git a/cupy_backends/hip/cupy_rccl.h b/cupy_backends/hip/cupy_rccl.h index 94dde35054b..428ade01321 100644 --- a/cupy_backends/hip/cupy_rccl.h +++ b/cupy_backends/hip/cupy_rccl.h @@ -1,7 +1,11 @@ #ifndef INCLUDE_GUARD_HIP_CUPY_RCCL_H #define INCLUDE_GUARD_HIP_CUPY_RCCL_H - +#include +#if HIP_VERSION >= 50530600 +#include +#else #include +#endif typedef hipStream_t cudaStream_t; #endif diff --git a/docker/rocm/Dockerfile b/docker/rocm/Dockerfile index 0bc2769a583..9457af1908d 100644 --- a/docker/rocm/Dockerfile +++ b/docker/rocm/Dockerfile @@ -2,6 +2,8 @@ FROM rocm/dev-ubuntu-20.04:5.0.1 LABEL maintainer="CuPy Team" RUN curl -qL https://repo.radeon.com/rocm/rocm.gpg.key | apt-key add - +# Workaround for missing amdgpu drivers +RUN echo "deb [arch=amd64] https://repo.radeon.com/amdgpu/23.20/amdgpu/ubuntu focal main" | tee /etc/apt/sources.list.d/amdgpu.list RUN apt-get update -y && \ DEBIAN_FRONTEND=noninteractive apt-get install -y --no-install-recommends \ hipblas hipsparse rocsparse rocrand rocthrust rocsolver rocfft hipfft hipcub rocprim rccl && \ diff --git a/install/cupy_builder/_features.py b/install/cupy_builder/_features.py index 6e473f194b3..4a7b0a1906b 100644 --- a/install/cupy_builder/_features.py +++ b/install/cupy_builder/_features.py @@ -157,6 +157,7 @@ def get_features(ctx: Context) -> Dict[str, Feature]: # the HIP stubs (hip/cupy_*.h) would cause many symbols # to leak into all these modules even if unused. It's easier for all of # them to link to the same set of shared libraries. + rocm_version = utils.get_rocm_version() HIP_cuda_nvtx_cusolver = { # TODO(leofang): call this "rocm" or "hip" to avoid confusion? 'name': 'cuda', @@ -169,12 +170,12 @@ def get_features(ctx: Context) -> Dict[str, Feature]: 'include': [ 'hip/hip_runtime_api.h', 'hip/hiprtc.h', - 'hipblas.h', + 'hipblas/hipblas.h' if rocm_version >= 560 else 'hipblas.h', 'hiprand/hiprand.h', - 'hipsparse.h', - 'hipfft.h', + 'hipsparse/hipsparse.h' if rocm_version >= 560 else 'hipsparse.h', + 'hipfft/hipfft.h' if rocm_version >= 560 else 'hipfft.h', 'roctx.h', - 'rocsolver.h', + 'rocsolver/rocsolver.h' if rocm_version >= 560 else 'rocsolver.h', ], 'libraries': [ 'amdhip64', # was hiprtc and hip_hcc before ROCm 3.8.0 @@ -366,7 +367,7 @@ def get_features(ctx: Context) -> Dict[str, Feature]: 'cupy_backends.cuda.libs.nccl', ], 'include': [ - 'rccl.h', + 'rccl/rccl.h' if rocm_version >= 560 else 'rccl.h', ], 'libraries': [ 'rccl', @@ -381,7 +382,9 @@ def get_features(ctx: Context) -> Dict[str, Feature]: ('cupy.cuda.thrust', ['cupy/cuda/cupy_thrust.cu']), ], 'include': [ - 'thrust/version.h', + # WAR #9098: + # rocThrust 3.3.0 (ROCm 6.4.0) cannot be compiled by host compiler + # 'thrust/version.h', ], 'libraries': [ 'amdhip64', # was hiprtc and hip_hcc before ROCm 3.8.0 diff --git a/install/cupy_builder/install_build.py b/install/cupy_builder/install_build.py index 40f75fd9349..ef374bb65a4 100644 --- a/install/cupy_builder/install_build.py +++ b/install/cupy_builder/install_build.py @@ -169,6 +169,8 @@ def get_compiler_setting(ctx: Context, use_hip): if use_hip: extra_compile_args.append('-std=c++11') + define_macros.append( + ('THRUST_DEVICE_SYSTEM', 'THRUST_DEVICE_SYSTEM_HIP')) if PLATFORM_WIN32: nvtx_path = _environment.get_nvtx_path() @@ -457,8 +459,13 @@ def check_nccl_version(compiler, settings): #ifndef CUPY_USE_HIP #include #else + #include + #if HIP_VERSION >= 50530600 + #include + #else #include #endif + #endif #include #ifdef NCCL_MAJOR #ifndef NCCL_VERSION_CODE diff --git a/install/cupy_builder/install_utils.py b/install/cupy_builder/install_utils.py index b7dd550b6bc..caf39c574f5 100644 --- a/install/cupy_builder/install_utils.py +++ b/install/cupy_builder/install_utils.py @@ -20,3 +20,13 @@ def search_on_path(filenames: List[str]) -> Optional[str]: if os.path.exists(full): return os.path.abspath(full) return None + + +def get_rocm_version() -> int: + rocm_version = -1 + if os.getenv("ROCM_HOME"): + rocm_home = str(os.getenv("ROCM_HOME")) + version_path = os.path.join(rocm_home, ".info", "version") + rocm_version = int( + open(version_path).read().split("-")[0].replace(".", "")) + return rocm_version diff --git a/tests/cupy_tests/core_tests/test_ndarray.py b/tests/cupy_tests/core_tests/test_ndarray.py index 594dde77431..77e30e3cafe 100644 --- a/tests/cupy_tests/core_tests/test_ndarray.py +++ b/tests/cupy_tests/core_tests/test_ndarray.py @@ -282,8 +282,6 @@ def test_shape_need_copy(self): assert 'incompatible shape' in str(e.value).lower() -@pytest.mark.skipif(cupy.cuda.runtime.is_hip, - reason='HIP does not support this') class TestNdarrayCudaInterface(unittest.TestCase): def test_cuda_array_interface(self): @@ -344,8 +342,6 @@ def test_cuda_array_interface_zero_size(self): 'stream': ('null', 'new', 'ptds'), 'ver': (2, 3), })) -@pytest.mark.skipif(cupy.cuda.runtime.is_hip, - reason='HIP does not support this') class TestNdarrayCudaInterfaceStream(unittest.TestCase): def setUp(self): if self.stream == 'null': @@ -387,20 +383,22 @@ def test_cuda_array_interface_stream(self): assert iface['stream'] == stream.ptr -@pytest.mark.skipif(not cupy.cuda.runtime.is_hip, - reason='This is supported on CUDA') -class TestNdarrayCudaInterfaceNoneCUDA(unittest.TestCase): +class TestNdarrayCudaInterfaceAttr(unittest.TestCase): def setUp(self): self.arr = cupy.zeros(shape=(2, 3), dtype=cupy.float64) def test_cuda_array_interface_hasattr(self): - assert not hasattr(self.arr, '__cuda_array_interface__') + assert hasattr(self.arr, '__cuda_array_interface__') def test_cuda_array_interface_getattr(self): - with pytest.raises(AttributeError) as e: - getattr(self.arr, '__cuda_array_interface__') - assert 'HIP' in str(e.value) + try: + interface = getattr(self.arr, '__cuda_array_interface__') + self.assertIsNotNone(interface) + except AttributeError as e: + self.fail( + "__cuda_array_interface__ should be present, " + f"but got an AttributeError: {str(e)}") @testing.parameterize( diff --git a/tests/cupy_tests/core_tests/test_ndarray_cuda_array_interface.py b/tests/cupy_tests/core_tests/test_ndarray_cuda_array_interface.py index 822f121f070..a25316abc42 100644 --- a/tests/cupy_tests/core_tests/test_ndarray_cuda_array_interface.py +++ b/tests/cupy_tests/core_tests/test_ndarray_cuda_array_interface.py @@ -1,8 +1,8 @@ import unittest -import pytest from cupy_backends.cuda import stream as stream_module import cupy +import cupy_backends from cupy import _core from cupy import testing @@ -28,7 +28,11 @@ def __cuda_array_interface__(self): } if self.ver == 3: stream = cupy.cuda.get_current_stream() - desc['stream'] = 1 if stream.ptr == 0 else stream.ptr + # Only non-default streams use their actual ptr values. (ROCm) + if cupy_backends.cuda.api.runtime.is_hip: + desc['stream'] = stream.ptr + else: + desc['stream'] = 1 if stream.ptr == 0 else stream.ptr # noqa: F821, E501 return desc @@ -37,8 +41,6 @@ def __cuda_array_interface__(self): 'ver': (2, 3), })) @testing.with_requires('numpy>=1.25') -@pytest.mark.skipif(cupy.cuda.runtime.is_hip, - reason='HIP does not support this') class TestArrayUfunc(unittest.TestCase): def setUp(self): @@ -73,8 +75,6 @@ def test_add_scalar_with_strides(self): 'stream': ('null', 'new'), 'ver': (2, 3), })) -@pytest.mark.skipif(cupy.cuda.runtime.is_hip, - reason='HIP does not support this') class TestElementwiseKernel(unittest.TestCase): def setUp(self): @@ -110,8 +110,6 @@ def test_add_scalar_with_strides(self): 'stream': ('null', 'new'), 'ver': (2, 3), })) -@pytest.mark.skipif(cupy.cuda.runtime.is_hip, - reason='HIP does not support this') class TestSimpleReductionFunction(unittest.TestCase): def setUp(self): @@ -149,8 +147,6 @@ def test_shape_with_strides(self): 'stream': ('null', 'new'), 'ver': (2, 3), })) -@pytest.mark.skipif(cupy.cuda.runtime.is_hip, - reason='HIP does not support this') class TestReductionKernel(unittest.TestCase): def setUp(self): @@ -193,8 +189,6 @@ def test_shape_with_strides(self): {'shape': (10, 10), 'slices': (slice(2, None), slice(2, None))}, {'shape': (10, 10), 'slices': (slice(2, None), slice(4, None))}, ) -@pytest.mark.skipif(cupy.cuda.runtime.is_hip, - reason='HIP does not support this') class TestSlicingMemoryPointer(unittest.TestCase): @testing.for_all_dtypes_combination(names=['dtype']) @@ -238,8 +232,6 @@ def test_shape_with_strides(self, dtype, order): @testing.parameterize(*test_cases_with_stream) -@pytest.mark.skipif(cupy.cuda.runtime.is_hip, - reason='HIP does not support this') class TestCUDAArrayInterfaceCompliance(unittest.TestCase): def setUp(self): @@ -285,8 +277,6 @@ def test_value_type(self, dtype, order): @testing.parameterize(*testing.product({ 'stream': ('null', 'new', 'ptds'), })) -@pytest.mark.skipif(cupy.cuda.runtime.is_hip, - reason='HIP does not support this') class TestCUDAArrayInterfaceStream(unittest.TestCase): def setUp(self): if self.stream == 'null': diff --git a/tests/cupy_tests/core_tests/test_ndarray_reduction.py b/tests/cupy_tests/core_tests/test_ndarray_reduction.py index 25f1b346d62..191cc51f8d3 100644 --- a/tests/cupy_tests/core_tests/test_ndarray_reduction.py +++ b/tests/cupy_tests/core_tests/test_ndarray_reduction.py @@ -5,6 +5,7 @@ import cupy._core._accelerator as _acc from cupy._core import _cub_reduction from cupy import testing +from cupy.cuda import runtime @testing.parameterize(*testing.product({ @@ -369,6 +370,7 @@ def setUp(self): @testing.for_all_dtypes(no_bool=True) @testing.numpy_cupy_allclose( contiguous_check=False, accept_error=ValueError) + @pytest.mark.skipif(runtime.is_hip, reason='ROCm/HIP may have a bug') def test_cub_min(self, xp, dtype, axis): a = testing.shaped_random(self.shape, xp, dtype, order=self.order) @@ -414,6 +416,7 @@ def test_cub_min_empty_axis(self, xp, dtype, contiguous_check=False): @testing.for_all_dtypes(no_bool=True) @testing.numpy_cupy_allclose( contiguous_check=False, accept_error=ValueError) + @pytest.mark.skipif(runtime.is_hip, reason='ROCm/HIP may have a bug') def test_cub_max(self, xp, dtype, axis): a = testing.shaped_random(self.shape, xp, dtype, order=self.order) diff --git a/tests/cupy_tests/creation_tests/test_from_data.py b/tests/cupy_tests/creation_tests/test_from_data.py index 77c10fa57ba..22539215f4f 100644 --- a/tests/cupy_tests/creation_tests/test_from_data.py +++ b/tests/cupy_tests/creation_tests/test_from_data.py @@ -586,8 +586,6 @@ def test_fromfile_big_endian(self, xp): 'ver': tuple(range(max_cuda_array_interface_version+1)), 'strides': (False, None, True), })) -@pytest.mark.skipif( - cupy.cuda.runtime.is_hip, reason='HIP does not support this') class TestCudaArrayInterface(unittest.TestCase): @testing.for_all_dtypes() def test_base(self, dtype): @@ -656,8 +654,6 @@ def test_big_endian(self): 'ver': tuple(range(1, max_cuda_array_interface_version+1)), 'strides': (False, None, True), })) -@pytest.mark.skipif( - cupy.cuda.runtime.is_hip, reason='HIP does not support this') class TestCudaArrayInterfaceMaskedArray(unittest.TestCase): # TODO(leofang): update this test when masked array is supported @testing.for_all_dtypes() @@ -672,8 +668,6 @@ def test_masked_array(self, dtype): # marked slow as either numpy or cupy could go OOM in this test @testing.slow -@pytest.mark.skipif( - cupy.cuda.runtime.is_hip, reason='HIP does not support this') class TestCudaArrayInterfaceBigArray(unittest.TestCase): def test_with_over_size_array(self): # real example from #3009 @@ -683,8 +677,6 @@ def test_with_over_size_array(self): testing.assert_array_equal(a, b) -@pytest.mark.skipif( - cupy.cuda.runtime.is_hip, reason='HIP does not support this') class DummyObjectWithCudaArrayInterface(object): def __init__(self, a, ver, include_strides=False, mask=None, stream=None): assert ver in tuple(range(max_cuda_array_interface_version+1)) diff --git a/tests/cupy_tests/cuda_tests/test_graph.py b/tests/cupy_tests/cuda_tests/test_graph.py index 0e0710ff25b..b0aea06f85a 100644 --- a/tests/cupy_tests/cuda_tests/test_graph.py +++ b/tests/cupy_tests/cuda_tests/test_graph.py @@ -3,11 +3,10 @@ import cupy from cupy import cuda from cupy import testing +from cupy_backends.cuda.api import runtime import cupyx -@pytest.mark.skipif(cuda.runtime.is_hip, - reason='HIP does not support this') class TestGraph: def _helper1(self, a): @@ -143,7 +142,8 @@ def test_null_stream_cannot_capture(self, upload): # cudaStreamLegacy is unhappy when a blocking stream is capturing with pytest.raises(cuda.runtime.CUDARuntimeError) as e: cuda.Stream.null.is_capturing() - assert 'cudaErrorStreamCaptureImplicit' in str(e.value) + assert ('hipErrorStreamCaptureImplicit' if runtime.is_hip + else 'cudaErrorStreamCaptureImplicit') in str(e.value) g = s.end_capture() assert not s.is_capturing() assert not cuda.Stream.null.is_capturing() @@ -155,6 +155,8 @@ def test_null_stream_cannot_capture(self, upload): s.synchronize() testing.assert_array_equal(b, a + 4) + @pytest.mark.skipif(cuda.runtime.is_hip, + reason='HIP does not support this') def test_stream_capture_failure1(self): s = cupy.cuda.Stream(non_blocking=True) @@ -162,16 +164,20 @@ def test_stream_capture_failure1(self): s.begin_capture() with pytest.raises(cuda.runtime.CUDARuntimeError) as e: s.synchronize() - assert 'cudaErrorStreamCaptureUnsupported' in str(e.value) + assert ('hipErrorStreamCaptureUnsupported' if runtime.is_hip + else 'cudaErrorStreamCaptureUnsupported') in str(e.value) # invalid operation causes the capture sequence to be invalidated with pytest.raises(cuda.runtime.CUDARuntimeError) as e: g = s.end_capture() # noqa - assert 'cudaErrorStreamCaptureInvalidated' in str(e.value) + assert ('hipErrorStreamCaptureInvalidated' if runtime.is_hip + else 'cudaErrorStreamCaptureInvalidated') in str(e.value) # check s left the capture mode and permits normal usage assert not s.is_capturing() s.synchronize() + @pytest.mark.skipif(cuda.runtime.is_hip, + reason='HIP does not support this') def test_stream_capture_failure2(self): s1 = cupy.cuda.Stream(non_blocking=True) s2 = cupy.cuda.Stream(non_blocking=True) @@ -182,18 +188,21 @@ def test_stream_capture_failure2(self): s1.begin_capture() with pytest.raises(cuda.runtime.CUDARuntimeError) as e: g = s2.end_capture() - assert 'cudaErrorIllegalState' in str(e.value) + assert ('hipErrorIllegalState' if runtime.is_hip + else 'cudaErrorIllegalState') in str(e.value) e2.record(s1) s2.wait_event(e2) with s2: b = a**3 # noqa with pytest.raises(cuda.runtime.CUDARuntimeError) as e: g = s2.end_capture() - assert 'cudaErrorStreamCaptureUnmatched' in str(e.value) + assert ('hipErrorStreamCaptureUnmatched' if runtime.is_hip + else 'cudaErrorStreamCaptureUnmatched') in str(e.value) # invalid operation causes the capture sequence to be invalidated with pytest.raises(cuda.runtime.CUDARuntimeError) as e: g = s1.end_capture() # noqa - assert 'cudaErrorStreamCaptureInvalidated' in str(e.value) + assert ('hipErrorStreamCaptureInvalidated' if runtime.is_hip + else 'cudaErrorStreamCaptureInvalidated') in str(e.value) # check both s1 and s2 left the capture mode and permit normal usage assert not s1.is_capturing() @@ -201,6 +210,8 @@ def test_stream_capture_failure2(self): s1.synchronize() s2.synchronize() + @pytest.mark.skipif(cuda.runtime.is_hip, + reason='HIP does not support this') def test_stream_capture_failure3(self): s1 = cupy.cuda.Stream(non_blocking=True) s2 = cupy.cuda.Stream(non_blocking=True) @@ -220,7 +231,8 @@ def test_stream_capture_failure3(self): # invalid operation causes the capture sequence to be invalidated with pytest.raises(cuda.runtime.CUDARuntimeError) as e: g = s1.end_capture() # noqa - assert 'cudaErrorStreamCaptureUnjoined' in str(e.value) + assert ('hipErrorStreamCaptureUnjoined' if runtime.is_hip + else 'cudaErrorStreamCaptureUnjoined') in str(e.value) # check both s1 and s2 left the capture mode and permit normal usage assert not s1.is_capturing() @@ -228,6 +240,8 @@ def test_stream_capture_failure3(self): s1.synchronize() s2.synchronize() + @pytest.mark.skipif(cuda.runtime.is_hip, + reason='HIP does not support this') def test_stream_capture_failure4(self): s = cupy.cuda.Stream(non_blocking=True) @@ -237,7 +251,8 @@ def test_stream_capture_failure4(self): s.done with pytest.raises(cuda.runtime.CUDARuntimeError) as e: s.end_capture() - assert 'cudaErrorStreamCaptureInvalidated' in str(e.value) + assert ('hipErrorStreamCaptureImplicit' if runtime.is_hip + else 'cudaErrorStreamCaptureInvalidated') in str(e.value) # check s left the capture mode and permits normal usage assert not s.is_capturing() @@ -261,6 +276,8 @@ def test_stream_capture_failure5(self): assert not s.is_capturing() s.synchronize() + @pytest.mark.skipif(cuda.runtime.is_hip, + reason='HIP does not support this') def test_stream_capture_failure6(self): s = cupy.cuda.Stream(non_blocking=True) @@ -269,10 +286,12 @@ def test_stream_capture_failure6(self): # synchronize the stream is illegal during capturing with pytest.raises(cuda.runtime.CUDARuntimeError) as e: s.synchronize() - assert 'cudaErrorStreamCaptureUnsupported' in str(e.value) + assert ('hipErrorStreamCaptureUnsupported' if runtime.is_hip + else 'cudaErrorStreamCaptureUnsupported') in str(e.value) with pytest.raises(cuda.runtime.CUDARuntimeError) as e: s.end_capture() - assert 'cudaErrorStreamCaptureInvalidated' in str(e.value) + assert ('hipErrorStreamCaptureInvalidated' if runtime.is_hip + else 'cudaErrorStreamCaptureInvalidated') in str(e.value) # check s left the capture mode and permits normal usage assert not s.is_capturing() diff --git a/tests/cupy_tests/cuda_tests/test_nccl.py b/tests/cupy_tests/cuda_tests/test_nccl.py index 99441edc2a0..5e649bc49aa 100644 --- a/tests/cupy_tests/cuda_tests/test_nccl.py +++ b/tests/cupy_tests/cuda_tests/test_nccl.py @@ -1,11 +1,12 @@ import pickle import unittest +import pytest import cupy from cupy import cuda from cupy.cuda import nccl from cupy import testing - +from cupy.cuda import runtime nccl_available = nccl.available @@ -64,6 +65,8 @@ def test_comm_size(self): @testing.multi_gpu(2) @unittest.skipUnless(nccl_version >= 2700, 'Using old NCCL') + @pytest.mark.skipif(runtime.is_hip, + reason="rccl doesn't support multi-GPU") def test_send_recv(self): devs = [0, 1] comms = nccl.NcclCommunicator.initAll(devs) diff --git a/tests/cupy_tests/fft_tests/test_cache.py b/tests/cupy_tests/fft_tests/test_cache.py index b80da868d11..4d5187f3a1d 100644 --- a/tests/cupy_tests/fft_tests/test_cache.py +++ b/tests/cupy_tests/fft_tests/test_cache.py @@ -146,6 +146,8 @@ def test_LRU_cache4(self): cache[next(iterator)[0]] @testing.multi_gpu(2) + @pytest.mark.skipif(runtime.is_hip, + reason="hipFFT doesn't support multi-GPU") def test_LRU_cache5(self): # test if the LRU cache is thread-local @@ -198,6 +200,8 @@ def thread_init_caches(gpus, queue): assert stdout.count('uninitialized') == n_devices - 2 @testing.multi_gpu(2) + @pytest.mark.skipif(runtime.is_hip, + reason="hipFFT doesn't support multi-GPU") def test_LRU_cache6(self): # test if each device has a separate cache cache0 = self.caches[0] diff --git a/tests/cupy_tests/manipulation_tests/test_basic.py b/tests/cupy_tests/manipulation_tests/test_basic.py index 314d3c0d49f..d3a431f2ff0 100644 --- a/tests/cupy_tests/manipulation_tests/test_basic.py +++ b/tests/cupy_tests/manipulation_tests/test_basic.py @@ -7,6 +7,7 @@ import cupy from cupy import cuda from cupy import testing +from cupy.cuda import runtime from cupy.exceptions import ComplexWarning @@ -143,22 +144,26 @@ def get_numpy(): @testing.multi_gpu(2) @testing.for_all_dtypes() + @pytest.mark.skipif(runtime.is_hip, reason='ROCm/HIP may have a bug') def test_copyto_where_multigpu_raises(self, dtype): self._check_copyto_where_multigpu_raises(dtype, 2) @testing.multi_gpu(4) @testing.for_all_dtypes() + @pytest.mark.skipif(runtime.is_hip, reason='ROCm/HIP may have a bug') def test_copyto_where_multigpu_raises_4(self, dtype): self._check_copyto_where_multigpu_raises(dtype, 4) @testing.multi_gpu(6) @testing.for_all_dtypes() + @pytest.mark.skipif(runtime.is_hip, reason='ROCm/HIP may have a bug') def test_copyto_where_multigpu_raises_6(self, dtype): self._check_copyto_where_multigpu_raises(dtype, 6) @testing.multi_gpu(2) @testing.for_all_dtypes() @testing.numpy_cupy_array_equal() + @pytest.mark.skipif(runtime.is_hip, reason='ROCm/HIP may have a bug') def test_copyto_multigpu(self, xp, dtype): with cuda.Device(0): a = testing.shaped_arange((2, 3, 4), xp, dtype) @@ -169,6 +174,7 @@ def test_copyto_multigpu(self, xp, dtype): @testing.multi_gpu(2) @testing.for_all_dtypes() + @pytest.mark.skipif(runtime.is_hip, reason='ROCm/HIP may have a bug') def test_copyto_multigpu_noncontinguous(self, dtype): with cuda.Device(0): src = testing.shaped_arange((2, 3, 4), cupy, dtype) diff --git a/tests/cupy_tests/manipulation_tests/test_join.py b/tests/cupy_tests/manipulation_tests/test_join.py index 75a3f9e36f4..eb085cf836a 100644 --- a/tests/cupy_tests/manipulation_tests/test_join.py +++ b/tests/cupy_tests/manipulation_tests/test_join.py @@ -4,6 +4,7 @@ import cupy from cupy import testing from cupy import cuda +from cupy.cuda import runtime from cupy.exceptions import AxisError @@ -38,6 +39,7 @@ def test_column_stack_wrong_shape(self): @testing.for_all_dtypes(name='dtype') @testing.numpy_cupy_array_equal() + @pytest.mark.skipif(runtime.is_hip, reason='ROCm/HIP may have a bug ') def test_concatenate1(self, xp, dtype): a = testing.shaped_arange((2, 3, 4), xp, dtype) b = testing.shaped_reverse_arange((2, 3, 2), xp, dtype) @@ -46,6 +48,7 @@ def test_concatenate1(self, xp, dtype): @testing.for_all_dtypes(name='dtype') @testing.numpy_cupy_array_equal() + @pytest.mark.skipif(runtime.is_hip, reason='ROCm/HIP may have a bug ') def test_concatenate2(self, xp, dtype): a = testing.shaped_arange((2, 3, 4), xp, dtype) b = testing.shaped_reverse_arange((2, 3, 2), xp, dtype) @@ -54,6 +57,7 @@ def test_concatenate2(self, xp, dtype): @testing.for_all_dtypes(name='dtype') @testing.numpy_cupy_array_equal() + @pytest.mark.skipif(runtime.is_hip, reason='ROCm/HIP may have a bug ') def test_concatenate_axis_none(self, xp, dtype): a = testing.shaped_arange((2, 3), xp, dtype) b = testing.shaped_reverse_arange((3, 5, 2), xp, dtype) @@ -62,6 +66,7 @@ def test_concatenate_axis_none(self, xp, dtype): @testing.for_all_dtypes(name='dtype') @testing.numpy_cupy_array_equal() + @pytest.mark.skipif(runtime.is_hip, reason='ROCm/HIP may have a bug ') def test_concatenate_large_2(self, xp, dtype): a = testing.shaped_arange((2, 3, 4), xp, dtype) b = testing.shaped_reverse_arange((2, 3, 2), xp, dtype) @@ -72,6 +77,7 @@ def test_concatenate_large_2(self, xp, dtype): @testing.for_all_dtypes(name='dtype') @testing.numpy_cupy_array_equal() + @pytest.mark.skipif(runtime.is_hip, reason='ROCm/HIP may have a bug ') def test_concatenate_large_3(self, xp, dtype): a = testing.shaped_arange((2, 3, 1), xp, dtype) b = testing.shaped_reverse_arange((2, 3, 1), xp, dtype) @@ -79,6 +85,7 @@ def test_concatenate_large_3(self, xp, dtype): @testing.for_all_dtypes(name='dtype') @testing.numpy_cupy_array_equal() + @pytest.mark.skipif(runtime.is_hip, reason='ROCm/HIP may have a bug ') def test_concatenate_large_4(self, xp, dtype): a = testing.shaped_arange((2, 3, 4), xp, dtype) b = testing.shaped_reverse_arange((2, 3, 4), xp, dtype) @@ -86,12 +93,14 @@ def test_concatenate_large_4(self, xp, dtype): @testing.for_all_dtypes(name='dtype') @testing.numpy_cupy_array_equal() + @pytest.mark.skipif(runtime.is_hip, reason='ROCm/HIP may have a bug ') def test_concatenate_large_5(self, xp, dtype): a = testing.shaped_arange((2, 3, 4), xp, dtype) b = testing.shaped_reverse_arange((2, 3, 4), xp, 'i') return xp.concatenate((a, b) * 10, axis=-1) @testing.multi_gpu(2) + @pytest.mark.skipif(runtime.is_hip, reason='ROCm/HIP may have a bug ') def test_concatenate_large_different_devices(self): arrs = [] for i in range(10): @@ -106,6 +115,7 @@ def test_concatenate_large_different_devices(self): @testing.for_all_dtypes(name='dtype') @testing.numpy_cupy_array_equal() + @pytest.mark.skipif(runtime.is_hip, reason='ROCm/HIP may have a bug ') def test_concatenate_f_contiguous(self, xp, dtype): a = testing.shaped_arange((2, 3, 4), xp, dtype) b = testing.shaped_arange((2, 3, 2), xp, dtype).T @@ -114,6 +124,7 @@ def test_concatenate_f_contiguous(self, xp, dtype): @testing.for_all_dtypes(name='dtype') @testing.numpy_cupy_array_equal() + @pytest.mark.skipif(runtime.is_hip, reason='ROCm/HIP may have a bug ') def test_concatenate_large_f_contiguous(self, xp, dtype): a = testing.shaped_arange((2, 3, 4), xp, dtype) b = testing.shaped_arange((2, 3, 2), xp, dtype).T @@ -123,12 +134,14 @@ def test_concatenate_large_f_contiguous(self, xp, dtype): return xp.concatenate((a, b, c, d, e) * 2, axis=-1) @testing.numpy_cupy_array_equal() + @pytest.mark.skipif(runtime.is_hip, reason='ROCm/HIP may have a bug ') def test_concatenate_many_multi_dtype(self, xp): a = testing.shaped_arange((2, 1), xp, 'i') b = testing.shaped_arange((2, 1), xp, 'f') return xp.concatenate((a, b) * 1024, axis=1) @testing.slow + @pytest.mark.skipif(runtime.is_hip, reason='ROCm/HIP may have a bug ') def test_concatenate_32bit_boundary(self): a = cupy.zeros((2 ** 30,), dtype=cupy.int8) b = cupy.zeros((2 ** 30,), dtype=cupy.int8) @@ -139,12 +152,14 @@ def test_concatenate_32bit_boundary(self): # Free huge memory for slow test cupy.get_default_memory_pool().free_all_blocks() + @pytest.mark.skipif(runtime.is_hip, reason='ROCm/HIP may have a bug ') def test_concatenate_wrong_ndim(self): a = cupy.empty((2, 3)) b = cupy.empty((2,)) with pytest.raises(ValueError): cupy.concatenate((a, b)) + @pytest.mark.skipif(runtime.is_hip, reason='ROCm/HIP may have a bug ') def test_concatenate_wrong_shape(self): a = cupy.empty((2, 3, 4)) b = cupy.empty((3, 3, 4)) @@ -154,6 +169,7 @@ def test_concatenate_wrong_shape(self): @testing.for_all_dtypes(name='dtype') @testing.numpy_cupy_array_equal() + @pytest.mark.skipif(runtime.is_hip, reason='ROCm/HIP may have a bug ') def test_concatenate_out(self, xp, dtype): a = testing.shaped_arange((3, 4), xp, dtype) b = testing.shaped_reverse_arange((3, 4), xp, dtype) @@ -163,6 +179,7 @@ def test_concatenate_out(self, xp, dtype): return out @testing.numpy_cupy_array_equal() + @pytest.mark.skipif(runtime.is_hip, reason='ROCm/HIP may have a bug ') def test_concatenate_out_same_kind(self, xp): a = testing.shaped_arange((3, 4), xp, xp.float64) b = testing.shaped_reverse_arange((3, 4), xp, xp.float64) @@ -171,6 +188,7 @@ def test_concatenate_out_same_kind(self, xp): xp.concatenate((a, b, c), axis=1, out=out) return out + @pytest.mark.skipif(runtime.is_hip, reason='ROCm/HIP may have a bug ') def test_concatenate_out_invalid_shape(self): for xp in (numpy, cupy): a = testing.shaped_arange((3, 4), xp, xp.float64) @@ -180,6 +198,7 @@ def test_concatenate_out_invalid_shape(self): with pytest.raises(ValueError): xp.concatenate((a, b, c), axis=1, out=out) + @pytest.mark.skipif(runtime.is_hip, reason='ROCm/HIP may have a bug ') def test_concatenate_out_invalid_shape_2(self): for xp in (numpy, cupy): a = testing.shaped_arange((3, 4), xp, xp.float64) @@ -189,6 +208,7 @@ def test_concatenate_out_invalid_shape_2(self): with pytest.raises(ValueError): xp.concatenate((a, b, c), axis=1, out=out) + @pytest.mark.skipif(runtime.is_hip, reason='ROCm/HIP may have a bug ') def test_concatenate_out_invalid_dtype(self): for xp in (numpy, cupy): a = testing.shaped_arange((3, 4), xp, xp.float64) @@ -200,6 +220,7 @@ def test_concatenate_out_invalid_dtype(self): @testing.for_all_dtypes_combination(names=['dtype1', 'dtype2']) @testing.numpy_cupy_array_equal() + @pytest.mark.skipif(runtime.is_hip, reason='ROCm/HIP may have a bug ') def test_concatenate_different_dtype(self, xp, dtype1, dtype2): a = testing.shaped_arange((3, 4), xp, dtype1) b = testing.shaped_arange((3, 4), xp, dtype2) @@ -207,6 +228,7 @@ def test_concatenate_different_dtype(self, xp, dtype1, dtype2): @testing.for_all_dtypes_combination(names=['dtype1', 'dtype2']) @testing.numpy_cupy_array_equal(accept_error=TypeError) + @pytest.mark.skipif(runtime.is_hip, reason='ROCm/HIP may have a bug ') def test_concatenate_out_different_dtype(self, xp, dtype1, dtype2): a = testing.shaped_arange((3, 4), xp, dtype1) b = testing.shaped_arange((3, 4), xp, dtype1) @@ -216,12 +238,14 @@ def test_concatenate_out_different_dtype(self, xp, dtype1, dtype2): @testing.with_requires('numpy>=1.20.0') @testing.for_all_dtypes_combination(names=['dtype1', 'dtype2']) @testing.numpy_cupy_array_equal(accept_error=TypeError) + @pytest.mark.skipif(runtime.is_hip, reason='ROCm/HIP may have a bug ') def test_concatenate_dtype(self, xp, dtype1, dtype2): a = testing.shaped_arange((3, 4), xp, dtype1) b = testing.shaped_arange((3, 4), xp, dtype1) return xp.concatenate((a, b), dtype=dtype2) @testing.with_requires('numpy>=1.20.0') + @pytest.mark.skipif(runtime.is_hip, reason='ROCm/HIP may have a bug ') def test_concatenate_dtype_invalid_out(self): for xp in (numpy, cupy): a = testing.shaped_arange((3, 4), xp, xp.float64) @@ -242,6 +266,7 @@ def test_concatenate_dtype_invalid_out(self): @testing.for_all_dtypes_combination(names=['dtype1', 'dtype2']) @testing.numpy_cupy_array_equal( accept_error=(TypeError, cupy.exceptions.ComplexWarning)) + @pytest.mark.skipif(runtime.is_hip, reason='ROCm/HIP may have a bug ') def test_concatenate_casting(self, xp, dtype1, dtype2, casting): a = testing.shaped_arange((3, 4), xp, dtype1) b = testing.shaped_arange((3, 4), xp, dtype1) diff --git a/tests/cupy_tests/math_tests/test_sumprod.py b/tests/cupy_tests/math_tests/test_sumprod.py index 5abc61a83d5..d8e80a52caa 100644 --- a/tests/cupy_tests/math_tests/test_sumprod.py +++ b/tests/cupy_tests/math_tests/test_sumprod.py @@ -8,6 +8,7 @@ import cupy.cuda.cutensor from cupy._core import _cub_reduction from cupy import testing +from cupy.cuda import runtime from cupy.exceptions import AxisError @@ -231,6 +232,7 @@ def setUp(self): # sum supports less dtypes; don't test float16 as it's not as accurate? @testing.for_dtypes('qQfdFD') @testing.numpy_cupy_allclose(rtol=1E-5) + @pytest.mark.skipif(runtime.is_hip, reason='ROCm/HIP may have a bug') def test_cub_sum(self, xp, dtype, axis): a = testing.shaped_random(self.shape, xp, dtype) if self.order in ('c', 'C'): @@ -281,6 +283,7 @@ def test_cub_sum_empty_axis(self, xp, dtype): # prod supports less dtypes; don't test float16 as it's not as accurate? @testing.for_dtypes('qQfdFD') @testing.numpy_cupy_allclose(rtol=1E-5) + @pytest.mark.skipif(runtime.is_hip, reason='ROCm/HIP may have a bug') def test_cub_prod(self, xp, dtype, axis): a = testing.shaped_random(self.shape, xp, dtype) if self.order in ('c', 'C'): @@ -717,6 +720,9 @@ def test_ndarray_cumprod_2dim_with_axis(self, xp, dtype): return a.cumprod(axis=1) @testing.slow + @pytest.mark.xfail( + runtime.is_hip, + reason='Workload size is bigger than what ROCm/CUDA supports') def test_cumprod_huge_array(self): size = 2 ** 32 # Free huge memory for slow test diff --git a/tests/cupy_tests/random_tests/test_generator.py b/tests/cupy_tests/random_tests/test_generator.py index a83b012afce..a5e98a0be91 100644 --- a/tests/cupy_tests/random_tests/test_generator.py +++ b/tests/cupy_tests/random_tests/test_generator.py @@ -13,6 +13,7 @@ from cupy import testing from cupy.testing import _condition from cupy.testing import _hypothesis +from cupy.cuda import driver from cupy_tests.random_tests import common_distributions @@ -141,7 +142,8 @@ def test_methods(self): for method in methods: if (runtime.is_hip and - method == cupy.cuda.curand.CURAND_RNG_PSEUDO_MT19937): + method == cupy.cuda.curand.CURAND_RNG_PSEUDO_MT19937 + and driver.get_build_version() < 50530201): # hipRAND fails for MT19937 with the status code 1000, # HIPRAND_STATUS_NOT_IMPLEMENTED. We use `pytest.raises` here # so that we will be able to find it once hipRAND implement diff --git a/tests/cupy_tests/random_tests/test_sample.py b/tests/cupy_tests/random_tests/test_sample.py index 83b1bcf7698..e27a4e0e391 100644 --- a/tests/cupy_tests/random_tests/test_sample.py +++ b/tests/cupy_tests/random_tests/test_sample.py @@ -2,11 +2,9 @@ from unittest import mock import numpy -import pytest import cupy from cupy import cuda -from cupy.cuda import runtime from cupy import random from cupy import testing from cupy.testing import _condition @@ -98,7 +96,6 @@ def test_goodness_of_fit(self): assert _hypothesis.chi_square_test(counts, expected) @_condition.repeat(3, 10) - @pytest.mark.xfail(runtime.is_hip, reason='ROCm/HIP may have a bug') def test_goodness_of_fit_2(self): mx = 5 vals = random.randint(mx, size=(5, 20)).get() @@ -192,7 +189,6 @@ def test_goodness_of_fit(self): assert _hypothesis.chi_square_test(counts, expected) @_condition.repeat(3, 10) - @pytest.mark.xfail(runtime.is_hip, reason='ROCm/HIP may have a bug') def test_goodness_of_fit_2(self): mx = 5 vals = random.randint(0, mx, (5, 20)).get() diff --git a/tests/cupy_tests/sorting_tests/test_search.py b/tests/cupy_tests/sorting_tests/test_search.py index a96c5e3dded..f0752111939 100644 --- a/tests/cupy_tests/sorting_tests/test_search.py +++ b/tests/cupy_tests/sorting_tests/test_search.py @@ -5,6 +5,7 @@ import cupy._core._accelerator as _acc from cupy._core import _cub_reduction from cupy import testing +from cupy.cuda import runtime class TestSearch: @@ -206,6 +207,7 @@ def setUp(self): @testing.for_dtypes('bhilBHILefdFD') @testing.numpy_cupy_allclose(rtol=1E-5, contiguous_check=False) + @pytest.mark.skipif(runtime.is_hip, reason='ROCm/HIP may have a bug') def test_cub_argmin(self, xp, dtype): _skip_cuda90(dtype) a = testing.shaped_random(self.shape, xp, dtype) @@ -241,6 +243,7 @@ def test_cub_argmin(self, xp, dtype): @testing.for_dtypes('bhilBHILefdFD') @testing.numpy_cupy_allclose(rtol=1E-5, contiguous_check=False) + @pytest.mark.skipif(runtime.is_hip, reason='ROCm/HIP may have a bug') def test_cub_argmax(self, xp, dtype): _skip_cuda90(dtype) a = testing.shaped_random(self.shape, xp, dtype) diff --git a/tests/cupyx_tests/scipy_tests/ndimage_tests/test_filters.py b/tests/cupyx_tests/scipy_tests/ndimage_tests/test_filters.py index 3a6eea3803f..16c48b3545e 100644 --- a/tests/cupyx_tests/scipy_tests/ndimage_tests/test_filters.py +++ b/tests/cupyx_tests/scipy_tests/ndimage_tests/test_filters.py @@ -230,6 +230,7 @@ def _kshape(self): ) )) @testing.with_requires('scipy') +@pytest.mark.skipif(runtime.is_hip, reason='ROCm/HIP may have a bug') class TestFilter(FilterTestCaseBase): def _hip_skip_invalid_condition(self): @@ -270,6 +271,7 @@ def test_filter(self, xp, scp): @testing.with_requires('scipy') +@pytest.mark.skipif(runtime.is_hip, reason='ROCm/HIP may have a bug') class TestNearestFilterEdgeCase: @testing.numpy_cupy_allclose(atol=1e-5, rtol=1e-5, scipy_name='scp') @@ -345,6 +347,7 @@ def dummy_deriv_func(input, axis, output, mode, cval, *args, **kwargs): ) )) @testing.with_requires('scipy') +@pytest.mark.skipif(runtime.is_hip, reason='ROCm/HIP may have a bug') class TestFilterFast(FilterTestCaseBase): def _hip_skip_invalid_condition(self): @@ -502,6 +505,7 @@ def test_filter(self, xp, scp): ) )) @testing.with_requires('scipy') +@pytest.mark.skipif(runtime.is_hip, reason='ROCm/HIP may have a bug') class TestFilterComplexFast(FilterTestCaseBase): @testing.numpy_cupy_allclose(atol=1e-5, rtol=1e-5, scipy_name='scp') @@ -582,6 +586,7 @@ def lt_pyfunc(x): }) )) @testing.with_requires('scipy') +@pytest.mark.skipif(runtime.is_hip, reason='ROCm/HIP may have a bug') class TestGenericFilter(FilterTestCaseBase): _func_or_kernels = { @@ -647,6 +652,7 @@ def shift_pyfunc(src, dst): ) )) @testing.with_requires('scipy') +@pytest.mark.skipif(runtime.is_hip, reason='ROCm/HIP may have a bug') class TestGeneric1DFilter(FilterTestCaseBase): _func_or_kernels = { 'shift_raw': shift_raw, @@ -688,6 +694,7 @@ def test_filter(self, xp, scp): )) # SciPy behavior fixed in 1.5.0: https://github.com/scipy/scipy/issues/11661 @testing.with_requires('scipy>=1.5.0') +@pytest.mark.skipif(runtime.is_hip, reason='ROCm/HIP may have a bug') class TestMirrorWithDim1(FilterTestCaseBase): @testing.numpy_cupy_allclose(atol=1e-5, rtol=1e-5, scipy_name='scp') def test_filter(self, xp, scp): @@ -714,6 +721,7 @@ def test_filter(self, xp, scp): ) )) @testing.with_requires('scipy') +@pytest.mark.skipif(runtime.is_hip, reason='ROCm/HIP may have a bug') class TestShellSort(FilterTestCaseBase): @testing.numpy_cupy_allclose(atol=1e-5, rtol=1e-5, scipy_name='scp') def test_filter(self, xp, scp): @@ -739,6 +747,7 @@ def test_filter(self, xp, scp): ) )) @testing.with_requires('scipy') +@pytest.mark.skipif(runtime.is_hip, reason='ROCm/HIP may have a bug') class TestFortranOrder(FilterTestCaseBase): @testing.numpy_cupy_allclose(atol=1e-5, rtol=1e-5, scipy_name='scp') def test_filter(self, xp, scp): @@ -765,6 +774,7 @@ def test_filter(self, xp, scp): ) )) @testing.with_requires('scipy') +@pytest.mark.skipif(runtime.is_hip, reason='ROCm/HIP may have a bug') class TestWeightDtype(FilterTestCaseBase): @testing.numpy_cupy_allclose(atol=1e-5, rtol=1e-5, scipy_name='scp') def test_filter(self, xp, scp): @@ -795,6 +805,7 @@ def test_filter(self, xp, scp): ) )) @testing.with_requires('scipy>=1.5.9') +@pytest.mark.skipif(runtime.is_hip, reason='ROCm/HIP may have a bug') class TestWeightComplexDtype(FilterTestCaseBase): def _skip_noncomplex(self): @@ -842,6 +853,7 @@ def test_filter_complex_output_dtype_warns(self): 'dtype': [numpy.uint8, numpy.float64], })) @testing.with_requires('scipy') +@pytest.mark.skipif(runtime.is_hip, reason='ROCm/HIP may have a bug') class TestSpecialWeightCases(FilterTestCaseBase): @testing.numpy_cupy_allclose(atol=1e-5, rtol=1e-5, scipy_name='scp', accept_error=ValueError) @@ -879,6 +891,7 @@ def test_replace_dim_with_0(self, xp, scp): 'dtype': [numpy.uint8, numpy.float64], })) @testing.with_requires('scipy') +@pytest.mark.skipif(runtime.is_hip, reason='ROCm/HIP may have a bug') class TestSpecialCases1D(FilterTestCaseBase): @testing.numpy_cupy_allclose(atol=1e-5, rtol=1e-5, scipy_name='scp', accept_error=RuntimeError) @@ -894,6 +907,7 @@ def test_0_dim(self, xp, scp): 'shape': [(4, 5), (3, 4, 5), (1, 3, 4, 5)], })) @testing.with_requires('scipy') +@pytest.mark.skipif(runtime.is_hip, reason='ROCm/HIP may have a bug') class TestInvalidAxis(FilterTestCaseBase): @testing.numpy_cupy_allclose(atol=1e-5, rtol=1e-5, scipy_name='scp', accept_error=ValueError) @@ -927,6 +941,7 @@ def test_invalid_axis_neg(self, xp, scp): 'shape': [(4, 5)], })) @testing.with_requires('scipy') +@pytest.mark.skipif(runtime.is_hip, reason='ROCm/HIP may have a bug') class TestInvalidMode(FilterTestCaseBase): @testing.numpy_cupy_allclose(atol=1e-5, rtol=1e-5, scipy_name='scp', accept_error=RuntimeError) @@ -945,6 +960,7 @@ def test_invalid_mode(self, xp, scp): })) # SciPy behavior fixed in 1.2.0: https://github.com/scipy/scipy/issues/822 @testing.with_requires('scipy>=1.2.0') +@pytest.mark.skipif(runtime.is_hip, reason='ROCm/HIP may have a bug') class TestInvalidOrigin(FilterTestCaseBase): @testing.numpy_cupy_allclose(atol=1e-5, rtol=1e-5, scipy_name='scp', accept_error=ValueError) diff --git a/tests/cupyx_tests/scipy_tests/ndimage_tests/test_fourier.py b/tests/cupyx_tests/scipy_tests/ndimage_tests/test_fourier.py index 5a897a4b8b7..a54a3e2adf4 100644 --- a/tests/cupyx_tests/scipy_tests/ndimage_tests/test_fourier.py +++ b/tests/cupyx_tests/scipy_tests/ndimage_tests/test_fourier.py @@ -8,6 +8,7 @@ import cupyx.scipy.fft # NOQA import cupyx.scipy.fftpack # NOQA import cupyx.scipy.ndimage # NOQA +from cupy.cuda import runtime try: # scipy.fft only available since SciPy 1.4.0 @@ -54,6 +55,7 @@ ) ) @testing.with_requires("scipy") +@pytest.mark.skipif(runtime.is_hip, reason='ROCm/HIP may have a bug') class TestFourierShift: def _test_real_nd(self, xp, scp, x, real_axis): @@ -141,6 +143,7 @@ def test_complex_fft_with_output(self, xp, scp): ) ) @testing.with_requires("scipy") +@pytest.mark.skipif(runtime.is_hip, reason='ROCm/HIP may have a bug') class TestFourierGaussian: def _test_real_nd(self, xp, scp, x, real_axis): @@ -228,6 +231,7 @@ def test_complex_fft_with_output(self, xp, scp): ) ) @testing.with_requires("scipy") +@pytest.mark.skipif(runtime.is_hip, reason='ROCm/HIP may have a bug') class TestFourierUniform: def _test_real_nd(self, xp, scp, x, real_axis): @@ -309,6 +313,7 @@ def test_complex_fft_with_output(self, xp, scp): ) ) @testing.with_requires('scipy') +@pytest.mark.skipif(runtime.is_hip, reason='ROCm/HIP may have a bug') class TestFourierEllipsoid(): def _test_real_nd(self, xp, scp, x, real_axis): if x.ndim == 1 and scipy_version < '1.5.3': @@ -379,6 +384,7 @@ def test_complex_fft_with_output(self, xp, scp, dtype): @testing.with_requires('scipy') +@pytest.mark.skipif(runtime.is_hip, reason='ROCm/HIP may have a bug') class TestFourierEllipsoidInvalid(): # SciPy < 1.5 raises ValueError instead of AxisError diff --git a/tests/cupyx_tests/scipy_tests/ndimage_tests/test_interpolation.py b/tests/cupyx_tests/scipy_tests/ndimage_tests/test_interpolation.py index 3308b28c427..5ac1e814c2d 100644 --- a/tests/cupyx_tests/scipy_tests/ndimage_tests/test_interpolation.py +++ b/tests/cupyx_tests/scipy_tests/ndimage_tests/test_interpolation.py @@ -41,6 +41,7 @@ def _conditional_scipy_version_skip(mode, order): 'prefilter': [True], })) @testing.with_requires('scipy') +@pytest.mark.skipif(runtime.is_hip, reason='ROCm/HIP may have a bug') class TestMapCoordinates: _multiprocess_can_split = True @@ -122,6 +123,7 @@ def test_map_coordinates_int(self, xp, scp, dtype): 'mode': ['constant', 'nearest', 'mirror'] + scipy16_modes, })) @testing.with_requires('scipy') +@pytest.mark.skipif(runtime.is_hip, reason='ROCm/HIP may have a bug') class TestMapCoordinatesHalfInteger: def _map_coordinates(self, xp, scp, a, coordinates): @@ -150,6 +152,7 @@ def test_map_coordinates_float(self, xp, scp, dtype): 'prefilter': [False, True], })) @testing.with_requires('scipy') +@pytest.mark.skipif(runtime.is_hip, reason='ROCm/HIP may have a bug') class TestAffineTransform: _multiprocess_can_split = True @@ -234,6 +237,7 @@ def test_affine_transform_int(self, xp, scp, dtype): @testing.with_requires('scipy') +@pytest.mark.skipif(runtime.is_hip, reason='ROCm/HIP may have a bug') class TestAffineExceptions: def test_invalid_affine_ndim(self): @@ -321,6 +325,7 @@ def test_invalid_texture_arguments(self): 'theta': [0, 90, 180, 270] })) @testing.with_requires('scipy') +@pytest.mark.skipif(runtime.is_hip, reason='ROCm/HIP may have a bug') class TestAffineTransformTextureMemory: _multiprocess_can_split = True @@ -406,6 +411,7 @@ def test_affine_transform_texture_memory(self, xp, scp): @testing.with_requires('opencv-python') +@pytest.mark.skipif(runtime.is_hip, reason='ROCm/HIP may have a bug') class TestAffineTransformOpenCV: _multiprocess_can_split = True @@ -446,6 +452,7 @@ def test_affine_transform_opencv(self, xp, dtype): }) )) @testing.with_requires('scipy') +@pytest.mark.skipif(runtime.is_hip, reason='ROCm/HIP may have a bug') class TestRotate: _multiprocess_can_split = True @@ -523,6 +530,7 @@ def test_rotate_int(self, xp, scp, dtype): # Scipy older than 1.3.0 raises IndexError instead of ValueError @testing.with_requires('scipy>=1.3.0') +@pytest.mark.skipif(runtime.is_hip, reason='ROCm/HIP may have a bug') class TestRotateExceptions: def test_rotate_invalid_plane(self): @@ -543,6 +551,7 @@ def test_rotate_invalid_plane(self): {'axes': (-2, 2)}, ) @testing.with_requires('scipy') +@pytest.mark.skipif(runtime.is_hip, reason='ROCm/HIP may have a bug') class TestRotateAxes: _multiprocess_can_split = True @@ -556,6 +565,7 @@ def test_rotate_axes(self, xp, scp, dtype): @testing.with_requires('opencv-python') +@pytest.mark.skipif(runtime.is_hip, reason='ROCm/HIP may have a bug') class TestRotateOpenCV: _multiprocess_can_split = True @@ -590,6 +600,7 @@ def test_rotate_opencv(self, xp, dtype): }) )) @testing.with_requires('scipy') +@pytest.mark.skipif(runtime.is_hip, reason='ROCm/HIP may have a bug') class TestShift: _multiprocess_can_split = True @@ -669,6 +680,7 @@ def test_shift_int(self, xp, scp, dtype): 'mode': ['constant', 'nearest'], 'cval': [cupy.nan, cupy.inf, -cupy.inf], })) +@pytest.mark.skipif(runtime.is_hip, reason='ROCm/HIP may have a bug') class TestInterpolationInvalidCval: def _prep_output(self, a): @@ -747,6 +759,7 @@ def test_map_coordinates(self, dtype): @testing.with_requires('opencv-python') +@pytest.mark.skipif(runtime.is_hip, reason='ROCm/HIP may have a bug') class TestShiftOpenCV: _multiprocess_can_split = True @@ -773,6 +786,7 @@ def test_shift_opencv(self, xp, dtype): 'prefilter': [True], })) @testing.with_requires('scipy') +@pytest.mark.skipif(runtime.is_hip, reason='ROCm/HIP may have a bug') class TestZoom: _multiprocess_can_split = True @@ -835,6 +849,7 @@ def test_zoom_int(self, xp, scp, dtype): 'zoom': [(1, 1), (3, 5), (8, 2), (8, 8)], 'mode': ['nearest', 'reflect', 'mirror', 'grid-wrap', 'grid-constant'], })) +@pytest.mark.skipif(runtime.is_hip, reason='ROCm/HIP may have a bug') class TestZoomOrder0IntegerGrid(): def test_zoom_grid_by_int_order0(self): @@ -858,6 +873,7 @@ def test_zoom_grid_by_int_order0(self): 'order': [0, 1, 2, 3, 4, 5], 'grid_mode': [False, True], })) +@pytest.mark.skipif(runtime.is_hip, reason='ROCm/HIP may have a bug') class TestZoomOutputSize1(): @testing.for_float_dtypes(no_float16=True) @@ -875,6 +891,7 @@ def test_zoom_output_size1(self, xp, scp, dtype): {'zoom': 0.3}, ) @testing.with_requires('opencv-python') +@pytest.mark.skipif(runtime.is_hip, reason='ROCm/HIP may have a bug') class TestZoomOpenCV: _multiprocess_can_split = True @@ -912,6 +929,7 @@ def test_zoom_opencv_output_size1(self, xp, dtype): 'axis': [0, 1, 2, -1], })) @testing.with_requires('scipy') +@pytest.mark.skipif(runtime.is_hip, reason='ROCm/HIP may have a bug') class TestSplineFilter1d: @testing.numpy_cupy_allclose(atol=1e-5, rtol=1e-5, scipy_name='scp') def test_spline_filter1d(self, xp, scp): @@ -937,6 +955,7 @@ def test_spline_filter1d_output(self, xp, scp, array_order): # See #5537 @testing.slow @testing.with_requires('scipy') +@pytest.mark.skipif(runtime.is_hip, reason='ROCm/HIP may have a bug') class TestSplineFilter1dLargeArray: @pytest.mark.parametrize('mode', ['mirror', 'grid-wrap', 'reflect']) @@ -962,6 +981,7 @@ def test_spline_filter1d_large_array(self, xp, scp, mode): 'output': [numpy.float64, numpy.float32], })) @testing.with_requires('scipy') +@pytest.mark.skipif(runtime.is_hip, reason='ROCm/HIP may have a bug') class TestSplineFilter: @testing.numpy_cupy_allclose(atol=1e-4, rtol=1e-4, scipy_name='scp') def test_spline_filter(self, xp, scp): @@ -1001,6 +1021,7 @@ def test_spline_filter_with_output(self, xp, scp, array_order): 'output': [numpy.complex64, numpy.complex128], })) @testing.with_requires('scipy') +@pytest.mark.skipif(runtime.is_hip, reason='ROCm/HIP may have a bug') class TestSplineFilterComplex: @testing.with_requires('scipy>=1.6') diff --git a/tests/cupyx_tests/scipy_tests/ndimage_tests/test_measurements.py b/tests/cupyx_tests/scipy_tests/ndimage_tests/test_measurements.py index e1c06db1553..ef76a199a7f 100644 --- a/tests/cupyx_tests/scipy_tests/ndimage_tests/test_measurements.py +++ b/tests/cupyx_tests/scipy_tests/ndimage_tests/test_measurements.py @@ -46,6 +46,7 @@ def _generate_binary_structure(rank, connectivity): 'o_type': [None, 'ndarray'] })) @testing.with_requires('scipy') +@pytest.mark.skipif(runtime.is_hip, reason='ROCm/HIP may have a bug') class TestLabel: @testing.numpy_cupy_array_equal(scipy_name='scp') @@ -70,6 +71,7 @@ def test_label(self, xp, scp): @testing.with_requires('scipy') +@pytest.mark.skipif(runtime.is_hip, reason='ROCm/HIP may have a bug') class TestLabelSpecialCases: @testing.numpy_cupy_array_equal(scipy_name='scp') @@ -114,6 +116,7 @@ def test_label_swirl(self, xp, scp): 'op': stats_ops, })) @testing.with_requires('scipy') +@pytest.mark.skipif(runtime.is_hip, reason='ROCm/HIP may have a bug') class TestStats: def _make_image(self, shape, xp, dtype): @@ -295,6 +298,7 @@ def test_no_values(self, xp, scp, dtype): 'enable_cub': [True, False], })) @testing.with_requires('scipy') +@pytest.mark.skipif(runtime.is_hip, reason='ROCm/HIP may have a bug') class TestMeasurementsSelect: @pytest.fixture(autouse=True) @@ -372,6 +376,7 @@ def test_measurements_select(self, xp, scp, dtype): 'shape': [(200,), (16, 20)], })) @testing.with_requires('scipy') +@pytest.mark.skipif(runtime.is_hip, reason='ROCm/HIP may have a bug') class TestHistogram(): def _make_image(self, shape, xp, dtype, scale): @@ -413,6 +418,7 @@ def test_histogram(self, xp, scp, dtype): 'pass_positions': [True, False], })) @testing.with_requires('scipy') +@pytest.mark.skipif(runtime.is_hip, reason='ROCm/HIP may have a bug') class TestLabeledComprehension(): def _make_image(self, shape, xp, dtype, scale): diff --git a/tests/cupyx_tests/scipy_tests/ndimage_tests/test_morphology.py b/tests/cupyx_tests/scipy_tests/ndimage_tests/test_morphology.py index b339f14a589..8ab8d5dd9fd 100644 --- a/tests/cupyx_tests/scipy_tests/ndimage_tests/test_morphology.py +++ b/tests/cupyx_tests/scipy_tests/ndimage_tests/test_morphology.py @@ -3,6 +3,7 @@ import pytest from cupy import testing +from cupy.cuda import runtime import cupyx.scipy.ndimage # NOQA try: @@ -24,6 +25,7 @@ {'rank': 3, 'connectivity': 0}, {'rank': 3, 'connectivity': 500}) @testing.with_requires('scipy') +@pytest.mark.skipif(runtime.is_hip, reason='ROCm/HIP may have a bug') class TestGenerateBinaryStructure: @testing.numpy_cupy_array_equal(scipy_name='scp') @@ -33,6 +35,7 @@ def test_generate_binary_structure(self, xp, scp): @testing.with_requires('scipy') +@pytest.mark.skipif(runtime.is_hip, reason='ROCm/HIP may have a bug') class TestIterateStructure: @testing.numpy_cupy_array_equal(scipy_name='scp') @@ -88,6 +91,7 @@ def test_iterate_structure3(self, xp, scp): ) )) @testing.with_requires('scipy') +@pytest.mark.skipif(runtime.is_hip, reason='ROCm/HIP may have a bug') class TestBinaryErosionAndDilation1d: def _filter(self, xp, scp, x): filter = getattr(scp.ndimage, self.filter) @@ -157,6 +161,7 @@ def test_binary_erosion_and_dilation_1d(self, xp, scp): ) )) @testing.with_requires('scipy>=1.1.0') +@pytest.mark.skipif(runtime.is_hip, reason='ROCm/HIP may have a bug') class TestBinaryOpeningAndClosing: def _filter(self, xp, scp, x): filter = getattr(scp.ndimage, self.filter) @@ -272,6 +277,7 @@ def test_binary_opening_and_closing(self): )) ) @testing.with_requires('scipy') +@pytest.mark.skipif(runtime.is_hip, reason='ROCm/HIP may have a bug') class TestBinaryFillHoles: def _filter(self, xp, scp, x): filter = scp.ndimage.binary_fill_holes @@ -334,6 +340,7 @@ def test_binary_fill_holes(self, xp, scp): ) )) @testing.with_requires('scipy') +@pytest.mark.skipif(runtime.is_hip, reason='ROCm/HIP may have a bug') class TestBinaryHitOrMiss: def _filter(self, xp, scp, x): filter = scp.ndimage.binary_hit_or_miss @@ -430,6 +437,7 @@ def test_binary_hit_or_miss(self, xp, scp): ) )) @testing.with_requires('scipy') +@pytest.mark.skipif(runtime.is_hip, reason='ROCm/HIP may have a bug') class TestBinaryPropagation: def _filter(self, xp, scp, x): filter = scp.ndimage.binary_propagation @@ -471,6 +479,7 @@ def test_binary_propagation(self, xp, scp): ) )) @testing.with_requires('scipy') +@pytest.mark.skipif(runtime.is_hip, reason='ROCm/HIP may have a bug') class TestBinaryErosionAndDilation: def _filter(self, xp, scp, x): filter = getattr(scp.ndimage, self.filter) @@ -507,6 +516,7 @@ def test_binary_erosion_and_dilation(self, xp, scp): )) ) @testing.with_requires('scipy') +@pytest.mark.skipif(runtime.is_hip, reason='ROCm/HIP may have a bug') class TestBinaryErosionAndDilationContiguity: def _filter(self, xp, scp, x): filter = getattr(scp.ndimage, self.filter) @@ -565,6 +575,7 @@ def test_binary_erosion_and_dilation_input_contiguity(self, xp, scp): }) )) @testing.with_requires('scipy') +@pytest.mark.skipif(runtime.is_hip, reason='ROCm/HIP may have a bug') class TestGreyErosionAndDilation: def _filter(self, xp, scp, x): @@ -610,6 +621,7 @@ def test_grey_erosion_and_dilation(self, xp, scp): 'filter': ['grey_closing', 'grey_opening'] })) @testing.with_requires('scipy') +@pytest.mark.skipif(runtime.is_hip, reason='ROCm/HIP may have a bug') class TestGreyClosingAndOpening: shape = (4, 5) @@ -659,6 +671,7 @@ def test_grey_closing_and_opening(self, xp, scp): )) ) @testing.with_requires('scipy') +@pytest.mark.skipif(runtime.is_hip, reason='ROCm/HIP may have a bug') class TestMorphologicalGradientAndLaplace: def _filter(self, xp, scp, x): @@ -721,6 +734,7 @@ def test_morphological_gradient_and_laplace(self, xp, scp): )) ) @testing.with_requires('scipy') +@pytest.mark.skipif(runtime.is_hip, reason='ROCm/HIP may have a bug') class TestWhiteTophatAndBlackTopHat: def _filter(self, xp, scp, x): diff --git a/tests/cupyx_tests/scipy_tests/signal_tests/test_signaltools.py b/tests/cupyx_tests/scipy_tests/signal_tests/test_signaltools.py index 083f7d94dc4..1411b8ea4a6 100644 --- a/tests/cupyx_tests/scipy_tests/signal_tests/test_signaltools.py +++ b/tests/cupyx_tests/scipy_tests/signal_tests/test_signaltools.py @@ -342,6 +342,7 @@ class TestOrderFilter: @testing.for_all_dtypes(no_float16=True, no_bool=True, no_complex=True) @testing.numpy_cupy_allclose(atol=1e-8, rtol=1e-8, scipy_name='scp', accept_error=ValueError) # for even kernels + @pytest.mark.skipif(runtime.is_hip, reason='ROCm/HIP may have a bug') def test_order_filter(self, xp, scp, dtype): a = testing.shaped_random(self.a, xp, dtype) d = self.domain @@ -374,6 +375,7 @@ def test_medfilt_no_complex(self, xp, scp, dtype): @testing.numpy_cupy_allclose( atol=1e-8, rtol=1e-8, scipy_name='scp', accept_error=ValueError) # for even kernels + @pytest.mark.skipif(runtime.is_hip, reason='ROCm/HIP may have a bug') def test_medfilt(self, xp, scp, dtype): if sys.platform == 'win32': pytest.xfail('medfilt broken for Scipy 1.7.0 in windows') @@ -405,6 +407,7 @@ def test_medfilt2d_no_complex(self, xp, scp, dtype): @testing.numpy_cupy_allclose( atol=1e-8, rtol=1e-8, scipy_name='scp', accept_error=ValueError) # for even kernels + @pytest.mark.skipif(runtime.is_hip, reason='ROCm/HIP may have a bug') def test_medfilt2d(self, xp, scp, dtype): if sys.platform == 'win32': pytest.xfail('medfilt2d broken for Scipy 1.7.0 in windows') diff --git a/tests/cupyx_tests/scipy_tests/sparse_tests/test_linalg.py b/tests/cupyx_tests/scipy_tests/sparse_tests/test_linalg.py index 8097b1c30c3..dacff163f53 100644 --- a/tests/cupyx_tests/scipy_tests/sparse_tests/test_linalg.py +++ b/tests/cupyx_tests/scipy_tests/sparse_tests/test_linalg.py @@ -1202,7 +1202,9 @@ def test_random_initial_float32(self, xp, sp): return eigvals, _eigen_vec_transform(vecs, xp) @pytest.mark.xfail( - runtime.is_hip and driver.get_build_version() >= 5_00_00000, + runtime.is_hip and + (driver.get_build_version() >= 5_00_00000 and + driver.get_build_version() < 50530201), reason='ROCm 5.0+ may have a bug') @pytest.mark.xfail( cupy.cuda.cusolver._getVersion() >= (11, 4, 5), # CUDA 12.1.1+ diff --git a/tests/install_tests/test_build.py b/tests/install_tests/test_build.py index 070aeb3a254..948b17c4df1 100644 --- a/tests/install_tests/test_build.py +++ b/tests/install_tests/test_build.py @@ -20,7 +20,7 @@ def setUp(self): sysconfig.customize_compiler(self.compiler) self.settings = build.get_compiler_setting(ctx, False) - @pytest.mark.skipif(not test_hip, reason='For ROCm/HIP environment') + @pytest.mark.skipif(test_hip, reason='For ROCm/HIP environment') def test_check_hip_version(self): with self.assertRaises(RuntimeError): build.get_hip_version() diff --git a/tests/run_tests_rocm.py b/tests/run_tests_rocm.py new file mode 100644 index 00000000000..cae5cbc1afb --- /dev/null +++ b/tests/run_tests_rocm.py @@ -0,0 +1,142 @@ +import os +import argparse +import re + +TEST_ROOT = os.path.dirname(os.path.abspath(__file__)) +CUPY_TESTS = [ + 'array_api_tests', + 'binary_tests', + 'core_tests', + 'creation_tests', + 'cuda_tests', + 'fft_tests', + 'functional_tests', + 'indexing_tests', + 'io_tests', + 'lib_tests', + 'linalg_tests', + 'logic_tests', + 'manipulation_tests', + 'math_tests', + 'misc_tests', + 'padding_tests', + 'polynomial_tests', + 'prof_tests', + 'random_tests', + 'sorting_tests', + 'statistics_tests', + 'test_cublas.py', + 'testing_tests', + 'test_init.py', + 'test_ndim.py', + 'test_numpy_interop.py', + 'test_type_routines.py', + 'test_typing.py', +] + +CUPYX_TESTS = [ + 'distributed_tests', + 'fallback_mode_tests', + 'jit_tests', + 'linalg_tests', + 'profiler_tests', + 'scipy_tests/fftpack_tests', + 'scipy_tests/fft_tests', + 'scipy_tests/interpolate_tests', + 'scipy_tests/linalg_tests', + 'scipy_tests/ndimage_tests', + 'scipy_tests/signal_tests', + 'scipy_tests/sparse_tests', + 'scipy_tests/spatial_tests', + 'scipy_tests/special_tests', + 'scipy_tests/stats_tests', + 'scipy_tests/test_get_array_module.py', + 'test_cudnn.py', + 'test_cupyx.py', + 'test_cusolver.py', + 'test_cusparse.py', + 'test_cutensor.py', + 'test_lapack.py', + 'test_optimize.py', + 'test_pinned_array.py', + 'test_rsqrt.py', + 'test_runtime.py', + 'test_time.py', + 'tools_tests', +] + +TEST_SUITES = [ + 'cupy_tests', + 'cupyx_tests', + 'example_tests', + 'install_tests', +] + + +def parse_test_log_and_get_summary(test_name): + fs = open("/root/.cache/cupy-tests/cupy_test.log", 'r') + lines = fs.readlines() + fs.close() + + count = "" + summary = "" + pattern = "^=*" + for j in range(len(lines)): + line = lines[j].rstrip() + if ("collecting ..." in line): + count = line.split("collected")[1].split("items")[0].strip() + if ("==" in line): + summary = re.split(pattern, line)[1].split("=")[0] + test_data = test_name + "|" + count + "|" + summary + return test_data + + +def run_all_tests(): + initial_cmd = 'CUPY_TEST_GPU_LIMIT=4 CUPY_INSTALL_USE_HIP=1 ' + \ + 'pytest -vvv -k "not compile_cuda and not fft_allocate" -m "not slow" ' + os.system("mkdir -p ~/.cache/cupy-tests") + test_summary = [] + for test_suite in TEST_SUITES: + if test_suite == "cupy_tests": + for cupy_test in CUPY_TESTS: + cmd = initial_cmd + TEST_ROOT + "/cupy_tests/" + \ + cupy_test + " | tee ~/.cache/cupy-tests/cupy_test.log" + print("Running : " + cmd) + os.system(cmd) + test_name = "tests/cupy_tests/" + cupy_test + test_summary.append(parse_test_log_and_get_summary(test_name)) + elif test_suite == "cupyx_tests": + for cupyx_test in CUPYX_TESTS: + cmd = initial_cmd + TEST_ROOT + "/cupyx_tests/" + \ + cupyx_test + " | tee ~/.cache/cupy-tests/cupy_test.log" + print("Running : " + cmd) + os.system(cmd) + test_name = "tests/cupyx_tests/" + cupyx_test + test_summary.append(parse_test_log_and_get_summary(test_name)) + else: + cmd = initial_cmd + TEST_ROOT + "/" + test_suite + \ + " | tee ~/.cache/cupy-tests/cupy_test.log" + print("Running : " + cmd) + os.system(cmd) + test_name = "tests/" + test_suite + test_summary.append(parse_test_log_and_get_summary(test_name)) + + return test_summary + + +def main(): + all_tests = args.all_tests + if all_tests: + test_summary = run_all_tests() + print("---------------------- TEST SUMMARY ------------------") + for j in range(len(test_summary)): + print(test_summary[j]) + + +if __name__ == '__main__': + parser = argparse.ArgumentParser() + parser.add_argument("--all-tests", action="store_true", + default=True, required=False, help="Run all tests") + args = parser.parse_args() + + main()