master: f72e969 Merge pull request #5474 from kmaehashi/add-ci-generator ============================= test session starts ============================== platform linux -- Python 3.8.7, pytest-6.2.2, py-1.10.0, pluggy-0.13.1 rootdir: /global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy, configfile: setup.cfg plugins: html-3.1.1, metadata-1.11.0 collected 97915 items / 661 deselected / 2 skipped / 97252 selected tests/cupy_tests/test_cublas.py ........................................ [ 0%] ........................................................................ [ 0%] ........................................................................ [ 0%] ........................................................................ [ 0%] ........................................................................ [ 0%] ........................................................................ [ 0%] ........................................................................ [ 0%] ........................................................................ [ 0%] ........................................................................ [ 0%] ........................................................................ [ 0%] ........................................................................ [ 0%] ........................................................................ [ 0%] ........................................................................ [ 0%] ........................................................................ [ 1%] ........................................................................ [ 1%] ........................................................................ [ 1%] ........................................................................ [ 1%] ........................................................................ [ 1%] ........................................................................ [ 1%] ........................................................................ [ 1%] ........................................................................ [ 1%] ........................................................................ [ 1%] ........................................................................ [ 1%] ........................................................................ [ 1%] ........................................................................ [ 1%] ........................................................................ [ 1%] ........................................................................ [ 1%] ........................................................................ [ 2%] ........................................................................ [ 2%] ........................................................................ [ 2%] ........................................................................ [ 2%] ........................................................................ [ 2%] ........................................................................ [ 2%] .................................................x.x....x.x............. [ 2%] .................... [ 2%] tests/cupy_tests/test_cudnn.py sssssssssssssssssssssssssssssssssssssssss [ 2%] ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 2%] ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 2%] ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 2%] ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 2%] ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 2%] ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 2%] ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 3%] ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 3%] ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 3%] ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 3%] ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 3%] ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 3%] ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 3%] ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 3%] ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 3%] ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 3%] ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 3%] ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 3%] ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 3%] ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 4%] ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 4%] ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 4%] ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 4%] ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 4%] ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 4%] tests/cupy_tests/test_cusolver.py ssssssssssssssssssssssssssssssssssssss [ 4%] ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 4%] ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 4%] ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 4%] ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 4%] ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 4%] ssssssssss................................ssssssssssssssssssssssssssssss [ 4%] ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 4%] ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 5%] ssssssssssssssssssssssssssssssssss [ 5%] tests/cupy_tests/test_cusparse.py ..x.x.x.x..x..x..x..x................. [ 5%] .......xxx.xxx....................x.x.x.xssssssssssssssss.....x.xx.x..x. [ 5%] x.x.xx.x..x.x.x.xx.x..x.x.x.xx.x..x.x.x.xx.x..x.x.x.xx.x..x.x.x.xx.x..x. [ 5%] x.x.xx.x..x.x...xxxx....xx..xxxx....xx..xxxx....xx..xxxx....xx..xxxx.... [ 5%] xx..xxxx....xx..xxxx....xx..xxxx....xx..xxxx....xx..xxxx....xx..xxxx.... [ 5%] xx..xxxx....xx..xxxx....xx..xxxx....xx..xxxx....xx..xxxx....xx.....xx... [ 5%] .xx....xx....xx....xx....xx....xx....xx....xx....xx....xx....xx....xx... [ 5%] .xx....xx....xx....xx....xx....xx....xx....xx....xx....xx....xx....xx... [ 5%] .xx....xx....xx....xx....xx....xx....xx....xx....xx....xx....xx....xx... [ 5%] .xx....xx....xx....xx....xx....xx....xx....xx....xx....xx....xx....xx... [ 5%] .xx....xx....xx....xx....xx....xx....xx....xx....xx....xx....xx....xx... [ 5%] .xx....xx....xx....xx....xx....xx....xx....xx....xx....xx....xx....xx... [ 5%] .xx....xx....xx....xx....xx....xx....xx....xx....xx....xx....xx....xx... [ 5%] .xx....xx....xx....xx....xx....xx....xx....xx....xx....xx....xx......... [ 6%] .................xxxxxxxxxxxxxxxxxx.................. [ 6%] tests/cupy_tests/test_cutensor.py ssssssssssssssssssssssssssssssssssssss [ 6%] sssssssssssssssssssssssssssssssssssss [ 6%] tests/cupy_tests/test_init.py ..s........ [ 6%] tests/cupy_tests/test_ndim.py ............ [ 6%] tests/cupy_tests/test_numpy_interop.py ..x...... [ 6%] tests/cupy_tests/test_type_routines.py ................................. [ 6%] . [ 6%] tests/cupy_tests/array_api_tests/test_array_object.py ... [ 6%] tests/cupy_tests/array_api_tests/test_creation_functions.py ............ [ 6%] . [ 6%] tests/cupy_tests/array_api_tests/test_elementwise_functions.py .. [ 6%] tests/cupy_tests/binary_tests/test_elementwise.py ...... [ 6%] tests/cupy_tests/binary_tests/test_packing.py ...... [ 6%] tests/cupy_tests/core_tests/test_array_function.py ..... [ 6%] tests/cupy_tests/core_tests/test_carray.py ..... [ 6%] tests/cupy_tests/core_tests/test_core.py ...... [ 6%] tests/cupy_tests/core_tests/test_cub_reduction.py ...................... [ 6%] ... [ 6%] tests/cupy_tests/core_tests/test_dlpack.py ...... [ 6%] tests/cupy_tests/core_tests/test_elementwise.py ............... [ 6%] tests/cupy_tests/core_tests/test_flags.py ............................. [ 6%] tests/cupy_tests/core_tests/test_function.py ..... [ 6%] tests/cupy_tests/core_tests/test_gufuncs.py ............................ [ 6%] ................................ [ 6%] tests/cupy_tests/core_tests/test_internal.py ........................... [ 6%] ............................................................... [ 6%] tests/cupy_tests/core_tests/test_iter.py ............ [ 6%] tests/cupy_tests/core_tests/test_ndarray.py ..........................ss [ 6%] ssss.................................................................... [ 6%] .......................F....sss............... [ 6%] tests/cupy_tests/core_tests/test_ndarray_adv_indexing.py ............... [ 6%] ........................................................................ [ 6%] ........................................................................ [ 6%] ........................................................................ [ 6%] ............................................. [ 6%] tests/cupy_tests/core_tests/test_ndarray_complex_ops.py ................ [ 6%] ..... [ 7%] tests/cupy_tests/core_tests/test_ndarray_contiguity.py . [ 7%] tests/cupy_tests/core_tests/test_ndarray_conversion.py .......xx.. [ 7%] tests/cupy_tests/core_tests/test_ndarray_copy_and_view.py .............. [ 7%] .......................F..sssssss [ 7%] tests/cupy_tests/core_tests/test_ndarray_cuda_array_interface.py sssssss [ 7%] sssssssssssssssssssssssssssssssssssssssssssssssss [ 7%] tests/cupy_tests/core_tests/test_ndarray_elementwise_op.py .F........... [ 7%] ........................................F............................... [ 7%] ................................................................... [ 7%] tests/cupy_tests/core_tests/test_ndarray_get.py ............ [ 7%] tests/cupy_tests/core_tests/test_ndarray_indexing.py xxxx..........x...x [ 7%] ........x..........x............................. [ 7%] tests/cupy_tests/core_tests/test_ndarray_math.py ......................s [ 7%] sssssssssss [ 7%] tests/cupy_tests/core_tests/test_ndarray_owndata.py ... [ 7%] tests/cupy_tests/core_tests/test_ndarray_reduction.py .................. [ 7%] ........................................................................ [ 7%] ........................................................................ [ 7%] ................................ [ 7%] tests/cupy_tests/core_tests/test_ndarray_scatter.py .................... [ 7%] ........................................................................ [ 7%] ............................................................... [ 7%] tests/cupy_tests/core_tests/test_ndarray_ufunc.py ......... [ 7%] tests/cupy_tests/core_tests/test_ndarray_unary_op.py ..................s [ 7%] ... [ 7%] tests/cupy_tests/core_tests/test_raw.py .ss.........sss..s.s.s........ss [ 7%] .........sss..s.s.s........ss.........sss..s.s.s......ssssssssssssssssss [ 7%] ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss.s [ 7%] s.......sssss....s.s......sssssssssssssssssssssssssssss [ 8%] tests/cupy_tests/core_tests/test_reduction.py ....................... [ 8%] tests/cupy_tests/core_tests/test_scan.py .... [ 8%] tests/cupy_tests/core_tests/test_syncdetect.py .... [ 8%] tests/cupy_tests/core_tests/test_userkernel.py ....sss............. [ 8%] tests/cupy_tests/core_tests/fusion_tests/test_array.py ................. [ 8%] ........................................................................ [ 8%] .................................................... [ 8%] tests/cupy_tests/core_tests/fusion_tests/test_indexing.py .............. [ 8%] .................s [ 8%] tests/cupy_tests/core_tests/fusion_tests/test_kernel_cache.py .... [ 8%] tests/cupy_tests/core_tests/fusion_tests/test_misc.py .......s.......... [ 8%] ........ [ 8%] tests/cupy_tests/core_tests/fusion_tests/test_optimization.py .......... [ 8%] .... [ 8%] tests/cupy_tests/core_tests/fusion_tests/test_reduction.py ............. [ 8%] ........................................................................ [ 8%] .........................................................s.s.s.s.s.s.s.s [ 8%] .s.ssssssssssssssssss.sss [ 8%] tests/cupy_tests/core_tests/fusion_tests/test_routines.py .............. [ 8%] ........................................................................ [ 8%] .............. [ 8%] tests/cupy_tests/core_tests/fusion_tests/test_ufunc.py ................. [ 8%] ...................................................... [ 8%] tests/cupy_tests/creation_tests/test_basic.py .......................... [ 8%] ...........................................................s............ [ 8%] ....s. [ 8%] tests/cupy_tests/creation_tests/test_from_data.py ssssssssssssssssssssss [ 8%] sssssssssssssssssssssssssssssssssssssssssssssssssssssssssss............. [ 8%] ................................................................... [ 8%] tests/cupy_tests/creation_tests/test_matrix.py ......................... [ 8%] ................................. [ 8%] tests/cupy_tests/creation_tests/test_ranges.py ......................... [ 9%] .............................................................. [ 9%] tests/cupy_tests/cuda_tests/test_compiler.py sssssss........ [ 9%] tests/cupy_tests/cuda_tests/test_cublas.py . [ 9%] tests/cupy_tests/cuda_tests/test_cudnn.py s [ 9%] tests/cupy_tests/cuda_tests/test_cufft.py ssssssssssssssss. [ 9%] tests/cupy_tests/cuda_tests/test_curand.py ..... [ 9%] tests/cupy_tests/cuda_tests/test_cusolver.py . [ 9%] tests/cupy_tests/cuda_tests/test_cusparse.py . [ 9%] tests/cupy_tests/cuda_tests/test_cutensor.py s [ 9%] tests/cupy_tests/cuda_tests/test_device.py ................ [ 9%] tests/cupy_tests/cuda_tests/test_driver.py ss. [ 9%] tests/cupy_tests/cuda_tests/test_memory.py ....ssssssss................. [ 9%] .ssssssssss......s.............................................sssssssss [ 9%] ssssssssssss [ 9%] tests/cupy_tests/cuda_tests/test_memory_hook.py . [ 9%] tests/cupy_tests/cuda_tests/test_nccl.py ........ [ 9%] tests/cupy_tests/cuda_tests/test_nvrtc.py . [ 9%] tests/cupy_tests/cuda_tests/test_nvtx.py .... [ 9%] tests/cupy_tests/cuda_tests/test_pinned_memory.py ............ [ 9%] tests/cupy_tests/cuda_tests/test_profile.py .. [ 9%] tests/cupy_tests/cuda_tests/test_runtime.py . [ 9%] tests/cupy_tests/cuda_tests/test_stream.py ..s..s......ssssssssssss.s [ 9%] tests/cupy_tests/cuda_tests/memory_hooks_tests/test_debug_print.py . [ 9%] tests/cupy_tests/cuda_tests/memory_hooks_tests/test_line_profile.py .. [ 9%] tests/cupy_tests/fft_tests/test_cache.py ..sss.....s.. [ 9%] tests/cupy_tests/fft_tests/test_callback.py ssssssssssssssssssssssssssss [ 9%] ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 9%] ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 9%] ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 9%] ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 9%] ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 9%] ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 9%] ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 9%] ssssssssssssssssssssssssssssssssssssssssssssssssssss [ 9%] tests/cupy_tests/fft_tests/test_fft.py ....ssss....ssss........ssss....s [ 9%] sss........ssss....ssss........ssss....ssss........ssss....ssss........s [ 10%] sss....ssss........ssss....ssss........ssss....ssss........ssss....ssss. [ 10%] .......ssss....ssss....................................................s [ 10%] ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 10%] sssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss. [ 10%] .s.s..s.s..s.s..s.s..s.s..s.s..s.s..s.s..s.s..s.s..s.s..s.s..s.s..s.s..s [ 10%] .s..s.s..s.s..s.s..s.s..s.s..s.s..s.s..s.s..s.s..s.s..s.s..s.s..s.s..s.s [ 10%] ..s.s..s.s..s.s..s.s..s.s..s.s..s.s..s.s..s.s..s.s..s.s..s.s..s.s..s.s.. [ 10%] s.s..s.s..s.s..s.s..s.s..s.s..s.s..s.s..s.s..s.s..s.s..s.s..s.s..s.s..s. [ 10%] s..s.s..s.s..s.s..s.s..s.s..s.s..s.s..s.s..s.s..s.s..s.s..s.s..s.s..s.s. [ 10%] .s.s..s.s..s.s..s.s..s.s..s.s..s.s..s.s..s.s..s.s..s.s..s.s..s.s..s.s..s [ 10%] .s..s.s..s.s..s.s..s.s..s.s..s.s..s.s..s.s.s.s.s.s.s.s.s.s.s.s.s.s.s.s.s [ 10%] .s.s.sssss.s.s.s.s.s.s.s.s.s.s.s.s.s.s.s.s.s.s.s.s.s.s.s.s.s.s.s.s.s.s.s [ 10%] .sssss.s.s.s.s.s.s.s.s.s.s.s.s.s.s.s.s.s.s.s.s.s.s.s.s.s.s.s.s.s.s.s.sss [ 11%] ss.s.s.s.s.s.s.s.s.s.s.s.s.s.s.s.s.s.s.s.s.s.s.s.s.s.s.s.s.s.s.s.s.s.s.s [ 11%] .sssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss.s.s..s.s. [ 11%] .s.s..s.s..s.s..s.s..s.s..s.s..s.s..s.s..s.s..s.s..s.s..s.s..s.s..s.s..s [ 11%] .s..s.s..s.s..s.s..s.s..s.s..s.s..s.s...ss..ss....ss..ss....ss..ss....ss [ 11%] ..ss....ss..ss....ss..ss....ss..ss....ss..ss...s.s.s.s.s.s.s.s.s.s.s.s.s [ 11%] .s.s.s.s.s.s.s.s.s.s.s.s.s..s.s..s.s..s.s..s.s..s.s..s.s..s.s..s.s..s.s. [ 11%] .s.s..s.s..s.s..s.s..s.s..s.s..s.s..s.s..s.s..s.s..s.s..s.s..s.s..s.s..s [ 11%] .s..s.s..s.s..s.s..s.s..s.s..s.s..s.s......s.s..s.s..s.s..s.s..s.s..s.s. [ 11%] .s.s..s.s..s.s..s.s..s.s..s.s..s.s..s.s..s.s..s.s..s.s..s.s..s.s..s.s..s [ 11%] .s..s.s..s.s..s.s..s.s..s.s..s.s..s.s..s.s..s.s..s.s..s.s..s.s..s.s..s.s [ 11%] ..s.s..s.s..s.s..s.s..s.s..s.s..s.s.sssss.s.s..s.s.sssssssssssssssssssss [ 11%] ssssssssssssssssssssssssssssssssssssssssssss.s.s..s.s..s.s..s.s..s.s..s. [ 11%] s..s.s..s.s..s.s..s.s..s.s..s.s..s.s..s.s..s.s..s.s..s.s..s.s..s.s..s.s. [ 11%] .s.s..s.s..s.s..s.s..s.s..s.s.......ss..ss....ss..ss....ss..ss....ss..ss [ 12%] ....ss..ss....ss..ss....ss..ss....ss..ss........................ [ 12%] tests/cupy_tests/functional_tests/test_piecewise.py ............. [ 12%] tests/cupy_tests/functional_tests/test_vectorize.py .................... [ 12%] .................................. [ 12%] tests/cupy_tests/indexing_tests/test_generate.py ....................... [ 12%] ............... [ 12%] tests/cupy_tests/indexing_tests/test_indexing.py ....................... [ 12%] ......................... [ 12%] tests/cupy_tests/indexing_tests/test_insert.py ......................... [ 12%] ........................................................................ [ 12%] ........................................................................ [ 12%] ......................................................... [ 12%] tests/cupy_tests/indexing_tests/test_iterate.py ........................ [ 12%] ............................................... [ 12%] tests/cupy_tests/io_tests/test_base_n.py ............................ [ 12%] tests/cupy_tests/io_tests/test_formatting.py .. [ 12%] tests/cupy_tests/io_tests/test_npz.py ....... [ 12%] tests/cupy_tests/lib_tests/test_polynomial.py ..............ssss........ [ 12%] ..........sssssssssssssss........ssssssssssssssssssssssssssssssssssssxxx [ 12%] xxxxxxxxxxxxxxxxx...ssssssssssssssssssssssssssssssssssssssssssssssss.... [ 12%] ssssssssssssssssssssssssssssssssssssssssssssssss......xxxxxxxxxxxx...... [ 12%] ..........................................xxxxxxxxxxxx.................. [ 12%] ..............................xxxxxxxx........xxxxxxxx.................. [ 13%] ..................................................x...x................. [ 13%] .. [ 13%] tests/cupy_tests/lib_tests/test_shape_base.py ...................... [ 13%] tests/cupy_tests/lib_tests/test_strided_tricks.py .. [ 13%] tests/cupy_tests/linalg_tests/test_decomposition.py .s....s....s....s... [ 13%] ................................... [ 13%] tests/cupy_tests/linalg_tests/test_eigenvalue.py .................. [ 13%] tests/cupy_tests/linalg_tests/test_einsum.py ........................... [ 13%] ........................................................................ [ 13%] ........................................................................ [ 13%] ........................................................................ [ 13%] ........................................................................ [ 13%] ........................................................................ [ 13%] ........................................................................ [ 13%] ........................................................................ [ 13%] ........................................................................ [ 13%] ........................................................................ [ 13%] ........................................................................ [ 13%] .......... [ 13%] tests/cupy_tests/linalg_tests/test_norms.py ............................ [ 13%] ........................................................................ [ 14%] ........................................................................ [ 14%] ................................................................. [ 14%] tests/cupy_tests/linalg_tests/test_product.py F.F.F.F.........F.F.F.F... [ 14%] ......F.F.F.F.F.F.F.F.F.F.F.F.F.F.F.F.F.F.F.F........................... [ 14%] ..............................F.F.F.F.F.F.F.F........................... [ 14%] .....FF.FFFF.F.FF.FFF.FF.F.FFF....FF....FF...... [ 14%] tests/cupy_tests/linalg_tests/test_solve.py ......................F..... [ 14%] .................................... [ 14%] tests/cupy_tests/logic_tests/test_comparison.py ........................ [ 14%] ..... [ 14%] tests/cupy_tests/logic_tests/test_content.py ... [ 14%] tests/cupy_tests/logic_tests/test_ops.py .... [ 14%] tests/cupy_tests/logic_tests/test_truth.py ............................. [ 14%] ........................................................................ [ 14%] ........................................................................ [ 14%] ........................................................................ [ 14%] ........................................................................ [ 14%] ........................................................................ [ 14%] ........................................................................ [ 15%] ........................................................................ [ 15%] ........................................................................ [ 15%] ........................................................................ [ 15%] ........................................... [ 15%] tests/cupy_tests/logic_tests/test_type_test.py ......................... [ 15%] ....... [ 15%] tests/cupy_tests/manipulation_tests/test_add_remove.py ................. [ 15%] ............................ [ 15%] tests/cupy_tests/manipulation_tests/test_basic.py ...................... [ 15%] ....................................................................... [ 15%] tests/cupy_tests/manipulation_tests/test_dims.py .....x.x............... [ 15%] ................................................ [ 15%] tests/cupy_tests/manipulation_tests/test_join.py ....................... [ 15%] ........................... [ 15%] tests/cupy_tests/manipulation_tests/test_kind.py ........... [ 15%] tests/cupy_tests/manipulation_tests/test_rearrange.py .................. [ 15%] ....................................................... [ 15%] tests/cupy_tests/manipulation_tests/test_shape.py ...................... [ 15%] ........................................ [ 15%] tests/cupy_tests/manipulation_tests/test_split.py .................. [ 15%] tests/cupy_tests/manipulation_tests/test_tiling.py ..................... [ 15%] ........... [ 15%] tests/cupy_tests/manipulation_tests/test_transpose.py .................. [ 15%] .......... [ 15%] tests/cupy_tests/math_tests/test_arithmetic.py ......................... [ 15%] ........................................................................ [ 15%] ........................................................................ [ 15%] ........................................................................ [ 16%] ........................................................................ [ 16%] ........................................................................ [ 16%] ........................................................................ [ 16%] ........................................................................ [ 16%] ........................................................................ [ 16%] ........................................................................ [ 16%] ........................................................................ [ 16%] ........................................................................ [ 16%] ........................................................................ [ 16%] ........................................................................ [ 16%] ........................................................................ [ 16%] ........................................................................ [ 16%] ........................................................................ [ 17%] ........................................................................ [ 17%] ........................................................................ [ 17%] ........................................................................ [ 17%] ........................................................................ [ 17%] ........................................................................ [ 17%] ........................................................................ [ 17%] ........................................................................ [ 17%] ........................................................................ [ 17%] ........................................................................ [ 17%] ........................................................................ [ 17%] ........................................................................ [ 17%] ........................................................................ [ 17%] ........................................................................ [ 17%] ........................................................................ [ 18%] ........................................................................ [ 18%] ........................................................................ [ 18%] ........................................................................ [ 18%] ........................................................................ [ 18%] ........................................................................ [ 18%] ........................................................................ [ 18%] ........................................................................ [ 18%] ........................................................................ [ 18%] ........................................................................ [ 18%] ........................................................................ [ 18%] ........................................................................ [ 18%] ........................................................................ [ 18%] ........................................................................ [ 19%] ........................................................................ [ 19%] ........................................................................ [ 19%] ........................................................................ [ 19%] ........................................................................ [ 19%] ........................................................................ [ 19%] ........................................................................ [ 19%] ........................................................................ [ 19%] ........................................................................ [ 19%] ........................................................................ [ 19%] ........................................................................ [ 19%] ........................................................................ [ 19%] ........................................................................ [ 19%] ........................................................................ [ 19%] ........................................................................ [ 20%] ........................................................................ [ 20%] ........................................................................ [ 20%] ........................................................................ [ 20%] ........................................................................ [ 20%] ........................................................................ [ 20%] ........................................................................ [ 20%] ........................................................................ [ 20%] ........................................................................ [ 20%] ........................................................................ [ 20%] ........................................................................ [ 20%] ........................................................................ [ 20%] ........................................................................ [ 20%] ........................................................................ [ 21%] ........................................................................ [ 21%] ........................................................................ [ 21%] ........................................................................ [ 21%] ........................................................................ [ 21%] ........................................................................ [ 21%] ........................................................................ [ 21%] ........................................................................ [ 21%] ........................................................................ [ 21%] ........................................................................ [ 21%] ........................................................................ [ 21%] ........................................................................ [ 21%] ........................................................................ [ 21%] ........................................................................ [ 21%] ........................................................................ [ 22%] ........................................................................ [ 22%] ........................................................................ [ 22%] ........................................................................ [ 22%] ........................................................................ [ 22%] ........................................................................ [ 22%] ........................................................................ [ 22%] ........................................................................ [ 22%] ........................................................................ [ 22%] ........................................................................ [ 22%] ........................................................................ [ 22%] ........................................................................ [ 22%] ........................................................................ [ 22%] ........................................................................ [ 23%] ........................................................................ [ 23%] ........................................................................ [ 23%] ........................................................................ [ 23%] ........................................................................ [ 23%] ........................................................................ [ 23%] ........................................................................ [ 23%] ........................................................................ [ 23%] ........................................................................ [ 23%] ........................................................................ [ 23%] ........................................................................ [ 23%] ........................................................................ [ 23%] ........................................................................ [ 23%] ........................................................................ [ 23%] ........................................................................ [ 24%] ........................................................................ [ 24%] ....................ss...xxsssssxxxxxxxxxxx........ [ 24%] tests/cupy_tests/math_tests/test_explog.py ......... [ 24%] tests/cupy_tests/math_tests/test_floating.py ....... [ 24%] tests/cupy_tests/math_tests/test_hyperbolic.py ...... [ 24%] tests/cupy_tests/math_tests/test_matmul.py FF......FF..FF............... [ 24%] ........................................................................ [ 24%] ........................................................................ [ 24%] ........................................................................ [ 24%] ................................................ [ 24%] tests/cupy_tests/math_tests/test_misc.py ............................... [ 24%] ............................................F..F..F..................... [ 24%] .............................. [ 24%] tests/cupy_tests/math_tests/test_rational.py .... [ 24%] tests/cupy_tests/math_tests/test_rounding.py ........................... [ 24%] .......................................... [ 24%] tests/cupy_tests/math_tests/test_special.py ... [ 24%] tests/cupy_tests/math_tests/test_sumprod.py ............................ [ 24%] ........................................................................ [ 24%] .......ssssssssssssssss................................................. [ 24%] ........................................................................ [ 24%] ........................................................................ [ 25%] ........................................................................ [ 25%] ........................................................................ [ 25%] ........................................................................ [ 25%] ........................................................................ [ 25%] ........................................................................ [ 25%] .............................................................. [ 25%] tests/cupy_tests/math_tests/test_trigonometric.py ..........sssssssss [ 25%] tests/cupy_tests/math_tests/test_window.py ............................. [ 25%] ..... [ 25%] tests/cupy_tests/misc_tests/test_memory_ranges.py ................ [ 25%] tests/cupy_tests/misc_tests/test_who.py ...... [ 25%] tests/cupy_tests/padding_tests/test_pad.py ............................. [ 25%] .............................................. [ 25%] tests/cupy_tests/polynomial_tests/test_polynomial.py ................... [ 25%] [ 25%] tests/cupy_tests/polynomial_tests/test_polyutils.py .................... [ 25%] ................ [ 25%] tests/cupy_tests/prof_tests/test_range.py ........ [ 25%] tests/cupy_tests/random_tests/test_bit_generator.py ssssssssssss [ 25%] tests/cupy_tests/random_tests/test_distributions.py .................... [ 25%] ........................................................................ [ 25%] ........................................................................ [ 25%] .............xxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxx....................... [ 25%] ........................................................................ [ 26%] ........................................................................ [ 26%] ........................................................................ [ 26%] ........................................................................ [ 26%] ........................................................................ [ 26%] ........................................................................ [ 26%] ........................................................................ [ 26%] ........................................................................ [ 26%] ........................................................................ [ 26%] ........................................................................ [ 26%] ........................................................................ [ 26%] .....xxxxxxx..................................xxxxxxx................... [ 26%] ........................................................................ [ 26%] ............................................. [ 26%] tests/cupy_tests/random_tests/test_generator.py .......s.s.s..........s. [ 26%] ..................s..................................s.......s.s........ [ 27%] ..........................s............................................. [ 27%] ..................................s.s.s......................s.s........ [ 27%] ...................................................................x.... [ 27%] ........................... [ 27%] tests/cupy_tests/random_tests/test_generator_api.py ssssssssssssssssssss [ 27%] ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 27%] ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 27%] ssssssssss [ 27%] tests/cupy_tests/random_tests/test_init.py . [ 27%] tests/cupy_tests/random_tests/test_permutations.py ...x...x............. [ 27%] ........................................................................ [ 27%] ........................................................................ [ 27%] ............. [ 27%] tests/cupy_tests/random_tests/test_random.py . [ 27%] tests/cupy_tests/random_tests/test_sample.py ...............x........x.. [ 27%] ............ [ 27%] tests/cupy_tests/sorting_tests/test_count.py .... [ 27%] tests/cupy_tests/sorting_tests/test_search.py ........................ss [ 27%] ss....ssss....................ssss....ssss.............................. [ 27%] ........................................................................ [ 27%] .................................................................... [ 27%] tests/cupy_tests/sorting_tests/test_sort.py ............................ [ 27%] ........................................................................ [ 28%] ..........................................................x........... [ 28%] tests/cupy_tests/statistics_tests/test_correlation.py .................. [ 28%] ..............................F..F.xF................... [ 28%] tests/cupy_tests/statistics_tests/test_histogram.py .................... [ 28%] ........................................................................ [ 28%] ........................................................................ [ 28%] ........................................................................ [ 28%] ........................................................................ [ 28%] ........................................................ [ 28%] tests/cupy_tests/statistics_tests/test_meanvar.py ...................... [ 28%] ........................................................................ [ 28%] ........................................................................ [ 28%] ...... [ 28%] tests/cupy_tests/statistics_tests/test_order.py sssss................... [ 28%] ......................... [ 28%] tests/cupy_tests/testing_tests/test_array.py ........................... [ 28%] ............................................... [ 28%] tests/cupy_tests/testing_tests/test_condition.py ................ [ 28%] tests/cupy_tests/testing_tests/test_helper.py .......................... [ 28%] ........................................................................ [ 29%] .....sssss..sssssss..ssxxxxxxxx................x..x................. [ 29%] tests/cupy_tests/testing_tests/test_parameterized.py ............s.s.s.s [ 29%] ....ssss...... [ 29%] tests/cupyx_tests/test_cupyx.py .. [ 29%] tests/cupyx_tests/test_lapack.py .s.s.s.s.....s.s.s.s.....s.s.s.s.....s. [ 29%] s.s.s.............sssssssssssss [ 29%] tests/cupyx_tests/test_optimize.py sssssssss [ 29%] tests/cupyx_tests/test_pinned_array.py ................................. [ 29%] ....................................... [ 29%] tests/cupyx_tests/test_rsqrt.py . [ 29%] tests/cupyx_tests/test_runtime.py s. [ 29%] tests/cupyx_tests/test_scatter.py ... [ 29%] tests/cupyx_tests/test_time.py ......... [ 29%] tests/cupyx_tests/fallback_mode_tests/test_fallback.py ................. [ 29%] ........................................................................ [ 29%] ......... [ 29%] tests/cupyx_tests/fallback_mode_tests/test_notifications.py ......... [ 29%] tests/cupyx_tests/jit_tests/test_raw.py ................................ [ 29%] ........... [ 29%] tests/cupyx_tests/linalg_tests/test_solve.py xxxxxxxxxxxxxxxx....ssss [ 29%] tests/cupyx_tests/scipy_tests/test_get_array_module.py .. [ 29%] tests/cupyx_tests/scipy_tests/fft_tests/test_fft.py .................... [ 29%] ........................................................................ [ 29%] ........................................................................ [ 29%] ........................................................................ [ 29%] ........................................................................ [ 29%] ........................................................................ [ 29%] ........................................................................ [ 29%] ........................................................................ [ 29%] ........................................................................ [ 30%] ........................................................................ [ 30%] ........................................................................ [ 30%] ........................................................................ [ 30%] ........................................................................ [ 30%] ........................................................................ [ 30%] ........................................................................ [ 30%] ........................................................................ [ 30%] ........................................................................ [ 30%] ........................................................................ [ 30%] ........................................................................ [ 30%] ........................................................................ [ 30%] ........................................................................ [ 30%] ........................................................................ [ 31%] ........................................................................ [ 31%] ........................................................................ [ 31%] ........................................................................ [ 31%] ........................................................................ [ 31%] ........................................................................ [ 31%] ........................................................................ [ 31%] ........................................................................ [ 31%] ........................................................................ [ 31%] ........................................................................ [ 31%] ........................................................................ [ 31%] ........................................................................ [ 31%] ........................................................................ [ 31%] ........................................................................ [ 31%] ........................................................................ [ 32%] ........................................................................ [ 32%] ........................................................................ [ 32%] ........................................................................ [ 32%] ........................................................................ [ 32%] ........................................................................ [ 32%] ........................................................................ [ 32%] ........................................................................ [ 32%] ........................................................................ [ 32%] ........................................................................ [ 32%] ........................................................................ [ 32%] ........................................................................ [ 32%] ........................................................................ [ 32%] ........................................................................ [ 33%] ........................................................................ [ 33%] ........................................................................ [ 33%] ........................................................................ [ 33%] ........................................................................ [ 33%] ........................................................................ [ 33%] ........................................................................ [ 33%] ........................................................................ [ 33%] ........................................................................ [ 33%] ........................................................................ [ 33%] ........................................................................ [ 33%] ........................................................................ [ 33%] ........................................................................ [ 33%] ........................................................................ [ 33%] ........................................................................ [ 34%] ........................................................................ [ 34%] ........................................................................ [ 34%] ........................................................................ [ 34%] ........................................................................ [ 34%] ........................................................................ [ 34%] ........................................................................ [ 34%] ........................................................................ [ 34%] ........................................................................ [ 34%] ........................................................................ [ 34%] ........................................................................ [ 34%] ........................................................................ [ 34%] ........................................................................ [ 34%] ........................................................................ [ 35%] ........................................................................ [ 35%] ........................................................................ [ 35%] ........................................................................ [ 35%] ........................................................................ [ 35%] ........................................................................ [ 35%] ........................................................................ [ 35%] ........................................................................ [ 35%] ........................................................................ [ 35%] ........................................................................ [ 35%] ........................................................................ [ 35%] ........................................................................ [ 35%] ........................................................................ [ 35%] ........................................................................ [ 35%] ........................................................................ [ 36%] ........................................................................ [ 36%] ........................................................................ [ 36%] ........................................................................ [ 36%] ........................................................................ [ 36%] ........................................................................ [ 36%] ........................................................................ [ 36%] ........................................................................ [ 36%] ........................................................................ [ 36%] ........................................................................ [ 36%] ........................................................................ [ 36%] ........................................................................ [ 36%] ........................................................................ [ 36%] ........................................................................ [ 37%] ........................................................................ [ 37%] ........................................................................ [ 37%] ........................................................................ [ 37%] ........................................................................ [ 37%] ........................................................................ [ 37%] ........................................................................ [ 37%] ........................................................................ [ 37%] ........................................................................ [ 37%] ......................................................................ss [ 37%] ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 37%] ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 37%] ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss........ [ 37%] ........................................................................ [ 37%] ........................................................................ [ 38%] ........................................................................ [ 38%] ........................................................................ [ 38%] ........................................................................ [ 38%] ........................................................................ [ 38%] ........................................................................ [ 38%] ........................................................................ [ 38%] ..............................................ssssssssssssssssssssssssss [ 38%] ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 38%] ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 38%] ssssssssssssssssssssssssssssssssssssssss................................ [ 38%] ........................................................................ [ 38%] ........................................................................ [ 38%] ........................................................................ [ 39%] ........................................................................ [ 39%] ........................................................................ [ 39%] ........................................................................ [ 39%] ........................................................................ [ 39%] ........................................................................ [ 39%] ........................................................................ [ 39%] ........................................................................ [ 39%] ........................................................................ [ 39%] ........................................................................ [ 39%] ........................................................................ [ 39%] ...................................................................... [ 39%] tests/cupyx_tests/scipy_tests/fft_tests/test_helper.py . [ 39%] tests/cupyx_tests/scipy_tests/fftpack_tests/test_fftpack.py ............ [ 39%] ........................................................................ [ 39%] ........................................................................ [ 40%] ........................................................................ [ 40%] ........................................................................ [ 40%] ........................................................................ [ 40%] ........................................................................ [ 40%] ........................................................................ [ 40%] ........................................................................ [ 40%] ........................................................................ [ 40%] ........................................................................ [ 40%] ........................................................................ [ 40%] ........................................................................ [ 40%] ......................... [ 40%] tests/cupyx_tests/scipy_tests/linalg_tests/test_decomp_lu.py ........... [ 40%] ................................................... [ 40%] tests/cupyx_tests/scipy_tests/linalg_tests/test_solve_triangular.py .... [ 40%] ........................................................................ [ 40%] ........................................................................ [ 40%] ........................................................................ [ 41%] .................................................................... [ 41%] tests/cupyx_tests/scipy_tests/linalg_tests/test_special_matrices.py .... [ 41%] ........................................................................ [ 41%] ................................................. [ 41%] tests/cupyx_tests/scipy_tests/ndimage_tests/test_filters.py ...sss...... [ 41%] sss......sss......sss......sss......sss.........sss......sss......sss... [ 41%] ...sss......sss......sss................................................ [ 41%] ...........................................................sss......sss. [ 41%] .....sss......sss......sss......sss.........sss......sss......sss......s [ 41%] ss......sss......sss.................................................... [ 41%] .......................................................sss......sss..... [ 41%] .sss......sss......sss......sss.........sss......sss......sss......sss.. [ 41%] ....sss......sss........................................................ [ 41%] ...................................................sss......sss......sss [ 41%] ......sss......sss......sss.........sss......sss......sss......sss...... [ 42%] sss......sss............................................................ [ 42%] ...............................................sss......sss......sss.... [ 42%] ..sss......sss......sss.........sss......sss......sss......sss......sss. [ 42%] .....sss................................................................ [ 42%] ...........................................sss......sss......sss......ss [ 42%] s......sss......sss.........sss......sss......sss......sss......sss..... [ 42%] .sss.................................................................... [ 42%] .......................................sss......sss......sss......sss... [ 42%] ...sss......sss.........sss......sss......sss......sss......sss......sss [ 42%] ........................................................................ [ 42%] ...................................sss......sss......sss......sss......s [ 42%] ss......sss.........sss......sss......sss......sss......sss......sss.... [ 42%] ........................................................................ [ 42%] ...............................sss......sss......sss......sss......sss.. [ 43%] ....sss.........sss......sss......sss......sss......sss......sss........ [ 43%] ........................................................................ [ 43%] ...........................sss......sss......sss......sss......sss...... [ 43%] sss.........sss......sss......sss......sss......sss......sss............ [ 43%] ........................................................................ [ 43%] .......................sss......sss......sss......sss......sss......sss. [ 43%] ........sss......sss......sss......sss......sss......sss................ [ 43%] ........................................................................ [ 43%] ...................sss......sss......sss......sss......sss......sss..... [ 43%] ....sss......sss......sss......sss......sss......sss.................... [ 43%] ........................................................................ [ 43%] ...............sss......sss......sss......sss......sss......sss......... [ 43%] sss......sss......sss......sss......sss......sss........................ [ 44%] ........................................................................ [ 44%] ...........sss......sss......sss......sss......sss......sss.........sss. [ 44%] .....sss......sss......sss......sss......sss............................ [ 44%] ........................................................................ [ 44%] xxx....sss......sss......sss......sss......sss......sss....x..x.sss.x..x [ 44%] .sss.x..x.sss......sss......sss......sss................................ [ 44%] .....................................................x..x.............x. [ 44%] ......................................x.....x....x.....x......x.....x... [ 44%] ........................................................................ [ 44%] ........................................................................ [ 44%] ........................................................................ [ 44%] ........................................................................ [ 44%] ........................................................................ [ 44%] ........................................................................ [ 45%] ........................................................................ [ 45%] ........................................................................ [ 45%] ........................................................................ [ 45%] ........................................................................ [ 45%] ........................................................................ [ 45%] ........................................................................ [ 45%] ........................................................................ [ 45%] ........................................................................ [ 45%] ........................................................................ [ 45%] ........................................................................ [ 45%] ........................................................................ [ 45%] ........................................................................ [ 45%] ........................................................................ [ 46%] ........................................................................ [ 46%] ........................................................................ [ 46%] ........................................................................ [ 46%] ........................................................................ [ 46%] ........................................................................ [ 46%] ........................................................................ [ 46%] ........................................................................ [ 46%] ........................s.s.s.s.s.s.s.s.s.s.s.s.s.s.s.s.s.s............. [ 46%] ........................................................................ [ 46%] .........................s.s.s.s.s.s.s.s.s.s.s.s.s.s.s.s.s.ss.s.s.s.s.s. [ 46%] s.s.s.s.s.s.s.s.s.s.s.s................................................. [ 46%] .............................................................s.s.s.s.s.s [ 46%] .s.s.s.s.s.s.s.s.s.s.s.ss.s.s.s.s.s.s.s.s.s.s.s.s.s.s.s.s.s............. [ 46%] ........................................................................ [ 47%] .........................s.s.s.s.s.s.s.s.s.s.s.s.s.s.s.s.s.ss.s.s.s.s.s. [ 47%] s.s.s.s.s.s.s.s.s.s.s.s................................................. [ 47%] .............................................................s.s.s.s.s.s [ 47%] .s.s.s.s.s.s.s.s.s.s.s.ss.s.s.s.s.s.s.s.s.s.s.s.s.s.s.s.s.s............. [ 47%] ........................................................................ [ 47%] .........................s.s.s.s.s.s.s.s.s.s.s.s.s.s.s.s.s.ss.s.s.s.s.s. [ 47%] s.s.s.s.s.s.s.s.s.s.s.s................................................. [ 47%] .............................................................s.s.s.s.s.s [ 47%] .s.s.s.s.s.s.s.s.s.s.s.ss.s.s.s.s.s.s.s.s.s.s.s.s.s.s.s.s.s............. [ 47%] ........................................................................ [ 47%] .........................s.s.s.s.s.s.s.s.s.s.s.s.s.s.s.s.s.ss.s.s.s.s.s. [ 47%] s.s.s.s.s.s.s.s.s.s.s.s................................................. [ 47%] .............................................................s.s.s.s.s.s [ 48%] .s.s.s.s.s.s.s.s.s.s.s.ssss..sss..sss............sss..sss..sss.......... [ 48%] ..sss..sss..sss............sss..sss..sss............sss..sss..sss....... [ 48%] .....sss..sss..sss............sss..sss..sss............sss..sss..sss.... [ 48%] ........sss..sss..sss............sss..sss..sss............sss..sss..sss. [ 48%] ...........sss..sss..sss................................................ [ 48%] ........................................................................ [ 48%] .................................... [ 48%] tests/cupyx_tests/scipy_tests/ndimage_tests/test_fourier.py ............ [ 48%] ........................................................................ [ 48%] ........................................................................ [ 48%] ........................................................................ [ 48%] ........................................................................ [ 48%] ........................................................................ [ 48%] ........................................................................ [ 48%] ........................................................................ [ 49%] ........................................ [ 49%] tests/cupyx_tests/scipy_tests/ndimage_tests/test_interpolation.py ...... [ 49%] ........................................................................ [ 49%] ........................................................................ [ 49%] ........................................................................ [ 49%] ........................................................................ [ 49%] ........................................................................ [ 49%] ........................................................................ [ 49%] ........................................................................ [ 49%] ........................................................................ [ 49%] ........................................................................ [ 49%] ........................................................................ [ 49%] ........................................................................ [ 49%] ........................................................................ [ 49%] ........................................................................ [ 50%] ........................................................................ [ 50%] ........................................................................ [ 50%] ........................................................................ [ 50%] ........................................................................ [ 50%] ................FF...................................................... [ 50%] ........................................................................ [ 50%] ........................................................................ [ 50%] ........................................................................ [ 50%] ........................................................................ [ 50%] ........................................................................ [ 50%] ........................................................................ [ 50%] ........................................................................ [ 50%] ........................................................................ [ 50%] ........................................................................ [ 51%] ........................................................................ [ 51%] ........................................................................ [ 51%] ........................................................................ [ 51%] ........................................................................ [ 51%] ........................................................................ [ 51%] ........................................................................ [ 51%] ........................................................................ [ 51%] ........................................................................ [ 51%] ........................................................................ [ 51%] ........................................................................ [ 51%] ........................................................................ [ 51%] ........................................................................ [ 51%] ........................................................................ [ 52%] ........................................................................ [ 52%] ........................................................................ [ 52%] ........................................................................ [ 52%] ........................................................................ [ 52%] ........................................................................ [ 52%] ........................................................................ [ 52%] ........................................................................ [ 52%] ........................................................................ [ 52%] ........................................................................ [ 52%] ........................................................................ [ 52%] ........................................................................ [ 52%] ........................................................................ [ 52%] ........................................................................ [ 52%] ........................................................................ [ 53%] ........................................................................ [ 53%] ........................................................................ [ 53%] ........................................................................ [ 53%] ........................................................................ [ 53%] ........................................................................ [ 53%] ........................................................................ [ 53%] ........................................................................ [ 53%] ........................................................................ [ 53%] ........................................................................ [ 53%] ........................................................................ [ 53%] ........................................................................ [ 53%] ........................................................................ [ 53%] ........................................................................ [ 54%] ........................................................................ [ 54%] ........................................................................ [ 54%] ........................................................................ [ 54%] ........................................................................ [ 54%] ........................................................................ [ 54%] ........................................................................ [ 54%] ........................................................................ [ 54%] ........................................................................ [ 54%] ........................................................................ [ 54%] ........................................................................ [ 54%] ........................................................................ [ 54%] ........................................................................ [ 54%] ........................................................................ [ 54%] ........................................................................ [ 55%] ........................................................................ [ 55%] ........................................................................ [ 55%] ........................................................................ [ 55%] ........................................................................ [ 55%] ........................................................................ [ 55%] ........................................................................ [ 55%] ........................................................................ [ 55%] ........................................................................ [ 55%] ........................................................................ [ 55%] ........................................................................ [ 55%] ........................................................................ [ 55%] ........................................................................ [ 55%] ........................................................................ [ 56%] ........................................................................ [ 56%] ........................................................................ [ 56%] ........................................................................ [ 56%] ........................................................................ [ 56%] ........................................................................ [ 56%] ........................................................................ [ 56%] ........................................................................ [ 56%] ........................................................................ [ 56%] ........................................................................ [ 56%] ........................................................................ [ 56%] ........................................................................ [ 56%] ........................................................................ [ 56%] ........................................................................ [ 56%] ........................................................................ [ 57%] ........................................................................ [ 57%] ........................................................................ [ 57%] ........................................................................ [ 57%] ........................................................................ [ 57%] ........................................................................ [ 57%] ...............................x...x.x...x.x...x.x...x.............x...x [ 57%] .x...x.x...x.x...x.............x...x.x...x.x...x.x...x.............x...x [ 57%] .x...x.x...x.x...x.............x...x.x...x.x...x.x...x.............x...x [ 57%] .x...x.x...x.x...x.............x...x.x...x.x...x.x...x.............x...x [ 57%] .x...x.x...x.x...x.............x...x.x...x.x...x.x...x.............x...x [ 57%] .x...x.x...x.x...x.............x...x.x...x.x...x.x...x.............x...x [ 57%] .x...x.x...x.x...x.............x...x.x...x.x...x.x...x.............x...x [ 57%] .x...x.x...x.x...x.............x...x.x...x.x...x.x...x.............x...x [ 58%] .x...x.x...x.x...x.............x...x.x...x.x...x.x...x.............x...x [ 58%] .x...x.x...x.x...x.............x...x.x...x.x...x.x...x.............x...x [ 58%] .x...x.x...x.x...x.............x...x.x...x.x...x.x...x.............x...x [ 58%] .x...x.x...x.x...x.............x...x.x...x.x...x.x...x.............x...x [ 58%] .x...x.x...x.x...x.............x...x.x...x.x...x.x...x.............x...x [ 58%] .x...x.x...x.x...x.............x...x.x...x.x...x.x...x.............x...x [ 58%] .x...x.x...x.x...x.............x...x.x...x.x...x.x...x.............x...x [ 58%] .x...x.x...x.x...x.............x...x.x...x.x...x.x...x.............x...x [ 58%] .x...x.x...x.x...x...................................................... [ 58%] ........................................................................ [ 58%] ........................................................................ [ 58%] ........................................................................ [ 58%] ........................................................................ [ 58%] ........................................................................ [ 59%] ........................................................................ [ 59%] ........................................................................ [ 59%] ........................................................................ [ 59%] ........................................................................ [ 59%] ........................................................................ [ 59%] ........................................................................ [ 59%] ........................................................................ [ 59%] ........................................................................ [ 59%] ........................................................................ [ 59%] ........................................................................ [ 59%] .....................sssssssssssssssssssssssssssssssssssssssssssssssssss [ 59%] ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 59%] ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 60%] ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 60%] ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 60%] sssssssssssssssssssssssssssssssssssssssssssssss......................... [ 60%] ........................................................................ [ 60%] ........................................................................ [ 60%] ........................................................................ [ 60%] ..................................................x.....x.....x.....x... [ 60%] ..x.....x.....x.....x.....x.....x.....x.....x.......x........x..x..x.... [ 60%] ....s................................................................... [ 60%] ........................................................................ [ 60%] ........................................................................ [ 60%] ........................................................................ [ 60%] ........................................................................ [ 60%] ........................................................................ [ 61%] ........................................................................ [ 61%] ........................................................................ [ 61%] ........................................................................ [ 61%] ........................................................................ [ 61%] ................................x.....x....................x.....x...... [ 61%] ..............x.....x....................x.....x....................x... [ 61%] ..x....................x.....x....................x.....x............... [ 61%] .....x.....x............................................................ [ 61%] ........................................................................ [ 61%] ........................................................................ [ 61%] .................................................................s...... [ 61%] ........................................................................ [ 61%] ........................................................................ [ 62%] ........................................................................ [ 62%] ......................................FF..........FF....FF.............. [ 62%] ........FF....FF..........FF....FF......................FFss............ [ 62%] ........................................................................ [ 62%] ........................................................................ [ 62%] ........................................................................ [ 62%] ........................................................................ [ 62%] ........................................................................ [ 62%] ........................................................................ [ 62%] ........................................................................ [ 62%] ........................................................................ [ 62%] ........................................................................ [ 62%] ........................................................................ [ 62%] ....................................ssssssssssssssssssssssssssssssssssss [ 63%] ssssssssssss [ 63%] tests/cupyx_tests/scipy_tests/ndimage_tests/test_measurements.py ....... [ 63%] ........................................................................ [ 63%] ........................................................................ [ 63%] ........................................................................ [ 63%] ........................................................................ [ 63%] ........................................................................ [ 63%] ........................................................................ [ 63%] ........................................................................ [ 63%] ........................................................................ [ 63%] ........................................................................ [ 63%] ........................................................................ [ 63%] ........................................................................ [ 63%] ........................................................................ [ 63%] ........................................................................ [ 64%] ........................................................................ [ 64%] ........................................................................ [ 64%] ........................................................................ [ 64%] ........................................................................ [ 64%] ........................................................................ [ 64%] ........................................................................ [ 64%] ........................................................................ [ 64%] ........................................................................ [ 64%] ........................................................................ [ 64%] ........................................................................ [ 64%] ........................................................................ [ 64%] ........................................................................ [ 64%] ........................................................................ [ 65%] ........................................................................ [ 65%] ........................................................................ [ 65%] ........................................................................ [ 65%] ........................................................................ [ 65%] ........................................................................ [ 65%] ........................................................................ [ 65%] ........................................................................ [ 65%] ........................................................................ [ 65%] ........................................................................ [ 65%] ........................................................................ [ 65%] ........................................................................ [ 65%] ...................xx..........xx..........xx..........xx..........xx... [ 65%] ...................xx..................................xx............... [ 65%] ...................xx..........xx..........xx..........xx..........xx... [ 66%] ...................xx..................................xx............... [ 66%] ........................................................................ [ 66%] ........................................................................ [ 66%] ................. [ 66%] tests/cupyx_tests/scipy_tests/ndimage_tests/test_morphology.py ......... [ 66%] ......................s....s....s..s....s....s.......................... [ 66%] ..........s....s....s..s....s....s....................................s. [ 66%] ...s....s..s....s....s....................................s....s....s..s [ 66%] ....s....s....................................s....s....s..s....s....s.. [ 66%] ..................................s....s....s..s....s....s.............. [ 66%] ......................s....s....s..s....s....s.......................... [ 66%] ..........s....s....s..s....s....s....................................s. [ 66%] ...s....s..s....s....s....................................s....s....s..s [ 66%] ....s....s....................................s....s....s..s....s....s.. [ 66%] ..................................s....s....s..s....s....s.............. [ 67%] ......................s....s....s..s....s....s.......................... [ 67%] ..........s....s....s..s....s....s....................................s. [ 67%] ...s....s..s....s....s....................................s....s....s..s [ 67%] ....s....s....................................s....s....s..s....s....s.. [ 67%] ..................................s....s....s..s....s....s.............. [ 67%] ......................s....s....s..s....s....s.......................... [ 67%] ..........s....s....s..s....s....s....................................s. [ 67%] ...s....s..s....s....s....................................s....s....s..s [ 67%] ....s....s....................................s....s....s..s....s....s.. [ 67%] ..................................s....s....s..s....s....s.............. [ 67%] ........................................................................ [ 67%] ........................................................................ [ 67%] ........................................................................ [ 67%] ........................................................................ [ 68%] ........................................................................ [ 68%] ........................................................................ [ 68%] ........................................................................ [ 68%] .......................................................................s [ 68%] .s.s.............s.s.s.............s.s.s.............s.s.s.............s [ 68%] .s.s.............s.s.s.............s.s.s.............s.s.s.............s [ 68%] .s.s.............s.s.s.............s.s.s.............s.s.s.............s [ 68%] .s.s.............s.s.s.............s.s.s.............s.s.s.............s [ 68%] .s.s.............s.s.s.............s.s.s.............s.s.s.............s [ 68%] .s.s.............s.s.s.............s.s.s.............s.s.s.............s [ 68%] .s.s.............s.s.s.............s.s.s.............s.s.s.............s [ 68%] .s.s.............s.s.s.............s.s.s.............s.s.s.............s [ 68%] .s.s.............s.s.s.............s.s.s.............s.s.s.............s [ 69%] .s.s.............s.s.s.............s.s.s.............s.s.s.............s [ 69%] .s.s.............s.s.s.............s.s.s.............s.s.s.............s [ 69%] .s.s.............s.s.s.............s.s.s.............s.s.s.............s [ 69%] .s.s.............s.s.s.............s.s.s.............s.s.s.............s [ 69%] .s.s.............s.s.s.............s.s.s.............s.s.s.............s [ 69%] .s.s.............s.s.s.............s.s.s.............s.s.s.............s [ 69%] .s.s.............s.s.s.............s.s.s.............s.s.s.............s [ 69%] .s.s.............s.s.s.............s.s.s.............s.s.s.............s [ 69%] .s.s.............s.s.s.............s.s.s.............s.s.s.............s [ 69%] .s.s.............s.s.s.............s.s.s.............s.s.s.............s [ 69%] .s.s.............s.s.s.............s.s.s.............s.s.s.............s [ 69%] .s.s.............s.s.s.............s.s.s.............s.s.s.............s [ 69%] .s.s.............s.s.s.............s.s.s.............s.s.s.............s [ 69%] .s.s.............s.s.s.............s.s.s.............s.s.s.............s [ 70%] .s.s.............s.s.s.............s.s.s.............s.s.s.............s [ 70%] .s.s.............s.s.s.............s.s.s.............s.s.s.............s [ 70%] .s.s.............s.s.s.............s.s.s.............s.s.s.............s [ 70%] .s.s.............s.s.s.............s.s.s.............s.s.s.............s [ 70%] .s.s.............s.s.s.............s.s.s.............s.s.s.............s [ 70%] .s.s.............s.s.s.............s.s.s.............s.s.s.............s [ 70%] .s.s.............s.s.s.............s.s.s.............s.s.s.............s [ 70%] .s.s.............s.s.s.............s.s.s.............s.s.s.............s [ 70%] .s.s.............s.s.s.............s.s.s.............s.s.s.............s [ 70%] .s.s.............s.s.s.............s.s.s.............s.s.s.............s [ 70%] .s.s.............s.s.s.............s.s.s.............s.s.s.............s [ 70%] .s.s.............s.s.s.............s.s.s.............s.s.s.............s [ 70%] .s.s.............s.s.s.............s.s.s.............s.s.s.............. [ 71%] ........................................................................ [ 71%] ......s....s....s....s....s....s....s....s....s....s....s....s......s... [ 71%] .s....s....s....s....s....s....s....s....s....s....s.................... [ 71%] ..........................................s....s....s....s....s....s.... [ 71%] s....s....s....s....s....s......s....s....s....s....s....s....s....s.... [ 71%] s....s....s....s........................................................ [ 71%] ......s....s....s....s....s....s....s....s....s....s....s....s......s... [ 71%] .s....s....s....s....s....s....s....s....s....s....s.................... [ 71%] ..........................................s....s....s....s....s....s.... [ 71%] s....s....s....s....s....s......s....s....s....s....s....s....s....s.... [ 71%] s....s....s....s........................................................ [ 71%] ......s....s....s....s....s....s....s....s....s....s....s....s......s... [ 71%] .s....s....s....s....s....s....s....s....s....s....s.................... [ 71%] ..........................................s....s....s....s....s....s.... [ 72%] s....s....s....s....s....s......s....s....s....s....s....s....s....s.... [ 72%] s....s....s....s........................................................ [ 72%] ......s....s....s....s....s....s....s....s....s....s....s....s......s... [ 72%] .s....s....s....s....s....s....s....s....s....s....s.................... [ 72%] ..........................................s....s....s....s....s....s.... [ 72%] s....s....s....s....s....s......s....s....s....s....s....s....s....s.... [ 72%] s....s....s....s........................................................ [ 72%] ......s....s....s....s....s....s....s....s....s....s....s....s......s... [ 72%] .s....s....s....s....s....s....s....s....s....s....s.................... [ 72%] ..........................................s....s....s....s....s....s.... [ 72%] s....s....s....s....s....s......s....s....s....s....s....s....s....s.... [ 72%] s....s....s....s........................................................ [ 72%] ......s....s....s....s....s....s....s....s....s....s....s....s......s... [ 73%] .s....s....s....s....s....s....s....s....s....s....s.................... [ 73%] ..........................................s....s....s....s....s....s.... [ 73%] s....s....s....s....s....s......s....s....s....s....s....s....s....s.... [ 73%] s....s....s....s........................................................ [ 73%] ........................................................................ [ 73%] ........................................................................ [ 73%] ........................................................................ [ 73%] ........................................................ssssssss........ [ 73%] ........................................................ssssssss........ [ 73%] ........................................................ssssssss........ [ 73%] ........................................................ssssssss........ [ 73%] ........................................................................ [ 73%] ........................................................................ [ 73%] ........................................................................ [ 74%] ........................................................................ [ 74%] ........................................................................ [ 74%] ........................................................................ [ 74%] ........................................................................ [ 74%] ........................................................................ [ 74%] ........................................................................ [ 74%] ........................................................................ [ 74%] ........................................................................ [ 74%] ........................................................................ [ 74%] ........................................................................ [ 74%] ........................................................................ [ 74%] [ 74%] tests/cupyx_tests/scipy_tests/signal_tests/test_bsplines.py ............ [ 74%] [ 74%] tests/cupyx_tests/scipy_tests/signal_tests/test_signaltools.py ......... [ 74%] ........................................................................ [ 74%] ............................xxx............xxxx............F..x......... [ 75%] ....xxx............xxxx...............x.............xxx............xxxx. [ 75%] ..............x....................................x.x..x.x..x.x..x.x..x [ 75%] .x..x.x..x.x..x.x..x.x..x.x..x.x..x.x................................... [ 75%] ........................................................................ [ 75%] ........................................................................ [ 75%] ........................................................................ [ 75%] ........................................................................ [ 75%] ....................................................ssssssssssssssssss [ 75%] tests/cupyx_tests/scipy_tests/sparse_tests/test_base.py ....... [ 75%] tests/cupyx_tests/scipy_tests/sparse_tests/test_construct.py ........... [ 75%] ........................................................................ [ 75%] ................................................................s.s.s.s. [ 75%] s.s..................................................................... [ 75%] ..............xxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxx [ 75%] xxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxx [ 76%] xxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxx.................................. [ 76%] ..xxxxxxxxxxxxxxxxxxxxxxxx.............................................. [ 76%] ......................................xxxxxxxxxxxx...................... [ 76%] ..............xxxxxxxxxxxxxxxxxxxxxxxx.................................. [ 76%] ..................................................xxxxxxxxxxxx.......... [ 76%] ..........................xxxxxxxxxxxxxxxxxxxxxxxx...................... [ 76%] ..............................................................xxxxxxxxxx [ 76%] xxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxx [ 76%] xxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxx [ 76%] xxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxx [ 76%] xxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxx [ 76%] xxxxxxxxxxxxxxxxxxxxxxxxxx....................................xxxxxxxxxx [ 76%] xxxxxxxxxxxxxx.......................................................... [ 76%] ..........................xxxxxxxxxxxx.................................. [ 77%] ..xxxxxxxxxxxxxxxxxxxxxxxx.............................................. [ 77%] ......................................xxxxxxxxxxxx...................... [ 77%] ..............xxxxxxxxxxxxxxxxxxxxxxxx.................................. [ 77%] ..................................................xxxxxxxxxxxx.......... [ 77%] ..........................xxxxxxxxxxxxxxxxxxxxxxxx...................... [ 77%] ..............................................................xxxxxxxxxx [ 77%] xx....................................xxxxxxxxxxxxxxxxxxxxxxxx.......... [ 77%] ........................................................................ [ 77%] ..xxxxxxxxxxxx....................................xxxxxxxxxxxxxxxxxxxxxx [ 77%] xx...................................................................... [ 77%] ..............xxxxxxxxxxxx....................................xxxxxxxxxx [ 77%] xxxxxxxxxxxxxx.......................................................... [ 77%] ..........................xxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxx [ 78%] xxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxx.......... [ 78%] ........................................................................ [ 78%] ..xxxxxxxxxxxx.......................................................... [ 78%] ..........................xxxxxxxxxxxx.................................. [ 78%] ..................................................xxxxxxxxxxxx.......... [ 78%] ........................................................................ [ 78%] ..xxxxxxxxxxxx.......................................................... [ 78%] ..........................xxxxxxxxxxxx.................................. [ 78%] ..................................................xxxxxxxxxxxx.......... [ 78%] ........................................................................ [ 78%] .. [ 78%] tests/cupyx_tests/scipy_tests/sparse_tests/test_coo.py ................. [ 78%] ........................................................................ [ 78%] ........................................................................ [ 78%] ........................................................................ [ 79%] ........................................................................ [ 79%] ........................................................................ [ 79%] ........................................................................ [ 79%] ........................................................................ [ 79%] ........................................................................ [ 79%] ........................................................................ [ 79%] ........................................................................ [ 79%] ........................................................................ [ 79%] ........................................................................ [ 79%] ........................................................................ [ 79%] ........................................................................ [ 79%] ........................................................................ [ 79%] ........................................................................ [ 79%] ........................................................................ [ 80%] ........................................................................ [ 80%] ........................................................................ [ 80%] ........................................................................ [ 80%] ........................................................................ [ 80%] ........................................................................ [ 80%] ........................................................................ [ 80%] ........................................................................ [ 80%] ........................................................................ [ 80%] ............................................. [ 80%] tests/cupyx_tests/scipy_tests/sparse_tests/test_csc.py ................. [ 80%] ...................xxxx................................................. [ 80%] ........................................................................ [ 80%] .............x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x. [ 80%] .x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x [ 80%] .x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x [ 81%] ..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x.. [ 81%] x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x. [ 81%] x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.xxxxxxxxxxxxxxxxxxxxxx [ 81%] xxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxx..x.x..x.x..x.x..x.x..x.x..x.x..x [ 81%] .x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x [ 81%] ..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x.. [ 81%] x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x. [ 81%] x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x. [ 81%] .x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x [ 81%] .x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x [ 81%] ..x.x..x.x..x.x..x.x..x.x..x.x..x.xxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxx [ 81%] xxx..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x. [ 81%] x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x. [ 82%] .x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x [ 82%] .x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x [ 82%] ..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x.. [ 82%] x.x..x.x..x.x..x.x..x.xssssssssssssssssssss..x.x..x.x..x.x..x.x..x.x..x. [ 82%] x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x. [ 82%] .x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x [ 82%] .x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x [ 82%] ..x.x..x.x..x.xxxxxxx......xxxxxxxxxxxx......xxxxxx......xxxxxxxxxxxx... [ 82%] ...................................................ss........ss......... [ 82%] ...........................xxxx....xxxxxxxx............................. [ 82%] ........................................................................ [ 82%] ................................................................. [ 82%] tests/cupyx_tests/scipy_tests/sparse_tests/test_csr.py ................. [ 82%] ...................xxxx................................................. [ 83%] ........................................................................ [ 83%] .............x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x. [ 83%] .x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x [ 83%] .x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x [ 83%] ..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x.. [ 83%] x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x. [ 83%] x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x. [ 83%] .x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x [ 83%] .x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x [ 83%] ..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x.. [ 83%] x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x. [ 83%] x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x. [ 83%] .x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x [ 83%] .x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x [ 84%] ..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x.. [ 84%] x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x. [ 84%] x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x. [ 84%] .x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x [ 84%] .x..x.x..x.x..x.x..x.x..x.xxxxxxxxxxxxxxxxxxxxx..x.x..x.x..x.x..x.x..x.x [ 84%] ..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x.. [ 84%] x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x. [ 84%] x..x.x..x.x..x.x..x.x..x.x..x.xssssssssssssssssssss..x.x..x.x..x.x..x.x. [ 84%] .x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x [ 84%] .x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x [ 84%] ..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x.. [ 84%] x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x. [ 84%] x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x. [ 85%] .x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x [ 85%] .x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x..x.x [ 85%] ..x.x..x.x..x.x......................................xxxxxx............x [ 85%] xxxxx......xxxxxx............xxxxxx......xxxxxx............xxxxxx......x [ 85%] xxxxx............xxxxxx................................................s [ 85%] s........ss........................................xxxx........xxxx..... [ 85%] ........................................................................ [ 85%] ........................................................................ [ 85%] .ssssssssssss........................................................... [ 85%] ........................................................................ [ 85%] ........................................................................ [ 85%] ........................................................................ [ 85%] ........................................................................ [ 85%] ........................................................................ [ 86%] ........................................................................ [ 86%] ........................................................................ [ 86%] ........................................................................ [ 86%] ........................................................................ [ 86%] ........................................................................ [ 86%] ........................................................................ [ 86%] ........................................................................ [ 86%] ........................................................................ [ 86%] ........................................................................ [ 86%] ........................................................................ [ 86%] ........................................................................ [ 86%] ........................................................................ [ 86%] ........................................................................ [ 87%] ........................................................................ [ 87%] ........................................................................ [ 87%] ........................................................................ [ 87%] ........................................................................ [ 87%] ........................................................................ [ 87%] ........................................................................ [ 87%] ........................................................................ [ 87%] ........................................................................ [ 87%] ........................................................................ [ 87%] ........................................................................ [ 87%] ........................................................................ [ 87%] ........................................................................ [ 87%] ........................................................................ [ 87%] ........................................................................ [ 88%] ........................................................................ [ 88%] ........................................................................ [ 88%] ........................................................................ [ 88%] ........................................................................ [ 88%] ........................................................................ [ 88%] ........................................................................ [ 88%] ........................................................................ [ 88%] ........................................................................ [ 88%] ........................................................................ [ 88%] ........................................................................ [ 88%] ........................................................................ [ 88%] ........................................................................ [ 88%] ........................................................................ [ 89%] ........................................................................ [ 89%] ........................................................................ [ 89%] .....................ssss........ssss........ssss........ssss........sss [ 89%] s........ssss........ssss........ssss........ssss........ssss........sss [ 89%] s........ssss........ssss........ssss........ssss........ssss........... [ 89%] ........................................................................ [ 89%] ........................................................................ [ 89%] ........................................................................ [ 89%] ........................................................................ [ 89%] ........................................................................ [ 89%] ......................... [ 89%] tests/cupyx_tests/scipy_tests/sparse_tests/test_dia.py ................. [ 89%] ........................................................................ [ 89%] ......xxxxxxxxxxxxxxx...............xxxxxxxxxxxxxxx...............xxxxxx [ 89%] xxxxxxxxx...............xxxxxxxxxxxxxxx............xxxxxxxxxxxx......... [ 90%] ...............xxxxxxxxxxxx.. [ 90%] tests/cupyx_tests/scipy_tests/sparse_tests/test_extract.py ............. [ 90%] ........................................................................ [ 90%] ....................... [ 90%] tests/cupyx_tests/scipy_tests/sparse_tests/test_index.py ............... [ 90%] ........................................................................ [ 90%] ........................................................................ [ 90%] ........................................................................ [ 90%] ........................................................................ [ 90%] ........................................................................ [ 90%] ........................................................................ [ 90%] ........................................................................ [ 90%] ........................................................................ [ 90%] ........................................................................ [ 90%] ........................................................................ [ 90%] ........................................................................ [ 90%] ........................................................................ [ 91%] ........................................................................ [ 91%] ........................................................................ [ 91%] ........................................................................ [ 91%] ........................................................................ [ 91%] ........................................................................ [ 91%] ........................................................................ [ 91%] ........................................................................ [ 91%] ........................................................................ [ 91%] ........................................................................ [ 91%] ........................................................................ [ 91%] ........................................................................ [ 91%] ........................................................................ [ 91%] ........................................................................ [ 92%] ........................................................................ [ 92%] ........................................................................ [ 92%] ........................................................................ [ 92%] ........................................................................ [ 92%] ........................................................................ [ 92%] ........................................................................ [ 92%] ........................................................................ [ 92%] ........................................................................ [ 92%] ........................................................................ [ 92%] ........................................................................ [ 92%] ........................................................................ [ 92%] ........................................................................ [ 92%] ....................................x...x.....................x...x..... [ 92%] ......x...x........................x...x.........xx..xx................. [ 93%] .................................x...x.................................. [ 93%] ...............x...x.....................x...x.x.........x...x.x........ [ 93%] ..............x...x.........x...x.....................................x. [ 93%] ..x.........x...x.........x...x.........x...x........................... [ 93%] ........................................................................ [ 93%] ........................................................................ [ 93%] ........................................................................ [ 93%] ........................................................................ [ 93%] ........................................................................ [ 93%] ..................................................x...x...........x...x. [ 93%] .......................x...x.........xx..xx............................. [ 93%] .....................x...x.............................................. [ 93%] .............................x...x.x.........x...x.x.................... [ 94%] ..x...x.........x...x.....................................x...x......... [ 94%] x...x.........x...x.........x...x....................................... [ 94%] ........................................................................ [ 94%] ........................................................................ [ 94%] ........................................................................ [ 94%] ........................................................................ [ 94%] ........................................................................ [ 94%] ......................................x...x...........x...x............. [ 94%] ...........x...x.........xx..xx......................................... [ 94%] .........x...x.......................................................... [ 94%] .................x...x.x.........x...x.x......................x...x..... [ 94%] ....x...x.....................................x...x.........x...x....... [ 94%] ..x...x.........x...x................................................... [ 94%] ........................................................................ [ 95%] ........................................................................ [ 95%] ................................. [ 95%] tests/cupyx_tests/scipy_tests/sparse_tests/test_linalg.py ssssssss.x.x.x [ 95%] ..x..x.x.x..x..x.x.x..x..x.x.x..x..x.x.x..x..x.x.x..x..x.x.x..x..x.x.x.. [ 95%] x..x.x.x..x..x.x.x..x..x.x.x..x..x.x.x..x.xx..xxxxxxxxxx....xx..xxxxxxxx [ 95%] xx....xx..xxxxxxxxxx....xx..xxxxxxxxxx.................................. [ 95%] ..........................................xx..xxxxxxxxxx....xx..xxxxxxxx [ 95%] xx....xx..xxxxxxxxxx....xx..xxxxxxxxxx.................................. [ 95%] ..................................................................xxxxxx [ 95%] xxxxxxxxxxxxxxxxxx................................................xxxxxx [ 95%] xxxxxxxxxxxxxxxxxxss..ss..ss..ss..ss..ss..xxxxxxxxxxxxxxxxxxxxxxxxxxxxxx [ 95%] xxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxx.............................. [ 95%] ..........................................xxxxxxxxxxxxxxxxxxxxxxxxxxxxxx [ 95%] xxxxxxs.s.s.s.s.s.s.s.s.s.s.s.s.s.s.s.s.s............................... [ 95%] ..................................xxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxx...... [ 96%] ........................................................................ [ 96%] ...................sss.sssssssssssssssssssssssssss...................... [ 96%] ........................................................................ [ 96%] ..................................xxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxx [ 96%] xxxxxxxxxxxxxxxxxxxxxxxxxx.............................................. [ 96%] ........................................................................ [ 96%] ........................................................................ [ 96%] ...sssssss.sssssssssssssssssssssssssssssssssssssssssssssssssssssss....xx [ 96%] .............xx.............xx.............xx.............xx............ [ 96%] .xx.............xx.............xx.............xx.............xx......... [ 96%] ....xx.............xx.............xx.............xx.............xx...... [ 96%] .......xx.............xx.............xx.............xx.............xx... [ 96%] ..........xx.............xx.............xx.............xx.............xx [ 97%] .............xx.............xx.............xx.............xx............ [ 97%] .xx.............xx.............xx............x..............x........... [ 97%] ...x..............x..............x..............x..............x........ [ 97%] ......x..............x..............x..............x..............x..... [ 97%] .........x..............x..............x..............x..............x.. [ 97%] ............x..............x..............x..............x.............. [ 97%] x..............x..............x..............x..............x........... [ 97%] ...x..............x..............x..............x..............x........ [ 97%] ......x...............xx.............xx.............xx.............xx... [ 97%] ..........xx.............xx.............xx.............xx.............xx [ 97%] .............xx.............xx.............xx.............xx............ [ 97%] .xx.............xx.............xx.............xx.............xx......... [ 97%] ....xx.............xx.............xx.............xx.............xx...... [ 97%] .......xx.............xx.............xx.............xx.............xx... [ 98%] ..........xx.............xx.............xx.............xx............... [ 98%] ..................................................................ssssss [ 98%] ss.sssssssssssssssssssssss.F............F............................... [ 98%] ..............................................s.s.....s.s.....s.sxxxxxxx [ 98%] xxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxx........................xxxxxxx [ 98%] xxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxx........................xxxxxxx [ 98%] xxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxx.....s.s.....s.s.....s.s.....s. [ 98%] s.....s.s.....s.s.....s.s.....s.s.....s.s............................... [ 98%] ........................................................................ [ 98%] ..........................................sss.sss.sss.sss.sss.sss.sss.ss [ 98%] s.sss.sss.sss.sss.sss.sss.sss.sss.sss.sss.sss.sss.sss.sss.sss.sss.sss.ss [ 98%] s.sss.sss.sss.sss.sss.sss.sss.sss.sss.sss.sss.sss.sss.sss.sss.sss.sss.ss [ 98%] s.sss.sss.sss.sss.sss.sss.sss.sss.sss.sss............................... [ 99%] .................................xxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxx....... [ 99%] ........................................................................ [ 99%] ..................sss.sssssssssssssssssssssssssss....................... [ 99%] .........xxxxxxxxxxxxxxxx................................xxxxxxxxxxxxxxx [ 99%] x....................................................................... [ 99%] ..........................sss.sss.sss.sss.sss.sss.sss.sss.sss.sss.sss.ss [ 99%] sssssssssssssssssssssssssssssssssssssssssssssssss.sss.sss.sss.sss.sss.ss [ 99%] s.sss.sss.sss.sss.sss.ssssssssssssssssssssssssssssssssssssssssssssssssss [ 99%] s [ 99%] tests/cupyx_tests/scipy_tests/sparse_tests/csgraph_tests/test_traversal.py s [ 99%] ssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 99%] sssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssssss [ 99%] tests/cupyx_tests/scipy_tests/special_tests/test_bessel.py ............ [ 99%] tests/cupyx_tests/scipy_tests/special_tests/test_convex_analysis.py .... [ 99%] . [ 99%] tests/cupyx_tests/scipy_tests/special_tests/test_digamma.py ..... [ 99%] tests/cupyx_tests/scipy_tests/special_tests/test_erf.py ............ [ 99%] tests/cupyx_tests/scipy_tests/special_tests/test_gamma.py .... [ 99%] tests/cupyx_tests/scipy_tests/special_tests/test_gammaln.py .... [ 99%] tests/cupyx_tests/scipy_tests/special_tests/test_polygamma.py .... [ 99%] tests/cupyx_tests/scipy_tests/special_tests/test_statistics.py .. [ 99%] tests/cupyx_tests/scipy_tests/special_tests/test_zeta.py .... [ 99%] tests/cupyx_tests/scipy_tests/stats_tests/test_distributions.py ........ [ 99%] ........................................................................ [ 99%] ........................................................................ [ 99%] [ 99%] tests/cupyx_tests/tools_tests/test_install_library.py ...... [ 99%] tests/example_tests/test_custom_struct.py ... [ 99%] tests/example_tests/test_finance.py ... [ 99%] tests/example_tests/test_gemm.py . [ 99%] tests/example_tests/test_gmm.py ss [ 99%] tests/example_tests/test_kmeans.py sss [ 99%] tests/install_tests/test_build.py sx. [ 99%] tests/install_tests/test_utils.py .. [ 99%] tests/install_tests/test_cupy_builder/test_context.py .... [100%] =================================== FAILURES =================================== ______________ TestNdarrayCopy.test_copy_multi_device_with_stream ______________ self = @testing.multi_gpu(2) def test_copy_multi_device_with_stream(self): # Kernel that takes long enough then finally writes values. kern = cupy.RawKernel( _test_copy_multi_device_with_stream_src, 'wait_and_write') # Allocates a memory and launches the kernel on a device with its # stream. with cuda.Device(0): # Keep this stream alive over the D2D copy below for HIP with cuda.Stream() as s1: # NOQA a = cupy.zeros((2,), dtype=numpy.uint64) kern((1,), (1,), a) # D2D copy to another device with another stream should get the # original values of the memory before the kernel on the first device # finally makes the write. with cuda.Device(1): with cuda.Stream(): b = a.copy() > testing.assert_array_equal( b, numpy.array([0, 0], dtype=numpy.uint64)) tests/cupy_tests/core_tests/test_ndarray.py:224: _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ x = array([13744632839234567870, 13744632839234567870], dtype=uint64) y = array([0, 0], dtype=uint64), err_msg = '', verbose = True strides_check = False def assert_array_equal(x, y, err_msg='', verbose=True, strides_check=False): """Raises an AssertionError if two array_like objects are not equal. Args: x(numpy.ndarray or cupy.ndarray): The actual object to check. y(numpy.ndarray or cupy.ndarray): The desired, expected object. strides_check(bool): If ``True``, consistency of strides is also checked. err_msg(str): The error message to be printed in case of failure. verbose(bool): If ``True``, the conflicting values are appended to the error message. .. seealso:: :func:`numpy.testing.assert_array_equal` """ > numpy.testing.assert_array_equal( cupy.asnumpy(x), cupy.asnumpy(y), err_msg=err_msg, verbose=verbose) E AssertionError: E Arrays are not equal E E Mismatched elements: 2 / 2 (100%) E Max absolute difference: 13744632839234567870 E Max relative difference: inf E x: array([13744632839234567870, 13744632839234567870], dtype=uint64) E y: array([0, 0], dtype=uint64) cupy/testing/_array.py:91: AssertionError ________________ TestArrayCopyAndView.test_astype_boolean_view _________________ args = (,) kw = {}, cupy_result = (array([0, 1, 2], dtype=int8),), cupy_error = None numpy_result = (array([0, 1, 1], dtype=int8),), numpy_error = None cupy_numpy_result_ndarrays = [(array([0, 1, 2], dtype=int8), array([0, 1, 1], dtype=int8))] cupy_r = array([0, 1, 2], dtype=int8), numpy_r = array([0, 1, 1], dtype=int8) masks = [None] @_wraps_partial_xp(impl, name, sp_name, scipy_name) def test_func(*args, **kw): # Run cupy and numpy ( cupy_result, cupy_error, numpy_result, numpy_error) = ( _call_func_numpy_cupy( impl, args, kw, name, sp_name, scipy_name)) assert cupy_result is not None or cupy_error is not None assert numpy_result is not None or numpy_error is not None # Check errors raised if cupy_error or numpy_error: _check_cupy_numpy_error(cupy_error, numpy_error, accept_error=accept_error) return # Check returned arrays if not isinstance(cupy_result, (tuple, list)): cupy_result = cupy_result, if not isinstance(numpy_result, (tuple, list)): numpy_result = numpy_result, assert len(cupy_result) == len(numpy_result) # Check types cupy_numpy_result_ndarrays = [ _convert_output_to_ndarray( cupy_r, numpy_r, sp_name, check_sparse_format) for cupy_r, numpy_r in zip(cupy_result, numpy_result)] # Check dtypes if type_check: for cupy_r, numpy_r in cupy_numpy_result_ndarrays: if cupy_r.dtype != numpy_r.dtype: raise AssertionError( '''ndarrays of different dtypes are returned. cupy: {} numpy: {}'''.format(cupy_r.dtype, numpy_r.dtype)) # Check contiguities if contiguous_check: for cupy_r, numpy_r in zip(cupy_result, numpy_result): if isinstance(numpy_r, numpy.ndarray): if (numpy_r.flags.c_contiguous and not cupy_r.flags.c_contiguous): raise AssertionError( 'The state of c_contiguous flag is false. ' '(cupy_result:{} numpy_result:{})'.format( cupy_r.flags.c_contiguous, numpy_r.flags.c_contiguous)) if (numpy_r.flags.f_contiguous and not cupy_r.flags.f_contiguous): raise AssertionError( 'The state of f_contiguous flag is false. ' '(cupy_result:{} numpy_result:{})'.format( cupy_r.flags.f_contiguous, numpy_r.flags.f_contiguous)) # Check shapes for cupy_r, numpy_r in cupy_numpy_result_ndarrays: assert cupy_r.shape == numpy_r.shape masks = [None] * len(cupy_result) if _contains_signed_and_unsigned(kw): needs_mask = [ cupy_r.dtype in _unsigned_dtypes for cupy_r in cupy_result] if any(needs_mask): masks = _make_positive_masks( impl, args, kw, name, sp_name, scipy_name) for i, flag in enumerate(needs_mask): if not flag: masks[i] = None # Check item values for (cupy_r, numpy_r), mask in zip( cupy_numpy_result_ndarrays, masks): # Behavior of assigning a negative value to an unsigned integer # variable is undefined. # nVidia GPUs and Intel CPUs behave differently. # To avoid this difference, we need to ignore dimensions whose # values are negative. skip = False if mask is not None: if cupy_r.shape == (): skip = (mask == 0).all() else: cupy_r = cupy_r[mask].get() numpy_r = numpy_r[mask] if not skip: > check_func(cupy_r, numpy_r) cupy/testing/_helper.py:355: _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ cupy/testing/_helper.py:660: in check_func _array.assert_array_equal(x, y, err_msg, verbose, strides_check) _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ x = array([0, 1, 2], dtype=int8), y = array([0, 1, 1], dtype=int8), err_msg = '' verbose = True, strides_check = False def assert_array_equal(x, y, err_msg='', verbose=True, strides_check=False): """Raises an AssertionError if two array_like objects are not equal. Args: x(numpy.ndarray or cupy.ndarray): The actual object to check. y(numpy.ndarray or cupy.ndarray): The desired, expected object. strides_check(bool): If ``True``, consistency of strides is also checked. err_msg(str): The error message to be printed in case of failure. verbose(bool): If ``True``, the conflicting values are appended to the error message. .. seealso:: :func:`numpy.testing.assert_array_equal` """ > numpy.testing.assert_array_equal( cupy.asnumpy(x), cupy.asnumpy(y), err_msg=err_msg, verbose=verbose) E AssertionError: E Arrays are not equal E E Mismatched elements: 1 / 3 (33.3%) E Max absolute difference: 1 E Max relative difference: 1. E x: array([0, 1, 2], dtype=int8) E y: array([0, 1, 1], dtype=int8) cupy/testing/_array.py:91: AssertionError _______________ TestArrayElementwiseOp.test_add_array_boolarray ________________ self = def test_add_array_boolarray(self): > self.check_array_boolarray_op(operator.add) tests/cupy_tests/core_tests/test_ndarray_elementwise_op.py:478: _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ cupy/testing/_helper.py:837: in test_func impl(*args, **kw) cupy/testing/_helper.py:355: in test_func check_func(cupy_r, numpy_r) cupy/testing/_helper.py:507: in check_func _array.assert_allclose(c, n, rtol1, atol1, err_msg, verbose) _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ actual = array([[ 5., 8., 5.], [263., 253., 255.]]) desired = array([[3., 8., 2.], [9., 3., 9.]]), rtol = 1e-07, atol = 0 err_msg = '', verbose = True def assert_allclose(actual, desired, rtol=1e-7, atol=0, err_msg='', verbose=True): """Raises an AssertionError if objects are not equal up to desired tolerance. Args: actual(numpy.ndarray or cupy.ndarray): The actual object to check. desired(numpy.ndarray or cupy.ndarray): The desired, expected object. rtol(float): Relative tolerance. atol(float): Absolute tolerance. err_msg(str): The error message to be printed in case of failure. verbose(bool): If ``True``, the conflicting values are appended to the error message. .. seealso:: :func:`numpy.testing.assert_allclose` """ # NOQA > numpy.testing.assert_allclose( cupy.asnumpy(actual), cupy.asnumpy(desired), rtol=rtol, atol=atol, err_msg=err_msg, verbose=verbose) E AssertionError: E Not equal to tolerance rtol=1e-07, atol=0 E E Mismatched elements: 5 / 6 (83.3%) E Max absolute difference: 254. E Max relative difference: 83.33333333 E x: array([[ 5., 8., 5.], E [263., 253., 255.]]) E y: array([[3., 8., 2.], E [9., 3., 9.]]) cupy/testing/_array.py:24: AssertionError ----------------------------- Captured stdout call ----------------------------- x_type is _______________ TestArrayElementwiseOp.test_iadd_array_boolarray _______________ self = def test_iadd_array_boolarray(self): > self.check_array_boolarray_op(operator.iadd) tests/cupy_tests/core_tests/test_ndarray_elementwise_op.py:481: _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ cupy/testing/_helper.py:837: in test_func impl(*args, **kw) cupy/testing/_helper.py:355: in test_func check_func(cupy_r, numpy_r) cupy/testing/_helper.py:507: in check_func _array.assert_allclose(c, n, rtol1, atol1, err_msg, verbose) _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ actual = array([[ 5., 8., 5.], [263., 253., 255.]]) desired = array([[3., 8., 2.], [9., 3., 9.]]), rtol = 1e-07, atol = 0 err_msg = '', verbose = True def assert_allclose(actual, desired, rtol=1e-7, atol=0, err_msg='', verbose=True): """Raises an AssertionError if objects are not equal up to desired tolerance. Args: actual(numpy.ndarray or cupy.ndarray): The actual object to check. desired(numpy.ndarray or cupy.ndarray): The desired, expected object. rtol(float): Relative tolerance. atol(float): Absolute tolerance. err_msg(str): The error message to be printed in case of failure. verbose(bool): If ``True``, the conflicting values are appended to the error message. .. seealso:: :func:`numpy.testing.assert_allclose` """ # NOQA > numpy.testing.assert_allclose( cupy.asnumpy(actual), cupy.asnumpy(desired), rtol=rtol, atol=atol, err_msg=err_msg, verbose=verbose) E AssertionError: E Not equal to tolerance rtol=1e-07, atol=0 E E Mismatched elements: 5 / 6 (83.3%) E Max absolute difference: 254. E Max relative difference: 83.33333333 E x: array([[ 5., 8., 5.], E [263., 253., 255.]]) E y: array([[3., 8., 2.], E [9., 3., 9.]]) cupy/testing/_array.py:24: AssertionError ----------------------------- Captured stdout call ----------------------------- x_type is _ TestDot_param_0_{shape=((2, 3, 4), (3, 4, 2)), trans_a=True, trans_b=True}.test_dot _ self = < parameter: {'shape': ((2, 3, 4), (3, 4, 2)), 'trans_a': True, 'trans_b': True}> xp = dtype_a = , dtype_b = @testing.for_all_dtypes_combination(['dtype_a', 'dtype_b']) @testing.numpy_cupy_allclose() def test_dot(self, xp, dtype_a, dtype_b): shape_a, shape_b = self.shape if self.trans_a: a = testing.shaped_arange(shape_a[::-1], xp, dtype_a).T else: a = testing.shaped_arange(shape_a, xp, dtype_a) if self.trans_b: b = testing.shaped_arange(shape_b[::-1], xp, dtype_b).T else: b = testing.shaped_arange(shape_b, xp, dtype_b) > return xp.dot(a, b) tests/cupy_tests/linalg_tests/test_product.py:49: _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ cupy/linalg/_product.py:65: in dot return a.dot(b, out) cupy/_core/core.pyx:1607: in cupy._core.core.ndarray.dot return _linalg.dot(self, b, out) cupy/_core/_routines_linalg.pyx:427: in cupy._core._routines_linalg.dot return tensordot_core(a, b, out, n, m, k, ret_shape) cupy/_core/_routines_linalg.pyx:501: in cupy._core._routines_linalg.tensordot_core return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) cupy/_core/_routines_linalg.pyx:327: in cupy._core._routines_linalg._integral_tensordot_core kern = _tensordot_core_int_kernel(config, dtype) cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/_routines_linalg.pyx:302: in cupy._core._routines_linalg._tensordot_core_int_kernel ker = mod.get_function( cupy/_core/raw.pyx:470: in cupy._core.raw.RawModule.get_function mangled_name = self.module.mapping.get(name) cupy/_core/raw.pyx:394: in cupy._core.raw.RawModule.module.__get__ return self._module() cupy/_core/raw.pyx:402: in cupy._core.raw.RawModule._module mod = _get_raw_module( cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/raw.pyx:547: in cupy._core.raw._get_raw_module mod = cupy._core.core.compile_with_cache( cupy/_core/core.pyx:2015: in cupy._core.core.compile_with_cache cpdef function.Module compile_with_cache( cupy/_core/core.pyx:2078: in cupy._core.core.compile_with_cache return cuda.compile_with_cache( cupy/cuda/compiler.py:453: in compile_with_cache return _compile_with_cache_hip( cupy/cuda/compiler.py:853: in _compile_with_cache_hip binary, mapping = compile_using_nvrtc( cupy/cuda/compiler.py:295: in compile_using_nvrtc return _compile(source, options, cu_path, cupy/cuda/compiler.py:279: in _compile compiled_obj, mapping = prog.compile(options, log_stream) _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ self = options = ('--std=c++11', '-I/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/_core/include', '-I/opt/rocm/include', '-fcuda-flush-denormals-to-zero') log_stream = None def compile(self, options=(), log_stream=None): try: if self.name_expressions: for ker in self.name_expressions: nvrtc.addAddNameExpression(self.ptr, ker) nvrtc.compileProgram(self.ptr, options) mapping = None if self.name_expressions: mapping = {} for ker in self.name_expressions: mapping[ker] = nvrtc.getLoweredName(self.ptr, ker) if log_stream is not None: log_stream.write(nvrtc.getProgramLog(self.ptr)) # This is to ensure backwards compatibility with nvrtc if self.method == 'cubin': return nvrtc.getCUBIN(self.ptr), mapping elif self.method == 'ptx': return nvrtc.getPTX(self.ptr), mapping else: raise RuntimeError('Unknown NVRTC compile method') except nvrtc.NVRTCError: log = nvrtc.getProgramLog(self.ptr) > raise CompileException(log, self.src, self.name, options, 'nvrtc' if not runtime.is_hip else 'hiprtc') E AssertionError: Only cupy raises error E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 650, in compile E nvrtc.compileProgram(self.ptr, options) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 136, in cupy_backends.cuda.libs.nvrtc.compileProgram E cpdef compileProgram(intptr_t prog, options): E File "cupy_backends/cuda/libs/nvrtc.pyx", line 148, in cupy_backends.cuda.libs.nvrtc.compileProgram E check_status(status) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 67, in cupy_backends.cuda.libs.nvrtc.check_status E raise NVRTCError(status) E cupy_backends.cuda.libs.nvrtc.NVRTCError: HIPRTC_ERROR_COMPILATION (6) E E During handling of the above exception, another exception occurred: E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/testing/_helper.py", line 47, in _call_func E result = impl(*args, **kw) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/tests/cupy_tests/linalg_tests/test_product.py", line 49, in test_dot E return xp.dot(a, b) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/linalg/_product.py", line 65, in dot E return a.dot(b, out) E File "cupy/_core/core.pyx", line 1607, in cupy._core.core.ndarray.dot E return _linalg.dot(self, b, out) E File "cupy/_core/_routines_linalg.pyx", line 427, in cupy._core._routines_linalg.dot E return tensordot_core(a, b, out, n, m, k, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 501, in cupy._core._routines_linalg.tensordot_core E return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 327, in cupy._core._routines_linalg._integral_tensordot_core E kern = _tensordot_core_int_kernel(config, dtype) E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/_routines_linalg.pyx", line 302, in cupy._core._routines_linalg._tensordot_core_int_kernel E ker = mod.get_function( E File "cupy/_core/raw.pyx", line 470, in cupy._core.raw.RawModule.get_function E mangled_name = self.module.mapping.get(name) E File "cupy/_core/raw.pyx", line 394, in cupy._core.raw.RawModule.module.__get__ E return self._module() E File "cupy/_core/raw.pyx", line 402, in cupy._core.raw.RawModule._module E mod = _get_raw_module( E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/raw.pyx", line 547, in cupy._core.raw._get_raw_module E mod = cupy._core.core.compile_with_cache( E File "cupy/_core/core.pyx", line 2015, in cupy._core.core.compile_with_cache E cpdef function.Module compile_with_cache( E File "cupy/_core/core.pyx", line 2078, in cupy._core.core.compile_with_cache E return cuda.compile_with_cache( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 453, in compile_with_cache E return _compile_with_cache_hip( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 853, in _compile_with_cache_hip E binary, mapping = compile_using_nvrtc( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 295, in compile_using_nvrtc E return _compile(source, options, cu_path, E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 279, in _compile E compiled_obj, mapping = prog.compile(options, log_stream) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 667, in compile E raise CompileException(log, self.src, self.name, options, E cupy.cuda.compiler.CompileException: /tmp/comgr-779aba/input/CompileSource:203:27: error: redeclaration of '__hiprtc_16' with a different type: 'void (*const)(int, int, int, const signed char *, const signed char *, signed char *)' vs 'void (*const)(int, int, int, const bool *, const bool *, bool *)' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-779aba/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-779aba/input/CompileSource:204:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-779aba/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-779aba/input/CompileSource:205:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-779aba/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-779aba/input/CompileSource:206:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-779aba/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-779aba/input/CompileSource:207:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-779aba/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-779aba/input/CompileSource:208:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-779aba/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-779aba/input/CompileSource:209:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-779aba/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-779aba/input/CompileSource:210:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-779aba/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-779aba/input/CompileSource:211:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-779aba/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-779aba/input/CompileSource:212:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-779aba/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E 10 errors generated when compiling for gfx908. E Error: Failed to compile opencl source (from CL or HIP source to LLVM IR). cupy/cuda/compiler.py:667: AssertionError ----------------------------- Captured stdout call ----------------------------- {'dtype_a': , 'dtype_b': } _ TestDot_param_1_{shape=((2, 3, 4), (3, 4, 2)), trans_a=True, trans_b=False}.test_dot _ self = < parameter: {'shape': ((2, 3, 4), (3, 4, 2)), 'trans_a': True, 'trans_b': False}> xp = dtype_a = , dtype_b = @testing.for_all_dtypes_combination(['dtype_a', 'dtype_b']) @testing.numpy_cupy_allclose() def test_dot(self, xp, dtype_a, dtype_b): shape_a, shape_b = self.shape if self.trans_a: a = testing.shaped_arange(shape_a[::-1], xp, dtype_a).T else: a = testing.shaped_arange(shape_a, xp, dtype_a) if self.trans_b: b = testing.shaped_arange(shape_b[::-1], xp, dtype_b).T else: b = testing.shaped_arange(shape_b, xp, dtype_b) > return xp.dot(a, b) tests/cupy_tests/linalg_tests/test_product.py:49: _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ cupy/linalg/_product.py:65: in dot return a.dot(b, out) cupy/_core/core.pyx:1607: in cupy._core.core.ndarray.dot return _linalg.dot(self, b, out) cupy/_core/_routines_linalg.pyx:427: in cupy._core._routines_linalg.dot return tensordot_core(a, b, out, n, m, k, ret_shape) cupy/_core/_routines_linalg.pyx:501: in cupy._core._routines_linalg.tensordot_core return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) cupy/_core/_routines_linalg.pyx:327: in cupy._core._routines_linalg._integral_tensordot_core kern = _tensordot_core_int_kernel(config, dtype) cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/_routines_linalg.pyx:302: in cupy._core._routines_linalg._tensordot_core_int_kernel ker = mod.get_function( cupy/_core/raw.pyx:470: in cupy._core.raw.RawModule.get_function mangled_name = self.module.mapping.get(name) cupy/_core/raw.pyx:394: in cupy._core.raw.RawModule.module.__get__ return self._module() cupy/_core/raw.pyx:402: in cupy._core.raw.RawModule._module mod = _get_raw_module( cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/raw.pyx:547: in cupy._core.raw._get_raw_module mod = cupy._core.core.compile_with_cache( cupy/_core/core.pyx:2015: in cupy._core.core.compile_with_cache cpdef function.Module compile_with_cache( cupy/_core/core.pyx:2078: in cupy._core.core.compile_with_cache return cuda.compile_with_cache( cupy/cuda/compiler.py:453: in compile_with_cache return _compile_with_cache_hip( cupy/cuda/compiler.py:853: in _compile_with_cache_hip binary, mapping = compile_using_nvrtc( cupy/cuda/compiler.py:295: in compile_using_nvrtc return _compile(source, options, cu_path, cupy/cuda/compiler.py:279: in _compile compiled_obj, mapping = prog.compile(options, log_stream) _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ self = options = ('--std=c++11', '-I/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/_core/include', '-I/opt/rocm/include', '-fcuda-flush-denormals-to-zero') log_stream = None def compile(self, options=(), log_stream=None): try: if self.name_expressions: for ker in self.name_expressions: nvrtc.addAddNameExpression(self.ptr, ker) nvrtc.compileProgram(self.ptr, options) mapping = None if self.name_expressions: mapping = {} for ker in self.name_expressions: mapping[ker] = nvrtc.getLoweredName(self.ptr, ker) if log_stream is not None: log_stream.write(nvrtc.getProgramLog(self.ptr)) # This is to ensure backwards compatibility with nvrtc if self.method == 'cubin': return nvrtc.getCUBIN(self.ptr), mapping elif self.method == 'ptx': return nvrtc.getPTX(self.ptr), mapping else: raise RuntimeError('Unknown NVRTC compile method') except nvrtc.NVRTCError: log = nvrtc.getProgramLog(self.ptr) > raise CompileException(log, self.src, self.name, options, 'nvrtc' if not runtime.is_hip else 'hiprtc') E AssertionError: Only cupy raises error E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 650, in compile E nvrtc.compileProgram(self.ptr, options) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 136, in cupy_backends.cuda.libs.nvrtc.compileProgram E cpdef compileProgram(intptr_t prog, options): E File "cupy_backends/cuda/libs/nvrtc.pyx", line 148, in cupy_backends.cuda.libs.nvrtc.compileProgram E check_status(status) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 67, in cupy_backends.cuda.libs.nvrtc.check_status E raise NVRTCError(status) E cupy_backends.cuda.libs.nvrtc.NVRTCError: HIPRTC_ERROR_COMPILATION (6) E E During handling of the above exception, another exception occurred: E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/testing/_helper.py", line 47, in _call_func E result = impl(*args, **kw) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/tests/cupy_tests/linalg_tests/test_product.py", line 49, in test_dot E return xp.dot(a, b) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/linalg/_product.py", line 65, in dot E return a.dot(b, out) E File "cupy/_core/core.pyx", line 1607, in cupy._core.core.ndarray.dot E return _linalg.dot(self, b, out) E File "cupy/_core/_routines_linalg.pyx", line 427, in cupy._core._routines_linalg.dot E return tensordot_core(a, b, out, n, m, k, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 501, in cupy._core._routines_linalg.tensordot_core E return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 327, in cupy._core._routines_linalg._integral_tensordot_core E kern = _tensordot_core_int_kernel(config, dtype) E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/_routines_linalg.pyx", line 302, in cupy._core._routines_linalg._tensordot_core_int_kernel E ker = mod.get_function( E File "cupy/_core/raw.pyx", line 470, in cupy._core.raw.RawModule.get_function E mangled_name = self.module.mapping.get(name) E File "cupy/_core/raw.pyx", line 394, in cupy._core.raw.RawModule.module.__get__ E return self._module() E File "cupy/_core/raw.pyx", line 402, in cupy._core.raw.RawModule._module E mod = _get_raw_module( E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/raw.pyx", line 547, in cupy._core.raw._get_raw_module E mod = cupy._core.core.compile_with_cache( E File "cupy/_core/core.pyx", line 2015, in cupy._core.core.compile_with_cache E cpdef function.Module compile_with_cache( E File "cupy/_core/core.pyx", line 2078, in cupy._core.core.compile_with_cache E return cuda.compile_with_cache( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 453, in compile_with_cache E return _compile_with_cache_hip( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 853, in _compile_with_cache_hip E binary, mapping = compile_using_nvrtc( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 295, in compile_using_nvrtc E return _compile(source, options, cu_path, E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 279, in _compile E compiled_obj, mapping = prog.compile(options, log_stream) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 667, in compile E raise CompileException(log, self.src, self.name, options, E cupy.cuda.compiler.CompileException: /tmp/comgr-e3e812/input/CompileSource:203:27: error: redeclaration of '__hiprtc_16' with a different type: 'void (*const)(int, int, int, const signed char *, const signed char *, signed char *)' vs 'void (*const)(int, int, int, const bool *, const bool *, bool *)' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-e3e812/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-e3e812/input/CompileSource:204:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-e3e812/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-e3e812/input/CompileSource:205:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-e3e812/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-e3e812/input/CompileSource:206:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-e3e812/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-e3e812/input/CompileSource:207:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-e3e812/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-e3e812/input/CompileSource:208:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-e3e812/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-e3e812/input/CompileSource:209:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-e3e812/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-e3e812/input/CompileSource:210:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-e3e812/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-e3e812/input/CompileSource:211:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-e3e812/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-e3e812/input/CompileSource:212:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-e3e812/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E 10 errors generated when compiling for gfx908. E Error: Failed to compile opencl source (from CL or HIP source to LLVM IR). cupy/cuda/compiler.py:667: AssertionError ----------------------------- Captured stdout call ----------------------------- {'dtype_a': , 'dtype_b': } _ TestDot_param_2_{shape=((2, 3, 4), (3, 4, 2)), trans_a=False, trans_b=True}.test_dot _ self = < parameter: {'shape': ((2, 3, 4), (3, 4, 2)), 'trans_a': False, 'trans_b': True}> xp = dtype_a = , dtype_b = @testing.for_all_dtypes_combination(['dtype_a', 'dtype_b']) @testing.numpy_cupy_allclose() def test_dot(self, xp, dtype_a, dtype_b): shape_a, shape_b = self.shape if self.trans_a: a = testing.shaped_arange(shape_a[::-1], xp, dtype_a).T else: a = testing.shaped_arange(shape_a, xp, dtype_a) if self.trans_b: b = testing.shaped_arange(shape_b[::-1], xp, dtype_b).T else: b = testing.shaped_arange(shape_b, xp, dtype_b) > return xp.dot(a, b) tests/cupy_tests/linalg_tests/test_product.py:49: _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ cupy/linalg/_product.py:65: in dot return a.dot(b, out) cupy/_core/core.pyx:1607: in cupy._core.core.ndarray.dot return _linalg.dot(self, b, out) cupy/_core/_routines_linalg.pyx:427: in cupy._core._routines_linalg.dot return tensordot_core(a, b, out, n, m, k, ret_shape) cupy/_core/_routines_linalg.pyx:501: in cupy._core._routines_linalg.tensordot_core return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) cupy/_core/_routines_linalg.pyx:327: in cupy._core._routines_linalg._integral_tensordot_core kern = _tensordot_core_int_kernel(config, dtype) cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/_routines_linalg.pyx:302: in cupy._core._routines_linalg._tensordot_core_int_kernel ker = mod.get_function( cupy/_core/raw.pyx:470: in cupy._core.raw.RawModule.get_function mangled_name = self.module.mapping.get(name) cupy/_core/raw.pyx:394: in cupy._core.raw.RawModule.module.__get__ return self._module() cupy/_core/raw.pyx:402: in cupy._core.raw.RawModule._module mod = _get_raw_module( cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/raw.pyx:547: in cupy._core.raw._get_raw_module mod = cupy._core.core.compile_with_cache( cupy/_core/core.pyx:2015: in cupy._core.core.compile_with_cache cpdef function.Module compile_with_cache( cupy/_core/core.pyx:2078: in cupy._core.core.compile_with_cache return cuda.compile_with_cache( cupy/cuda/compiler.py:453: in compile_with_cache return _compile_with_cache_hip( cupy/cuda/compiler.py:853: in _compile_with_cache_hip binary, mapping = compile_using_nvrtc( cupy/cuda/compiler.py:295: in compile_using_nvrtc return _compile(source, options, cu_path, cupy/cuda/compiler.py:279: in _compile compiled_obj, mapping = prog.compile(options, log_stream) _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ self = options = ('--std=c++11', '-I/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/_core/include', '-I/opt/rocm/include', '-fcuda-flush-denormals-to-zero') log_stream = None def compile(self, options=(), log_stream=None): try: if self.name_expressions: for ker in self.name_expressions: nvrtc.addAddNameExpression(self.ptr, ker) nvrtc.compileProgram(self.ptr, options) mapping = None if self.name_expressions: mapping = {} for ker in self.name_expressions: mapping[ker] = nvrtc.getLoweredName(self.ptr, ker) if log_stream is not None: log_stream.write(nvrtc.getProgramLog(self.ptr)) # This is to ensure backwards compatibility with nvrtc if self.method == 'cubin': return nvrtc.getCUBIN(self.ptr), mapping elif self.method == 'ptx': return nvrtc.getPTX(self.ptr), mapping else: raise RuntimeError('Unknown NVRTC compile method') except nvrtc.NVRTCError: log = nvrtc.getProgramLog(self.ptr) > raise CompileException(log, self.src, self.name, options, 'nvrtc' if not runtime.is_hip else 'hiprtc') E AssertionError: Only cupy raises error E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 650, in compile E nvrtc.compileProgram(self.ptr, options) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 136, in cupy_backends.cuda.libs.nvrtc.compileProgram E cpdef compileProgram(intptr_t prog, options): E File "cupy_backends/cuda/libs/nvrtc.pyx", line 148, in cupy_backends.cuda.libs.nvrtc.compileProgram E check_status(status) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 67, in cupy_backends.cuda.libs.nvrtc.check_status E raise NVRTCError(status) E cupy_backends.cuda.libs.nvrtc.NVRTCError: HIPRTC_ERROR_COMPILATION (6) E E During handling of the above exception, another exception occurred: E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/testing/_helper.py", line 47, in _call_func E result = impl(*args, **kw) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/tests/cupy_tests/linalg_tests/test_product.py", line 49, in test_dot E return xp.dot(a, b) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/linalg/_product.py", line 65, in dot E return a.dot(b, out) E File "cupy/_core/core.pyx", line 1607, in cupy._core.core.ndarray.dot E return _linalg.dot(self, b, out) E File "cupy/_core/_routines_linalg.pyx", line 427, in cupy._core._routines_linalg.dot E return tensordot_core(a, b, out, n, m, k, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 501, in cupy._core._routines_linalg.tensordot_core E return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 327, in cupy._core._routines_linalg._integral_tensordot_core E kern = _tensordot_core_int_kernel(config, dtype) E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/_routines_linalg.pyx", line 302, in cupy._core._routines_linalg._tensordot_core_int_kernel E ker = mod.get_function( E File "cupy/_core/raw.pyx", line 470, in cupy._core.raw.RawModule.get_function E mangled_name = self.module.mapping.get(name) E File "cupy/_core/raw.pyx", line 394, in cupy._core.raw.RawModule.module.__get__ E return self._module() E File "cupy/_core/raw.pyx", line 402, in cupy._core.raw.RawModule._module E mod = _get_raw_module( E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/raw.pyx", line 547, in cupy._core.raw._get_raw_module E mod = cupy._core.core.compile_with_cache( E File "cupy/_core/core.pyx", line 2015, in cupy._core.core.compile_with_cache E cpdef function.Module compile_with_cache( E File "cupy/_core/core.pyx", line 2078, in cupy._core.core.compile_with_cache E return cuda.compile_with_cache( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 453, in compile_with_cache E return _compile_with_cache_hip( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 853, in _compile_with_cache_hip E binary, mapping = compile_using_nvrtc( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 295, in compile_using_nvrtc E return _compile(source, options, cu_path, E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 279, in _compile E compiled_obj, mapping = prog.compile(options, log_stream) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 667, in compile E raise CompileException(log, self.src, self.name, options, E cupy.cuda.compiler.CompileException: /tmp/comgr-917c3a/input/CompileSource:203:27: error: redeclaration of '__hiprtc_16' with a different type: 'void (*const)(int, int, int, const signed char *, const signed char *, signed char *)' vs 'void (*const)(int, int, int, const bool *, const bool *, bool *)' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-917c3a/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-917c3a/input/CompileSource:204:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-917c3a/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-917c3a/input/CompileSource:205:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-917c3a/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-917c3a/input/CompileSource:206:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-917c3a/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-917c3a/input/CompileSource:207:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-917c3a/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-917c3a/input/CompileSource:208:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-917c3a/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-917c3a/input/CompileSource:209:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-917c3a/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-917c3a/input/CompileSource:210:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-917c3a/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-917c3a/input/CompileSource:211:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-917c3a/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-917c3a/input/CompileSource:212:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-917c3a/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E 10 errors generated when compiling for gfx908. E Error: Failed to compile opencl source (from CL or HIP source to LLVM IR). cupy/cuda/compiler.py:667: AssertionError ----------------------------- Captured stdout call ----------------------------- {'dtype_a': , 'dtype_b': } _ TestDot_param_3_{shape=((2, 3, 4), (3, 4, 2)), trans_a=False, trans_b=False}.test_dot _ self = < parameter: {'shape': ((2, 3, 4), (3, 4, 2)), 'trans_a': False, 'trans_b': False}> xp = dtype_a = , dtype_b = @testing.for_all_dtypes_combination(['dtype_a', 'dtype_b']) @testing.numpy_cupy_allclose() def test_dot(self, xp, dtype_a, dtype_b): shape_a, shape_b = self.shape if self.trans_a: a = testing.shaped_arange(shape_a[::-1], xp, dtype_a).T else: a = testing.shaped_arange(shape_a, xp, dtype_a) if self.trans_b: b = testing.shaped_arange(shape_b[::-1], xp, dtype_b).T else: b = testing.shaped_arange(shape_b, xp, dtype_b) > return xp.dot(a, b) tests/cupy_tests/linalg_tests/test_product.py:49: _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ cupy/linalg/_product.py:65: in dot return a.dot(b, out) cupy/_core/core.pyx:1607: in cupy._core.core.ndarray.dot return _linalg.dot(self, b, out) cupy/_core/_routines_linalg.pyx:427: in cupy._core._routines_linalg.dot return tensordot_core(a, b, out, n, m, k, ret_shape) cupy/_core/_routines_linalg.pyx:501: in cupy._core._routines_linalg.tensordot_core return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) cupy/_core/_routines_linalg.pyx:327: in cupy._core._routines_linalg._integral_tensordot_core kern = _tensordot_core_int_kernel(config, dtype) cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/_routines_linalg.pyx:302: in cupy._core._routines_linalg._tensordot_core_int_kernel ker = mod.get_function( cupy/_core/raw.pyx:470: in cupy._core.raw.RawModule.get_function mangled_name = self.module.mapping.get(name) cupy/_core/raw.pyx:394: in cupy._core.raw.RawModule.module.__get__ return self._module() cupy/_core/raw.pyx:402: in cupy._core.raw.RawModule._module mod = _get_raw_module( cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/raw.pyx:547: in cupy._core.raw._get_raw_module mod = cupy._core.core.compile_with_cache( cupy/_core/core.pyx:2015: in cupy._core.core.compile_with_cache cpdef function.Module compile_with_cache( cupy/_core/core.pyx:2078: in cupy._core.core.compile_with_cache return cuda.compile_with_cache( cupy/cuda/compiler.py:453: in compile_with_cache return _compile_with_cache_hip( cupy/cuda/compiler.py:853: in _compile_with_cache_hip binary, mapping = compile_using_nvrtc( cupy/cuda/compiler.py:295: in compile_using_nvrtc return _compile(source, options, cu_path, cupy/cuda/compiler.py:279: in _compile compiled_obj, mapping = prog.compile(options, log_stream) _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ self = options = ('--std=c++11', '-I/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/_core/include', '-I/opt/rocm/include', '-fcuda-flush-denormals-to-zero') log_stream = None def compile(self, options=(), log_stream=None): try: if self.name_expressions: for ker in self.name_expressions: nvrtc.addAddNameExpression(self.ptr, ker) nvrtc.compileProgram(self.ptr, options) mapping = None if self.name_expressions: mapping = {} for ker in self.name_expressions: mapping[ker] = nvrtc.getLoweredName(self.ptr, ker) if log_stream is not None: log_stream.write(nvrtc.getProgramLog(self.ptr)) # This is to ensure backwards compatibility with nvrtc if self.method == 'cubin': return nvrtc.getCUBIN(self.ptr), mapping elif self.method == 'ptx': return nvrtc.getPTX(self.ptr), mapping else: raise RuntimeError('Unknown NVRTC compile method') except nvrtc.NVRTCError: log = nvrtc.getProgramLog(self.ptr) > raise CompileException(log, self.src, self.name, options, 'nvrtc' if not runtime.is_hip else 'hiprtc') E AssertionError: Only cupy raises error E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 650, in compile E nvrtc.compileProgram(self.ptr, options) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 136, in cupy_backends.cuda.libs.nvrtc.compileProgram E cpdef compileProgram(intptr_t prog, options): E File "cupy_backends/cuda/libs/nvrtc.pyx", line 148, in cupy_backends.cuda.libs.nvrtc.compileProgram E check_status(status) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 67, in cupy_backends.cuda.libs.nvrtc.check_status E raise NVRTCError(status) E cupy_backends.cuda.libs.nvrtc.NVRTCError: HIPRTC_ERROR_COMPILATION (6) E E During handling of the above exception, another exception occurred: E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/testing/_helper.py", line 47, in _call_func E result = impl(*args, **kw) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/tests/cupy_tests/linalg_tests/test_product.py", line 49, in test_dot E return xp.dot(a, b) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/linalg/_product.py", line 65, in dot E return a.dot(b, out) E File "cupy/_core/core.pyx", line 1607, in cupy._core.core.ndarray.dot E return _linalg.dot(self, b, out) E File "cupy/_core/_routines_linalg.pyx", line 427, in cupy._core._routines_linalg.dot E return tensordot_core(a, b, out, n, m, k, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 501, in cupy._core._routines_linalg.tensordot_core E return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 327, in cupy._core._routines_linalg._integral_tensordot_core E kern = _tensordot_core_int_kernel(config, dtype) E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/_routines_linalg.pyx", line 302, in cupy._core._routines_linalg._tensordot_core_int_kernel E ker = mod.get_function( E File "cupy/_core/raw.pyx", line 470, in cupy._core.raw.RawModule.get_function E mangled_name = self.module.mapping.get(name) E File "cupy/_core/raw.pyx", line 394, in cupy._core.raw.RawModule.module.__get__ E return self._module() E File "cupy/_core/raw.pyx", line 402, in cupy._core.raw.RawModule._module E mod = _get_raw_module( E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/raw.pyx", line 547, in cupy._core.raw._get_raw_module E mod = cupy._core.core.compile_with_cache( E File "cupy/_core/core.pyx", line 2015, in cupy._core.core.compile_with_cache E cpdef function.Module compile_with_cache( E File "cupy/_core/core.pyx", line 2078, in cupy._core.core.compile_with_cache E return cuda.compile_with_cache( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 453, in compile_with_cache E return _compile_with_cache_hip( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 853, in _compile_with_cache_hip E binary, mapping = compile_using_nvrtc( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 295, in compile_using_nvrtc E return _compile(source, options, cu_path, E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 279, in _compile E compiled_obj, mapping = prog.compile(options, log_stream) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 667, in compile E raise CompileException(log, self.src, self.name, options, E cupy.cuda.compiler.CompileException: /tmp/comgr-966fd7/input/CompileSource:203:27: error: redeclaration of '__hiprtc_16' with a different type: 'void (*const)(int, int, int, const signed char *, const signed char *, signed char *)' vs 'void (*const)(int, int, int, const bool *, const bool *, bool *)' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-966fd7/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-966fd7/input/CompileSource:204:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-966fd7/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-966fd7/input/CompileSource:205:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-966fd7/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-966fd7/input/CompileSource:206:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-966fd7/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-966fd7/input/CompileSource:207:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-966fd7/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-966fd7/input/CompileSource:208:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-966fd7/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-966fd7/input/CompileSource:209:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-966fd7/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-966fd7/input/CompileSource:210:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-966fd7/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-966fd7/input/CompileSource:211:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-966fd7/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-966fd7/input/CompileSource:212:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-966fd7/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E 10 errors generated when compiling for gfx908. E Error: Failed to compile opencl source (from CL or HIP source to LLVM IR). cupy/cuda/compiler.py:667: AssertionError ----------------------------- Captured stdout call ----------------------------- {'dtype_a': , 'dtype_b': } _ TestDot_param_8_{shape=((1, 1), (1, 2)), trans_a=True, trans_b=True}.test_dot _ self = < parameter: {'shape': ((1, 1), (1, 2)), 'trans_a': True, 'trans_b': True}> xp = dtype_a = , dtype_b = @testing.for_all_dtypes_combination(['dtype_a', 'dtype_b']) @testing.numpy_cupy_allclose() def test_dot(self, xp, dtype_a, dtype_b): shape_a, shape_b = self.shape if self.trans_a: a = testing.shaped_arange(shape_a[::-1], xp, dtype_a).T else: a = testing.shaped_arange(shape_a, xp, dtype_a) if self.trans_b: b = testing.shaped_arange(shape_b[::-1], xp, dtype_b).T else: b = testing.shaped_arange(shape_b, xp, dtype_b) > return xp.dot(a, b) tests/cupy_tests/linalg_tests/test_product.py:49: _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ cupy/linalg/_product.py:65: in dot return a.dot(b, out) cupy/_core/core.pyx:1607: in cupy._core.core.ndarray.dot return _linalg.dot(self, b, out) cupy/_core/_routines_linalg.pyx:427: in cupy._core._routines_linalg.dot return tensordot_core(a, b, out, n, m, k, ret_shape) cupy/_core/_routines_linalg.pyx:501: in cupy._core._routines_linalg.tensordot_core return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) cupy/_core/_routines_linalg.pyx:327: in cupy._core._routines_linalg._integral_tensordot_core kern = _tensordot_core_int_kernel(config, dtype) cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/_routines_linalg.pyx:302: in cupy._core._routines_linalg._tensordot_core_int_kernel ker = mod.get_function( cupy/_core/raw.pyx:470: in cupy._core.raw.RawModule.get_function mangled_name = self.module.mapping.get(name) cupy/_core/raw.pyx:394: in cupy._core.raw.RawModule.module.__get__ return self._module() cupy/_core/raw.pyx:402: in cupy._core.raw.RawModule._module mod = _get_raw_module( cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/raw.pyx:547: in cupy._core.raw._get_raw_module mod = cupy._core.core.compile_with_cache( cupy/_core/core.pyx:2015: in cupy._core.core.compile_with_cache cpdef function.Module compile_with_cache( cupy/_core/core.pyx:2078: in cupy._core.core.compile_with_cache return cuda.compile_with_cache( cupy/cuda/compiler.py:453: in compile_with_cache return _compile_with_cache_hip( cupy/cuda/compiler.py:853: in _compile_with_cache_hip binary, mapping = compile_using_nvrtc( cupy/cuda/compiler.py:295: in compile_using_nvrtc return _compile(source, options, cu_path, cupy/cuda/compiler.py:279: in _compile compiled_obj, mapping = prog.compile(options, log_stream) _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ self = options = ('--std=c++11', '-I/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/_core/include', '-I/opt/rocm/include', '-fcuda-flush-denormals-to-zero') log_stream = None def compile(self, options=(), log_stream=None): try: if self.name_expressions: for ker in self.name_expressions: nvrtc.addAddNameExpression(self.ptr, ker) nvrtc.compileProgram(self.ptr, options) mapping = None if self.name_expressions: mapping = {} for ker in self.name_expressions: mapping[ker] = nvrtc.getLoweredName(self.ptr, ker) if log_stream is not None: log_stream.write(nvrtc.getProgramLog(self.ptr)) # This is to ensure backwards compatibility with nvrtc if self.method == 'cubin': return nvrtc.getCUBIN(self.ptr), mapping elif self.method == 'ptx': return nvrtc.getPTX(self.ptr), mapping else: raise RuntimeError('Unknown NVRTC compile method') except nvrtc.NVRTCError: log = nvrtc.getProgramLog(self.ptr) > raise CompileException(log, self.src, self.name, options, 'nvrtc' if not runtime.is_hip else 'hiprtc') E AssertionError: Only cupy raises error E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 650, in compile E nvrtc.compileProgram(self.ptr, options) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 136, in cupy_backends.cuda.libs.nvrtc.compileProgram E cpdef compileProgram(intptr_t prog, options): E File "cupy_backends/cuda/libs/nvrtc.pyx", line 148, in cupy_backends.cuda.libs.nvrtc.compileProgram E check_status(status) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 67, in cupy_backends.cuda.libs.nvrtc.check_status E raise NVRTCError(status) E cupy_backends.cuda.libs.nvrtc.NVRTCError: HIPRTC_ERROR_COMPILATION (6) E E During handling of the above exception, another exception occurred: E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/testing/_helper.py", line 47, in _call_func E result = impl(*args, **kw) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/tests/cupy_tests/linalg_tests/test_product.py", line 49, in test_dot E return xp.dot(a, b) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/linalg/_product.py", line 65, in dot E return a.dot(b, out) E File "cupy/_core/core.pyx", line 1607, in cupy._core.core.ndarray.dot E return _linalg.dot(self, b, out) E File "cupy/_core/_routines_linalg.pyx", line 427, in cupy._core._routines_linalg.dot E return tensordot_core(a, b, out, n, m, k, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 501, in cupy._core._routines_linalg.tensordot_core E return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 327, in cupy._core._routines_linalg._integral_tensordot_core E kern = _tensordot_core_int_kernel(config, dtype) E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/_routines_linalg.pyx", line 302, in cupy._core._routines_linalg._tensordot_core_int_kernel E ker = mod.get_function( E File "cupy/_core/raw.pyx", line 470, in cupy._core.raw.RawModule.get_function E mangled_name = self.module.mapping.get(name) E File "cupy/_core/raw.pyx", line 394, in cupy._core.raw.RawModule.module.__get__ E return self._module() E File "cupy/_core/raw.pyx", line 402, in cupy._core.raw.RawModule._module E mod = _get_raw_module( E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/raw.pyx", line 547, in cupy._core.raw._get_raw_module E mod = cupy._core.core.compile_with_cache( E File "cupy/_core/core.pyx", line 2015, in cupy._core.core.compile_with_cache E cpdef function.Module compile_with_cache( E File "cupy/_core/core.pyx", line 2078, in cupy._core.core.compile_with_cache E return cuda.compile_with_cache( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 453, in compile_with_cache E return _compile_with_cache_hip( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 853, in _compile_with_cache_hip E binary, mapping = compile_using_nvrtc( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 295, in compile_using_nvrtc E return _compile(source, options, cu_path, E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 279, in _compile E compiled_obj, mapping = prog.compile(options, log_stream) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 667, in compile E raise CompileException(log, self.src, self.name, options, E cupy.cuda.compiler.CompileException: /tmp/comgr-dc3984/input/CompileSource:203:27: error: redeclaration of '__hiprtc_16' with a different type: 'void (*const)(int, int, int, const signed char *, const signed char *, signed char *)' vs 'void (*const)(int, int, int, const bool *, const bool *, bool *)' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-dc3984/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-dc3984/input/CompileSource:204:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-dc3984/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-dc3984/input/CompileSource:205:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-dc3984/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-dc3984/input/CompileSource:206:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-dc3984/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-dc3984/input/CompileSource:207:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-dc3984/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-dc3984/input/CompileSource:208:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-dc3984/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-dc3984/input/CompileSource:209:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-dc3984/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-dc3984/input/CompileSource:210:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-dc3984/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-dc3984/input/CompileSource:211:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-dc3984/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-dc3984/input/CompileSource:212:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-dc3984/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E 10 errors generated when compiling for gfx908. E Error: Failed to compile opencl source (from CL or HIP source to LLVM IR). cupy/cuda/compiler.py:667: AssertionError ----------------------------- Captured stdout call ----------------------------- {'dtype_a': , 'dtype_b': } _ TestDot_param_9_{shape=((1, 1), (1, 2)), trans_a=True, trans_b=False}.test_dot _ self = < parameter: {'shape': ((1, 1), (1, 2)), 'trans_a': True, 'trans_b': False}> xp = dtype_a = , dtype_b = @testing.for_all_dtypes_combination(['dtype_a', 'dtype_b']) @testing.numpy_cupy_allclose() def test_dot(self, xp, dtype_a, dtype_b): shape_a, shape_b = self.shape if self.trans_a: a = testing.shaped_arange(shape_a[::-1], xp, dtype_a).T else: a = testing.shaped_arange(shape_a, xp, dtype_a) if self.trans_b: b = testing.shaped_arange(shape_b[::-1], xp, dtype_b).T else: b = testing.shaped_arange(shape_b, xp, dtype_b) > return xp.dot(a, b) tests/cupy_tests/linalg_tests/test_product.py:49: _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ cupy/linalg/_product.py:65: in dot return a.dot(b, out) cupy/_core/core.pyx:1607: in cupy._core.core.ndarray.dot return _linalg.dot(self, b, out) cupy/_core/_routines_linalg.pyx:427: in cupy._core._routines_linalg.dot return tensordot_core(a, b, out, n, m, k, ret_shape) cupy/_core/_routines_linalg.pyx:501: in cupy._core._routines_linalg.tensordot_core return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) cupy/_core/_routines_linalg.pyx:327: in cupy._core._routines_linalg._integral_tensordot_core kern = _tensordot_core_int_kernel(config, dtype) cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/_routines_linalg.pyx:302: in cupy._core._routines_linalg._tensordot_core_int_kernel ker = mod.get_function( cupy/_core/raw.pyx:470: in cupy._core.raw.RawModule.get_function mangled_name = self.module.mapping.get(name) cupy/_core/raw.pyx:394: in cupy._core.raw.RawModule.module.__get__ return self._module() cupy/_core/raw.pyx:402: in cupy._core.raw.RawModule._module mod = _get_raw_module( cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/raw.pyx:547: in cupy._core.raw._get_raw_module mod = cupy._core.core.compile_with_cache( cupy/_core/core.pyx:2015: in cupy._core.core.compile_with_cache cpdef function.Module compile_with_cache( cupy/_core/core.pyx:2078: in cupy._core.core.compile_with_cache return cuda.compile_with_cache( cupy/cuda/compiler.py:453: in compile_with_cache return _compile_with_cache_hip( cupy/cuda/compiler.py:853: in _compile_with_cache_hip binary, mapping = compile_using_nvrtc( cupy/cuda/compiler.py:295: in compile_using_nvrtc return _compile(source, options, cu_path, cupy/cuda/compiler.py:279: in _compile compiled_obj, mapping = prog.compile(options, log_stream) _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ self = options = ('--std=c++11', '-I/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/_core/include', '-I/opt/rocm/include', '-fcuda-flush-denormals-to-zero') log_stream = None def compile(self, options=(), log_stream=None): try: if self.name_expressions: for ker in self.name_expressions: nvrtc.addAddNameExpression(self.ptr, ker) nvrtc.compileProgram(self.ptr, options) mapping = None if self.name_expressions: mapping = {} for ker in self.name_expressions: mapping[ker] = nvrtc.getLoweredName(self.ptr, ker) if log_stream is not None: log_stream.write(nvrtc.getProgramLog(self.ptr)) # This is to ensure backwards compatibility with nvrtc if self.method == 'cubin': return nvrtc.getCUBIN(self.ptr), mapping elif self.method == 'ptx': return nvrtc.getPTX(self.ptr), mapping else: raise RuntimeError('Unknown NVRTC compile method') except nvrtc.NVRTCError: log = nvrtc.getProgramLog(self.ptr) > raise CompileException(log, self.src, self.name, options, 'nvrtc' if not runtime.is_hip else 'hiprtc') E AssertionError: Only cupy raises error E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 650, in compile E nvrtc.compileProgram(self.ptr, options) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 136, in cupy_backends.cuda.libs.nvrtc.compileProgram E cpdef compileProgram(intptr_t prog, options): E File "cupy_backends/cuda/libs/nvrtc.pyx", line 148, in cupy_backends.cuda.libs.nvrtc.compileProgram E check_status(status) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 67, in cupy_backends.cuda.libs.nvrtc.check_status E raise NVRTCError(status) E cupy_backends.cuda.libs.nvrtc.NVRTCError: HIPRTC_ERROR_COMPILATION (6) E E During handling of the above exception, another exception occurred: E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/testing/_helper.py", line 47, in _call_func E result = impl(*args, **kw) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/tests/cupy_tests/linalg_tests/test_product.py", line 49, in test_dot E return xp.dot(a, b) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/linalg/_product.py", line 65, in dot E return a.dot(b, out) E File "cupy/_core/core.pyx", line 1607, in cupy._core.core.ndarray.dot E return _linalg.dot(self, b, out) E File "cupy/_core/_routines_linalg.pyx", line 427, in cupy._core._routines_linalg.dot E return tensordot_core(a, b, out, n, m, k, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 501, in cupy._core._routines_linalg.tensordot_core E return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 327, in cupy._core._routines_linalg._integral_tensordot_core E kern = _tensordot_core_int_kernel(config, dtype) E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/_routines_linalg.pyx", line 302, in cupy._core._routines_linalg._tensordot_core_int_kernel E ker = mod.get_function( E File "cupy/_core/raw.pyx", line 470, in cupy._core.raw.RawModule.get_function E mangled_name = self.module.mapping.get(name) E File "cupy/_core/raw.pyx", line 394, in cupy._core.raw.RawModule.module.__get__ E return self._module() E File "cupy/_core/raw.pyx", line 402, in cupy._core.raw.RawModule._module E mod = _get_raw_module( E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/raw.pyx", line 547, in cupy._core.raw._get_raw_module E mod = cupy._core.core.compile_with_cache( E File "cupy/_core/core.pyx", line 2015, in cupy._core.core.compile_with_cache E cpdef function.Module compile_with_cache( E File "cupy/_core/core.pyx", line 2078, in cupy._core.core.compile_with_cache E return cuda.compile_with_cache( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 453, in compile_with_cache E return _compile_with_cache_hip( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 853, in _compile_with_cache_hip E binary, mapping = compile_using_nvrtc( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 295, in compile_using_nvrtc E return _compile(source, options, cu_path, E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 279, in _compile E compiled_obj, mapping = prog.compile(options, log_stream) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 667, in compile E raise CompileException(log, self.src, self.name, options, E cupy.cuda.compiler.CompileException: /tmp/comgr-7293c6/input/CompileSource:203:27: error: redeclaration of '__hiprtc_16' with a different type: 'void (*const)(int, int, int, const signed char *, const signed char *, signed char *)' vs 'void (*const)(int, int, int, const bool *, const bool *, bool *)' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-7293c6/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-7293c6/input/CompileSource:204:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-7293c6/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-7293c6/input/CompileSource:205:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-7293c6/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-7293c6/input/CompileSource:206:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-7293c6/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-7293c6/input/CompileSource:207:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-7293c6/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-7293c6/input/CompileSource:208:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-7293c6/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-7293c6/input/CompileSource:209:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-7293c6/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-7293c6/input/CompileSource:210:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-7293c6/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-7293c6/input/CompileSource:211:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-7293c6/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-7293c6/input/CompileSource:212:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-7293c6/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E 10 errors generated when compiling for gfx908. E Error: Failed to compile opencl source (from CL or HIP source to LLVM IR). cupy/cuda/compiler.py:667: AssertionError ----------------------------- Captured stdout call ----------------------------- {'dtype_a': , 'dtype_b': } _ TestDot_param_10_{shape=((1, 1), (1, 2)), trans_a=False, trans_b=True}.test_dot _ self = < parameter: {'shape': ((1, 1), (1, 2)), 'trans_a': False, 'trans_b': True}> xp = dtype_a = , dtype_b = @testing.for_all_dtypes_combination(['dtype_a', 'dtype_b']) @testing.numpy_cupy_allclose() def test_dot(self, xp, dtype_a, dtype_b): shape_a, shape_b = self.shape if self.trans_a: a = testing.shaped_arange(shape_a[::-1], xp, dtype_a).T else: a = testing.shaped_arange(shape_a, xp, dtype_a) if self.trans_b: b = testing.shaped_arange(shape_b[::-1], xp, dtype_b).T else: b = testing.shaped_arange(shape_b, xp, dtype_b) > return xp.dot(a, b) tests/cupy_tests/linalg_tests/test_product.py:49: _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ cupy/linalg/_product.py:65: in dot return a.dot(b, out) cupy/_core/core.pyx:1607: in cupy._core.core.ndarray.dot return _linalg.dot(self, b, out) cupy/_core/_routines_linalg.pyx:427: in cupy._core._routines_linalg.dot return tensordot_core(a, b, out, n, m, k, ret_shape) cupy/_core/_routines_linalg.pyx:501: in cupy._core._routines_linalg.tensordot_core return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) cupy/_core/_routines_linalg.pyx:327: in cupy._core._routines_linalg._integral_tensordot_core kern = _tensordot_core_int_kernel(config, dtype) cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/_routines_linalg.pyx:302: in cupy._core._routines_linalg._tensordot_core_int_kernel ker = mod.get_function( cupy/_core/raw.pyx:470: in cupy._core.raw.RawModule.get_function mangled_name = self.module.mapping.get(name) cupy/_core/raw.pyx:394: in cupy._core.raw.RawModule.module.__get__ return self._module() cupy/_core/raw.pyx:402: in cupy._core.raw.RawModule._module mod = _get_raw_module( cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/raw.pyx:547: in cupy._core.raw._get_raw_module mod = cupy._core.core.compile_with_cache( cupy/_core/core.pyx:2015: in cupy._core.core.compile_with_cache cpdef function.Module compile_with_cache( cupy/_core/core.pyx:2078: in cupy._core.core.compile_with_cache return cuda.compile_with_cache( cupy/cuda/compiler.py:453: in compile_with_cache return _compile_with_cache_hip( cupy/cuda/compiler.py:853: in _compile_with_cache_hip binary, mapping = compile_using_nvrtc( cupy/cuda/compiler.py:295: in compile_using_nvrtc return _compile(source, options, cu_path, cupy/cuda/compiler.py:279: in _compile compiled_obj, mapping = prog.compile(options, log_stream) _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ self = options = ('--std=c++11', '-I/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/_core/include', '-I/opt/rocm/include', '-fcuda-flush-denormals-to-zero') log_stream = None def compile(self, options=(), log_stream=None): try: if self.name_expressions: for ker in self.name_expressions: nvrtc.addAddNameExpression(self.ptr, ker) nvrtc.compileProgram(self.ptr, options) mapping = None if self.name_expressions: mapping = {} for ker in self.name_expressions: mapping[ker] = nvrtc.getLoweredName(self.ptr, ker) if log_stream is not None: log_stream.write(nvrtc.getProgramLog(self.ptr)) # This is to ensure backwards compatibility with nvrtc if self.method == 'cubin': return nvrtc.getCUBIN(self.ptr), mapping elif self.method == 'ptx': return nvrtc.getPTX(self.ptr), mapping else: raise RuntimeError('Unknown NVRTC compile method') except nvrtc.NVRTCError: log = nvrtc.getProgramLog(self.ptr) > raise CompileException(log, self.src, self.name, options, 'nvrtc' if not runtime.is_hip else 'hiprtc') E AssertionError: Only cupy raises error E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 650, in compile E nvrtc.compileProgram(self.ptr, options) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 136, in cupy_backends.cuda.libs.nvrtc.compileProgram E cpdef compileProgram(intptr_t prog, options): E File "cupy_backends/cuda/libs/nvrtc.pyx", line 148, in cupy_backends.cuda.libs.nvrtc.compileProgram E check_status(status) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 67, in cupy_backends.cuda.libs.nvrtc.check_status E raise NVRTCError(status) E cupy_backends.cuda.libs.nvrtc.NVRTCError: HIPRTC_ERROR_COMPILATION (6) E E During handling of the above exception, another exception occurred: E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/testing/_helper.py", line 47, in _call_func E result = impl(*args, **kw) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/tests/cupy_tests/linalg_tests/test_product.py", line 49, in test_dot E return xp.dot(a, b) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/linalg/_product.py", line 65, in dot E return a.dot(b, out) E File "cupy/_core/core.pyx", line 1607, in cupy._core.core.ndarray.dot E return _linalg.dot(self, b, out) E File "cupy/_core/_routines_linalg.pyx", line 427, in cupy._core._routines_linalg.dot E return tensordot_core(a, b, out, n, m, k, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 501, in cupy._core._routines_linalg.tensordot_core E return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 327, in cupy._core._routines_linalg._integral_tensordot_core E kern = _tensordot_core_int_kernel(config, dtype) E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/_routines_linalg.pyx", line 302, in cupy._core._routines_linalg._tensordot_core_int_kernel E ker = mod.get_function( E File "cupy/_core/raw.pyx", line 470, in cupy._core.raw.RawModule.get_function E mangled_name = self.module.mapping.get(name) E File "cupy/_core/raw.pyx", line 394, in cupy._core.raw.RawModule.module.__get__ E return self._module() E File "cupy/_core/raw.pyx", line 402, in cupy._core.raw.RawModule._module E mod = _get_raw_module( E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/raw.pyx", line 547, in cupy._core.raw._get_raw_module E mod = cupy._core.core.compile_with_cache( E File "cupy/_core/core.pyx", line 2015, in cupy._core.core.compile_with_cache E cpdef function.Module compile_with_cache( E File "cupy/_core/core.pyx", line 2078, in cupy._core.core.compile_with_cache E return cuda.compile_with_cache( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 453, in compile_with_cache E return _compile_with_cache_hip( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 853, in _compile_with_cache_hip E binary, mapping = compile_using_nvrtc( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 295, in compile_using_nvrtc E return _compile(source, options, cu_path, E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 279, in _compile E compiled_obj, mapping = prog.compile(options, log_stream) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 667, in compile E raise CompileException(log, self.src, self.name, options, E cupy.cuda.compiler.CompileException: /tmp/comgr-95c61f/input/CompileSource:203:27: error: redeclaration of '__hiprtc_16' with a different type: 'void (*const)(int, int, int, const signed char *, const signed char *, signed char *)' vs 'void (*const)(int, int, int, const bool *, const bool *, bool *)' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-95c61f/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-95c61f/input/CompileSource:204:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-95c61f/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-95c61f/input/CompileSource:205:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-95c61f/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-95c61f/input/CompileSource:206:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-95c61f/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-95c61f/input/CompileSource:207:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-95c61f/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-95c61f/input/CompileSource:208:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-95c61f/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-95c61f/input/CompileSource:209:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-95c61f/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-95c61f/input/CompileSource:210:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-95c61f/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-95c61f/input/CompileSource:211:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-95c61f/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-95c61f/input/CompileSource:212:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-95c61f/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E 10 errors generated when compiling for gfx908. E Error: Failed to compile opencl source (from CL or HIP source to LLVM IR). cupy/cuda/compiler.py:667: AssertionError ----------------------------- Captured stdout call ----------------------------- {'dtype_a': , 'dtype_b': } _ TestDot_param_11_{shape=((1, 1), (1, 2)), trans_a=False, trans_b=False}.test_dot _ self = < parameter: {'shape': ((1, 1), (1, 2)), 'trans_a': False, 'trans_b': False}> xp = dtype_a = , dtype_b = @testing.for_all_dtypes_combination(['dtype_a', 'dtype_b']) @testing.numpy_cupy_allclose() def test_dot(self, xp, dtype_a, dtype_b): shape_a, shape_b = self.shape if self.trans_a: a = testing.shaped_arange(shape_a[::-1], xp, dtype_a).T else: a = testing.shaped_arange(shape_a, xp, dtype_a) if self.trans_b: b = testing.shaped_arange(shape_b[::-1], xp, dtype_b).T else: b = testing.shaped_arange(shape_b, xp, dtype_b) > return xp.dot(a, b) tests/cupy_tests/linalg_tests/test_product.py:49: _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ cupy/linalg/_product.py:65: in dot return a.dot(b, out) cupy/_core/core.pyx:1607: in cupy._core.core.ndarray.dot return _linalg.dot(self, b, out) cupy/_core/_routines_linalg.pyx:427: in cupy._core._routines_linalg.dot return tensordot_core(a, b, out, n, m, k, ret_shape) cupy/_core/_routines_linalg.pyx:501: in cupy._core._routines_linalg.tensordot_core return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) cupy/_core/_routines_linalg.pyx:327: in cupy._core._routines_linalg._integral_tensordot_core kern = _tensordot_core_int_kernel(config, dtype) cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/_routines_linalg.pyx:302: in cupy._core._routines_linalg._tensordot_core_int_kernel ker = mod.get_function( cupy/_core/raw.pyx:470: in cupy._core.raw.RawModule.get_function mangled_name = self.module.mapping.get(name) cupy/_core/raw.pyx:394: in cupy._core.raw.RawModule.module.__get__ return self._module() cupy/_core/raw.pyx:402: in cupy._core.raw.RawModule._module mod = _get_raw_module( cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/raw.pyx:547: in cupy._core.raw._get_raw_module mod = cupy._core.core.compile_with_cache( cupy/_core/core.pyx:2015: in cupy._core.core.compile_with_cache cpdef function.Module compile_with_cache( cupy/_core/core.pyx:2078: in cupy._core.core.compile_with_cache return cuda.compile_with_cache( cupy/cuda/compiler.py:453: in compile_with_cache return _compile_with_cache_hip( cupy/cuda/compiler.py:853: in _compile_with_cache_hip binary, mapping = compile_using_nvrtc( cupy/cuda/compiler.py:295: in compile_using_nvrtc return _compile(source, options, cu_path, cupy/cuda/compiler.py:279: in _compile compiled_obj, mapping = prog.compile(options, log_stream) _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ self = options = ('--std=c++11', '-I/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/_core/include', '-I/opt/rocm/include', '-fcuda-flush-denormals-to-zero') log_stream = None def compile(self, options=(), log_stream=None): try: if self.name_expressions: for ker in self.name_expressions: nvrtc.addAddNameExpression(self.ptr, ker) nvrtc.compileProgram(self.ptr, options) mapping = None if self.name_expressions: mapping = {} for ker in self.name_expressions: mapping[ker] = nvrtc.getLoweredName(self.ptr, ker) if log_stream is not None: log_stream.write(nvrtc.getProgramLog(self.ptr)) # This is to ensure backwards compatibility with nvrtc if self.method == 'cubin': return nvrtc.getCUBIN(self.ptr), mapping elif self.method == 'ptx': return nvrtc.getPTX(self.ptr), mapping else: raise RuntimeError('Unknown NVRTC compile method') except nvrtc.NVRTCError: log = nvrtc.getProgramLog(self.ptr) > raise CompileException(log, self.src, self.name, options, 'nvrtc' if not runtime.is_hip else 'hiprtc') E AssertionError: Only cupy raises error E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 650, in compile E nvrtc.compileProgram(self.ptr, options) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 136, in cupy_backends.cuda.libs.nvrtc.compileProgram E cpdef compileProgram(intptr_t prog, options): E File "cupy_backends/cuda/libs/nvrtc.pyx", line 148, in cupy_backends.cuda.libs.nvrtc.compileProgram E check_status(status) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 67, in cupy_backends.cuda.libs.nvrtc.check_status E raise NVRTCError(status) E cupy_backends.cuda.libs.nvrtc.NVRTCError: HIPRTC_ERROR_COMPILATION (6) E E During handling of the above exception, another exception occurred: E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/testing/_helper.py", line 47, in _call_func E result = impl(*args, **kw) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/tests/cupy_tests/linalg_tests/test_product.py", line 49, in test_dot E return xp.dot(a, b) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/linalg/_product.py", line 65, in dot E return a.dot(b, out) E File "cupy/_core/core.pyx", line 1607, in cupy._core.core.ndarray.dot E return _linalg.dot(self, b, out) E File "cupy/_core/_routines_linalg.pyx", line 427, in cupy._core._routines_linalg.dot E return tensordot_core(a, b, out, n, m, k, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 501, in cupy._core._routines_linalg.tensordot_core E return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 327, in cupy._core._routines_linalg._integral_tensordot_core E kern = _tensordot_core_int_kernel(config, dtype) E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/_routines_linalg.pyx", line 302, in cupy._core._routines_linalg._tensordot_core_int_kernel E ker = mod.get_function( E File "cupy/_core/raw.pyx", line 470, in cupy._core.raw.RawModule.get_function E mangled_name = self.module.mapping.get(name) E File "cupy/_core/raw.pyx", line 394, in cupy._core.raw.RawModule.module.__get__ E return self._module() E File "cupy/_core/raw.pyx", line 402, in cupy._core.raw.RawModule._module E mod = _get_raw_module( E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/raw.pyx", line 547, in cupy._core.raw._get_raw_module E mod = cupy._core.core.compile_with_cache( E File "cupy/_core/core.pyx", line 2015, in cupy._core.core.compile_with_cache E cpdef function.Module compile_with_cache( E File "cupy/_core/core.pyx", line 2078, in cupy._core.core.compile_with_cache E return cuda.compile_with_cache( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 453, in compile_with_cache E return _compile_with_cache_hip( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 853, in _compile_with_cache_hip E binary, mapping = compile_using_nvrtc( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 295, in compile_using_nvrtc E return _compile(source, options, cu_path, E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 279, in _compile E compiled_obj, mapping = prog.compile(options, log_stream) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 667, in compile E raise CompileException(log, self.src, self.name, options, E cupy.cuda.compiler.CompileException: /tmp/comgr-c19433/input/CompileSource:203:27: error: redeclaration of '__hiprtc_16' with a different type: 'void (*const)(int, int, int, const signed char *, const signed char *, signed char *)' vs 'void (*const)(int, int, int, const bool *, const bool *, bool *)' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-c19433/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-c19433/input/CompileSource:204:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-c19433/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-c19433/input/CompileSource:205:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-c19433/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-c19433/input/CompileSource:206:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-c19433/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-c19433/input/CompileSource:207:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-c19433/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-c19433/input/CompileSource:208:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-c19433/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-c19433/input/CompileSource:209:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-c19433/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-c19433/input/CompileSource:210:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-c19433/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-c19433/input/CompileSource:211:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-c19433/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-c19433/input/CompileSource:212:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-c19433/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E 10 errors generated when compiling for gfx908. E Error: Failed to compile opencl source (from CL or HIP source to LLVM IR). cupy/cuda/compiler.py:667: AssertionError ----------------------------- Captured stdout call ----------------------------- {'dtype_a': , 'dtype_b': } _ TestDot_param_16_{shape=((2, 1), (1, 1)), trans_a=True, trans_b=True}.test_dot _ self = < parameter: {'shape': ((2, 1), (1, 1)), 'trans_a': True, 'trans_b': True}> xp = dtype_a = , dtype_b = @testing.for_all_dtypes_combination(['dtype_a', 'dtype_b']) @testing.numpy_cupy_allclose() def test_dot(self, xp, dtype_a, dtype_b): shape_a, shape_b = self.shape if self.trans_a: a = testing.shaped_arange(shape_a[::-1], xp, dtype_a).T else: a = testing.shaped_arange(shape_a, xp, dtype_a) if self.trans_b: b = testing.shaped_arange(shape_b[::-1], xp, dtype_b).T else: b = testing.shaped_arange(shape_b, xp, dtype_b) > return xp.dot(a, b) tests/cupy_tests/linalg_tests/test_product.py:49: _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ cupy/linalg/_product.py:65: in dot return a.dot(b, out) cupy/_core/core.pyx:1607: in cupy._core.core.ndarray.dot return _linalg.dot(self, b, out) cupy/_core/_routines_linalg.pyx:427: in cupy._core._routines_linalg.dot return tensordot_core(a, b, out, n, m, k, ret_shape) cupy/_core/_routines_linalg.pyx:501: in cupy._core._routines_linalg.tensordot_core return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) cupy/_core/_routines_linalg.pyx:327: in cupy._core._routines_linalg._integral_tensordot_core kern = _tensordot_core_int_kernel(config, dtype) cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/_routines_linalg.pyx:302: in cupy._core._routines_linalg._tensordot_core_int_kernel ker = mod.get_function( cupy/_core/raw.pyx:470: in cupy._core.raw.RawModule.get_function mangled_name = self.module.mapping.get(name) cupy/_core/raw.pyx:394: in cupy._core.raw.RawModule.module.__get__ return self._module() cupy/_core/raw.pyx:402: in cupy._core.raw.RawModule._module mod = _get_raw_module( cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/raw.pyx:547: in cupy._core.raw._get_raw_module mod = cupy._core.core.compile_with_cache( cupy/_core/core.pyx:2015: in cupy._core.core.compile_with_cache cpdef function.Module compile_with_cache( cupy/_core/core.pyx:2078: in cupy._core.core.compile_with_cache return cuda.compile_with_cache( cupy/cuda/compiler.py:453: in compile_with_cache return _compile_with_cache_hip( cupy/cuda/compiler.py:853: in _compile_with_cache_hip binary, mapping = compile_using_nvrtc( cupy/cuda/compiler.py:295: in compile_using_nvrtc return _compile(source, options, cu_path, cupy/cuda/compiler.py:279: in _compile compiled_obj, mapping = prog.compile(options, log_stream) _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ self = options = ('--std=c++11', '-I/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/_core/include', '-I/opt/rocm/include', '-fcuda-flush-denormals-to-zero') log_stream = None def compile(self, options=(), log_stream=None): try: if self.name_expressions: for ker in self.name_expressions: nvrtc.addAddNameExpression(self.ptr, ker) nvrtc.compileProgram(self.ptr, options) mapping = None if self.name_expressions: mapping = {} for ker in self.name_expressions: mapping[ker] = nvrtc.getLoweredName(self.ptr, ker) if log_stream is not None: log_stream.write(nvrtc.getProgramLog(self.ptr)) # This is to ensure backwards compatibility with nvrtc if self.method == 'cubin': return nvrtc.getCUBIN(self.ptr), mapping elif self.method == 'ptx': return nvrtc.getPTX(self.ptr), mapping else: raise RuntimeError('Unknown NVRTC compile method') except nvrtc.NVRTCError: log = nvrtc.getProgramLog(self.ptr) > raise CompileException(log, self.src, self.name, options, 'nvrtc' if not runtime.is_hip else 'hiprtc') E AssertionError: Only cupy raises error E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 650, in compile E nvrtc.compileProgram(self.ptr, options) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 136, in cupy_backends.cuda.libs.nvrtc.compileProgram E cpdef compileProgram(intptr_t prog, options): E File "cupy_backends/cuda/libs/nvrtc.pyx", line 148, in cupy_backends.cuda.libs.nvrtc.compileProgram E check_status(status) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 67, in cupy_backends.cuda.libs.nvrtc.check_status E raise NVRTCError(status) E cupy_backends.cuda.libs.nvrtc.NVRTCError: HIPRTC_ERROR_COMPILATION (6) E E During handling of the above exception, another exception occurred: E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/testing/_helper.py", line 47, in _call_func E result = impl(*args, **kw) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/tests/cupy_tests/linalg_tests/test_product.py", line 49, in test_dot E return xp.dot(a, b) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/linalg/_product.py", line 65, in dot E return a.dot(b, out) E File "cupy/_core/core.pyx", line 1607, in cupy._core.core.ndarray.dot E return _linalg.dot(self, b, out) E File "cupy/_core/_routines_linalg.pyx", line 427, in cupy._core._routines_linalg.dot E return tensordot_core(a, b, out, n, m, k, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 501, in cupy._core._routines_linalg.tensordot_core E return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 327, in cupy._core._routines_linalg._integral_tensordot_core E kern = _tensordot_core_int_kernel(config, dtype) E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/_routines_linalg.pyx", line 302, in cupy._core._routines_linalg._tensordot_core_int_kernel E ker = mod.get_function( E File "cupy/_core/raw.pyx", line 470, in cupy._core.raw.RawModule.get_function E mangled_name = self.module.mapping.get(name) E File "cupy/_core/raw.pyx", line 394, in cupy._core.raw.RawModule.module.__get__ E return self._module() E File "cupy/_core/raw.pyx", line 402, in cupy._core.raw.RawModule._module E mod = _get_raw_module( E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/raw.pyx", line 547, in cupy._core.raw._get_raw_module E mod = cupy._core.core.compile_with_cache( E File "cupy/_core/core.pyx", line 2015, in cupy._core.core.compile_with_cache E cpdef function.Module compile_with_cache( E File "cupy/_core/core.pyx", line 2078, in cupy._core.core.compile_with_cache E return cuda.compile_with_cache( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 453, in compile_with_cache E return _compile_with_cache_hip( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 853, in _compile_with_cache_hip E binary, mapping = compile_using_nvrtc( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 295, in compile_using_nvrtc E return _compile(source, options, cu_path, E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 279, in _compile E compiled_obj, mapping = prog.compile(options, log_stream) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 667, in compile E raise CompileException(log, self.src, self.name, options, E cupy.cuda.compiler.CompileException: /tmp/comgr-312c6e/input/CompileSource:203:27: error: redeclaration of '__hiprtc_16' with a different type: 'void (*const)(int, int, int, const signed char *, const signed char *, signed char *)' vs 'void (*const)(int, int, int, const bool *, const bool *, bool *)' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-312c6e/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-312c6e/input/CompileSource:204:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-312c6e/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-312c6e/input/CompileSource:205:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-312c6e/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-312c6e/input/CompileSource:206:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-312c6e/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-312c6e/input/CompileSource:207:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-312c6e/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-312c6e/input/CompileSource:208:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-312c6e/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-312c6e/input/CompileSource:209:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-312c6e/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-312c6e/input/CompileSource:210:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-312c6e/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-312c6e/input/CompileSource:211:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-312c6e/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-312c6e/input/CompileSource:212:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-312c6e/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E 10 errors generated when compiling for gfx908. E Error: Failed to compile opencl source (from CL or HIP source to LLVM IR). cupy/cuda/compiler.py:667: AssertionError ----------------------------- Captured stdout call ----------------------------- {'dtype_a': , 'dtype_b': } _ TestDot_param_17_{shape=((2, 1), (1, 1)), trans_a=True, trans_b=False}.test_dot _ self = < parameter: {'shape': ((2, 1), (1, 1)), 'trans_a': True, 'trans_b': False}> xp = dtype_a = , dtype_b = @testing.for_all_dtypes_combination(['dtype_a', 'dtype_b']) @testing.numpy_cupy_allclose() def test_dot(self, xp, dtype_a, dtype_b): shape_a, shape_b = self.shape if self.trans_a: a = testing.shaped_arange(shape_a[::-1], xp, dtype_a).T else: a = testing.shaped_arange(shape_a, xp, dtype_a) if self.trans_b: b = testing.shaped_arange(shape_b[::-1], xp, dtype_b).T else: b = testing.shaped_arange(shape_b, xp, dtype_b) > return xp.dot(a, b) tests/cupy_tests/linalg_tests/test_product.py:49: _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ cupy/linalg/_product.py:65: in dot return a.dot(b, out) cupy/_core/core.pyx:1607: in cupy._core.core.ndarray.dot return _linalg.dot(self, b, out) cupy/_core/_routines_linalg.pyx:427: in cupy._core._routines_linalg.dot return tensordot_core(a, b, out, n, m, k, ret_shape) cupy/_core/_routines_linalg.pyx:501: in cupy._core._routines_linalg.tensordot_core return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) cupy/_core/_routines_linalg.pyx:327: in cupy._core._routines_linalg._integral_tensordot_core kern = _tensordot_core_int_kernel(config, dtype) cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/_routines_linalg.pyx:302: in cupy._core._routines_linalg._tensordot_core_int_kernel ker = mod.get_function( cupy/_core/raw.pyx:470: in cupy._core.raw.RawModule.get_function mangled_name = self.module.mapping.get(name) cupy/_core/raw.pyx:394: in cupy._core.raw.RawModule.module.__get__ return self._module() cupy/_core/raw.pyx:402: in cupy._core.raw.RawModule._module mod = _get_raw_module( cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/raw.pyx:547: in cupy._core.raw._get_raw_module mod = cupy._core.core.compile_with_cache( cupy/_core/core.pyx:2015: in cupy._core.core.compile_with_cache cpdef function.Module compile_with_cache( cupy/_core/core.pyx:2078: in cupy._core.core.compile_with_cache return cuda.compile_with_cache( cupy/cuda/compiler.py:453: in compile_with_cache return _compile_with_cache_hip( cupy/cuda/compiler.py:853: in _compile_with_cache_hip binary, mapping = compile_using_nvrtc( cupy/cuda/compiler.py:295: in compile_using_nvrtc return _compile(source, options, cu_path, cupy/cuda/compiler.py:279: in _compile compiled_obj, mapping = prog.compile(options, log_stream) _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ self = options = ('--std=c++11', '-I/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/_core/include', '-I/opt/rocm/include', '-fcuda-flush-denormals-to-zero') log_stream = None def compile(self, options=(), log_stream=None): try: if self.name_expressions: for ker in self.name_expressions: nvrtc.addAddNameExpression(self.ptr, ker) nvrtc.compileProgram(self.ptr, options) mapping = None if self.name_expressions: mapping = {} for ker in self.name_expressions: mapping[ker] = nvrtc.getLoweredName(self.ptr, ker) if log_stream is not None: log_stream.write(nvrtc.getProgramLog(self.ptr)) # This is to ensure backwards compatibility with nvrtc if self.method == 'cubin': return nvrtc.getCUBIN(self.ptr), mapping elif self.method == 'ptx': return nvrtc.getPTX(self.ptr), mapping else: raise RuntimeError('Unknown NVRTC compile method') except nvrtc.NVRTCError: log = nvrtc.getProgramLog(self.ptr) > raise CompileException(log, self.src, self.name, options, 'nvrtc' if not runtime.is_hip else 'hiprtc') E AssertionError: Only cupy raises error E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 650, in compile E nvrtc.compileProgram(self.ptr, options) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 136, in cupy_backends.cuda.libs.nvrtc.compileProgram E cpdef compileProgram(intptr_t prog, options): E File "cupy_backends/cuda/libs/nvrtc.pyx", line 148, in cupy_backends.cuda.libs.nvrtc.compileProgram E check_status(status) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 67, in cupy_backends.cuda.libs.nvrtc.check_status E raise NVRTCError(status) E cupy_backends.cuda.libs.nvrtc.NVRTCError: HIPRTC_ERROR_COMPILATION (6) E E During handling of the above exception, another exception occurred: E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/testing/_helper.py", line 47, in _call_func E result = impl(*args, **kw) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/tests/cupy_tests/linalg_tests/test_product.py", line 49, in test_dot E return xp.dot(a, b) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/linalg/_product.py", line 65, in dot E return a.dot(b, out) E File "cupy/_core/core.pyx", line 1607, in cupy._core.core.ndarray.dot E return _linalg.dot(self, b, out) E File "cupy/_core/_routines_linalg.pyx", line 427, in cupy._core._routines_linalg.dot E return tensordot_core(a, b, out, n, m, k, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 501, in cupy._core._routines_linalg.tensordot_core E return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 327, in cupy._core._routines_linalg._integral_tensordot_core E kern = _tensordot_core_int_kernel(config, dtype) E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/_routines_linalg.pyx", line 302, in cupy._core._routines_linalg._tensordot_core_int_kernel E ker = mod.get_function( E File "cupy/_core/raw.pyx", line 470, in cupy._core.raw.RawModule.get_function E mangled_name = self.module.mapping.get(name) E File "cupy/_core/raw.pyx", line 394, in cupy._core.raw.RawModule.module.__get__ E return self._module() E File "cupy/_core/raw.pyx", line 402, in cupy._core.raw.RawModule._module E mod = _get_raw_module( E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/raw.pyx", line 547, in cupy._core.raw._get_raw_module E mod = cupy._core.core.compile_with_cache( E File "cupy/_core/core.pyx", line 2015, in cupy._core.core.compile_with_cache E cpdef function.Module compile_with_cache( E File "cupy/_core/core.pyx", line 2078, in cupy._core.core.compile_with_cache E return cuda.compile_with_cache( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 453, in compile_with_cache E return _compile_with_cache_hip( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 853, in _compile_with_cache_hip E binary, mapping = compile_using_nvrtc( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 295, in compile_using_nvrtc E return _compile(source, options, cu_path, E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 279, in _compile E compiled_obj, mapping = prog.compile(options, log_stream) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 667, in compile E raise CompileException(log, self.src, self.name, options, E cupy.cuda.compiler.CompileException: /tmp/comgr-c908aa/input/CompileSource:203:27: error: redeclaration of '__hiprtc_16' with a different type: 'void (*const)(int, int, int, const signed char *, const signed char *, signed char *)' vs 'void (*const)(int, int, int, const bool *, const bool *, bool *)' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-c908aa/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-c908aa/input/CompileSource:204:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-c908aa/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-c908aa/input/CompileSource:205:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-c908aa/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-c908aa/input/CompileSource:206:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-c908aa/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-c908aa/input/CompileSource:207:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-c908aa/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-c908aa/input/CompileSource:208:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-c908aa/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-c908aa/input/CompileSource:209:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-c908aa/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-c908aa/input/CompileSource:210:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-c908aa/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-c908aa/input/CompileSource:211:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-c908aa/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-c908aa/input/CompileSource:212:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-c908aa/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E 10 errors generated when compiling for gfx908. E Error: Failed to compile opencl source (from CL or HIP source to LLVM IR). cupy/cuda/compiler.py:667: AssertionError ----------------------------- Captured stdout call ----------------------------- {'dtype_a': , 'dtype_b': } _ TestDot_param_18_{shape=((2, 1), (1, 1)), trans_a=False, trans_b=True}.test_dot _ self = < parameter: {'shape': ((2, 1), (1, 1)), 'trans_a': False, 'trans_b': True}> xp = dtype_a = , dtype_b = @testing.for_all_dtypes_combination(['dtype_a', 'dtype_b']) @testing.numpy_cupy_allclose() def test_dot(self, xp, dtype_a, dtype_b): shape_a, shape_b = self.shape if self.trans_a: a = testing.shaped_arange(shape_a[::-1], xp, dtype_a).T else: a = testing.shaped_arange(shape_a, xp, dtype_a) if self.trans_b: b = testing.shaped_arange(shape_b[::-1], xp, dtype_b).T else: b = testing.shaped_arange(shape_b, xp, dtype_b) > return xp.dot(a, b) tests/cupy_tests/linalg_tests/test_product.py:49: _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ cupy/linalg/_product.py:65: in dot return a.dot(b, out) cupy/_core/core.pyx:1607: in cupy._core.core.ndarray.dot return _linalg.dot(self, b, out) cupy/_core/_routines_linalg.pyx:427: in cupy._core._routines_linalg.dot return tensordot_core(a, b, out, n, m, k, ret_shape) cupy/_core/_routines_linalg.pyx:501: in cupy._core._routines_linalg.tensordot_core return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) cupy/_core/_routines_linalg.pyx:327: in cupy._core._routines_linalg._integral_tensordot_core kern = _tensordot_core_int_kernel(config, dtype) cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/_routines_linalg.pyx:302: in cupy._core._routines_linalg._tensordot_core_int_kernel ker = mod.get_function( cupy/_core/raw.pyx:470: in cupy._core.raw.RawModule.get_function mangled_name = self.module.mapping.get(name) cupy/_core/raw.pyx:394: in cupy._core.raw.RawModule.module.__get__ return self._module() cupy/_core/raw.pyx:402: in cupy._core.raw.RawModule._module mod = _get_raw_module( cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/raw.pyx:547: in cupy._core.raw._get_raw_module mod = cupy._core.core.compile_with_cache( cupy/_core/core.pyx:2015: in cupy._core.core.compile_with_cache cpdef function.Module compile_with_cache( cupy/_core/core.pyx:2078: in cupy._core.core.compile_with_cache return cuda.compile_with_cache( cupy/cuda/compiler.py:453: in compile_with_cache return _compile_with_cache_hip( cupy/cuda/compiler.py:853: in _compile_with_cache_hip binary, mapping = compile_using_nvrtc( cupy/cuda/compiler.py:295: in compile_using_nvrtc return _compile(source, options, cu_path, cupy/cuda/compiler.py:279: in _compile compiled_obj, mapping = prog.compile(options, log_stream) _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ self = options = ('--std=c++11', '-I/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/_core/include', '-I/opt/rocm/include', '-fcuda-flush-denormals-to-zero') log_stream = None def compile(self, options=(), log_stream=None): try: if self.name_expressions: for ker in self.name_expressions: nvrtc.addAddNameExpression(self.ptr, ker) nvrtc.compileProgram(self.ptr, options) mapping = None if self.name_expressions: mapping = {} for ker in self.name_expressions: mapping[ker] = nvrtc.getLoweredName(self.ptr, ker) if log_stream is not None: log_stream.write(nvrtc.getProgramLog(self.ptr)) # This is to ensure backwards compatibility with nvrtc if self.method == 'cubin': return nvrtc.getCUBIN(self.ptr), mapping elif self.method == 'ptx': return nvrtc.getPTX(self.ptr), mapping else: raise RuntimeError('Unknown NVRTC compile method') except nvrtc.NVRTCError: log = nvrtc.getProgramLog(self.ptr) > raise CompileException(log, self.src, self.name, options, 'nvrtc' if not runtime.is_hip else 'hiprtc') E AssertionError: Only cupy raises error E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 650, in compile E nvrtc.compileProgram(self.ptr, options) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 136, in cupy_backends.cuda.libs.nvrtc.compileProgram E cpdef compileProgram(intptr_t prog, options): E File "cupy_backends/cuda/libs/nvrtc.pyx", line 148, in cupy_backends.cuda.libs.nvrtc.compileProgram E check_status(status) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 67, in cupy_backends.cuda.libs.nvrtc.check_status E raise NVRTCError(status) E cupy_backends.cuda.libs.nvrtc.NVRTCError: HIPRTC_ERROR_COMPILATION (6) E E During handling of the above exception, another exception occurred: E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/testing/_helper.py", line 47, in _call_func E result = impl(*args, **kw) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/tests/cupy_tests/linalg_tests/test_product.py", line 49, in test_dot E return xp.dot(a, b) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/linalg/_product.py", line 65, in dot E return a.dot(b, out) E File "cupy/_core/core.pyx", line 1607, in cupy._core.core.ndarray.dot E return _linalg.dot(self, b, out) E File "cupy/_core/_routines_linalg.pyx", line 427, in cupy._core._routines_linalg.dot E return tensordot_core(a, b, out, n, m, k, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 501, in cupy._core._routines_linalg.tensordot_core E return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 327, in cupy._core._routines_linalg._integral_tensordot_core E kern = _tensordot_core_int_kernel(config, dtype) E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/_routines_linalg.pyx", line 302, in cupy._core._routines_linalg._tensordot_core_int_kernel E ker = mod.get_function( E File "cupy/_core/raw.pyx", line 470, in cupy._core.raw.RawModule.get_function E mangled_name = self.module.mapping.get(name) E File "cupy/_core/raw.pyx", line 394, in cupy._core.raw.RawModule.module.__get__ E return self._module() E File "cupy/_core/raw.pyx", line 402, in cupy._core.raw.RawModule._module E mod = _get_raw_module( E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/raw.pyx", line 547, in cupy._core.raw._get_raw_module E mod = cupy._core.core.compile_with_cache( E File "cupy/_core/core.pyx", line 2015, in cupy._core.core.compile_with_cache E cpdef function.Module compile_with_cache( E File "cupy/_core/core.pyx", line 2078, in cupy._core.core.compile_with_cache E return cuda.compile_with_cache( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 453, in compile_with_cache E return _compile_with_cache_hip( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 853, in _compile_with_cache_hip E binary, mapping = compile_using_nvrtc( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 295, in compile_using_nvrtc E return _compile(source, options, cu_path, E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 279, in _compile E compiled_obj, mapping = prog.compile(options, log_stream) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 667, in compile E raise CompileException(log, self.src, self.name, options, E cupy.cuda.compiler.CompileException: /tmp/comgr-232545/input/CompileSource:203:27: error: redeclaration of '__hiprtc_16' with a different type: 'void (*const)(int, int, int, const signed char *, const signed char *, signed char *)' vs 'void (*const)(int, int, int, const bool *, const bool *, bool *)' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-232545/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-232545/input/CompileSource:204:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-232545/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-232545/input/CompileSource:205:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-232545/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-232545/input/CompileSource:206:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-232545/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-232545/input/CompileSource:207:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-232545/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-232545/input/CompileSource:208:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-232545/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-232545/input/CompileSource:209:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-232545/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-232545/input/CompileSource:210:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-232545/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-232545/input/CompileSource:211:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-232545/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-232545/input/CompileSource:212:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-232545/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E 10 errors generated when compiling for gfx908. E Error: Failed to compile opencl source (from CL or HIP source to LLVM IR). cupy/cuda/compiler.py:667: AssertionError ----------------------------- Captured stdout call ----------------------------- {'dtype_a': , 'dtype_b': } _ TestDot_param_19_{shape=((2, 1), (1, 1)), trans_a=False, trans_b=False}.test_dot _ self = < parameter: {'shape': ((2, 1), (1, 1)), 'trans_a': False, 'trans_b': False}> xp = dtype_a = , dtype_b = @testing.for_all_dtypes_combination(['dtype_a', 'dtype_b']) @testing.numpy_cupy_allclose() def test_dot(self, xp, dtype_a, dtype_b): shape_a, shape_b = self.shape if self.trans_a: a = testing.shaped_arange(shape_a[::-1], xp, dtype_a).T else: a = testing.shaped_arange(shape_a, xp, dtype_a) if self.trans_b: b = testing.shaped_arange(shape_b[::-1], xp, dtype_b).T else: b = testing.shaped_arange(shape_b, xp, dtype_b) > return xp.dot(a, b) tests/cupy_tests/linalg_tests/test_product.py:49: _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ cupy/linalg/_product.py:65: in dot return a.dot(b, out) cupy/_core/core.pyx:1607: in cupy._core.core.ndarray.dot return _linalg.dot(self, b, out) cupy/_core/_routines_linalg.pyx:427: in cupy._core._routines_linalg.dot return tensordot_core(a, b, out, n, m, k, ret_shape) cupy/_core/_routines_linalg.pyx:501: in cupy._core._routines_linalg.tensordot_core return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) cupy/_core/_routines_linalg.pyx:327: in cupy._core._routines_linalg._integral_tensordot_core kern = _tensordot_core_int_kernel(config, dtype) cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/_routines_linalg.pyx:302: in cupy._core._routines_linalg._tensordot_core_int_kernel ker = mod.get_function( cupy/_core/raw.pyx:470: in cupy._core.raw.RawModule.get_function mangled_name = self.module.mapping.get(name) cupy/_core/raw.pyx:394: in cupy._core.raw.RawModule.module.__get__ return self._module() cupy/_core/raw.pyx:402: in cupy._core.raw.RawModule._module mod = _get_raw_module( cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/raw.pyx:547: in cupy._core.raw._get_raw_module mod = cupy._core.core.compile_with_cache( cupy/_core/core.pyx:2015: in cupy._core.core.compile_with_cache cpdef function.Module compile_with_cache( cupy/_core/core.pyx:2078: in cupy._core.core.compile_with_cache return cuda.compile_with_cache( cupy/cuda/compiler.py:453: in compile_with_cache return _compile_with_cache_hip( cupy/cuda/compiler.py:853: in _compile_with_cache_hip binary, mapping = compile_using_nvrtc( cupy/cuda/compiler.py:295: in compile_using_nvrtc return _compile(source, options, cu_path, cupy/cuda/compiler.py:279: in _compile compiled_obj, mapping = prog.compile(options, log_stream) _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ self = options = ('--std=c++11', '-I/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/_core/include', '-I/opt/rocm/include', '-fcuda-flush-denormals-to-zero') log_stream = None def compile(self, options=(), log_stream=None): try: if self.name_expressions: for ker in self.name_expressions: nvrtc.addAddNameExpression(self.ptr, ker) nvrtc.compileProgram(self.ptr, options) mapping = None if self.name_expressions: mapping = {} for ker in self.name_expressions: mapping[ker] = nvrtc.getLoweredName(self.ptr, ker) if log_stream is not None: log_stream.write(nvrtc.getProgramLog(self.ptr)) # This is to ensure backwards compatibility with nvrtc if self.method == 'cubin': return nvrtc.getCUBIN(self.ptr), mapping elif self.method == 'ptx': return nvrtc.getPTX(self.ptr), mapping else: raise RuntimeError('Unknown NVRTC compile method') except nvrtc.NVRTCError: log = nvrtc.getProgramLog(self.ptr) > raise CompileException(log, self.src, self.name, options, 'nvrtc' if not runtime.is_hip else 'hiprtc') E AssertionError: Only cupy raises error E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 650, in compile E nvrtc.compileProgram(self.ptr, options) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 136, in cupy_backends.cuda.libs.nvrtc.compileProgram E cpdef compileProgram(intptr_t prog, options): E File "cupy_backends/cuda/libs/nvrtc.pyx", line 148, in cupy_backends.cuda.libs.nvrtc.compileProgram E check_status(status) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 67, in cupy_backends.cuda.libs.nvrtc.check_status E raise NVRTCError(status) E cupy_backends.cuda.libs.nvrtc.NVRTCError: HIPRTC_ERROR_COMPILATION (6) E E During handling of the above exception, another exception occurred: E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/testing/_helper.py", line 47, in _call_func E result = impl(*args, **kw) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/tests/cupy_tests/linalg_tests/test_product.py", line 49, in test_dot E return xp.dot(a, b) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/linalg/_product.py", line 65, in dot E return a.dot(b, out) E File "cupy/_core/core.pyx", line 1607, in cupy._core.core.ndarray.dot E return _linalg.dot(self, b, out) E File "cupy/_core/_routines_linalg.pyx", line 427, in cupy._core._routines_linalg.dot E return tensordot_core(a, b, out, n, m, k, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 501, in cupy._core._routines_linalg.tensordot_core E return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 327, in cupy._core._routines_linalg._integral_tensordot_core E kern = _tensordot_core_int_kernel(config, dtype) E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/_routines_linalg.pyx", line 302, in cupy._core._routines_linalg._tensordot_core_int_kernel E ker = mod.get_function( E File "cupy/_core/raw.pyx", line 470, in cupy._core.raw.RawModule.get_function E mangled_name = self.module.mapping.get(name) E File "cupy/_core/raw.pyx", line 394, in cupy._core.raw.RawModule.module.__get__ E return self._module() E File "cupy/_core/raw.pyx", line 402, in cupy._core.raw.RawModule._module E mod = _get_raw_module( E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/raw.pyx", line 547, in cupy._core.raw._get_raw_module E mod = cupy._core.core.compile_with_cache( E File "cupy/_core/core.pyx", line 2015, in cupy._core.core.compile_with_cache E cpdef function.Module compile_with_cache( E File "cupy/_core/core.pyx", line 2078, in cupy._core.core.compile_with_cache E return cuda.compile_with_cache( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 453, in compile_with_cache E return _compile_with_cache_hip( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 853, in _compile_with_cache_hip E binary, mapping = compile_using_nvrtc( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 295, in compile_using_nvrtc E return _compile(source, options, cu_path, E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 279, in _compile E compiled_obj, mapping = prog.compile(options, log_stream) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 667, in compile E raise CompileException(log, self.src, self.name, options, E cupy.cuda.compiler.CompileException: /tmp/comgr-3ccf5d/input/CompileSource:203:27: error: redeclaration of '__hiprtc_16' with a different type: 'void (*const)(int, int, int, const signed char *, const signed char *, signed char *)' vs 'void (*const)(int, int, int, const bool *, const bool *, bool *)' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-3ccf5d/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-3ccf5d/input/CompileSource:204:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-3ccf5d/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-3ccf5d/input/CompileSource:205:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-3ccf5d/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-3ccf5d/input/CompileSource:206:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-3ccf5d/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-3ccf5d/input/CompileSource:207:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-3ccf5d/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-3ccf5d/input/CompileSource:208:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-3ccf5d/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-3ccf5d/input/CompileSource:209:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-3ccf5d/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-3ccf5d/input/CompileSource:210:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-3ccf5d/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-3ccf5d/input/CompileSource:211:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-3ccf5d/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-3ccf5d/input/CompileSource:212:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-3ccf5d/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E 10 errors generated when compiling for gfx908. E Error: Failed to compile opencl source (from CL or HIP source to LLVM IR). cupy/cuda/compiler.py:667: AssertionError ----------------------------- Captured stdout call ----------------------------- {'dtype_a': , 'dtype_b': } _ TestDot_param_20_{shape=((1, 2), (2, 3)), trans_a=True, trans_b=True}.test_dot _ self = < parameter: {'shape': ((1, 2), (2, 3)), 'trans_a': True, 'trans_b': True}> xp = dtype_a = , dtype_b = @testing.for_all_dtypes_combination(['dtype_a', 'dtype_b']) @testing.numpy_cupy_allclose() def test_dot(self, xp, dtype_a, dtype_b): shape_a, shape_b = self.shape if self.trans_a: a = testing.shaped_arange(shape_a[::-1], xp, dtype_a).T else: a = testing.shaped_arange(shape_a, xp, dtype_a) if self.trans_b: b = testing.shaped_arange(shape_b[::-1], xp, dtype_b).T else: b = testing.shaped_arange(shape_b, xp, dtype_b) > return xp.dot(a, b) tests/cupy_tests/linalg_tests/test_product.py:49: _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ cupy/linalg/_product.py:65: in dot return a.dot(b, out) cupy/_core/core.pyx:1607: in cupy._core.core.ndarray.dot return _linalg.dot(self, b, out) cupy/_core/_routines_linalg.pyx:427: in cupy._core._routines_linalg.dot return tensordot_core(a, b, out, n, m, k, ret_shape) cupy/_core/_routines_linalg.pyx:501: in cupy._core._routines_linalg.tensordot_core return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) cupy/_core/_routines_linalg.pyx:327: in cupy._core._routines_linalg._integral_tensordot_core kern = _tensordot_core_int_kernel(config, dtype) cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/_routines_linalg.pyx:302: in cupy._core._routines_linalg._tensordot_core_int_kernel ker = mod.get_function( cupy/_core/raw.pyx:470: in cupy._core.raw.RawModule.get_function mangled_name = self.module.mapping.get(name) cupy/_core/raw.pyx:394: in cupy._core.raw.RawModule.module.__get__ return self._module() cupy/_core/raw.pyx:402: in cupy._core.raw.RawModule._module mod = _get_raw_module( cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/raw.pyx:547: in cupy._core.raw._get_raw_module mod = cupy._core.core.compile_with_cache( cupy/_core/core.pyx:2015: in cupy._core.core.compile_with_cache cpdef function.Module compile_with_cache( cupy/_core/core.pyx:2078: in cupy._core.core.compile_with_cache return cuda.compile_with_cache( cupy/cuda/compiler.py:453: in compile_with_cache return _compile_with_cache_hip( cupy/cuda/compiler.py:853: in _compile_with_cache_hip binary, mapping = compile_using_nvrtc( cupy/cuda/compiler.py:295: in compile_using_nvrtc return _compile(source, options, cu_path, cupy/cuda/compiler.py:279: in _compile compiled_obj, mapping = prog.compile(options, log_stream) _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ self = options = ('--std=c++11', '-I/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/_core/include', '-I/opt/rocm/include', '-fcuda-flush-denormals-to-zero') log_stream = None def compile(self, options=(), log_stream=None): try: if self.name_expressions: for ker in self.name_expressions: nvrtc.addAddNameExpression(self.ptr, ker) nvrtc.compileProgram(self.ptr, options) mapping = None if self.name_expressions: mapping = {} for ker in self.name_expressions: mapping[ker] = nvrtc.getLoweredName(self.ptr, ker) if log_stream is not None: log_stream.write(nvrtc.getProgramLog(self.ptr)) # This is to ensure backwards compatibility with nvrtc if self.method == 'cubin': return nvrtc.getCUBIN(self.ptr), mapping elif self.method == 'ptx': return nvrtc.getPTX(self.ptr), mapping else: raise RuntimeError('Unknown NVRTC compile method') except nvrtc.NVRTCError: log = nvrtc.getProgramLog(self.ptr) > raise CompileException(log, self.src, self.name, options, 'nvrtc' if not runtime.is_hip else 'hiprtc') E AssertionError: Only cupy raises error E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 650, in compile E nvrtc.compileProgram(self.ptr, options) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 136, in cupy_backends.cuda.libs.nvrtc.compileProgram E cpdef compileProgram(intptr_t prog, options): E File "cupy_backends/cuda/libs/nvrtc.pyx", line 148, in cupy_backends.cuda.libs.nvrtc.compileProgram E check_status(status) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 67, in cupy_backends.cuda.libs.nvrtc.check_status E raise NVRTCError(status) E cupy_backends.cuda.libs.nvrtc.NVRTCError: HIPRTC_ERROR_COMPILATION (6) E E During handling of the above exception, another exception occurred: E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/testing/_helper.py", line 47, in _call_func E result = impl(*args, **kw) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/tests/cupy_tests/linalg_tests/test_product.py", line 49, in test_dot E return xp.dot(a, b) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/linalg/_product.py", line 65, in dot E return a.dot(b, out) E File "cupy/_core/core.pyx", line 1607, in cupy._core.core.ndarray.dot E return _linalg.dot(self, b, out) E File "cupy/_core/_routines_linalg.pyx", line 427, in cupy._core._routines_linalg.dot E return tensordot_core(a, b, out, n, m, k, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 501, in cupy._core._routines_linalg.tensordot_core E return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 327, in cupy._core._routines_linalg._integral_tensordot_core E kern = _tensordot_core_int_kernel(config, dtype) E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/_routines_linalg.pyx", line 302, in cupy._core._routines_linalg._tensordot_core_int_kernel E ker = mod.get_function( E File "cupy/_core/raw.pyx", line 470, in cupy._core.raw.RawModule.get_function E mangled_name = self.module.mapping.get(name) E File "cupy/_core/raw.pyx", line 394, in cupy._core.raw.RawModule.module.__get__ E return self._module() E File "cupy/_core/raw.pyx", line 402, in cupy._core.raw.RawModule._module E mod = _get_raw_module( E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/raw.pyx", line 547, in cupy._core.raw._get_raw_module E mod = cupy._core.core.compile_with_cache( E File "cupy/_core/core.pyx", line 2015, in cupy._core.core.compile_with_cache E cpdef function.Module compile_with_cache( E File "cupy/_core/core.pyx", line 2078, in cupy._core.core.compile_with_cache E return cuda.compile_with_cache( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 453, in compile_with_cache E return _compile_with_cache_hip( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 853, in _compile_with_cache_hip E binary, mapping = compile_using_nvrtc( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 295, in compile_using_nvrtc E return _compile(source, options, cu_path, E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 279, in _compile E compiled_obj, mapping = prog.compile(options, log_stream) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 667, in compile E raise CompileException(log, self.src, self.name, options, E cupy.cuda.compiler.CompileException: /tmp/comgr-aafdd2/input/CompileSource:203:27: error: redeclaration of '__hiprtc_16' with a different type: 'void (*const)(int, int, int, const signed char *, const signed char *, signed char *)' vs 'void (*const)(int, int, int, const bool *, const bool *, bool *)' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-aafdd2/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-aafdd2/input/CompileSource:204:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-aafdd2/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-aafdd2/input/CompileSource:205:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-aafdd2/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-aafdd2/input/CompileSource:206:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-aafdd2/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-aafdd2/input/CompileSource:207:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-aafdd2/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-aafdd2/input/CompileSource:208:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-aafdd2/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-aafdd2/input/CompileSource:209:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-aafdd2/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-aafdd2/input/CompileSource:210:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-aafdd2/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-aafdd2/input/CompileSource:211:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-aafdd2/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-aafdd2/input/CompileSource:212:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-aafdd2/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E 10 errors generated when compiling for gfx908. E Error: Failed to compile opencl source (from CL or HIP source to LLVM IR). cupy/cuda/compiler.py:667: AssertionError ----------------------------- Captured stdout call ----------------------------- {'dtype_a': , 'dtype_b': } _ TestDot_param_21_{shape=((1, 2), (2, 3)), trans_a=True, trans_b=False}.test_dot _ self = < parameter: {'shape': ((1, 2), (2, 3)), 'trans_a': True, 'trans_b': False}> xp = dtype_a = , dtype_b = @testing.for_all_dtypes_combination(['dtype_a', 'dtype_b']) @testing.numpy_cupy_allclose() def test_dot(self, xp, dtype_a, dtype_b): shape_a, shape_b = self.shape if self.trans_a: a = testing.shaped_arange(shape_a[::-1], xp, dtype_a).T else: a = testing.shaped_arange(shape_a, xp, dtype_a) if self.trans_b: b = testing.shaped_arange(shape_b[::-1], xp, dtype_b).T else: b = testing.shaped_arange(shape_b, xp, dtype_b) > return xp.dot(a, b) tests/cupy_tests/linalg_tests/test_product.py:49: _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ cupy/linalg/_product.py:65: in dot return a.dot(b, out) cupy/_core/core.pyx:1607: in cupy._core.core.ndarray.dot return _linalg.dot(self, b, out) cupy/_core/_routines_linalg.pyx:427: in cupy._core._routines_linalg.dot return tensordot_core(a, b, out, n, m, k, ret_shape) cupy/_core/_routines_linalg.pyx:501: in cupy._core._routines_linalg.tensordot_core return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) cupy/_core/_routines_linalg.pyx:327: in cupy._core._routines_linalg._integral_tensordot_core kern = _tensordot_core_int_kernel(config, dtype) cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/_routines_linalg.pyx:302: in cupy._core._routines_linalg._tensordot_core_int_kernel ker = mod.get_function( cupy/_core/raw.pyx:470: in cupy._core.raw.RawModule.get_function mangled_name = self.module.mapping.get(name) cupy/_core/raw.pyx:394: in cupy._core.raw.RawModule.module.__get__ return self._module() cupy/_core/raw.pyx:402: in cupy._core.raw.RawModule._module mod = _get_raw_module( cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/raw.pyx:547: in cupy._core.raw._get_raw_module mod = cupy._core.core.compile_with_cache( cupy/_core/core.pyx:2015: in cupy._core.core.compile_with_cache cpdef function.Module compile_with_cache( cupy/_core/core.pyx:2078: in cupy._core.core.compile_with_cache return cuda.compile_with_cache( cupy/cuda/compiler.py:453: in compile_with_cache return _compile_with_cache_hip( cupy/cuda/compiler.py:853: in _compile_with_cache_hip binary, mapping = compile_using_nvrtc( cupy/cuda/compiler.py:295: in compile_using_nvrtc return _compile(source, options, cu_path, cupy/cuda/compiler.py:279: in _compile compiled_obj, mapping = prog.compile(options, log_stream) _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ self = options = ('--std=c++11', '-I/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/_core/include', '-I/opt/rocm/include', '-fcuda-flush-denormals-to-zero') log_stream = None def compile(self, options=(), log_stream=None): try: if self.name_expressions: for ker in self.name_expressions: nvrtc.addAddNameExpression(self.ptr, ker) nvrtc.compileProgram(self.ptr, options) mapping = None if self.name_expressions: mapping = {} for ker in self.name_expressions: mapping[ker] = nvrtc.getLoweredName(self.ptr, ker) if log_stream is not None: log_stream.write(nvrtc.getProgramLog(self.ptr)) # This is to ensure backwards compatibility with nvrtc if self.method == 'cubin': return nvrtc.getCUBIN(self.ptr), mapping elif self.method == 'ptx': return nvrtc.getPTX(self.ptr), mapping else: raise RuntimeError('Unknown NVRTC compile method') except nvrtc.NVRTCError: log = nvrtc.getProgramLog(self.ptr) > raise CompileException(log, self.src, self.name, options, 'nvrtc' if not runtime.is_hip else 'hiprtc') E AssertionError: Only cupy raises error E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 650, in compile E nvrtc.compileProgram(self.ptr, options) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 136, in cupy_backends.cuda.libs.nvrtc.compileProgram E cpdef compileProgram(intptr_t prog, options): E File "cupy_backends/cuda/libs/nvrtc.pyx", line 148, in cupy_backends.cuda.libs.nvrtc.compileProgram E check_status(status) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 67, in cupy_backends.cuda.libs.nvrtc.check_status E raise NVRTCError(status) E cupy_backends.cuda.libs.nvrtc.NVRTCError: HIPRTC_ERROR_COMPILATION (6) E E During handling of the above exception, another exception occurred: E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/testing/_helper.py", line 47, in _call_func E result = impl(*args, **kw) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/tests/cupy_tests/linalg_tests/test_product.py", line 49, in test_dot E return xp.dot(a, b) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/linalg/_product.py", line 65, in dot E return a.dot(b, out) E File "cupy/_core/core.pyx", line 1607, in cupy._core.core.ndarray.dot E return _linalg.dot(self, b, out) E File "cupy/_core/_routines_linalg.pyx", line 427, in cupy._core._routines_linalg.dot E return tensordot_core(a, b, out, n, m, k, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 501, in cupy._core._routines_linalg.tensordot_core E return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 327, in cupy._core._routines_linalg._integral_tensordot_core E kern = _tensordot_core_int_kernel(config, dtype) E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/_routines_linalg.pyx", line 302, in cupy._core._routines_linalg._tensordot_core_int_kernel E ker = mod.get_function( E File "cupy/_core/raw.pyx", line 470, in cupy._core.raw.RawModule.get_function E mangled_name = self.module.mapping.get(name) E File "cupy/_core/raw.pyx", line 394, in cupy._core.raw.RawModule.module.__get__ E return self._module() E File "cupy/_core/raw.pyx", line 402, in cupy._core.raw.RawModule._module E mod = _get_raw_module( E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/raw.pyx", line 547, in cupy._core.raw._get_raw_module E mod = cupy._core.core.compile_with_cache( E File "cupy/_core/core.pyx", line 2015, in cupy._core.core.compile_with_cache E cpdef function.Module compile_with_cache( E File "cupy/_core/core.pyx", line 2078, in cupy._core.core.compile_with_cache E return cuda.compile_with_cache( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 453, in compile_with_cache E return _compile_with_cache_hip( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 853, in _compile_with_cache_hip E binary, mapping = compile_using_nvrtc( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 295, in compile_using_nvrtc E return _compile(source, options, cu_path, E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 279, in _compile E compiled_obj, mapping = prog.compile(options, log_stream) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 667, in compile E raise CompileException(log, self.src, self.name, options, E cupy.cuda.compiler.CompileException: /tmp/comgr-cc2f9e/input/CompileSource:203:27: error: redeclaration of '__hiprtc_16' with a different type: 'void (*const)(int, int, int, const signed char *, const signed char *, signed char *)' vs 'void (*const)(int, int, int, const bool *, const bool *, bool *)' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-cc2f9e/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-cc2f9e/input/CompileSource:204:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-cc2f9e/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-cc2f9e/input/CompileSource:205:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-cc2f9e/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-cc2f9e/input/CompileSource:206:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-cc2f9e/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-cc2f9e/input/CompileSource:207:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-cc2f9e/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-cc2f9e/input/CompileSource:208:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-cc2f9e/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-cc2f9e/input/CompileSource:209:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-cc2f9e/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-cc2f9e/input/CompileSource:210:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-cc2f9e/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-cc2f9e/input/CompileSource:211:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-cc2f9e/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-cc2f9e/input/CompileSource:212:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-cc2f9e/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E 10 errors generated when compiling for gfx908. E Error: Failed to compile opencl source (from CL or HIP source to LLVM IR). cupy/cuda/compiler.py:667: AssertionError ----------------------------- Captured stdout call ----------------------------- {'dtype_a': , 'dtype_b': } _ TestDot_param_22_{shape=((1, 2), (2, 3)), trans_a=False, trans_b=True}.test_dot _ self = < parameter: {'shape': ((1, 2), (2, 3)), 'trans_a': False, 'trans_b': True}> xp = dtype_a = , dtype_b = @testing.for_all_dtypes_combination(['dtype_a', 'dtype_b']) @testing.numpy_cupy_allclose() def test_dot(self, xp, dtype_a, dtype_b): shape_a, shape_b = self.shape if self.trans_a: a = testing.shaped_arange(shape_a[::-1], xp, dtype_a).T else: a = testing.shaped_arange(shape_a, xp, dtype_a) if self.trans_b: b = testing.shaped_arange(shape_b[::-1], xp, dtype_b).T else: b = testing.shaped_arange(shape_b, xp, dtype_b) > return xp.dot(a, b) tests/cupy_tests/linalg_tests/test_product.py:49: _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ cupy/linalg/_product.py:65: in dot return a.dot(b, out) cupy/_core/core.pyx:1607: in cupy._core.core.ndarray.dot return _linalg.dot(self, b, out) cupy/_core/_routines_linalg.pyx:427: in cupy._core._routines_linalg.dot return tensordot_core(a, b, out, n, m, k, ret_shape) cupy/_core/_routines_linalg.pyx:501: in cupy._core._routines_linalg.tensordot_core return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) cupy/_core/_routines_linalg.pyx:327: in cupy._core._routines_linalg._integral_tensordot_core kern = _tensordot_core_int_kernel(config, dtype) cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/_routines_linalg.pyx:302: in cupy._core._routines_linalg._tensordot_core_int_kernel ker = mod.get_function( cupy/_core/raw.pyx:470: in cupy._core.raw.RawModule.get_function mangled_name = self.module.mapping.get(name) cupy/_core/raw.pyx:394: in cupy._core.raw.RawModule.module.__get__ return self._module() cupy/_core/raw.pyx:402: in cupy._core.raw.RawModule._module mod = _get_raw_module( cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/raw.pyx:547: in cupy._core.raw._get_raw_module mod = cupy._core.core.compile_with_cache( cupy/_core/core.pyx:2015: in cupy._core.core.compile_with_cache cpdef function.Module compile_with_cache( cupy/_core/core.pyx:2078: in cupy._core.core.compile_with_cache return cuda.compile_with_cache( cupy/cuda/compiler.py:453: in compile_with_cache return _compile_with_cache_hip( cupy/cuda/compiler.py:853: in _compile_with_cache_hip binary, mapping = compile_using_nvrtc( cupy/cuda/compiler.py:295: in compile_using_nvrtc return _compile(source, options, cu_path, cupy/cuda/compiler.py:279: in _compile compiled_obj, mapping = prog.compile(options, log_stream) _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ self = options = ('--std=c++11', '-I/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/_core/include', '-I/opt/rocm/include', '-fcuda-flush-denormals-to-zero') log_stream = None def compile(self, options=(), log_stream=None): try: if self.name_expressions: for ker in self.name_expressions: nvrtc.addAddNameExpression(self.ptr, ker) nvrtc.compileProgram(self.ptr, options) mapping = None if self.name_expressions: mapping = {} for ker in self.name_expressions: mapping[ker] = nvrtc.getLoweredName(self.ptr, ker) if log_stream is not None: log_stream.write(nvrtc.getProgramLog(self.ptr)) # This is to ensure backwards compatibility with nvrtc if self.method == 'cubin': return nvrtc.getCUBIN(self.ptr), mapping elif self.method == 'ptx': return nvrtc.getPTX(self.ptr), mapping else: raise RuntimeError('Unknown NVRTC compile method') except nvrtc.NVRTCError: log = nvrtc.getProgramLog(self.ptr) > raise CompileException(log, self.src, self.name, options, 'nvrtc' if not runtime.is_hip else 'hiprtc') E AssertionError: Only cupy raises error E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 650, in compile E nvrtc.compileProgram(self.ptr, options) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 136, in cupy_backends.cuda.libs.nvrtc.compileProgram E cpdef compileProgram(intptr_t prog, options): E File "cupy_backends/cuda/libs/nvrtc.pyx", line 148, in cupy_backends.cuda.libs.nvrtc.compileProgram E check_status(status) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 67, in cupy_backends.cuda.libs.nvrtc.check_status E raise NVRTCError(status) E cupy_backends.cuda.libs.nvrtc.NVRTCError: HIPRTC_ERROR_COMPILATION (6) E E During handling of the above exception, another exception occurred: E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/testing/_helper.py", line 47, in _call_func E result = impl(*args, **kw) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/tests/cupy_tests/linalg_tests/test_product.py", line 49, in test_dot E return xp.dot(a, b) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/linalg/_product.py", line 65, in dot E return a.dot(b, out) E File "cupy/_core/core.pyx", line 1607, in cupy._core.core.ndarray.dot E return _linalg.dot(self, b, out) E File "cupy/_core/_routines_linalg.pyx", line 427, in cupy._core._routines_linalg.dot E return tensordot_core(a, b, out, n, m, k, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 501, in cupy._core._routines_linalg.tensordot_core E return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 327, in cupy._core._routines_linalg._integral_tensordot_core E kern = _tensordot_core_int_kernel(config, dtype) E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/_routines_linalg.pyx", line 302, in cupy._core._routines_linalg._tensordot_core_int_kernel E ker = mod.get_function( E File "cupy/_core/raw.pyx", line 470, in cupy._core.raw.RawModule.get_function E mangled_name = self.module.mapping.get(name) E File "cupy/_core/raw.pyx", line 394, in cupy._core.raw.RawModule.module.__get__ E return self._module() E File "cupy/_core/raw.pyx", line 402, in cupy._core.raw.RawModule._module E mod = _get_raw_module( E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/raw.pyx", line 547, in cupy._core.raw._get_raw_module E mod = cupy._core.core.compile_with_cache( E File "cupy/_core/core.pyx", line 2015, in cupy._core.core.compile_with_cache E cpdef function.Module compile_with_cache( E File "cupy/_core/core.pyx", line 2078, in cupy._core.core.compile_with_cache E return cuda.compile_with_cache( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 453, in compile_with_cache E return _compile_with_cache_hip( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 853, in _compile_with_cache_hip E binary, mapping = compile_using_nvrtc( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 295, in compile_using_nvrtc E return _compile(source, options, cu_path, E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 279, in _compile E compiled_obj, mapping = prog.compile(options, log_stream) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 667, in compile E raise CompileException(log, self.src, self.name, options, E cupy.cuda.compiler.CompileException: /tmp/comgr-6131b3/input/CompileSource:203:27: error: redeclaration of '__hiprtc_16' with a different type: 'void (*const)(int, int, int, const signed char *, const signed char *, signed char *)' vs 'void (*const)(int, int, int, const bool *, const bool *, bool *)' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-6131b3/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-6131b3/input/CompileSource:204:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-6131b3/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-6131b3/input/CompileSource:205:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-6131b3/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-6131b3/input/CompileSource:206:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-6131b3/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-6131b3/input/CompileSource:207:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-6131b3/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-6131b3/input/CompileSource:208:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-6131b3/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-6131b3/input/CompileSource:209:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-6131b3/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-6131b3/input/CompileSource:210:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-6131b3/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-6131b3/input/CompileSource:211:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-6131b3/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-6131b3/input/CompileSource:212:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-6131b3/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E 10 errors generated when compiling for gfx908. E Error: Failed to compile opencl source (from CL or HIP source to LLVM IR). cupy/cuda/compiler.py:667: AssertionError ----------------------------- Captured stdout call ----------------------------- {'dtype_a': , 'dtype_b': } _ TestDot_param_23_{shape=((1, 2), (2, 3)), trans_a=False, trans_b=False}.test_dot _ self = < parameter: {'shape': ((1, 2), (2, 3)), 'trans_a': False, 'trans_b': False}> xp = dtype_a = , dtype_b = @testing.for_all_dtypes_combination(['dtype_a', 'dtype_b']) @testing.numpy_cupy_allclose() def test_dot(self, xp, dtype_a, dtype_b): shape_a, shape_b = self.shape if self.trans_a: a = testing.shaped_arange(shape_a[::-1], xp, dtype_a).T else: a = testing.shaped_arange(shape_a, xp, dtype_a) if self.trans_b: b = testing.shaped_arange(shape_b[::-1], xp, dtype_b).T else: b = testing.shaped_arange(shape_b, xp, dtype_b) > return xp.dot(a, b) tests/cupy_tests/linalg_tests/test_product.py:49: _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ cupy/linalg/_product.py:65: in dot return a.dot(b, out) cupy/_core/core.pyx:1607: in cupy._core.core.ndarray.dot return _linalg.dot(self, b, out) cupy/_core/_routines_linalg.pyx:427: in cupy._core._routines_linalg.dot return tensordot_core(a, b, out, n, m, k, ret_shape) cupy/_core/_routines_linalg.pyx:501: in cupy._core._routines_linalg.tensordot_core return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) cupy/_core/_routines_linalg.pyx:327: in cupy._core._routines_linalg._integral_tensordot_core kern = _tensordot_core_int_kernel(config, dtype) cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/_routines_linalg.pyx:302: in cupy._core._routines_linalg._tensordot_core_int_kernel ker = mod.get_function( cupy/_core/raw.pyx:470: in cupy._core.raw.RawModule.get_function mangled_name = self.module.mapping.get(name) cupy/_core/raw.pyx:394: in cupy._core.raw.RawModule.module.__get__ return self._module() cupy/_core/raw.pyx:402: in cupy._core.raw.RawModule._module mod = _get_raw_module( cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/raw.pyx:547: in cupy._core.raw._get_raw_module mod = cupy._core.core.compile_with_cache( cupy/_core/core.pyx:2015: in cupy._core.core.compile_with_cache cpdef function.Module compile_with_cache( cupy/_core/core.pyx:2078: in cupy._core.core.compile_with_cache return cuda.compile_with_cache( cupy/cuda/compiler.py:453: in compile_with_cache return _compile_with_cache_hip( cupy/cuda/compiler.py:853: in _compile_with_cache_hip binary, mapping = compile_using_nvrtc( cupy/cuda/compiler.py:295: in compile_using_nvrtc return _compile(source, options, cu_path, cupy/cuda/compiler.py:279: in _compile compiled_obj, mapping = prog.compile(options, log_stream) _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ self = options = ('--std=c++11', '-I/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/_core/include', '-I/opt/rocm/include', '-fcuda-flush-denormals-to-zero') log_stream = None def compile(self, options=(), log_stream=None): try: if self.name_expressions: for ker in self.name_expressions: nvrtc.addAddNameExpression(self.ptr, ker) nvrtc.compileProgram(self.ptr, options) mapping = None if self.name_expressions: mapping = {} for ker in self.name_expressions: mapping[ker] = nvrtc.getLoweredName(self.ptr, ker) if log_stream is not None: log_stream.write(nvrtc.getProgramLog(self.ptr)) # This is to ensure backwards compatibility with nvrtc if self.method == 'cubin': return nvrtc.getCUBIN(self.ptr), mapping elif self.method == 'ptx': return nvrtc.getPTX(self.ptr), mapping else: raise RuntimeError('Unknown NVRTC compile method') except nvrtc.NVRTCError: log = nvrtc.getProgramLog(self.ptr) > raise CompileException(log, self.src, self.name, options, 'nvrtc' if not runtime.is_hip else 'hiprtc') E AssertionError: Only cupy raises error E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 650, in compile E nvrtc.compileProgram(self.ptr, options) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 136, in cupy_backends.cuda.libs.nvrtc.compileProgram E cpdef compileProgram(intptr_t prog, options): E File "cupy_backends/cuda/libs/nvrtc.pyx", line 148, in cupy_backends.cuda.libs.nvrtc.compileProgram E check_status(status) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 67, in cupy_backends.cuda.libs.nvrtc.check_status E raise NVRTCError(status) E cupy_backends.cuda.libs.nvrtc.NVRTCError: HIPRTC_ERROR_COMPILATION (6) E E During handling of the above exception, another exception occurred: E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/testing/_helper.py", line 47, in _call_func E result = impl(*args, **kw) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/tests/cupy_tests/linalg_tests/test_product.py", line 49, in test_dot E return xp.dot(a, b) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/linalg/_product.py", line 65, in dot E return a.dot(b, out) E File "cupy/_core/core.pyx", line 1607, in cupy._core.core.ndarray.dot E return _linalg.dot(self, b, out) E File "cupy/_core/_routines_linalg.pyx", line 427, in cupy._core._routines_linalg.dot E return tensordot_core(a, b, out, n, m, k, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 501, in cupy._core._routines_linalg.tensordot_core E return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 327, in cupy._core._routines_linalg._integral_tensordot_core E kern = _tensordot_core_int_kernel(config, dtype) E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/_routines_linalg.pyx", line 302, in cupy._core._routines_linalg._tensordot_core_int_kernel E ker = mod.get_function( E File "cupy/_core/raw.pyx", line 470, in cupy._core.raw.RawModule.get_function E mangled_name = self.module.mapping.get(name) E File "cupy/_core/raw.pyx", line 394, in cupy._core.raw.RawModule.module.__get__ E return self._module() E File "cupy/_core/raw.pyx", line 402, in cupy._core.raw.RawModule._module E mod = _get_raw_module( E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/raw.pyx", line 547, in cupy._core.raw._get_raw_module E mod = cupy._core.core.compile_with_cache( E File "cupy/_core/core.pyx", line 2015, in cupy._core.core.compile_with_cache E cpdef function.Module compile_with_cache( E File "cupy/_core/core.pyx", line 2078, in cupy._core.core.compile_with_cache E return cuda.compile_with_cache( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 453, in compile_with_cache E return _compile_with_cache_hip( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 853, in _compile_with_cache_hip E binary, mapping = compile_using_nvrtc( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 295, in compile_using_nvrtc E return _compile(source, options, cu_path, E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 279, in _compile E compiled_obj, mapping = prog.compile(options, log_stream) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 667, in compile E raise CompileException(log, self.src, self.name, options, E cupy.cuda.compiler.CompileException: /tmp/comgr-6c6288/input/CompileSource:203:27: error: redeclaration of '__hiprtc_16' with a different type: 'void (*const)(int, int, int, const signed char *, const signed char *, signed char *)' vs 'void (*const)(int, int, int, const bool *, const bool *, bool *)' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-6c6288/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-6c6288/input/CompileSource:204:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-6c6288/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-6c6288/input/CompileSource:205:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-6c6288/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-6c6288/input/CompileSource:206:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-6c6288/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-6c6288/input/CompileSource:207:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-6c6288/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-6c6288/input/CompileSource:208:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-6c6288/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-6c6288/input/CompileSource:209:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-6c6288/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-6c6288/input/CompileSource:210:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-6c6288/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-6c6288/input/CompileSource:211:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-6c6288/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-6c6288/input/CompileSource:212:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-6c6288/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E 10 errors generated when compiling for gfx908. E Error: Failed to compile opencl source (from CL or HIP source to LLVM IR). cupy/cuda/compiler.py:667: AssertionError ----------------------------- Captured stdout call ----------------------------- {'dtype_a': , 'dtype_b': } _ TestDot_param_24_{shape=((2, 1), (1, 3)), trans_a=True, trans_b=True}.test_dot _ self = < parameter: {'shape': ((2, 1), (1, 3)), 'trans_a': True, 'trans_b': True}> xp = dtype_a = , dtype_b = @testing.for_all_dtypes_combination(['dtype_a', 'dtype_b']) @testing.numpy_cupy_allclose() def test_dot(self, xp, dtype_a, dtype_b): shape_a, shape_b = self.shape if self.trans_a: a = testing.shaped_arange(shape_a[::-1], xp, dtype_a).T else: a = testing.shaped_arange(shape_a, xp, dtype_a) if self.trans_b: b = testing.shaped_arange(shape_b[::-1], xp, dtype_b).T else: b = testing.shaped_arange(shape_b, xp, dtype_b) > return xp.dot(a, b) tests/cupy_tests/linalg_tests/test_product.py:49: _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ cupy/linalg/_product.py:65: in dot return a.dot(b, out) cupy/_core/core.pyx:1607: in cupy._core.core.ndarray.dot return _linalg.dot(self, b, out) cupy/_core/_routines_linalg.pyx:427: in cupy._core._routines_linalg.dot return tensordot_core(a, b, out, n, m, k, ret_shape) cupy/_core/_routines_linalg.pyx:501: in cupy._core._routines_linalg.tensordot_core return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) cupy/_core/_routines_linalg.pyx:327: in cupy._core._routines_linalg._integral_tensordot_core kern = _tensordot_core_int_kernel(config, dtype) cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/_routines_linalg.pyx:302: in cupy._core._routines_linalg._tensordot_core_int_kernel ker = mod.get_function( cupy/_core/raw.pyx:470: in cupy._core.raw.RawModule.get_function mangled_name = self.module.mapping.get(name) cupy/_core/raw.pyx:394: in cupy._core.raw.RawModule.module.__get__ return self._module() cupy/_core/raw.pyx:402: in cupy._core.raw.RawModule._module mod = _get_raw_module( cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/raw.pyx:547: in cupy._core.raw._get_raw_module mod = cupy._core.core.compile_with_cache( cupy/_core/core.pyx:2015: in cupy._core.core.compile_with_cache cpdef function.Module compile_with_cache( cupy/_core/core.pyx:2078: in cupy._core.core.compile_with_cache return cuda.compile_with_cache( cupy/cuda/compiler.py:453: in compile_with_cache return _compile_with_cache_hip( cupy/cuda/compiler.py:853: in _compile_with_cache_hip binary, mapping = compile_using_nvrtc( cupy/cuda/compiler.py:295: in compile_using_nvrtc return _compile(source, options, cu_path, cupy/cuda/compiler.py:279: in _compile compiled_obj, mapping = prog.compile(options, log_stream) _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ self = options = ('--std=c++11', '-I/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/_core/include', '-I/opt/rocm/include', '-fcuda-flush-denormals-to-zero') log_stream = None def compile(self, options=(), log_stream=None): try: if self.name_expressions: for ker in self.name_expressions: nvrtc.addAddNameExpression(self.ptr, ker) nvrtc.compileProgram(self.ptr, options) mapping = None if self.name_expressions: mapping = {} for ker in self.name_expressions: mapping[ker] = nvrtc.getLoweredName(self.ptr, ker) if log_stream is not None: log_stream.write(nvrtc.getProgramLog(self.ptr)) # This is to ensure backwards compatibility with nvrtc if self.method == 'cubin': return nvrtc.getCUBIN(self.ptr), mapping elif self.method == 'ptx': return nvrtc.getPTX(self.ptr), mapping else: raise RuntimeError('Unknown NVRTC compile method') except nvrtc.NVRTCError: log = nvrtc.getProgramLog(self.ptr) > raise CompileException(log, self.src, self.name, options, 'nvrtc' if not runtime.is_hip else 'hiprtc') E AssertionError: Only cupy raises error E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 650, in compile E nvrtc.compileProgram(self.ptr, options) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 136, in cupy_backends.cuda.libs.nvrtc.compileProgram E cpdef compileProgram(intptr_t prog, options): E File "cupy_backends/cuda/libs/nvrtc.pyx", line 148, in cupy_backends.cuda.libs.nvrtc.compileProgram E check_status(status) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 67, in cupy_backends.cuda.libs.nvrtc.check_status E raise NVRTCError(status) E cupy_backends.cuda.libs.nvrtc.NVRTCError: HIPRTC_ERROR_COMPILATION (6) E E During handling of the above exception, another exception occurred: E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/testing/_helper.py", line 47, in _call_func E result = impl(*args, **kw) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/tests/cupy_tests/linalg_tests/test_product.py", line 49, in test_dot E return xp.dot(a, b) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/linalg/_product.py", line 65, in dot E return a.dot(b, out) E File "cupy/_core/core.pyx", line 1607, in cupy._core.core.ndarray.dot E return _linalg.dot(self, b, out) E File "cupy/_core/_routines_linalg.pyx", line 427, in cupy._core._routines_linalg.dot E return tensordot_core(a, b, out, n, m, k, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 501, in cupy._core._routines_linalg.tensordot_core E return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 327, in cupy._core._routines_linalg._integral_tensordot_core E kern = _tensordot_core_int_kernel(config, dtype) E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/_routines_linalg.pyx", line 302, in cupy._core._routines_linalg._tensordot_core_int_kernel E ker = mod.get_function( E File "cupy/_core/raw.pyx", line 470, in cupy._core.raw.RawModule.get_function E mangled_name = self.module.mapping.get(name) E File "cupy/_core/raw.pyx", line 394, in cupy._core.raw.RawModule.module.__get__ E return self._module() E File "cupy/_core/raw.pyx", line 402, in cupy._core.raw.RawModule._module E mod = _get_raw_module( E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/raw.pyx", line 547, in cupy._core.raw._get_raw_module E mod = cupy._core.core.compile_with_cache( E File "cupy/_core/core.pyx", line 2015, in cupy._core.core.compile_with_cache E cpdef function.Module compile_with_cache( E File "cupy/_core/core.pyx", line 2078, in cupy._core.core.compile_with_cache E return cuda.compile_with_cache( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 453, in compile_with_cache E return _compile_with_cache_hip( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 853, in _compile_with_cache_hip E binary, mapping = compile_using_nvrtc( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 295, in compile_using_nvrtc E return _compile(source, options, cu_path, E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 279, in _compile E compiled_obj, mapping = prog.compile(options, log_stream) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 667, in compile E raise CompileException(log, self.src, self.name, options, E cupy.cuda.compiler.CompileException: /tmp/comgr-4e8b0c/input/CompileSource:203:27: error: redeclaration of '__hiprtc_16' with a different type: 'void (*const)(int, int, int, const signed char *, const signed char *, signed char *)' vs 'void (*const)(int, int, int, const bool *, const bool *, bool *)' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-4e8b0c/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-4e8b0c/input/CompileSource:204:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-4e8b0c/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-4e8b0c/input/CompileSource:205:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-4e8b0c/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-4e8b0c/input/CompileSource:206:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-4e8b0c/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-4e8b0c/input/CompileSource:207:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-4e8b0c/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-4e8b0c/input/CompileSource:208:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-4e8b0c/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-4e8b0c/input/CompileSource:209:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-4e8b0c/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-4e8b0c/input/CompileSource:210:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-4e8b0c/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-4e8b0c/input/CompileSource:211:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-4e8b0c/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-4e8b0c/input/CompileSource:212:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-4e8b0c/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E 10 errors generated when compiling for gfx908. E Error: Failed to compile opencl source (from CL or HIP source to LLVM IR). cupy/cuda/compiler.py:667: AssertionError ----------------------------- Captured stdout call ----------------------------- {'dtype_a': , 'dtype_b': } _ TestDot_param_25_{shape=((2, 1), (1, 3)), trans_a=True, trans_b=False}.test_dot _ self = < parameter: {'shape': ((2, 1), (1, 3)), 'trans_a': True, 'trans_b': False}> xp = dtype_a = , dtype_b = @testing.for_all_dtypes_combination(['dtype_a', 'dtype_b']) @testing.numpy_cupy_allclose() def test_dot(self, xp, dtype_a, dtype_b): shape_a, shape_b = self.shape if self.trans_a: a = testing.shaped_arange(shape_a[::-1], xp, dtype_a).T else: a = testing.shaped_arange(shape_a, xp, dtype_a) if self.trans_b: b = testing.shaped_arange(shape_b[::-1], xp, dtype_b).T else: b = testing.shaped_arange(shape_b, xp, dtype_b) > return xp.dot(a, b) tests/cupy_tests/linalg_tests/test_product.py:49: _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ cupy/linalg/_product.py:65: in dot return a.dot(b, out) cupy/_core/core.pyx:1607: in cupy._core.core.ndarray.dot return _linalg.dot(self, b, out) cupy/_core/_routines_linalg.pyx:427: in cupy._core._routines_linalg.dot return tensordot_core(a, b, out, n, m, k, ret_shape) cupy/_core/_routines_linalg.pyx:501: in cupy._core._routines_linalg.tensordot_core return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) cupy/_core/_routines_linalg.pyx:327: in cupy._core._routines_linalg._integral_tensordot_core kern = _tensordot_core_int_kernel(config, dtype) cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/_routines_linalg.pyx:302: in cupy._core._routines_linalg._tensordot_core_int_kernel ker = mod.get_function( cupy/_core/raw.pyx:470: in cupy._core.raw.RawModule.get_function mangled_name = self.module.mapping.get(name) cupy/_core/raw.pyx:394: in cupy._core.raw.RawModule.module.__get__ return self._module() cupy/_core/raw.pyx:402: in cupy._core.raw.RawModule._module mod = _get_raw_module( cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/raw.pyx:547: in cupy._core.raw._get_raw_module mod = cupy._core.core.compile_with_cache( cupy/_core/core.pyx:2015: in cupy._core.core.compile_with_cache cpdef function.Module compile_with_cache( cupy/_core/core.pyx:2078: in cupy._core.core.compile_with_cache return cuda.compile_with_cache( cupy/cuda/compiler.py:453: in compile_with_cache return _compile_with_cache_hip( cupy/cuda/compiler.py:853: in _compile_with_cache_hip binary, mapping = compile_using_nvrtc( cupy/cuda/compiler.py:295: in compile_using_nvrtc return _compile(source, options, cu_path, cupy/cuda/compiler.py:279: in _compile compiled_obj, mapping = prog.compile(options, log_stream) _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ self = options = ('--std=c++11', '-I/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/_core/include', '-I/opt/rocm/include', '-fcuda-flush-denormals-to-zero') log_stream = None def compile(self, options=(), log_stream=None): try: if self.name_expressions: for ker in self.name_expressions: nvrtc.addAddNameExpression(self.ptr, ker) nvrtc.compileProgram(self.ptr, options) mapping = None if self.name_expressions: mapping = {} for ker in self.name_expressions: mapping[ker] = nvrtc.getLoweredName(self.ptr, ker) if log_stream is not None: log_stream.write(nvrtc.getProgramLog(self.ptr)) # This is to ensure backwards compatibility with nvrtc if self.method == 'cubin': return nvrtc.getCUBIN(self.ptr), mapping elif self.method == 'ptx': return nvrtc.getPTX(self.ptr), mapping else: raise RuntimeError('Unknown NVRTC compile method') except nvrtc.NVRTCError: log = nvrtc.getProgramLog(self.ptr) > raise CompileException(log, self.src, self.name, options, 'nvrtc' if not runtime.is_hip else 'hiprtc') E AssertionError: Only cupy raises error E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 650, in compile E nvrtc.compileProgram(self.ptr, options) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 136, in cupy_backends.cuda.libs.nvrtc.compileProgram E cpdef compileProgram(intptr_t prog, options): E File "cupy_backends/cuda/libs/nvrtc.pyx", line 148, in cupy_backends.cuda.libs.nvrtc.compileProgram E check_status(status) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 67, in cupy_backends.cuda.libs.nvrtc.check_status E raise NVRTCError(status) E cupy_backends.cuda.libs.nvrtc.NVRTCError: HIPRTC_ERROR_COMPILATION (6) E E During handling of the above exception, another exception occurred: E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/testing/_helper.py", line 47, in _call_func E result = impl(*args, **kw) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/tests/cupy_tests/linalg_tests/test_product.py", line 49, in test_dot E return xp.dot(a, b) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/linalg/_product.py", line 65, in dot E return a.dot(b, out) E File "cupy/_core/core.pyx", line 1607, in cupy._core.core.ndarray.dot E return _linalg.dot(self, b, out) E File "cupy/_core/_routines_linalg.pyx", line 427, in cupy._core._routines_linalg.dot E return tensordot_core(a, b, out, n, m, k, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 501, in cupy._core._routines_linalg.tensordot_core E return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 327, in cupy._core._routines_linalg._integral_tensordot_core E kern = _tensordot_core_int_kernel(config, dtype) E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/_routines_linalg.pyx", line 302, in cupy._core._routines_linalg._tensordot_core_int_kernel E ker = mod.get_function( E File "cupy/_core/raw.pyx", line 470, in cupy._core.raw.RawModule.get_function E mangled_name = self.module.mapping.get(name) E File "cupy/_core/raw.pyx", line 394, in cupy._core.raw.RawModule.module.__get__ E return self._module() E File "cupy/_core/raw.pyx", line 402, in cupy._core.raw.RawModule._module E mod = _get_raw_module( E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/raw.pyx", line 547, in cupy._core.raw._get_raw_module E mod = cupy._core.core.compile_with_cache( E File "cupy/_core/core.pyx", line 2015, in cupy._core.core.compile_with_cache E cpdef function.Module compile_with_cache( E File "cupy/_core/core.pyx", line 2078, in cupy._core.core.compile_with_cache E return cuda.compile_with_cache( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 453, in compile_with_cache E return _compile_with_cache_hip( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 853, in _compile_with_cache_hip E binary, mapping = compile_using_nvrtc( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 295, in compile_using_nvrtc E return _compile(source, options, cu_path, E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 279, in _compile E compiled_obj, mapping = prog.compile(options, log_stream) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 667, in compile E raise CompileException(log, self.src, self.name, options, E cupy.cuda.compiler.CompileException: /tmp/comgr-0f76bd/input/CompileSource:203:27: error: redeclaration of '__hiprtc_16' with a different type: 'void (*const)(int, int, int, const signed char *, const signed char *, signed char *)' vs 'void (*const)(int, int, int, const bool *, const bool *, bool *)' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-0f76bd/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-0f76bd/input/CompileSource:204:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-0f76bd/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-0f76bd/input/CompileSource:205:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-0f76bd/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-0f76bd/input/CompileSource:206:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-0f76bd/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-0f76bd/input/CompileSource:207:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-0f76bd/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-0f76bd/input/CompileSource:208:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-0f76bd/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-0f76bd/input/CompileSource:209:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-0f76bd/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-0f76bd/input/CompileSource:210:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-0f76bd/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-0f76bd/input/CompileSource:211:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-0f76bd/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-0f76bd/input/CompileSource:212:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-0f76bd/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E 10 errors generated when compiling for gfx908. E Error: Failed to compile opencl source (from CL or HIP source to LLVM IR). cupy/cuda/compiler.py:667: AssertionError ----------------------------- Captured stdout call ----------------------------- {'dtype_a': , 'dtype_b': } _ TestDot_param_26_{shape=((2, 1), (1, 3)), trans_a=False, trans_b=True}.test_dot _ self = < parameter: {'shape': ((2, 1), (1, 3)), 'trans_a': False, 'trans_b': True}> xp = dtype_a = , dtype_b = @testing.for_all_dtypes_combination(['dtype_a', 'dtype_b']) @testing.numpy_cupy_allclose() def test_dot(self, xp, dtype_a, dtype_b): shape_a, shape_b = self.shape if self.trans_a: a = testing.shaped_arange(shape_a[::-1], xp, dtype_a).T else: a = testing.shaped_arange(shape_a, xp, dtype_a) if self.trans_b: b = testing.shaped_arange(shape_b[::-1], xp, dtype_b).T else: b = testing.shaped_arange(shape_b, xp, dtype_b) > return xp.dot(a, b) tests/cupy_tests/linalg_tests/test_product.py:49: _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ cupy/linalg/_product.py:65: in dot return a.dot(b, out) cupy/_core/core.pyx:1607: in cupy._core.core.ndarray.dot return _linalg.dot(self, b, out) cupy/_core/_routines_linalg.pyx:427: in cupy._core._routines_linalg.dot return tensordot_core(a, b, out, n, m, k, ret_shape) cupy/_core/_routines_linalg.pyx:501: in cupy._core._routines_linalg.tensordot_core return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) cupy/_core/_routines_linalg.pyx:327: in cupy._core._routines_linalg._integral_tensordot_core kern = _tensordot_core_int_kernel(config, dtype) cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/_routines_linalg.pyx:302: in cupy._core._routines_linalg._tensordot_core_int_kernel ker = mod.get_function( cupy/_core/raw.pyx:470: in cupy._core.raw.RawModule.get_function mangled_name = self.module.mapping.get(name) cupy/_core/raw.pyx:394: in cupy._core.raw.RawModule.module.__get__ return self._module() cupy/_core/raw.pyx:402: in cupy._core.raw.RawModule._module mod = _get_raw_module( cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/raw.pyx:547: in cupy._core.raw._get_raw_module mod = cupy._core.core.compile_with_cache( cupy/_core/core.pyx:2015: in cupy._core.core.compile_with_cache cpdef function.Module compile_with_cache( cupy/_core/core.pyx:2078: in cupy._core.core.compile_with_cache return cuda.compile_with_cache( cupy/cuda/compiler.py:453: in compile_with_cache return _compile_with_cache_hip( cupy/cuda/compiler.py:853: in _compile_with_cache_hip binary, mapping = compile_using_nvrtc( cupy/cuda/compiler.py:295: in compile_using_nvrtc return _compile(source, options, cu_path, cupy/cuda/compiler.py:279: in _compile compiled_obj, mapping = prog.compile(options, log_stream) _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ self = options = ('--std=c++11', '-I/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/_core/include', '-I/opt/rocm/include', '-fcuda-flush-denormals-to-zero') log_stream = None def compile(self, options=(), log_stream=None): try: if self.name_expressions: for ker in self.name_expressions: nvrtc.addAddNameExpression(self.ptr, ker) nvrtc.compileProgram(self.ptr, options) mapping = None if self.name_expressions: mapping = {} for ker in self.name_expressions: mapping[ker] = nvrtc.getLoweredName(self.ptr, ker) if log_stream is not None: log_stream.write(nvrtc.getProgramLog(self.ptr)) # This is to ensure backwards compatibility with nvrtc if self.method == 'cubin': return nvrtc.getCUBIN(self.ptr), mapping elif self.method == 'ptx': return nvrtc.getPTX(self.ptr), mapping else: raise RuntimeError('Unknown NVRTC compile method') except nvrtc.NVRTCError: log = nvrtc.getProgramLog(self.ptr) > raise CompileException(log, self.src, self.name, options, 'nvrtc' if not runtime.is_hip else 'hiprtc') E AssertionError: Only cupy raises error E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 650, in compile E nvrtc.compileProgram(self.ptr, options) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 136, in cupy_backends.cuda.libs.nvrtc.compileProgram E cpdef compileProgram(intptr_t prog, options): E File "cupy_backends/cuda/libs/nvrtc.pyx", line 148, in cupy_backends.cuda.libs.nvrtc.compileProgram E check_status(status) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 67, in cupy_backends.cuda.libs.nvrtc.check_status E raise NVRTCError(status) E cupy_backends.cuda.libs.nvrtc.NVRTCError: HIPRTC_ERROR_COMPILATION (6) E E During handling of the above exception, another exception occurred: E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/testing/_helper.py", line 47, in _call_func E result = impl(*args, **kw) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/tests/cupy_tests/linalg_tests/test_product.py", line 49, in test_dot E return xp.dot(a, b) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/linalg/_product.py", line 65, in dot E return a.dot(b, out) E File "cupy/_core/core.pyx", line 1607, in cupy._core.core.ndarray.dot E return _linalg.dot(self, b, out) E File "cupy/_core/_routines_linalg.pyx", line 427, in cupy._core._routines_linalg.dot E return tensordot_core(a, b, out, n, m, k, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 501, in cupy._core._routines_linalg.tensordot_core E return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 327, in cupy._core._routines_linalg._integral_tensordot_core E kern = _tensordot_core_int_kernel(config, dtype) E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/_routines_linalg.pyx", line 302, in cupy._core._routines_linalg._tensordot_core_int_kernel E ker = mod.get_function( E File "cupy/_core/raw.pyx", line 470, in cupy._core.raw.RawModule.get_function E mangled_name = self.module.mapping.get(name) E File "cupy/_core/raw.pyx", line 394, in cupy._core.raw.RawModule.module.__get__ E return self._module() E File "cupy/_core/raw.pyx", line 402, in cupy._core.raw.RawModule._module E mod = _get_raw_module( E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/raw.pyx", line 547, in cupy._core.raw._get_raw_module E mod = cupy._core.core.compile_with_cache( E File "cupy/_core/core.pyx", line 2015, in cupy._core.core.compile_with_cache E cpdef function.Module compile_with_cache( E File "cupy/_core/core.pyx", line 2078, in cupy._core.core.compile_with_cache E return cuda.compile_with_cache( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 453, in compile_with_cache E return _compile_with_cache_hip( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 853, in _compile_with_cache_hip E binary, mapping = compile_using_nvrtc( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 295, in compile_using_nvrtc E return _compile(source, options, cu_path, E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 279, in _compile E compiled_obj, mapping = prog.compile(options, log_stream) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 667, in compile E raise CompileException(log, self.src, self.name, options, E cupy.cuda.compiler.CompileException: /tmp/comgr-df32db/input/CompileSource:203:27: error: redeclaration of '__hiprtc_16' with a different type: 'void (*const)(int, int, int, const signed char *, const signed char *, signed char *)' vs 'void (*const)(int, int, int, const bool *, const bool *, bool *)' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-df32db/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-df32db/input/CompileSource:204:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-df32db/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-df32db/input/CompileSource:205:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-df32db/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-df32db/input/CompileSource:206:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-df32db/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-df32db/input/CompileSource:207:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-df32db/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-df32db/input/CompileSource:208:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-df32db/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-df32db/input/CompileSource:209:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-df32db/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-df32db/input/CompileSource:210:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-df32db/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-df32db/input/CompileSource:211:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-df32db/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-df32db/input/CompileSource:212:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-df32db/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E 10 errors generated when compiling for gfx908. E Error: Failed to compile opencl source (from CL or HIP source to LLVM IR). cupy/cuda/compiler.py:667: AssertionError ----------------------------- Captured stdout call ----------------------------- {'dtype_a': , 'dtype_b': } _ TestDot_param_27_{shape=((2, 1), (1, 3)), trans_a=False, trans_b=False}.test_dot _ self = < parameter: {'shape': ((2, 1), (1, 3)), 'trans_a': False, 'trans_b': False}> xp = dtype_a = , dtype_b = @testing.for_all_dtypes_combination(['dtype_a', 'dtype_b']) @testing.numpy_cupy_allclose() def test_dot(self, xp, dtype_a, dtype_b): shape_a, shape_b = self.shape if self.trans_a: a = testing.shaped_arange(shape_a[::-1], xp, dtype_a).T else: a = testing.shaped_arange(shape_a, xp, dtype_a) if self.trans_b: b = testing.shaped_arange(shape_b[::-1], xp, dtype_b).T else: b = testing.shaped_arange(shape_b, xp, dtype_b) > return xp.dot(a, b) tests/cupy_tests/linalg_tests/test_product.py:49: _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ cupy/linalg/_product.py:65: in dot return a.dot(b, out) cupy/_core/core.pyx:1607: in cupy._core.core.ndarray.dot return _linalg.dot(self, b, out) cupy/_core/_routines_linalg.pyx:427: in cupy._core._routines_linalg.dot return tensordot_core(a, b, out, n, m, k, ret_shape) cupy/_core/_routines_linalg.pyx:501: in cupy._core._routines_linalg.tensordot_core return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) cupy/_core/_routines_linalg.pyx:327: in cupy._core._routines_linalg._integral_tensordot_core kern = _tensordot_core_int_kernel(config, dtype) cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/_routines_linalg.pyx:302: in cupy._core._routines_linalg._tensordot_core_int_kernel ker = mod.get_function( cupy/_core/raw.pyx:470: in cupy._core.raw.RawModule.get_function mangled_name = self.module.mapping.get(name) cupy/_core/raw.pyx:394: in cupy._core.raw.RawModule.module.__get__ return self._module() cupy/_core/raw.pyx:402: in cupy._core.raw.RawModule._module mod = _get_raw_module( cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/raw.pyx:547: in cupy._core.raw._get_raw_module mod = cupy._core.core.compile_with_cache( cupy/_core/core.pyx:2015: in cupy._core.core.compile_with_cache cpdef function.Module compile_with_cache( cupy/_core/core.pyx:2078: in cupy._core.core.compile_with_cache return cuda.compile_with_cache( cupy/cuda/compiler.py:453: in compile_with_cache return _compile_with_cache_hip( cupy/cuda/compiler.py:853: in _compile_with_cache_hip binary, mapping = compile_using_nvrtc( cupy/cuda/compiler.py:295: in compile_using_nvrtc return _compile(source, options, cu_path, cupy/cuda/compiler.py:279: in _compile compiled_obj, mapping = prog.compile(options, log_stream) _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ self = options = ('--std=c++11', '-I/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/_core/include', '-I/opt/rocm/include', '-fcuda-flush-denormals-to-zero') log_stream = None def compile(self, options=(), log_stream=None): try: if self.name_expressions: for ker in self.name_expressions: nvrtc.addAddNameExpression(self.ptr, ker) nvrtc.compileProgram(self.ptr, options) mapping = None if self.name_expressions: mapping = {} for ker in self.name_expressions: mapping[ker] = nvrtc.getLoweredName(self.ptr, ker) if log_stream is not None: log_stream.write(nvrtc.getProgramLog(self.ptr)) # This is to ensure backwards compatibility with nvrtc if self.method == 'cubin': return nvrtc.getCUBIN(self.ptr), mapping elif self.method == 'ptx': return nvrtc.getPTX(self.ptr), mapping else: raise RuntimeError('Unknown NVRTC compile method') except nvrtc.NVRTCError: log = nvrtc.getProgramLog(self.ptr) > raise CompileException(log, self.src, self.name, options, 'nvrtc' if not runtime.is_hip else 'hiprtc') E AssertionError: Only cupy raises error E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 650, in compile E nvrtc.compileProgram(self.ptr, options) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 136, in cupy_backends.cuda.libs.nvrtc.compileProgram E cpdef compileProgram(intptr_t prog, options): E File "cupy_backends/cuda/libs/nvrtc.pyx", line 148, in cupy_backends.cuda.libs.nvrtc.compileProgram E check_status(status) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 67, in cupy_backends.cuda.libs.nvrtc.check_status E raise NVRTCError(status) E cupy_backends.cuda.libs.nvrtc.NVRTCError: HIPRTC_ERROR_COMPILATION (6) E E During handling of the above exception, another exception occurred: E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/testing/_helper.py", line 47, in _call_func E result = impl(*args, **kw) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/tests/cupy_tests/linalg_tests/test_product.py", line 49, in test_dot E return xp.dot(a, b) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/linalg/_product.py", line 65, in dot E return a.dot(b, out) E File "cupy/_core/core.pyx", line 1607, in cupy._core.core.ndarray.dot E return _linalg.dot(self, b, out) E File "cupy/_core/_routines_linalg.pyx", line 427, in cupy._core._routines_linalg.dot E return tensordot_core(a, b, out, n, m, k, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 501, in cupy._core._routines_linalg.tensordot_core E return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 327, in cupy._core._routines_linalg._integral_tensordot_core E kern = _tensordot_core_int_kernel(config, dtype) E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/_routines_linalg.pyx", line 302, in cupy._core._routines_linalg._tensordot_core_int_kernel E ker = mod.get_function( E File "cupy/_core/raw.pyx", line 470, in cupy._core.raw.RawModule.get_function E mangled_name = self.module.mapping.get(name) E File "cupy/_core/raw.pyx", line 394, in cupy._core.raw.RawModule.module.__get__ E return self._module() E File "cupy/_core/raw.pyx", line 402, in cupy._core.raw.RawModule._module E mod = _get_raw_module( E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/raw.pyx", line 547, in cupy._core.raw._get_raw_module E mod = cupy._core.core.compile_with_cache( E File "cupy/_core/core.pyx", line 2015, in cupy._core.core.compile_with_cache E cpdef function.Module compile_with_cache( E File "cupy/_core/core.pyx", line 2078, in cupy._core.core.compile_with_cache E return cuda.compile_with_cache( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 453, in compile_with_cache E return _compile_with_cache_hip( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 853, in _compile_with_cache_hip E binary, mapping = compile_using_nvrtc( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 295, in compile_using_nvrtc E return _compile(source, options, cu_path, E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 279, in _compile E compiled_obj, mapping = prog.compile(options, log_stream) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 667, in compile E raise CompileException(log, self.src, self.name, options, E cupy.cuda.compiler.CompileException: /tmp/comgr-addac4/input/CompileSource:203:27: error: redeclaration of '__hiprtc_16' with a different type: 'void (*const)(int, int, int, const signed char *, const signed char *, signed char *)' vs 'void (*const)(int, int, int, const bool *, const bool *, bool *)' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-addac4/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-addac4/input/CompileSource:204:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-addac4/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-addac4/input/CompileSource:205:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-addac4/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-addac4/input/CompileSource:206:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-addac4/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-addac4/input/CompileSource:207:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-addac4/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-addac4/input/CompileSource:208:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-addac4/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-addac4/input/CompileSource:209:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-addac4/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-addac4/input/CompileSource:210:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-addac4/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-addac4/input/CompileSource:211:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-addac4/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-addac4/input/CompileSource:212:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-addac4/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E 10 errors generated when compiling for gfx908. E Error: Failed to compile opencl source (from CL or HIP source to LLVM IR). cupy/cuda/compiler.py:667: AssertionError ----------------------------- Captured stdout call ----------------------------- {'dtype_a': , 'dtype_b': } _ TestDot_param_28_{shape=((2, 3), (3, 1)), trans_a=True, trans_b=True}.test_dot _ self = < parameter: {'shape': ((2, 3), (3, 1)), 'trans_a': True, 'trans_b': True}> xp = dtype_a = , dtype_b = @testing.for_all_dtypes_combination(['dtype_a', 'dtype_b']) @testing.numpy_cupy_allclose() def test_dot(self, xp, dtype_a, dtype_b): shape_a, shape_b = self.shape if self.trans_a: a = testing.shaped_arange(shape_a[::-1], xp, dtype_a).T else: a = testing.shaped_arange(shape_a, xp, dtype_a) if self.trans_b: b = testing.shaped_arange(shape_b[::-1], xp, dtype_b).T else: b = testing.shaped_arange(shape_b, xp, dtype_b) > return xp.dot(a, b) tests/cupy_tests/linalg_tests/test_product.py:49: _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ cupy/linalg/_product.py:65: in dot return a.dot(b, out) cupy/_core/core.pyx:1607: in cupy._core.core.ndarray.dot return _linalg.dot(self, b, out) cupy/_core/_routines_linalg.pyx:427: in cupy._core._routines_linalg.dot return tensordot_core(a, b, out, n, m, k, ret_shape) cupy/_core/_routines_linalg.pyx:501: in cupy._core._routines_linalg.tensordot_core return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) cupy/_core/_routines_linalg.pyx:327: in cupy._core._routines_linalg._integral_tensordot_core kern = _tensordot_core_int_kernel(config, dtype) cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/_routines_linalg.pyx:302: in cupy._core._routines_linalg._tensordot_core_int_kernel ker = mod.get_function( cupy/_core/raw.pyx:470: in cupy._core.raw.RawModule.get_function mangled_name = self.module.mapping.get(name) cupy/_core/raw.pyx:394: in cupy._core.raw.RawModule.module.__get__ return self._module() cupy/_core/raw.pyx:402: in cupy._core.raw.RawModule._module mod = _get_raw_module( cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/raw.pyx:547: in cupy._core.raw._get_raw_module mod = cupy._core.core.compile_with_cache( cupy/_core/core.pyx:2015: in cupy._core.core.compile_with_cache cpdef function.Module compile_with_cache( cupy/_core/core.pyx:2078: in cupy._core.core.compile_with_cache return cuda.compile_with_cache( cupy/cuda/compiler.py:453: in compile_with_cache return _compile_with_cache_hip( cupy/cuda/compiler.py:853: in _compile_with_cache_hip binary, mapping = compile_using_nvrtc( cupy/cuda/compiler.py:295: in compile_using_nvrtc return _compile(source, options, cu_path, cupy/cuda/compiler.py:279: in _compile compiled_obj, mapping = prog.compile(options, log_stream) _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ self = options = ('--std=c++11', '-I/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/_core/include', '-I/opt/rocm/include', '-fcuda-flush-denormals-to-zero') log_stream = None def compile(self, options=(), log_stream=None): try: if self.name_expressions: for ker in self.name_expressions: nvrtc.addAddNameExpression(self.ptr, ker) nvrtc.compileProgram(self.ptr, options) mapping = None if self.name_expressions: mapping = {} for ker in self.name_expressions: mapping[ker] = nvrtc.getLoweredName(self.ptr, ker) if log_stream is not None: log_stream.write(nvrtc.getProgramLog(self.ptr)) # This is to ensure backwards compatibility with nvrtc if self.method == 'cubin': return nvrtc.getCUBIN(self.ptr), mapping elif self.method == 'ptx': return nvrtc.getPTX(self.ptr), mapping else: raise RuntimeError('Unknown NVRTC compile method') except nvrtc.NVRTCError: log = nvrtc.getProgramLog(self.ptr) > raise CompileException(log, self.src, self.name, options, 'nvrtc' if not runtime.is_hip else 'hiprtc') E AssertionError: Only cupy raises error E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 650, in compile E nvrtc.compileProgram(self.ptr, options) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 136, in cupy_backends.cuda.libs.nvrtc.compileProgram E cpdef compileProgram(intptr_t prog, options): E File "cupy_backends/cuda/libs/nvrtc.pyx", line 148, in cupy_backends.cuda.libs.nvrtc.compileProgram E check_status(status) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 67, in cupy_backends.cuda.libs.nvrtc.check_status E raise NVRTCError(status) E cupy_backends.cuda.libs.nvrtc.NVRTCError: HIPRTC_ERROR_COMPILATION (6) E E During handling of the above exception, another exception occurred: E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/testing/_helper.py", line 47, in _call_func E result = impl(*args, **kw) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/tests/cupy_tests/linalg_tests/test_product.py", line 49, in test_dot E return xp.dot(a, b) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/linalg/_product.py", line 65, in dot E return a.dot(b, out) E File "cupy/_core/core.pyx", line 1607, in cupy._core.core.ndarray.dot E return _linalg.dot(self, b, out) E File "cupy/_core/_routines_linalg.pyx", line 427, in cupy._core._routines_linalg.dot E return tensordot_core(a, b, out, n, m, k, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 501, in cupy._core._routines_linalg.tensordot_core E return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 327, in cupy._core._routines_linalg._integral_tensordot_core E kern = _tensordot_core_int_kernel(config, dtype) E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/_routines_linalg.pyx", line 302, in cupy._core._routines_linalg._tensordot_core_int_kernel E ker = mod.get_function( E File "cupy/_core/raw.pyx", line 470, in cupy._core.raw.RawModule.get_function E mangled_name = self.module.mapping.get(name) E File "cupy/_core/raw.pyx", line 394, in cupy._core.raw.RawModule.module.__get__ E return self._module() E File "cupy/_core/raw.pyx", line 402, in cupy._core.raw.RawModule._module E mod = _get_raw_module( E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/raw.pyx", line 547, in cupy._core.raw._get_raw_module E mod = cupy._core.core.compile_with_cache( E File "cupy/_core/core.pyx", line 2015, in cupy._core.core.compile_with_cache E cpdef function.Module compile_with_cache( E File "cupy/_core/core.pyx", line 2078, in cupy._core.core.compile_with_cache E return cuda.compile_with_cache( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 453, in compile_with_cache E return _compile_with_cache_hip( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 853, in _compile_with_cache_hip E binary, mapping = compile_using_nvrtc( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 295, in compile_using_nvrtc E return _compile(source, options, cu_path, E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 279, in _compile E compiled_obj, mapping = prog.compile(options, log_stream) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 667, in compile E raise CompileException(log, self.src, self.name, options, E cupy.cuda.compiler.CompileException: /tmp/comgr-1e5fd8/input/CompileSource:203:27: error: redeclaration of '__hiprtc_16' with a different type: 'void (*const)(int, int, int, const signed char *, const signed char *, signed char *)' vs 'void (*const)(int, int, int, const bool *, const bool *, bool *)' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-1e5fd8/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-1e5fd8/input/CompileSource:204:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-1e5fd8/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-1e5fd8/input/CompileSource:205:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-1e5fd8/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-1e5fd8/input/CompileSource:206:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-1e5fd8/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-1e5fd8/input/CompileSource:207:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-1e5fd8/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-1e5fd8/input/CompileSource:208:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-1e5fd8/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-1e5fd8/input/CompileSource:209:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-1e5fd8/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-1e5fd8/input/CompileSource:210:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-1e5fd8/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-1e5fd8/input/CompileSource:211:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-1e5fd8/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-1e5fd8/input/CompileSource:212:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-1e5fd8/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E 10 errors generated when compiling for gfx908. E Error: Failed to compile opencl source (from CL or HIP source to LLVM IR). cupy/cuda/compiler.py:667: AssertionError ----------------------------- Captured stdout call ----------------------------- {'dtype_a': , 'dtype_b': } _ TestDot_param_29_{shape=((2, 3), (3, 1)), trans_a=True, trans_b=False}.test_dot _ self = < parameter: {'shape': ((2, 3), (3, 1)), 'trans_a': True, 'trans_b': False}> xp = dtype_a = , dtype_b = @testing.for_all_dtypes_combination(['dtype_a', 'dtype_b']) @testing.numpy_cupy_allclose() def test_dot(self, xp, dtype_a, dtype_b): shape_a, shape_b = self.shape if self.trans_a: a = testing.shaped_arange(shape_a[::-1], xp, dtype_a).T else: a = testing.shaped_arange(shape_a, xp, dtype_a) if self.trans_b: b = testing.shaped_arange(shape_b[::-1], xp, dtype_b).T else: b = testing.shaped_arange(shape_b, xp, dtype_b) > return xp.dot(a, b) tests/cupy_tests/linalg_tests/test_product.py:49: _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ cupy/linalg/_product.py:65: in dot return a.dot(b, out) cupy/_core/core.pyx:1607: in cupy._core.core.ndarray.dot return _linalg.dot(self, b, out) cupy/_core/_routines_linalg.pyx:427: in cupy._core._routines_linalg.dot return tensordot_core(a, b, out, n, m, k, ret_shape) cupy/_core/_routines_linalg.pyx:501: in cupy._core._routines_linalg.tensordot_core return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) cupy/_core/_routines_linalg.pyx:327: in cupy._core._routines_linalg._integral_tensordot_core kern = _tensordot_core_int_kernel(config, dtype) cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/_routines_linalg.pyx:302: in cupy._core._routines_linalg._tensordot_core_int_kernel ker = mod.get_function( cupy/_core/raw.pyx:470: in cupy._core.raw.RawModule.get_function mangled_name = self.module.mapping.get(name) cupy/_core/raw.pyx:394: in cupy._core.raw.RawModule.module.__get__ return self._module() cupy/_core/raw.pyx:402: in cupy._core.raw.RawModule._module mod = _get_raw_module( cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/raw.pyx:547: in cupy._core.raw._get_raw_module mod = cupy._core.core.compile_with_cache( cupy/_core/core.pyx:2015: in cupy._core.core.compile_with_cache cpdef function.Module compile_with_cache( cupy/_core/core.pyx:2078: in cupy._core.core.compile_with_cache return cuda.compile_with_cache( cupy/cuda/compiler.py:453: in compile_with_cache return _compile_with_cache_hip( cupy/cuda/compiler.py:853: in _compile_with_cache_hip binary, mapping = compile_using_nvrtc( cupy/cuda/compiler.py:295: in compile_using_nvrtc return _compile(source, options, cu_path, cupy/cuda/compiler.py:279: in _compile compiled_obj, mapping = prog.compile(options, log_stream) _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ self = options = ('--std=c++11', '-I/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/_core/include', '-I/opt/rocm/include', '-fcuda-flush-denormals-to-zero') log_stream = None def compile(self, options=(), log_stream=None): try: if self.name_expressions: for ker in self.name_expressions: nvrtc.addAddNameExpression(self.ptr, ker) nvrtc.compileProgram(self.ptr, options) mapping = None if self.name_expressions: mapping = {} for ker in self.name_expressions: mapping[ker] = nvrtc.getLoweredName(self.ptr, ker) if log_stream is not None: log_stream.write(nvrtc.getProgramLog(self.ptr)) # This is to ensure backwards compatibility with nvrtc if self.method == 'cubin': return nvrtc.getCUBIN(self.ptr), mapping elif self.method == 'ptx': return nvrtc.getPTX(self.ptr), mapping else: raise RuntimeError('Unknown NVRTC compile method') except nvrtc.NVRTCError: log = nvrtc.getProgramLog(self.ptr) > raise CompileException(log, self.src, self.name, options, 'nvrtc' if not runtime.is_hip else 'hiprtc') E AssertionError: Only cupy raises error E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 650, in compile E nvrtc.compileProgram(self.ptr, options) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 136, in cupy_backends.cuda.libs.nvrtc.compileProgram E cpdef compileProgram(intptr_t prog, options): E File "cupy_backends/cuda/libs/nvrtc.pyx", line 148, in cupy_backends.cuda.libs.nvrtc.compileProgram E check_status(status) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 67, in cupy_backends.cuda.libs.nvrtc.check_status E raise NVRTCError(status) E cupy_backends.cuda.libs.nvrtc.NVRTCError: HIPRTC_ERROR_COMPILATION (6) E E During handling of the above exception, another exception occurred: E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/testing/_helper.py", line 47, in _call_func E result = impl(*args, **kw) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/tests/cupy_tests/linalg_tests/test_product.py", line 49, in test_dot E return xp.dot(a, b) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/linalg/_product.py", line 65, in dot E return a.dot(b, out) E File "cupy/_core/core.pyx", line 1607, in cupy._core.core.ndarray.dot E return _linalg.dot(self, b, out) E File "cupy/_core/_routines_linalg.pyx", line 427, in cupy._core._routines_linalg.dot E return tensordot_core(a, b, out, n, m, k, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 501, in cupy._core._routines_linalg.tensordot_core E return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 327, in cupy._core._routines_linalg._integral_tensordot_core E kern = _tensordot_core_int_kernel(config, dtype) E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/_routines_linalg.pyx", line 302, in cupy._core._routines_linalg._tensordot_core_int_kernel E ker = mod.get_function( E File "cupy/_core/raw.pyx", line 470, in cupy._core.raw.RawModule.get_function E mangled_name = self.module.mapping.get(name) E File "cupy/_core/raw.pyx", line 394, in cupy._core.raw.RawModule.module.__get__ E return self._module() E File "cupy/_core/raw.pyx", line 402, in cupy._core.raw.RawModule._module E mod = _get_raw_module( E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/raw.pyx", line 547, in cupy._core.raw._get_raw_module E mod = cupy._core.core.compile_with_cache( E File "cupy/_core/core.pyx", line 2015, in cupy._core.core.compile_with_cache E cpdef function.Module compile_with_cache( E File "cupy/_core/core.pyx", line 2078, in cupy._core.core.compile_with_cache E return cuda.compile_with_cache( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 453, in compile_with_cache E return _compile_with_cache_hip( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 853, in _compile_with_cache_hip E binary, mapping = compile_using_nvrtc( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 295, in compile_using_nvrtc E return _compile(source, options, cu_path, E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 279, in _compile E compiled_obj, mapping = prog.compile(options, log_stream) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 667, in compile E raise CompileException(log, self.src, self.name, options, E cupy.cuda.compiler.CompileException: /tmp/comgr-1afc7c/input/CompileSource:203:27: error: redeclaration of '__hiprtc_16' with a different type: 'void (*const)(int, int, int, const signed char *, const signed char *, signed char *)' vs 'void (*const)(int, int, int, const bool *, const bool *, bool *)' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-1afc7c/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-1afc7c/input/CompileSource:204:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-1afc7c/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-1afc7c/input/CompileSource:205:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-1afc7c/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-1afc7c/input/CompileSource:206:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-1afc7c/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-1afc7c/input/CompileSource:207:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-1afc7c/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-1afc7c/input/CompileSource:208:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-1afc7c/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-1afc7c/input/CompileSource:209:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-1afc7c/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-1afc7c/input/CompileSource:210:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-1afc7c/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-1afc7c/input/CompileSource:211:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-1afc7c/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-1afc7c/input/CompileSource:212:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-1afc7c/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E 10 errors generated when compiling for gfx908. E Error: Failed to compile opencl source (from CL or HIP source to LLVM IR). cupy/cuda/compiler.py:667: AssertionError ----------------------------- Captured stdout call ----------------------------- {'dtype_a': , 'dtype_b': } _ TestDot_param_30_{shape=((2, 3), (3, 1)), trans_a=False, trans_b=True}.test_dot _ self = < parameter: {'shape': ((2, 3), (3, 1)), 'trans_a': False, 'trans_b': True}> xp = dtype_a = , dtype_b = @testing.for_all_dtypes_combination(['dtype_a', 'dtype_b']) @testing.numpy_cupy_allclose() def test_dot(self, xp, dtype_a, dtype_b): shape_a, shape_b = self.shape if self.trans_a: a = testing.shaped_arange(shape_a[::-1], xp, dtype_a).T else: a = testing.shaped_arange(shape_a, xp, dtype_a) if self.trans_b: b = testing.shaped_arange(shape_b[::-1], xp, dtype_b).T else: b = testing.shaped_arange(shape_b, xp, dtype_b) > return xp.dot(a, b) tests/cupy_tests/linalg_tests/test_product.py:49: _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ cupy/linalg/_product.py:65: in dot return a.dot(b, out) cupy/_core/core.pyx:1607: in cupy._core.core.ndarray.dot return _linalg.dot(self, b, out) cupy/_core/_routines_linalg.pyx:427: in cupy._core._routines_linalg.dot return tensordot_core(a, b, out, n, m, k, ret_shape) cupy/_core/_routines_linalg.pyx:501: in cupy._core._routines_linalg.tensordot_core return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) cupy/_core/_routines_linalg.pyx:327: in cupy._core._routines_linalg._integral_tensordot_core kern = _tensordot_core_int_kernel(config, dtype) cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/_routines_linalg.pyx:302: in cupy._core._routines_linalg._tensordot_core_int_kernel ker = mod.get_function( cupy/_core/raw.pyx:470: in cupy._core.raw.RawModule.get_function mangled_name = self.module.mapping.get(name) cupy/_core/raw.pyx:394: in cupy._core.raw.RawModule.module.__get__ return self._module() cupy/_core/raw.pyx:402: in cupy._core.raw.RawModule._module mod = _get_raw_module( cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/raw.pyx:547: in cupy._core.raw._get_raw_module mod = cupy._core.core.compile_with_cache( cupy/_core/core.pyx:2015: in cupy._core.core.compile_with_cache cpdef function.Module compile_with_cache( cupy/_core/core.pyx:2078: in cupy._core.core.compile_with_cache return cuda.compile_with_cache( cupy/cuda/compiler.py:453: in compile_with_cache return _compile_with_cache_hip( cupy/cuda/compiler.py:853: in _compile_with_cache_hip binary, mapping = compile_using_nvrtc( cupy/cuda/compiler.py:295: in compile_using_nvrtc return _compile(source, options, cu_path, cupy/cuda/compiler.py:279: in _compile compiled_obj, mapping = prog.compile(options, log_stream) _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ self = options = ('--std=c++11', '-I/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/_core/include', '-I/opt/rocm/include', '-fcuda-flush-denormals-to-zero') log_stream = None def compile(self, options=(), log_stream=None): try: if self.name_expressions: for ker in self.name_expressions: nvrtc.addAddNameExpression(self.ptr, ker) nvrtc.compileProgram(self.ptr, options) mapping = None if self.name_expressions: mapping = {} for ker in self.name_expressions: mapping[ker] = nvrtc.getLoweredName(self.ptr, ker) if log_stream is not None: log_stream.write(nvrtc.getProgramLog(self.ptr)) # This is to ensure backwards compatibility with nvrtc if self.method == 'cubin': return nvrtc.getCUBIN(self.ptr), mapping elif self.method == 'ptx': return nvrtc.getPTX(self.ptr), mapping else: raise RuntimeError('Unknown NVRTC compile method') except nvrtc.NVRTCError: log = nvrtc.getProgramLog(self.ptr) > raise CompileException(log, self.src, self.name, options, 'nvrtc' if not runtime.is_hip else 'hiprtc') E AssertionError: Only cupy raises error E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 650, in compile E nvrtc.compileProgram(self.ptr, options) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 136, in cupy_backends.cuda.libs.nvrtc.compileProgram E cpdef compileProgram(intptr_t prog, options): E File "cupy_backends/cuda/libs/nvrtc.pyx", line 148, in cupy_backends.cuda.libs.nvrtc.compileProgram E check_status(status) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 67, in cupy_backends.cuda.libs.nvrtc.check_status E raise NVRTCError(status) E cupy_backends.cuda.libs.nvrtc.NVRTCError: HIPRTC_ERROR_COMPILATION (6) E E During handling of the above exception, another exception occurred: E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/testing/_helper.py", line 47, in _call_func E result = impl(*args, **kw) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/tests/cupy_tests/linalg_tests/test_product.py", line 49, in test_dot E return xp.dot(a, b) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/linalg/_product.py", line 65, in dot E return a.dot(b, out) E File "cupy/_core/core.pyx", line 1607, in cupy._core.core.ndarray.dot E return _linalg.dot(self, b, out) E File "cupy/_core/_routines_linalg.pyx", line 427, in cupy._core._routines_linalg.dot E return tensordot_core(a, b, out, n, m, k, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 501, in cupy._core._routines_linalg.tensordot_core E return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 327, in cupy._core._routines_linalg._integral_tensordot_core E kern = _tensordot_core_int_kernel(config, dtype) E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/_routines_linalg.pyx", line 302, in cupy._core._routines_linalg._tensordot_core_int_kernel E ker = mod.get_function( E File "cupy/_core/raw.pyx", line 470, in cupy._core.raw.RawModule.get_function E mangled_name = self.module.mapping.get(name) E File "cupy/_core/raw.pyx", line 394, in cupy._core.raw.RawModule.module.__get__ E return self._module() E File "cupy/_core/raw.pyx", line 402, in cupy._core.raw.RawModule._module E mod = _get_raw_module( E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/raw.pyx", line 547, in cupy._core.raw._get_raw_module E mod = cupy._core.core.compile_with_cache( E File "cupy/_core/core.pyx", line 2015, in cupy._core.core.compile_with_cache E cpdef function.Module compile_with_cache( E File "cupy/_core/core.pyx", line 2078, in cupy._core.core.compile_with_cache E return cuda.compile_with_cache( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 453, in compile_with_cache E return _compile_with_cache_hip( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 853, in _compile_with_cache_hip E binary, mapping = compile_using_nvrtc( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 295, in compile_using_nvrtc E return _compile(source, options, cu_path, E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 279, in _compile E compiled_obj, mapping = prog.compile(options, log_stream) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 667, in compile E raise CompileException(log, self.src, self.name, options, E cupy.cuda.compiler.CompileException: /tmp/comgr-76684b/input/CompileSource:203:27: error: redeclaration of '__hiprtc_16' with a different type: 'void (*const)(int, int, int, const signed char *, const signed char *, signed char *)' vs 'void (*const)(int, int, int, const bool *, const bool *, bool *)' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-76684b/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-76684b/input/CompileSource:204:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-76684b/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-76684b/input/CompileSource:205:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-76684b/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-76684b/input/CompileSource:206:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-76684b/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-76684b/input/CompileSource:207:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-76684b/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-76684b/input/CompileSource:208:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-76684b/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-76684b/input/CompileSource:209:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-76684b/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-76684b/input/CompileSource:210:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-76684b/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-76684b/input/CompileSource:211:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-76684b/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-76684b/input/CompileSource:212:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-76684b/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E 10 errors generated when compiling for gfx908. E Error: Failed to compile opencl source (from CL or HIP source to LLVM IR). cupy/cuda/compiler.py:667: AssertionError ----------------------------- Captured stdout call ----------------------------- {'dtype_a': , 'dtype_b': } _ TestDot_param_31_{shape=((2, 3), (3, 1)), trans_a=False, trans_b=False}.test_dot _ self = < parameter: {'shape': ((2, 3), (3, 1)), 'trans_a': False, 'trans_b': False}> xp = dtype_a = , dtype_b = @testing.for_all_dtypes_combination(['dtype_a', 'dtype_b']) @testing.numpy_cupy_allclose() def test_dot(self, xp, dtype_a, dtype_b): shape_a, shape_b = self.shape if self.trans_a: a = testing.shaped_arange(shape_a[::-1], xp, dtype_a).T else: a = testing.shaped_arange(shape_a, xp, dtype_a) if self.trans_b: b = testing.shaped_arange(shape_b[::-1], xp, dtype_b).T else: b = testing.shaped_arange(shape_b, xp, dtype_b) > return xp.dot(a, b) tests/cupy_tests/linalg_tests/test_product.py:49: _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ cupy/linalg/_product.py:65: in dot return a.dot(b, out) cupy/_core/core.pyx:1607: in cupy._core.core.ndarray.dot return _linalg.dot(self, b, out) cupy/_core/_routines_linalg.pyx:427: in cupy._core._routines_linalg.dot return tensordot_core(a, b, out, n, m, k, ret_shape) cupy/_core/_routines_linalg.pyx:501: in cupy._core._routines_linalg.tensordot_core return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) cupy/_core/_routines_linalg.pyx:327: in cupy._core._routines_linalg._integral_tensordot_core kern = _tensordot_core_int_kernel(config, dtype) cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/_routines_linalg.pyx:302: in cupy._core._routines_linalg._tensordot_core_int_kernel ker = mod.get_function( cupy/_core/raw.pyx:470: in cupy._core.raw.RawModule.get_function mangled_name = self.module.mapping.get(name) cupy/_core/raw.pyx:394: in cupy._core.raw.RawModule.module.__get__ return self._module() cupy/_core/raw.pyx:402: in cupy._core.raw.RawModule._module mod = _get_raw_module( cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/raw.pyx:547: in cupy._core.raw._get_raw_module mod = cupy._core.core.compile_with_cache( cupy/_core/core.pyx:2015: in cupy._core.core.compile_with_cache cpdef function.Module compile_with_cache( cupy/_core/core.pyx:2078: in cupy._core.core.compile_with_cache return cuda.compile_with_cache( cupy/cuda/compiler.py:453: in compile_with_cache return _compile_with_cache_hip( cupy/cuda/compiler.py:853: in _compile_with_cache_hip binary, mapping = compile_using_nvrtc( cupy/cuda/compiler.py:295: in compile_using_nvrtc return _compile(source, options, cu_path, cupy/cuda/compiler.py:279: in _compile compiled_obj, mapping = prog.compile(options, log_stream) _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ self = options = ('--std=c++11', '-I/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/_core/include', '-I/opt/rocm/include', '-fcuda-flush-denormals-to-zero') log_stream = None def compile(self, options=(), log_stream=None): try: if self.name_expressions: for ker in self.name_expressions: nvrtc.addAddNameExpression(self.ptr, ker) nvrtc.compileProgram(self.ptr, options) mapping = None if self.name_expressions: mapping = {} for ker in self.name_expressions: mapping[ker] = nvrtc.getLoweredName(self.ptr, ker) if log_stream is not None: log_stream.write(nvrtc.getProgramLog(self.ptr)) # This is to ensure backwards compatibility with nvrtc if self.method == 'cubin': return nvrtc.getCUBIN(self.ptr), mapping elif self.method == 'ptx': return nvrtc.getPTX(self.ptr), mapping else: raise RuntimeError('Unknown NVRTC compile method') except nvrtc.NVRTCError: log = nvrtc.getProgramLog(self.ptr) > raise CompileException(log, self.src, self.name, options, 'nvrtc' if not runtime.is_hip else 'hiprtc') E AssertionError: Only cupy raises error E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 650, in compile E nvrtc.compileProgram(self.ptr, options) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 136, in cupy_backends.cuda.libs.nvrtc.compileProgram E cpdef compileProgram(intptr_t prog, options): E File "cupy_backends/cuda/libs/nvrtc.pyx", line 148, in cupy_backends.cuda.libs.nvrtc.compileProgram E check_status(status) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 67, in cupy_backends.cuda.libs.nvrtc.check_status E raise NVRTCError(status) E cupy_backends.cuda.libs.nvrtc.NVRTCError: HIPRTC_ERROR_COMPILATION (6) E E During handling of the above exception, another exception occurred: E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/testing/_helper.py", line 47, in _call_func E result = impl(*args, **kw) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/tests/cupy_tests/linalg_tests/test_product.py", line 49, in test_dot E return xp.dot(a, b) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/linalg/_product.py", line 65, in dot E return a.dot(b, out) E File "cupy/_core/core.pyx", line 1607, in cupy._core.core.ndarray.dot E return _linalg.dot(self, b, out) E File "cupy/_core/_routines_linalg.pyx", line 427, in cupy._core._routines_linalg.dot E return tensordot_core(a, b, out, n, m, k, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 501, in cupy._core._routines_linalg.tensordot_core E return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 327, in cupy._core._routines_linalg._integral_tensordot_core E kern = _tensordot_core_int_kernel(config, dtype) E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/_routines_linalg.pyx", line 302, in cupy._core._routines_linalg._tensordot_core_int_kernel E ker = mod.get_function( E File "cupy/_core/raw.pyx", line 470, in cupy._core.raw.RawModule.get_function E mangled_name = self.module.mapping.get(name) E File "cupy/_core/raw.pyx", line 394, in cupy._core.raw.RawModule.module.__get__ E return self._module() E File "cupy/_core/raw.pyx", line 402, in cupy._core.raw.RawModule._module E mod = _get_raw_module( E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/raw.pyx", line 547, in cupy._core.raw._get_raw_module E mod = cupy._core.core.compile_with_cache( E File "cupy/_core/core.pyx", line 2015, in cupy._core.core.compile_with_cache E cpdef function.Module compile_with_cache( E File "cupy/_core/core.pyx", line 2078, in cupy._core.core.compile_with_cache E return cuda.compile_with_cache( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 453, in compile_with_cache E return _compile_with_cache_hip( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 853, in _compile_with_cache_hip E binary, mapping = compile_using_nvrtc( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 295, in compile_using_nvrtc E return _compile(source, options, cu_path, E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 279, in _compile E compiled_obj, mapping = prog.compile(options, log_stream) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 667, in compile E raise CompileException(log, self.src, self.name, options, E cupy.cuda.compiler.CompileException: /tmp/comgr-2d7376/input/CompileSource:203:27: error: redeclaration of '__hiprtc_16' with a different type: 'void (*const)(int, int, int, const signed char *, const signed char *, signed char *)' vs 'void (*const)(int, int, int, const bool *, const bool *, bool *)' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-2d7376/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-2d7376/input/CompileSource:204:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-2d7376/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-2d7376/input/CompileSource:205:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-2d7376/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-2d7376/input/CompileSource:206:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-2d7376/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-2d7376/input/CompileSource:207:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-2d7376/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-2d7376/input/CompileSource:208:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-2d7376/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-2d7376/input/CompileSource:209:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-2d7376/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-2d7376/input/CompileSource:210:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-2d7376/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-2d7376/input/CompileSource:211:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-2d7376/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-2d7376/input/CompileSource:212:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-2d7376/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E 10 errors generated when compiling for gfx908. E Error: Failed to compile opencl source (from CL or HIP source to LLVM IR). cupy/cuda/compiler.py:667: AssertionError ----------------------------- Captured stdout call ----------------------------- {'dtype_a': , 'dtype_b': } _ TestDot_param_32_{shape=((2, 3), (3, 4)), trans_a=True, trans_b=True}.test_dot _ self = < parameter: {'shape': ((2, 3), (3, 4)), 'trans_a': True, 'trans_b': True}> xp = dtype_a = , dtype_b = @testing.for_all_dtypes_combination(['dtype_a', 'dtype_b']) @testing.numpy_cupy_allclose() def test_dot(self, xp, dtype_a, dtype_b): shape_a, shape_b = self.shape if self.trans_a: a = testing.shaped_arange(shape_a[::-1], xp, dtype_a).T else: a = testing.shaped_arange(shape_a, xp, dtype_a) if self.trans_b: b = testing.shaped_arange(shape_b[::-1], xp, dtype_b).T else: b = testing.shaped_arange(shape_b, xp, dtype_b) > return xp.dot(a, b) tests/cupy_tests/linalg_tests/test_product.py:49: _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ cupy/linalg/_product.py:65: in dot return a.dot(b, out) cupy/_core/core.pyx:1607: in cupy._core.core.ndarray.dot return _linalg.dot(self, b, out) cupy/_core/_routines_linalg.pyx:427: in cupy._core._routines_linalg.dot return tensordot_core(a, b, out, n, m, k, ret_shape) cupy/_core/_routines_linalg.pyx:501: in cupy._core._routines_linalg.tensordot_core return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) cupy/_core/_routines_linalg.pyx:327: in cupy._core._routines_linalg._integral_tensordot_core kern = _tensordot_core_int_kernel(config, dtype) cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/_routines_linalg.pyx:302: in cupy._core._routines_linalg._tensordot_core_int_kernel ker = mod.get_function( cupy/_core/raw.pyx:470: in cupy._core.raw.RawModule.get_function mangled_name = self.module.mapping.get(name) cupy/_core/raw.pyx:394: in cupy._core.raw.RawModule.module.__get__ return self._module() cupy/_core/raw.pyx:402: in cupy._core.raw.RawModule._module mod = _get_raw_module( cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/raw.pyx:547: in cupy._core.raw._get_raw_module mod = cupy._core.core.compile_with_cache( cupy/_core/core.pyx:2015: in cupy._core.core.compile_with_cache cpdef function.Module compile_with_cache( cupy/_core/core.pyx:2078: in cupy._core.core.compile_with_cache return cuda.compile_with_cache( cupy/cuda/compiler.py:453: in compile_with_cache return _compile_with_cache_hip( cupy/cuda/compiler.py:853: in _compile_with_cache_hip binary, mapping = compile_using_nvrtc( cupy/cuda/compiler.py:295: in compile_using_nvrtc return _compile(source, options, cu_path, cupy/cuda/compiler.py:279: in _compile compiled_obj, mapping = prog.compile(options, log_stream) _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ self = options = ('--std=c++11', '-I/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/_core/include', '-I/opt/rocm/include', '-fcuda-flush-denormals-to-zero') log_stream = None def compile(self, options=(), log_stream=None): try: if self.name_expressions: for ker in self.name_expressions: nvrtc.addAddNameExpression(self.ptr, ker) nvrtc.compileProgram(self.ptr, options) mapping = None if self.name_expressions: mapping = {} for ker in self.name_expressions: mapping[ker] = nvrtc.getLoweredName(self.ptr, ker) if log_stream is not None: log_stream.write(nvrtc.getProgramLog(self.ptr)) # This is to ensure backwards compatibility with nvrtc if self.method == 'cubin': return nvrtc.getCUBIN(self.ptr), mapping elif self.method == 'ptx': return nvrtc.getPTX(self.ptr), mapping else: raise RuntimeError('Unknown NVRTC compile method') except nvrtc.NVRTCError: log = nvrtc.getProgramLog(self.ptr) > raise CompileException(log, self.src, self.name, options, 'nvrtc' if not runtime.is_hip else 'hiprtc') E AssertionError: Only cupy raises error E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 650, in compile E nvrtc.compileProgram(self.ptr, options) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 136, in cupy_backends.cuda.libs.nvrtc.compileProgram E cpdef compileProgram(intptr_t prog, options): E File "cupy_backends/cuda/libs/nvrtc.pyx", line 148, in cupy_backends.cuda.libs.nvrtc.compileProgram E check_status(status) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 67, in cupy_backends.cuda.libs.nvrtc.check_status E raise NVRTCError(status) E cupy_backends.cuda.libs.nvrtc.NVRTCError: HIPRTC_ERROR_COMPILATION (6) E E During handling of the above exception, another exception occurred: E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/testing/_helper.py", line 47, in _call_func E result = impl(*args, **kw) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/tests/cupy_tests/linalg_tests/test_product.py", line 49, in test_dot E return xp.dot(a, b) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/linalg/_product.py", line 65, in dot E return a.dot(b, out) E File "cupy/_core/core.pyx", line 1607, in cupy._core.core.ndarray.dot E return _linalg.dot(self, b, out) E File "cupy/_core/_routines_linalg.pyx", line 427, in cupy._core._routines_linalg.dot E return tensordot_core(a, b, out, n, m, k, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 501, in cupy._core._routines_linalg.tensordot_core E return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 327, in cupy._core._routines_linalg._integral_tensordot_core E kern = _tensordot_core_int_kernel(config, dtype) E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/_routines_linalg.pyx", line 302, in cupy._core._routines_linalg._tensordot_core_int_kernel E ker = mod.get_function( E File "cupy/_core/raw.pyx", line 470, in cupy._core.raw.RawModule.get_function E mangled_name = self.module.mapping.get(name) E File "cupy/_core/raw.pyx", line 394, in cupy._core.raw.RawModule.module.__get__ E return self._module() E File "cupy/_core/raw.pyx", line 402, in cupy._core.raw.RawModule._module E mod = _get_raw_module( E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/raw.pyx", line 547, in cupy._core.raw._get_raw_module E mod = cupy._core.core.compile_with_cache( E File "cupy/_core/core.pyx", line 2015, in cupy._core.core.compile_with_cache E cpdef function.Module compile_with_cache( E File "cupy/_core/core.pyx", line 2078, in cupy._core.core.compile_with_cache E return cuda.compile_with_cache( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 453, in compile_with_cache E return _compile_with_cache_hip( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 853, in _compile_with_cache_hip E binary, mapping = compile_using_nvrtc( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 295, in compile_using_nvrtc E return _compile(source, options, cu_path, E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 279, in _compile E compiled_obj, mapping = prog.compile(options, log_stream) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 667, in compile E raise CompileException(log, self.src, self.name, options, E cupy.cuda.compiler.CompileException: /tmp/comgr-3f5a5b/input/CompileSource:203:27: error: redeclaration of '__hiprtc_16' with a different type: 'void (*const)(int, int, int, const signed char *, const signed char *, signed char *)' vs 'void (*const)(int, int, int, const bool *, const bool *, bool *)' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-3f5a5b/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-3f5a5b/input/CompileSource:204:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-3f5a5b/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-3f5a5b/input/CompileSource:205:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-3f5a5b/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-3f5a5b/input/CompileSource:206:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-3f5a5b/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-3f5a5b/input/CompileSource:207:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-3f5a5b/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-3f5a5b/input/CompileSource:208:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-3f5a5b/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-3f5a5b/input/CompileSource:209:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-3f5a5b/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-3f5a5b/input/CompileSource:210:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-3f5a5b/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-3f5a5b/input/CompileSource:211:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-3f5a5b/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-3f5a5b/input/CompileSource:212:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-3f5a5b/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E 10 errors generated when compiling for gfx908. E Error: Failed to compile opencl source (from CL or HIP source to LLVM IR). cupy/cuda/compiler.py:667: AssertionError ----------------------------- Captured stdout call ----------------------------- {'dtype_a': , 'dtype_b': } _ TestDot_param_33_{shape=((2, 3), (3, 4)), trans_a=True, trans_b=False}.test_dot _ self = < parameter: {'shape': ((2, 3), (3, 4)), 'trans_a': True, 'trans_b': False}> xp = dtype_a = , dtype_b = @testing.for_all_dtypes_combination(['dtype_a', 'dtype_b']) @testing.numpy_cupy_allclose() def test_dot(self, xp, dtype_a, dtype_b): shape_a, shape_b = self.shape if self.trans_a: a = testing.shaped_arange(shape_a[::-1], xp, dtype_a).T else: a = testing.shaped_arange(shape_a, xp, dtype_a) if self.trans_b: b = testing.shaped_arange(shape_b[::-1], xp, dtype_b).T else: b = testing.shaped_arange(shape_b, xp, dtype_b) > return xp.dot(a, b) tests/cupy_tests/linalg_tests/test_product.py:49: _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ cupy/linalg/_product.py:65: in dot return a.dot(b, out) cupy/_core/core.pyx:1607: in cupy._core.core.ndarray.dot return _linalg.dot(self, b, out) cupy/_core/_routines_linalg.pyx:427: in cupy._core._routines_linalg.dot return tensordot_core(a, b, out, n, m, k, ret_shape) cupy/_core/_routines_linalg.pyx:501: in cupy._core._routines_linalg.tensordot_core return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) cupy/_core/_routines_linalg.pyx:327: in cupy._core._routines_linalg._integral_tensordot_core kern = _tensordot_core_int_kernel(config, dtype) cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/_routines_linalg.pyx:302: in cupy._core._routines_linalg._tensordot_core_int_kernel ker = mod.get_function( cupy/_core/raw.pyx:470: in cupy._core.raw.RawModule.get_function mangled_name = self.module.mapping.get(name) cupy/_core/raw.pyx:394: in cupy._core.raw.RawModule.module.__get__ return self._module() cupy/_core/raw.pyx:402: in cupy._core.raw.RawModule._module mod = _get_raw_module( cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/raw.pyx:547: in cupy._core.raw._get_raw_module mod = cupy._core.core.compile_with_cache( cupy/_core/core.pyx:2015: in cupy._core.core.compile_with_cache cpdef function.Module compile_with_cache( cupy/_core/core.pyx:2078: in cupy._core.core.compile_with_cache return cuda.compile_with_cache( cupy/cuda/compiler.py:453: in compile_with_cache return _compile_with_cache_hip( cupy/cuda/compiler.py:853: in _compile_with_cache_hip binary, mapping = compile_using_nvrtc( cupy/cuda/compiler.py:295: in compile_using_nvrtc return _compile(source, options, cu_path, cupy/cuda/compiler.py:279: in _compile compiled_obj, mapping = prog.compile(options, log_stream) _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ self = options = ('--std=c++11', '-I/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/_core/include', '-I/opt/rocm/include', '-fcuda-flush-denormals-to-zero') log_stream = None def compile(self, options=(), log_stream=None): try: if self.name_expressions: for ker in self.name_expressions: nvrtc.addAddNameExpression(self.ptr, ker) nvrtc.compileProgram(self.ptr, options) mapping = None if self.name_expressions: mapping = {} for ker in self.name_expressions: mapping[ker] = nvrtc.getLoweredName(self.ptr, ker) if log_stream is not None: log_stream.write(nvrtc.getProgramLog(self.ptr)) # This is to ensure backwards compatibility with nvrtc if self.method == 'cubin': return nvrtc.getCUBIN(self.ptr), mapping elif self.method == 'ptx': return nvrtc.getPTX(self.ptr), mapping else: raise RuntimeError('Unknown NVRTC compile method') except nvrtc.NVRTCError: log = nvrtc.getProgramLog(self.ptr) > raise CompileException(log, self.src, self.name, options, 'nvrtc' if not runtime.is_hip else 'hiprtc') E AssertionError: Only cupy raises error E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 650, in compile E nvrtc.compileProgram(self.ptr, options) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 136, in cupy_backends.cuda.libs.nvrtc.compileProgram E cpdef compileProgram(intptr_t prog, options): E File "cupy_backends/cuda/libs/nvrtc.pyx", line 148, in cupy_backends.cuda.libs.nvrtc.compileProgram E check_status(status) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 67, in cupy_backends.cuda.libs.nvrtc.check_status E raise NVRTCError(status) E cupy_backends.cuda.libs.nvrtc.NVRTCError: HIPRTC_ERROR_COMPILATION (6) E E During handling of the above exception, another exception occurred: E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/testing/_helper.py", line 47, in _call_func E result = impl(*args, **kw) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/tests/cupy_tests/linalg_tests/test_product.py", line 49, in test_dot E return xp.dot(a, b) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/linalg/_product.py", line 65, in dot E return a.dot(b, out) E File "cupy/_core/core.pyx", line 1607, in cupy._core.core.ndarray.dot E return _linalg.dot(self, b, out) E File "cupy/_core/_routines_linalg.pyx", line 427, in cupy._core._routines_linalg.dot E return tensordot_core(a, b, out, n, m, k, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 501, in cupy._core._routines_linalg.tensordot_core E return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 327, in cupy._core._routines_linalg._integral_tensordot_core E kern = _tensordot_core_int_kernel(config, dtype) E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/_routines_linalg.pyx", line 302, in cupy._core._routines_linalg._tensordot_core_int_kernel E ker = mod.get_function( E File "cupy/_core/raw.pyx", line 470, in cupy._core.raw.RawModule.get_function E mangled_name = self.module.mapping.get(name) E File "cupy/_core/raw.pyx", line 394, in cupy._core.raw.RawModule.module.__get__ E return self._module() E File "cupy/_core/raw.pyx", line 402, in cupy._core.raw.RawModule._module E mod = _get_raw_module( E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/raw.pyx", line 547, in cupy._core.raw._get_raw_module E mod = cupy._core.core.compile_with_cache( E File "cupy/_core/core.pyx", line 2015, in cupy._core.core.compile_with_cache E cpdef function.Module compile_with_cache( E File "cupy/_core/core.pyx", line 2078, in cupy._core.core.compile_with_cache E return cuda.compile_with_cache( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 453, in compile_with_cache E return _compile_with_cache_hip( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 853, in _compile_with_cache_hip E binary, mapping = compile_using_nvrtc( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 295, in compile_using_nvrtc E return _compile(source, options, cu_path, E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 279, in _compile E compiled_obj, mapping = prog.compile(options, log_stream) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 667, in compile E raise CompileException(log, self.src, self.name, options, E cupy.cuda.compiler.CompileException: /tmp/comgr-19e365/input/CompileSource:203:27: error: redeclaration of '__hiprtc_16' with a different type: 'void (*const)(int, int, int, const signed char *, const signed char *, signed char *)' vs 'void (*const)(int, int, int, const bool *, const bool *, bool *)' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-19e365/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-19e365/input/CompileSource:204:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-19e365/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-19e365/input/CompileSource:205:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-19e365/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-19e365/input/CompileSource:206:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-19e365/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-19e365/input/CompileSource:207:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-19e365/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-19e365/input/CompileSource:208:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-19e365/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-19e365/input/CompileSource:209:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-19e365/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-19e365/input/CompileSource:210:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-19e365/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-19e365/input/CompileSource:211:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-19e365/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-19e365/input/CompileSource:212:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-19e365/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E 10 errors generated when compiling for gfx908. E Error: Failed to compile opencl source (from CL or HIP source to LLVM IR). cupy/cuda/compiler.py:667: AssertionError ----------------------------- Captured stdout call ----------------------------- {'dtype_a': , 'dtype_b': } _ TestDot_param_34_{shape=((2, 3), (3, 4)), trans_a=False, trans_b=True}.test_dot _ self = < parameter: {'shape': ((2, 3), (3, 4)), 'trans_a': False, 'trans_b': True}> xp = dtype_a = , dtype_b = @testing.for_all_dtypes_combination(['dtype_a', 'dtype_b']) @testing.numpy_cupy_allclose() def test_dot(self, xp, dtype_a, dtype_b): shape_a, shape_b = self.shape if self.trans_a: a = testing.shaped_arange(shape_a[::-1], xp, dtype_a).T else: a = testing.shaped_arange(shape_a, xp, dtype_a) if self.trans_b: b = testing.shaped_arange(shape_b[::-1], xp, dtype_b).T else: b = testing.shaped_arange(shape_b, xp, dtype_b) > return xp.dot(a, b) tests/cupy_tests/linalg_tests/test_product.py:49: _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ cupy/linalg/_product.py:65: in dot return a.dot(b, out) cupy/_core/core.pyx:1607: in cupy._core.core.ndarray.dot return _linalg.dot(self, b, out) cupy/_core/_routines_linalg.pyx:427: in cupy._core._routines_linalg.dot return tensordot_core(a, b, out, n, m, k, ret_shape) cupy/_core/_routines_linalg.pyx:501: in cupy._core._routines_linalg.tensordot_core return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) cupy/_core/_routines_linalg.pyx:327: in cupy._core._routines_linalg._integral_tensordot_core kern = _tensordot_core_int_kernel(config, dtype) cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/_routines_linalg.pyx:302: in cupy._core._routines_linalg._tensordot_core_int_kernel ker = mod.get_function( cupy/_core/raw.pyx:470: in cupy._core.raw.RawModule.get_function mangled_name = self.module.mapping.get(name) cupy/_core/raw.pyx:394: in cupy._core.raw.RawModule.module.__get__ return self._module() cupy/_core/raw.pyx:402: in cupy._core.raw.RawModule._module mod = _get_raw_module( cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/raw.pyx:547: in cupy._core.raw._get_raw_module mod = cupy._core.core.compile_with_cache( cupy/_core/core.pyx:2015: in cupy._core.core.compile_with_cache cpdef function.Module compile_with_cache( cupy/_core/core.pyx:2078: in cupy._core.core.compile_with_cache return cuda.compile_with_cache( cupy/cuda/compiler.py:453: in compile_with_cache return _compile_with_cache_hip( cupy/cuda/compiler.py:853: in _compile_with_cache_hip binary, mapping = compile_using_nvrtc( cupy/cuda/compiler.py:295: in compile_using_nvrtc return _compile(source, options, cu_path, cupy/cuda/compiler.py:279: in _compile compiled_obj, mapping = prog.compile(options, log_stream) _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ self = options = ('--std=c++11', '-I/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/_core/include', '-I/opt/rocm/include', '-fcuda-flush-denormals-to-zero') log_stream = None def compile(self, options=(), log_stream=None): try: if self.name_expressions: for ker in self.name_expressions: nvrtc.addAddNameExpression(self.ptr, ker) nvrtc.compileProgram(self.ptr, options) mapping = None if self.name_expressions: mapping = {} for ker in self.name_expressions: mapping[ker] = nvrtc.getLoweredName(self.ptr, ker) if log_stream is not None: log_stream.write(nvrtc.getProgramLog(self.ptr)) # This is to ensure backwards compatibility with nvrtc if self.method == 'cubin': return nvrtc.getCUBIN(self.ptr), mapping elif self.method == 'ptx': return nvrtc.getPTX(self.ptr), mapping else: raise RuntimeError('Unknown NVRTC compile method') except nvrtc.NVRTCError: log = nvrtc.getProgramLog(self.ptr) > raise CompileException(log, self.src, self.name, options, 'nvrtc' if not runtime.is_hip else 'hiprtc') E AssertionError: Only cupy raises error E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 650, in compile E nvrtc.compileProgram(self.ptr, options) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 136, in cupy_backends.cuda.libs.nvrtc.compileProgram E cpdef compileProgram(intptr_t prog, options): E File "cupy_backends/cuda/libs/nvrtc.pyx", line 148, in cupy_backends.cuda.libs.nvrtc.compileProgram E check_status(status) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 67, in cupy_backends.cuda.libs.nvrtc.check_status E raise NVRTCError(status) E cupy_backends.cuda.libs.nvrtc.NVRTCError: HIPRTC_ERROR_COMPILATION (6) E E During handling of the above exception, another exception occurred: E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/testing/_helper.py", line 47, in _call_func E result = impl(*args, **kw) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/tests/cupy_tests/linalg_tests/test_product.py", line 49, in test_dot E return xp.dot(a, b) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/linalg/_product.py", line 65, in dot E return a.dot(b, out) E File "cupy/_core/core.pyx", line 1607, in cupy._core.core.ndarray.dot E return _linalg.dot(self, b, out) E File "cupy/_core/_routines_linalg.pyx", line 427, in cupy._core._routines_linalg.dot E return tensordot_core(a, b, out, n, m, k, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 501, in cupy._core._routines_linalg.tensordot_core E return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 327, in cupy._core._routines_linalg._integral_tensordot_core E kern = _tensordot_core_int_kernel(config, dtype) E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/_routines_linalg.pyx", line 302, in cupy._core._routines_linalg._tensordot_core_int_kernel E ker = mod.get_function( E File "cupy/_core/raw.pyx", line 470, in cupy._core.raw.RawModule.get_function E mangled_name = self.module.mapping.get(name) E File "cupy/_core/raw.pyx", line 394, in cupy._core.raw.RawModule.module.__get__ E return self._module() E File "cupy/_core/raw.pyx", line 402, in cupy._core.raw.RawModule._module E mod = _get_raw_module( E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/raw.pyx", line 547, in cupy._core.raw._get_raw_module E mod = cupy._core.core.compile_with_cache( E File "cupy/_core/core.pyx", line 2015, in cupy._core.core.compile_with_cache E cpdef function.Module compile_with_cache( E File "cupy/_core/core.pyx", line 2078, in cupy._core.core.compile_with_cache E return cuda.compile_with_cache( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 453, in compile_with_cache E return _compile_with_cache_hip( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 853, in _compile_with_cache_hip E binary, mapping = compile_using_nvrtc( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 295, in compile_using_nvrtc E return _compile(source, options, cu_path, E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 279, in _compile E compiled_obj, mapping = prog.compile(options, log_stream) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 667, in compile E raise CompileException(log, self.src, self.name, options, E cupy.cuda.compiler.CompileException: /tmp/comgr-fae293/input/CompileSource:203:27: error: redeclaration of '__hiprtc_16' with a different type: 'void (*const)(int, int, int, const signed char *, const signed char *, signed char *)' vs 'void (*const)(int, int, int, const bool *, const bool *, bool *)' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-fae293/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-fae293/input/CompileSource:204:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-fae293/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-fae293/input/CompileSource:205:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-fae293/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-fae293/input/CompileSource:206:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-fae293/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-fae293/input/CompileSource:207:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-fae293/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-fae293/input/CompileSource:208:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-fae293/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-fae293/input/CompileSource:209:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-fae293/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-fae293/input/CompileSource:210:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-fae293/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-fae293/input/CompileSource:211:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-fae293/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-fae293/input/CompileSource:212:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-fae293/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E 10 errors generated when compiling for gfx908. E Error: Failed to compile opencl source (from CL or HIP source to LLVM IR). cupy/cuda/compiler.py:667: AssertionError ----------------------------- Captured stdout call ----------------------------- {'dtype_a': , 'dtype_b': } _ TestDot_param_35_{shape=((2, 3), (3, 4)), trans_a=False, trans_b=False}.test_dot _ self = < parameter: {'shape': ((2, 3), (3, 4)), 'trans_a': False, 'trans_b': False}> xp = dtype_a = , dtype_b = @testing.for_all_dtypes_combination(['dtype_a', 'dtype_b']) @testing.numpy_cupy_allclose() def test_dot(self, xp, dtype_a, dtype_b): shape_a, shape_b = self.shape if self.trans_a: a = testing.shaped_arange(shape_a[::-1], xp, dtype_a).T else: a = testing.shaped_arange(shape_a, xp, dtype_a) if self.trans_b: b = testing.shaped_arange(shape_b[::-1], xp, dtype_b).T else: b = testing.shaped_arange(shape_b, xp, dtype_b) > return xp.dot(a, b) tests/cupy_tests/linalg_tests/test_product.py:49: _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ cupy/linalg/_product.py:65: in dot return a.dot(b, out) cupy/_core/core.pyx:1607: in cupy._core.core.ndarray.dot return _linalg.dot(self, b, out) cupy/_core/_routines_linalg.pyx:427: in cupy._core._routines_linalg.dot return tensordot_core(a, b, out, n, m, k, ret_shape) cupy/_core/_routines_linalg.pyx:501: in cupy._core._routines_linalg.tensordot_core return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) cupy/_core/_routines_linalg.pyx:327: in cupy._core._routines_linalg._integral_tensordot_core kern = _tensordot_core_int_kernel(config, dtype) cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/_routines_linalg.pyx:302: in cupy._core._routines_linalg._tensordot_core_int_kernel ker = mod.get_function( cupy/_core/raw.pyx:470: in cupy._core.raw.RawModule.get_function mangled_name = self.module.mapping.get(name) cupy/_core/raw.pyx:394: in cupy._core.raw.RawModule.module.__get__ return self._module() cupy/_core/raw.pyx:402: in cupy._core.raw.RawModule._module mod = _get_raw_module( cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/raw.pyx:547: in cupy._core.raw._get_raw_module mod = cupy._core.core.compile_with_cache( cupy/_core/core.pyx:2015: in cupy._core.core.compile_with_cache cpdef function.Module compile_with_cache( cupy/_core/core.pyx:2078: in cupy._core.core.compile_with_cache return cuda.compile_with_cache( cupy/cuda/compiler.py:453: in compile_with_cache return _compile_with_cache_hip( cupy/cuda/compiler.py:853: in _compile_with_cache_hip binary, mapping = compile_using_nvrtc( cupy/cuda/compiler.py:295: in compile_using_nvrtc return _compile(source, options, cu_path, cupy/cuda/compiler.py:279: in _compile compiled_obj, mapping = prog.compile(options, log_stream) _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ self = options = ('--std=c++11', '-I/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/_core/include', '-I/opt/rocm/include', '-fcuda-flush-denormals-to-zero') log_stream = None def compile(self, options=(), log_stream=None): try: if self.name_expressions: for ker in self.name_expressions: nvrtc.addAddNameExpression(self.ptr, ker) nvrtc.compileProgram(self.ptr, options) mapping = None if self.name_expressions: mapping = {} for ker in self.name_expressions: mapping[ker] = nvrtc.getLoweredName(self.ptr, ker) if log_stream is not None: log_stream.write(nvrtc.getProgramLog(self.ptr)) # This is to ensure backwards compatibility with nvrtc if self.method == 'cubin': return nvrtc.getCUBIN(self.ptr), mapping elif self.method == 'ptx': return nvrtc.getPTX(self.ptr), mapping else: raise RuntimeError('Unknown NVRTC compile method') except nvrtc.NVRTCError: log = nvrtc.getProgramLog(self.ptr) > raise CompileException(log, self.src, self.name, options, 'nvrtc' if not runtime.is_hip else 'hiprtc') E AssertionError: Only cupy raises error E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 650, in compile E nvrtc.compileProgram(self.ptr, options) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 136, in cupy_backends.cuda.libs.nvrtc.compileProgram E cpdef compileProgram(intptr_t prog, options): E File "cupy_backends/cuda/libs/nvrtc.pyx", line 148, in cupy_backends.cuda.libs.nvrtc.compileProgram E check_status(status) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 67, in cupy_backends.cuda.libs.nvrtc.check_status E raise NVRTCError(status) E cupy_backends.cuda.libs.nvrtc.NVRTCError: HIPRTC_ERROR_COMPILATION (6) E E During handling of the above exception, another exception occurred: E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/testing/_helper.py", line 47, in _call_func E result = impl(*args, **kw) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/tests/cupy_tests/linalg_tests/test_product.py", line 49, in test_dot E return xp.dot(a, b) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/linalg/_product.py", line 65, in dot E return a.dot(b, out) E File "cupy/_core/core.pyx", line 1607, in cupy._core.core.ndarray.dot E return _linalg.dot(self, b, out) E File "cupy/_core/_routines_linalg.pyx", line 427, in cupy._core._routines_linalg.dot E return tensordot_core(a, b, out, n, m, k, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 501, in cupy._core._routines_linalg.tensordot_core E return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 327, in cupy._core._routines_linalg._integral_tensordot_core E kern = _tensordot_core_int_kernel(config, dtype) E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/_routines_linalg.pyx", line 302, in cupy._core._routines_linalg._tensordot_core_int_kernel E ker = mod.get_function( E File "cupy/_core/raw.pyx", line 470, in cupy._core.raw.RawModule.get_function E mangled_name = self.module.mapping.get(name) E File "cupy/_core/raw.pyx", line 394, in cupy._core.raw.RawModule.module.__get__ E return self._module() E File "cupy/_core/raw.pyx", line 402, in cupy._core.raw.RawModule._module E mod = _get_raw_module( E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/raw.pyx", line 547, in cupy._core.raw._get_raw_module E mod = cupy._core.core.compile_with_cache( E File "cupy/_core/core.pyx", line 2015, in cupy._core.core.compile_with_cache E cpdef function.Module compile_with_cache( E File "cupy/_core/core.pyx", line 2078, in cupy._core.core.compile_with_cache E return cuda.compile_with_cache( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 453, in compile_with_cache E return _compile_with_cache_hip( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 853, in _compile_with_cache_hip E binary, mapping = compile_using_nvrtc( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 295, in compile_using_nvrtc E return _compile(source, options, cu_path, E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 279, in _compile E compiled_obj, mapping = prog.compile(options, log_stream) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 667, in compile E raise CompileException(log, self.src, self.name, options, E cupy.cuda.compiler.CompileException: /tmp/comgr-0bb249/input/CompileSource:203:27: error: redeclaration of '__hiprtc_16' with a different type: 'void (*const)(int, int, int, const signed char *, const signed char *, signed char *)' vs 'void (*const)(int, int, int, const bool *, const bool *, bool *)' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-0bb249/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-0bb249/input/CompileSource:204:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-0bb249/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-0bb249/input/CompileSource:205:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-0bb249/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-0bb249/input/CompileSource:206:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-0bb249/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-0bb249/input/CompileSource:207:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-0bb249/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-0bb249/input/CompileSource:208:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-0bb249/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-0bb249/input/CompileSource:209:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-0bb249/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-0bb249/input/CompileSource:210:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-0bb249/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-0bb249/input/CompileSource:211:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-0bb249/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-0bb249/input/CompileSource:212:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-0bb249/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E 10 errors generated when compiling for gfx908. E Error: Failed to compile opencl source (from CL or HIP source to LLVM IR). cupy/cuda/compiler.py:667: AssertionError ----------------------------- Captured stdout call ----------------------------- {'dtype_a': , 'dtype_b': } _ TestDot_param_64_{shape=((2,), (2, 4)), trans_a=True, trans_b=True}.test_dot _ self = < parameter: {'shape': ((2,), (2, 4)), 'trans_a': True, 'trans_b': True}> xp = dtype_a = , dtype_b = @testing.for_all_dtypes_combination(['dtype_a', 'dtype_b']) @testing.numpy_cupy_allclose() def test_dot(self, xp, dtype_a, dtype_b): shape_a, shape_b = self.shape if self.trans_a: a = testing.shaped_arange(shape_a[::-1], xp, dtype_a).T else: a = testing.shaped_arange(shape_a, xp, dtype_a) if self.trans_b: b = testing.shaped_arange(shape_b[::-1], xp, dtype_b).T else: b = testing.shaped_arange(shape_b, xp, dtype_b) > return xp.dot(a, b) tests/cupy_tests/linalg_tests/test_product.py:49: _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ cupy/linalg/_product.py:65: in dot return a.dot(b, out) cupy/_core/core.pyx:1607: in cupy._core.core.ndarray.dot return _linalg.dot(self, b, out) cupy/_core/_routines_linalg.pyx:427: in cupy._core._routines_linalg.dot return tensordot_core(a, b, out, n, m, k, ret_shape) cupy/_core/_routines_linalg.pyx:501: in cupy._core._routines_linalg.tensordot_core return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) cupy/_core/_routines_linalg.pyx:327: in cupy._core._routines_linalg._integral_tensordot_core kern = _tensordot_core_int_kernel(config, dtype) cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/_routines_linalg.pyx:302: in cupy._core._routines_linalg._tensordot_core_int_kernel ker = mod.get_function( cupy/_core/raw.pyx:470: in cupy._core.raw.RawModule.get_function mangled_name = self.module.mapping.get(name) cupy/_core/raw.pyx:394: in cupy._core.raw.RawModule.module.__get__ return self._module() cupy/_core/raw.pyx:402: in cupy._core.raw.RawModule._module mod = _get_raw_module( cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/raw.pyx:547: in cupy._core.raw._get_raw_module mod = cupy._core.core.compile_with_cache( cupy/_core/core.pyx:2015: in cupy._core.core.compile_with_cache cpdef function.Module compile_with_cache( cupy/_core/core.pyx:2078: in cupy._core.core.compile_with_cache return cuda.compile_with_cache( cupy/cuda/compiler.py:453: in compile_with_cache return _compile_with_cache_hip( cupy/cuda/compiler.py:853: in _compile_with_cache_hip binary, mapping = compile_using_nvrtc( cupy/cuda/compiler.py:295: in compile_using_nvrtc return _compile(source, options, cu_path, cupy/cuda/compiler.py:279: in _compile compiled_obj, mapping = prog.compile(options, log_stream) _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ self = options = ('--std=c++11', '-I/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/_core/include', '-I/opt/rocm/include', '-fcuda-flush-denormals-to-zero') log_stream = None def compile(self, options=(), log_stream=None): try: if self.name_expressions: for ker in self.name_expressions: nvrtc.addAddNameExpression(self.ptr, ker) nvrtc.compileProgram(self.ptr, options) mapping = None if self.name_expressions: mapping = {} for ker in self.name_expressions: mapping[ker] = nvrtc.getLoweredName(self.ptr, ker) if log_stream is not None: log_stream.write(nvrtc.getProgramLog(self.ptr)) # This is to ensure backwards compatibility with nvrtc if self.method == 'cubin': return nvrtc.getCUBIN(self.ptr), mapping elif self.method == 'ptx': return nvrtc.getPTX(self.ptr), mapping else: raise RuntimeError('Unknown NVRTC compile method') except nvrtc.NVRTCError: log = nvrtc.getProgramLog(self.ptr) > raise CompileException(log, self.src, self.name, options, 'nvrtc' if not runtime.is_hip else 'hiprtc') E AssertionError: Only cupy raises error E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 650, in compile E nvrtc.compileProgram(self.ptr, options) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 136, in cupy_backends.cuda.libs.nvrtc.compileProgram E cpdef compileProgram(intptr_t prog, options): E File "cupy_backends/cuda/libs/nvrtc.pyx", line 148, in cupy_backends.cuda.libs.nvrtc.compileProgram E check_status(status) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 67, in cupy_backends.cuda.libs.nvrtc.check_status E raise NVRTCError(status) E cupy_backends.cuda.libs.nvrtc.NVRTCError: HIPRTC_ERROR_COMPILATION (6) E E During handling of the above exception, another exception occurred: E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/testing/_helper.py", line 47, in _call_func E result = impl(*args, **kw) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/tests/cupy_tests/linalg_tests/test_product.py", line 49, in test_dot E return xp.dot(a, b) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/linalg/_product.py", line 65, in dot E return a.dot(b, out) E File "cupy/_core/core.pyx", line 1607, in cupy._core.core.ndarray.dot E return _linalg.dot(self, b, out) E File "cupy/_core/_routines_linalg.pyx", line 427, in cupy._core._routines_linalg.dot E return tensordot_core(a, b, out, n, m, k, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 501, in cupy._core._routines_linalg.tensordot_core E return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 327, in cupy._core._routines_linalg._integral_tensordot_core E kern = _tensordot_core_int_kernel(config, dtype) E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/_routines_linalg.pyx", line 302, in cupy._core._routines_linalg._tensordot_core_int_kernel E ker = mod.get_function( E File "cupy/_core/raw.pyx", line 470, in cupy._core.raw.RawModule.get_function E mangled_name = self.module.mapping.get(name) E File "cupy/_core/raw.pyx", line 394, in cupy._core.raw.RawModule.module.__get__ E return self._module() E File "cupy/_core/raw.pyx", line 402, in cupy._core.raw.RawModule._module E mod = _get_raw_module( E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/raw.pyx", line 547, in cupy._core.raw._get_raw_module E mod = cupy._core.core.compile_with_cache( E File "cupy/_core/core.pyx", line 2015, in cupy._core.core.compile_with_cache E cpdef function.Module compile_with_cache( E File "cupy/_core/core.pyx", line 2078, in cupy._core.core.compile_with_cache E return cuda.compile_with_cache( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 453, in compile_with_cache E return _compile_with_cache_hip( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 853, in _compile_with_cache_hip E binary, mapping = compile_using_nvrtc( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 295, in compile_using_nvrtc E return _compile(source, options, cu_path, E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 279, in _compile E compiled_obj, mapping = prog.compile(options, log_stream) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 667, in compile E raise CompileException(log, self.src, self.name, options, E cupy.cuda.compiler.CompileException: /tmp/comgr-6f5594/input/CompileSource:203:27: error: redeclaration of '__hiprtc_16' with a different type: 'void (*const)(int, int, int, const signed char *, const signed char *, signed char *)' vs 'void (*const)(int, int, int, const bool *, const bool *, bool *)' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-6f5594/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-6f5594/input/CompileSource:204:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-6f5594/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-6f5594/input/CompileSource:205:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-6f5594/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-6f5594/input/CompileSource:206:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-6f5594/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-6f5594/input/CompileSource:207:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-6f5594/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-6f5594/input/CompileSource:208:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-6f5594/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-6f5594/input/CompileSource:209:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-6f5594/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-6f5594/input/CompileSource:210:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-6f5594/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-6f5594/input/CompileSource:211:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-6f5594/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-6f5594/input/CompileSource:212:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-6f5594/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E 10 errors generated when compiling for gfx908. E Error: Failed to compile opencl source (from CL or HIP source to LLVM IR). cupy/cuda/compiler.py:667: AssertionError ----------------------------- Captured stdout call ----------------------------- {'dtype_a': , 'dtype_b': } _ TestDot_param_65_{shape=((2,), (2, 4)), trans_a=True, trans_b=False}.test_dot _ self = < parameter: {'shape': ((2,), (2, 4)), 'trans_a': True, 'trans_b': False}> xp = dtype_a = , dtype_b = @testing.for_all_dtypes_combination(['dtype_a', 'dtype_b']) @testing.numpy_cupy_allclose() def test_dot(self, xp, dtype_a, dtype_b): shape_a, shape_b = self.shape if self.trans_a: a = testing.shaped_arange(shape_a[::-1], xp, dtype_a).T else: a = testing.shaped_arange(shape_a, xp, dtype_a) if self.trans_b: b = testing.shaped_arange(shape_b[::-1], xp, dtype_b).T else: b = testing.shaped_arange(shape_b, xp, dtype_b) > return xp.dot(a, b) tests/cupy_tests/linalg_tests/test_product.py:49: _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ cupy/linalg/_product.py:65: in dot return a.dot(b, out) cupy/_core/core.pyx:1607: in cupy._core.core.ndarray.dot return _linalg.dot(self, b, out) cupy/_core/_routines_linalg.pyx:427: in cupy._core._routines_linalg.dot return tensordot_core(a, b, out, n, m, k, ret_shape) cupy/_core/_routines_linalg.pyx:501: in cupy._core._routines_linalg.tensordot_core return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) cupy/_core/_routines_linalg.pyx:327: in cupy._core._routines_linalg._integral_tensordot_core kern = _tensordot_core_int_kernel(config, dtype) cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/_routines_linalg.pyx:302: in cupy._core._routines_linalg._tensordot_core_int_kernel ker = mod.get_function( cupy/_core/raw.pyx:470: in cupy._core.raw.RawModule.get_function mangled_name = self.module.mapping.get(name) cupy/_core/raw.pyx:394: in cupy._core.raw.RawModule.module.__get__ return self._module() cupy/_core/raw.pyx:402: in cupy._core.raw.RawModule._module mod = _get_raw_module( cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/raw.pyx:547: in cupy._core.raw._get_raw_module mod = cupy._core.core.compile_with_cache( cupy/_core/core.pyx:2015: in cupy._core.core.compile_with_cache cpdef function.Module compile_with_cache( cupy/_core/core.pyx:2078: in cupy._core.core.compile_with_cache return cuda.compile_with_cache( cupy/cuda/compiler.py:453: in compile_with_cache return _compile_with_cache_hip( cupy/cuda/compiler.py:853: in _compile_with_cache_hip binary, mapping = compile_using_nvrtc( cupy/cuda/compiler.py:295: in compile_using_nvrtc return _compile(source, options, cu_path, cupy/cuda/compiler.py:279: in _compile compiled_obj, mapping = prog.compile(options, log_stream) _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ self = options = ('--std=c++11', '-I/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/_core/include', '-I/opt/rocm/include', '-fcuda-flush-denormals-to-zero') log_stream = None def compile(self, options=(), log_stream=None): try: if self.name_expressions: for ker in self.name_expressions: nvrtc.addAddNameExpression(self.ptr, ker) nvrtc.compileProgram(self.ptr, options) mapping = None if self.name_expressions: mapping = {} for ker in self.name_expressions: mapping[ker] = nvrtc.getLoweredName(self.ptr, ker) if log_stream is not None: log_stream.write(nvrtc.getProgramLog(self.ptr)) # This is to ensure backwards compatibility with nvrtc if self.method == 'cubin': return nvrtc.getCUBIN(self.ptr), mapping elif self.method == 'ptx': return nvrtc.getPTX(self.ptr), mapping else: raise RuntimeError('Unknown NVRTC compile method') except nvrtc.NVRTCError: log = nvrtc.getProgramLog(self.ptr) > raise CompileException(log, self.src, self.name, options, 'nvrtc' if not runtime.is_hip else 'hiprtc') E AssertionError: Only cupy raises error E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 650, in compile E nvrtc.compileProgram(self.ptr, options) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 136, in cupy_backends.cuda.libs.nvrtc.compileProgram E cpdef compileProgram(intptr_t prog, options): E File "cupy_backends/cuda/libs/nvrtc.pyx", line 148, in cupy_backends.cuda.libs.nvrtc.compileProgram E check_status(status) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 67, in cupy_backends.cuda.libs.nvrtc.check_status E raise NVRTCError(status) E cupy_backends.cuda.libs.nvrtc.NVRTCError: HIPRTC_ERROR_COMPILATION (6) E E During handling of the above exception, another exception occurred: E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/testing/_helper.py", line 47, in _call_func E result = impl(*args, **kw) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/tests/cupy_tests/linalg_tests/test_product.py", line 49, in test_dot E return xp.dot(a, b) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/linalg/_product.py", line 65, in dot E return a.dot(b, out) E File "cupy/_core/core.pyx", line 1607, in cupy._core.core.ndarray.dot E return _linalg.dot(self, b, out) E File "cupy/_core/_routines_linalg.pyx", line 427, in cupy._core._routines_linalg.dot E return tensordot_core(a, b, out, n, m, k, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 501, in cupy._core._routines_linalg.tensordot_core E return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 327, in cupy._core._routines_linalg._integral_tensordot_core E kern = _tensordot_core_int_kernel(config, dtype) E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/_routines_linalg.pyx", line 302, in cupy._core._routines_linalg._tensordot_core_int_kernel E ker = mod.get_function( E File "cupy/_core/raw.pyx", line 470, in cupy._core.raw.RawModule.get_function E mangled_name = self.module.mapping.get(name) E File "cupy/_core/raw.pyx", line 394, in cupy._core.raw.RawModule.module.__get__ E return self._module() E File "cupy/_core/raw.pyx", line 402, in cupy._core.raw.RawModule._module E mod = _get_raw_module( E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/raw.pyx", line 547, in cupy._core.raw._get_raw_module E mod = cupy._core.core.compile_with_cache( E File "cupy/_core/core.pyx", line 2015, in cupy._core.core.compile_with_cache E cpdef function.Module compile_with_cache( E File "cupy/_core/core.pyx", line 2078, in cupy._core.core.compile_with_cache E return cuda.compile_with_cache( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 453, in compile_with_cache E return _compile_with_cache_hip( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 853, in _compile_with_cache_hip E binary, mapping = compile_using_nvrtc( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 295, in compile_using_nvrtc E return _compile(source, options, cu_path, E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 279, in _compile E compiled_obj, mapping = prog.compile(options, log_stream) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 667, in compile E raise CompileException(log, self.src, self.name, options, E cupy.cuda.compiler.CompileException: /tmp/comgr-cabd56/input/CompileSource:203:27: error: redeclaration of '__hiprtc_16' with a different type: 'void (*const)(int, int, int, const signed char *, const signed char *, signed char *)' vs 'void (*const)(int, int, int, const bool *, const bool *, bool *)' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-cabd56/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-cabd56/input/CompileSource:204:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-cabd56/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-cabd56/input/CompileSource:205:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-cabd56/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-cabd56/input/CompileSource:206:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-cabd56/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-cabd56/input/CompileSource:207:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-cabd56/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-cabd56/input/CompileSource:208:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-cabd56/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-cabd56/input/CompileSource:209:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-cabd56/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-cabd56/input/CompileSource:210:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-cabd56/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-cabd56/input/CompileSource:211:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-cabd56/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-cabd56/input/CompileSource:212:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-cabd56/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E 10 errors generated when compiling for gfx908. E Error: Failed to compile opencl source (from CL or HIP source to LLVM IR). cupy/cuda/compiler.py:667: AssertionError ----------------------------- Captured stdout call ----------------------------- {'dtype_a': , 'dtype_b': } _ TestDot_param_66_{shape=((2,), (2, 4)), trans_a=False, trans_b=True}.test_dot _ self = < parameter: {'shape': ((2,), (2, 4)), 'trans_a': False, 'trans_b': True}> xp = dtype_a = , dtype_b = @testing.for_all_dtypes_combination(['dtype_a', 'dtype_b']) @testing.numpy_cupy_allclose() def test_dot(self, xp, dtype_a, dtype_b): shape_a, shape_b = self.shape if self.trans_a: a = testing.shaped_arange(shape_a[::-1], xp, dtype_a).T else: a = testing.shaped_arange(shape_a, xp, dtype_a) if self.trans_b: b = testing.shaped_arange(shape_b[::-1], xp, dtype_b).T else: b = testing.shaped_arange(shape_b, xp, dtype_b) > return xp.dot(a, b) tests/cupy_tests/linalg_tests/test_product.py:49: _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ cupy/linalg/_product.py:65: in dot return a.dot(b, out) cupy/_core/core.pyx:1607: in cupy._core.core.ndarray.dot return _linalg.dot(self, b, out) cupy/_core/_routines_linalg.pyx:427: in cupy._core._routines_linalg.dot return tensordot_core(a, b, out, n, m, k, ret_shape) cupy/_core/_routines_linalg.pyx:501: in cupy._core._routines_linalg.tensordot_core return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) cupy/_core/_routines_linalg.pyx:327: in cupy._core._routines_linalg._integral_tensordot_core kern = _tensordot_core_int_kernel(config, dtype) cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/_routines_linalg.pyx:302: in cupy._core._routines_linalg._tensordot_core_int_kernel ker = mod.get_function( cupy/_core/raw.pyx:470: in cupy._core.raw.RawModule.get_function mangled_name = self.module.mapping.get(name) cupy/_core/raw.pyx:394: in cupy._core.raw.RawModule.module.__get__ return self._module() cupy/_core/raw.pyx:402: in cupy._core.raw.RawModule._module mod = _get_raw_module( cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/raw.pyx:547: in cupy._core.raw._get_raw_module mod = cupy._core.core.compile_with_cache( cupy/_core/core.pyx:2015: in cupy._core.core.compile_with_cache cpdef function.Module compile_with_cache( cupy/_core/core.pyx:2078: in cupy._core.core.compile_with_cache return cuda.compile_with_cache( cupy/cuda/compiler.py:453: in compile_with_cache return _compile_with_cache_hip( cupy/cuda/compiler.py:853: in _compile_with_cache_hip binary, mapping = compile_using_nvrtc( cupy/cuda/compiler.py:295: in compile_using_nvrtc return _compile(source, options, cu_path, cupy/cuda/compiler.py:279: in _compile compiled_obj, mapping = prog.compile(options, log_stream) _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ self = options = ('--std=c++11', '-I/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/_core/include', '-I/opt/rocm/include', '-fcuda-flush-denormals-to-zero') log_stream = None def compile(self, options=(), log_stream=None): try: if self.name_expressions: for ker in self.name_expressions: nvrtc.addAddNameExpression(self.ptr, ker) nvrtc.compileProgram(self.ptr, options) mapping = None if self.name_expressions: mapping = {} for ker in self.name_expressions: mapping[ker] = nvrtc.getLoweredName(self.ptr, ker) if log_stream is not None: log_stream.write(nvrtc.getProgramLog(self.ptr)) # This is to ensure backwards compatibility with nvrtc if self.method == 'cubin': return nvrtc.getCUBIN(self.ptr), mapping elif self.method == 'ptx': return nvrtc.getPTX(self.ptr), mapping else: raise RuntimeError('Unknown NVRTC compile method') except nvrtc.NVRTCError: log = nvrtc.getProgramLog(self.ptr) > raise CompileException(log, self.src, self.name, options, 'nvrtc' if not runtime.is_hip else 'hiprtc') E AssertionError: Only cupy raises error E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 650, in compile E nvrtc.compileProgram(self.ptr, options) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 136, in cupy_backends.cuda.libs.nvrtc.compileProgram E cpdef compileProgram(intptr_t prog, options): E File "cupy_backends/cuda/libs/nvrtc.pyx", line 148, in cupy_backends.cuda.libs.nvrtc.compileProgram E check_status(status) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 67, in cupy_backends.cuda.libs.nvrtc.check_status E raise NVRTCError(status) E cupy_backends.cuda.libs.nvrtc.NVRTCError: HIPRTC_ERROR_COMPILATION (6) E E During handling of the above exception, another exception occurred: E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/testing/_helper.py", line 47, in _call_func E result = impl(*args, **kw) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/tests/cupy_tests/linalg_tests/test_product.py", line 49, in test_dot E return xp.dot(a, b) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/linalg/_product.py", line 65, in dot E return a.dot(b, out) E File "cupy/_core/core.pyx", line 1607, in cupy._core.core.ndarray.dot E return _linalg.dot(self, b, out) E File "cupy/_core/_routines_linalg.pyx", line 427, in cupy._core._routines_linalg.dot E return tensordot_core(a, b, out, n, m, k, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 501, in cupy._core._routines_linalg.tensordot_core E return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 327, in cupy._core._routines_linalg._integral_tensordot_core E kern = _tensordot_core_int_kernel(config, dtype) E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/_routines_linalg.pyx", line 302, in cupy._core._routines_linalg._tensordot_core_int_kernel E ker = mod.get_function( E File "cupy/_core/raw.pyx", line 470, in cupy._core.raw.RawModule.get_function E mangled_name = self.module.mapping.get(name) E File "cupy/_core/raw.pyx", line 394, in cupy._core.raw.RawModule.module.__get__ E return self._module() E File "cupy/_core/raw.pyx", line 402, in cupy._core.raw.RawModule._module E mod = _get_raw_module( E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/raw.pyx", line 547, in cupy._core.raw._get_raw_module E mod = cupy._core.core.compile_with_cache( E File "cupy/_core/core.pyx", line 2015, in cupy._core.core.compile_with_cache E cpdef function.Module compile_with_cache( E File "cupy/_core/core.pyx", line 2078, in cupy._core.core.compile_with_cache E return cuda.compile_with_cache( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 453, in compile_with_cache E return _compile_with_cache_hip( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 853, in _compile_with_cache_hip E binary, mapping = compile_using_nvrtc( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 295, in compile_using_nvrtc E return _compile(source, options, cu_path, E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 279, in _compile E compiled_obj, mapping = prog.compile(options, log_stream) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 667, in compile E raise CompileException(log, self.src, self.name, options, E cupy.cuda.compiler.CompileException: /tmp/comgr-842e37/input/CompileSource:203:27: error: redeclaration of '__hiprtc_16' with a different type: 'void (*const)(int, int, int, const signed char *, const signed char *, signed char *)' vs 'void (*const)(int, int, int, const bool *, const bool *, bool *)' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-842e37/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-842e37/input/CompileSource:204:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-842e37/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-842e37/input/CompileSource:205:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-842e37/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-842e37/input/CompileSource:206:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-842e37/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-842e37/input/CompileSource:207:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-842e37/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-842e37/input/CompileSource:208:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-842e37/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-842e37/input/CompileSource:209:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-842e37/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-842e37/input/CompileSource:210:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-842e37/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-842e37/input/CompileSource:211:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-842e37/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-842e37/input/CompileSource:212:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-842e37/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E 10 errors generated when compiling for gfx908. E Error: Failed to compile opencl source (from CL or HIP source to LLVM IR). cupy/cuda/compiler.py:667: AssertionError ----------------------------- Captured stdout call ----------------------------- {'dtype_a': , 'dtype_b': } _ TestDot_param_67_{shape=((2,), (2, 4)), trans_a=False, trans_b=False}.test_dot _ self = < parameter: {'shape': ((2,), (2, 4)), 'trans_a': False, 'trans_b': False}> xp = dtype_a = , dtype_b = @testing.for_all_dtypes_combination(['dtype_a', 'dtype_b']) @testing.numpy_cupy_allclose() def test_dot(self, xp, dtype_a, dtype_b): shape_a, shape_b = self.shape if self.trans_a: a = testing.shaped_arange(shape_a[::-1], xp, dtype_a).T else: a = testing.shaped_arange(shape_a, xp, dtype_a) if self.trans_b: b = testing.shaped_arange(shape_b[::-1], xp, dtype_b).T else: b = testing.shaped_arange(shape_b, xp, dtype_b) > return xp.dot(a, b) tests/cupy_tests/linalg_tests/test_product.py:49: _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ cupy/linalg/_product.py:65: in dot return a.dot(b, out) cupy/_core/core.pyx:1607: in cupy._core.core.ndarray.dot return _linalg.dot(self, b, out) cupy/_core/_routines_linalg.pyx:427: in cupy._core._routines_linalg.dot return tensordot_core(a, b, out, n, m, k, ret_shape) cupy/_core/_routines_linalg.pyx:501: in cupy._core._routines_linalg.tensordot_core return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) cupy/_core/_routines_linalg.pyx:327: in cupy._core._routines_linalg._integral_tensordot_core kern = _tensordot_core_int_kernel(config, dtype) cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/_routines_linalg.pyx:302: in cupy._core._routines_linalg._tensordot_core_int_kernel ker = mod.get_function( cupy/_core/raw.pyx:470: in cupy._core.raw.RawModule.get_function mangled_name = self.module.mapping.get(name) cupy/_core/raw.pyx:394: in cupy._core.raw.RawModule.module.__get__ return self._module() cupy/_core/raw.pyx:402: in cupy._core.raw.RawModule._module mod = _get_raw_module( cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/raw.pyx:547: in cupy._core.raw._get_raw_module mod = cupy._core.core.compile_with_cache( cupy/_core/core.pyx:2015: in cupy._core.core.compile_with_cache cpdef function.Module compile_with_cache( cupy/_core/core.pyx:2078: in cupy._core.core.compile_with_cache return cuda.compile_with_cache( cupy/cuda/compiler.py:453: in compile_with_cache return _compile_with_cache_hip( cupy/cuda/compiler.py:853: in _compile_with_cache_hip binary, mapping = compile_using_nvrtc( cupy/cuda/compiler.py:295: in compile_using_nvrtc return _compile(source, options, cu_path, cupy/cuda/compiler.py:279: in _compile compiled_obj, mapping = prog.compile(options, log_stream) _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ self = options = ('--std=c++11', '-I/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/_core/include', '-I/opt/rocm/include', '-fcuda-flush-denormals-to-zero') log_stream = None def compile(self, options=(), log_stream=None): try: if self.name_expressions: for ker in self.name_expressions: nvrtc.addAddNameExpression(self.ptr, ker) nvrtc.compileProgram(self.ptr, options) mapping = None if self.name_expressions: mapping = {} for ker in self.name_expressions: mapping[ker] = nvrtc.getLoweredName(self.ptr, ker) if log_stream is not None: log_stream.write(nvrtc.getProgramLog(self.ptr)) # This is to ensure backwards compatibility with nvrtc if self.method == 'cubin': return nvrtc.getCUBIN(self.ptr), mapping elif self.method == 'ptx': return nvrtc.getPTX(self.ptr), mapping else: raise RuntimeError('Unknown NVRTC compile method') except nvrtc.NVRTCError: log = nvrtc.getProgramLog(self.ptr) > raise CompileException(log, self.src, self.name, options, 'nvrtc' if not runtime.is_hip else 'hiprtc') E AssertionError: Only cupy raises error E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 650, in compile E nvrtc.compileProgram(self.ptr, options) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 136, in cupy_backends.cuda.libs.nvrtc.compileProgram E cpdef compileProgram(intptr_t prog, options): E File "cupy_backends/cuda/libs/nvrtc.pyx", line 148, in cupy_backends.cuda.libs.nvrtc.compileProgram E check_status(status) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 67, in cupy_backends.cuda.libs.nvrtc.check_status E raise NVRTCError(status) E cupy_backends.cuda.libs.nvrtc.NVRTCError: HIPRTC_ERROR_COMPILATION (6) E E During handling of the above exception, another exception occurred: E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/testing/_helper.py", line 47, in _call_func E result = impl(*args, **kw) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/tests/cupy_tests/linalg_tests/test_product.py", line 49, in test_dot E return xp.dot(a, b) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/linalg/_product.py", line 65, in dot E return a.dot(b, out) E File "cupy/_core/core.pyx", line 1607, in cupy._core.core.ndarray.dot E return _linalg.dot(self, b, out) E File "cupy/_core/_routines_linalg.pyx", line 427, in cupy._core._routines_linalg.dot E return tensordot_core(a, b, out, n, m, k, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 501, in cupy._core._routines_linalg.tensordot_core E return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 327, in cupy._core._routines_linalg._integral_tensordot_core E kern = _tensordot_core_int_kernel(config, dtype) E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/_routines_linalg.pyx", line 302, in cupy._core._routines_linalg._tensordot_core_int_kernel E ker = mod.get_function( E File "cupy/_core/raw.pyx", line 470, in cupy._core.raw.RawModule.get_function E mangled_name = self.module.mapping.get(name) E File "cupy/_core/raw.pyx", line 394, in cupy._core.raw.RawModule.module.__get__ E return self._module() E File "cupy/_core/raw.pyx", line 402, in cupy._core.raw.RawModule._module E mod = _get_raw_module( E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/raw.pyx", line 547, in cupy._core.raw._get_raw_module E mod = cupy._core.core.compile_with_cache( E File "cupy/_core/core.pyx", line 2015, in cupy._core.core.compile_with_cache E cpdef function.Module compile_with_cache( E File "cupy/_core/core.pyx", line 2078, in cupy._core.core.compile_with_cache E return cuda.compile_with_cache( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 453, in compile_with_cache E return _compile_with_cache_hip( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 853, in _compile_with_cache_hip E binary, mapping = compile_using_nvrtc( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 295, in compile_using_nvrtc E return _compile(source, options, cu_path, E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 279, in _compile E compiled_obj, mapping = prog.compile(options, log_stream) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 667, in compile E raise CompileException(log, self.src, self.name, options, E cupy.cuda.compiler.CompileException: /tmp/comgr-9186b3/input/CompileSource:203:27: error: redeclaration of '__hiprtc_16' with a different type: 'void (*const)(int, int, int, const signed char *, const signed char *, signed char *)' vs 'void (*const)(int, int, int, const bool *, const bool *, bool *)' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-9186b3/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-9186b3/input/CompileSource:204:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-9186b3/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-9186b3/input/CompileSource:205:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-9186b3/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-9186b3/input/CompileSource:206:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-9186b3/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-9186b3/input/CompileSource:207:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-9186b3/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-9186b3/input/CompileSource:208:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-9186b3/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-9186b3/input/CompileSource:209:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-9186b3/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-9186b3/input/CompileSource:210:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-9186b3/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-9186b3/input/CompileSource:211:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-9186b3/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-9186b3/input/CompileSource:212:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-9186b3/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E 10 errors generated when compiling for gfx908. E Error: Failed to compile opencl source (from CL or HIP source to LLVM IR). cupy/cuda/compiler.py:667: AssertionError ----------------------------- Captured stdout call ----------------------------- {'dtype_a': , 'dtype_b': } _ TestDot_param_68_{shape=((4, 2), (2,)), trans_a=True, trans_b=True}.test_dot _ self = < parameter: {'shape': ((4, 2), (2,)), 'trans_a': True, 'trans_b': True}> xp = dtype_a = , dtype_b = @testing.for_all_dtypes_combination(['dtype_a', 'dtype_b']) @testing.numpy_cupy_allclose() def test_dot(self, xp, dtype_a, dtype_b): shape_a, shape_b = self.shape if self.trans_a: a = testing.shaped_arange(shape_a[::-1], xp, dtype_a).T else: a = testing.shaped_arange(shape_a, xp, dtype_a) if self.trans_b: b = testing.shaped_arange(shape_b[::-1], xp, dtype_b).T else: b = testing.shaped_arange(shape_b, xp, dtype_b) > return xp.dot(a, b) tests/cupy_tests/linalg_tests/test_product.py:49: _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ cupy/linalg/_product.py:65: in dot return a.dot(b, out) cupy/_core/core.pyx:1607: in cupy._core.core.ndarray.dot return _linalg.dot(self, b, out) cupy/_core/_routines_linalg.pyx:427: in cupy._core._routines_linalg.dot return tensordot_core(a, b, out, n, m, k, ret_shape) cupy/_core/_routines_linalg.pyx:501: in cupy._core._routines_linalg.tensordot_core return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) cupy/_core/_routines_linalg.pyx:327: in cupy._core._routines_linalg._integral_tensordot_core kern = _tensordot_core_int_kernel(config, dtype) cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/_routines_linalg.pyx:302: in cupy._core._routines_linalg._tensordot_core_int_kernel ker = mod.get_function( cupy/_core/raw.pyx:470: in cupy._core.raw.RawModule.get_function mangled_name = self.module.mapping.get(name) cupy/_core/raw.pyx:394: in cupy._core.raw.RawModule.module.__get__ return self._module() cupy/_core/raw.pyx:402: in cupy._core.raw.RawModule._module mod = _get_raw_module( cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/raw.pyx:547: in cupy._core.raw._get_raw_module mod = cupy._core.core.compile_with_cache( cupy/_core/core.pyx:2015: in cupy._core.core.compile_with_cache cpdef function.Module compile_with_cache( cupy/_core/core.pyx:2078: in cupy._core.core.compile_with_cache return cuda.compile_with_cache( cupy/cuda/compiler.py:453: in compile_with_cache return _compile_with_cache_hip( cupy/cuda/compiler.py:853: in _compile_with_cache_hip binary, mapping = compile_using_nvrtc( cupy/cuda/compiler.py:295: in compile_using_nvrtc return _compile(source, options, cu_path, cupy/cuda/compiler.py:279: in _compile compiled_obj, mapping = prog.compile(options, log_stream) _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ self = options = ('--std=c++11', '-I/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/_core/include', '-I/opt/rocm/include', '-fcuda-flush-denormals-to-zero') log_stream = None def compile(self, options=(), log_stream=None): try: if self.name_expressions: for ker in self.name_expressions: nvrtc.addAddNameExpression(self.ptr, ker) nvrtc.compileProgram(self.ptr, options) mapping = None if self.name_expressions: mapping = {} for ker in self.name_expressions: mapping[ker] = nvrtc.getLoweredName(self.ptr, ker) if log_stream is not None: log_stream.write(nvrtc.getProgramLog(self.ptr)) # This is to ensure backwards compatibility with nvrtc if self.method == 'cubin': return nvrtc.getCUBIN(self.ptr), mapping elif self.method == 'ptx': return nvrtc.getPTX(self.ptr), mapping else: raise RuntimeError('Unknown NVRTC compile method') except nvrtc.NVRTCError: log = nvrtc.getProgramLog(self.ptr) > raise CompileException(log, self.src, self.name, options, 'nvrtc' if not runtime.is_hip else 'hiprtc') E AssertionError: Only cupy raises error E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 650, in compile E nvrtc.compileProgram(self.ptr, options) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 136, in cupy_backends.cuda.libs.nvrtc.compileProgram E cpdef compileProgram(intptr_t prog, options): E File "cupy_backends/cuda/libs/nvrtc.pyx", line 148, in cupy_backends.cuda.libs.nvrtc.compileProgram E check_status(status) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 67, in cupy_backends.cuda.libs.nvrtc.check_status E raise NVRTCError(status) E cupy_backends.cuda.libs.nvrtc.NVRTCError: HIPRTC_ERROR_COMPILATION (6) E E During handling of the above exception, another exception occurred: E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/testing/_helper.py", line 47, in _call_func E result = impl(*args, **kw) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/tests/cupy_tests/linalg_tests/test_product.py", line 49, in test_dot E return xp.dot(a, b) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/linalg/_product.py", line 65, in dot E return a.dot(b, out) E File "cupy/_core/core.pyx", line 1607, in cupy._core.core.ndarray.dot E return _linalg.dot(self, b, out) E File "cupy/_core/_routines_linalg.pyx", line 427, in cupy._core._routines_linalg.dot E return tensordot_core(a, b, out, n, m, k, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 501, in cupy._core._routines_linalg.tensordot_core E return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 327, in cupy._core._routines_linalg._integral_tensordot_core E kern = _tensordot_core_int_kernel(config, dtype) E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/_routines_linalg.pyx", line 302, in cupy._core._routines_linalg._tensordot_core_int_kernel E ker = mod.get_function( E File "cupy/_core/raw.pyx", line 470, in cupy._core.raw.RawModule.get_function E mangled_name = self.module.mapping.get(name) E File "cupy/_core/raw.pyx", line 394, in cupy._core.raw.RawModule.module.__get__ E return self._module() E File "cupy/_core/raw.pyx", line 402, in cupy._core.raw.RawModule._module E mod = _get_raw_module( E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/raw.pyx", line 547, in cupy._core.raw._get_raw_module E mod = cupy._core.core.compile_with_cache( E File "cupy/_core/core.pyx", line 2015, in cupy._core.core.compile_with_cache E cpdef function.Module compile_with_cache( E File "cupy/_core/core.pyx", line 2078, in cupy._core.core.compile_with_cache E return cuda.compile_with_cache( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 453, in compile_with_cache E return _compile_with_cache_hip( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 853, in _compile_with_cache_hip E binary, mapping = compile_using_nvrtc( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 295, in compile_using_nvrtc E return _compile(source, options, cu_path, E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 279, in _compile E compiled_obj, mapping = prog.compile(options, log_stream) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 667, in compile E raise CompileException(log, self.src, self.name, options, E cupy.cuda.compiler.CompileException: /tmp/comgr-87b0bd/input/CompileSource:203:27: error: redeclaration of '__hiprtc_16' with a different type: 'void (*const)(int, int, int, const signed char *, const signed char *, signed char *)' vs 'void (*const)(int, int, int, const bool *, const bool *, bool *)' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-87b0bd/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-87b0bd/input/CompileSource:204:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-87b0bd/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-87b0bd/input/CompileSource:205:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-87b0bd/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-87b0bd/input/CompileSource:206:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-87b0bd/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-87b0bd/input/CompileSource:207:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-87b0bd/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-87b0bd/input/CompileSource:208:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-87b0bd/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-87b0bd/input/CompileSource:209:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-87b0bd/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-87b0bd/input/CompileSource:210:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-87b0bd/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-87b0bd/input/CompileSource:211:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-87b0bd/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-87b0bd/input/CompileSource:212:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-87b0bd/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E 10 errors generated when compiling for gfx908. E Error: Failed to compile opencl source (from CL or HIP source to LLVM IR). cupy/cuda/compiler.py:667: AssertionError ----------------------------- Captured stdout call ----------------------------- {'dtype_a': , 'dtype_b': } _ TestDot_param_69_{shape=((4, 2), (2,)), trans_a=True, trans_b=False}.test_dot _ self = < parameter: {'shape': ((4, 2), (2,)), 'trans_a': True, 'trans_b': False}> xp = dtype_a = , dtype_b = @testing.for_all_dtypes_combination(['dtype_a', 'dtype_b']) @testing.numpy_cupy_allclose() def test_dot(self, xp, dtype_a, dtype_b): shape_a, shape_b = self.shape if self.trans_a: a = testing.shaped_arange(shape_a[::-1], xp, dtype_a).T else: a = testing.shaped_arange(shape_a, xp, dtype_a) if self.trans_b: b = testing.shaped_arange(shape_b[::-1], xp, dtype_b).T else: b = testing.shaped_arange(shape_b, xp, dtype_b) > return xp.dot(a, b) tests/cupy_tests/linalg_tests/test_product.py:49: _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ cupy/linalg/_product.py:65: in dot return a.dot(b, out) cupy/_core/core.pyx:1607: in cupy._core.core.ndarray.dot return _linalg.dot(self, b, out) cupy/_core/_routines_linalg.pyx:427: in cupy._core._routines_linalg.dot return tensordot_core(a, b, out, n, m, k, ret_shape) cupy/_core/_routines_linalg.pyx:501: in cupy._core._routines_linalg.tensordot_core return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) cupy/_core/_routines_linalg.pyx:327: in cupy._core._routines_linalg._integral_tensordot_core kern = _tensordot_core_int_kernel(config, dtype) cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/_routines_linalg.pyx:302: in cupy._core._routines_linalg._tensordot_core_int_kernel ker = mod.get_function( cupy/_core/raw.pyx:470: in cupy._core.raw.RawModule.get_function mangled_name = self.module.mapping.get(name) cupy/_core/raw.pyx:394: in cupy._core.raw.RawModule.module.__get__ return self._module() cupy/_core/raw.pyx:402: in cupy._core.raw.RawModule._module mod = _get_raw_module( cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/raw.pyx:547: in cupy._core.raw._get_raw_module mod = cupy._core.core.compile_with_cache( cupy/_core/core.pyx:2015: in cupy._core.core.compile_with_cache cpdef function.Module compile_with_cache( cupy/_core/core.pyx:2078: in cupy._core.core.compile_with_cache return cuda.compile_with_cache( cupy/cuda/compiler.py:453: in compile_with_cache return _compile_with_cache_hip( cupy/cuda/compiler.py:853: in _compile_with_cache_hip binary, mapping = compile_using_nvrtc( cupy/cuda/compiler.py:295: in compile_using_nvrtc return _compile(source, options, cu_path, cupy/cuda/compiler.py:279: in _compile compiled_obj, mapping = prog.compile(options, log_stream) _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ self = options = ('--std=c++11', '-I/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/_core/include', '-I/opt/rocm/include', '-fcuda-flush-denormals-to-zero') log_stream = None def compile(self, options=(), log_stream=None): try: if self.name_expressions: for ker in self.name_expressions: nvrtc.addAddNameExpression(self.ptr, ker) nvrtc.compileProgram(self.ptr, options) mapping = None if self.name_expressions: mapping = {} for ker in self.name_expressions: mapping[ker] = nvrtc.getLoweredName(self.ptr, ker) if log_stream is not None: log_stream.write(nvrtc.getProgramLog(self.ptr)) # This is to ensure backwards compatibility with nvrtc if self.method == 'cubin': return nvrtc.getCUBIN(self.ptr), mapping elif self.method == 'ptx': return nvrtc.getPTX(self.ptr), mapping else: raise RuntimeError('Unknown NVRTC compile method') except nvrtc.NVRTCError: log = nvrtc.getProgramLog(self.ptr) > raise CompileException(log, self.src, self.name, options, 'nvrtc' if not runtime.is_hip else 'hiprtc') E AssertionError: Only cupy raises error E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 650, in compile E nvrtc.compileProgram(self.ptr, options) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 136, in cupy_backends.cuda.libs.nvrtc.compileProgram E cpdef compileProgram(intptr_t prog, options): E File "cupy_backends/cuda/libs/nvrtc.pyx", line 148, in cupy_backends.cuda.libs.nvrtc.compileProgram E check_status(status) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 67, in cupy_backends.cuda.libs.nvrtc.check_status E raise NVRTCError(status) E cupy_backends.cuda.libs.nvrtc.NVRTCError: HIPRTC_ERROR_COMPILATION (6) E E During handling of the above exception, another exception occurred: E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/testing/_helper.py", line 47, in _call_func E result = impl(*args, **kw) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/tests/cupy_tests/linalg_tests/test_product.py", line 49, in test_dot E return xp.dot(a, b) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/linalg/_product.py", line 65, in dot E return a.dot(b, out) E File "cupy/_core/core.pyx", line 1607, in cupy._core.core.ndarray.dot E return _linalg.dot(self, b, out) E File "cupy/_core/_routines_linalg.pyx", line 427, in cupy._core._routines_linalg.dot E return tensordot_core(a, b, out, n, m, k, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 501, in cupy._core._routines_linalg.tensordot_core E return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 327, in cupy._core._routines_linalg._integral_tensordot_core E kern = _tensordot_core_int_kernel(config, dtype) E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/_routines_linalg.pyx", line 302, in cupy._core._routines_linalg._tensordot_core_int_kernel E ker = mod.get_function( E File "cupy/_core/raw.pyx", line 470, in cupy._core.raw.RawModule.get_function E mangled_name = self.module.mapping.get(name) E File "cupy/_core/raw.pyx", line 394, in cupy._core.raw.RawModule.module.__get__ E return self._module() E File "cupy/_core/raw.pyx", line 402, in cupy._core.raw.RawModule._module E mod = _get_raw_module( E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/raw.pyx", line 547, in cupy._core.raw._get_raw_module E mod = cupy._core.core.compile_with_cache( E File "cupy/_core/core.pyx", line 2015, in cupy._core.core.compile_with_cache E cpdef function.Module compile_with_cache( E File "cupy/_core/core.pyx", line 2078, in cupy._core.core.compile_with_cache E return cuda.compile_with_cache( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 453, in compile_with_cache E return _compile_with_cache_hip( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 853, in _compile_with_cache_hip E binary, mapping = compile_using_nvrtc( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 295, in compile_using_nvrtc E return _compile(source, options, cu_path, E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 279, in _compile E compiled_obj, mapping = prog.compile(options, log_stream) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 667, in compile E raise CompileException(log, self.src, self.name, options, E cupy.cuda.compiler.CompileException: /tmp/comgr-3dbce3/input/CompileSource:203:27: error: redeclaration of '__hiprtc_16' with a different type: 'void (*const)(int, int, int, const signed char *, const signed char *, signed char *)' vs 'void (*const)(int, int, int, const bool *, const bool *, bool *)' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-3dbce3/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-3dbce3/input/CompileSource:204:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-3dbce3/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-3dbce3/input/CompileSource:205:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-3dbce3/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-3dbce3/input/CompileSource:206:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-3dbce3/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-3dbce3/input/CompileSource:207:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-3dbce3/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-3dbce3/input/CompileSource:208:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-3dbce3/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-3dbce3/input/CompileSource:209:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-3dbce3/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-3dbce3/input/CompileSource:210:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-3dbce3/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-3dbce3/input/CompileSource:211:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-3dbce3/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-3dbce3/input/CompileSource:212:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-3dbce3/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E 10 errors generated when compiling for gfx908. E Error: Failed to compile opencl source (from CL or HIP source to LLVM IR). cupy/cuda/compiler.py:667: AssertionError ----------------------------- Captured stdout call ----------------------------- {'dtype_a': , 'dtype_b': } _ TestDot_param_70_{shape=((4, 2), (2,)), trans_a=False, trans_b=True}.test_dot _ self = < parameter: {'shape': ((4, 2), (2,)), 'trans_a': False, 'trans_b': True}> xp = dtype_a = , dtype_b = @testing.for_all_dtypes_combination(['dtype_a', 'dtype_b']) @testing.numpy_cupy_allclose() def test_dot(self, xp, dtype_a, dtype_b): shape_a, shape_b = self.shape if self.trans_a: a = testing.shaped_arange(shape_a[::-1], xp, dtype_a).T else: a = testing.shaped_arange(shape_a, xp, dtype_a) if self.trans_b: b = testing.shaped_arange(shape_b[::-1], xp, dtype_b).T else: b = testing.shaped_arange(shape_b, xp, dtype_b) > return xp.dot(a, b) tests/cupy_tests/linalg_tests/test_product.py:49: _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ cupy/linalg/_product.py:65: in dot return a.dot(b, out) cupy/_core/core.pyx:1607: in cupy._core.core.ndarray.dot return _linalg.dot(self, b, out) cupy/_core/_routines_linalg.pyx:427: in cupy._core._routines_linalg.dot return tensordot_core(a, b, out, n, m, k, ret_shape) cupy/_core/_routines_linalg.pyx:501: in cupy._core._routines_linalg.tensordot_core return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) cupy/_core/_routines_linalg.pyx:327: in cupy._core._routines_linalg._integral_tensordot_core kern = _tensordot_core_int_kernel(config, dtype) cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/_routines_linalg.pyx:302: in cupy._core._routines_linalg._tensordot_core_int_kernel ker = mod.get_function( cupy/_core/raw.pyx:470: in cupy._core.raw.RawModule.get_function mangled_name = self.module.mapping.get(name) cupy/_core/raw.pyx:394: in cupy._core.raw.RawModule.module.__get__ return self._module() cupy/_core/raw.pyx:402: in cupy._core.raw.RawModule._module mod = _get_raw_module( cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/raw.pyx:547: in cupy._core.raw._get_raw_module mod = cupy._core.core.compile_with_cache( cupy/_core/core.pyx:2015: in cupy._core.core.compile_with_cache cpdef function.Module compile_with_cache( cupy/_core/core.pyx:2078: in cupy._core.core.compile_with_cache return cuda.compile_with_cache( cupy/cuda/compiler.py:453: in compile_with_cache return _compile_with_cache_hip( cupy/cuda/compiler.py:853: in _compile_with_cache_hip binary, mapping = compile_using_nvrtc( cupy/cuda/compiler.py:295: in compile_using_nvrtc return _compile(source, options, cu_path, cupy/cuda/compiler.py:279: in _compile compiled_obj, mapping = prog.compile(options, log_stream) _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ self = options = ('--std=c++11', '-I/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/_core/include', '-I/opt/rocm/include', '-fcuda-flush-denormals-to-zero') log_stream = None def compile(self, options=(), log_stream=None): try: if self.name_expressions: for ker in self.name_expressions: nvrtc.addAddNameExpression(self.ptr, ker) nvrtc.compileProgram(self.ptr, options) mapping = None if self.name_expressions: mapping = {} for ker in self.name_expressions: mapping[ker] = nvrtc.getLoweredName(self.ptr, ker) if log_stream is not None: log_stream.write(nvrtc.getProgramLog(self.ptr)) # This is to ensure backwards compatibility with nvrtc if self.method == 'cubin': return nvrtc.getCUBIN(self.ptr), mapping elif self.method == 'ptx': return nvrtc.getPTX(self.ptr), mapping else: raise RuntimeError('Unknown NVRTC compile method') except nvrtc.NVRTCError: log = nvrtc.getProgramLog(self.ptr) > raise CompileException(log, self.src, self.name, options, 'nvrtc' if not runtime.is_hip else 'hiprtc') E AssertionError: Only cupy raises error E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 650, in compile E nvrtc.compileProgram(self.ptr, options) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 136, in cupy_backends.cuda.libs.nvrtc.compileProgram E cpdef compileProgram(intptr_t prog, options): E File "cupy_backends/cuda/libs/nvrtc.pyx", line 148, in cupy_backends.cuda.libs.nvrtc.compileProgram E check_status(status) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 67, in cupy_backends.cuda.libs.nvrtc.check_status E raise NVRTCError(status) E cupy_backends.cuda.libs.nvrtc.NVRTCError: HIPRTC_ERROR_COMPILATION (6) E E During handling of the above exception, another exception occurred: E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/testing/_helper.py", line 47, in _call_func E result = impl(*args, **kw) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/tests/cupy_tests/linalg_tests/test_product.py", line 49, in test_dot E return xp.dot(a, b) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/linalg/_product.py", line 65, in dot E return a.dot(b, out) E File "cupy/_core/core.pyx", line 1607, in cupy._core.core.ndarray.dot E return _linalg.dot(self, b, out) E File "cupy/_core/_routines_linalg.pyx", line 427, in cupy._core._routines_linalg.dot E return tensordot_core(a, b, out, n, m, k, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 501, in cupy._core._routines_linalg.tensordot_core E return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 327, in cupy._core._routines_linalg._integral_tensordot_core E kern = _tensordot_core_int_kernel(config, dtype) E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/_routines_linalg.pyx", line 302, in cupy._core._routines_linalg._tensordot_core_int_kernel E ker = mod.get_function( E File "cupy/_core/raw.pyx", line 470, in cupy._core.raw.RawModule.get_function E mangled_name = self.module.mapping.get(name) E File "cupy/_core/raw.pyx", line 394, in cupy._core.raw.RawModule.module.__get__ E return self._module() E File "cupy/_core/raw.pyx", line 402, in cupy._core.raw.RawModule._module E mod = _get_raw_module( E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/raw.pyx", line 547, in cupy._core.raw._get_raw_module E mod = cupy._core.core.compile_with_cache( E File "cupy/_core/core.pyx", line 2015, in cupy._core.core.compile_with_cache E cpdef function.Module compile_with_cache( E File "cupy/_core/core.pyx", line 2078, in cupy._core.core.compile_with_cache E return cuda.compile_with_cache( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 453, in compile_with_cache E return _compile_with_cache_hip( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 853, in _compile_with_cache_hip E binary, mapping = compile_using_nvrtc( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 295, in compile_using_nvrtc E return _compile(source, options, cu_path, E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 279, in _compile E compiled_obj, mapping = prog.compile(options, log_stream) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 667, in compile E raise CompileException(log, self.src, self.name, options, E cupy.cuda.compiler.CompileException: /tmp/comgr-9e315e/input/CompileSource:203:27: error: redeclaration of '__hiprtc_16' with a different type: 'void (*const)(int, int, int, const signed char *, const signed char *, signed char *)' vs 'void (*const)(int, int, int, const bool *, const bool *, bool *)' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-9e315e/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-9e315e/input/CompileSource:204:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-9e315e/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-9e315e/input/CompileSource:205:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-9e315e/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-9e315e/input/CompileSource:206:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-9e315e/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-9e315e/input/CompileSource:207:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-9e315e/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-9e315e/input/CompileSource:208:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-9e315e/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-9e315e/input/CompileSource:209:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-9e315e/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-9e315e/input/CompileSource:210:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-9e315e/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-9e315e/input/CompileSource:211:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-9e315e/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-9e315e/input/CompileSource:212:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-9e315e/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E 10 errors generated when compiling for gfx908. E Error: Failed to compile opencl source (from CL or HIP source to LLVM IR). cupy/cuda/compiler.py:667: AssertionError ----------------------------- Captured stdout call ----------------------------- {'dtype_a': , 'dtype_b': } _ TestDot_param_71_{shape=((4, 2), (2,)), trans_a=False, trans_b=False}.test_dot _ self = < parameter: {'shape': ((4, 2), (2,)), 'trans_a': False, 'trans_b': False}> xp = dtype_a = , dtype_b = @testing.for_all_dtypes_combination(['dtype_a', 'dtype_b']) @testing.numpy_cupy_allclose() def test_dot(self, xp, dtype_a, dtype_b): shape_a, shape_b = self.shape if self.trans_a: a = testing.shaped_arange(shape_a[::-1], xp, dtype_a).T else: a = testing.shaped_arange(shape_a, xp, dtype_a) if self.trans_b: b = testing.shaped_arange(shape_b[::-1], xp, dtype_b).T else: b = testing.shaped_arange(shape_b, xp, dtype_b) > return xp.dot(a, b) tests/cupy_tests/linalg_tests/test_product.py:49: _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ cupy/linalg/_product.py:65: in dot return a.dot(b, out) cupy/_core/core.pyx:1607: in cupy._core.core.ndarray.dot return _linalg.dot(self, b, out) cupy/_core/_routines_linalg.pyx:427: in cupy._core._routines_linalg.dot return tensordot_core(a, b, out, n, m, k, ret_shape) cupy/_core/_routines_linalg.pyx:501: in cupy._core._routines_linalg.tensordot_core return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) cupy/_core/_routines_linalg.pyx:327: in cupy._core._routines_linalg._integral_tensordot_core kern = _tensordot_core_int_kernel(config, dtype) cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/_routines_linalg.pyx:302: in cupy._core._routines_linalg._tensordot_core_int_kernel ker = mod.get_function( cupy/_core/raw.pyx:470: in cupy._core.raw.RawModule.get_function mangled_name = self.module.mapping.get(name) cupy/_core/raw.pyx:394: in cupy._core.raw.RawModule.module.__get__ return self._module() cupy/_core/raw.pyx:402: in cupy._core.raw.RawModule._module mod = _get_raw_module( cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/raw.pyx:547: in cupy._core.raw._get_raw_module mod = cupy._core.core.compile_with_cache( cupy/_core/core.pyx:2015: in cupy._core.core.compile_with_cache cpdef function.Module compile_with_cache( cupy/_core/core.pyx:2078: in cupy._core.core.compile_with_cache return cuda.compile_with_cache( cupy/cuda/compiler.py:453: in compile_with_cache return _compile_with_cache_hip( cupy/cuda/compiler.py:853: in _compile_with_cache_hip binary, mapping = compile_using_nvrtc( cupy/cuda/compiler.py:295: in compile_using_nvrtc return _compile(source, options, cu_path, cupy/cuda/compiler.py:279: in _compile compiled_obj, mapping = prog.compile(options, log_stream) _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ self = options = ('--std=c++11', '-I/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/_core/include', '-I/opt/rocm/include', '-fcuda-flush-denormals-to-zero') log_stream = None def compile(self, options=(), log_stream=None): try: if self.name_expressions: for ker in self.name_expressions: nvrtc.addAddNameExpression(self.ptr, ker) nvrtc.compileProgram(self.ptr, options) mapping = None if self.name_expressions: mapping = {} for ker in self.name_expressions: mapping[ker] = nvrtc.getLoweredName(self.ptr, ker) if log_stream is not None: log_stream.write(nvrtc.getProgramLog(self.ptr)) # This is to ensure backwards compatibility with nvrtc if self.method == 'cubin': return nvrtc.getCUBIN(self.ptr), mapping elif self.method == 'ptx': return nvrtc.getPTX(self.ptr), mapping else: raise RuntimeError('Unknown NVRTC compile method') except nvrtc.NVRTCError: log = nvrtc.getProgramLog(self.ptr) > raise CompileException(log, self.src, self.name, options, 'nvrtc' if not runtime.is_hip else 'hiprtc') E AssertionError: Only cupy raises error E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 650, in compile E nvrtc.compileProgram(self.ptr, options) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 136, in cupy_backends.cuda.libs.nvrtc.compileProgram E cpdef compileProgram(intptr_t prog, options): E File "cupy_backends/cuda/libs/nvrtc.pyx", line 148, in cupy_backends.cuda.libs.nvrtc.compileProgram E check_status(status) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 67, in cupy_backends.cuda.libs.nvrtc.check_status E raise NVRTCError(status) E cupy_backends.cuda.libs.nvrtc.NVRTCError: HIPRTC_ERROR_COMPILATION (6) E E During handling of the above exception, another exception occurred: E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/testing/_helper.py", line 47, in _call_func E result = impl(*args, **kw) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/tests/cupy_tests/linalg_tests/test_product.py", line 49, in test_dot E return xp.dot(a, b) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/linalg/_product.py", line 65, in dot E return a.dot(b, out) E File "cupy/_core/core.pyx", line 1607, in cupy._core.core.ndarray.dot E return _linalg.dot(self, b, out) E File "cupy/_core/_routines_linalg.pyx", line 427, in cupy._core._routines_linalg.dot E return tensordot_core(a, b, out, n, m, k, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 501, in cupy._core._routines_linalg.tensordot_core E return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 327, in cupy._core._routines_linalg._integral_tensordot_core E kern = _tensordot_core_int_kernel(config, dtype) E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/_routines_linalg.pyx", line 302, in cupy._core._routines_linalg._tensordot_core_int_kernel E ker = mod.get_function( E File "cupy/_core/raw.pyx", line 470, in cupy._core.raw.RawModule.get_function E mangled_name = self.module.mapping.get(name) E File "cupy/_core/raw.pyx", line 394, in cupy._core.raw.RawModule.module.__get__ E return self._module() E File "cupy/_core/raw.pyx", line 402, in cupy._core.raw.RawModule._module E mod = _get_raw_module( E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/raw.pyx", line 547, in cupy._core.raw._get_raw_module E mod = cupy._core.core.compile_with_cache( E File "cupy/_core/core.pyx", line 2015, in cupy._core.core.compile_with_cache E cpdef function.Module compile_with_cache( E File "cupy/_core/core.pyx", line 2078, in cupy._core.core.compile_with_cache E return cuda.compile_with_cache( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 453, in compile_with_cache E return _compile_with_cache_hip( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 853, in _compile_with_cache_hip E binary, mapping = compile_using_nvrtc( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 295, in compile_using_nvrtc E return _compile(source, options, cu_path, E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 279, in _compile E compiled_obj, mapping = prog.compile(options, log_stream) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 667, in compile E raise CompileException(log, self.src, self.name, options, E cupy.cuda.compiler.CompileException: /tmp/comgr-0073d3/input/CompileSource:203:27: error: redeclaration of '__hiprtc_16' with a different type: 'void (*const)(int, int, int, const signed char *, const signed char *, signed char *)' vs 'void (*const)(int, int, int, const bool *, const bool *, bool *)' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-0073d3/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-0073d3/input/CompileSource:204:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-0073d3/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-0073d3/input/CompileSource:205:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-0073d3/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-0073d3/input/CompileSource:206:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-0073d3/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-0073d3/input/CompileSource:207:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-0073d3/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-0073d3/input/CompileSource:208:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-0073d3/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-0073d3/input/CompileSource:209:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-0073d3/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-0073d3/input/CompileSource:210:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-0073d3/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-0073d3/input/CompileSource:211:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-0073d3/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-0073d3/input/CompileSource:212:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-0073d3/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E 10 errors generated when compiling for gfx908. E Error: Failed to compile opencl source (from CL or HIP source to LLVM IR). cupy/cuda/compiler.py:667: AssertionError ----------------------------- Captured stdout call ----------------------------- {'dtype_a': , 'dtype_b': } _________________ TestProduct.test_dot_with_single_elem_array1 _________________ self = xp = dtype = @testing.for_all_dtypes() @testing.numpy_cupy_allclose() def test_dot_with_single_elem_array1(self, xp, dtype): a = testing.shaped_arange((3, 1), xp, dtype) b = xp.array([[2]], dtype=dtype) > return xp.dot(a, b) tests/cupy_tests/linalg_tests/test_product.py:189: _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ cupy/linalg/_product.py:65: in dot return a.dot(b, out) cupy/_core/core.pyx:1607: in cupy._core.core.ndarray.dot return _linalg.dot(self, b, out) cupy/_core/_routines_linalg.pyx:427: in cupy._core._routines_linalg.dot return tensordot_core(a, b, out, n, m, k, ret_shape) cupy/_core/_routines_linalg.pyx:501: in cupy._core._routines_linalg.tensordot_core return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) cupy/_core/_routines_linalg.pyx:327: in cupy._core._routines_linalg._integral_tensordot_core kern = _tensordot_core_int_kernel(config, dtype) cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/_routines_linalg.pyx:302: in cupy._core._routines_linalg._tensordot_core_int_kernel ker = mod.get_function( cupy/_core/raw.pyx:470: in cupy._core.raw.RawModule.get_function mangled_name = self.module.mapping.get(name) cupy/_core/raw.pyx:394: in cupy._core.raw.RawModule.module.__get__ return self._module() cupy/_core/raw.pyx:402: in cupy._core.raw.RawModule._module mod = _get_raw_module( cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/raw.pyx:547: in cupy._core.raw._get_raw_module mod = cupy._core.core.compile_with_cache( cupy/_core/core.pyx:2015: in cupy._core.core.compile_with_cache cpdef function.Module compile_with_cache( cupy/_core/core.pyx:2078: in cupy._core.core.compile_with_cache return cuda.compile_with_cache( cupy/cuda/compiler.py:453: in compile_with_cache return _compile_with_cache_hip( cupy/cuda/compiler.py:853: in _compile_with_cache_hip binary, mapping = compile_using_nvrtc( cupy/cuda/compiler.py:295: in compile_using_nvrtc return _compile(source, options, cu_path, cupy/cuda/compiler.py:279: in _compile compiled_obj, mapping = prog.compile(options, log_stream) _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ self = options = ('--std=c++11', '-I/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/_core/include', '-I/opt/rocm/include', '-fcuda-flush-denormals-to-zero') log_stream = None def compile(self, options=(), log_stream=None): try: if self.name_expressions: for ker in self.name_expressions: nvrtc.addAddNameExpression(self.ptr, ker) nvrtc.compileProgram(self.ptr, options) mapping = None if self.name_expressions: mapping = {} for ker in self.name_expressions: mapping[ker] = nvrtc.getLoweredName(self.ptr, ker) if log_stream is not None: log_stream.write(nvrtc.getProgramLog(self.ptr)) # This is to ensure backwards compatibility with nvrtc if self.method == 'cubin': return nvrtc.getCUBIN(self.ptr), mapping elif self.method == 'ptx': return nvrtc.getPTX(self.ptr), mapping else: raise RuntimeError('Unknown NVRTC compile method') except nvrtc.NVRTCError: log = nvrtc.getProgramLog(self.ptr) > raise CompileException(log, self.src, self.name, options, 'nvrtc' if not runtime.is_hip else 'hiprtc') E AssertionError: Only cupy raises error E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 650, in compile E nvrtc.compileProgram(self.ptr, options) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 136, in cupy_backends.cuda.libs.nvrtc.compileProgram E cpdef compileProgram(intptr_t prog, options): E File "cupy_backends/cuda/libs/nvrtc.pyx", line 148, in cupy_backends.cuda.libs.nvrtc.compileProgram E check_status(status) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 67, in cupy_backends.cuda.libs.nvrtc.check_status E raise NVRTCError(status) E cupy_backends.cuda.libs.nvrtc.NVRTCError: HIPRTC_ERROR_COMPILATION (6) E E During handling of the above exception, another exception occurred: E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/testing/_helper.py", line 47, in _call_func E result = impl(*args, **kw) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/tests/cupy_tests/linalg_tests/test_product.py", line 189, in test_dot_with_single_elem_array1 E return xp.dot(a, b) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/linalg/_product.py", line 65, in dot E return a.dot(b, out) E File "cupy/_core/core.pyx", line 1607, in cupy._core.core.ndarray.dot E return _linalg.dot(self, b, out) E File "cupy/_core/_routines_linalg.pyx", line 427, in cupy._core._routines_linalg.dot E return tensordot_core(a, b, out, n, m, k, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 501, in cupy._core._routines_linalg.tensordot_core E return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 327, in cupy._core._routines_linalg._integral_tensordot_core E kern = _tensordot_core_int_kernel(config, dtype) E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/_routines_linalg.pyx", line 302, in cupy._core._routines_linalg._tensordot_core_int_kernel E ker = mod.get_function( E File "cupy/_core/raw.pyx", line 470, in cupy._core.raw.RawModule.get_function E mangled_name = self.module.mapping.get(name) E File "cupy/_core/raw.pyx", line 394, in cupy._core.raw.RawModule.module.__get__ E return self._module() E File "cupy/_core/raw.pyx", line 402, in cupy._core.raw.RawModule._module E mod = _get_raw_module( E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/raw.pyx", line 547, in cupy._core.raw._get_raw_module E mod = cupy._core.core.compile_with_cache( E File "cupy/_core/core.pyx", line 2015, in cupy._core.core.compile_with_cache E cpdef function.Module compile_with_cache( E File "cupy/_core/core.pyx", line 2078, in cupy._core.core.compile_with_cache E return cuda.compile_with_cache( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 453, in compile_with_cache E return _compile_with_cache_hip( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 853, in _compile_with_cache_hip E binary, mapping = compile_using_nvrtc( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 295, in compile_using_nvrtc E return _compile(source, options, cu_path, E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 279, in _compile E compiled_obj, mapping = prog.compile(options, log_stream) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 667, in compile E raise CompileException(log, self.src, self.name, options, E cupy.cuda.compiler.CompileException: /tmp/comgr-3b6f3d/input/CompileSource:203:27: error: redeclaration of '__hiprtc_16' with a different type: 'void (*const)(int, int, int, const signed char *, const signed char *, signed char *)' vs 'void (*const)(int, int, int, const bool *, const bool *, bool *)' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-3b6f3d/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-3b6f3d/input/CompileSource:204:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-3b6f3d/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-3b6f3d/input/CompileSource:205:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-3b6f3d/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-3b6f3d/input/CompileSource:206:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-3b6f3d/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-3b6f3d/input/CompileSource:207:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-3b6f3d/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-3b6f3d/input/CompileSource:208:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-3b6f3d/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-3b6f3d/input/CompileSource:209:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-3b6f3d/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-3b6f3d/input/CompileSource:210:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-3b6f3d/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-3b6f3d/input/CompileSource:211:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-3b6f3d/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-3b6f3d/input/CompileSource:212:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-3b6f3d/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E 10 errors generated when compiling for gfx908. E Error: Failed to compile opencl source (from CL or HIP source to LLVM IR). cupy/cuda/compiler.py:667: AssertionError ----------------------------- Captured stdout call ----------------------------- dtype is _________________ TestProduct.test_dot_with_single_elem_array2 _________________ self = xp = dtype = @testing.for_all_dtypes() @testing.numpy_cupy_allclose() def test_dot_with_single_elem_array2(self, xp, dtype): a = xp.array([[2]], dtype=dtype) b = testing.shaped_arange((1, 3), xp, dtype) > return xp.dot(a, b) tests/cupy_tests/linalg_tests/test_product.py:196: _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ cupy/linalg/_product.py:65: in dot return a.dot(b, out) cupy/_core/core.pyx:1607: in cupy._core.core.ndarray.dot return _linalg.dot(self, b, out) cupy/_core/_routines_linalg.pyx:427: in cupy._core._routines_linalg.dot return tensordot_core(a, b, out, n, m, k, ret_shape) cupy/_core/_routines_linalg.pyx:501: in cupy._core._routines_linalg.tensordot_core return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) cupy/_core/_routines_linalg.pyx:327: in cupy._core._routines_linalg._integral_tensordot_core kern = _tensordot_core_int_kernel(config, dtype) cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/_routines_linalg.pyx:302: in cupy._core._routines_linalg._tensordot_core_int_kernel ker = mod.get_function( cupy/_core/raw.pyx:470: in cupy._core.raw.RawModule.get_function mangled_name = self.module.mapping.get(name) cupy/_core/raw.pyx:394: in cupy._core.raw.RawModule.module.__get__ return self._module() cupy/_core/raw.pyx:402: in cupy._core.raw.RawModule._module mod = _get_raw_module( cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/raw.pyx:547: in cupy._core.raw._get_raw_module mod = cupy._core.core.compile_with_cache( cupy/_core/core.pyx:2015: in cupy._core.core.compile_with_cache cpdef function.Module compile_with_cache( cupy/_core/core.pyx:2078: in cupy._core.core.compile_with_cache return cuda.compile_with_cache( cupy/cuda/compiler.py:453: in compile_with_cache return _compile_with_cache_hip( cupy/cuda/compiler.py:853: in _compile_with_cache_hip binary, mapping = compile_using_nvrtc( cupy/cuda/compiler.py:295: in compile_using_nvrtc return _compile(source, options, cu_path, cupy/cuda/compiler.py:279: in _compile compiled_obj, mapping = prog.compile(options, log_stream) _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ self = options = ('--std=c++11', '-I/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/_core/include', '-I/opt/rocm/include', '-fcuda-flush-denormals-to-zero') log_stream = None def compile(self, options=(), log_stream=None): try: if self.name_expressions: for ker in self.name_expressions: nvrtc.addAddNameExpression(self.ptr, ker) nvrtc.compileProgram(self.ptr, options) mapping = None if self.name_expressions: mapping = {} for ker in self.name_expressions: mapping[ker] = nvrtc.getLoweredName(self.ptr, ker) if log_stream is not None: log_stream.write(nvrtc.getProgramLog(self.ptr)) # This is to ensure backwards compatibility with nvrtc if self.method == 'cubin': return nvrtc.getCUBIN(self.ptr), mapping elif self.method == 'ptx': return nvrtc.getPTX(self.ptr), mapping else: raise RuntimeError('Unknown NVRTC compile method') except nvrtc.NVRTCError: log = nvrtc.getProgramLog(self.ptr) > raise CompileException(log, self.src, self.name, options, 'nvrtc' if not runtime.is_hip else 'hiprtc') E AssertionError: Only cupy raises error E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 650, in compile E nvrtc.compileProgram(self.ptr, options) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 136, in cupy_backends.cuda.libs.nvrtc.compileProgram E cpdef compileProgram(intptr_t prog, options): E File "cupy_backends/cuda/libs/nvrtc.pyx", line 148, in cupy_backends.cuda.libs.nvrtc.compileProgram E check_status(status) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 67, in cupy_backends.cuda.libs.nvrtc.check_status E raise NVRTCError(status) E cupy_backends.cuda.libs.nvrtc.NVRTCError: HIPRTC_ERROR_COMPILATION (6) E E During handling of the above exception, another exception occurred: E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/testing/_helper.py", line 47, in _call_func E result = impl(*args, **kw) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/tests/cupy_tests/linalg_tests/test_product.py", line 196, in test_dot_with_single_elem_array2 E return xp.dot(a, b) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/linalg/_product.py", line 65, in dot E return a.dot(b, out) E File "cupy/_core/core.pyx", line 1607, in cupy._core.core.ndarray.dot E return _linalg.dot(self, b, out) E File "cupy/_core/_routines_linalg.pyx", line 427, in cupy._core._routines_linalg.dot E return tensordot_core(a, b, out, n, m, k, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 501, in cupy._core._routines_linalg.tensordot_core E return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 327, in cupy._core._routines_linalg._integral_tensordot_core E kern = _tensordot_core_int_kernel(config, dtype) E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/_routines_linalg.pyx", line 302, in cupy._core._routines_linalg._tensordot_core_int_kernel E ker = mod.get_function( E File "cupy/_core/raw.pyx", line 470, in cupy._core.raw.RawModule.get_function E mangled_name = self.module.mapping.get(name) E File "cupy/_core/raw.pyx", line 394, in cupy._core.raw.RawModule.module.__get__ E return self._module() E File "cupy/_core/raw.pyx", line 402, in cupy._core.raw.RawModule._module E mod = _get_raw_module( E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/raw.pyx", line 547, in cupy._core.raw._get_raw_module E mod = cupy._core.core.compile_with_cache( E File "cupy/_core/core.pyx", line 2015, in cupy._core.core.compile_with_cache E cpdef function.Module compile_with_cache( E File "cupy/_core/core.pyx", line 2078, in cupy._core.core.compile_with_cache E return cuda.compile_with_cache( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 453, in compile_with_cache E return _compile_with_cache_hip( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 853, in _compile_with_cache_hip E binary, mapping = compile_using_nvrtc( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 295, in compile_using_nvrtc E return _compile(source, options, cu_path, E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 279, in _compile E compiled_obj, mapping = prog.compile(options, log_stream) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 667, in compile E raise CompileException(log, self.src, self.name, options, E cupy.cuda.compiler.CompileException: /tmp/comgr-950dc1/input/CompileSource:203:27: error: redeclaration of '__hiprtc_16' with a different type: 'void (*const)(int, int, int, const signed char *, const signed char *, signed char *)' vs 'void (*const)(int, int, int, const bool *, const bool *, bool *)' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-950dc1/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-950dc1/input/CompileSource:204:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-950dc1/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-950dc1/input/CompileSource:205:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-950dc1/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-950dc1/input/CompileSource:206:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-950dc1/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-950dc1/input/CompileSource:207:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-950dc1/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-950dc1/input/CompileSource:208:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-950dc1/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-950dc1/input/CompileSource:209:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-950dc1/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-950dc1/input/CompileSource:210:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-950dc1/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-950dc1/input/CompileSource:211:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-950dc1/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-950dc1/input/CompileSource:212:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-950dc1/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E 10 errors generated when compiling for gfx908. E Error: Failed to compile opencl source (from CL or HIP source to LLVM IR). cupy/cuda/compiler.py:667: AssertionError ----------------------------- Captured stdout call ----------------------------- dtype is ____________________________ TestProduct.test_kron _____________________________ self = xp = dtype = @testing.for_all_dtypes() @testing.numpy_cupy_allclose() def test_kron(self, xp, dtype): a = testing.shaped_arange((4,), xp, dtype) b = testing.shaped_arange((5,), xp, dtype) > return xp.kron(a, b) tests/cupy_tests/linalg_tests/test_product.py:361: _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ cupy/linalg/_product.py:426: in kron out = _core.tensordot_core( cupy/_core/_routines_linalg.pyx:430: in cupy._core._routines_linalg.tensordot_core cpdef ndarray tensordot_core( cupy/_core/_routines_linalg.pyx:501: in cupy._core._routines_linalg.tensordot_core return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) cupy/_core/_routines_linalg.pyx:327: in cupy._core._routines_linalg._integral_tensordot_core kern = _tensordot_core_int_kernel(config, dtype) cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/_routines_linalg.pyx:302: in cupy._core._routines_linalg._tensordot_core_int_kernel ker = mod.get_function( cupy/_core/raw.pyx:470: in cupy._core.raw.RawModule.get_function mangled_name = self.module.mapping.get(name) cupy/_core/raw.pyx:394: in cupy._core.raw.RawModule.module.__get__ return self._module() cupy/_core/raw.pyx:402: in cupy._core.raw.RawModule._module mod = _get_raw_module( cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/raw.pyx:547: in cupy._core.raw._get_raw_module mod = cupy._core.core.compile_with_cache( cupy/_core/core.pyx:2015: in cupy._core.core.compile_with_cache cpdef function.Module compile_with_cache( cupy/_core/core.pyx:2078: in cupy._core.core.compile_with_cache return cuda.compile_with_cache( cupy/cuda/compiler.py:453: in compile_with_cache return _compile_with_cache_hip( cupy/cuda/compiler.py:853: in _compile_with_cache_hip binary, mapping = compile_using_nvrtc( cupy/cuda/compiler.py:295: in compile_using_nvrtc return _compile(source, options, cu_path, cupy/cuda/compiler.py:279: in _compile compiled_obj, mapping = prog.compile(options, log_stream) _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ self = options = ('--std=c++11', '-I/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/_core/include', '-I/opt/rocm/include', '-fcuda-flush-denormals-to-zero') log_stream = None def compile(self, options=(), log_stream=None): try: if self.name_expressions: for ker in self.name_expressions: nvrtc.addAddNameExpression(self.ptr, ker) nvrtc.compileProgram(self.ptr, options) mapping = None if self.name_expressions: mapping = {} for ker in self.name_expressions: mapping[ker] = nvrtc.getLoweredName(self.ptr, ker) if log_stream is not None: log_stream.write(nvrtc.getProgramLog(self.ptr)) # This is to ensure backwards compatibility with nvrtc if self.method == 'cubin': return nvrtc.getCUBIN(self.ptr), mapping elif self.method == 'ptx': return nvrtc.getPTX(self.ptr), mapping else: raise RuntimeError('Unknown NVRTC compile method') except nvrtc.NVRTCError: log = nvrtc.getProgramLog(self.ptr) > raise CompileException(log, self.src, self.name, options, 'nvrtc' if not runtime.is_hip else 'hiprtc') E AssertionError: Only cupy raises error E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 650, in compile E nvrtc.compileProgram(self.ptr, options) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 136, in cupy_backends.cuda.libs.nvrtc.compileProgram E cpdef compileProgram(intptr_t prog, options): E File "cupy_backends/cuda/libs/nvrtc.pyx", line 148, in cupy_backends.cuda.libs.nvrtc.compileProgram E check_status(status) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 67, in cupy_backends.cuda.libs.nvrtc.check_status E raise NVRTCError(status) E cupy_backends.cuda.libs.nvrtc.NVRTCError: HIPRTC_ERROR_COMPILATION (6) E E During handling of the above exception, another exception occurred: E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/testing/_helper.py", line 47, in _call_func E result = impl(*args, **kw) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/tests/cupy_tests/linalg_tests/test_product.py", line 361, in test_kron E return xp.kron(a, b) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/linalg/_product.py", line 426, in kron E out = _core.tensordot_core( E File "cupy/_core/_routines_linalg.pyx", line 430, in cupy._core._routines_linalg.tensordot_core E cpdef ndarray tensordot_core( E File "cupy/_core/_routines_linalg.pyx", line 501, in cupy._core._routines_linalg.tensordot_core E return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 327, in cupy._core._routines_linalg._integral_tensordot_core E kern = _tensordot_core_int_kernel(config, dtype) E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/_routines_linalg.pyx", line 302, in cupy._core._routines_linalg._tensordot_core_int_kernel E ker = mod.get_function( E File "cupy/_core/raw.pyx", line 470, in cupy._core.raw.RawModule.get_function E mangled_name = self.module.mapping.get(name) E File "cupy/_core/raw.pyx", line 394, in cupy._core.raw.RawModule.module.__get__ E return self._module() E File "cupy/_core/raw.pyx", line 402, in cupy._core.raw.RawModule._module E mod = _get_raw_module( E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/raw.pyx", line 547, in cupy._core.raw._get_raw_module E mod = cupy._core.core.compile_with_cache( E File "cupy/_core/core.pyx", line 2015, in cupy._core.core.compile_with_cache E cpdef function.Module compile_with_cache( E File "cupy/_core/core.pyx", line 2078, in cupy._core.core.compile_with_cache E return cuda.compile_with_cache( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 453, in compile_with_cache E return _compile_with_cache_hip( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 853, in _compile_with_cache_hip E binary, mapping = compile_using_nvrtc( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 295, in compile_using_nvrtc E return _compile(source, options, cu_path, E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 279, in _compile E compiled_obj, mapping = prog.compile(options, log_stream) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 667, in compile E raise CompileException(log, self.src, self.name, options, E cupy.cuda.compiler.CompileException: /tmp/comgr-784a4a/input/CompileSource:203:27: error: redeclaration of '__hiprtc_16' with a different type: 'void (*const)(int, int, int, const signed char *, const signed char *, signed char *)' vs 'void (*const)(int, int, int, const bool *, const bool *, bool *)' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-784a4a/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-784a4a/input/CompileSource:204:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-784a4a/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-784a4a/input/CompileSource:205:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-784a4a/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-784a4a/input/CompileSource:206:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-784a4a/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-784a4a/input/CompileSource:207:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-784a4a/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-784a4a/input/CompileSource:208:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-784a4a/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-784a4a/input/CompileSource:209:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-784a4a/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-784a4a/input/CompileSource:210:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-784a4a/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-784a4a/input/CompileSource:211:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-784a4a/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-784a4a/input/CompileSource:212:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-784a4a/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E 10 errors generated when compiling for gfx908. E Error: Failed to compile opencl source (from CL or HIP source to LLVM IR). cupy/cuda/compiler.py:667: AssertionError ----------------------------- Captured stdout call ----------------------------- dtype is _______________________ TestProduct.test_multidim_inner ________________________ self = xp = dtype = @testing.for_all_dtypes() @testing.numpy_cupy_allclose() def test_multidim_inner(self, xp, dtype): a = testing.shaped_arange((2, 3, 4), xp, dtype) b = testing.shaped_arange((3, 2, 4), xp, dtype) > return xp.inner(a, b) tests/cupy_tests/linalg_tests/test_product.py:246: _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ cupy/linalg/_product.py:251: in inner return _core.tensordot_core(a, b, None, n, m, k, ret_shape) cupy/_core/_routines_linalg.pyx:430: in cupy._core._routines_linalg.tensordot_core cpdef ndarray tensordot_core( cupy/_core/_routines_linalg.pyx:501: in cupy._core._routines_linalg.tensordot_core return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) cupy/_core/_routines_linalg.pyx:327: in cupy._core._routines_linalg._integral_tensordot_core kern = _tensordot_core_int_kernel(config, dtype) cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/_routines_linalg.pyx:302: in cupy._core._routines_linalg._tensordot_core_int_kernel ker = mod.get_function( cupy/_core/raw.pyx:470: in cupy._core.raw.RawModule.get_function mangled_name = self.module.mapping.get(name) cupy/_core/raw.pyx:394: in cupy._core.raw.RawModule.module.__get__ return self._module() cupy/_core/raw.pyx:402: in cupy._core.raw.RawModule._module mod = _get_raw_module( cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/raw.pyx:547: in cupy._core.raw._get_raw_module mod = cupy._core.core.compile_with_cache( cupy/_core/core.pyx:2015: in cupy._core.core.compile_with_cache cpdef function.Module compile_with_cache( cupy/_core/core.pyx:2078: in cupy._core.core.compile_with_cache return cuda.compile_with_cache( cupy/cuda/compiler.py:453: in compile_with_cache return _compile_with_cache_hip( cupy/cuda/compiler.py:853: in _compile_with_cache_hip binary, mapping = compile_using_nvrtc( cupy/cuda/compiler.py:295: in compile_using_nvrtc return _compile(source, options, cu_path, cupy/cuda/compiler.py:279: in _compile compiled_obj, mapping = prog.compile(options, log_stream) _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ self = options = ('--std=c++11', '-I/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/_core/include', '-I/opt/rocm/include', '-fcuda-flush-denormals-to-zero') log_stream = None def compile(self, options=(), log_stream=None): try: if self.name_expressions: for ker in self.name_expressions: nvrtc.addAddNameExpression(self.ptr, ker) nvrtc.compileProgram(self.ptr, options) mapping = None if self.name_expressions: mapping = {} for ker in self.name_expressions: mapping[ker] = nvrtc.getLoweredName(self.ptr, ker) if log_stream is not None: log_stream.write(nvrtc.getProgramLog(self.ptr)) # This is to ensure backwards compatibility with nvrtc if self.method == 'cubin': return nvrtc.getCUBIN(self.ptr), mapping elif self.method == 'ptx': return nvrtc.getPTX(self.ptr), mapping else: raise RuntimeError('Unknown NVRTC compile method') except nvrtc.NVRTCError: log = nvrtc.getProgramLog(self.ptr) > raise CompileException(log, self.src, self.name, options, 'nvrtc' if not runtime.is_hip else 'hiprtc') E AssertionError: Only cupy raises error E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 650, in compile E nvrtc.compileProgram(self.ptr, options) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 136, in cupy_backends.cuda.libs.nvrtc.compileProgram E cpdef compileProgram(intptr_t prog, options): E File "cupy_backends/cuda/libs/nvrtc.pyx", line 148, in cupy_backends.cuda.libs.nvrtc.compileProgram E check_status(status) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 67, in cupy_backends.cuda.libs.nvrtc.check_status E raise NVRTCError(status) E cupy_backends.cuda.libs.nvrtc.NVRTCError: HIPRTC_ERROR_COMPILATION (6) E E During handling of the above exception, another exception occurred: E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/testing/_helper.py", line 47, in _call_func E result = impl(*args, **kw) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/tests/cupy_tests/linalg_tests/test_product.py", line 246, in test_multidim_inner E return xp.inner(a, b) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/linalg/_product.py", line 251, in inner E return _core.tensordot_core(a, b, None, n, m, k, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 430, in cupy._core._routines_linalg.tensordot_core E cpdef ndarray tensordot_core( E File "cupy/_core/_routines_linalg.pyx", line 501, in cupy._core._routines_linalg.tensordot_core E return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 327, in cupy._core._routines_linalg._integral_tensordot_core E kern = _tensordot_core_int_kernel(config, dtype) E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/_routines_linalg.pyx", line 302, in cupy._core._routines_linalg._tensordot_core_int_kernel E ker = mod.get_function( E File "cupy/_core/raw.pyx", line 470, in cupy._core.raw.RawModule.get_function E mangled_name = self.module.mapping.get(name) E File "cupy/_core/raw.pyx", line 394, in cupy._core.raw.RawModule.module.__get__ E return self._module() E File "cupy/_core/raw.pyx", line 402, in cupy._core.raw.RawModule._module E mod = _get_raw_module( E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/raw.pyx", line 547, in cupy._core.raw._get_raw_module E mod = cupy._core.core.compile_with_cache( E File "cupy/_core/core.pyx", line 2015, in cupy._core.core.compile_with_cache E cpdef function.Module compile_with_cache( E File "cupy/_core/core.pyx", line 2078, in cupy._core.core.compile_with_cache E return cuda.compile_with_cache( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 453, in compile_with_cache E return _compile_with_cache_hip( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 853, in _compile_with_cache_hip E binary, mapping = compile_using_nvrtc( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 295, in compile_using_nvrtc E return _compile(source, options, cu_path, E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 279, in _compile E compiled_obj, mapping = prog.compile(options, log_stream) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 667, in compile E raise CompileException(log, self.src, self.name, options, E cupy.cuda.compiler.CompileException: /tmp/comgr-d667b6/input/CompileSource:203:27: error: redeclaration of '__hiprtc_16' with a different type: 'void (*const)(int, int, int, const signed char *, const signed char *, signed char *)' vs 'void (*const)(int, int, int, const bool *, const bool *, bool *)' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-d667b6/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-d667b6/input/CompileSource:204:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-d667b6/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-d667b6/input/CompileSource:205:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-d667b6/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-d667b6/input/CompileSource:206:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-d667b6/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-d667b6/input/CompileSource:207:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-d667b6/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-d667b6/input/CompileSource:208:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-d667b6/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-d667b6/input/CompileSource:209:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-d667b6/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-d667b6/input/CompileSource:210:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-d667b6/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-d667b6/input/CompileSource:211:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-d667b6/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-d667b6/input/CompileSource:212:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-d667b6/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E 10 errors generated when compiling for gfx908. E Error: Failed to compile opencl source (from CL or HIP source to LLVM IR). cupy/cuda/compiler.py:667: AssertionError ----------------------------- Captured stdout call ----------------------------- dtype is ________________________ TestProduct.test_multidim_kron ________________________ self = xp = dtype = @testing.for_all_dtypes() @testing.numpy_cupy_allclose() def test_multidim_kron(self, xp, dtype): a = testing.shaped_arange((2, 3, 4), xp, dtype) b = testing.shaped_arange((4, 2, 3), xp, dtype) > return xp.kron(a, b) tests/cupy_tests/linalg_tests/test_product.py:375: _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ cupy/linalg/_product.py:426: in kron out = _core.tensordot_core( cupy/_core/_routines_linalg.pyx:430: in cupy._core._routines_linalg.tensordot_core cpdef ndarray tensordot_core( cupy/_core/_routines_linalg.pyx:501: in cupy._core._routines_linalg.tensordot_core return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) cupy/_core/_routines_linalg.pyx:327: in cupy._core._routines_linalg._integral_tensordot_core kern = _tensordot_core_int_kernel(config, dtype) cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/_routines_linalg.pyx:302: in cupy._core._routines_linalg._tensordot_core_int_kernel ker = mod.get_function( cupy/_core/raw.pyx:470: in cupy._core.raw.RawModule.get_function mangled_name = self.module.mapping.get(name) cupy/_core/raw.pyx:394: in cupy._core.raw.RawModule.module.__get__ return self._module() cupy/_core/raw.pyx:402: in cupy._core.raw.RawModule._module mod = _get_raw_module( cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/raw.pyx:547: in cupy._core.raw._get_raw_module mod = cupy._core.core.compile_with_cache( cupy/_core/core.pyx:2015: in cupy._core.core.compile_with_cache cpdef function.Module compile_with_cache( cupy/_core/core.pyx:2078: in cupy._core.core.compile_with_cache return cuda.compile_with_cache( cupy/cuda/compiler.py:453: in compile_with_cache return _compile_with_cache_hip( cupy/cuda/compiler.py:853: in _compile_with_cache_hip binary, mapping = compile_using_nvrtc( cupy/cuda/compiler.py:295: in compile_using_nvrtc return _compile(source, options, cu_path, cupy/cuda/compiler.py:279: in _compile compiled_obj, mapping = prog.compile(options, log_stream) _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ self = options = ('--std=c++11', '-I/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/_core/include', '-I/opt/rocm/include', '-fcuda-flush-denormals-to-zero') log_stream = None def compile(self, options=(), log_stream=None): try: if self.name_expressions: for ker in self.name_expressions: nvrtc.addAddNameExpression(self.ptr, ker) nvrtc.compileProgram(self.ptr, options) mapping = None if self.name_expressions: mapping = {} for ker in self.name_expressions: mapping[ker] = nvrtc.getLoweredName(self.ptr, ker) if log_stream is not None: log_stream.write(nvrtc.getProgramLog(self.ptr)) # This is to ensure backwards compatibility with nvrtc if self.method == 'cubin': return nvrtc.getCUBIN(self.ptr), mapping elif self.method == 'ptx': return nvrtc.getPTX(self.ptr), mapping else: raise RuntimeError('Unknown NVRTC compile method') except nvrtc.NVRTCError: log = nvrtc.getProgramLog(self.ptr) > raise CompileException(log, self.src, self.name, options, 'nvrtc' if not runtime.is_hip else 'hiprtc') E AssertionError: Only cupy raises error E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 650, in compile E nvrtc.compileProgram(self.ptr, options) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 136, in cupy_backends.cuda.libs.nvrtc.compileProgram E cpdef compileProgram(intptr_t prog, options): E File "cupy_backends/cuda/libs/nvrtc.pyx", line 148, in cupy_backends.cuda.libs.nvrtc.compileProgram E check_status(status) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 67, in cupy_backends.cuda.libs.nvrtc.check_status E raise NVRTCError(status) E cupy_backends.cuda.libs.nvrtc.NVRTCError: HIPRTC_ERROR_COMPILATION (6) E E During handling of the above exception, another exception occurred: E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/testing/_helper.py", line 47, in _call_func E result = impl(*args, **kw) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/tests/cupy_tests/linalg_tests/test_product.py", line 375, in test_multidim_kron E return xp.kron(a, b) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/linalg/_product.py", line 426, in kron E out = _core.tensordot_core( E File "cupy/_core/_routines_linalg.pyx", line 430, in cupy._core._routines_linalg.tensordot_core E cpdef ndarray tensordot_core( E File "cupy/_core/_routines_linalg.pyx", line 501, in cupy._core._routines_linalg.tensordot_core E return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 327, in cupy._core._routines_linalg._integral_tensordot_core E kern = _tensordot_core_int_kernel(config, dtype) E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/_routines_linalg.pyx", line 302, in cupy._core._routines_linalg._tensordot_core_int_kernel E ker = mod.get_function( E File "cupy/_core/raw.pyx", line 470, in cupy._core.raw.RawModule.get_function E mangled_name = self.module.mapping.get(name) E File "cupy/_core/raw.pyx", line 394, in cupy._core.raw.RawModule.module.__get__ E return self._module() E File "cupy/_core/raw.pyx", line 402, in cupy._core.raw.RawModule._module E mod = _get_raw_module( E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/raw.pyx", line 547, in cupy._core.raw._get_raw_module E mod = cupy._core.core.compile_with_cache( E File "cupy/_core/core.pyx", line 2015, in cupy._core.core.compile_with_cache E cpdef function.Module compile_with_cache( E File "cupy/_core/core.pyx", line 2078, in cupy._core.core.compile_with_cache E return cuda.compile_with_cache( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 453, in compile_with_cache E return _compile_with_cache_hip( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 853, in _compile_with_cache_hip E binary, mapping = compile_using_nvrtc( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 295, in compile_using_nvrtc E return _compile(source, options, cu_path, E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 279, in _compile E compiled_obj, mapping = prog.compile(options, log_stream) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 667, in compile E raise CompileException(log, self.src, self.name, options, E cupy.cuda.compiler.CompileException: /tmp/comgr-583c17/input/CompileSource:203:27: error: redeclaration of '__hiprtc_16' with a different type: 'void (*const)(int, int, int, const signed char *, const signed char *, signed char *)' vs 'void (*const)(int, int, int, const bool *, const bool *, bool *)' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-583c17/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-583c17/input/CompileSource:204:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-583c17/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-583c17/input/CompileSource:205:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-583c17/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-583c17/input/CompileSource:206:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-583c17/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-583c17/input/CompileSource:207:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-583c17/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-583c17/input/CompileSource:208:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-583c17/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-583c17/input/CompileSource:209:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-583c17/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-583c17/input/CompileSource:210:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-583c17/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-583c17/input/CompileSource:211:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-583c17/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-583c17/input/CompileSource:212:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-583c17/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E 10 errors generated when compiling for gfx908. E Error: Failed to compile opencl source (from CL or HIP source to LLVM IR). cupy/cuda/compiler.py:667: AssertionError ----------------------------- Captured stdout call ----------------------------- dtype is _______________________ TestProduct.test_multidim_outer ________________________ self = xp = dtype = @testing.for_all_dtypes() @testing.numpy_cupy_allclose() def test_multidim_outer(self, xp, dtype): a = testing.shaped_arange((2, 3), xp, dtype) b = testing.shaped_arange((4, 5), xp, dtype) > return xp.outer(a, b) tests/cupy_tests/linalg_tests/test_product.py:274: _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ cupy/linalg/_product.py:276: in outer return _core.tensordot_core(a, b, None, n, m, 1, ret_shape) cupy/_core/_routines_linalg.pyx:430: in cupy._core._routines_linalg.tensordot_core cpdef ndarray tensordot_core( cupy/_core/_routines_linalg.pyx:501: in cupy._core._routines_linalg.tensordot_core return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) cupy/_core/_routines_linalg.pyx:327: in cupy._core._routines_linalg._integral_tensordot_core kern = _tensordot_core_int_kernel(config, dtype) cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/_routines_linalg.pyx:302: in cupy._core._routines_linalg._tensordot_core_int_kernel ker = mod.get_function( cupy/_core/raw.pyx:470: in cupy._core.raw.RawModule.get_function mangled_name = self.module.mapping.get(name) cupy/_core/raw.pyx:394: in cupy._core.raw.RawModule.module.__get__ return self._module() cupy/_core/raw.pyx:402: in cupy._core.raw.RawModule._module mod = _get_raw_module( cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/raw.pyx:547: in cupy._core.raw._get_raw_module mod = cupy._core.core.compile_with_cache( cupy/_core/core.pyx:2015: in cupy._core.core.compile_with_cache cpdef function.Module compile_with_cache( cupy/_core/core.pyx:2078: in cupy._core.core.compile_with_cache return cuda.compile_with_cache( cupy/cuda/compiler.py:453: in compile_with_cache return _compile_with_cache_hip( cupy/cuda/compiler.py:853: in _compile_with_cache_hip binary, mapping = compile_using_nvrtc( cupy/cuda/compiler.py:295: in compile_using_nvrtc return _compile(source, options, cu_path, cupy/cuda/compiler.py:279: in _compile compiled_obj, mapping = prog.compile(options, log_stream) _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ self = options = ('--std=c++11', '-I/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/_core/include', '-I/opt/rocm/include', '-fcuda-flush-denormals-to-zero') log_stream = None def compile(self, options=(), log_stream=None): try: if self.name_expressions: for ker in self.name_expressions: nvrtc.addAddNameExpression(self.ptr, ker) nvrtc.compileProgram(self.ptr, options) mapping = None if self.name_expressions: mapping = {} for ker in self.name_expressions: mapping[ker] = nvrtc.getLoweredName(self.ptr, ker) if log_stream is not None: log_stream.write(nvrtc.getProgramLog(self.ptr)) # This is to ensure backwards compatibility with nvrtc if self.method == 'cubin': return nvrtc.getCUBIN(self.ptr), mapping elif self.method == 'ptx': return nvrtc.getPTX(self.ptr), mapping else: raise RuntimeError('Unknown NVRTC compile method') except nvrtc.NVRTCError: log = nvrtc.getProgramLog(self.ptr) > raise CompileException(log, self.src, self.name, options, 'nvrtc' if not runtime.is_hip else 'hiprtc') E AssertionError: Only cupy raises error E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 650, in compile E nvrtc.compileProgram(self.ptr, options) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 136, in cupy_backends.cuda.libs.nvrtc.compileProgram E cpdef compileProgram(intptr_t prog, options): E File "cupy_backends/cuda/libs/nvrtc.pyx", line 148, in cupy_backends.cuda.libs.nvrtc.compileProgram E check_status(status) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 67, in cupy_backends.cuda.libs.nvrtc.check_status E raise NVRTCError(status) E cupy_backends.cuda.libs.nvrtc.NVRTCError: HIPRTC_ERROR_COMPILATION (6) E E During handling of the above exception, another exception occurred: E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/testing/_helper.py", line 47, in _call_func E result = impl(*args, **kw) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/tests/cupy_tests/linalg_tests/test_product.py", line 274, in test_multidim_outer E return xp.outer(a, b) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/linalg/_product.py", line 276, in outer E return _core.tensordot_core(a, b, None, n, m, 1, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 430, in cupy._core._routines_linalg.tensordot_core E cpdef ndarray tensordot_core( E File "cupy/_core/_routines_linalg.pyx", line 501, in cupy._core._routines_linalg.tensordot_core E return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 327, in cupy._core._routines_linalg._integral_tensordot_core E kern = _tensordot_core_int_kernel(config, dtype) E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/_routines_linalg.pyx", line 302, in cupy._core._routines_linalg._tensordot_core_int_kernel E ker = mod.get_function( E File "cupy/_core/raw.pyx", line 470, in cupy._core.raw.RawModule.get_function E mangled_name = self.module.mapping.get(name) E File "cupy/_core/raw.pyx", line 394, in cupy._core.raw.RawModule.module.__get__ E return self._module() E File "cupy/_core/raw.pyx", line 402, in cupy._core.raw.RawModule._module E mod = _get_raw_module( E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/raw.pyx", line 547, in cupy._core.raw._get_raw_module E mod = cupy._core.core.compile_with_cache( E File "cupy/_core/core.pyx", line 2015, in cupy._core.core.compile_with_cache E cpdef function.Module compile_with_cache( E File "cupy/_core/core.pyx", line 2078, in cupy._core.core.compile_with_cache E return cuda.compile_with_cache( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 453, in compile_with_cache E return _compile_with_cache_hip( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 853, in _compile_with_cache_hip E binary, mapping = compile_using_nvrtc( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 295, in compile_using_nvrtc E return _compile(source, options, cu_path, E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 279, in _compile E compiled_obj, mapping = prog.compile(options, log_stream) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 667, in compile E raise CompileException(log, self.src, self.name, options, E cupy.cuda.compiler.CompileException: /tmp/comgr-183fe9/input/CompileSource:203:27: error: redeclaration of '__hiprtc_16' with a different type: 'void (*const)(int, int, int, const signed char *, const signed char *, signed char *)' vs 'void (*const)(int, int, int, const bool *, const bool *, bool *)' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-183fe9/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-183fe9/input/CompileSource:204:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-183fe9/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-183fe9/input/CompileSource:205:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-183fe9/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-183fe9/input/CompileSource:206:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-183fe9/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-183fe9/input/CompileSource:207:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-183fe9/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-183fe9/input/CompileSource:208:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-183fe9/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-183fe9/input/CompileSource:209:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-183fe9/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-183fe9/input/CompileSource:210:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-183fe9/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-183fe9/input/CompileSource:211:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-183fe9/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-183fe9/input/CompileSource:212:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-183fe9/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E 10 errors generated when compiling for gfx908. E Error: Failed to compile opencl source (from CL or HIP source to LLVM IR). cupy/cuda/compiler.py:667: AssertionError ----------------------------- Captured stdout call ----------------------------- dtype is ____________________________ TestProduct.test_outer ____________________________ self = xp = dtype = @testing.for_all_dtypes() @testing.numpy_cupy_allclose() def test_outer(self, xp, dtype): a = testing.shaped_arange((5,), xp, dtype) b = testing.shaped_arange((4,), xp, dtype) > return xp.outer(a, b) tests/cupy_tests/linalg_tests/test_product.py:260: _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ cupy/linalg/_product.py:276: in outer return _core.tensordot_core(a, b, None, n, m, 1, ret_shape) cupy/_core/_routines_linalg.pyx:430: in cupy._core._routines_linalg.tensordot_core cpdef ndarray tensordot_core( cupy/_core/_routines_linalg.pyx:501: in cupy._core._routines_linalg.tensordot_core return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) cupy/_core/_routines_linalg.pyx:327: in cupy._core._routines_linalg._integral_tensordot_core kern = _tensordot_core_int_kernel(config, dtype) cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/_routines_linalg.pyx:302: in cupy._core._routines_linalg._tensordot_core_int_kernel ker = mod.get_function( cupy/_core/raw.pyx:470: in cupy._core.raw.RawModule.get_function mangled_name = self.module.mapping.get(name) cupy/_core/raw.pyx:394: in cupy._core.raw.RawModule.module.__get__ return self._module() cupy/_core/raw.pyx:402: in cupy._core.raw.RawModule._module mod = _get_raw_module( cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/raw.pyx:547: in cupy._core.raw._get_raw_module mod = cupy._core.core.compile_with_cache( cupy/_core/core.pyx:2015: in cupy._core.core.compile_with_cache cpdef function.Module compile_with_cache( cupy/_core/core.pyx:2078: in cupy._core.core.compile_with_cache return cuda.compile_with_cache( cupy/cuda/compiler.py:453: in compile_with_cache return _compile_with_cache_hip( cupy/cuda/compiler.py:853: in _compile_with_cache_hip binary, mapping = compile_using_nvrtc( cupy/cuda/compiler.py:295: in compile_using_nvrtc return _compile(source, options, cu_path, cupy/cuda/compiler.py:279: in _compile compiled_obj, mapping = prog.compile(options, log_stream) _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ self = options = ('--std=c++11', '-I/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/_core/include', '-I/opt/rocm/include', '-fcuda-flush-denormals-to-zero') log_stream = None def compile(self, options=(), log_stream=None): try: if self.name_expressions: for ker in self.name_expressions: nvrtc.addAddNameExpression(self.ptr, ker) nvrtc.compileProgram(self.ptr, options) mapping = None if self.name_expressions: mapping = {} for ker in self.name_expressions: mapping[ker] = nvrtc.getLoweredName(self.ptr, ker) if log_stream is not None: log_stream.write(nvrtc.getProgramLog(self.ptr)) # This is to ensure backwards compatibility with nvrtc if self.method == 'cubin': return nvrtc.getCUBIN(self.ptr), mapping elif self.method == 'ptx': return nvrtc.getPTX(self.ptr), mapping else: raise RuntimeError('Unknown NVRTC compile method') except nvrtc.NVRTCError: log = nvrtc.getProgramLog(self.ptr) > raise CompileException(log, self.src, self.name, options, 'nvrtc' if not runtime.is_hip else 'hiprtc') E AssertionError: Only cupy raises error E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 650, in compile E nvrtc.compileProgram(self.ptr, options) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 136, in cupy_backends.cuda.libs.nvrtc.compileProgram E cpdef compileProgram(intptr_t prog, options): E File "cupy_backends/cuda/libs/nvrtc.pyx", line 148, in cupy_backends.cuda.libs.nvrtc.compileProgram E check_status(status) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 67, in cupy_backends.cuda.libs.nvrtc.check_status E raise NVRTCError(status) E cupy_backends.cuda.libs.nvrtc.NVRTCError: HIPRTC_ERROR_COMPILATION (6) E E During handling of the above exception, another exception occurred: E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/testing/_helper.py", line 47, in _call_func E result = impl(*args, **kw) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/tests/cupy_tests/linalg_tests/test_product.py", line 260, in test_outer E return xp.outer(a, b) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/linalg/_product.py", line 276, in outer E return _core.tensordot_core(a, b, None, n, m, 1, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 430, in cupy._core._routines_linalg.tensordot_core E cpdef ndarray tensordot_core( E File "cupy/_core/_routines_linalg.pyx", line 501, in cupy._core._routines_linalg.tensordot_core E return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 327, in cupy._core._routines_linalg._integral_tensordot_core E kern = _tensordot_core_int_kernel(config, dtype) E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/_routines_linalg.pyx", line 302, in cupy._core._routines_linalg._tensordot_core_int_kernel E ker = mod.get_function( E File "cupy/_core/raw.pyx", line 470, in cupy._core.raw.RawModule.get_function E mangled_name = self.module.mapping.get(name) E File "cupy/_core/raw.pyx", line 394, in cupy._core.raw.RawModule.module.__get__ E return self._module() E File "cupy/_core/raw.pyx", line 402, in cupy._core.raw.RawModule._module E mod = _get_raw_module( E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/raw.pyx", line 547, in cupy._core.raw._get_raw_module E mod = cupy._core.core.compile_with_cache( E File "cupy/_core/core.pyx", line 2015, in cupy._core.core.compile_with_cache E cpdef function.Module compile_with_cache( E File "cupy/_core/core.pyx", line 2078, in cupy._core.core.compile_with_cache E return cuda.compile_with_cache( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 453, in compile_with_cache E return _compile_with_cache_hip( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 853, in _compile_with_cache_hip E binary, mapping = compile_using_nvrtc( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 295, in compile_using_nvrtc E return _compile(source, options, cu_path, E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 279, in _compile E compiled_obj, mapping = prog.compile(options, log_stream) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 667, in compile E raise CompileException(log, self.src, self.name, options, E cupy.cuda.compiler.CompileException: /tmp/comgr-f4a4cd/input/CompileSource:203:27: error: redeclaration of '__hiprtc_16' with a different type: 'void (*const)(int, int, int, const signed char *, const signed char *, signed char *)' vs 'void (*const)(int, int, int, const bool *, const bool *, bool *)' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-f4a4cd/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-f4a4cd/input/CompileSource:204:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-f4a4cd/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-f4a4cd/input/CompileSource:205:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-f4a4cd/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-f4a4cd/input/CompileSource:206:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-f4a4cd/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-f4a4cd/input/CompileSource:207:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-f4a4cd/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-f4a4cd/input/CompileSource:208:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-f4a4cd/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-f4a4cd/input/CompileSource:209:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-f4a4cd/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-f4a4cd/input/CompileSource:210:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-f4a4cd/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-f4a4cd/input/CompileSource:211:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-f4a4cd/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-f4a4cd/input/CompileSource:212:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-f4a4cd/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E 10 errors generated when compiling for gfx908. E Error: Failed to compile opencl source (from CL or HIP source to LLVM IR). cupy/cuda/compiler.py:667: AssertionError ----------------------------- Captured stdout call ----------------------------- dtype is ________________________ TestProduct.test_reversed_kron ________________________ self = xp = dtype = @testing.for_all_dtypes() @testing.numpy_cupy_allclose() def test_reversed_kron(self, xp, dtype): a = testing.shaped_arange((4,), xp, dtype) b = testing.shaped_arange((5,), xp, dtype) > return xp.kron(a[::-1], b[::-1]) tests/cupy_tests/linalg_tests/test_product.py:368: _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ cupy/linalg/_product.py:426: in kron out = _core.tensordot_core( cupy/_core/_routines_linalg.pyx:430: in cupy._core._routines_linalg.tensordot_core cpdef ndarray tensordot_core( cupy/_core/_routines_linalg.pyx:501: in cupy._core._routines_linalg.tensordot_core return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) cupy/_core/_routines_linalg.pyx:327: in cupy._core._routines_linalg._integral_tensordot_core kern = _tensordot_core_int_kernel(config, dtype) cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/_routines_linalg.pyx:302: in cupy._core._routines_linalg._tensordot_core_int_kernel ker = mod.get_function( cupy/_core/raw.pyx:470: in cupy._core.raw.RawModule.get_function mangled_name = self.module.mapping.get(name) cupy/_core/raw.pyx:394: in cupy._core.raw.RawModule.module.__get__ return self._module() cupy/_core/raw.pyx:402: in cupy._core.raw.RawModule._module mod = _get_raw_module( cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/raw.pyx:547: in cupy._core.raw._get_raw_module mod = cupy._core.core.compile_with_cache( cupy/_core/core.pyx:2015: in cupy._core.core.compile_with_cache cpdef function.Module compile_with_cache( cupy/_core/core.pyx:2078: in cupy._core.core.compile_with_cache return cuda.compile_with_cache( cupy/cuda/compiler.py:453: in compile_with_cache return _compile_with_cache_hip( cupy/cuda/compiler.py:853: in _compile_with_cache_hip binary, mapping = compile_using_nvrtc( cupy/cuda/compiler.py:295: in compile_using_nvrtc return _compile(source, options, cu_path, cupy/cuda/compiler.py:279: in _compile compiled_obj, mapping = prog.compile(options, log_stream) _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ self = options = ('--std=c++11', '-I/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/_core/include', '-I/opt/rocm/include', '-fcuda-flush-denormals-to-zero') log_stream = None def compile(self, options=(), log_stream=None): try: if self.name_expressions: for ker in self.name_expressions: nvrtc.addAddNameExpression(self.ptr, ker) nvrtc.compileProgram(self.ptr, options) mapping = None if self.name_expressions: mapping = {} for ker in self.name_expressions: mapping[ker] = nvrtc.getLoweredName(self.ptr, ker) if log_stream is not None: log_stream.write(nvrtc.getProgramLog(self.ptr)) # This is to ensure backwards compatibility with nvrtc if self.method == 'cubin': return nvrtc.getCUBIN(self.ptr), mapping elif self.method == 'ptx': return nvrtc.getPTX(self.ptr), mapping else: raise RuntimeError('Unknown NVRTC compile method') except nvrtc.NVRTCError: log = nvrtc.getProgramLog(self.ptr) > raise CompileException(log, self.src, self.name, options, 'nvrtc' if not runtime.is_hip else 'hiprtc') E AssertionError: Only cupy raises error E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 650, in compile E nvrtc.compileProgram(self.ptr, options) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 136, in cupy_backends.cuda.libs.nvrtc.compileProgram E cpdef compileProgram(intptr_t prog, options): E File "cupy_backends/cuda/libs/nvrtc.pyx", line 148, in cupy_backends.cuda.libs.nvrtc.compileProgram E check_status(status) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 67, in cupy_backends.cuda.libs.nvrtc.check_status E raise NVRTCError(status) E cupy_backends.cuda.libs.nvrtc.NVRTCError: HIPRTC_ERROR_COMPILATION (6) E E During handling of the above exception, another exception occurred: E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/testing/_helper.py", line 47, in _call_func E result = impl(*args, **kw) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/tests/cupy_tests/linalg_tests/test_product.py", line 368, in test_reversed_kron E return xp.kron(a[::-1], b[::-1]) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/linalg/_product.py", line 426, in kron E out = _core.tensordot_core( E File "cupy/_core/_routines_linalg.pyx", line 430, in cupy._core._routines_linalg.tensordot_core E cpdef ndarray tensordot_core( E File "cupy/_core/_routines_linalg.pyx", line 501, in cupy._core._routines_linalg.tensordot_core E return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 327, in cupy._core._routines_linalg._integral_tensordot_core E kern = _tensordot_core_int_kernel(config, dtype) E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/_routines_linalg.pyx", line 302, in cupy._core._routines_linalg._tensordot_core_int_kernel E ker = mod.get_function( E File "cupy/_core/raw.pyx", line 470, in cupy._core.raw.RawModule.get_function E mangled_name = self.module.mapping.get(name) E File "cupy/_core/raw.pyx", line 394, in cupy._core.raw.RawModule.module.__get__ E return self._module() E File "cupy/_core/raw.pyx", line 402, in cupy._core.raw.RawModule._module E mod = _get_raw_module( E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/raw.pyx", line 547, in cupy._core.raw._get_raw_module E mod = cupy._core.core.compile_with_cache( E File "cupy/_core/core.pyx", line 2015, in cupy._core.core.compile_with_cache E cpdef function.Module compile_with_cache( E File "cupy/_core/core.pyx", line 2078, in cupy._core.core.compile_with_cache E return cuda.compile_with_cache( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 453, in compile_with_cache E return _compile_with_cache_hip( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 853, in _compile_with_cache_hip E binary, mapping = compile_using_nvrtc( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 295, in compile_using_nvrtc E return _compile(source, options, cu_path, E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 279, in _compile E compiled_obj, mapping = prog.compile(options, log_stream) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 667, in compile E raise CompileException(log, self.src, self.name, options, E cupy.cuda.compiler.CompileException: /tmp/comgr-f4cb39/input/CompileSource:203:27: error: redeclaration of '__hiprtc_16' with a different type: 'void (*const)(int, int, int, const signed char *, const signed char *, signed char *)' vs 'void (*const)(int, int, int, const bool *, const bool *, bool *)' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-f4cb39/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-f4cb39/input/CompileSource:204:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-f4cb39/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-f4cb39/input/CompileSource:205:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-f4cb39/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-f4cb39/input/CompileSource:206:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-f4cb39/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-f4cb39/input/CompileSource:207:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-f4cb39/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-f4cb39/input/CompileSource:208:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-f4cb39/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-f4cb39/input/CompileSource:209:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-f4cb39/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-f4cb39/input/CompileSource:210:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-f4cb39/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-f4cb39/input/CompileSource:211:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-f4cb39/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-f4cb39/input/CompileSource:212:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-f4cb39/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E 10 errors generated when compiling for gfx908. E Error: Failed to compile opencl source (from CL or HIP source to LLVM IR). cupy/cuda/compiler.py:667: AssertionError ----------------------------- Captured stdout call ----------------------------- dtype is _______________________ TestProduct.test_reversed_outer ________________________ self = xp = dtype = @testing.for_all_dtypes() @testing.numpy_cupy_allclose() def test_reversed_outer(self, xp, dtype): a = testing.shaped_arange((5,), xp, dtype) b = testing.shaped_arange((4,), xp, dtype) > return xp.outer(a[::-1], b[::-1]) tests/cupy_tests/linalg_tests/test_product.py:267: _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ cupy/linalg/_product.py:276: in outer return _core.tensordot_core(a, b, None, n, m, 1, ret_shape) cupy/_core/_routines_linalg.pyx:430: in cupy._core._routines_linalg.tensordot_core cpdef ndarray tensordot_core( cupy/_core/_routines_linalg.pyx:501: in cupy._core._routines_linalg.tensordot_core return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) cupy/_core/_routines_linalg.pyx:327: in cupy._core._routines_linalg._integral_tensordot_core kern = _tensordot_core_int_kernel(config, dtype) cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/_routines_linalg.pyx:302: in cupy._core._routines_linalg._tensordot_core_int_kernel ker = mod.get_function( cupy/_core/raw.pyx:470: in cupy._core.raw.RawModule.get_function mangled_name = self.module.mapping.get(name) cupy/_core/raw.pyx:394: in cupy._core.raw.RawModule.module.__get__ return self._module() cupy/_core/raw.pyx:402: in cupy._core.raw.RawModule._module mod = _get_raw_module( cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/raw.pyx:547: in cupy._core.raw._get_raw_module mod = cupy._core.core.compile_with_cache( cupy/_core/core.pyx:2015: in cupy._core.core.compile_with_cache cpdef function.Module compile_with_cache( cupy/_core/core.pyx:2078: in cupy._core.core.compile_with_cache return cuda.compile_with_cache( cupy/cuda/compiler.py:453: in compile_with_cache return _compile_with_cache_hip( cupy/cuda/compiler.py:853: in _compile_with_cache_hip binary, mapping = compile_using_nvrtc( cupy/cuda/compiler.py:295: in compile_using_nvrtc return _compile(source, options, cu_path, cupy/cuda/compiler.py:279: in _compile compiled_obj, mapping = prog.compile(options, log_stream) _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ self = options = ('--std=c++11', '-I/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/_core/include', '-I/opt/rocm/include', '-fcuda-flush-denormals-to-zero') log_stream = None def compile(self, options=(), log_stream=None): try: if self.name_expressions: for ker in self.name_expressions: nvrtc.addAddNameExpression(self.ptr, ker) nvrtc.compileProgram(self.ptr, options) mapping = None if self.name_expressions: mapping = {} for ker in self.name_expressions: mapping[ker] = nvrtc.getLoweredName(self.ptr, ker) if log_stream is not None: log_stream.write(nvrtc.getProgramLog(self.ptr)) # This is to ensure backwards compatibility with nvrtc if self.method == 'cubin': return nvrtc.getCUBIN(self.ptr), mapping elif self.method == 'ptx': return nvrtc.getPTX(self.ptr), mapping else: raise RuntimeError('Unknown NVRTC compile method') except nvrtc.NVRTCError: log = nvrtc.getProgramLog(self.ptr) > raise CompileException(log, self.src, self.name, options, 'nvrtc' if not runtime.is_hip else 'hiprtc') E AssertionError: Only cupy raises error E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 650, in compile E nvrtc.compileProgram(self.ptr, options) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 136, in cupy_backends.cuda.libs.nvrtc.compileProgram E cpdef compileProgram(intptr_t prog, options): E File "cupy_backends/cuda/libs/nvrtc.pyx", line 148, in cupy_backends.cuda.libs.nvrtc.compileProgram E check_status(status) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 67, in cupy_backends.cuda.libs.nvrtc.check_status E raise NVRTCError(status) E cupy_backends.cuda.libs.nvrtc.NVRTCError: HIPRTC_ERROR_COMPILATION (6) E E During handling of the above exception, another exception occurred: E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/testing/_helper.py", line 47, in _call_func E result = impl(*args, **kw) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/tests/cupy_tests/linalg_tests/test_product.py", line 267, in test_reversed_outer E return xp.outer(a[::-1], b[::-1]) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/linalg/_product.py", line 276, in outer E return _core.tensordot_core(a, b, None, n, m, 1, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 430, in cupy._core._routines_linalg.tensordot_core E cpdef ndarray tensordot_core( E File "cupy/_core/_routines_linalg.pyx", line 501, in cupy._core._routines_linalg.tensordot_core E return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 327, in cupy._core._routines_linalg._integral_tensordot_core E kern = _tensordot_core_int_kernel(config, dtype) E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/_routines_linalg.pyx", line 302, in cupy._core._routines_linalg._tensordot_core_int_kernel E ker = mod.get_function( E File "cupy/_core/raw.pyx", line 470, in cupy._core.raw.RawModule.get_function E mangled_name = self.module.mapping.get(name) E File "cupy/_core/raw.pyx", line 394, in cupy._core.raw.RawModule.module.__get__ E return self._module() E File "cupy/_core/raw.pyx", line 402, in cupy._core.raw.RawModule._module E mod = _get_raw_module( E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/raw.pyx", line 547, in cupy._core.raw._get_raw_module E mod = cupy._core.core.compile_with_cache( E File "cupy/_core/core.pyx", line 2015, in cupy._core.core.compile_with_cache E cpdef function.Module compile_with_cache( E File "cupy/_core/core.pyx", line 2078, in cupy._core.core.compile_with_cache E return cuda.compile_with_cache( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 453, in compile_with_cache E return _compile_with_cache_hip( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 853, in _compile_with_cache_hip E binary, mapping = compile_using_nvrtc( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 295, in compile_using_nvrtc E return _compile(source, options, cu_path, E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 279, in _compile E compiled_obj, mapping = prog.compile(options, log_stream) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 667, in compile E raise CompileException(log, self.src, self.name, options, E cupy.cuda.compiler.CompileException: /tmp/comgr-62b995/input/CompileSource:203:27: error: redeclaration of '__hiprtc_16' with a different type: 'void (*const)(int, int, int, const signed char *, const signed char *, signed char *)' vs 'void (*const)(int, int, int, const bool *, const bool *, bool *)' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-62b995/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-62b995/input/CompileSource:204:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-62b995/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-62b995/input/CompileSource:205:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-62b995/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-62b995/input/CompileSource:206:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-62b995/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-62b995/input/CompileSource:207:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-62b995/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-62b995/input/CompileSource:208:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-62b995/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-62b995/input/CompileSource:209:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-62b995/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-62b995/input/CompileSource:210:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-62b995/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-62b995/input/CompileSource:211:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-62b995/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-62b995/input/CompileSource:212:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-62b995/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E 10 errors generated when compiling for gfx908. E Error: Failed to compile opencl source (from CL or HIP source to LLVM IR). cupy/cuda/compiler.py:667: AssertionError ----------------------------- Captured stdout call ----------------------------- dtype is __________________________ TestProduct.test_tensordot __________________________ self = xp = dtype = @testing.for_all_dtypes() @testing.numpy_cupy_allclose() def test_tensordot(self, xp, dtype): a = testing.shaped_arange((2, 3, 4), xp, dtype) b = testing.shaped_arange((3, 4, 5), xp, dtype) > return xp.tensordot(a, b) tests/cupy_tests/linalg_tests/test_product.py:281: _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ cupy/linalg/_product.py:349: in tensordot return _core.tensordot_core(a, b, None, n, m, k, ret_shape) cupy/_core/_routines_linalg.pyx:430: in cupy._core._routines_linalg.tensordot_core cpdef ndarray tensordot_core( cupy/_core/_routines_linalg.pyx:501: in cupy._core._routines_linalg.tensordot_core return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) cupy/_core/_routines_linalg.pyx:327: in cupy._core._routines_linalg._integral_tensordot_core kern = _tensordot_core_int_kernel(config, dtype) cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/_routines_linalg.pyx:302: in cupy._core._routines_linalg._tensordot_core_int_kernel ker = mod.get_function( cupy/_core/raw.pyx:470: in cupy._core.raw.RawModule.get_function mangled_name = self.module.mapping.get(name) cupy/_core/raw.pyx:394: in cupy._core.raw.RawModule.module.__get__ return self._module() cupy/_core/raw.pyx:402: in cupy._core.raw.RawModule._module mod = _get_raw_module( cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/raw.pyx:547: in cupy._core.raw._get_raw_module mod = cupy._core.core.compile_with_cache( cupy/_core/core.pyx:2015: in cupy._core.core.compile_with_cache cpdef function.Module compile_with_cache( cupy/_core/core.pyx:2078: in cupy._core.core.compile_with_cache return cuda.compile_with_cache( cupy/cuda/compiler.py:453: in compile_with_cache return _compile_with_cache_hip( cupy/cuda/compiler.py:853: in _compile_with_cache_hip binary, mapping = compile_using_nvrtc( cupy/cuda/compiler.py:295: in compile_using_nvrtc return _compile(source, options, cu_path, cupy/cuda/compiler.py:279: in _compile compiled_obj, mapping = prog.compile(options, log_stream) _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ self = options = ('--std=c++11', '-I/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/_core/include', '-I/opt/rocm/include', '-fcuda-flush-denormals-to-zero') log_stream = None def compile(self, options=(), log_stream=None): try: if self.name_expressions: for ker in self.name_expressions: nvrtc.addAddNameExpression(self.ptr, ker) nvrtc.compileProgram(self.ptr, options) mapping = None if self.name_expressions: mapping = {} for ker in self.name_expressions: mapping[ker] = nvrtc.getLoweredName(self.ptr, ker) if log_stream is not None: log_stream.write(nvrtc.getProgramLog(self.ptr)) # This is to ensure backwards compatibility with nvrtc if self.method == 'cubin': return nvrtc.getCUBIN(self.ptr), mapping elif self.method == 'ptx': return nvrtc.getPTX(self.ptr), mapping else: raise RuntimeError('Unknown NVRTC compile method') except nvrtc.NVRTCError: log = nvrtc.getProgramLog(self.ptr) > raise CompileException(log, self.src, self.name, options, 'nvrtc' if not runtime.is_hip else 'hiprtc') E AssertionError: Only cupy raises error E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 650, in compile E nvrtc.compileProgram(self.ptr, options) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 136, in cupy_backends.cuda.libs.nvrtc.compileProgram E cpdef compileProgram(intptr_t prog, options): E File "cupy_backends/cuda/libs/nvrtc.pyx", line 148, in cupy_backends.cuda.libs.nvrtc.compileProgram E check_status(status) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 67, in cupy_backends.cuda.libs.nvrtc.check_status E raise NVRTCError(status) E cupy_backends.cuda.libs.nvrtc.NVRTCError: HIPRTC_ERROR_COMPILATION (6) E E During handling of the above exception, another exception occurred: E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/testing/_helper.py", line 47, in _call_func E result = impl(*args, **kw) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/tests/cupy_tests/linalg_tests/test_product.py", line 281, in test_tensordot E return xp.tensordot(a, b) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/linalg/_product.py", line 349, in tensordot E return _core.tensordot_core(a, b, None, n, m, k, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 430, in cupy._core._routines_linalg.tensordot_core E cpdef ndarray tensordot_core( E File "cupy/_core/_routines_linalg.pyx", line 501, in cupy._core._routines_linalg.tensordot_core E return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 327, in cupy._core._routines_linalg._integral_tensordot_core E kern = _tensordot_core_int_kernel(config, dtype) E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/_routines_linalg.pyx", line 302, in cupy._core._routines_linalg._tensordot_core_int_kernel E ker = mod.get_function( E File "cupy/_core/raw.pyx", line 470, in cupy._core.raw.RawModule.get_function E mangled_name = self.module.mapping.get(name) E File "cupy/_core/raw.pyx", line 394, in cupy._core.raw.RawModule.module.__get__ E return self._module() E File "cupy/_core/raw.pyx", line 402, in cupy._core.raw.RawModule._module E mod = _get_raw_module( E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/raw.pyx", line 547, in cupy._core.raw._get_raw_module E mod = cupy._core.core.compile_with_cache( E File "cupy/_core/core.pyx", line 2015, in cupy._core.core.compile_with_cache E cpdef function.Module compile_with_cache( E File "cupy/_core/core.pyx", line 2078, in cupy._core.core.compile_with_cache E return cuda.compile_with_cache( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 453, in compile_with_cache E return _compile_with_cache_hip( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 853, in _compile_with_cache_hip E binary, mapping = compile_using_nvrtc( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 295, in compile_using_nvrtc E return _compile(source, options, cu_path, E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 279, in _compile E compiled_obj, mapping = prog.compile(options, log_stream) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 667, in compile E raise CompileException(log, self.src, self.name, options, E cupy.cuda.compiler.CompileException: /tmp/comgr-1d51a3/input/CompileSource:203:27: error: redeclaration of '__hiprtc_16' with a different type: 'void (*const)(int, int, int, const signed char *, const signed char *, signed char *)' vs 'void (*const)(int, int, int, const bool *, const bool *, bool *)' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-1d51a3/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-1d51a3/input/CompileSource:204:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-1d51a3/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-1d51a3/input/CompileSource:205:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-1d51a3/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-1d51a3/input/CompileSource:206:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-1d51a3/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-1d51a3/input/CompileSource:207:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-1d51a3/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-1d51a3/input/CompileSource:208:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-1d51a3/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-1d51a3/input/CompileSource:209:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-1d51a3/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-1d51a3/input/CompileSource:210:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-1d51a3/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-1d51a3/input/CompileSource:211:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-1d51a3/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-1d51a3/input/CompileSource:212:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-1d51a3/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E 10 errors generated when compiling for gfx908. E Error: Failed to compile opencl source (from CL or HIP source to LLVM IR). cupy/cuda/compiler.py:667: AssertionError ----------------------------- Captured stdout call ----------------------------- dtype is ___________________ TestProduct.test_tensordot_with_int_axes ___________________ self = xp = dtype = @testing.for_all_dtypes() @testing.numpy_cupy_allclose() def test_tensordot_with_int_axes(self, xp, dtype): if dtype in (numpy.uint8, numpy.int8, numpy.uint16, numpy.int16): a = testing.shaped_arange((1, 2, 3), xp, dtype) b = testing.shaped_arange((2, 3, 1), xp, dtype) return xp.tensordot(a, b, axes=2) else: a = testing.shaped_arange((2, 3, 4, 5), xp, dtype) b = testing.shaped_arange((3, 4, 5, 2), xp, dtype) > return xp.tensordot(a, b, axes=3) tests/cupy_tests/linalg_tests/test_product.py:300: _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ cupy/linalg/_product.py:349: in tensordot return _core.tensordot_core(a, b, None, n, m, k, ret_shape) cupy/_core/_routines_linalg.pyx:430: in cupy._core._routines_linalg.tensordot_core cpdef ndarray tensordot_core( cupy/_core/_routines_linalg.pyx:501: in cupy._core._routines_linalg.tensordot_core return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) cupy/_core/_routines_linalg.pyx:327: in cupy._core._routines_linalg._integral_tensordot_core kern = _tensordot_core_int_kernel(config, dtype) cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/_routines_linalg.pyx:302: in cupy._core._routines_linalg._tensordot_core_int_kernel ker = mod.get_function( cupy/_core/raw.pyx:470: in cupy._core.raw.RawModule.get_function mangled_name = self.module.mapping.get(name) cupy/_core/raw.pyx:394: in cupy._core.raw.RawModule.module.__get__ return self._module() cupy/_core/raw.pyx:402: in cupy._core.raw.RawModule._module mod = _get_raw_module( cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/raw.pyx:547: in cupy._core.raw._get_raw_module mod = cupy._core.core.compile_with_cache( cupy/_core/core.pyx:2015: in cupy._core.core.compile_with_cache cpdef function.Module compile_with_cache( cupy/_core/core.pyx:2078: in cupy._core.core.compile_with_cache return cuda.compile_with_cache( cupy/cuda/compiler.py:453: in compile_with_cache return _compile_with_cache_hip( cupy/cuda/compiler.py:853: in _compile_with_cache_hip binary, mapping = compile_using_nvrtc( cupy/cuda/compiler.py:295: in compile_using_nvrtc return _compile(source, options, cu_path, cupy/cuda/compiler.py:279: in _compile compiled_obj, mapping = prog.compile(options, log_stream) _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ self = options = ('--std=c++11', '-I/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/_core/include', '-I/opt/rocm/include', '-fcuda-flush-denormals-to-zero') log_stream = None def compile(self, options=(), log_stream=None): try: if self.name_expressions: for ker in self.name_expressions: nvrtc.addAddNameExpression(self.ptr, ker) nvrtc.compileProgram(self.ptr, options) mapping = None if self.name_expressions: mapping = {} for ker in self.name_expressions: mapping[ker] = nvrtc.getLoweredName(self.ptr, ker) if log_stream is not None: log_stream.write(nvrtc.getProgramLog(self.ptr)) # This is to ensure backwards compatibility with nvrtc if self.method == 'cubin': return nvrtc.getCUBIN(self.ptr), mapping elif self.method == 'ptx': return nvrtc.getPTX(self.ptr), mapping else: raise RuntimeError('Unknown NVRTC compile method') except nvrtc.NVRTCError: log = nvrtc.getProgramLog(self.ptr) > raise CompileException(log, self.src, self.name, options, 'nvrtc' if not runtime.is_hip else 'hiprtc') E AssertionError: Only cupy raises error E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 650, in compile E nvrtc.compileProgram(self.ptr, options) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 136, in cupy_backends.cuda.libs.nvrtc.compileProgram E cpdef compileProgram(intptr_t prog, options): E File "cupy_backends/cuda/libs/nvrtc.pyx", line 148, in cupy_backends.cuda.libs.nvrtc.compileProgram E check_status(status) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 67, in cupy_backends.cuda.libs.nvrtc.check_status E raise NVRTCError(status) E cupy_backends.cuda.libs.nvrtc.NVRTCError: HIPRTC_ERROR_COMPILATION (6) E E During handling of the above exception, another exception occurred: E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/testing/_helper.py", line 47, in _call_func E result = impl(*args, **kw) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/tests/cupy_tests/linalg_tests/test_product.py", line 300, in test_tensordot_with_int_axes E return xp.tensordot(a, b, axes=3) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/linalg/_product.py", line 349, in tensordot E return _core.tensordot_core(a, b, None, n, m, k, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 430, in cupy._core._routines_linalg.tensordot_core E cpdef ndarray tensordot_core( E File "cupy/_core/_routines_linalg.pyx", line 501, in cupy._core._routines_linalg.tensordot_core E return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 327, in cupy._core._routines_linalg._integral_tensordot_core E kern = _tensordot_core_int_kernel(config, dtype) E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/_routines_linalg.pyx", line 302, in cupy._core._routines_linalg._tensordot_core_int_kernel E ker = mod.get_function( E File "cupy/_core/raw.pyx", line 470, in cupy._core.raw.RawModule.get_function E mangled_name = self.module.mapping.get(name) E File "cupy/_core/raw.pyx", line 394, in cupy._core.raw.RawModule.module.__get__ E return self._module() E File "cupy/_core/raw.pyx", line 402, in cupy._core.raw.RawModule._module E mod = _get_raw_module( E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/raw.pyx", line 547, in cupy._core.raw._get_raw_module E mod = cupy._core.core.compile_with_cache( E File "cupy/_core/core.pyx", line 2015, in cupy._core.core.compile_with_cache E cpdef function.Module compile_with_cache( E File "cupy/_core/core.pyx", line 2078, in cupy._core.core.compile_with_cache E return cuda.compile_with_cache( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 453, in compile_with_cache E return _compile_with_cache_hip( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 853, in _compile_with_cache_hip E binary, mapping = compile_using_nvrtc( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 295, in compile_using_nvrtc E return _compile(source, options, cu_path, E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 279, in _compile E compiled_obj, mapping = prog.compile(options, log_stream) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 667, in compile E raise CompileException(log, self.src, self.name, options, E cupy.cuda.compiler.CompileException: /tmp/comgr-5e9835/input/CompileSource:203:27: error: redeclaration of '__hiprtc_16' with a different type: 'void (*const)(int, int, int, const signed char *, const signed char *, signed char *)' vs 'void (*const)(int, int, int, const bool *, const bool *, bool *)' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-5e9835/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-5e9835/input/CompileSource:204:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-5e9835/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-5e9835/input/CompileSource:205:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-5e9835/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-5e9835/input/CompileSource:206:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-5e9835/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-5e9835/input/CompileSource:207:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-5e9835/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-5e9835/input/CompileSource:208:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-5e9835/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-5e9835/input/CompileSource:209:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-5e9835/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-5e9835/input/CompileSource:210:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-5e9835/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-5e9835/input/CompileSource:211:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-5e9835/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-5e9835/input/CompileSource:212:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-5e9835/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E 10 errors generated when compiling for gfx908. E Error: Failed to compile opencl source (from CL or HIP source to LLVM IR). cupy/cuda/compiler.py:667: AssertionError ----------------------------- Captured stdout call ----------------------------- dtype is __________________ TestProduct.test_tensordot_with_list_axes ___________________ self = xp = dtype = @testing.for_all_dtypes() @testing.numpy_cupy_allclose() def test_tensordot_with_list_axes(self, xp, dtype): if dtype in (numpy.uint8, numpy.int8, numpy.uint16, numpy.int16): # Avoid overflow a = testing.shaped_arange((1, 2, 3), xp, dtype) b = testing.shaped_arange((3, 1, 2), xp, dtype) return xp.tensordot(a, b, axes=([2, 1], [0, 2])) else: a = testing.shaped_arange((2, 3, 4, 5), xp, dtype) b = testing.shaped_arange((3, 5, 4, 2), xp, dtype) > return xp.tensordot(a, b, axes=([3, 2, 1], [1, 2, 0])) tests/cupy_tests/linalg_tests/test_product.py:330: _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ cupy/linalg/_product.py:349: in tensordot return _core.tensordot_core(a, b, None, n, m, k, ret_shape) cupy/_core/_routines_linalg.pyx:430: in cupy._core._routines_linalg.tensordot_core cpdef ndarray tensordot_core( cupy/_core/_routines_linalg.pyx:501: in cupy._core._routines_linalg.tensordot_core return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) cupy/_core/_routines_linalg.pyx:327: in cupy._core._routines_linalg._integral_tensordot_core kern = _tensordot_core_int_kernel(config, dtype) cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/_routines_linalg.pyx:302: in cupy._core._routines_linalg._tensordot_core_int_kernel ker = mod.get_function( cupy/_core/raw.pyx:470: in cupy._core.raw.RawModule.get_function mangled_name = self.module.mapping.get(name) cupy/_core/raw.pyx:394: in cupy._core.raw.RawModule.module.__get__ return self._module() cupy/_core/raw.pyx:402: in cupy._core.raw.RawModule._module mod = _get_raw_module( cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/raw.pyx:547: in cupy._core.raw._get_raw_module mod = cupy._core.core.compile_with_cache( cupy/_core/core.pyx:2015: in cupy._core.core.compile_with_cache cpdef function.Module compile_with_cache( cupy/_core/core.pyx:2078: in cupy._core.core.compile_with_cache return cuda.compile_with_cache( cupy/cuda/compiler.py:453: in compile_with_cache return _compile_with_cache_hip( cupy/cuda/compiler.py:853: in _compile_with_cache_hip binary, mapping = compile_using_nvrtc( cupy/cuda/compiler.py:295: in compile_using_nvrtc return _compile(source, options, cu_path, cupy/cuda/compiler.py:279: in _compile compiled_obj, mapping = prog.compile(options, log_stream) _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ self = options = ('--std=c++11', '-I/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/_core/include', '-I/opt/rocm/include', '-fcuda-flush-denormals-to-zero') log_stream = None def compile(self, options=(), log_stream=None): try: if self.name_expressions: for ker in self.name_expressions: nvrtc.addAddNameExpression(self.ptr, ker) nvrtc.compileProgram(self.ptr, options) mapping = None if self.name_expressions: mapping = {} for ker in self.name_expressions: mapping[ker] = nvrtc.getLoweredName(self.ptr, ker) if log_stream is not None: log_stream.write(nvrtc.getProgramLog(self.ptr)) # This is to ensure backwards compatibility with nvrtc if self.method == 'cubin': return nvrtc.getCUBIN(self.ptr), mapping elif self.method == 'ptx': return nvrtc.getPTX(self.ptr), mapping else: raise RuntimeError('Unknown NVRTC compile method') except nvrtc.NVRTCError: log = nvrtc.getProgramLog(self.ptr) > raise CompileException(log, self.src, self.name, options, 'nvrtc' if not runtime.is_hip else 'hiprtc') E AssertionError: Only cupy raises error E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 650, in compile E nvrtc.compileProgram(self.ptr, options) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 136, in cupy_backends.cuda.libs.nvrtc.compileProgram E cpdef compileProgram(intptr_t prog, options): E File "cupy_backends/cuda/libs/nvrtc.pyx", line 148, in cupy_backends.cuda.libs.nvrtc.compileProgram E check_status(status) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 67, in cupy_backends.cuda.libs.nvrtc.check_status E raise NVRTCError(status) E cupy_backends.cuda.libs.nvrtc.NVRTCError: HIPRTC_ERROR_COMPILATION (6) E E During handling of the above exception, another exception occurred: E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/testing/_helper.py", line 47, in _call_func E result = impl(*args, **kw) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/tests/cupy_tests/linalg_tests/test_product.py", line 330, in test_tensordot_with_list_axes E return xp.tensordot(a, b, axes=([3, 2, 1], [1, 2, 0])) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/linalg/_product.py", line 349, in tensordot E return _core.tensordot_core(a, b, None, n, m, k, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 430, in cupy._core._routines_linalg.tensordot_core E cpdef ndarray tensordot_core( E File "cupy/_core/_routines_linalg.pyx", line 501, in cupy._core._routines_linalg.tensordot_core E return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 327, in cupy._core._routines_linalg._integral_tensordot_core E kern = _tensordot_core_int_kernel(config, dtype) E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/_routines_linalg.pyx", line 302, in cupy._core._routines_linalg._tensordot_core_int_kernel E ker = mod.get_function( E File "cupy/_core/raw.pyx", line 470, in cupy._core.raw.RawModule.get_function E mangled_name = self.module.mapping.get(name) E File "cupy/_core/raw.pyx", line 394, in cupy._core.raw.RawModule.module.__get__ E return self._module() E File "cupy/_core/raw.pyx", line 402, in cupy._core.raw.RawModule._module E mod = _get_raw_module( E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/raw.pyx", line 547, in cupy._core.raw._get_raw_module E mod = cupy._core.core.compile_with_cache( E File "cupy/_core/core.pyx", line 2015, in cupy._core.core.compile_with_cache E cpdef function.Module compile_with_cache( E File "cupy/_core/core.pyx", line 2078, in cupy._core.core.compile_with_cache E return cuda.compile_with_cache( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 453, in compile_with_cache E return _compile_with_cache_hip( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 853, in _compile_with_cache_hip E binary, mapping = compile_using_nvrtc( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 295, in compile_using_nvrtc E return _compile(source, options, cu_path, E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 279, in _compile E compiled_obj, mapping = prog.compile(options, log_stream) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 667, in compile E raise CompileException(log, self.src, self.name, options, E cupy.cuda.compiler.CompileException: /tmp/comgr-a3db02/input/CompileSource:203:27: error: redeclaration of '__hiprtc_16' with a different type: 'void (*const)(int, int, int, const signed char *, const signed char *, signed char *)' vs 'void (*const)(int, int, int, const bool *, const bool *, bool *)' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-a3db02/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-a3db02/input/CompileSource:204:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-a3db02/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-a3db02/input/CompileSource:205:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-a3db02/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-a3db02/input/CompileSource:206:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-a3db02/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-a3db02/input/CompileSource:207:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-a3db02/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-a3db02/input/CompileSource:208:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-a3db02/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-a3db02/input/CompileSource:209:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-a3db02/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-a3db02/input/CompileSource:210:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-a3db02/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-a3db02/input/CompileSource:211:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-a3db02/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-a3db02/input/CompileSource:212:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-a3db02/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E 10 errors generated when compiling for gfx908. E Error: Failed to compile opencl source (from CL or HIP source to LLVM IR). cupy/cuda/compiler.py:667: AssertionError ----------------------------- Captured stdout call ----------------------------- dtype is _______________________ TestProduct.test_transposed_dot ________________________ self = xp = dtype = @testing.for_all_dtypes() @testing.numpy_cupy_allclose() def test_transposed_dot(self, xp, dtype): a = testing.shaped_arange((2, 3, 4), xp, dtype).transpose(1, 0, 2) b = testing.shaped_arange((2, 3, 4), xp, dtype).transpose(0, 2, 1) > return xp.dot(a, b) tests/cupy_tests/linalg_tests/test_product.py:163: _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ cupy/linalg/_product.py:65: in dot return a.dot(b, out) cupy/_core/core.pyx:1607: in cupy._core.core.ndarray.dot return _linalg.dot(self, b, out) cupy/_core/_routines_linalg.pyx:427: in cupy._core._routines_linalg.dot return tensordot_core(a, b, out, n, m, k, ret_shape) cupy/_core/_routines_linalg.pyx:501: in cupy._core._routines_linalg.tensordot_core return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) cupy/_core/_routines_linalg.pyx:327: in cupy._core._routines_linalg._integral_tensordot_core kern = _tensordot_core_int_kernel(config, dtype) cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/_routines_linalg.pyx:302: in cupy._core._routines_linalg._tensordot_core_int_kernel ker = mod.get_function( cupy/_core/raw.pyx:470: in cupy._core.raw.RawModule.get_function mangled_name = self.module.mapping.get(name) cupy/_core/raw.pyx:394: in cupy._core.raw.RawModule.module.__get__ return self._module() cupy/_core/raw.pyx:402: in cupy._core.raw.RawModule._module mod = _get_raw_module( cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/raw.pyx:547: in cupy._core.raw._get_raw_module mod = cupy._core.core.compile_with_cache( cupy/_core/core.pyx:2015: in cupy._core.core.compile_with_cache cpdef function.Module compile_with_cache( cupy/_core/core.pyx:2078: in cupy._core.core.compile_with_cache return cuda.compile_with_cache( cupy/cuda/compiler.py:453: in compile_with_cache return _compile_with_cache_hip( cupy/cuda/compiler.py:853: in _compile_with_cache_hip binary, mapping = compile_using_nvrtc( cupy/cuda/compiler.py:295: in compile_using_nvrtc return _compile(source, options, cu_path, cupy/cuda/compiler.py:279: in _compile compiled_obj, mapping = prog.compile(options, log_stream) _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ self = options = ('--std=c++11', '-I/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/_core/include', '-I/opt/rocm/include', '-fcuda-flush-denormals-to-zero') log_stream = None def compile(self, options=(), log_stream=None): try: if self.name_expressions: for ker in self.name_expressions: nvrtc.addAddNameExpression(self.ptr, ker) nvrtc.compileProgram(self.ptr, options) mapping = None if self.name_expressions: mapping = {} for ker in self.name_expressions: mapping[ker] = nvrtc.getLoweredName(self.ptr, ker) if log_stream is not None: log_stream.write(nvrtc.getProgramLog(self.ptr)) # This is to ensure backwards compatibility with nvrtc if self.method == 'cubin': return nvrtc.getCUBIN(self.ptr), mapping elif self.method == 'ptx': return nvrtc.getPTX(self.ptr), mapping else: raise RuntimeError('Unknown NVRTC compile method') except nvrtc.NVRTCError: log = nvrtc.getProgramLog(self.ptr) > raise CompileException(log, self.src, self.name, options, 'nvrtc' if not runtime.is_hip else 'hiprtc') E AssertionError: Only cupy raises error E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 650, in compile E nvrtc.compileProgram(self.ptr, options) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 136, in cupy_backends.cuda.libs.nvrtc.compileProgram E cpdef compileProgram(intptr_t prog, options): E File "cupy_backends/cuda/libs/nvrtc.pyx", line 148, in cupy_backends.cuda.libs.nvrtc.compileProgram E check_status(status) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 67, in cupy_backends.cuda.libs.nvrtc.check_status E raise NVRTCError(status) E cupy_backends.cuda.libs.nvrtc.NVRTCError: HIPRTC_ERROR_COMPILATION (6) E E During handling of the above exception, another exception occurred: E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/testing/_helper.py", line 47, in _call_func E result = impl(*args, **kw) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/tests/cupy_tests/linalg_tests/test_product.py", line 163, in test_transposed_dot E return xp.dot(a, b) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/linalg/_product.py", line 65, in dot E return a.dot(b, out) E File "cupy/_core/core.pyx", line 1607, in cupy._core.core.ndarray.dot E return _linalg.dot(self, b, out) E File "cupy/_core/_routines_linalg.pyx", line 427, in cupy._core._routines_linalg.dot E return tensordot_core(a, b, out, n, m, k, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 501, in cupy._core._routines_linalg.tensordot_core E return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 327, in cupy._core._routines_linalg._integral_tensordot_core E kern = _tensordot_core_int_kernel(config, dtype) E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/_routines_linalg.pyx", line 302, in cupy._core._routines_linalg._tensordot_core_int_kernel E ker = mod.get_function( E File "cupy/_core/raw.pyx", line 470, in cupy._core.raw.RawModule.get_function E mangled_name = self.module.mapping.get(name) E File "cupy/_core/raw.pyx", line 394, in cupy._core.raw.RawModule.module.__get__ E return self._module() E File "cupy/_core/raw.pyx", line 402, in cupy._core.raw.RawModule._module E mod = _get_raw_module( E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/raw.pyx", line 547, in cupy._core.raw._get_raw_module E mod = cupy._core.core.compile_with_cache( E File "cupy/_core/core.pyx", line 2015, in cupy._core.core.compile_with_cache E cpdef function.Module compile_with_cache( E File "cupy/_core/core.pyx", line 2078, in cupy._core.core.compile_with_cache E return cuda.compile_with_cache( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 453, in compile_with_cache E return _compile_with_cache_hip( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 853, in _compile_with_cache_hip E binary, mapping = compile_using_nvrtc( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 295, in compile_using_nvrtc E return _compile(source, options, cu_path, E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 279, in _compile E compiled_obj, mapping = prog.compile(options, log_stream) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 667, in compile E raise CompileException(log, self.src, self.name, options, E cupy.cuda.compiler.CompileException: /tmp/comgr-266740/input/CompileSource:203:27: error: redeclaration of '__hiprtc_16' with a different type: 'void (*const)(int, int, int, const signed char *, const signed char *, signed char *)' vs 'void (*const)(int, int, int, const bool *, const bool *, bool *)' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-266740/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-266740/input/CompileSource:204:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-266740/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-266740/input/CompileSource:205:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-266740/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-266740/input/CompileSource:206:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-266740/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-266740/input/CompileSource:207:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-266740/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-266740/input/CompileSource:208:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-266740/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-266740/input/CompileSource:209:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-266740/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-266740/input/CompileSource:210:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-266740/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-266740/input/CompileSource:211:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-266740/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-266740/input/CompileSource:212:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-266740/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E 10 errors generated when compiling for gfx908. E Error: Failed to compile opencl source (from CL or HIP source to LLVM IR). cupy/cuda/compiler.py:667: AssertionError ----------------------------- Captured stdout call ----------------------------- dtype is ___________________ TestProduct.test_transposed_dot_with_out ___________________ self = xp = dtype = @testing.for_all_dtypes() @testing.numpy_cupy_allclose() def test_transposed_dot_with_out(self, xp, dtype): a = testing.shaped_arange((2, 3, 4), xp, dtype).transpose(1, 0, 2) b = testing.shaped_arange((4, 2, 3), xp, dtype).transpose(2, 0, 1) c = xp.ndarray((3, 2, 3, 2), dtype=dtype) > xp.dot(a, b, out=c) tests/cupy_tests/linalg_tests/test_product.py:171: _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ cupy/linalg/_product.py:65: in dot return a.dot(b, out) cupy/_core/core.pyx:1607: in cupy._core.core.ndarray.dot return _linalg.dot(self, b, out) cupy/_core/_routines_linalg.pyx:427: in cupy._core._routines_linalg.dot return tensordot_core(a, b, out, n, m, k, ret_shape) cupy/_core/_routines_linalg.pyx:501: in cupy._core._routines_linalg.tensordot_core return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) cupy/_core/_routines_linalg.pyx:327: in cupy._core._routines_linalg._integral_tensordot_core kern = _tensordot_core_int_kernel(config, dtype) cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/_routines_linalg.pyx:302: in cupy._core._routines_linalg._tensordot_core_int_kernel ker = mod.get_function( cupy/_core/raw.pyx:470: in cupy._core.raw.RawModule.get_function mangled_name = self.module.mapping.get(name) cupy/_core/raw.pyx:394: in cupy._core.raw.RawModule.module.__get__ return self._module() cupy/_core/raw.pyx:402: in cupy._core.raw.RawModule._module mod = _get_raw_module( cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/raw.pyx:547: in cupy._core.raw._get_raw_module mod = cupy._core.core.compile_with_cache( cupy/_core/core.pyx:2015: in cupy._core.core.compile_with_cache cpdef function.Module compile_with_cache( cupy/_core/core.pyx:2078: in cupy._core.core.compile_with_cache return cuda.compile_with_cache( cupy/cuda/compiler.py:453: in compile_with_cache return _compile_with_cache_hip( cupy/cuda/compiler.py:853: in _compile_with_cache_hip binary, mapping = compile_using_nvrtc( cupy/cuda/compiler.py:295: in compile_using_nvrtc return _compile(source, options, cu_path, cupy/cuda/compiler.py:279: in _compile compiled_obj, mapping = prog.compile(options, log_stream) _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ self = options = ('--std=c++11', '-I/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/_core/include', '-I/opt/rocm/include', '-fcuda-flush-denormals-to-zero') log_stream = None def compile(self, options=(), log_stream=None): try: if self.name_expressions: for ker in self.name_expressions: nvrtc.addAddNameExpression(self.ptr, ker) nvrtc.compileProgram(self.ptr, options) mapping = None if self.name_expressions: mapping = {} for ker in self.name_expressions: mapping[ker] = nvrtc.getLoweredName(self.ptr, ker) if log_stream is not None: log_stream.write(nvrtc.getProgramLog(self.ptr)) # This is to ensure backwards compatibility with nvrtc if self.method == 'cubin': return nvrtc.getCUBIN(self.ptr), mapping elif self.method == 'ptx': return nvrtc.getPTX(self.ptr), mapping else: raise RuntimeError('Unknown NVRTC compile method') except nvrtc.NVRTCError: log = nvrtc.getProgramLog(self.ptr) > raise CompileException(log, self.src, self.name, options, 'nvrtc' if not runtime.is_hip else 'hiprtc') E AssertionError: Only cupy raises error E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 650, in compile E nvrtc.compileProgram(self.ptr, options) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 136, in cupy_backends.cuda.libs.nvrtc.compileProgram E cpdef compileProgram(intptr_t prog, options): E File "cupy_backends/cuda/libs/nvrtc.pyx", line 148, in cupy_backends.cuda.libs.nvrtc.compileProgram E check_status(status) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 67, in cupy_backends.cuda.libs.nvrtc.check_status E raise NVRTCError(status) E cupy_backends.cuda.libs.nvrtc.NVRTCError: HIPRTC_ERROR_COMPILATION (6) E E During handling of the above exception, another exception occurred: E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/testing/_helper.py", line 47, in _call_func E result = impl(*args, **kw) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/tests/cupy_tests/linalg_tests/test_product.py", line 171, in test_transposed_dot_with_out E xp.dot(a, b, out=c) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/linalg/_product.py", line 65, in dot E return a.dot(b, out) E File "cupy/_core/core.pyx", line 1607, in cupy._core.core.ndarray.dot E return _linalg.dot(self, b, out) E File "cupy/_core/_routines_linalg.pyx", line 427, in cupy._core._routines_linalg.dot E return tensordot_core(a, b, out, n, m, k, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 501, in cupy._core._routines_linalg.tensordot_core E return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 327, in cupy._core._routines_linalg._integral_tensordot_core E kern = _tensordot_core_int_kernel(config, dtype) E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/_routines_linalg.pyx", line 302, in cupy._core._routines_linalg._tensordot_core_int_kernel E ker = mod.get_function( E File "cupy/_core/raw.pyx", line 470, in cupy._core.raw.RawModule.get_function E mangled_name = self.module.mapping.get(name) E File "cupy/_core/raw.pyx", line 394, in cupy._core.raw.RawModule.module.__get__ E return self._module() E File "cupy/_core/raw.pyx", line 402, in cupy._core.raw.RawModule._module E mod = _get_raw_module( E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/raw.pyx", line 547, in cupy._core.raw._get_raw_module E mod = cupy._core.core.compile_with_cache( E File "cupy/_core/core.pyx", line 2015, in cupy._core.core.compile_with_cache E cpdef function.Module compile_with_cache( E File "cupy/_core/core.pyx", line 2078, in cupy._core.core.compile_with_cache E return cuda.compile_with_cache( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 453, in compile_with_cache E return _compile_with_cache_hip( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 853, in _compile_with_cache_hip E binary, mapping = compile_using_nvrtc( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 295, in compile_using_nvrtc E return _compile(source, options, cu_path, E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 279, in _compile E compiled_obj, mapping = prog.compile(options, log_stream) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 667, in compile E raise CompileException(log, self.src, self.name, options, E cupy.cuda.compiler.CompileException: /tmp/comgr-7711ae/input/CompileSource:203:27: error: redeclaration of '__hiprtc_16' with a different type: 'void (*const)(int, int, int, const signed char *, const signed char *, signed char *)' vs 'void (*const)(int, int, int, const bool *, const bool *, bool *)' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-7711ae/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-7711ae/input/CompileSource:204:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-7711ae/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-7711ae/input/CompileSource:205:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-7711ae/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-7711ae/input/CompileSource:206:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-7711ae/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-7711ae/input/CompileSource:207:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-7711ae/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-7711ae/input/CompileSource:208:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-7711ae/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-7711ae/input/CompileSource:209:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-7711ae/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-7711ae/input/CompileSource:210:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-7711ae/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-7711ae/input/CompileSource:211:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-7711ae/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-7711ae/input/CompileSource:212:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-7711ae/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E 10 errors generated when compiling for gfx908. E Error: Failed to compile opencl source (from CL or HIP source to LLVM IR). cupy/cuda/compiler.py:667: AssertionError ----------------------------- Captured stdout call ----------------------------- dtype is ________________ TestProduct.test_transposed_higher_order_inner ________________ self = xp = dtype = @testing.for_all_dtypes() @testing.numpy_cupy_allclose() def test_transposed_higher_order_inner(self, xp, dtype): a = testing.shaped_arange((2, 4, 3), xp, dtype).transpose(2, 0, 1) b = testing.shaped_arange((4, 2, 3), xp, dtype).transpose(1, 2, 0) > return xp.inner(a, b) tests/cupy_tests/linalg_tests/test_product.py:253: _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ cupy/linalg/_product.py:251: in inner return _core.tensordot_core(a, b, None, n, m, k, ret_shape) cupy/_core/_routines_linalg.pyx:430: in cupy._core._routines_linalg.tensordot_core cpdef ndarray tensordot_core( cupy/_core/_routines_linalg.pyx:501: in cupy._core._routines_linalg.tensordot_core return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) cupy/_core/_routines_linalg.pyx:327: in cupy._core._routines_linalg._integral_tensordot_core kern = _tensordot_core_int_kernel(config, dtype) cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/_routines_linalg.pyx:302: in cupy._core._routines_linalg._tensordot_core_int_kernel ker = mod.get_function( cupy/_core/raw.pyx:470: in cupy._core.raw.RawModule.get_function mangled_name = self.module.mapping.get(name) cupy/_core/raw.pyx:394: in cupy._core.raw.RawModule.module.__get__ return self._module() cupy/_core/raw.pyx:402: in cupy._core.raw.RawModule._module mod = _get_raw_module( cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/raw.pyx:547: in cupy._core.raw._get_raw_module mod = cupy._core.core.compile_with_cache( cupy/_core/core.pyx:2015: in cupy._core.core.compile_with_cache cpdef function.Module compile_with_cache( cupy/_core/core.pyx:2078: in cupy._core.core.compile_with_cache return cuda.compile_with_cache( cupy/cuda/compiler.py:453: in compile_with_cache return _compile_with_cache_hip( cupy/cuda/compiler.py:853: in _compile_with_cache_hip binary, mapping = compile_using_nvrtc( cupy/cuda/compiler.py:295: in compile_using_nvrtc return _compile(source, options, cu_path, cupy/cuda/compiler.py:279: in _compile compiled_obj, mapping = prog.compile(options, log_stream) _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ self = options = ('--std=c++11', '-I/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/_core/include', '-I/opt/rocm/include', '-fcuda-flush-denormals-to-zero') log_stream = None def compile(self, options=(), log_stream=None): try: if self.name_expressions: for ker in self.name_expressions: nvrtc.addAddNameExpression(self.ptr, ker) nvrtc.compileProgram(self.ptr, options) mapping = None if self.name_expressions: mapping = {} for ker in self.name_expressions: mapping[ker] = nvrtc.getLoweredName(self.ptr, ker) if log_stream is not None: log_stream.write(nvrtc.getProgramLog(self.ptr)) # This is to ensure backwards compatibility with nvrtc if self.method == 'cubin': return nvrtc.getCUBIN(self.ptr), mapping elif self.method == 'ptx': return nvrtc.getPTX(self.ptr), mapping else: raise RuntimeError('Unknown NVRTC compile method') except nvrtc.NVRTCError: log = nvrtc.getProgramLog(self.ptr) > raise CompileException(log, self.src, self.name, options, 'nvrtc' if not runtime.is_hip else 'hiprtc') E AssertionError: Only cupy raises error E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 650, in compile E nvrtc.compileProgram(self.ptr, options) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 136, in cupy_backends.cuda.libs.nvrtc.compileProgram E cpdef compileProgram(intptr_t prog, options): E File "cupy_backends/cuda/libs/nvrtc.pyx", line 148, in cupy_backends.cuda.libs.nvrtc.compileProgram E check_status(status) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 67, in cupy_backends.cuda.libs.nvrtc.check_status E raise NVRTCError(status) E cupy_backends.cuda.libs.nvrtc.NVRTCError: HIPRTC_ERROR_COMPILATION (6) E E During handling of the above exception, another exception occurred: E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/testing/_helper.py", line 47, in _call_func E result = impl(*args, **kw) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/tests/cupy_tests/linalg_tests/test_product.py", line 253, in test_transposed_higher_order_inner E return xp.inner(a, b) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/linalg/_product.py", line 251, in inner E return _core.tensordot_core(a, b, None, n, m, k, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 430, in cupy._core._routines_linalg.tensordot_core E cpdef ndarray tensordot_core( E File "cupy/_core/_routines_linalg.pyx", line 501, in cupy._core._routines_linalg.tensordot_core E return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 327, in cupy._core._routines_linalg._integral_tensordot_core E kern = _tensordot_core_int_kernel(config, dtype) E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/_routines_linalg.pyx", line 302, in cupy._core._routines_linalg._tensordot_core_int_kernel E ker = mod.get_function( E File "cupy/_core/raw.pyx", line 470, in cupy._core.raw.RawModule.get_function E mangled_name = self.module.mapping.get(name) E File "cupy/_core/raw.pyx", line 394, in cupy._core.raw.RawModule.module.__get__ E return self._module() E File "cupy/_core/raw.pyx", line 402, in cupy._core.raw.RawModule._module E mod = _get_raw_module( E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/raw.pyx", line 547, in cupy._core.raw._get_raw_module E mod = cupy._core.core.compile_with_cache( E File "cupy/_core/core.pyx", line 2015, in cupy._core.core.compile_with_cache E cpdef function.Module compile_with_cache( E File "cupy/_core/core.pyx", line 2078, in cupy._core.core.compile_with_cache E return cuda.compile_with_cache( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 453, in compile_with_cache E return _compile_with_cache_hip( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 853, in _compile_with_cache_hip E binary, mapping = compile_using_nvrtc( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 295, in compile_using_nvrtc E return _compile(source, options, cu_path, E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 279, in _compile E compiled_obj, mapping = prog.compile(options, log_stream) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 667, in compile E raise CompileException(log, self.src, self.name, options, E cupy.cuda.compiler.CompileException: /tmp/comgr-dbbf11/input/CompileSource:203:27: error: redeclaration of '__hiprtc_16' with a different type: 'void (*const)(int, int, int, const signed char *, const signed char *, signed char *)' vs 'void (*const)(int, int, int, const bool *, const bool *, bool *)' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-dbbf11/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-dbbf11/input/CompileSource:204:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-dbbf11/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-dbbf11/input/CompileSource:205:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-dbbf11/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-dbbf11/input/CompileSource:206:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-dbbf11/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-dbbf11/input/CompileSource:207:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-dbbf11/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-dbbf11/input/CompileSource:208:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-dbbf11/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-dbbf11/input/CompileSource:209:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-dbbf11/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-dbbf11/input/CompileSource:210:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-dbbf11/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-dbbf11/input/CompileSource:211:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-dbbf11/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-dbbf11/input/CompileSource:212:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-dbbf11/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E 10 errors generated when compiling for gfx908. E Error: Failed to compile opencl source (from CL or HIP source to LLVM IR). cupy/cuda/compiler.py:667: AssertionError ----------------------------- Captured stdout call ----------------------------- dtype is ____________________ TestProduct.test_transposed_tensordot _____________________ self = xp = dtype = @testing.for_all_dtypes() @testing.numpy_cupy_allclose() def test_transposed_tensordot(self, xp, dtype): a = testing.shaped_arange((2, 3, 4), xp, dtype).transpose(1, 0, 2) b = testing.shaped_arange((4, 3, 2), xp, dtype).transpose(2, 0, 1) > return xp.tensordot(a, b) tests/cupy_tests/linalg_tests/test_product.py:288: _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ cupy/linalg/_product.py:349: in tensordot return _core.tensordot_core(a, b, None, n, m, k, ret_shape) cupy/_core/_routines_linalg.pyx:430: in cupy._core._routines_linalg.tensordot_core cpdef ndarray tensordot_core( cupy/_core/_routines_linalg.pyx:501: in cupy._core._routines_linalg.tensordot_core return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) cupy/_core/_routines_linalg.pyx:327: in cupy._core._routines_linalg._integral_tensordot_core kern = _tensordot_core_int_kernel(config, dtype) cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/_routines_linalg.pyx:302: in cupy._core._routines_linalg._tensordot_core_int_kernel ker = mod.get_function( cupy/_core/raw.pyx:470: in cupy._core.raw.RawModule.get_function mangled_name = self.module.mapping.get(name) cupy/_core/raw.pyx:394: in cupy._core.raw.RawModule.module.__get__ return self._module() cupy/_core/raw.pyx:402: in cupy._core.raw.RawModule._module mod = _get_raw_module( cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/raw.pyx:547: in cupy._core.raw._get_raw_module mod = cupy._core.core.compile_with_cache( cupy/_core/core.pyx:2015: in cupy._core.core.compile_with_cache cpdef function.Module compile_with_cache( cupy/_core/core.pyx:2078: in cupy._core.core.compile_with_cache return cuda.compile_with_cache( cupy/cuda/compiler.py:453: in compile_with_cache return _compile_with_cache_hip( cupy/cuda/compiler.py:853: in _compile_with_cache_hip binary, mapping = compile_using_nvrtc( cupy/cuda/compiler.py:295: in compile_using_nvrtc return _compile(source, options, cu_path, cupy/cuda/compiler.py:279: in _compile compiled_obj, mapping = prog.compile(options, log_stream) _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ self = options = ('--std=c++11', '-I/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/_core/include', '-I/opt/rocm/include', '-fcuda-flush-denormals-to-zero') log_stream = None def compile(self, options=(), log_stream=None): try: if self.name_expressions: for ker in self.name_expressions: nvrtc.addAddNameExpression(self.ptr, ker) nvrtc.compileProgram(self.ptr, options) mapping = None if self.name_expressions: mapping = {} for ker in self.name_expressions: mapping[ker] = nvrtc.getLoweredName(self.ptr, ker) if log_stream is not None: log_stream.write(nvrtc.getProgramLog(self.ptr)) # This is to ensure backwards compatibility with nvrtc if self.method == 'cubin': return nvrtc.getCUBIN(self.ptr), mapping elif self.method == 'ptx': return nvrtc.getPTX(self.ptr), mapping else: raise RuntimeError('Unknown NVRTC compile method') except nvrtc.NVRTCError: log = nvrtc.getProgramLog(self.ptr) > raise CompileException(log, self.src, self.name, options, 'nvrtc' if not runtime.is_hip else 'hiprtc') E AssertionError: Only cupy raises error E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 650, in compile E nvrtc.compileProgram(self.ptr, options) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 136, in cupy_backends.cuda.libs.nvrtc.compileProgram E cpdef compileProgram(intptr_t prog, options): E File "cupy_backends/cuda/libs/nvrtc.pyx", line 148, in cupy_backends.cuda.libs.nvrtc.compileProgram E check_status(status) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 67, in cupy_backends.cuda.libs.nvrtc.check_status E raise NVRTCError(status) E cupy_backends.cuda.libs.nvrtc.NVRTCError: HIPRTC_ERROR_COMPILATION (6) E E During handling of the above exception, another exception occurred: E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/testing/_helper.py", line 47, in _call_func E result = impl(*args, **kw) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/tests/cupy_tests/linalg_tests/test_product.py", line 288, in test_transposed_tensordot E return xp.tensordot(a, b) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/linalg/_product.py", line 349, in tensordot E return _core.tensordot_core(a, b, None, n, m, k, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 430, in cupy._core._routines_linalg.tensordot_core E cpdef ndarray tensordot_core( E File "cupy/_core/_routines_linalg.pyx", line 501, in cupy._core._routines_linalg.tensordot_core E return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 327, in cupy._core._routines_linalg._integral_tensordot_core E kern = _tensordot_core_int_kernel(config, dtype) E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/_routines_linalg.pyx", line 302, in cupy._core._routines_linalg._tensordot_core_int_kernel E ker = mod.get_function( E File "cupy/_core/raw.pyx", line 470, in cupy._core.raw.RawModule.get_function E mangled_name = self.module.mapping.get(name) E File "cupy/_core/raw.pyx", line 394, in cupy._core.raw.RawModule.module.__get__ E return self._module() E File "cupy/_core/raw.pyx", line 402, in cupy._core.raw.RawModule._module E mod = _get_raw_module( E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/raw.pyx", line 547, in cupy._core.raw._get_raw_module E mod = cupy._core.core.compile_with_cache( E File "cupy/_core/core.pyx", line 2015, in cupy._core.core.compile_with_cache E cpdef function.Module compile_with_cache( E File "cupy/_core/core.pyx", line 2078, in cupy._core.core.compile_with_cache E return cuda.compile_with_cache( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 453, in compile_with_cache E return _compile_with_cache_hip( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 853, in _compile_with_cache_hip E binary, mapping = compile_using_nvrtc( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 295, in compile_using_nvrtc E return _compile(source, options, cu_path, E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 279, in _compile E compiled_obj, mapping = prog.compile(options, log_stream) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 667, in compile E raise CompileException(log, self.src, self.name, options, E cupy.cuda.compiler.CompileException: /tmp/comgr-81df9e/input/CompileSource:203:27: error: redeclaration of '__hiprtc_16' with a different type: 'void (*const)(int, int, int, const signed char *, const signed char *, signed char *)' vs 'void (*const)(int, int, int, const bool *, const bool *, bool *)' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-81df9e/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-81df9e/input/CompileSource:204:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-81df9e/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-81df9e/input/CompileSource:205:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-81df9e/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-81df9e/input/CompileSource:206:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-81df9e/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-81df9e/input/CompileSource:207:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-81df9e/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-81df9e/input/CompileSource:208:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-81df9e/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-81df9e/input/CompileSource:209:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-81df9e/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-81df9e/input/CompileSource:210:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-81df9e/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-81df9e/input/CompileSource:211:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-81df9e/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-81df9e/input/CompileSource:212:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-81df9e/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E 10 errors generated when compiling for gfx908. E Error: Failed to compile opencl source (from CL or HIP source to LLVM IR). cupy/cuda/compiler.py:667: AssertionError ----------------------------- Captured stdout call ----------------------------- dtype is _____________ TestProduct.test_transposed_tensordot_with_int_axes ______________ self = xp = dtype = @testing.for_all_dtypes() @testing.numpy_cupy_allclose() def test_transposed_tensordot_with_int_axes(self, xp, dtype): if dtype in (numpy.uint8, numpy.int8, numpy.uint16, numpy.int16): # Avoid overflow a = testing.shaped_arange( (1, 2, 3), xp, dtype).transpose(2, 0, 1) b = testing.shaped_arange( (3, 2, 1), xp, dtype).transpose(2, 1, 0) > return xp.tensordot(a, b, axes=2) tests/cupy_tests/linalg_tests/test_product.py:311: _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ cupy/linalg/_product.py:349: in tensordot return _core.tensordot_core(a, b, None, n, m, k, ret_shape) cupy/_core/_routines_linalg.pyx:430: in cupy._core._routines_linalg.tensordot_core cpdef ndarray tensordot_core( cupy/_core/_routines_linalg.pyx:501: in cupy._core._routines_linalg.tensordot_core return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) cupy/_core/_routines_linalg.pyx:327: in cupy._core._routines_linalg._integral_tensordot_core kern = _tensordot_core_int_kernel(config, dtype) cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/_routines_linalg.pyx:302: in cupy._core._routines_linalg._tensordot_core_int_kernel ker = mod.get_function( cupy/_core/raw.pyx:470: in cupy._core.raw.RawModule.get_function mangled_name = self.module.mapping.get(name) cupy/_core/raw.pyx:394: in cupy._core.raw.RawModule.module.__get__ return self._module() cupy/_core/raw.pyx:402: in cupy._core.raw.RawModule._module mod = _get_raw_module( cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/raw.pyx:547: in cupy._core.raw._get_raw_module mod = cupy._core.core.compile_with_cache( cupy/_core/core.pyx:2015: in cupy._core.core.compile_with_cache cpdef function.Module compile_with_cache( cupy/_core/core.pyx:2078: in cupy._core.core.compile_with_cache return cuda.compile_with_cache( cupy/cuda/compiler.py:453: in compile_with_cache return _compile_with_cache_hip( cupy/cuda/compiler.py:853: in _compile_with_cache_hip binary, mapping = compile_using_nvrtc( cupy/cuda/compiler.py:295: in compile_using_nvrtc return _compile(source, options, cu_path, cupy/cuda/compiler.py:279: in _compile compiled_obj, mapping = prog.compile(options, log_stream) _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ self = options = ('--std=c++11', '-I/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/_core/include', '-I/opt/rocm/include', '-fcuda-flush-denormals-to-zero') log_stream = None def compile(self, options=(), log_stream=None): try: if self.name_expressions: for ker in self.name_expressions: nvrtc.addAddNameExpression(self.ptr, ker) nvrtc.compileProgram(self.ptr, options) mapping = None if self.name_expressions: mapping = {} for ker in self.name_expressions: mapping[ker] = nvrtc.getLoweredName(self.ptr, ker) if log_stream is not None: log_stream.write(nvrtc.getProgramLog(self.ptr)) # This is to ensure backwards compatibility with nvrtc if self.method == 'cubin': return nvrtc.getCUBIN(self.ptr), mapping elif self.method == 'ptx': return nvrtc.getPTX(self.ptr), mapping else: raise RuntimeError('Unknown NVRTC compile method') except nvrtc.NVRTCError: log = nvrtc.getProgramLog(self.ptr) > raise CompileException(log, self.src, self.name, options, 'nvrtc' if not runtime.is_hip else 'hiprtc') E AssertionError: Only cupy raises error E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 650, in compile E nvrtc.compileProgram(self.ptr, options) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 136, in cupy_backends.cuda.libs.nvrtc.compileProgram E cpdef compileProgram(intptr_t prog, options): E File "cupy_backends/cuda/libs/nvrtc.pyx", line 148, in cupy_backends.cuda.libs.nvrtc.compileProgram E check_status(status) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 67, in cupy_backends.cuda.libs.nvrtc.check_status E raise NVRTCError(status) E cupy_backends.cuda.libs.nvrtc.NVRTCError: HIPRTC_ERROR_COMPILATION (6) E E During handling of the above exception, another exception occurred: E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/testing/_helper.py", line 47, in _call_func E result = impl(*args, **kw) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/tests/cupy_tests/linalg_tests/test_product.py", line 311, in test_transposed_tensordot_with_int_axes E return xp.tensordot(a, b, axes=2) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/linalg/_product.py", line 349, in tensordot E return _core.tensordot_core(a, b, None, n, m, k, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 430, in cupy._core._routines_linalg.tensordot_core E cpdef ndarray tensordot_core( E File "cupy/_core/_routines_linalg.pyx", line 501, in cupy._core._routines_linalg.tensordot_core E return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 327, in cupy._core._routines_linalg._integral_tensordot_core E kern = _tensordot_core_int_kernel(config, dtype) E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/_routines_linalg.pyx", line 302, in cupy._core._routines_linalg._tensordot_core_int_kernel E ker = mod.get_function( E File "cupy/_core/raw.pyx", line 470, in cupy._core.raw.RawModule.get_function E mangled_name = self.module.mapping.get(name) E File "cupy/_core/raw.pyx", line 394, in cupy._core.raw.RawModule.module.__get__ E return self._module() E File "cupy/_core/raw.pyx", line 402, in cupy._core.raw.RawModule._module E mod = _get_raw_module( E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/raw.pyx", line 547, in cupy._core.raw._get_raw_module E mod = cupy._core.core.compile_with_cache( E File "cupy/_core/core.pyx", line 2015, in cupy._core.core.compile_with_cache E cpdef function.Module compile_with_cache( E File "cupy/_core/core.pyx", line 2078, in cupy._core.core.compile_with_cache E return cuda.compile_with_cache( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 453, in compile_with_cache E return _compile_with_cache_hip( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 853, in _compile_with_cache_hip E binary, mapping = compile_using_nvrtc( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 295, in compile_using_nvrtc E return _compile(source, options, cu_path, E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 279, in _compile E compiled_obj, mapping = prog.compile(options, log_stream) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 667, in compile E raise CompileException(log, self.src, self.name, options, E cupy.cuda.compiler.CompileException: /tmp/comgr-2f7fa2/input/CompileSource:203:27: error: redeclaration of '__hiprtc_16' with a different type: 'void (*const)(int, int, int, const signed char *, const signed char *, signed char *)' vs 'void (*const)(int, int, int, const bool *, const bool *, bool *)' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-2f7fa2/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-2f7fa2/input/CompileSource:204:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-2f7fa2/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-2f7fa2/input/CompileSource:205:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-2f7fa2/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-2f7fa2/input/CompileSource:206:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-2f7fa2/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-2f7fa2/input/CompileSource:207:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-2f7fa2/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-2f7fa2/input/CompileSource:208:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-2f7fa2/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-2f7fa2/input/CompileSource:209:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-2f7fa2/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-2f7fa2/input/CompileSource:210:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-2f7fa2/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-2f7fa2/input/CompileSource:211:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-2f7fa2/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-2f7fa2/input/CompileSource:212:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-2f7fa2/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E 10 errors generated when compiling for gfx908. E Error: Failed to compile opencl source (from CL or HIP source to LLVM IR). cupy/cuda/compiler.py:667: AssertionError ----------------------------- Captured stdout call ----------------------------- dtype is _____________ TestProduct.test_transposed_tensordot_with_list_axes _____________ self = xp = dtype = @testing.for_all_dtypes() @testing.numpy_cupy_allclose() def test_transposed_tensordot_with_list_axes(self, xp, dtype): if dtype in (numpy.uint8, numpy.int8, numpy.uint16, numpy.int16): # Avoid overflow a = testing.shaped_arange( (1, 2, 3), xp, dtype).transpose(2, 0, 1) b = testing.shaped_arange( (2, 3, 1), xp, dtype).transpose(0, 2, 1) return xp.tensordot(a, b, axes=([2, 0], [0, 2])) else: a = testing.shaped_arange( (2, 3, 4, 5), xp, dtype).transpose(2, 0, 3, 1) b = testing.shaped_arange( (3, 5, 4, 2), xp, dtype).transpose(3, 0, 2, 1) > return xp.tensordot(a, b, axes=([2, 0, 3], [3, 2, 1])) tests/cupy_tests/linalg_tests/test_product.py:347: _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ cupy/linalg/_product.py:349: in tensordot return _core.tensordot_core(a, b, None, n, m, k, ret_shape) cupy/_core/_routines_linalg.pyx:430: in cupy._core._routines_linalg.tensordot_core cpdef ndarray tensordot_core( cupy/_core/_routines_linalg.pyx:501: in cupy._core._routines_linalg.tensordot_core return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) cupy/_core/_routines_linalg.pyx:327: in cupy._core._routines_linalg._integral_tensordot_core kern = _tensordot_core_int_kernel(config, dtype) cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/_routines_linalg.pyx:302: in cupy._core._routines_linalg._tensordot_core_int_kernel ker = mod.get_function( cupy/_core/raw.pyx:470: in cupy._core.raw.RawModule.get_function mangled_name = self.module.mapping.get(name) cupy/_core/raw.pyx:394: in cupy._core.raw.RawModule.module.__get__ return self._module() cupy/_core/raw.pyx:402: in cupy._core.raw.RawModule._module mod = _get_raw_module( cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/raw.pyx:547: in cupy._core.raw._get_raw_module mod = cupy._core.core.compile_with_cache( cupy/_core/core.pyx:2015: in cupy._core.core.compile_with_cache cpdef function.Module compile_with_cache( cupy/_core/core.pyx:2078: in cupy._core.core.compile_with_cache return cuda.compile_with_cache( cupy/cuda/compiler.py:453: in compile_with_cache return _compile_with_cache_hip( cupy/cuda/compiler.py:853: in _compile_with_cache_hip binary, mapping = compile_using_nvrtc( cupy/cuda/compiler.py:295: in compile_using_nvrtc return _compile(source, options, cu_path, cupy/cuda/compiler.py:279: in _compile compiled_obj, mapping = prog.compile(options, log_stream) _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ self = options = ('--std=c++11', '-I/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/_core/include', '-I/opt/rocm/include', '-fcuda-flush-denormals-to-zero') log_stream = None def compile(self, options=(), log_stream=None): try: if self.name_expressions: for ker in self.name_expressions: nvrtc.addAddNameExpression(self.ptr, ker) nvrtc.compileProgram(self.ptr, options) mapping = None if self.name_expressions: mapping = {} for ker in self.name_expressions: mapping[ker] = nvrtc.getLoweredName(self.ptr, ker) if log_stream is not None: log_stream.write(nvrtc.getProgramLog(self.ptr)) # This is to ensure backwards compatibility with nvrtc if self.method == 'cubin': return nvrtc.getCUBIN(self.ptr), mapping elif self.method == 'ptx': return nvrtc.getPTX(self.ptr), mapping else: raise RuntimeError('Unknown NVRTC compile method') except nvrtc.NVRTCError: log = nvrtc.getProgramLog(self.ptr) > raise CompileException(log, self.src, self.name, options, 'nvrtc' if not runtime.is_hip else 'hiprtc') E AssertionError: Only cupy raises error E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 650, in compile E nvrtc.compileProgram(self.ptr, options) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 136, in cupy_backends.cuda.libs.nvrtc.compileProgram E cpdef compileProgram(intptr_t prog, options): E File "cupy_backends/cuda/libs/nvrtc.pyx", line 148, in cupy_backends.cuda.libs.nvrtc.compileProgram E check_status(status) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 67, in cupy_backends.cuda.libs.nvrtc.check_status E raise NVRTCError(status) E cupy_backends.cuda.libs.nvrtc.NVRTCError: HIPRTC_ERROR_COMPILATION (6) E E During handling of the above exception, another exception occurred: E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/testing/_helper.py", line 47, in _call_func E result = impl(*args, **kw) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/tests/cupy_tests/linalg_tests/test_product.py", line 347, in test_transposed_tensordot_with_list_axes E return xp.tensordot(a, b, axes=([2, 0, 3], [3, 2, 1])) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/linalg/_product.py", line 349, in tensordot E return _core.tensordot_core(a, b, None, n, m, k, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 430, in cupy._core._routines_linalg.tensordot_core E cpdef ndarray tensordot_core( E File "cupy/_core/_routines_linalg.pyx", line 501, in cupy._core._routines_linalg.tensordot_core E return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 327, in cupy._core._routines_linalg._integral_tensordot_core E kern = _tensordot_core_int_kernel(config, dtype) E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/_routines_linalg.pyx", line 302, in cupy._core._routines_linalg._tensordot_core_int_kernel E ker = mod.get_function( E File "cupy/_core/raw.pyx", line 470, in cupy._core.raw.RawModule.get_function E mangled_name = self.module.mapping.get(name) E File "cupy/_core/raw.pyx", line 394, in cupy._core.raw.RawModule.module.__get__ E return self._module() E File "cupy/_core/raw.pyx", line 402, in cupy._core.raw.RawModule._module E mod = _get_raw_module( E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/raw.pyx", line 547, in cupy._core.raw._get_raw_module E mod = cupy._core.core.compile_with_cache( E File "cupy/_core/core.pyx", line 2015, in cupy._core.core.compile_with_cache E cpdef function.Module compile_with_cache( E File "cupy/_core/core.pyx", line 2078, in cupy._core.core.compile_with_cache E return cuda.compile_with_cache( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 453, in compile_with_cache E return _compile_with_cache_hip( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 853, in _compile_with_cache_hip E binary, mapping = compile_using_nvrtc( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 295, in compile_using_nvrtc E return _compile(source, options, cu_path, E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 279, in _compile E compiled_obj, mapping = prog.compile(options, log_stream) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 667, in compile E raise CompileException(log, self.src, self.name, options, E cupy.cuda.compiler.CompileException: /tmp/comgr-cb04dd/input/CompileSource:203:27: error: redeclaration of '__hiprtc_16' with a different type: 'void (*const)(int, int, int, const signed char *, const signed char *, signed char *)' vs 'void (*const)(int, int, int, const bool *, const bool *, bool *)' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-cb04dd/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-cb04dd/input/CompileSource:204:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-cb04dd/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-cb04dd/input/CompileSource:205:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-cb04dd/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-cb04dd/input/CompileSource:206:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-cb04dd/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-cb04dd/input/CompileSource:207:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-cb04dd/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-cb04dd/input/CompileSource:208:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-cb04dd/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-cb04dd/input/CompileSource:209:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-cb04dd/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-cb04dd/input/CompileSource:210:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-cb04dd/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-cb04dd/input/CompileSource:211:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-cb04dd/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-cb04dd/input/CompileSource:212:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-cb04dd/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E 10 errors generated when compiling for gfx908. E Error: Failed to compile opencl source (from CL or HIP source to LLVM IR). cupy/cuda/compiler.py:667: AssertionError ----------------------------- Captured stdout call ----------------------------- dtype is _____________________ TestMatrixPower.test_matrix_power_2 ______________________ self = xp = dtype = @testing.for_all_dtypes() @testing.numpy_cupy_allclose() def test_matrix_power_2(self, xp, dtype): a = testing.shaped_arange((3, 3), xp, dtype) > return xp.linalg.matrix_power(a, 2) tests/cupy_tests/linalg_tests/test_product.py:423: _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ cupy/linalg/_product.py:382: in matrix_power return cupy.matmul(M, M) cupy/linalg/_product.py:42: in matmul return _gu_func_matmul(x1, x2, out=out, axes=axes) cupy/_core/_gufuncs.py:676: in __call__ self._apply_func_to_inputs( cupy/_core/_gufuncs.py:393: in _apply_func_to_inputs fouts = func(*args) cupy/_core/_routines_linalg.pyx:691: in cupy._core._routines_linalg.matmul cpdef ndarray matmul(ndarray a, ndarray b, ndarray out=None): cupy/_core/_routines_linalg.pyx:732: in cupy._core._routines_linalg.matmul return dot(a, b, out) cupy/_core/_routines_linalg.pyx:427: in cupy._core._routines_linalg.dot return tensordot_core(a, b, out, n, m, k, ret_shape) cupy/_core/_routines_linalg.pyx:501: in cupy._core._routines_linalg.tensordot_core return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) cupy/_core/_routines_linalg.pyx:327: in cupy._core._routines_linalg._integral_tensordot_core kern = _tensordot_core_int_kernel(config, dtype) cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/_routines_linalg.pyx:302: in cupy._core._routines_linalg._tensordot_core_int_kernel ker = mod.get_function( cupy/_core/raw.pyx:470: in cupy._core.raw.RawModule.get_function mangled_name = self.module.mapping.get(name) cupy/_core/raw.pyx:394: in cupy._core.raw.RawModule.module.__get__ return self._module() cupy/_core/raw.pyx:402: in cupy._core.raw.RawModule._module mod = _get_raw_module( cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/raw.pyx:547: in cupy._core.raw._get_raw_module mod = cupy._core.core.compile_with_cache( cupy/_core/core.pyx:2015: in cupy._core.core.compile_with_cache cpdef function.Module compile_with_cache( cupy/_core/core.pyx:2078: in cupy._core.core.compile_with_cache return cuda.compile_with_cache( cupy/cuda/compiler.py:453: in compile_with_cache return _compile_with_cache_hip( cupy/cuda/compiler.py:853: in _compile_with_cache_hip binary, mapping = compile_using_nvrtc( cupy/cuda/compiler.py:295: in compile_using_nvrtc return _compile(source, options, cu_path, cupy/cuda/compiler.py:279: in _compile compiled_obj, mapping = prog.compile(options, log_stream) _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ self = options = ('--std=c++11', '-I/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/_core/include', '-I/opt/rocm/include', '-fcuda-flush-denormals-to-zero') log_stream = None def compile(self, options=(), log_stream=None): try: if self.name_expressions: for ker in self.name_expressions: nvrtc.addAddNameExpression(self.ptr, ker) nvrtc.compileProgram(self.ptr, options) mapping = None if self.name_expressions: mapping = {} for ker in self.name_expressions: mapping[ker] = nvrtc.getLoweredName(self.ptr, ker) if log_stream is not None: log_stream.write(nvrtc.getProgramLog(self.ptr)) # This is to ensure backwards compatibility with nvrtc if self.method == 'cubin': return nvrtc.getCUBIN(self.ptr), mapping elif self.method == 'ptx': return nvrtc.getPTX(self.ptr), mapping else: raise RuntimeError('Unknown NVRTC compile method') except nvrtc.NVRTCError: log = nvrtc.getProgramLog(self.ptr) > raise CompileException(log, self.src, self.name, options, 'nvrtc' if not runtime.is_hip else 'hiprtc') E AssertionError: Only cupy raises error E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 650, in compile E nvrtc.compileProgram(self.ptr, options) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 136, in cupy_backends.cuda.libs.nvrtc.compileProgram E cpdef compileProgram(intptr_t prog, options): E File "cupy_backends/cuda/libs/nvrtc.pyx", line 148, in cupy_backends.cuda.libs.nvrtc.compileProgram E check_status(status) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 67, in cupy_backends.cuda.libs.nvrtc.check_status E raise NVRTCError(status) E cupy_backends.cuda.libs.nvrtc.NVRTCError: HIPRTC_ERROR_COMPILATION (6) E E During handling of the above exception, another exception occurred: E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/testing/_helper.py", line 47, in _call_func E result = impl(*args, **kw) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/tests/cupy_tests/linalg_tests/test_product.py", line 423, in test_matrix_power_2 E return xp.linalg.matrix_power(a, 2) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/linalg/_product.py", line 382, in matrix_power E return cupy.matmul(M, M) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/linalg/_product.py", line 42, in matmul E return _gu_func_matmul(x1, x2, out=out, axes=axes) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/_core/_gufuncs.py", line 676, in __call__ E self._apply_func_to_inputs( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/_core/_gufuncs.py", line 393, in _apply_func_to_inputs E fouts = func(*args) E File "cupy/_core/_routines_linalg.pyx", line 691, in cupy._core._routines_linalg.matmul E cpdef ndarray matmul(ndarray a, ndarray b, ndarray out=None): E File "cupy/_core/_routines_linalg.pyx", line 732, in cupy._core._routines_linalg.matmul E return dot(a, b, out) E File "cupy/_core/_routines_linalg.pyx", line 427, in cupy._core._routines_linalg.dot E return tensordot_core(a, b, out, n, m, k, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 501, in cupy._core._routines_linalg.tensordot_core E return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 327, in cupy._core._routines_linalg._integral_tensordot_core E kern = _tensordot_core_int_kernel(config, dtype) E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/_routines_linalg.pyx", line 302, in cupy._core._routines_linalg._tensordot_core_int_kernel E ker = mod.get_function( E File "cupy/_core/raw.pyx", line 470, in cupy._core.raw.RawModule.get_function E mangled_name = self.module.mapping.get(name) E File "cupy/_core/raw.pyx", line 394, in cupy._core.raw.RawModule.module.__get__ E return self._module() E File "cupy/_core/raw.pyx", line 402, in cupy._core.raw.RawModule._module E mod = _get_raw_module( E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/raw.pyx", line 547, in cupy._core.raw._get_raw_module E mod = cupy._core.core.compile_with_cache( E File "cupy/_core/core.pyx", line 2015, in cupy._core.core.compile_with_cache E cpdef function.Module compile_with_cache( E File "cupy/_core/core.pyx", line 2078, in cupy._core.core.compile_with_cache E return cuda.compile_with_cache( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 453, in compile_with_cache E return _compile_with_cache_hip( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 853, in _compile_with_cache_hip E binary, mapping = compile_using_nvrtc( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 295, in compile_using_nvrtc E return _compile(source, options, cu_path, E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 279, in _compile E compiled_obj, mapping = prog.compile(options, log_stream) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 667, in compile E raise CompileException(log, self.src, self.name, options, E cupy.cuda.compiler.CompileException: /tmp/comgr-57a762/input/CompileSource:203:27: error: redeclaration of '__hiprtc_16' with a different type: 'void (*const)(int, int, int, const signed char *, const signed char *, signed char *)' vs 'void (*const)(int, int, int, const bool *, const bool *, bool *)' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-57a762/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-57a762/input/CompileSource:204:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-57a762/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-57a762/input/CompileSource:205:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-57a762/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-57a762/input/CompileSource:206:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-57a762/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-57a762/input/CompileSource:207:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-57a762/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-57a762/input/CompileSource:208:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-57a762/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-57a762/input/CompileSource:209:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-57a762/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-57a762/input/CompileSource:210:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-57a762/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-57a762/input/CompileSource:211:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-57a762/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-57a762/input/CompileSource:212:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-57a762/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E 10 errors generated when compiling for gfx908. E Error: Failed to compile opencl source (from CL or HIP source to LLVM IR). cupy/cuda/compiler.py:667: AssertionError ----------------------------- Captured stdout call ----------------------------- dtype is _____________________ TestMatrixPower.test_matrix_power_3 ______________________ self = xp = dtype = @testing.for_all_dtypes() @testing.numpy_cupy_allclose() def test_matrix_power_3(self, xp, dtype): a = testing.shaped_arange((3, 3), xp, dtype) > return xp.linalg.matrix_power(a, 3) tests/cupy_tests/linalg_tests/test_product.py:429: _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ cupy/linalg/_product.py:384: in matrix_power return cupy.matmul(cupy.matmul(M, M), M) cupy/linalg/_product.py:42: in matmul return _gu_func_matmul(x1, x2, out=out, axes=axes) cupy/_core/_gufuncs.py:676: in __call__ self._apply_func_to_inputs( cupy/_core/_gufuncs.py:393: in _apply_func_to_inputs fouts = func(*args) cupy/_core/_routines_linalg.pyx:691: in cupy._core._routines_linalg.matmul cpdef ndarray matmul(ndarray a, ndarray b, ndarray out=None): cupy/_core/_routines_linalg.pyx:732: in cupy._core._routines_linalg.matmul return dot(a, b, out) cupy/_core/_routines_linalg.pyx:427: in cupy._core._routines_linalg.dot return tensordot_core(a, b, out, n, m, k, ret_shape) cupy/_core/_routines_linalg.pyx:501: in cupy._core._routines_linalg.tensordot_core return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) cupy/_core/_routines_linalg.pyx:327: in cupy._core._routines_linalg._integral_tensordot_core kern = _tensordot_core_int_kernel(config, dtype) cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/_routines_linalg.pyx:302: in cupy._core._routines_linalg._tensordot_core_int_kernel ker = mod.get_function( cupy/_core/raw.pyx:470: in cupy._core.raw.RawModule.get_function mangled_name = self.module.mapping.get(name) cupy/_core/raw.pyx:394: in cupy._core.raw.RawModule.module.__get__ return self._module() cupy/_core/raw.pyx:402: in cupy._core.raw.RawModule._module mod = _get_raw_module( cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/raw.pyx:547: in cupy._core.raw._get_raw_module mod = cupy._core.core.compile_with_cache( cupy/_core/core.pyx:2015: in cupy._core.core.compile_with_cache cpdef function.Module compile_with_cache( cupy/_core/core.pyx:2078: in cupy._core.core.compile_with_cache return cuda.compile_with_cache( cupy/cuda/compiler.py:453: in compile_with_cache return _compile_with_cache_hip( cupy/cuda/compiler.py:853: in _compile_with_cache_hip binary, mapping = compile_using_nvrtc( cupy/cuda/compiler.py:295: in compile_using_nvrtc return _compile(source, options, cu_path, cupy/cuda/compiler.py:279: in _compile compiled_obj, mapping = prog.compile(options, log_stream) _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ self = options = ('--std=c++11', '-I/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/_core/include', '-I/opt/rocm/include', '-fcuda-flush-denormals-to-zero') log_stream = None def compile(self, options=(), log_stream=None): try: if self.name_expressions: for ker in self.name_expressions: nvrtc.addAddNameExpression(self.ptr, ker) nvrtc.compileProgram(self.ptr, options) mapping = None if self.name_expressions: mapping = {} for ker in self.name_expressions: mapping[ker] = nvrtc.getLoweredName(self.ptr, ker) if log_stream is not None: log_stream.write(nvrtc.getProgramLog(self.ptr)) # This is to ensure backwards compatibility with nvrtc if self.method == 'cubin': return nvrtc.getCUBIN(self.ptr), mapping elif self.method == 'ptx': return nvrtc.getPTX(self.ptr), mapping else: raise RuntimeError('Unknown NVRTC compile method') except nvrtc.NVRTCError: log = nvrtc.getProgramLog(self.ptr) > raise CompileException(log, self.src, self.name, options, 'nvrtc' if not runtime.is_hip else 'hiprtc') E AssertionError: Only cupy raises error E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 650, in compile E nvrtc.compileProgram(self.ptr, options) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 136, in cupy_backends.cuda.libs.nvrtc.compileProgram E cpdef compileProgram(intptr_t prog, options): E File "cupy_backends/cuda/libs/nvrtc.pyx", line 148, in cupy_backends.cuda.libs.nvrtc.compileProgram E check_status(status) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 67, in cupy_backends.cuda.libs.nvrtc.check_status E raise NVRTCError(status) E cupy_backends.cuda.libs.nvrtc.NVRTCError: HIPRTC_ERROR_COMPILATION (6) E E During handling of the above exception, another exception occurred: E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/testing/_helper.py", line 47, in _call_func E result = impl(*args, **kw) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/tests/cupy_tests/linalg_tests/test_product.py", line 429, in test_matrix_power_3 E return xp.linalg.matrix_power(a, 3) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/linalg/_product.py", line 384, in matrix_power E return cupy.matmul(cupy.matmul(M, M), M) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/linalg/_product.py", line 42, in matmul E return _gu_func_matmul(x1, x2, out=out, axes=axes) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/_core/_gufuncs.py", line 676, in __call__ E self._apply_func_to_inputs( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/_core/_gufuncs.py", line 393, in _apply_func_to_inputs E fouts = func(*args) E File "cupy/_core/_routines_linalg.pyx", line 691, in cupy._core._routines_linalg.matmul E cpdef ndarray matmul(ndarray a, ndarray b, ndarray out=None): E File "cupy/_core/_routines_linalg.pyx", line 732, in cupy._core._routines_linalg.matmul E return dot(a, b, out) E File "cupy/_core/_routines_linalg.pyx", line 427, in cupy._core._routines_linalg.dot E return tensordot_core(a, b, out, n, m, k, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 501, in cupy._core._routines_linalg.tensordot_core E return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 327, in cupy._core._routines_linalg._integral_tensordot_core E kern = _tensordot_core_int_kernel(config, dtype) E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/_routines_linalg.pyx", line 302, in cupy._core._routines_linalg._tensordot_core_int_kernel E ker = mod.get_function( E File "cupy/_core/raw.pyx", line 470, in cupy._core.raw.RawModule.get_function E mangled_name = self.module.mapping.get(name) E File "cupy/_core/raw.pyx", line 394, in cupy._core.raw.RawModule.module.__get__ E return self._module() E File "cupy/_core/raw.pyx", line 402, in cupy._core.raw.RawModule._module E mod = _get_raw_module( E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/raw.pyx", line 547, in cupy._core.raw._get_raw_module E mod = cupy._core.core.compile_with_cache( E File "cupy/_core/core.pyx", line 2015, in cupy._core.core.compile_with_cache E cpdef function.Module compile_with_cache( E File "cupy/_core/core.pyx", line 2078, in cupy._core.core.compile_with_cache E return cuda.compile_with_cache( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 453, in compile_with_cache E return _compile_with_cache_hip( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 853, in _compile_with_cache_hip E binary, mapping = compile_using_nvrtc( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 295, in compile_using_nvrtc E return _compile(source, options, cu_path, E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 279, in _compile E compiled_obj, mapping = prog.compile(options, log_stream) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 667, in compile E raise CompileException(log, self.src, self.name, options, E cupy.cuda.compiler.CompileException: /tmp/comgr-e19ac9/input/CompileSource:203:27: error: redeclaration of '__hiprtc_16' with a different type: 'void (*const)(int, int, int, const signed char *, const signed char *, signed char *)' vs 'void (*const)(int, int, int, const bool *, const bool *, bool *)' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-e19ac9/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-e19ac9/input/CompileSource:204:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-e19ac9/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-e19ac9/input/CompileSource:205:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-e19ac9/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-e19ac9/input/CompileSource:206:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-e19ac9/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-e19ac9/input/CompileSource:207:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-e19ac9/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-e19ac9/input/CompileSource:208:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-e19ac9/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-e19ac9/input/CompileSource:209:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-e19ac9/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-e19ac9/input/CompileSource:210:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-e19ac9/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-e19ac9/input/CompileSource:211:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-e19ac9/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-e19ac9/input/CompileSource:212:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-e19ac9/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E 10 errors generated when compiling for gfx908. E Error: Failed to compile opencl source (from CL or HIP source to LLVM IR). cupy/cuda/compiler.py:667: AssertionError ----------------------------- Captured stdout call ----------------------------- dtype is ___________________ TestMatrixPower.test_matrix_power_large ____________________ self = xp = dtype = @testing.for_all_dtypes() @testing.numpy_cupy_allclose() def test_matrix_power_large(self, xp, dtype): a = xp.eye(23, k=17, dtype=dtype) + xp.eye(23, k=-6, dtype=dtype) > return xp.linalg.matrix_power(a, 123456789123456789) tests/cupy_tests/linalg_tests/test_product.py:462: _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ cupy/linalg/_product.py:390: in matrix_power Z = M if Z is None else cupy.matmul(Z, Z) cupy/linalg/_product.py:42: in matmul return _gu_func_matmul(x1, x2, out=out, axes=axes) cupy/_core/_gufuncs.py:676: in __call__ self._apply_func_to_inputs( cupy/_core/_gufuncs.py:393: in _apply_func_to_inputs fouts = func(*args) cupy/_core/_routines_linalg.pyx:691: in cupy._core._routines_linalg.matmul cpdef ndarray matmul(ndarray a, ndarray b, ndarray out=None): cupy/_core/_routines_linalg.pyx:732: in cupy._core._routines_linalg.matmul return dot(a, b, out) cupy/_core/_routines_linalg.pyx:427: in cupy._core._routines_linalg.dot return tensordot_core(a, b, out, n, m, k, ret_shape) cupy/_core/_routines_linalg.pyx:501: in cupy._core._routines_linalg.tensordot_core return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) cupy/_core/_routines_linalg.pyx:327: in cupy._core._routines_linalg._integral_tensordot_core kern = _tensordot_core_int_kernel(config, dtype) cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/_routines_linalg.pyx:302: in cupy._core._routines_linalg._tensordot_core_int_kernel ker = mod.get_function( cupy/_core/raw.pyx:470: in cupy._core.raw.RawModule.get_function mangled_name = self.module.mapping.get(name) cupy/_core/raw.pyx:394: in cupy._core.raw.RawModule.module.__get__ return self._module() cupy/_core/raw.pyx:402: in cupy._core.raw.RawModule._module mod = _get_raw_module( cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/raw.pyx:547: in cupy._core.raw._get_raw_module mod = cupy._core.core.compile_with_cache( cupy/_core/core.pyx:2015: in cupy._core.core.compile_with_cache cpdef function.Module compile_with_cache( cupy/_core/core.pyx:2078: in cupy._core.core.compile_with_cache return cuda.compile_with_cache( cupy/cuda/compiler.py:453: in compile_with_cache return _compile_with_cache_hip( cupy/cuda/compiler.py:853: in _compile_with_cache_hip binary, mapping = compile_using_nvrtc( cupy/cuda/compiler.py:295: in compile_using_nvrtc return _compile(source, options, cu_path, cupy/cuda/compiler.py:279: in _compile compiled_obj, mapping = prog.compile(options, log_stream) _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ self = options = ('--std=c++11', '-I/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/_core/include', '-I/opt/rocm/include', '-fcuda-flush-denormals-to-zero') log_stream = None def compile(self, options=(), log_stream=None): try: if self.name_expressions: for ker in self.name_expressions: nvrtc.addAddNameExpression(self.ptr, ker) nvrtc.compileProgram(self.ptr, options) mapping = None if self.name_expressions: mapping = {} for ker in self.name_expressions: mapping[ker] = nvrtc.getLoweredName(self.ptr, ker) if log_stream is not None: log_stream.write(nvrtc.getProgramLog(self.ptr)) # This is to ensure backwards compatibility with nvrtc if self.method == 'cubin': return nvrtc.getCUBIN(self.ptr), mapping elif self.method == 'ptx': return nvrtc.getPTX(self.ptr), mapping else: raise RuntimeError('Unknown NVRTC compile method') except nvrtc.NVRTCError: log = nvrtc.getProgramLog(self.ptr) > raise CompileException(log, self.src, self.name, options, 'nvrtc' if not runtime.is_hip else 'hiprtc') E AssertionError: Only cupy raises error E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 650, in compile E nvrtc.compileProgram(self.ptr, options) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 136, in cupy_backends.cuda.libs.nvrtc.compileProgram E cpdef compileProgram(intptr_t prog, options): E File "cupy_backends/cuda/libs/nvrtc.pyx", line 148, in cupy_backends.cuda.libs.nvrtc.compileProgram E check_status(status) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 67, in cupy_backends.cuda.libs.nvrtc.check_status E raise NVRTCError(status) E cupy_backends.cuda.libs.nvrtc.NVRTCError: HIPRTC_ERROR_COMPILATION (6) E E During handling of the above exception, another exception occurred: E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/testing/_helper.py", line 47, in _call_func E result = impl(*args, **kw) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/tests/cupy_tests/linalg_tests/test_product.py", line 462, in test_matrix_power_large E return xp.linalg.matrix_power(a, 123456789123456789) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/linalg/_product.py", line 390, in matrix_power E Z = M if Z is None else cupy.matmul(Z, Z) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/linalg/_product.py", line 42, in matmul E return _gu_func_matmul(x1, x2, out=out, axes=axes) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/_core/_gufuncs.py", line 676, in __call__ E self._apply_func_to_inputs( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/_core/_gufuncs.py", line 393, in _apply_func_to_inputs E fouts = func(*args) E File "cupy/_core/_routines_linalg.pyx", line 691, in cupy._core._routines_linalg.matmul E cpdef ndarray matmul(ndarray a, ndarray b, ndarray out=None): E File "cupy/_core/_routines_linalg.pyx", line 732, in cupy._core._routines_linalg.matmul E return dot(a, b, out) E File "cupy/_core/_routines_linalg.pyx", line 427, in cupy._core._routines_linalg.dot E return tensordot_core(a, b, out, n, m, k, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 501, in cupy._core._routines_linalg.tensordot_core E return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 327, in cupy._core._routines_linalg._integral_tensordot_core E kern = _tensordot_core_int_kernel(config, dtype) E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/_routines_linalg.pyx", line 302, in cupy._core._routines_linalg._tensordot_core_int_kernel E ker = mod.get_function( E File "cupy/_core/raw.pyx", line 470, in cupy._core.raw.RawModule.get_function E mangled_name = self.module.mapping.get(name) E File "cupy/_core/raw.pyx", line 394, in cupy._core.raw.RawModule.module.__get__ E return self._module() E File "cupy/_core/raw.pyx", line 402, in cupy._core.raw.RawModule._module E mod = _get_raw_module( E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/raw.pyx", line 547, in cupy._core.raw._get_raw_module E mod = cupy._core.core.compile_with_cache( E File "cupy/_core/core.pyx", line 2015, in cupy._core.core.compile_with_cache E cpdef function.Module compile_with_cache( E File "cupy/_core/core.pyx", line 2078, in cupy._core.core.compile_with_cache E return cuda.compile_with_cache( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 453, in compile_with_cache E return _compile_with_cache_hip( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 853, in _compile_with_cache_hip E binary, mapping = compile_using_nvrtc( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 295, in compile_using_nvrtc E return _compile(source, options, cu_path, E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 279, in _compile E compiled_obj, mapping = prog.compile(options, log_stream) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 667, in compile E raise CompileException(log, self.src, self.name, options, E cupy.cuda.compiler.CompileException: /tmp/comgr-b8000b/input/CompileSource:203:27: error: redeclaration of '__hiprtc_16' with a different type: 'void (*const)(int, int, int, const signed char *, const signed char *, signed char *)' vs 'void (*const)(int, int, int, const bool *, const bool *, bool *)' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-b8000b/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-b8000b/input/CompileSource:204:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-b8000b/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-b8000b/input/CompileSource:205:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-b8000b/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-b8000b/input/CompileSource:206:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-b8000b/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-b8000b/input/CompileSource:207:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-b8000b/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-b8000b/input/CompileSource:208:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-b8000b/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-b8000b/input/CompileSource:209:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-b8000b/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-b8000b/input/CompileSource:210:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-b8000b/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-b8000b/input/CompileSource:211:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-b8000b/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-b8000b/input/CompileSource:212:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-b8000b/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E 10 errors generated when compiling for gfx908. E Error: Failed to compile opencl source (from CL or HIP source to LLVM IR). cupy/cuda/compiler.py:667: AssertionError ----------------------------- Captured stdout call ----------------------------- dtype is ___________________ TestMatrixPower.test_matrix_power_of_two ___________________ self = xp = dtype = @testing.for_all_dtypes() @testing.numpy_cupy_allclose() def test_matrix_power_of_two(self, xp, dtype): a = xp.eye(23, k=17, dtype=dtype) + xp.eye(23, k=-6, dtype=dtype) > return xp.linalg.matrix_power(a, 1 << 50) tests/cupy_tests/linalg_tests/test_product.py:456: _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ cupy/linalg/_product.py:390: in matrix_power Z = M if Z is None else cupy.matmul(Z, Z) cupy/linalg/_product.py:42: in matmul return _gu_func_matmul(x1, x2, out=out, axes=axes) cupy/_core/_gufuncs.py:676: in __call__ self._apply_func_to_inputs( cupy/_core/_gufuncs.py:393: in _apply_func_to_inputs fouts = func(*args) cupy/_core/_routines_linalg.pyx:691: in cupy._core._routines_linalg.matmul cpdef ndarray matmul(ndarray a, ndarray b, ndarray out=None): cupy/_core/_routines_linalg.pyx:732: in cupy._core._routines_linalg.matmul return dot(a, b, out) cupy/_core/_routines_linalg.pyx:427: in cupy._core._routines_linalg.dot return tensordot_core(a, b, out, n, m, k, ret_shape) cupy/_core/_routines_linalg.pyx:501: in cupy._core._routines_linalg.tensordot_core return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) cupy/_core/_routines_linalg.pyx:327: in cupy._core._routines_linalg._integral_tensordot_core kern = _tensordot_core_int_kernel(config, dtype) cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/_routines_linalg.pyx:302: in cupy._core._routines_linalg._tensordot_core_int_kernel ker = mod.get_function( cupy/_core/raw.pyx:470: in cupy._core.raw.RawModule.get_function mangled_name = self.module.mapping.get(name) cupy/_core/raw.pyx:394: in cupy._core.raw.RawModule.module.__get__ return self._module() cupy/_core/raw.pyx:402: in cupy._core.raw.RawModule._module mod = _get_raw_module( cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/raw.pyx:547: in cupy._core.raw._get_raw_module mod = cupy._core.core.compile_with_cache( cupy/_core/core.pyx:2015: in cupy._core.core.compile_with_cache cpdef function.Module compile_with_cache( cupy/_core/core.pyx:2078: in cupy._core.core.compile_with_cache return cuda.compile_with_cache( cupy/cuda/compiler.py:453: in compile_with_cache return _compile_with_cache_hip( cupy/cuda/compiler.py:853: in _compile_with_cache_hip binary, mapping = compile_using_nvrtc( cupy/cuda/compiler.py:295: in compile_using_nvrtc return _compile(source, options, cu_path, cupy/cuda/compiler.py:279: in _compile compiled_obj, mapping = prog.compile(options, log_stream) _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ self = options = ('--std=c++11', '-I/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/_core/include', '-I/opt/rocm/include', '-fcuda-flush-denormals-to-zero') log_stream = None def compile(self, options=(), log_stream=None): try: if self.name_expressions: for ker in self.name_expressions: nvrtc.addAddNameExpression(self.ptr, ker) nvrtc.compileProgram(self.ptr, options) mapping = None if self.name_expressions: mapping = {} for ker in self.name_expressions: mapping[ker] = nvrtc.getLoweredName(self.ptr, ker) if log_stream is not None: log_stream.write(nvrtc.getProgramLog(self.ptr)) # This is to ensure backwards compatibility with nvrtc if self.method == 'cubin': return nvrtc.getCUBIN(self.ptr), mapping elif self.method == 'ptx': return nvrtc.getPTX(self.ptr), mapping else: raise RuntimeError('Unknown NVRTC compile method') except nvrtc.NVRTCError: log = nvrtc.getProgramLog(self.ptr) > raise CompileException(log, self.src, self.name, options, 'nvrtc' if not runtime.is_hip else 'hiprtc') E AssertionError: Only cupy raises error E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 650, in compile E nvrtc.compileProgram(self.ptr, options) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 136, in cupy_backends.cuda.libs.nvrtc.compileProgram E cpdef compileProgram(intptr_t prog, options): E File "cupy_backends/cuda/libs/nvrtc.pyx", line 148, in cupy_backends.cuda.libs.nvrtc.compileProgram E check_status(status) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 67, in cupy_backends.cuda.libs.nvrtc.check_status E raise NVRTCError(status) E cupy_backends.cuda.libs.nvrtc.NVRTCError: HIPRTC_ERROR_COMPILATION (6) E E During handling of the above exception, another exception occurred: E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/testing/_helper.py", line 47, in _call_func E result = impl(*args, **kw) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/tests/cupy_tests/linalg_tests/test_product.py", line 456, in test_matrix_power_of_two E return xp.linalg.matrix_power(a, 1 << 50) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/linalg/_product.py", line 390, in matrix_power E Z = M if Z is None else cupy.matmul(Z, Z) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/linalg/_product.py", line 42, in matmul E return _gu_func_matmul(x1, x2, out=out, axes=axes) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/_core/_gufuncs.py", line 676, in __call__ E self._apply_func_to_inputs( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/_core/_gufuncs.py", line 393, in _apply_func_to_inputs E fouts = func(*args) E File "cupy/_core/_routines_linalg.pyx", line 691, in cupy._core._routines_linalg.matmul E cpdef ndarray matmul(ndarray a, ndarray b, ndarray out=None): E File "cupy/_core/_routines_linalg.pyx", line 732, in cupy._core._routines_linalg.matmul E return dot(a, b, out) E File "cupy/_core/_routines_linalg.pyx", line 427, in cupy._core._routines_linalg.dot E return tensordot_core(a, b, out, n, m, k, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 501, in cupy._core._routines_linalg.tensordot_core E return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 327, in cupy._core._routines_linalg._integral_tensordot_core E kern = _tensordot_core_int_kernel(config, dtype) E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/_routines_linalg.pyx", line 302, in cupy._core._routines_linalg._tensordot_core_int_kernel E ker = mod.get_function( E File "cupy/_core/raw.pyx", line 470, in cupy._core.raw.RawModule.get_function E mangled_name = self.module.mapping.get(name) E File "cupy/_core/raw.pyx", line 394, in cupy._core.raw.RawModule.module.__get__ E return self._module() E File "cupy/_core/raw.pyx", line 402, in cupy._core.raw.RawModule._module E mod = _get_raw_module( E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/raw.pyx", line 547, in cupy._core.raw._get_raw_module E mod = cupy._core.core.compile_with_cache( E File "cupy/_core/core.pyx", line 2015, in cupy._core.core.compile_with_cache E cpdef function.Module compile_with_cache( E File "cupy/_core/core.pyx", line 2078, in cupy._core.core.compile_with_cache E return cuda.compile_with_cache( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 453, in compile_with_cache E return _compile_with_cache_hip( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 853, in _compile_with_cache_hip E binary, mapping = compile_using_nvrtc( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 295, in compile_using_nvrtc E return _compile(source, options, cu_path, E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 279, in _compile E compiled_obj, mapping = prog.compile(options, log_stream) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 667, in compile E raise CompileException(log, self.src, self.name, options, E cupy.cuda.compiler.CompileException: /tmp/comgr-a9e8a8/input/CompileSource:203:27: error: redeclaration of '__hiprtc_16' with a different type: 'void (*const)(int, int, int, const signed char *, const signed char *, signed char *)' vs 'void (*const)(int, int, int, const bool *, const bool *, bool *)' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-a9e8a8/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-a9e8a8/input/CompileSource:204:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-a9e8a8/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-a9e8a8/input/CompileSource:205:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-a9e8a8/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-a9e8a8/input/CompileSource:206:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-a9e8a8/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-a9e8a8/input/CompileSource:207:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-a9e8a8/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-a9e8a8/input/CompileSource:208:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-a9e8a8/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-a9e8a8/input/CompileSource:209:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-a9e8a8/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-a9e8a8/input/CompileSource:210:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-a9e8a8/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-a9e8a8/input/CompileSource:211:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-a9e8a8/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-a9e8a8/input/CompileSource:212:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-a9e8a8/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E 10 errors generated when compiling for gfx908. E Error: Failed to compile opencl source (from CL or HIP source to LLVM IR). cupy/cuda/compiler.py:667: AssertionError ----------------------------- Captured stdout call ----------------------------- dtype is ________________________ TestLstsq.test_lstsq_solutions ________________________ self = def test_lstsq_solutions(self): # Comapres numpy.linalg.lstsq and cupy.linalg.lstsq solutions for: # a shapes range from (3, 3) to (5, 3) and (3, 5) # b shapes range from (i, 3) to (i, ) # sets a random seed for deterministic testing for i in range(3, 6): for j in range(3, 6): for k in range(2, 4): seed = i + j + k # check when b has shape (i, k) self.check_lstsq_solution((i, j), (i, k), seed, rcond=-1) self.check_lstsq_solution((i, j), (i, k), seed, rcond=None) self.check_lstsq_solution((i, j), (i, k), seed, rcond=0.5) > self.check_lstsq_solution((i, j), (i, k), seed, rcond=1e-6, singular=True) tests/cupy_tests/linalg_tests/test_solve.py:242: _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ cupy/testing/_helper.py:837: in test_func impl(*args, **kw) cupy/testing/_helper.py:273: in test_func _check_cupy_numpy_error(cupy_error, cupy/testing/_helper.py:179: in _check_cupy_numpy_error _fail_test_with_unexpected_errors( cupy/testing/_helper.py:134: in _fail_test_with_unexpected_errors raise AssertionError(msg).with_traceback(tb) cupy/testing/_helper.py:47: in _call_func result = impl(*args, **kw) tests/cupy_tests/linalg_tests/test_solve.py:203: in check_lstsq_solution a = xp.matmul( cupy/linalg/_product.py:42: in matmul return _gu_func_matmul(x1, x2, out=out, axes=axes) cupy/_core/_gufuncs.py:676: in __call__ self._apply_func_to_inputs( cupy/_core/_gufuncs.py:393: in _apply_func_to_inputs fouts = func(*args) cupy/_core/_routines_linalg.pyx:691: in cupy._core._routines_linalg.matmul cpdef ndarray matmul(ndarray a, ndarray b, ndarray out=None): cupy/_core/_routines_linalg.pyx:732: in cupy._core._routines_linalg.matmul return dot(a, b, out) cupy/_core/_routines_linalg.pyx:427: in cupy._core._routines_linalg.dot return tensordot_core(a, b, out, n, m, k, ret_shape) cupy/_core/_routines_linalg.pyx:501: in cupy._core._routines_linalg.tensordot_core return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) cupy/_core/_routines_linalg.pyx:327: in cupy._core._routines_linalg._integral_tensordot_core kern = _tensordot_core_int_kernel(config, dtype) cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/_routines_linalg.pyx:302: in cupy._core._routines_linalg._tensordot_core_int_kernel ker = mod.get_function( cupy/_core/raw.pyx:470: in cupy._core.raw.RawModule.get_function mangled_name = self.module.mapping.get(name) cupy/_core/raw.pyx:394: in cupy._core.raw.RawModule.module.__get__ return self._module() cupy/_core/raw.pyx:402: in cupy._core.raw.RawModule._module mod = _get_raw_module( cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/raw.pyx:547: in cupy._core.raw._get_raw_module mod = cupy._core.core.compile_with_cache( cupy/_core/core.pyx:2015: in cupy._core.core.compile_with_cache cpdef function.Module compile_with_cache( cupy/_core/core.pyx:2078: in cupy._core.core.compile_with_cache return cuda.compile_with_cache( cupy/cuda/compiler.py:453: in compile_with_cache return _compile_with_cache_hip( cupy/cuda/compiler.py:853: in _compile_with_cache_hip binary, mapping = compile_using_nvrtc( cupy/cuda/compiler.py:295: in compile_using_nvrtc return _compile(source, options, cu_path, cupy/cuda/compiler.py:279: in _compile compiled_obj, mapping = prog.compile(options, log_stream) _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ self = options = ('--std=c++11', '-I/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/_core/include', '-I/opt/rocm/include', '-fcuda-flush-denormals-to-zero') log_stream = None def compile(self, options=(), log_stream=None): try: if self.name_expressions: for ker in self.name_expressions: nvrtc.addAddNameExpression(self.ptr, ker) nvrtc.compileProgram(self.ptr, options) mapping = None if self.name_expressions: mapping = {} for ker in self.name_expressions: mapping[ker] = nvrtc.getLoweredName(self.ptr, ker) if log_stream is not None: log_stream.write(nvrtc.getProgramLog(self.ptr)) # This is to ensure backwards compatibility with nvrtc if self.method == 'cubin': return nvrtc.getCUBIN(self.ptr), mapping elif self.method == 'ptx': return nvrtc.getPTX(self.ptr), mapping else: raise RuntimeError('Unknown NVRTC compile method') except nvrtc.NVRTCError: log = nvrtc.getProgramLog(self.ptr) > raise CompileException(log, self.src, self.name, options, 'nvrtc' if not runtime.is_hip else 'hiprtc') E AssertionError: Only cupy raises error E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 650, in compile E nvrtc.compileProgram(self.ptr, options) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 136, in cupy_backends.cuda.libs.nvrtc.compileProgram E cpdef compileProgram(intptr_t prog, options): E File "cupy_backends/cuda/libs/nvrtc.pyx", line 148, in cupy_backends.cuda.libs.nvrtc.compileProgram E check_status(status) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 67, in cupy_backends.cuda.libs.nvrtc.check_status E raise NVRTCError(status) E cupy_backends.cuda.libs.nvrtc.NVRTCError: HIPRTC_ERROR_COMPILATION (6) E E During handling of the above exception, another exception occurred: E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/testing/_helper.py", line 47, in _call_func E result = impl(*args, **kw) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/tests/cupy_tests/linalg_tests/test_solve.py", line 203, in check_lstsq_solution E a = xp.matmul( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/linalg/_product.py", line 42, in matmul E return _gu_func_matmul(x1, x2, out=out, axes=axes) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/_core/_gufuncs.py", line 676, in __call__ E self._apply_func_to_inputs( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/_core/_gufuncs.py", line 393, in _apply_func_to_inputs E fouts = func(*args) E File "cupy/_core/_routines_linalg.pyx", line 691, in cupy._core._routines_linalg.matmul E cpdef ndarray matmul(ndarray a, ndarray b, ndarray out=None): E File "cupy/_core/_routines_linalg.pyx", line 732, in cupy._core._routines_linalg.matmul E return dot(a, b, out) E File "cupy/_core/_routines_linalg.pyx", line 427, in cupy._core._routines_linalg.dot E return tensordot_core(a, b, out, n, m, k, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 501, in cupy._core._routines_linalg.tensordot_core E return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 327, in cupy._core._routines_linalg._integral_tensordot_core E kern = _tensordot_core_int_kernel(config, dtype) E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/_routines_linalg.pyx", line 302, in cupy._core._routines_linalg._tensordot_core_int_kernel E ker = mod.get_function( E File "cupy/_core/raw.pyx", line 470, in cupy._core.raw.RawModule.get_function E mangled_name = self.module.mapping.get(name) E File "cupy/_core/raw.pyx", line 394, in cupy._core.raw.RawModule.module.__get__ E return self._module() E File "cupy/_core/raw.pyx", line 402, in cupy._core.raw.RawModule._module E mod = _get_raw_module( E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/raw.pyx", line 547, in cupy._core.raw._get_raw_module E mod = cupy._core.core.compile_with_cache( E File "cupy/_core/core.pyx", line 2015, in cupy._core.core.compile_with_cache E cpdef function.Module compile_with_cache( E File "cupy/_core/core.pyx", line 2078, in cupy._core.core.compile_with_cache E return cuda.compile_with_cache( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 453, in compile_with_cache E return _compile_with_cache_hip( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 853, in _compile_with_cache_hip E binary, mapping = compile_using_nvrtc( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 295, in compile_using_nvrtc E return _compile(source, options, cu_path, E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 279, in _compile E compiled_obj, mapping = prog.compile(options, log_stream) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 667, in compile E raise CompileException(log, self.src, self.name, options, E cupy.cuda.compiler.CompileException: /tmp/comgr-f7ebcd/input/CompileSource:203:27: error: redeclaration of '__hiprtc_16' with a different type: 'void (*const)(int, int, int, const signed char *, const signed char *, signed char *)' vs 'void (*const)(int, int, int, const bool *, const bool *, bool *)' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-f7ebcd/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-f7ebcd/input/CompileSource:204:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-f7ebcd/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-f7ebcd/input/CompileSource:205:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-f7ebcd/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-f7ebcd/input/CompileSource:206:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-f7ebcd/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-f7ebcd/input/CompileSource:207:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-f7ebcd/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-f7ebcd/input/CompileSource:208:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-f7ebcd/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-f7ebcd/input/CompileSource:209:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-f7ebcd/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-f7ebcd/input/CompileSource:210:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-f7ebcd/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-f7ebcd/input/CompileSource:211:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-f7ebcd/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-f7ebcd/input/CompileSource:212:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-f7ebcd/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E 10 errors generated when compiling for gfx908. E Error: Failed to compile opencl source (from CL or HIP source to LLVM IR). cupy/cuda/compiler.py:667: AssertionError ----------------------------- Captured stdout call ----------------------------- dtype is i ______ TestMatmul_param_0_{shape_pair=((3, 2), (2, 4))}.test_cupy_matmul _______ self = < parameter: {'shape_pair': ((3, 2), (2, 4))}> xp = dtype1 = , dtype2 = @testing.for_all_dtypes(name='dtype1') @testing.for_all_dtypes(name='dtype2') @testing.numpy_cupy_allclose(rtol=1e-3, atol=1e-3) # required for uint8 def test_cupy_matmul(self, xp, dtype1, dtype2): x1 = testing.shaped_arange(self.shape_pair[0], xp, dtype1) x2 = testing.shaped_arange(self.shape_pair[1], xp, dtype2) > return xp.matmul(x1, x2) tests/cupy_tests/math_tests/test_matmul.py:75: _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ cupy/linalg/_product.py:42: in matmul return _gu_func_matmul(x1, x2, out=out, axes=axes) cupy/_core/_gufuncs.py:676: in __call__ self._apply_func_to_inputs( cupy/_core/_gufuncs.py:393: in _apply_func_to_inputs fouts = func(*args) cupy/_core/_routines_linalg.pyx:691: in cupy._core._routines_linalg.matmul cpdef ndarray matmul(ndarray a, ndarray b, ndarray out=None): cupy/_core/_routines_linalg.pyx:732: in cupy._core._routines_linalg.matmul return dot(a, b, out) cupy/_core/_routines_linalg.pyx:427: in cupy._core._routines_linalg.dot return tensordot_core(a, b, out, n, m, k, ret_shape) cupy/_core/_routines_linalg.pyx:501: in cupy._core._routines_linalg.tensordot_core return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) cupy/_core/_routines_linalg.pyx:327: in cupy._core._routines_linalg._integral_tensordot_core kern = _tensordot_core_int_kernel(config, dtype) cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/_routines_linalg.pyx:302: in cupy._core._routines_linalg._tensordot_core_int_kernel ker = mod.get_function( cupy/_core/raw.pyx:470: in cupy._core.raw.RawModule.get_function mangled_name = self.module.mapping.get(name) cupy/_core/raw.pyx:394: in cupy._core.raw.RawModule.module.__get__ return self._module() cupy/_core/raw.pyx:402: in cupy._core.raw.RawModule._module mod = _get_raw_module( cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/raw.pyx:547: in cupy._core.raw._get_raw_module mod = cupy._core.core.compile_with_cache( cupy/_core/core.pyx:2015: in cupy._core.core.compile_with_cache cpdef function.Module compile_with_cache( cupy/_core/core.pyx:2078: in cupy._core.core.compile_with_cache return cuda.compile_with_cache( cupy/cuda/compiler.py:453: in compile_with_cache return _compile_with_cache_hip( cupy/cuda/compiler.py:853: in _compile_with_cache_hip binary, mapping = compile_using_nvrtc( cupy/cuda/compiler.py:295: in compile_using_nvrtc return _compile(source, options, cu_path, cupy/cuda/compiler.py:279: in _compile compiled_obj, mapping = prog.compile(options, log_stream) _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ self = options = ('--std=c++11', '-I/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/_core/include', '-I/opt/rocm/include', '-fcuda-flush-denormals-to-zero') log_stream = None def compile(self, options=(), log_stream=None): try: if self.name_expressions: for ker in self.name_expressions: nvrtc.addAddNameExpression(self.ptr, ker) nvrtc.compileProgram(self.ptr, options) mapping = None if self.name_expressions: mapping = {} for ker in self.name_expressions: mapping[ker] = nvrtc.getLoweredName(self.ptr, ker) if log_stream is not None: log_stream.write(nvrtc.getProgramLog(self.ptr)) # This is to ensure backwards compatibility with nvrtc if self.method == 'cubin': return nvrtc.getCUBIN(self.ptr), mapping elif self.method == 'ptx': return nvrtc.getPTX(self.ptr), mapping else: raise RuntimeError('Unknown NVRTC compile method') except nvrtc.NVRTCError: log = nvrtc.getProgramLog(self.ptr) > raise CompileException(log, self.src, self.name, options, 'nvrtc' if not runtime.is_hip else 'hiprtc') E AssertionError: Only cupy raises error E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 650, in compile E nvrtc.compileProgram(self.ptr, options) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 136, in cupy_backends.cuda.libs.nvrtc.compileProgram E cpdef compileProgram(intptr_t prog, options): E File "cupy_backends/cuda/libs/nvrtc.pyx", line 148, in cupy_backends.cuda.libs.nvrtc.compileProgram E check_status(status) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 67, in cupy_backends.cuda.libs.nvrtc.check_status E raise NVRTCError(status) E cupy_backends.cuda.libs.nvrtc.NVRTCError: HIPRTC_ERROR_COMPILATION (6) E E During handling of the above exception, another exception occurred: E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/testing/_helper.py", line 47, in _call_func E result = impl(*args, **kw) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/tests/cupy_tests/math_tests/test_matmul.py", line 75, in test_cupy_matmul E return xp.matmul(x1, x2) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/linalg/_product.py", line 42, in matmul E return _gu_func_matmul(x1, x2, out=out, axes=axes) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/_core/_gufuncs.py", line 676, in __call__ E self._apply_func_to_inputs( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/_core/_gufuncs.py", line 393, in _apply_func_to_inputs E fouts = func(*args) E File "cupy/_core/_routines_linalg.pyx", line 691, in cupy._core._routines_linalg.matmul E cpdef ndarray matmul(ndarray a, ndarray b, ndarray out=None): E File "cupy/_core/_routines_linalg.pyx", line 732, in cupy._core._routines_linalg.matmul E return dot(a, b, out) E File "cupy/_core/_routines_linalg.pyx", line 427, in cupy._core._routines_linalg.dot E return tensordot_core(a, b, out, n, m, k, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 501, in cupy._core._routines_linalg.tensordot_core E return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 327, in cupy._core._routines_linalg._integral_tensordot_core E kern = _tensordot_core_int_kernel(config, dtype) E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/_routines_linalg.pyx", line 302, in cupy._core._routines_linalg._tensordot_core_int_kernel E ker = mod.get_function( E File "cupy/_core/raw.pyx", line 470, in cupy._core.raw.RawModule.get_function E mangled_name = self.module.mapping.get(name) E File "cupy/_core/raw.pyx", line 394, in cupy._core.raw.RawModule.module.__get__ E return self._module() E File "cupy/_core/raw.pyx", line 402, in cupy._core.raw.RawModule._module E mod = _get_raw_module( E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/raw.pyx", line 547, in cupy._core.raw._get_raw_module E mod = cupy._core.core.compile_with_cache( E File "cupy/_core/core.pyx", line 2015, in cupy._core.core.compile_with_cache E cpdef function.Module compile_with_cache( E File "cupy/_core/core.pyx", line 2078, in cupy._core.core.compile_with_cache E return cuda.compile_with_cache( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 453, in compile_with_cache E return _compile_with_cache_hip( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 853, in _compile_with_cache_hip E binary, mapping = compile_using_nvrtc( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 295, in compile_using_nvrtc E return _compile(source, options, cu_path, E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 279, in _compile E compiled_obj, mapping = prog.compile(options, log_stream) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 667, in compile E raise CompileException(log, self.src, self.name, options, E cupy.cuda.compiler.CompileException: /tmp/comgr-b9b843/input/CompileSource:203:27: error: redeclaration of '__hiprtc_16' with a different type: 'void (*const)(int, int, int, const signed char *, const signed char *, signed char *)' vs 'void (*const)(int, int, int, const bool *, const bool *, bool *)' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-b9b843/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-b9b843/input/CompileSource:204:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-b9b843/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-b9b843/input/CompileSource:205:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-b9b843/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-b9b843/input/CompileSource:206:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-b9b843/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-b9b843/input/CompileSource:207:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-b9b843/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-b9b843/input/CompileSource:208:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-b9b843/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-b9b843/input/CompileSource:209:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-b9b843/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-b9b843/input/CompileSource:210:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-b9b843/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-b9b843/input/CompileSource:211:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-b9b843/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-b9b843/input/CompileSource:212:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-b9b843/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E 10 errors generated when compiling for gfx908. E Error: Failed to compile opencl source (from CL or HIP source to LLVM IR). cupy/cuda/compiler.py:667: AssertionError ----------------------------- Captured stdout call ----------------------------- dtype2 is dtype1 is ____ TestMatmul_param_0_{shape_pair=((3, 2), (2, 4))}.test_operator_matmul _____ self = < parameter: {'shape_pair': ((3, 2), (2, 4))}> xp = dtype1 = , dtype2 = @testing.for_all_dtypes(name='dtype1') @testing.for_all_dtypes(name='dtype2') @testing.numpy_cupy_allclose(rtol=1e-3, atol=1e-3) # required for uint8 def test_operator_matmul(self, xp, dtype1, dtype2): x1 = testing.shaped_arange(self.shape_pair[0], xp, dtype1) x2 = testing.shaped_arange(self.shape_pair[1], xp, dtype2) > return operator.matmul(x1, x2) tests/cupy_tests/math_tests/test_matmul.py:67: _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ cupy/_core/core.pyx:1171: in cupy._core.core.ndarray.__matmul__ return cupy.linalg._product.matmul(x, y) cupy/linalg/_product.py:42: in matmul return _gu_func_matmul(x1, x2, out=out, axes=axes) cupy/_core/_gufuncs.py:676: in __call__ self._apply_func_to_inputs( cupy/_core/_gufuncs.py:393: in _apply_func_to_inputs fouts = func(*args) cupy/_core/_routines_linalg.pyx:691: in cupy._core._routines_linalg.matmul cpdef ndarray matmul(ndarray a, ndarray b, ndarray out=None): cupy/_core/_routines_linalg.pyx:732: in cupy._core._routines_linalg.matmul return dot(a, b, out) cupy/_core/_routines_linalg.pyx:427: in cupy._core._routines_linalg.dot return tensordot_core(a, b, out, n, m, k, ret_shape) cupy/_core/_routines_linalg.pyx:501: in cupy._core._routines_linalg.tensordot_core return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) cupy/_core/_routines_linalg.pyx:327: in cupy._core._routines_linalg._integral_tensordot_core kern = _tensordot_core_int_kernel(config, dtype) cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/_routines_linalg.pyx:302: in cupy._core._routines_linalg._tensordot_core_int_kernel ker = mod.get_function( cupy/_core/raw.pyx:470: in cupy._core.raw.RawModule.get_function mangled_name = self.module.mapping.get(name) cupy/_core/raw.pyx:394: in cupy._core.raw.RawModule.module.__get__ return self._module() cupy/_core/raw.pyx:402: in cupy._core.raw.RawModule._module mod = _get_raw_module( cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/raw.pyx:547: in cupy._core.raw._get_raw_module mod = cupy._core.core.compile_with_cache( cupy/_core/core.pyx:2015: in cupy._core.core.compile_with_cache cpdef function.Module compile_with_cache( cupy/_core/core.pyx:2078: in cupy._core.core.compile_with_cache return cuda.compile_with_cache( cupy/cuda/compiler.py:453: in compile_with_cache return _compile_with_cache_hip( cupy/cuda/compiler.py:853: in _compile_with_cache_hip binary, mapping = compile_using_nvrtc( cupy/cuda/compiler.py:295: in compile_using_nvrtc return _compile(source, options, cu_path, cupy/cuda/compiler.py:279: in _compile compiled_obj, mapping = prog.compile(options, log_stream) _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ self = options = ('--std=c++11', '-I/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/_core/include', '-I/opt/rocm/include', '-fcuda-flush-denormals-to-zero') log_stream = None def compile(self, options=(), log_stream=None): try: if self.name_expressions: for ker in self.name_expressions: nvrtc.addAddNameExpression(self.ptr, ker) nvrtc.compileProgram(self.ptr, options) mapping = None if self.name_expressions: mapping = {} for ker in self.name_expressions: mapping[ker] = nvrtc.getLoweredName(self.ptr, ker) if log_stream is not None: log_stream.write(nvrtc.getProgramLog(self.ptr)) # This is to ensure backwards compatibility with nvrtc if self.method == 'cubin': return nvrtc.getCUBIN(self.ptr), mapping elif self.method == 'ptx': return nvrtc.getPTX(self.ptr), mapping else: raise RuntimeError('Unknown NVRTC compile method') except nvrtc.NVRTCError: log = nvrtc.getProgramLog(self.ptr) > raise CompileException(log, self.src, self.name, options, 'nvrtc' if not runtime.is_hip else 'hiprtc') E AssertionError: Only cupy raises error E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 650, in compile E nvrtc.compileProgram(self.ptr, options) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 136, in cupy_backends.cuda.libs.nvrtc.compileProgram E cpdef compileProgram(intptr_t prog, options): E File "cupy_backends/cuda/libs/nvrtc.pyx", line 148, in cupy_backends.cuda.libs.nvrtc.compileProgram E check_status(status) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 67, in cupy_backends.cuda.libs.nvrtc.check_status E raise NVRTCError(status) E cupy_backends.cuda.libs.nvrtc.NVRTCError: HIPRTC_ERROR_COMPILATION (6) E E During handling of the above exception, another exception occurred: E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/testing/_helper.py", line 47, in _call_func E result = impl(*args, **kw) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/tests/cupy_tests/math_tests/test_matmul.py", line 67, in test_operator_matmul E return operator.matmul(x1, x2) E File "cupy/_core/core.pyx", line 1171, in cupy._core.core.ndarray.__matmul__ E return cupy.linalg._product.matmul(x, y) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/linalg/_product.py", line 42, in matmul E return _gu_func_matmul(x1, x2, out=out, axes=axes) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/_core/_gufuncs.py", line 676, in __call__ E self._apply_func_to_inputs( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/_core/_gufuncs.py", line 393, in _apply_func_to_inputs E fouts = func(*args) E File "cupy/_core/_routines_linalg.pyx", line 691, in cupy._core._routines_linalg.matmul E cpdef ndarray matmul(ndarray a, ndarray b, ndarray out=None): E File "cupy/_core/_routines_linalg.pyx", line 732, in cupy._core._routines_linalg.matmul E return dot(a, b, out) E File "cupy/_core/_routines_linalg.pyx", line 427, in cupy._core._routines_linalg.dot E return tensordot_core(a, b, out, n, m, k, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 501, in cupy._core._routines_linalg.tensordot_core E return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 327, in cupy._core._routines_linalg._integral_tensordot_core E kern = _tensordot_core_int_kernel(config, dtype) E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/_routines_linalg.pyx", line 302, in cupy._core._routines_linalg._tensordot_core_int_kernel E ker = mod.get_function( E File "cupy/_core/raw.pyx", line 470, in cupy._core.raw.RawModule.get_function E mangled_name = self.module.mapping.get(name) E File "cupy/_core/raw.pyx", line 394, in cupy._core.raw.RawModule.module.__get__ E return self._module() E File "cupy/_core/raw.pyx", line 402, in cupy._core.raw.RawModule._module E mod = _get_raw_module( E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/raw.pyx", line 547, in cupy._core.raw._get_raw_module E mod = cupy._core.core.compile_with_cache( E File "cupy/_core/core.pyx", line 2015, in cupy._core.core.compile_with_cache E cpdef function.Module compile_with_cache( E File "cupy/_core/core.pyx", line 2078, in cupy._core.core.compile_with_cache E return cuda.compile_with_cache( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 453, in compile_with_cache E return _compile_with_cache_hip( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 853, in _compile_with_cache_hip E binary, mapping = compile_using_nvrtc( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 295, in compile_using_nvrtc E return _compile(source, options, cu_path, E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 279, in _compile E compiled_obj, mapping = prog.compile(options, log_stream) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 667, in compile E raise CompileException(log, self.src, self.name, options, E cupy.cuda.compiler.CompileException: /tmp/comgr-44ff30/input/CompileSource:203:27: error: redeclaration of '__hiprtc_16' with a different type: 'void (*const)(int, int, int, const signed char *, const signed char *, signed char *)' vs 'void (*const)(int, int, int, const bool *, const bool *, bool *)' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-44ff30/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-44ff30/input/CompileSource:204:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-44ff30/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-44ff30/input/CompileSource:205:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-44ff30/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-44ff30/input/CompileSource:206:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-44ff30/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-44ff30/input/CompileSource:207:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-44ff30/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-44ff30/input/CompileSource:208:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-44ff30/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-44ff30/input/CompileSource:209:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-44ff30/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-44ff30/input/CompileSource:210:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-44ff30/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-44ff30/input/CompileSource:211:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-44ff30/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-44ff30/input/CompileSource:212:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-44ff30/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E 10 errors generated when compiling for gfx908. E Error: Failed to compile opencl source (from CL or HIP source to LLVM IR). cupy/cuda/compiler.py:667: AssertionError ----------------------------- Captured stdout call ----------------------------- dtype2 is dtype1 is _______ TestMatmul_param_4_{shape_pair=((2,), (2, 4))}.test_cupy_matmul ________ self = < parameter: {'shape_pair': ((2,), (2, 4))}> xp = dtype1 = , dtype2 = @testing.for_all_dtypes(name='dtype1') @testing.for_all_dtypes(name='dtype2') @testing.numpy_cupy_allclose(rtol=1e-3, atol=1e-3) # required for uint8 def test_cupy_matmul(self, xp, dtype1, dtype2): x1 = testing.shaped_arange(self.shape_pair[0], xp, dtype1) x2 = testing.shaped_arange(self.shape_pair[1], xp, dtype2) > return xp.matmul(x1, x2) tests/cupy_tests/math_tests/test_matmul.py:75: _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ cupy/linalg/_product.py:42: in matmul return _gu_func_matmul(x1, x2, out=out, axes=axes) cupy/_core/_gufuncs.py:676: in __call__ self._apply_func_to_inputs( cupy/_core/_gufuncs.py:393: in _apply_func_to_inputs fouts = func(*args) cupy/_core/_routines_linalg.pyx:691: in cupy._core._routines_linalg.matmul cpdef ndarray matmul(ndarray a, ndarray b, ndarray out=None): cupy/_core/_routines_linalg.pyx:732: in cupy._core._routines_linalg.matmul return dot(a, b, out) cupy/_core/_routines_linalg.pyx:427: in cupy._core._routines_linalg.dot return tensordot_core(a, b, out, n, m, k, ret_shape) cupy/_core/_routines_linalg.pyx:501: in cupy._core._routines_linalg.tensordot_core return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) cupy/_core/_routines_linalg.pyx:327: in cupy._core._routines_linalg._integral_tensordot_core kern = _tensordot_core_int_kernel(config, dtype) cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/_routines_linalg.pyx:302: in cupy._core._routines_linalg._tensordot_core_int_kernel ker = mod.get_function( cupy/_core/raw.pyx:470: in cupy._core.raw.RawModule.get_function mangled_name = self.module.mapping.get(name) cupy/_core/raw.pyx:394: in cupy._core.raw.RawModule.module.__get__ return self._module() cupy/_core/raw.pyx:402: in cupy._core.raw.RawModule._module mod = _get_raw_module( cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/raw.pyx:547: in cupy._core.raw._get_raw_module mod = cupy._core.core.compile_with_cache( cupy/_core/core.pyx:2015: in cupy._core.core.compile_with_cache cpdef function.Module compile_with_cache( cupy/_core/core.pyx:2078: in cupy._core.core.compile_with_cache return cuda.compile_with_cache( cupy/cuda/compiler.py:453: in compile_with_cache return _compile_with_cache_hip( cupy/cuda/compiler.py:853: in _compile_with_cache_hip binary, mapping = compile_using_nvrtc( cupy/cuda/compiler.py:295: in compile_using_nvrtc return _compile(source, options, cu_path, cupy/cuda/compiler.py:279: in _compile compiled_obj, mapping = prog.compile(options, log_stream) _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ self = options = ('--std=c++11', '-I/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/_core/include', '-I/opt/rocm/include', '-fcuda-flush-denormals-to-zero') log_stream = None def compile(self, options=(), log_stream=None): try: if self.name_expressions: for ker in self.name_expressions: nvrtc.addAddNameExpression(self.ptr, ker) nvrtc.compileProgram(self.ptr, options) mapping = None if self.name_expressions: mapping = {} for ker in self.name_expressions: mapping[ker] = nvrtc.getLoweredName(self.ptr, ker) if log_stream is not None: log_stream.write(nvrtc.getProgramLog(self.ptr)) # This is to ensure backwards compatibility with nvrtc if self.method == 'cubin': return nvrtc.getCUBIN(self.ptr), mapping elif self.method == 'ptx': return nvrtc.getPTX(self.ptr), mapping else: raise RuntimeError('Unknown NVRTC compile method') except nvrtc.NVRTCError: log = nvrtc.getProgramLog(self.ptr) > raise CompileException(log, self.src, self.name, options, 'nvrtc' if not runtime.is_hip else 'hiprtc') E AssertionError: Only cupy raises error E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 650, in compile E nvrtc.compileProgram(self.ptr, options) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 136, in cupy_backends.cuda.libs.nvrtc.compileProgram E cpdef compileProgram(intptr_t prog, options): E File "cupy_backends/cuda/libs/nvrtc.pyx", line 148, in cupy_backends.cuda.libs.nvrtc.compileProgram E check_status(status) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 67, in cupy_backends.cuda.libs.nvrtc.check_status E raise NVRTCError(status) E cupy_backends.cuda.libs.nvrtc.NVRTCError: HIPRTC_ERROR_COMPILATION (6) E E During handling of the above exception, another exception occurred: E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/testing/_helper.py", line 47, in _call_func E result = impl(*args, **kw) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/tests/cupy_tests/math_tests/test_matmul.py", line 75, in test_cupy_matmul E return xp.matmul(x1, x2) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/linalg/_product.py", line 42, in matmul E return _gu_func_matmul(x1, x2, out=out, axes=axes) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/_core/_gufuncs.py", line 676, in __call__ E self._apply_func_to_inputs( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/_core/_gufuncs.py", line 393, in _apply_func_to_inputs E fouts = func(*args) E File "cupy/_core/_routines_linalg.pyx", line 691, in cupy._core._routines_linalg.matmul E cpdef ndarray matmul(ndarray a, ndarray b, ndarray out=None): E File "cupy/_core/_routines_linalg.pyx", line 732, in cupy._core._routines_linalg.matmul E return dot(a, b, out) E File "cupy/_core/_routines_linalg.pyx", line 427, in cupy._core._routines_linalg.dot E return tensordot_core(a, b, out, n, m, k, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 501, in cupy._core._routines_linalg.tensordot_core E return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 327, in cupy._core._routines_linalg._integral_tensordot_core E kern = _tensordot_core_int_kernel(config, dtype) E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/_routines_linalg.pyx", line 302, in cupy._core._routines_linalg._tensordot_core_int_kernel E ker = mod.get_function( E File "cupy/_core/raw.pyx", line 470, in cupy._core.raw.RawModule.get_function E mangled_name = self.module.mapping.get(name) E File "cupy/_core/raw.pyx", line 394, in cupy._core.raw.RawModule.module.__get__ E return self._module() E File "cupy/_core/raw.pyx", line 402, in cupy._core.raw.RawModule._module E mod = _get_raw_module( E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/raw.pyx", line 547, in cupy._core.raw._get_raw_module E mod = cupy._core.core.compile_with_cache( E File "cupy/_core/core.pyx", line 2015, in cupy._core.core.compile_with_cache E cpdef function.Module compile_with_cache( E File "cupy/_core/core.pyx", line 2078, in cupy._core.core.compile_with_cache E return cuda.compile_with_cache( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 453, in compile_with_cache E return _compile_with_cache_hip( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 853, in _compile_with_cache_hip E binary, mapping = compile_using_nvrtc( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 295, in compile_using_nvrtc E return _compile(source, options, cu_path, E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 279, in _compile E compiled_obj, mapping = prog.compile(options, log_stream) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 667, in compile E raise CompileException(log, self.src, self.name, options, E cupy.cuda.compiler.CompileException: /tmp/comgr-fc1b5d/input/CompileSource:203:27: error: redeclaration of '__hiprtc_16' with a different type: 'void (*const)(int, int, int, const signed char *, const signed char *, signed char *)' vs 'void (*const)(int, int, int, const bool *, const bool *, bool *)' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-fc1b5d/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-fc1b5d/input/CompileSource:204:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-fc1b5d/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-fc1b5d/input/CompileSource:205:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-fc1b5d/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-fc1b5d/input/CompileSource:206:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-fc1b5d/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-fc1b5d/input/CompileSource:207:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-fc1b5d/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-fc1b5d/input/CompileSource:208:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-fc1b5d/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-fc1b5d/input/CompileSource:209:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-fc1b5d/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-fc1b5d/input/CompileSource:210:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-fc1b5d/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-fc1b5d/input/CompileSource:211:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-fc1b5d/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-fc1b5d/input/CompileSource:212:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-fc1b5d/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E 10 errors generated when compiling for gfx908. E Error: Failed to compile opencl source (from CL or HIP source to LLVM IR). cupy/cuda/compiler.py:667: AssertionError ----------------------------- Captured stdout call ----------------------------- dtype2 is dtype1 is _____ TestMatmul_param_4_{shape_pair=((2,), (2, 4))}.test_operator_matmul ______ self = < parameter: {'shape_pair': ((2,), (2, 4))}> xp = dtype1 = , dtype2 = @testing.for_all_dtypes(name='dtype1') @testing.for_all_dtypes(name='dtype2') @testing.numpy_cupy_allclose(rtol=1e-3, atol=1e-3) # required for uint8 def test_operator_matmul(self, xp, dtype1, dtype2): x1 = testing.shaped_arange(self.shape_pair[0], xp, dtype1) x2 = testing.shaped_arange(self.shape_pair[1], xp, dtype2) > return operator.matmul(x1, x2) tests/cupy_tests/math_tests/test_matmul.py:67: _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ cupy/_core/core.pyx:1171: in cupy._core.core.ndarray.__matmul__ return cupy.linalg._product.matmul(x, y) cupy/linalg/_product.py:42: in matmul return _gu_func_matmul(x1, x2, out=out, axes=axes) cupy/_core/_gufuncs.py:676: in __call__ self._apply_func_to_inputs( cupy/_core/_gufuncs.py:393: in _apply_func_to_inputs fouts = func(*args) cupy/_core/_routines_linalg.pyx:691: in cupy._core._routines_linalg.matmul cpdef ndarray matmul(ndarray a, ndarray b, ndarray out=None): cupy/_core/_routines_linalg.pyx:732: in cupy._core._routines_linalg.matmul return dot(a, b, out) cupy/_core/_routines_linalg.pyx:427: in cupy._core._routines_linalg.dot return tensordot_core(a, b, out, n, m, k, ret_shape) cupy/_core/_routines_linalg.pyx:501: in cupy._core._routines_linalg.tensordot_core return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) cupy/_core/_routines_linalg.pyx:327: in cupy._core._routines_linalg._integral_tensordot_core kern = _tensordot_core_int_kernel(config, dtype) cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/_routines_linalg.pyx:302: in cupy._core._routines_linalg._tensordot_core_int_kernel ker = mod.get_function( cupy/_core/raw.pyx:470: in cupy._core.raw.RawModule.get_function mangled_name = self.module.mapping.get(name) cupy/_core/raw.pyx:394: in cupy._core.raw.RawModule.module.__get__ return self._module() cupy/_core/raw.pyx:402: in cupy._core.raw.RawModule._module mod = _get_raw_module( cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/raw.pyx:547: in cupy._core.raw._get_raw_module mod = cupy._core.core.compile_with_cache( cupy/_core/core.pyx:2015: in cupy._core.core.compile_with_cache cpdef function.Module compile_with_cache( cupy/_core/core.pyx:2078: in cupy._core.core.compile_with_cache return cuda.compile_with_cache( cupy/cuda/compiler.py:453: in compile_with_cache return _compile_with_cache_hip( cupy/cuda/compiler.py:853: in _compile_with_cache_hip binary, mapping = compile_using_nvrtc( cupy/cuda/compiler.py:295: in compile_using_nvrtc return _compile(source, options, cu_path, cupy/cuda/compiler.py:279: in _compile compiled_obj, mapping = prog.compile(options, log_stream) _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ self = options = ('--std=c++11', '-I/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/_core/include', '-I/opt/rocm/include', '-fcuda-flush-denormals-to-zero') log_stream = None def compile(self, options=(), log_stream=None): try: if self.name_expressions: for ker in self.name_expressions: nvrtc.addAddNameExpression(self.ptr, ker) nvrtc.compileProgram(self.ptr, options) mapping = None if self.name_expressions: mapping = {} for ker in self.name_expressions: mapping[ker] = nvrtc.getLoweredName(self.ptr, ker) if log_stream is not None: log_stream.write(nvrtc.getProgramLog(self.ptr)) # This is to ensure backwards compatibility with nvrtc if self.method == 'cubin': return nvrtc.getCUBIN(self.ptr), mapping elif self.method == 'ptx': return nvrtc.getPTX(self.ptr), mapping else: raise RuntimeError('Unknown NVRTC compile method') except nvrtc.NVRTCError: log = nvrtc.getProgramLog(self.ptr) > raise CompileException(log, self.src, self.name, options, 'nvrtc' if not runtime.is_hip else 'hiprtc') E AssertionError: Only cupy raises error E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 650, in compile E nvrtc.compileProgram(self.ptr, options) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 136, in cupy_backends.cuda.libs.nvrtc.compileProgram E cpdef compileProgram(intptr_t prog, options): E File "cupy_backends/cuda/libs/nvrtc.pyx", line 148, in cupy_backends.cuda.libs.nvrtc.compileProgram E check_status(status) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 67, in cupy_backends.cuda.libs.nvrtc.check_status E raise NVRTCError(status) E cupy_backends.cuda.libs.nvrtc.NVRTCError: HIPRTC_ERROR_COMPILATION (6) E E During handling of the above exception, another exception occurred: E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/testing/_helper.py", line 47, in _call_func E result = impl(*args, **kw) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/tests/cupy_tests/math_tests/test_matmul.py", line 67, in test_operator_matmul E return operator.matmul(x1, x2) E File "cupy/_core/core.pyx", line 1171, in cupy._core.core.ndarray.__matmul__ E return cupy.linalg._product.matmul(x, y) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/linalg/_product.py", line 42, in matmul E return _gu_func_matmul(x1, x2, out=out, axes=axes) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/_core/_gufuncs.py", line 676, in __call__ E self._apply_func_to_inputs( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/_core/_gufuncs.py", line 393, in _apply_func_to_inputs E fouts = func(*args) E File "cupy/_core/_routines_linalg.pyx", line 691, in cupy._core._routines_linalg.matmul E cpdef ndarray matmul(ndarray a, ndarray b, ndarray out=None): E File "cupy/_core/_routines_linalg.pyx", line 732, in cupy._core._routines_linalg.matmul E return dot(a, b, out) E File "cupy/_core/_routines_linalg.pyx", line 427, in cupy._core._routines_linalg.dot E return tensordot_core(a, b, out, n, m, k, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 501, in cupy._core._routines_linalg.tensordot_core E return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 327, in cupy._core._routines_linalg._integral_tensordot_core E kern = _tensordot_core_int_kernel(config, dtype) E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/_routines_linalg.pyx", line 302, in cupy._core._routines_linalg._tensordot_core_int_kernel E ker = mod.get_function( E File "cupy/_core/raw.pyx", line 470, in cupy._core.raw.RawModule.get_function E mangled_name = self.module.mapping.get(name) E File "cupy/_core/raw.pyx", line 394, in cupy._core.raw.RawModule.module.__get__ E return self._module() E File "cupy/_core/raw.pyx", line 402, in cupy._core.raw.RawModule._module E mod = _get_raw_module( E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/raw.pyx", line 547, in cupy._core.raw._get_raw_module E mod = cupy._core.core.compile_with_cache( E File "cupy/_core/core.pyx", line 2015, in cupy._core.core.compile_with_cache E cpdef function.Module compile_with_cache( E File "cupy/_core/core.pyx", line 2078, in cupy._core.core.compile_with_cache E return cuda.compile_with_cache( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 453, in compile_with_cache E return _compile_with_cache_hip( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 853, in _compile_with_cache_hip E binary, mapping = compile_using_nvrtc( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 295, in compile_using_nvrtc E return _compile(source, options, cu_path, E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 279, in _compile E compiled_obj, mapping = prog.compile(options, log_stream) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 667, in compile E raise CompileException(log, self.src, self.name, options, E cupy.cuda.compiler.CompileException: /tmp/comgr-bedf2c/input/CompileSource:203:27: error: redeclaration of '__hiprtc_16' with a different type: 'void (*const)(int, int, int, const signed char *, const signed char *, signed char *)' vs 'void (*const)(int, int, int, const bool *, const bool *, bool *)' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-bedf2c/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-bedf2c/input/CompileSource:204:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-bedf2c/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-bedf2c/input/CompileSource:205:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-bedf2c/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-bedf2c/input/CompileSource:206:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-bedf2c/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-bedf2c/input/CompileSource:207:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-bedf2c/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-bedf2c/input/CompileSource:208:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-bedf2c/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-bedf2c/input/CompileSource:209:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-bedf2c/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-bedf2c/input/CompileSource:210:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-bedf2c/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-bedf2c/input/CompileSource:211:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-bedf2c/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-bedf2c/input/CompileSource:212:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-bedf2c/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E 10 errors generated when compiling for gfx908. E Error: Failed to compile opencl source (from CL or HIP source to LLVM IR). cupy/cuda/compiler.py:667: AssertionError ----------------------------- Captured stdout call ----------------------------- dtype2 is dtype1 is _______ TestMatmul_param_6_{shape_pair=((3, 2), (2,))}.test_cupy_matmul ________ self = < parameter: {'shape_pair': ((3, 2), (2,))}> xp = dtype1 = , dtype2 = @testing.for_all_dtypes(name='dtype1') @testing.for_all_dtypes(name='dtype2') @testing.numpy_cupy_allclose(rtol=1e-3, atol=1e-3) # required for uint8 def test_cupy_matmul(self, xp, dtype1, dtype2): x1 = testing.shaped_arange(self.shape_pair[0], xp, dtype1) x2 = testing.shaped_arange(self.shape_pair[1], xp, dtype2) > return xp.matmul(x1, x2) tests/cupy_tests/math_tests/test_matmul.py:75: _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ cupy/linalg/_product.py:42: in matmul return _gu_func_matmul(x1, x2, out=out, axes=axes) cupy/_core/_gufuncs.py:676: in __call__ self._apply_func_to_inputs( cupy/_core/_gufuncs.py:393: in _apply_func_to_inputs fouts = func(*args) cupy/_core/_routines_linalg.pyx:691: in cupy._core._routines_linalg.matmul cpdef ndarray matmul(ndarray a, ndarray b, ndarray out=None): cupy/_core/_routines_linalg.pyx:732: in cupy._core._routines_linalg.matmul return dot(a, b, out) cupy/_core/_routines_linalg.pyx:427: in cupy._core._routines_linalg.dot return tensordot_core(a, b, out, n, m, k, ret_shape) cupy/_core/_routines_linalg.pyx:501: in cupy._core._routines_linalg.tensordot_core return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) cupy/_core/_routines_linalg.pyx:327: in cupy._core._routines_linalg._integral_tensordot_core kern = _tensordot_core_int_kernel(config, dtype) cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/_routines_linalg.pyx:302: in cupy._core._routines_linalg._tensordot_core_int_kernel ker = mod.get_function( cupy/_core/raw.pyx:470: in cupy._core.raw.RawModule.get_function mangled_name = self.module.mapping.get(name) cupy/_core/raw.pyx:394: in cupy._core.raw.RawModule.module.__get__ return self._module() cupy/_core/raw.pyx:402: in cupy._core.raw.RawModule._module mod = _get_raw_module( cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/raw.pyx:547: in cupy._core.raw._get_raw_module mod = cupy._core.core.compile_with_cache( cupy/_core/core.pyx:2015: in cupy._core.core.compile_with_cache cpdef function.Module compile_with_cache( cupy/_core/core.pyx:2078: in cupy._core.core.compile_with_cache return cuda.compile_with_cache( cupy/cuda/compiler.py:453: in compile_with_cache return _compile_with_cache_hip( cupy/cuda/compiler.py:853: in _compile_with_cache_hip binary, mapping = compile_using_nvrtc( cupy/cuda/compiler.py:295: in compile_using_nvrtc return _compile(source, options, cu_path, cupy/cuda/compiler.py:279: in _compile compiled_obj, mapping = prog.compile(options, log_stream) _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ self = options = ('--std=c++11', '-I/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/_core/include', '-I/opt/rocm/include', '-fcuda-flush-denormals-to-zero') log_stream = None def compile(self, options=(), log_stream=None): try: if self.name_expressions: for ker in self.name_expressions: nvrtc.addAddNameExpression(self.ptr, ker) nvrtc.compileProgram(self.ptr, options) mapping = None if self.name_expressions: mapping = {} for ker in self.name_expressions: mapping[ker] = nvrtc.getLoweredName(self.ptr, ker) if log_stream is not None: log_stream.write(nvrtc.getProgramLog(self.ptr)) # This is to ensure backwards compatibility with nvrtc if self.method == 'cubin': return nvrtc.getCUBIN(self.ptr), mapping elif self.method == 'ptx': return nvrtc.getPTX(self.ptr), mapping else: raise RuntimeError('Unknown NVRTC compile method') except nvrtc.NVRTCError: log = nvrtc.getProgramLog(self.ptr) > raise CompileException(log, self.src, self.name, options, 'nvrtc' if not runtime.is_hip else 'hiprtc') E AssertionError: Only cupy raises error E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 650, in compile E nvrtc.compileProgram(self.ptr, options) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 136, in cupy_backends.cuda.libs.nvrtc.compileProgram E cpdef compileProgram(intptr_t prog, options): E File "cupy_backends/cuda/libs/nvrtc.pyx", line 148, in cupy_backends.cuda.libs.nvrtc.compileProgram E check_status(status) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 67, in cupy_backends.cuda.libs.nvrtc.check_status E raise NVRTCError(status) E cupy_backends.cuda.libs.nvrtc.NVRTCError: HIPRTC_ERROR_COMPILATION (6) E E During handling of the above exception, another exception occurred: E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/testing/_helper.py", line 47, in _call_func E result = impl(*args, **kw) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/tests/cupy_tests/math_tests/test_matmul.py", line 75, in test_cupy_matmul E return xp.matmul(x1, x2) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/linalg/_product.py", line 42, in matmul E return _gu_func_matmul(x1, x2, out=out, axes=axes) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/_core/_gufuncs.py", line 676, in __call__ E self._apply_func_to_inputs( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/_core/_gufuncs.py", line 393, in _apply_func_to_inputs E fouts = func(*args) E File "cupy/_core/_routines_linalg.pyx", line 691, in cupy._core._routines_linalg.matmul E cpdef ndarray matmul(ndarray a, ndarray b, ndarray out=None): E File "cupy/_core/_routines_linalg.pyx", line 732, in cupy._core._routines_linalg.matmul E return dot(a, b, out) E File "cupy/_core/_routines_linalg.pyx", line 427, in cupy._core._routines_linalg.dot E return tensordot_core(a, b, out, n, m, k, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 501, in cupy._core._routines_linalg.tensordot_core E return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 327, in cupy._core._routines_linalg._integral_tensordot_core E kern = _tensordot_core_int_kernel(config, dtype) E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/_routines_linalg.pyx", line 302, in cupy._core._routines_linalg._tensordot_core_int_kernel E ker = mod.get_function( E File "cupy/_core/raw.pyx", line 470, in cupy._core.raw.RawModule.get_function E mangled_name = self.module.mapping.get(name) E File "cupy/_core/raw.pyx", line 394, in cupy._core.raw.RawModule.module.__get__ E return self._module() E File "cupy/_core/raw.pyx", line 402, in cupy._core.raw.RawModule._module E mod = _get_raw_module( E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/raw.pyx", line 547, in cupy._core.raw._get_raw_module E mod = cupy._core.core.compile_with_cache( E File "cupy/_core/core.pyx", line 2015, in cupy._core.core.compile_with_cache E cpdef function.Module compile_with_cache( E File "cupy/_core/core.pyx", line 2078, in cupy._core.core.compile_with_cache E return cuda.compile_with_cache( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 453, in compile_with_cache E return _compile_with_cache_hip( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 853, in _compile_with_cache_hip E binary, mapping = compile_using_nvrtc( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 295, in compile_using_nvrtc E return _compile(source, options, cu_path, E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 279, in _compile E compiled_obj, mapping = prog.compile(options, log_stream) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 667, in compile E raise CompileException(log, self.src, self.name, options, E cupy.cuda.compiler.CompileException: /tmp/comgr-b4eb0f/input/CompileSource:203:27: error: redeclaration of '__hiprtc_16' with a different type: 'void (*const)(int, int, int, const signed char *, const signed char *, signed char *)' vs 'void (*const)(int, int, int, const bool *, const bool *, bool *)' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-b4eb0f/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-b4eb0f/input/CompileSource:204:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-b4eb0f/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-b4eb0f/input/CompileSource:205:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-b4eb0f/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-b4eb0f/input/CompileSource:206:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-b4eb0f/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-b4eb0f/input/CompileSource:207:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-b4eb0f/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-b4eb0f/input/CompileSource:208:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-b4eb0f/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-b4eb0f/input/CompileSource:209:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-b4eb0f/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-b4eb0f/input/CompileSource:210:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-b4eb0f/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-b4eb0f/input/CompileSource:211:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-b4eb0f/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-b4eb0f/input/CompileSource:212:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-b4eb0f/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E 10 errors generated when compiling for gfx908. E Error: Failed to compile opencl source (from CL or HIP source to LLVM IR). cupy/cuda/compiler.py:667: AssertionError ----------------------------- Captured stdout call ----------------------------- dtype2 is dtype1 is _____ TestMatmul_param_6_{shape_pair=((3, 2), (2,))}.test_operator_matmul ______ self = < parameter: {'shape_pair': ((3, 2), (2,))}> xp = dtype1 = , dtype2 = @testing.for_all_dtypes(name='dtype1') @testing.for_all_dtypes(name='dtype2') @testing.numpy_cupy_allclose(rtol=1e-3, atol=1e-3) # required for uint8 def test_operator_matmul(self, xp, dtype1, dtype2): x1 = testing.shaped_arange(self.shape_pair[0], xp, dtype1) x2 = testing.shaped_arange(self.shape_pair[1], xp, dtype2) > return operator.matmul(x1, x2) tests/cupy_tests/math_tests/test_matmul.py:67: _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ cupy/_core/core.pyx:1171: in cupy._core.core.ndarray.__matmul__ return cupy.linalg._product.matmul(x, y) cupy/linalg/_product.py:42: in matmul return _gu_func_matmul(x1, x2, out=out, axes=axes) cupy/_core/_gufuncs.py:676: in __call__ self._apply_func_to_inputs( cupy/_core/_gufuncs.py:393: in _apply_func_to_inputs fouts = func(*args) cupy/_core/_routines_linalg.pyx:691: in cupy._core._routines_linalg.matmul cpdef ndarray matmul(ndarray a, ndarray b, ndarray out=None): cupy/_core/_routines_linalg.pyx:732: in cupy._core._routines_linalg.matmul return dot(a, b, out) cupy/_core/_routines_linalg.pyx:427: in cupy._core._routines_linalg.dot return tensordot_core(a, b, out, n, m, k, ret_shape) cupy/_core/_routines_linalg.pyx:501: in cupy._core._routines_linalg.tensordot_core return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) cupy/_core/_routines_linalg.pyx:327: in cupy._core._routines_linalg._integral_tensordot_core kern = _tensordot_core_int_kernel(config, dtype) cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/_routines_linalg.pyx:302: in cupy._core._routines_linalg._tensordot_core_int_kernel ker = mod.get_function( cupy/_core/raw.pyx:470: in cupy._core.raw.RawModule.get_function mangled_name = self.module.mapping.get(name) cupy/_core/raw.pyx:394: in cupy._core.raw.RawModule.module.__get__ return self._module() cupy/_core/raw.pyx:402: in cupy._core.raw.RawModule._module mod = _get_raw_module( cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/raw.pyx:547: in cupy._core.raw._get_raw_module mod = cupy._core.core.compile_with_cache( cupy/_core/core.pyx:2015: in cupy._core.core.compile_with_cache cpdef function.Module compile_with_cache( cupy/_core/core.pyx:2078: in cupy._core.core.compile_with_cache return cuda.compile_with_cache( cupy/cuda/compiler.py:453: in compile_with_cache return _compile_with_cache_hip( cupy/cuda/compiler.py:853: in _compile_with_cache_hip binary, mapping = compile_using_nvrtc( cupy/cuda/compiler.py:295: in compile_using_nvrtc return _compile(source, options, cu_path, cupy/cuda/compiler.py:279: in _compile compiled_obj, mapping = prog.compile(options, log_stream) _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ self = options = ('--std=c++11', '-I/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/_core/include', '-I/opt/rocm/include', '-fcuda-flush-denormals-to-zero') log_stream = None def compile(self, options=(), log_stream=None): try: if self.name_expressions: for ker in self.name_expressions: nvrtc.addAddNameExpression(self.ptr, ker) nvrtc.compileProgram(self.ptr, options) mapping = None if self.name_expressions: mapping = {} for ker in self.name_expressions: mapping[ker] = nvrtc.getLoweredName(self.ptr, ker) if log_stream is not None: log_stream.write(nvrtc.getProgramLog(self.ptr)) # This is to ensure backwards compatibility with nvrtc if self.method == 'cubin': return nvrtc.getCUBIN(self.ptr), mapping elif self.method == 'ptx': return nvrtc.getPTX(self.ptr), mapping else: raise RuntimeError('Unknown NVRTC compile method') except nvrtc.NVRTCError: log = nvrtc.getProgramLog(self.ptr) > raise CompileException(log, self.src, self.name, options, 'nvrtc' if not runtime.is_hip else 'hiprtc') E AssertionError: Only cupy raises error E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 650, in compile E nvrtc.compileProgram(self.ptr, options) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 136, in cupy_backends.cuda.libs.nvrtc.compileProgram E cpdef compileProgram(intptr_t prog, options): E File "cupy_backends/cuda/libs/nvrtc.pyx", line 148, in cupy_backends.cuda.libs.nvrtc.compileProgram E check_status(status) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 67, in cupy_backends.cuda.libs.nvrtc.check_status E raise NVRTCError(status) E cupy_backends.cuda.libs.nvrtc.NVRTCError: HIPRTC_ERROR_COMPILATION (6) E E During handling of the above exception, another exception occurred: E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/testing/_helper.py", line 47, in _call_func E result = impl(*args, **kw) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/tests/cupy_tests/math_tests/test_matmul.py", line 67, in test_operator_matmul E return operator.matmul(x1, x2) E File "cupy/_core/core.pyx", line 1171, in cupy._core.core.ndarray.__matmul__ E return cupy.linalg._product.matmul(x, y) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/linalg/_product.py", line 42, in matmul E return _gu_func_matmul(x1, x2, out=out, axes=axes) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/_core/_gufuncs.py", line 676, in __call__ E self._apply_func_to_inputs( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/_core/_gufuncs.py", line 393, in _apply_func_to_inputs E fouts = func(*args) E File "cupy/_core/_routines_linalg.pyx", line 691, in cupy._core._routines_linalg.matmul E cpdef ndarray matmul(ndarray a, ndarray b, ndarray out=None): E File "cupy/_core/_routines_linalg.pyx", line 732, in cupy._core._routines_linalg.matmul E return dot(a, b, out) E File "cupy/_core/_routines_linalg.pyx", line 427, in cupy._core._routines_linalg.dot E return tensordot_core(a, b, out, n, m, k, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 501, in cupy._core._routines_linalg.tensordot_core E return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 327, in cupy._core._routines_linalg._integral_tensordot_core E kern = _tensordot_core_int_kernel(config, dtype) E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/_routines_linalg.pyx", line 302, in cupy._core._routines_linalg._tensordot_core_int_kernel E ker = mod.get_function( E File "cupy/_core/raw.pyx", line 470, in cupy._core.raw.RawModule.get_function E mangled_name = self.module.mapping.get(name) E File "cupy/_core/raw.pyx", line 394, in cupy._core.raw.RawModule.module.__get__ E return self._module() E File "cupy/_core/raw.pyx", line 402, in cupy._core.raw.RawModule._module E mod = _get_raw_module( E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/raw.pyx", line 547, in cupy._core.raw._get_raw_module E mod = cupy._core.core.compile_with_cache( E File "cupy/_core/core.pyx", line 2015, in cupy._core.core.compile_with_cache E cpdef function.Module compile_with_cache( E File "cupy/_core/core.pyx", line 2078, in cupy._core.core.compile_with_cache E return cuda.compile_with_cache( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 453, in compile_with_cache E return _compile_with_cache_hip( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 853, in _compile_with_cache_hip E binary, mapping = compile_using_nvrtc( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 295, in compile_using_nvrtc E return _compile(source, options, cu_path, E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 279, in _compile E compiled_obj, mapping = prog.compile(options, log_stream) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 667, in compile E raise CompileException(log, self.src, self.name, options, E cupy.cuda.compiler.CompileException: /tmp/comgr-6c11af/input/CompileSource:203:27: error: redeclaration of '__hiprtc_16' with a different type: 'void (*const)(int, int, int, const signed char *, const signed char *, signed char *)' vs 'void (*const)(int, int, int, const bool *, const bool *, bool *)' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-6c11af/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-6c11af/input/CompileSource:204:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-6c11af/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-6c11af/input/CompileSource:205:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-6c11af/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-6c11af/input/CompileSource:206:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-6c11af/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-6c11af/input/CompileSource:207:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-6c11af/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-6c11af/input/CompileSource:208:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-6c11af/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-6c11af/input/CompileSource:209:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-6c11af/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-6c11af/input/CompileSource:210:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-6c11af/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-6c11af/input/CompileSource:211:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-6c11af/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-6c11af/input/CompileSource:212:27: error: redefinition of '__hiprtc_16' E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-6c11af/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_16 = _tensordot_core_int_kernel; E ^ E 10 errors generated when compiling for gfx908. E Error: Failed to compile opencl source (from CL or HIP source to LLVM IR). cupy/cuda/compiler.py:667: AssertionError ----------------------------- Captured stdout call ----------------------------- dtype2 is dtype1 is _________ TestConvolve_param_0_{mode='valid'}.test_convolve_diff_types _________ [XPASS(strict)] HIP/ROCm may have a bug with larger `b` _________ TestConvolve_param_1_{mode='same'}.test_convolve_diff_types __________ [XPASS(strict)] HIP/ROCm may have a bug with larger `b` _________ TestConvolve_param_2_{mode='full'}.test_convolve_diff_types __________ [XPASS(strict)] HIP/ROCm may have a bug with larger `b` ________ TestCorrelate_param_0_{mode='valid'}.test_correlate_diff_types ________ [XPASS(strict)] ROCm/HIP may have a bug ________ TestCorrelate_param_1_{mode='full'}.test_correlate_diff_types _________ [XPASS(strict)] ROCm/HIP may have a bug ________ TestCorrelate_param_2_{mode='same'}.test_correlate_diff_types _________ [XPASS(strict)] ROCm/HIP may have a bug _ TestMapCoordinatesHalfInteger.test_map_coordinates_float[_param_46_{mode='grid-constant', order=4}] _ args = (< parameter: {'mode': 'grid-constant', 'order': 4}>,) kw = {'dtype': }, dtype = @_wraps_partial(impl, name) def test_func(*args, **kw): for dtype in dtypes: try: kw[name] = numpy.dtype(dtype).type > impl(*args, **kw) cupy/testing/_helper.py:837: _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ cupy/testing/_helper.py:355: in test_func check_func(cupy_r, numpy_r) cupy/testing/_helper.py:507: in check_func _array.assert_allclose(c, n, rtol1, atol1, err_msg, verbose) _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ actual = array([229.76997843, 598.96953239]) desired = array([3.06993892, 8. ]), rtol = 1e-07, atol = 0.0001 err_msg = '', verbose = True def assert_allclose(actual, desired, rtol=1e-7, atol=0, err_msg='', verbose=True): """Raises an AssertionError if objects are not equal up to desired tolerance. Args: actual(numpy.ndarray or cupy.ndarray): The actual object to check. desired(numpy.ndarray or cupy.ndarray): The desired, expected object. rtol(float): Relative tolerance. atol(float): Absolute tolerance. err_msg(str): The error message to be printed in case of failure. verbose(bool): If ``True``, the conflicting values are appended to the error message. .. seealso:: :func:`numpy.testing.assert_allclose` """ # NOQA > numpy.testing.assert_allclose( cupy.asnumpy(actual), cupy.asnumpy(desired), rtol=rtol, atol=atol, err_msg=err_msg, verbose=verbose) E AssertionError: E Not equal to tolerance rtol=1e-07, atol=0.0001 E E Mismatched elements: 2 / 2 (100%) E Max absolute difference: 590.96953239 E Max relative difference: 73.87119155 E x: array([229.769978, 598.969532]) E y: array([3.069939, 8. ]) cupy/testing/_array.py:24: AssertionError ----------------------------- Captured stdout call ----------------------------- dtype is _ TestMapCoordinatesHalfInteger.test_map_coordinates_float[_param_47_{mode='grid-constant', order=5}] _ args = (< parameter: {'mode': 'grid-constant', 'order': 5}>,) kw = {'dtype': }, dtype = @_wraps_partial(impl, name) def test_func(*args, **kw): for dtype in dtypes: try: kw[name] = numpy.dtype(dtype).type > impl(*args, **kw) cupy/testing/_helper.py:837: _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ cupy/testing/_helper.py:355: in test_func check_func(cupy_r, numpy_r) cupy/testing/_helper.py:507: in check_func _array.assert_allclose(c, n, rtol1, atol1, err_msg, verbose) _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ actual = array([ 78.42728447, 201.97560613]) desired = array([3.12122783, 8. ]), rtol = 1e-07, atol = 0.0001 err_msg = '', verbose = True def assert_allclose(actual, desired, rtol=1e-7, atol=0, err_msg='', verbose=True): """Raises an AssertionError if objects are not equal up to desired tolerance. Args: actual(numpy.ndarray or cupy.ndarray): The actual object to check. desired(numpy.ndarray or cupy.ndarray): The desired, expected object. rtol(float): Relative tolerance. atol(float): Absolute tolerance. err_msg(str): The error message to be printed in case of failure. verbose(bool): If ``True``, the conflicting values are appended to the error message. .. seealso:: :func:`numpy.testing.assert_allclose` """ # NOQA > numpy.testing.assert_allclose( cupy.asnumpy(actual), cupy.asnumpy(desired), rtol=rtol, atol=atol, err_msg=err_msg, verbose=verbose) E AssertionError: E Not equal to tolerance rtol=1e-07, atol=0.0001 E E Mismatched elements: 2 / 2 (100%) E Max absolute difference: 193.97560613 E Max relative difference: 24.24695077 E x: array([ 78.427284, 201.975606]) E y: array([3.121228, 8. ]) cupy/testing/_array.py:24: AssertionError ----------------------------- Captured stdout call ----------------------------- dtype is _ TestZoomOutputSize1.test_zoom_output_size1[_param_4_{grid_mode=False, mode='constant', order=4, shape=(5, 5, 2), zoom=(2, 2, 0.5)}] _ args = (< parameter: {'grid_mode': False, 'mode': 'constant', 'order': 4, 'shape': (5, 5, 2), 'zoom': (2, 2, 0.5)}>,) kw = {'dtype': }, dtype = @_wraps_partial(impl, name) def test_func(*args, **kw): for dtype in dtypes: try: kw[name] = numpy.dtype(dtype).type > impl(*args, **kw) cupy/testing/_helper.py:837: _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ cupy/testing/_helper.py:355: in test_func check_func(cupy_r, numpy_r) cupy/testing/_helper.py:507: in check_func _array.assert_allclose(c, n, rtol1, atol1, err_msg, verbose) _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ actual = array([[[-9.26215027e-04], [-7.12910551e-05], [-1.04548770e-03], [ 6.02170185e-04], [ ... [-1.58413204e-05], [-2.26292832e-05], [-1.52809316e-05], [-2.62873400e-05]]], dtype=float32) desired = array([[[-1.91677706e-16], [ 6.02007280e-16], [ 7.39644969e-18], [ 9.42839520e-17], [-... [ 2.49078413e-18], [-6.86374620e-18], [-2.01807717e-17], [-1.36889723e-17]]], dtype=float32) rtol = 1e-07, atol = 1e-05, err_msg = '', verbose = True def assert_allclose(actual, desired, rtol=1e-7, atol=0, err_msg='', verbose=True): """Raises an AssertionError if objects are not equal up to desired tolerance. Args: actual(numpy.ndarray or cupy.ndarray): The actual object to check. desired(numpy.ndarray or cupy.ndarray): The desired, expected object. rtol(float): Relative tolerance. atol(float): Absolute tolerance. err_msg(str): The error message to be printed in case of failure. verbose(bool): If ``True``, the conflicting values are appended to the error message. .. seealso:: :func:`numpy.testing.assert_allclose` """ # NOQA > numpy.testing.assert_allclose( cupy.asnumpy(actual), cupy.asnumpy(desired), rtol=rtol, atol=atol, err_msg=err_msg, verbose=verbose) E AssertionError: E Not equal to tolerance rtol=1e-07, atol=1e-05 E E Mismatched elements: 95 / 100 (95%) E Max absolute difference: 0.00124037 E Max relative difference: 1.6367563e+14 E x: array([[[-9.262150e-04], E [-7.129106e-05], E [-1.045488e-03],... E y: array([[[-1.916777e-16], E [ 6.020073e-16], E [ 7.396450e-18],... cupy/testing/_array.py:24: AssertionError ----------------------------- Captured stdout call ----------------------------- dtype is _ TestZoomOutputSize1.test_zoom_output_size1[_param_5_{grid_mode=False, mode='constant', order=5, shape=(5, 5, 2), zoom=(2, 2, 0.5)}] _ args = (< parameter: {'grid_mode': False, 'mode': 'constant', 'order': 5, 'shape': (5, 5, 2), 'zoom': (2, 2, 0.5)}>,) kw = {'dtype': }, dtype = @_wraps_partial(impl, name) def test_func(*args, **kw): for dtype in dtypes: try: kw[name] = numpy.dtype(dtype).type > impl(*args, **kw) cupy/testing/_helper.py:837: _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ cupy/testing/_helper.py:355: in test_func check_func(cupy_r, numpy_r) cupy/testing/_helper.py:507: in check_func _array.assert_allclose(c, n, rtol1, atol1, err_msg, verbose) _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ actual = array([[[-6.23873319e-04], [-6.52701594e-04], [-5.34185849e-04], [-7.05319864e-04], [-... [-2.50156190e-05], [ 9.55036171e-07], [-1.27298597e-06], [-4.17136789e-06]]], dtype=float32) desired = array([[[ 2.1000416e-16], [-7.7505151e-16], [-6.6172874e-17], [-6.5140899e-16], [ 7.95... [-1.5922411e-17], [ 3.1236140e-17], [-5.3157008e-17], [-1.3124647e-17]]], dtype=float32) rtol = 1e-07, atol = 1e-05, err_msg = '', verbose = True def assert_allclose(actual, desired, rtol=1e-7, atol=0, err_msg='', verbose=True): """Raises an AssertionError if objects are not equal up to desired tolerance. Args: actual(numpy.ndarray or cupy.ndarray): The actual object to check. desired(numpy.ndarray or cupy.ndarray): The desired, expected object. rtol(float): Relative tolerance. atol(float): Absolute tolerance. err_msg(str): The error message to be printed in case of failure. verbose(bool): If ``True``, the conflicting values are appended to the error message. .. seealso:: :func:`numpy.testing.assert_allclose` """ # NOQA > numpy.testing.assert_allclose( cupy.asnumpy(actual), cupy.asnumpy(desired), rtol=rtol, atol=atol, err_msg=err_msg, verbose=verbose) E AssertionError: E Not equal to tolerance rtol=1e-07, atol=1e-05 E E Mismatched elements: 85 / 100 (85%) E Max absolute difference: 0.00084527 E Max relative difference: 1.881553e+13 E x: array([[[-6.238733e-04], E [-6.527016e-04], E [-5.341858e-04],... E y: array([[[ 2.100042e-16], E [-7.750515e-16], E [-6.617287e-17],... cupy/testing/_array.py:24: AssertionError ----------------------------- Captured stdout call ----------------------------- dtype is _ TestZoomOutputSize1.test_zoom_output_size1[_param_16_{grid_mode=False, mode='mirror', order=4, shape=(5, 5, 2), zoom=(2, 2, 0.5)}] _ args = (< parameter: {'grid_mode': False, 'mode': 'mirror', 'order': 4, 'shape': (5, 5, 2), 'zoom': (2, 2, 0.5)}>,) kw = {'dtype': }, dtype = @_wraps_partial(impl, name) def test_func(*args, **kw): for dtype in dtypes: try: kw[name] = numpy.dtype(dtype).type > impl(*args, **kw) cupy/testing/_helper.py:837: _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ cupy/testing/_helper.py:355: in test_func check_func(cupy_r, numpy_r) cupy/testing/_helper.py:507: in check_func _array.assert_allclose(c, n, rtol1, atol1, err_msg, verbose) _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ actual = array([[[-9.26215027e-04], [-7.12910551e-05], [-1.04548770e-03], [ 6.02170185e-04], [ ... [-1.58413204e-05], [-2.26292832e-05], [-1.52809316e-05], [-2.62873400e-05]]], dtype=float32) desired = array([[[-1.91677706e-16], [ 6.02007280e-16], [ 7.39644969e-18], [ 9.42839520e-17], [-... [ 2.49078413e-18], [-6.86374620e-18], [-2.01807717e-17], [-1.36889723e-17]]], dtype=float32) rtol = 1e-07, atol = 1e-05, err_msg = '', verbose = True def assert_allclose(actual, desired, rtol=1e-7, atol=0, err_msg='', verbose=True): """Raises an AssertionError if objects are not equal up to desired tolerance. Args: actual(numpy.ndarray or cupy.ndarray): The actual object to check. desired(numpy.ndarray or cupy.ndarray): The desired, expected object. rtol(float): Relative tolerance. atol(float): Absolute tolerance. err_msg(str): The error message to be printed in case of failure. verbose(bool): If ``True``, the conflicting values are appended to the error message. .. seealso:: :func:`numpy.testing.assert_allclose` """ # NOQA > numpy.testing.assert_allclose( cupy.asnumpy(actual), cupy.asnumpy(desired), rtol=rtol, atol=atol, err_msg=err_msg, verbose=verbose) E AssertionError: E Not equal to tolerance rtol=1e-07, atol=1e-05 E E Mismatched elements: 95 / 100 (95%) E Max absolute difference: 0.00124037 E Max relative difference: 1.6367563e+14 E x: array([[[-9.262150e-04], E [-7.129106e-05], E [-1.045488e-03],... E y: array([[[-1.916777e-16], E [ 6.020073e-16], E [ 7.396450e-18],... cupy/testing/_array.py:24: AssertionError ----------------------------- Captured stdout call ----------------------------- dtype is _ TestZoomOutputSize1.test_zoom_output_size1[_param_17_{grid_mode=False, mode='mirror', order=5, shape=(5, 5, 2), zoom=(2, 2, 0.5)}] _ args = (< parameter: {'grid_mode': False, 'mode': 'mirror', 'order': 5, 'shape': (5, 5, 2), 'zoom': (2, 2, 0.5)}>,) kw = {'dtype': }, dtype = @_wraps_partial(impl, name) def test_func(*args, **kw): for dtype in dtypes: try: kw[name] = numpy.dtype(dtype).type > impl(*args, **kw) cupy/testing/_helper.py:837: _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ cupy/testing/_helper.py:355: in test_func check_func(cupy_r, numpy_r) cupy/testing/_helper.py:507: in check_func _array.assert_allclose(c, n, rtol1, atol1, err_msg, verbose) _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ actual = array([[[-6.23873319e-04], [-6.52701594e-04], [-5.34185849e-04], [-7.05319864e-04], [-... [-2.50156190e-05], [ 9.55036171e-07], [-1.27298597e-06], [-4.17136789e-06]]], dtype=float32) desired = array([[[ 2.1000416e-16], [-7.7505151e-16], [-6.6172874e-17], [-6.5140899e-16], [ 7.95... [-1.5922411e-17], [ 3.1236140e-17], [-5.3157008e-17], [-1.3124647e-17]]], dtype=float32) rtol = 1e-07, atol = 1e-05, err_msg = '', verbose = True def assert_allclose(actual, desired, rtol=1e-7, atol=0, err_msg='', verbose=True): """Raises an AssertionError if objects are not equal up to desired tolerance. Args: actual(numpy.ndarray or cupy.ndarray): The actual object to check. desired(numpy.ndarray or cupy.ndarray): The desired, expected object. rtol(float): Relative tolerance. atol(float): Absolute tolerance. err_msg(str): The error message to be printed in case of failure. verbose(bool): If ``True``, the conflicting values are appended to the error message. .. seealso:: :func:`numpy.testing.assert_allclose` """ # NOQA > numpy.testing.assert_allclose( cupy.asnumpy(actual), cupy.asnumpy(desired), rtol=rtol, atol=atol, err_msg=err_msg, verbose=verbose) E AssertionError: E Not equal to tolerance rtol=1e-07, atol=1e-05 E E Mismatched elements: 85 / 100 (85%) E Max absolute difference: 0.00084527 E Max relative difference: 1.881553e+13 E x: array([[[-6.238733e-04], E [-6.527016e-04], E [-5.341858e-04],... E y: array([[[ 2.100042e-16], E [-7.750515e-16], E [-6.617287e-17],... cupy/testing/_array.py:24: AssertionError ----------------------------- Captured stdout call ----------------------------- dtype is _ TestZoomOutputSize1.test_zoom_output_size1[_param_22_{grid_mode=False, mode='wrap', order=4, shape=(5, 5, 2), zoom=(2, 2, 0.5)}] _ args = (< parameter: {'grid_mode': False, 'mode': 'wrap', 'order': 4, 'shape': (5, 5, 2), 'zoom': (2, 2, 0.5)}>,) kw = {'dtype': }, dtype = @_wraps_partial(impl, name) def test_func(*args, **kw): for dtype in dtypes: try: kw[name] = numpy.dtype(dtype).type > impl(*args, **kw) cupy/testing/_helper.py:837: _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ cupy/testing/_helper.py:355: in test_func check_func(cupy_r, numpy_r) cupy/testing/_helper.py:507: in check_func _array.assert_allclose(c, n, rtol1, atol1, err_msg, verbose) _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ actual = array([[[-9.26215027e-04], [-7.12910551e-05], [-1.04548770e-03], [ 6.02170185e-04], [ ... [-1.58413204e-05], [-2.26292832e-05], [-1.52809316e-05], [-2.62873400e-05]]], dtype=float32) desired = array([[[-1.91677706e-16], [ 6.02007280e-16], [ 7.39644969e-18], [ 9.42839520e-17], [-... [ 2.49078413e-18], [-6.86374620e-18], [-2.01807717e-17], [-1.36889723e-17]]], dtype=float32) rtol = 1e-07, atol = 1e-05, err_msg = '', verbose = True def assert_allclose(actual, desired, rtol=1e-7, atol=0, err_msg='', verbose=True): """Raises an AssertionError if objects are not equal up to desired tolerance. Args: actual(numpy.ndarray or cupy.ndarray): The actual object to check. desired(numpy.ndarray or cupy.ndarray): The desired, expected object. rtol(float): Relative tolerance. atol(float): Absolute tolerance. err_msg(str): The error message to be printed in case of failure. verbose(bool): If ``True``, the conflicting values are appended to the error message. .. seealso:: :func:`numpy.testing.assert_allclose` """ # NOQA > numpy.testing.assert_allclose( cupy.asnumpy(actual), cupy.asnumpy(desired), rtol=rtol, atol=atol, err_msg=err_msg, verbose=verbose) E AssertionError: E Not equal to tolerance rtol=1e-07, atol=1e-05 E E Mismatched elements: 95 / 100 (95%) E Max absolute difference: 0.00124037 E Max relative difference: 1.6367563e+14 E x: array([[[-9.262150e-04], E [-7.129106e-05], E [-1.045488e-03],... E y: array([[[-1.916777e-16], E [ 6.020073e-16], E [ 7.396450e-18],... cupy/testing/_array.py:24: AssertionError ----------------------------- Captured stdout call ----------------------------- dtype is _ TestZoomOutputSize1.test_zoom_output_size1[_param_23_{grid_mode=False, mode='wrap', order=5, shape=(5, 5, 2), zoom=(2, 2, 0.5)}] _ args = (< parameter: {'grid_mode': False, 'mode': 'wrap', 'order': 5, 'shape': (5, 5, 2), 'zoom': (2, 2, 0.5)}>,) kw = {'dtype': }, dtype = @_wraps_partial(impl, name) def test_func(*args, **kw): for dtype in dtypes: try: kw[name] = numpy.dtype(dtype).type > impl(*args, **kw) cupy/testing/_helper.py:837: _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ cupy/testing/_helper.py:355: in test_func check_func(cupy_r, numpy_r) cupy/testing/_helper.py:507: in check_func _array.assert_allclose(c, n, rtol1, atol1, err_msg, verbose) _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ actual = array([[[-6.23873319e-04], [-6.52701594e-04], [-5.34185849e-04], [-7.05319864e-04], [-... [-2.50156190e-05], [ 9.55036171e-07], [-1.27298597e-06], [-4.17136789e-06]]], dtype=float32) desired = array([[[ 2.1000416e-16], [-7.7505151e-16], [-6.6172874e-17], [-6.5140899e-16], [ 7.95... [-1.5922411e-17], [ 3.1236140e-17], [-5.3157008e-17], [-1.3124647e-17]]], dtype=float32) rtol = 1e-07, atol = 1e-05, err_msg = '', verbose = True def assert_allclose(actual, desired, rtol=1e-7, atol=0, err_msg='', verbose=True): """Raises an AssertionError if objects are not equal up to desired tolerance. Args: actual(numpy.ndarray or cupy.ndarray): The actual object to check. desired(numpy.ndarray or cupy.ndarray): The desired, expected object. rtol(float): Relative tolerance. atol(float): Absolute tolerance. err_msg(str): The error message to be printed in case of failure. verbose(bool): If ``True``, the conflicting values are appended to the error message. .. seealso:: :func:`numpy.testing.assert_allclose` """ # NOQA > numpy.testing.assert_allclose( cupy.asnumpy(actual), cupy.asnumpy(desired), rtol=rtol, atol=atol, err_msg=err_msg, verbose=verbose) E AssertionError: E Not equal to tolerance rtol=1e-07, atol=1e-05 E E Mismatched elements: 85 / 100 (85%) E Max absolute difference: 0.00084527 E Max relative difference: 1.881553e+13 E x: array([[[-6.238733e-04], E [-6.527016e-04], E [-5.341858e-04],... E y: array([[[ 2.100042e-16], E [-7.750515e-16], E [-6.617287e-17],... cupy/testing/_array.py:24: AssertionError ----------------------------- Captured stdout call ----------------------------- dtype is _ TestZoomOutputSize1.test_zoom_output_size1[_param_46_{grid_mode=False, mode='grid-constant', order=4, shape=(5, 5, 2), zoom=(2, 2, 0.5)}] _ args = (< parameter: {'grid_mode': False, 'mode': 'grid-constant', 'order': 4, 'shape': (5, 5, 2), 'zoom': (2, 2, 0.5)}>,) kw = {'dtype': }, dtype = @_wraps_partial(impl, name) def test_func(*args, **kw): for dtype in dtypes: try: kw[name] = numpy.dtype(dtype).type > impl(*args, **kw) cupy/testing/_helper.py:837: _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ cupy/testing/_helper.py:355: in test_func check_func(cupy_r, numpy_r) cupy/testing/_helper.py:507: in check_func _array.assert_allclose(c, n, rtol1, atol1, err_msg, verbose) _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ actual = array([[[-1.6466675e-04], [-3.1765160e-06], [-2.8256103e-04], [ 1.0465207e-04], [-5.02... [-1.4928388e-06], [-5.8970027e-06], [-2.4932410e-06], [-3.5395026e-06]]], dtype=float32) desired = array([[[-8.12850668e-17], [-9.62678620e-17], [-7.64445911e-17], [ 6.90388762e-17], [-... [-1.31306278e-18], [-2.34584224e-18], [-1.16198641e-18], [-3.16820747e-19]]], dtype=float32) rtol = 1e-07, atol = 1e-05, err_msg = '', verbose = True def assert_allclose(actual, desired, rtol=1e-7, atol=0, err_msg='', verbose=True): """Raises an AssertionError if objects are not equal up to desired tolerance. Args: actual(numpy.ndarray or cupy.ndarray): The actual object to check. desired(numpy.ndarray or cupy.ndarray): The desired, expected object. rtol(float): Relative tolerance. atol(float): Absolute tolerance. err_msg(str): The error message to be printed in case of failure. verbose(bool): If ``True``, the conflicting values are appended to the error message. .. seealso:: :func:`numpy.testing.assert_allclose` """ # NOQA > numpy.testing.assert_allclose( cupy.asnumpy(actual), cupy.asnumpy(desired), rtol=rtol, atol=atol, err_msg=err_msg, verbose=verbose) E AssertionError: E Not equal to tolerance rtol=1e-07, atol=1e-05 E E Mismatched elements: 75 / 100 (75%) E Max absolute difference: 0.00034115 E Max relative difference: 2.5044335e+14 E x: array([[[-1.646668e-04], E [-3.176516e-06], E [-2.825610e-04],... E y: array([[[-8.128507e-17], E [-9.626786e-17], E [-7.644459e-17],... cupy/testing/_array.py:24: AssertionError ----------------------------- Captured stdout call ----------------------------- dtype is _ TestZoomOutputSize1.test_zoom_output_size1[_param_47_{grid_mode=False, mode='grid-constant', order=5, shape=(5, 5, 2), zoom=(2, 2, 0.5)}] _ args = (< parameter: {'grid_mode': False, 'mode': 'grid-constant', 'order': 5, 'shape': (5, 5, 2), 'zoom': (2, 2, 0.5)}>,) kw = {'dtype': }, dtype = @_wraps_partial(impl, name) def test_func(*args, **kw): for dtype in dtypes: try: kw[name] = numpy.dtype(dtype).type > impl(*args, **kw) cupy/testing/_helper.py:837: _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ cupy/testing/_helper.py:355: in test_func check_func(cupy_r, numpy_r) cupy/testing/_helper.py:507: in check_func _array.assert_allclose(c, n, rtol1, atol1, err_msg, verbose) _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ actual = array([[[ 1.46576144e-06], [-1.07652804e-05], [ 1.48250510e-05], [-2.88788979e-05], [-... [ 3.59045725e-06], [ 3.90473633e-06], [ 2.02454498e-06], [ 1.08106917e-06]]], dtype=float32) desired = array([[[-5.28892732e-17], [-4.68497761e-17], [-1.70810335e-16], [-1.91986781e-16], [-... [-3.94138989e-18], [-8.65396191e-18], [ 1.07858407e-18], [-2.29544853e-18]]], dtype=float32) rtol = 1e-07, atol = 1e-05, err_msg = '', verbose = True def assert_allclose(actual, desired, rtol=1e-7, atol=0, err_msg='', verbose=True): """Raises an AssertionError if objects are not equal up to desired tolerance. Args: actual(numpy.ndarray or cupy.ndarray): The actual object to check. desired(numpy.ndarray or cupy.ndarray): The desired, expected object. rtol(float): Relative tolerance. atol(float): Absolute tolerance. err_msg(str): The error message to be printed in case of failure. verbose(bool): If ``True``, the conflicting values are appended to the error message. .. seealso:: :func:`numpy.testing.assert_allclose` """ # NOQA > numpy.testing.assert_allclose( cupy.asnumpy(actual), cupy.asnumpy(desired), rtol=rtol, atol=atol, err_msg=err_msg, verbose=verbose) E AssertionError: E Not equal to tolerance rtol=1e-07, atol=1e-05 E E Mismatched elements: 40 / 100 (40%) E Max absolute difference: 0.00011814 E Max relative difference: 9.6834375e+13 E x: array([[[ 1.465761e-06], E [-1.076528e-05], E [ 1.482505e-05],... E y: array([[[-5.288927e-17], E [-4.684978e-17], E [-1.708103e-16],... cupy/testing/_array.py:24: AssertionError ----------------------------- Captured stdout call ----------------------------- dtype is _ TestZoomOutputSize1.test_zoom_output_size1[_param_52_{grid_mode=True, mode='constant', order=4, shape=(5, 5, 2), zoom=(2, 2, 0.5)}] _ args = (< parameter: {'grid_mode': True, 'mode': 'constant', 'order': 4, 'shape': (5, 5, 2), 'zoom': (2, 2, 0.5)}>,) kw = {'dtype': }, dtype = @_wraps_partial(impl, name) def test_func(*args, **kw): for dtype in dtypes: try: kw[name] = numpy.dtype(dtype).type > impl(*args, **kw) cupy/testing/_helper.py:837: _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ cupy/testing/_helper.py:355: in test_func check_func(cupy_r, numpy_r) cupy/testing/_helper.py:507: in check_func _array.assert_allclose(c, n, rtol1, atol1, err_msg, verbose) _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ actual = array([[[ 0. ], [ 0. ], [ 0. ], [ 0. ], [ 0. ... [ 0. ], [ 0. ], [ 0. ], [ 0. ], [ 0. ]]]) desired = array([[[ 0.00000000e+00], [ 0.00000000e+00], [ 0.00000000e+00], [ 0.00000000e+00], [ ...000e+00], [ 0.00000000e+00], [ 0.00000000e+00], [ 0.00000000e+00], [ 0.00000000e+00]]]) rtol = 1e-07, atol = 1e-05, err_msg = '', verbose = True def assert_allclose(actual, desired, rtol=1e-7, atol=0, err_msg='', verbose=True): """Raises an AssertionError if objects are not equal up to desired tolerance. Args: actual(numpy.ndarray or cupy.ndarray): The actual object to check. desired(numpy.ndarray or cupy.ndarray): The desired, expected object. rtol(float): Relative tolerance. atol(float): Absolute tolerance. err_msg(str): The error message to be printed in case of failure. verbose(bool): If ``True``, the conflicting values are appended to the error message. .. seealso:: :func:`numpy.testing.assert_allclose` """ # NOQA > numpy.testing.assert_allclose( cupy.asnumpy(actual), cupy.asnumpy(desired), rtol=rtol, atol=atol, err_msg=err_msg, verbose=verbose) E AssertionError: E Not equal to tolerance rtol=1e-07, atol=1e-05 E E Mismatched elements: 64 / 100 (64%) E Max absolute difference: 2336.52632552 E Max relative difference: 629887.8879561 E x: array([[[ 0. ], E [ 0. ], E [ 0. ],... E y: array([[[ 0.000000e+00], E [ 0.000000e+00], E [ 0.000000e+00],... cupy/testing/_array.py:24: AssertionError ----------------------------- Captured stdout call ----------------------------- dtype is _ TestZoomOutputSize1.test_zoom_output_size1[_param_53_{grid_mode=True, mode='constant', order=5, shape=(5, 5, 2), zoom=(2, 2, 0.5)}] _ args = (< parameter: {'grid_mode': True, 'mode': 'constant', 'order': 5, 'shape': (5, 5, 2), 'zoom': (2, 2, 0.5)}>,) kw = {'dtype': }, dtype = @_wraps_partial(impl, name) def test_func(*args, **kw): for dtype in dtypes: try: kw[name] = numpy.dtype(dtype).type > impl(*args, **kw) cupy/testing/_helper.py:837: _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ cupy/testing/_helper.py:355: in test_func check_func(cupy_r, numpy_r) cupy/testing/_helper.py:507: in check_func _array.assert_allclose(c, n, rtol1, atol1, err_msg, verbose) _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ actual = array([[[ 0.00000000e+00], [ 0.00000000e+00], [ 0.00000000e+00], [ 0.00000000e+00], [ ...000e+00], [ 0.00000000e+00], [ 0.00000000e+00], [ 0.00000000e+00], [ 0.00000000e+00]]]) desired = array([[[ 0.00000000e+00], [ 0.00000000e+00], [ 0.00000000e+00], [ 0.00000000e+00], [ ...000e+00], [ 0.00000000e+00], [ 0.00000000e+00], [ 0.00000000e+00], [ 0.00000000e+00]]]) rtol = 1e-07, atol = 1e-05, err_msg = '', verbose = True def assert_allclose(actual, desired, rtol=1e-7, atol=0, err_msg='', verbose=True): """Raises an AssertionError if objects are not equal up to desired tolerance. Args: actual(numpy.ndarray or cupy.ndarray): The actual object to check. desired(numpy.ndarray or cupy.ndarray): The desired, expected object. rtol(float): Relative tolerance. atol(float): Absolute tolerance. err_msg(str): The error message to be printed in case of failure. verbose(bool): If ``True``, the conflicting values are appended to the error message. .. seealso:: :func:`numpy.testing.assert_allclose` """ # NOQA > numpy.testing.assert_allclose( cupy.asnumpy(actual), cupy.asnumpy(desired), rtol=rtol, atol=atol, err_msg=err_msg, verbose=verbose) E AssertionError: E Not equal to tolerance rtol=1e-07, atol=1e-05 E E Mismatched elements: 64 / 100 (64%) E Max absolute difference: 233.83963068 E Max relative difference: 148867.79248851 E x: array([[[ 0.000000e+00], E [ 0.000000e+00], E [ 0.000000e+00],... E y: array([[[ 0.000000e+00], E [ 0.000000e+00], E [ 0.000000e+00],... cupy/testing/_array.py:24: AssertionError ----------------------------- Captured stdout call ----------------------------- dtype is _ TestZoomOutputSize1.test_zoom_output_size1[_param_64_{grid_mode=True, mode='mirror', order=4, shape=(5, 5, 2), zoom=(2, 2, 0.5)}] _ args = (< parameter: {'grid_mode': True, 'mode': 'mirror', 'order': 4, 'shape': (5, 5, 2), 'zoom': (2, 2, 0.5)}>,) kw = {'dtype': }, dtype = @_wraps_partial(impl, name) def test_func(*args, **kw): for dtype in dtypes: try: kw[name] = numpy.dtype(dtype).type > impl(*args, **kw) cupy/testing/_helper.py:837: _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ cupy/testing/_helper.py:355: in test_func check_func(cupy_r, numpy_r) cupy/testing/_helper.py:507: in check_func _array.assert_allclose(c, n, rtol1, atol1, err_msg, verbose) _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ actual = array([[[ 48.82195577], [ 48.82195577], [ 290.39318135], [ 337.78017528], [ 112.1240... [ -53.24969236], [ -21.88859247], [ 44.15090669], [ 53.29762289], [ 53.29762289]]]) desired = array([[[ 8.54663183e-03], [ 8.54663183e-03], [ 5.30645857e-02], [ 6.15797757e-02], [ ...126e-03], [-5.57818230e-04], [ 2.89445841e-04], [ 8.46143247e-05], [ 8.46143247e-05]]]) rtol = 1e-07, atol = 1e-05, err_msg = '', verbose = True def assert_allclose(actual, desired, rtol=1e-7, atol=0, err_msg='', verbose=True): """Raises an AssertionError if objects are not equal up to desired tolerance. Args: actual(numpy.ndarray or cupy.ndarray): The actual object to check. desired(numpy.ndarray or cupy.ndarray): The desired, expected object. rtol(float): Relative tolerance. atol(float): Absolute tolerance. err_msg(str): The error message to be printed in case of failure. verbose(bool): If ``True``, the conflicting values are appended to the error message. .. seealso:: :func:`numpy.testing.assert_allclose` """ # NOQA > numpy.testing.assert_allclose( cupy.asnumpy(actual), cupy.asnumpy(desired), rtol=rtol, atol=atol, err_msg=err_msg, verbose=verbose) E AssertionError: E Not equal to tolerance rtol=1e-07, atol=1e-05 E E Mismatched elements: 100 / 100 (100%) E Max absolute difference: 2336.52632552 E Max relative difference: 629887.88795612 E x: array([[[ 48.821956], E [ 48.821956], E [ 290.393181],... E y: array([[[ 8.546632e-03], E [ 8.546632e-03], E [ 5.306459e-02],... cupy/testing/_array.py:24: AssertionError ----------------------------- Captured stdout call ----------------------------- dtype is _ TestZoomOutputSize1.test_zoom_output_size1[_param_65_{grid_mode=True, mode='mirror', order=5, shape=(5, 5, 2), zoom=(2, 2, 0.5)}] _ args = (< parameter: {'grid_mode': True, 'mode': 'mirror', 'order': 5, 'shape': (5, 5, 2), 'zoom': (2, 2, 0.5)}>,) kw = {'dtype': }, dtype = @_wraps_partial(impl, name) def test_func(*args, **kw): for dtype in dtypes: try: kw[name] = numpy.dtype(dtype).type > impl(*args, **kw) cupy/testing/_helper.py:837: _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ cupy/testing/_helper.py:355: in test_func check_func(cupy_r, numpy_r) cupy/testing/_helper.py:507: in check_func _array.assert_allclose(c, n, rtol1, atol1, err_msg, verbose) _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ actual = array([[[ 4.71685805e+00], [ 4.71685805e+00], [ 2.76221015e+01], [ 3.32433453e+01], [ ...958e+00], [-1.62349834e+00], [ 1.50933488e+01], [ 1.97919669e+01], [ 1.97919669e+01]]]) desired = array([[[ 8.60075856e-03], [ 8.60075856e-03], [ 5.25533883e-02], [ 6.23449044e-02], [ ...600e-03], [-8.00397905e-04], [ 4.16976616e-04], [ 1.32949066e-04], [ 1.32949066e-04]]]) rtol = 1e-07, atol = 1e-05, err_msg = '', verbose = True def assert_allclose(actual, desired, rtol=1e-7, atol=0, err_msg='', verbose=True): """Raises an AssertionError if objects are not equal up to desired tolerance. Args: actual(numpy.ndarray or cupy.ndarray): The actual object to check. desired(numpy.ndarray or cupy.ndarray): The desired, expected object. rtol(float): Relative tolerance. atol(float): Absolute tolerance. err_msg(str): The error message to be printed in case of failure. verbose(bool): If ``True``, the conflicting values are appended to the error message. .. seealso:: :func:`numpy.testing.assert_allclose` """ # NOQA > numpy.testing.assert_allclose( cupy.asnumpy(actual), cupy.asnumpy(desired), rtol=rtol, atol=atol, err_msg=err_msg, verbose=verbose) E AssertionError: E Not equal to tolerance rtol=1e-07, atol=1e-05 E E Mismatched elements: 100 / 100 (100%) E Max absolute difference: 233.83963068 E Max relative difference: 148867.79248854 E x: array([[[ 4.716858e+00], E [ 4.716858e+00], E [ 2.762210e+01],... E y: array([[[ 8.600759e-03], E [ 8.600759e-03], E [ 5.255339e-02],... cupy/testing/_array.py:24: AssertionError ----------------------------- Captured stdout call ----------------------------- dtype is _ TestZoomOutputSize1.test_zoom_output_size1[_param_70_{grid_mode=True, mode='wrap', order=4, shape=(5, 5, 2), zoom=(2, 2, 0.5)}] _ args = (< parameter: {'grid_mode': True, 'mode': 'wrap', 'order': 4, 'shape': (5, 5, 2), 'zoom': (2, 2, 0.5)}>,) kw = {'dtype': }, dtype = @_wraps_partial(impl, name) def test_func(*args, **kw): for dtype in dtypes: try: kw[name] = numpy.dtype(dtype).type > impl(*args, **kw) cupy/testing/_helper.py:837: _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ cupy/testing/_helper.py:355: in test_func check_func(cupy_r, numpy_r) cupy/testing/_helper.py:507: in check_func _array.assert_allclose(c, n, rtol1, atol1, err_msg, verbose) _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ actual = array([[[ 53.29762289], [ 51.0107262 ], [ 303.41199632], [ 352.92342892], [ 117.1507... [ -50.96485228], [ -20.9493958 ], [ 42.25647769], [ 51.0107262 ], [ 48.82195577]]]) desired = array([[[ 8.46143247e-05], [ 8.50392545e-04], [ 5.27994291e-03], [ 6.12720698e-03], [ ...636e-02], [-5.60619571e-03], [ 2.90899427e-03], [ 8.50392545e-04], [ 8.54663183e-03]]]) rtol = 1e-07, atol = 1e-05, err_msg = '', verbose = True def assert_allclose(actual, desired, rtol=1e-7, atol=0, err_msg='', verbose=True): """Raises an AssertionError if objects are not equal up to desired tolerance. Args: actual(numpy.ndarray or cupy.ndarray): The actual object to check. desired(numpy.ndarray or cupy.ndarray): The desired, expected object. rtol(float): Relative tolerance. atol(float): Absolute tolerance. err_msg(str): The error message to be printed in case of failure. verbose(bool): If ``True``, the conflicting values are appended to the error message. .. seealso:: :func:`numpy.testing.assert_allclose` """ # NOQA > numpy.testing.assert_allclose( cupy.asnumpy(actual), cupy.asnumpy(desired), rtol=rtol, atol=atol, err_msg=err_msg, verbose=verbose) E AssertionError: E Not equal to tolerance rtol=1e-07, atol=1e-05 E E Mismatched elements: 100 / 100 (100%) E Max absolute difference: 2336.52632552 E Max relative difference: 629887.8879561 E x: array([[[ 53.297623], E [ 51.010726], E [ 303.411996],... E y: array([[[ 8.461432e-05], E [ 8.503925e-04], E [ 5.279943e-03],... cupy/testing/_array.py:24: AssertionError ----------------------------- Captured stdout call ----------------------------- dtype is _ TestZoomOutputSize1.test_zoom_output_size1[_param_71_{grid_mode=True, mode='wrap', order=5, shape=(5, 5, 2), zoom=(2, 2, 0.5)}] _ args = (< parameter: {'grid_mode': True, 'mode': 'wrap', 'order': 5, 'shape': (5, 5, 2), 'zoom': (2, 2, 0.5)}>,) kw = {'dtype': }, dtype = @_wraps_partial(impl, name) def test_func(*args, **kw): for dtype in dtypes: try: kw[name] = numpy.dtype(dtype).type > impl(*args, **kw) cupy/testing/_helper.py:837: _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ cupy/testing/_helper.py:355: in test_func check_func(cupy_r, numpy_r) cupy/testing/_helper.py:507: in check_func _array.assert_allclose(c, n, rtol1, atol1, err_msg, verbose) _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ actual = array([[[ 1.97919669e+01], [ 9.66208561e+00], [ 5.65815436e+01], [ 6.80961870e+01], [ ...577e+00], [-7.92562963e-01], [ 7.36830399e+00], [ 9.66208561e+00], [ 4.71685805e+00]]]) desired = array([[[ 1.32949066e-04], [ 1.06932821e-03], [ 6.53393769e-03], [ 7.75131221e-03], [ ...718e-02], [-6.43771396e-03], [ 3.35380211e-03], [ 1.06932821e-03], [ 8.60075856e-03]]]) rtol = 1e-07, atol = 1e-05, err_msg = '', verbose = True def assert_allclose(actual, desired, rtol=1e-7, atol=0, err_msg='', verbose=True): """Raises an AssertionError if objects are not equal up to desired tolerance. Args: actual(numpy.ndarray or cupy.ndarray): The actual object to check. desired(numpy.ndarray or cupy.ndarray): The desired, expected object. rtol(float): Relative tolerance. atol(float): Absolute tolerance. err_msg(str): The error message to be printed in case of failure. verbose(bool): If ``True``, the conflicting values are appended to the error message. .. seealso:: :func:`numpy.testing.assert_allclose` """ # NOQA > numpy.testing.assert_allclose( cupy.asnumpy(actual), cupy.asnumpy(desired), rtol=rtol, atol=atol, err_msg=err_msg, verbose=verbose) E AssertionError: E Not equal to tolerance rtol=1e-07, atol=1e-05 E E Mismatched elements: 100 / 100 (100%) E Max absolute difference: 233.83963068 E Max relative difference: 148867.79248851 E x: array([[[ 1.979197e+01], E [ 9.662086e+00], E [ 5.658154e+01],... E y: array([[[ 1.329491e-04], E [ 1.069328e-03], E [ 6.533938e-03],... cupy/testing/_array.py:24: AssertionError ----------------------------- Captured stdout call ----------------------------- dtype is _ TestZoomOutputSize1.test_zoom_output_size1[_param_94_{grid_mode=True, mode='grid-constant', order=4, shape=(5, 5, 2), zoom=(2, 2, 0.5)}] _ args = (< parameter: {'grid_mode': True, 'mode': 'grid-constant', 'order': 4, 'shape': (5, 5, 2), 'zoom': (2, 2, 0.5)}>,) kw = {'dtype': }, dtype = @_wraps_partial(impl, name) def test_func(*args, **kw): for dtype in dtypes: try: kw[name] = numpy.dtype(dtype).type > impl(*args, **kw) cupy/testing/_helper.py:837: _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ cupy/testing/_helper.py:355: in test_func check_func(cupy_r, numpy_r) cupy/testing/_helper.py:507: in check_func _array.assert_allclose(c, n, rtol1, atol1, err_msg, verbose) _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ actual = array([[[ 5.63055219e+01], [-1.22528405e+02], [-3.83476709e+02], [-3.83476709e+02], [-...021e+00], [ 4.40535983e+00], [-2.69124172e+00], [-1.59183812e+00], [ 9.72456581e-01]]]) desired = array([[[ 1.27949283e-02], [-2.43606830e-02], [-7.91692940e-02], [-7.91692940e-02], [-...333e-03], [ 9.13158865e-04], [-5.54269262e-04], [-3.27391850e-04], [ 1.99955080e-04]]]) rtol = 1e-07, atol = 1e-05, err_msg = '', verbose = True def assert_allclose(actual, desired, rtol=1e-7, atol=0, err_msg='', verbose=True): """Raises an AssertionError if objects are not equal up to desired tolerance. Args: actual(numpy.ndarray or cupy.ndarray): The actual object to check. desired(numpy.ndarray or cupy.ndarray): The desired, expected object. rtol(float): Relative tolerance. atol(float): Absolute tolerance. err_msg(str): The error message to be printed in case of failure. verbose(bool): If ``True``, the conflicting values are appended to the error message. .. seealso:: :func:`numpy.testing.assert_allclose` """ # NOQA > numpy.testing.assert_allclose( cupy.asnumpy(actual), cupy.asnumpy(desired), rtol=rtol, atol=atol, err_msg=err_msg, verbose=verbose) E AssertionError: E Not equal to tolerance rtol=1e-07, atol=1e-05 E E Mismatched elements: 100 / 100 (100%) E Max absolute difference: 2611.23242921 E Max relative difference: 5747.85704324 E x: array([[[ 5.630552e+01], E [-1.225284e+02], E [-3.834767e+02],... E y: array([[[ 1.279493e-02], E [-2.436068e-02], E [-7.916929e-02],... cupy/testing/_array.py:24: AssertionError ----------------------------- Captured stdout call ----------------------------- dtype is _ TestZoomOutputSize1.test_zoom_output_size1[_param_95_{grid_mode=True, mode='grid-constant', order=5, shape=(5, 5, 2), zoom=(2, 2, 0.5)}] _ args = (< parameter: {'grid_mode': True, 'mode': 'grid-constant', 'order': 5, 'shape': (5, 5, 2), 'zoom': (2, 2, 0.5)}>,) kw = {'dtype': }, dtype = @_wraps_partial(impl, name) def test_func(*args, **kw): for dtype in dtypes: try: kw[name] = numpy.dtype(dtype).type > impl(*args, **kw) cupy/testing/_helper.py:837: _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ cupy/testing/_helper.py:355: in test_func check_func(cupy_r, numpy_r) cupy/testing/_helper.py:507: in check_func _array.assert_allclose(c, n, rtol1, atol1, err_msg, verbose) _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ actual = array([[[ 4.16965077e+00], [-1.18699209e+01], [-3.36177788e+01], [-3.36177788e+01], [-...671e-01], [ 6.07624141e-01], [-3.96545295e-01], [-2.61764771e-01], [ 1.70742628e-01]]]) desired = array([[[ 1.44354197e-02], [-2.67893098e-02], [-8.45564959e-02], [-8.45564959e-02], [-...699e-03], [ 1.49724992e-03], [-9.55056276e-04], [-6.25982935e-04], [ 4.07383783e-04]]]) rtol = 1e-07, atol = 1e-05, err_msg = '', verbose = True def assert_allclose(actual, desired, rtol=1e-7, atol=0, err_msg='', verbose=True): """Raises an AssertionError if objects are not equal up to desired tolerance. Args: actual(numpy.ndarray or cupy.ndarray): The actual object to check. desired(numpy.ndarray or cupy.ndarray): The desired, expected object. rtol(float): Relative tolerance. atol(float): Absolute tolerance. err_msg(str): The error message to be printed in case of failure. verbose(bool): If ``True``, the conflicting values are appended to the error message. .. seealso:: :func:`numpy.testing.assert_allclose` """ # NOQA > numpy.testing.assert_allclose( cupy.asnumpy(actual), cupy.asnumpy(desired), rtol=rtol, atol=atol, err_msg=err_msg, verbose=verbose) E AssertionError: E Not equal to tolerance rtol=1e-07, atol=1e-05 E E Mismatched elements: 100 / 100 (100%) E Max absolute difference: 270.54779997 E Max relative difference: 678.67662434 E x: array([[[ 4.169651e+00], E [-1.186992e+01], E [-3.361778e+01],... E y: array([[[ 1.443542e-02], E [-2.678931e-02], E [-8.455650e-02],... cupy/testing/_array.py:24: AssertionError ----------------------------- Captured stdout call ----------------------------- dtype is _ TestFFTConvolve.test_fftconvolve[_param_44_{mode='valid', size1=(3, 4, 10), size2=3}] _ args = (< parameter: {'mode': 'valid', 'size1': (3, 4, 10), 'size2': 3}>,) kw = {'dtype': }, dtype = @_wraps_partial(impl, name) def test_func(*args, **kw): for dtype in dtypes: try: kw[name] = numpy.dtype(dtype).type > impl(*args, **kw) cupy/testing/_helper.py:837: _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ cupy/testing/_helper.py:355: in test_func check_func(cupy_r, numpy_r) cupy/testing/_helper.py:507: in check_func _array.assert_allclose(c, n, rtol1, atol1, err_msg, verbose) _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ actual = array([[[ -336., 2067., 552., -1533., 1806., -1719., 2414., -199.], [ -563., 1198., 105., -2202., -81., 1175., 1246., 1177.]]]) desired = array([[[648., 719., 544., 505., 439., 576., 586., 661.], [710., 755., 643., 555., 464., 581., 669., 710.]]]) rtol = 1e-08, atol = 1e-08, err_msg = '', verbose = True def assert_allclose(actual, desired, rtol=1e-7, atol=0, err_msg='', verbose=True): """Raises an AssertionError if objects are not equal up to desired tolerance. Args: actual(numpy.ndarray or cupy.ndarray): The actual object to check. desired(numpy.ndarray or cupy.ndarray): The desired, expected object. rtol(float): Relative tolerance. atol(float): Absolute tolerance. err_msg(str): The error message to be printed in case of failure. verbose(bool): If ``True``, the conflicting values are appended to the error message. .. seealso:: :func:`numpy.testing.assert_allclose` """ # NOQA > numpy.testing.assert_allclose( cupy.asnumpy(actual), cupy.asnumpy(desired), rtol=rtol, atol=atol, err_msg=err_msg, verbose=verbose) E AssertionError: E Not equal to tolerance rtol=1e-08, atol=1e-08 E E Mismatched elements: 16 / 16 (100%) E Max absolute difference: 2757. E Max relative difference: 4.96756757 E x: array([[[ -336., 2067., 552., -1533., 1806., -1719., 2414., -199.], E [ -563., 1198., 105., -2202., -81., 1175., 1246., 1177.]]]) E y: array([[[648., 719., 544., 505., 439., 576., 586., 661.], E [710., 755., 643., 555., 464., 581., 669., 710.]]]) cupy/testing/_array.py:24: AssertionError ----------------------------- Captured stdout call ----------------------------- dtype is _____________ TestLOBPCG.test_small_generate_input_for_mikota_pair _____________ self = xp = sp = @testing.numpy_cupy_allclose(rtol=1e-5, atol=1e-5, sp_name='sp', contiguous_check=False) def test_small_generate_input_for_mikota_pair(self, xp, sp): A, B = self._generate_input_for_mikota_pair(10, xp) n = A.shape[0] X = self._generate_random_initial_ortho_eigvec(n, 10, xp) > eigvals, eigvecs = sp.linalg.lobpcg(A, X, B=B, tol=1e-5, maxiter=30, largest=False) tests/cupyx_tests/scipy_tests/sparse_tests/test_linalg.py:964: _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ cupyx/scipy/sparse/linalg/_lobpcg.py:298: in lobpcg A_dense = A(cupy.eye(n, dtype=A.dtype)) cupyx/scipy/sparse/linalg/_interface.py:170: in __call__ return self*x cupyx/scipy/sparse/linalg/_interface.py:173: in __mul__ return self.dot(x) cupyx/scipy/sparse/linalg/_interface.py:186: in dot return self.matmat(x) cupyx/scipy/sparse/linalg/_interface.py:142: in matmat Y = self._matmat(X) cupyx/scipy/sparse/linalg/_interface.py:489: in _matmat return self.A.dot(X) cupy/_core/core.pyx:1607: in cupy._core.core.ndarray.dot return _linalg.dot(self, b, out) cupy/_core/_routines_linalg.pyx:427: in cupy._core._routines_linalg.dot return tensordot_core(a, b, out, n, m, k, ret_shape) cupy/_core/_routines_linalg.pyx:501: in cupy._core._routines_linalg.tensordot_core return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) cupy/_core/_routines_linalg.pyx:327: in cupy._core._routines_linalg._integral_tensordot_core kern = _tensordot_core_int_kernel(config, dtype) cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/_routines_linalg.pyx:302: in cupy._core._routines_linalg._tensordot_core_int_kernel ker = mod.get_function( cupy/_core/raw.pyx:470: in cupy._core.raw.RawModule.get_function mangled_name = self.module.mapping.get(name) cupy/_core/raw.pyx:394: in cupy._core.raw.RawModule.module.__get__ return self._module() cupy/_core/raw.pyx:402: in cupy._core.raw.RawModule._module mod = _get_raw_module( cupy/_util.pyx:59: in cupy._util.memoize.decorator.ret result = f(*args, **kwargs) cupy/_core/raw.pyx:547: in cupy._core.raw._get_raw_module mod = cupy._core.core.compile_with_cache( cupy/_core/core.pyx:2015: in cupy._core.core.compile_with_cache cpdef function.Module compile_with_cache( cupy/_core/core.pyx:2078: in cupy._core.core.compile_with_cache return cuda.compile_with_cache( cupy/cuda/compiler.py:453: in compile_with_cache return _compile_with_cache_hip( cupy/cuda/compiler.py:853: in _compile_with_cache_hip binary, mapping = compile_using_nvrtc( cupy/cuda/compiler.py:295: in compile_using_nvrtc return _compile(source, options, cu_path, cupy/cuda/compiler.py:279: in _compile compiled_obj, mapping = prog.compile(options, log_stream) _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ self = options = ('--std=c++11', '-I/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/_core/include', '-I/opt/rocm/include', '-fcuda-flush-denormals-to-zero') log_stream = None def compile(self, options=(), log_stream=None): try: if self.name_expressions: for ker in self.name_expressions: nvrtc.addAddNameExpression(self.ptr, ker) nvrtc.compileProgram(self.ptr, options) mapping = None if self.name_expressions: mapping = {} for ker in self.name_expressions: mapping[ker] = nvrtc.getLoweredName(self.ptr, ker) if log_stream is not None: log_stream.write(nvrtc.getProgramLog(self.ptr)) # This is to ensure backwards compatibility with nvrtc if self.method == 'cubin': return nvrtc.getCUBIN(self.ptr), mapping elif self.method == 'ptx': return nvrtc.getPTX(self.ptr), mapping else: raise RuntimeError('Unknown NVRTC compile method') except nvrtc.NVRTCError: log = nvrtc.getProgramLog(self.ptr) > raise CompileException(log, self.src, self.name, options, 'nvrtc' if not runtime.is_hip else 'hiprtc') E AssertionError: Only cupy raises error E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 650, in compile E nvrtc.compileProgram(self.ptr, options) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 136, in cupy_backends.cuda.libs.nvrtc.compileProgram E cpdef compileProgram(intptr_t prog, options): E File "cupy_backends/cuda/libs/nvrtc.pyx", line 148, in cupy_backends.cuda.libs.nvrtc.compileProgram E check_status(status) E File "cupy_backends/cuda/libs/nvrtc.pyx", line 67, in cupy_backends.cuda.libs.nvrtc.check_status E raise NVRTCError(status) E cupy_backends.cuda.libs.nvrtc.NVRTCError: HIPRTC_ERROR_COMPILATION (6) E E During handling of the above exception, another exception occurred: E E Traceback (most recent call last): E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/testing/_helper.py", line 47, in _call_func E result = impl(*args, **kw) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/tests/cupyx_tests/scipy_tests/sparse_tests/test_linalg.py", line 964, in test_small_generate_input_for_mikota_pair E eigvals, eigvecs = sp.linalg.lobpcg(A, E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupyx/scipy/sparse/linalg/_lobpcg.py", line 298, in lobpcg E A_dense = A(cupy.eye(n, dtype=A.dtype)) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupyx/scipy/sparse/linalg/_interface.py", line 170, in __call__ E return self*x E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupyx/scipy/sparse/linalg/_interface.py", line 173, in __mul__ E return self.dot(x) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupyx/scipy/sparse/linalg/_interface.py", line 186, in dot E return self.matmat(x) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupyx/scipy/sparse/linalg/_interface.py", line 142, in matmat E Y = self._matmat(X) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupyx/scipy/sparse/linalg/_interface.py", line 489, in _matmat E return self.A.dot(X) E File "cupy/_core/core.pyx", line 1607, in cupy._core.core.ndarray.dot E return _linalg.dot(self, b, out) E File "cupy/_core/_routines_linalg.pyx", line 427, in cupy._core._routines_linalg.dot E return tensordot_core(a, b, out, n, m, k, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 501, in cupy._core._routines_linalg.tensordot_core E return _integral_tensordot_core(b, a, out, m, n, k, dtype, ret_shape) E File "cupy/_core/_routines_linalg.pyx", line 327, in cupy._core._routines_linalg._integral_tensordot_core E kern = _tensordot_core_int_kernel(config, dtype) E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/_routines_linalg.pyx", line 302, in cupy._core._routines_linalg._tensordot_core_int_kernel E ker = mod.get_function( E File "cupy/_core/raw.pyx", line 470, in cupy._core.raw.RawModule.get_function E mangled_name = self.module.mapping.get(name) E File "cupy/_core/raw.pyx", line 394, in cupy._core.raw.RawModule.module.__get__ E return self._module() E File "cupy/_core/raw.pyx", line 402, in cupy._core.raw.RawModule._module E mod = _get_raw_module( E File "cupy/_util.pyx", line 59, in cupy._util.memoize.decorator.ret E result = f(*args, **kwargs) E File "cupy/_core/raw.pyx", line 547, in cupy._core.raw._get_raw_module E mod = cupy._core.core.compile_with_cache( E File "cupy/_core/core.pyx", line 2015, in cupy._core.core.compile_with_cache E cpdef function.Module compile_with_cache( E File "cupy/_core/core.pyx", line 2078, in cupy._core.core.compile_with_cache E return cuda.compile_with_cache( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 453, in compile_with_cache E return _compile_with_cache_hip( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 853, in _compile_with_cache_hip E binary, mapping = compile_using_nvrtc( E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 295, in compile_using_nvrtc E return _compile(source, options, cu_path, E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 279, in _compile E compiled_obj, mapping = prog.compile(options, log_stream) E File "/global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupy/cuda/compiler.py", line 667, in compile E raise CompileException(log, self.src, self.name, options, E cupy.cuda.compiler.CompileException: /tmp/comgr-be9423/input/CompileSource:203:27: error: redeclaration of '__hiprtc_24' with a different type: 'void (*const)(int, int, int, const signed char *, const signed char *, signed char *)' vs 'void (*const)(int, int, int, const bool *, const bool *, bool *)' E extern "C" constexpr auto __hiprtc_24 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-be9423/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_24 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-be9423/input/CompileSource:204:27: error: redefinition of '__hiprtc_24' E extern "C" constexpr auto __hiprtc_24 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-be9423/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_24 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-be9423/input/CompileSource:205:27: error: redefinition of '__hiprtc_24' E extern "C" constexpr auto __hiprtc_24 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-be9423/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_24 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-be9423/input/CompileSource:206:27: error: redefinition of '__hiprtc_24' E extern "C" constexpr auto __hiprtc_24 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-be9423/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_24 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-be9423/input/CompileSource:207:27: error: redefinition of '__hiprtc_24' E extern "C" constexpr auto __hiprtc_24 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-be9423/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_24 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-be9423/input/CompileSource:208:27: error: redefinition of '__hiprtc_24' E extern "C" constexpr auto __hiprtc_24 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-be9423/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_24 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-be9423/input/CompileSource:209:27: error: redefinition of '__hiprtc_24' E extern "C" constexpr auto __hiprtc_24 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-be9423/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_24 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-be9423/input/CompileSource:210:27: error: redefinition of '__hiprtc_24' E extern "C" constexpr auto __hiprtc_24 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-be9423/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_24 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-be9423/input/CompileSource:211:27: error: redefinition of '__hiprtc_24' E extern "C" constexpr auto __hiprtc_24 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-be9423/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_24 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-be9423/input/CompileSource:212:27: error: redefinition of '__hiprtc_24' E extern "C" constexpr auto __hiprtc_24 = _tensordot_core_int_kernel; E ^ E /tmp/comgr-be9423/input/CompileSource:202:27: note: previous definition is here E extern "C" constexpr auto __hiprtc_24 = _tensordot_core_int_kernel; E ^ E 10 errors generated when compiling for gfx908. E Error: Failed to compile opencl source (from CL or HIP source to LLVM IR). cupy/cuda/compiler.py:667: AssertionError __________________________ TestLOBPCG.test_maxit_None __________________________ self = def test_maxit_None(self): """Check lobpcg if maxit=None runs 20 iterations (the default) by checking the size of the iteration history output, which should be the number of iterations plus 2 (initial and final values). """ n = 50 m = 4 vals = -cupy.arange(1, n + 1) A = sparse.diags([vals], [0], (n, n)) A = A.astype(cupy.float32) X = testing.shaped_random((n, m), xp=cupy, seed=1566950023) _, _, l_h = sparse.linalg.lobpcg(A, X, tol=1e-8, maxiter=None, retLambdaHistory=True) > assert len(l_h) == 22 E assert 20 == 22 E + where 20 = len([array([-18.735699, -23.57251 , -27.328299, -33.428257], dtype=float32), array([ -7.7822046, -10.007652 , -13.157776 ,..., -4.689918 , -5.93529 ], dtype=float32), array([-1.1894418, -3.1553519, -3.572582 , -4.400922 ], dtype=float32), ...]) tests/cupyx_tests/scipy_tests/sparse_tests/test_linalg.py:1186: AssertionError =============================== warnings summary =============================== tests/cupy_tests/core_tests/fusion_tests/test_array.py: 1 warning tests/cupy_tests/math_tests/test_sumprod.py: 19 warnings /home/kmaeh/.pyenv/versions/rocm-ci/lib/python3.8/site-packages/numpy/core/_methods.py:51: RuntimeWarning: overflow encountered in reduce return umr_prod(a, axis, dtype, out, keepdims, initial, where) tests/cupy_tests/core_tests/fusion_tests/test_routines.py::TestFusionNumericalReduction_param_1_{func='prod'}::test_reduction tests/cupy_tests/math_tests/test_sumprod.py::TestNansumNanprodLong::test_nansum_all[_param_11_{axis=0, func='nanprod', keepdims=True, shape=(20, 30, 40), transpose_axes=False}] tests/cupy_tests/math_tests/test_sumprod.py::TestNansumNanprodLong::test_nansum_all[_param_15_{axis=0, func='nanprod', keepdims=False, shape=(20, 30, 40), transpose_axes=False}] tests/cupy_tests/math_tests/test_sumprod.py::TestNansumNanprodLong::test_nansum_axis_transposed[_param_11_{axis=0, func='nanprod', keepdims=True, shape=(20, 30, 40), transpose_axes=False}] tests/cupy_tests/math_tests/test_sumprod.py::TestNansumNanprodLong::test_nansum_axis_transposed[_param_15_{axis=0, func='nanprod', keepdims=False, shape=(20, 30, 40), transpose_axes=False}] /home/kmaeh/.pyenv/versions/rocm-ci/lib/python3.8/site-packages/numpy/core/fromnumeric.py:87: RuntimeWarning: overflow encountered in reduce return ufunc.reduce(obj, axis, dtype, out, **passkwargs) tests/cupy_tests/lib_tests/test_polynomial.py::TestPolyval::test_polyval[_param_0_{type_l='poly1d', type_r='ndarray'}] tests/cupy_tests/lib_tests/test_polynomial.py::TestPolyval::test_polyval[_param_3_{type_l='ndarray', type_r='ndarray'}] tests/cupy_tests/lib_tests/test_polynomial.py::TestPolyvalDtypesCombination::test_polyval_diff_types_array_array /home/kmaeh/.pyenv/versions/rocm-ci/lib/python3.8/site-packages/numpy/lib/polynomial.py:734: RuntimeWarning: overflow encountered in multiply y = y * x + p[i] tests/cupy_tests/lib_tests/test_polynomial.py::TestPolyval::test_polyval[_param_1_{type_l='poly1d', type_r='numpy_scalar'}] tests/cupy_tests/lib_tests/test_polynomial.py::TestPolyval::test_polyval[_param_4_{type_l='ndarray', type_r='numpy_scalar'}] tests/cupy_tests/lib_tests/test_polynomial.py::TestPolyvalDtypesCombination::test_polyval_diff_types_array_scalar /home/kmaeh/.pyenv/versions/rocm-ci/lib/python3.8/site-packages/numpy/lib/polynomial.py:734: RuntimeWarning: overflow encountered in half_scalars y = y * x + p[i] tests/cupy_tests/lib_tests/test_polynomial.py::TestPolyval::test_polyval[_param_1_{type_l='poly1d', type_r='numpy_scalar'}] tests/cupy_tests/lib_tests/test_polynomial.py::TestPolyval::test_polyval[_param_4_{type_l='ndarray', type_r='numpy_scalar'}] tests/cupy_tests/lib_tests/test_polynomial.py::TestPolyvalDtypesCombination::test_polyval_diff_types_array_scalar /home/kmaeh/.pyenv/versions/rocm-ci/lib/python3.8/site-packages/numpy/lib/polynomial.py:734: RuntimeWarning: overflow encountered in byte_scalars y = y * x + p[i] tests/cupy_tests/lib_tests/test_polynomial.py::TestPolyval::test_polyval[_param_1_{type_l='poly1d', type_r='numpy_scalar'}] tests/cupy_tests/lib_tests/test_polynomial.py::TestPolyval::test_polyval[_param_4_{type_l='ndarray', type_r='numpy_scalar'}] tests/cupy_tests/lib_tests/test_polynomial.py::TestPolyvalDtypesCombination::test_polyval_diff_types_array_scalar /home/kmaeh/.pyenv/versions/rocm-ci/lib/python3.8/site-packages/numpy/lib/polynomial.py:734: RuntimeWarning: overflow encountered in short_scalars y = y * x + p[i] tests/cupy_tests/lib_tests/test_polynomial.py::TestPolyval::test_polyval[_param_1_{type_l='poly1d', type_r='numpy_scalar'}] tests/cupy_tests/lib_tests/test_polynomial.py::TestPolyval::test_polyval[_param_4_{type_l='ndarray', type_r='numpy_scalar'}] tests/cupy_tests/lib_tests/test_polynomial.py::TestPolyvalDtypesCombination::test_polyval_diff_types_array_scalar /home/kmaeh/.pyenv/versions/rocm-ci/lib/python3.8/site-packages/numpy/lib/polynomial.py:734: RuntimeWarning: overflow encountered in ubyte_scalars y = y * x + p[i] tests/cupy_tests/lib_tests/test_polynomial.py::TestPolyval::test_polyval[_param_1_{type_l='poly1d', type_r='numpy_scalar'}] tests/cupy_tests/lib_tests/test_polynomial.py::TestPolyval::test_polyval[_param_4_{type_l='ndarray', type_r='numpy_scalar'}] tests/cupy_tests/lib_tests/test_polynomial.py::TestPolyvalDtypesCombination::test_polyval_diff_types_array_scalar /home/kmaeh/.pyenv/versions/rocm-ci/lib/python3.8/site-packages/numpy/lib/polynomial.py:734: RuntimeWarning: overflow encountered in ushort_scalars y = y * x + p[i] tests/cupy_tests/linalg_tests/test_norms.py: 14 warnings /home/kmaeh/.pyenv/versions/rocm-ci/lib/python3.8/site-packages/numpy/linalg/linalg.py:2568: RuntimeWarning: divide by zero encountered in power absx **= ord tests/cupy_tests/linalg_tests/test_norms.py: 14 warnings /home/kmaeh/.pyenv/versions/rocm-ci/lib/python3.8/site-packages/numpy/linalg/linalg.py:2568: RuntimeWarning: divide by zero encountered in reciprocal absx **= ord tests/cupy_tests/math_tests/test_misc.py::TestMisc::test_nan_to_num_inf tests/cupy_tests/math_tests/test_misc.py::TestMisc::test_nan_to_num_inf_nan /home/kmaeh/.pyenv/versions/rocm-ci/lib/python3.8/site-packages/numpy/core/numeric.py:2276: RuntimeWarning: overflow encountered in absolute return less_equal(abs(x-y), atol + rtol * abs(y)) tests/cupy_tests/math_tests/test_sumprod.py::TestSumprod::test_sum_all2 tests/cupy_tests/math_tests/test_sumprod.py::TestSumprod::test_sum_all_transposed2 tests/cupy_tests/math_tests/test_sumprod.py::TestSumprod::test_sum_axis_transposed2 tests/cupy_tests/math_tests/test_sumprod.py::TestSumprod::test_sum_axes2 tests/cupy_tests/math_tests/test_sumprod.py::TestSumprod::test_sum_axes4 tests/cupyx_tests/scipy_tests/ndimage_tests/test_measurements.py::TestStats::test_only_input[_param_2_{op='variance'}] tests/cupyx_tests/scipy_tests/ndimage_tests/test_measurements.py::TestStats::test_only_input[_param_3_{op='standard_deviation'}] /home/kmaeh/.pyenv/versions/rocm-ci/lib/python3.8/site-packages/numpy/core/_methods.py:47: RuntimeWarning: overflow encountered in reduce return umr_sum(a, axis, dtype, out, keepdims, initial, where) tests/cupy_tests/math_tests/test_sumprod.py: 15 warnings /home/kmaeh/.pyenv/versions/rocm-ci/lib/python3.8/site-packages/numpy/core/_methods.py:51: RuntimeWarning: invalid value encountered in reduce return umr_prod(a, axis, dtype, out, keepdims, initial, where) tests/cupy_tests/math_tests/test_sumprod.py::TestNansumNanprodLong::test_nansum_all[_param_11_{axis=0, func='nanprod', keepdims=True, shape=(20, 30, 40), transpose_axes=False}] tests/cupy_tests/math_tests/test_sumprod.py::TestNansumNanprodLong::test_nansum_all[_param_15_{axis=0, func='nanprod', keepdims=False, shape=(20, 30, 40), transpose_axes=False}] tests/cupy_tests/math_tests/test_sumprod.py::TestNansumNanprodLong::test_nansum_axis_transposed[_param_11_{axis=0, func='nanprod', keepdims=True, shape=(20, 30, 40), transpose_axes=False}] tests/cupy_tests/math_tests/test_sumprod.py::TestNansumNanprodLong::test_nansum_axis_transposed[_param_15_{axis=0, func='nanprod', keepdims=False, shape=(20, 30, 40), transpose_axes=False}] /home/kmaeh/.pyenv/versions/rocm-ci/lib/python3.8/site-packages/numpy/core/fromnumeric.py:87: RuntimeWarning: invalid value encountered in reduce return ufunc.reduce(obj, axis, dtype, out, **passkwargs) tests/cupy_tests/math_tests/test_sumprod.py: 32 warnings /global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/tests/cupy_tests/math_tests/test_sumprod.py:803: RuntimeWarning: invalid value encountered in true_divide a = a / a tests/cupy_tests/statistics_tests/test_meanvar.py: 32 warnings /global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/tests/cupy_tests/statistics_tests/test_meanvar.py:113: RuntimeWarning: invalid value encountered in true_divide a = a / a tests/cupy_tests/statistics_tests/test_meanvar.py::TestNanMedian::test_nanmedian[_param_4_{axis=0, keepdims=True, overwrite_input=True, shape=(3, 4, 5)}] tests/cupy_tests/statistics_tests/test_meanvar.py::TestNanMedian::test_nanmedian[_param_5_{axis=0, keepdims=True, overwrite_input=False, shape=(3, 4, 5)}] tests/cupy_tests/statistics_tests/test_meanvar.py::TestNanMedian::test_nanmedian[_param_6_{axis=0, keepdims=False, overwrite_input=True, shape=(3, 4, 5)}] tests/cupy_tests/statistics_tests/test_meanvar.py::TestNanMedian::test_nanmedian[_param_7_{axis=0, keepdims=False, overwrite_input=False, shape=(3, 4, 5)}] tests/cupy_tests/statistics_tests/test_meanvar.py::TestNanMedian::test_nanmedian[_param_8_{axis=1, keepdims=True, overwrite_input=True, shape=(3, 4, 5)}] tests/cupy_tests/statistics_tests/test_meanvar.py::TestNanMedian::test_nanmedian[_param_9_{axis=1, keepdims=True, overwrite_input=False, shape=(3, 4, 5)}] tests/cupy_tests/statistics_tests/test_meanvar.py::TestNanMedian::test_nanmedian[_param_10_{axis=1, keepdims=False, overwrite_input=True, shape=(3, 4, 5)}] tests/cupy_tests/statistics_tests/test_meanvar.py::TestNanMedian::test_nanmedian[_param_11_{axis=1, keepdims=False, overwrite_input=False, shape=(3, 4, 5)}] /home/kmaeh/.pyenv/versions/rocm-ci/lib/python3.8/site-packages/numpy/lib/nanfunctions.py:1113: RuntimeWarning: All-NaN slice encountered r, k = function_base._ureduce(a, func=_nanmedian, axis=axis, out=out, tests/cupy_tests/statistics_tests/test_meanvar.py::TestProductZeroLength::test_external_mean_zero_len[_param_1_{func='mean', params=((0,), None)}] tests/cupy_tests/statistics_tests/test_meanvar.py::TestProductZeroLength::test_external_mean_zero_len[_param_2_{func='mean', params=((0, 0), None)}] tests/cupy_tests/statistics_tests/test_meanvar.py::TestProductZeroLength::test_external_mean_zero_len[_param_3_{func='mean', params=((0, 0), 1)}] tests/cupy_tests/statistics_tests/test_meanvar.py::TestProductZeroLength::test_external_mean_zero_len[_param_4_{func='mean', params=((0, 0, 0), None)}] tests/cupy_tests/statistics_tests/test_meanvar.py::TestProductZeroLength::test_external_mean_zero_len[_param_5_{func='mean', params=((0, 0, 0), (0, 2))}] /home/kmaeh/.pyenv/versions/rocm-ci/lib/python3.8/site-packages/numpy/core/fromnumeric.py:3372: RuntimeWarning: Mean of empty slice. return _methods._mean(a, axis=axis, dtype=dtype, tests/cupy_tests/statistics_tests/test_meanvar.py::TestProductZeroLength::test_external_mean_zero_len[_param_1_{func='mean', params=((0,), None)}] tests/cupy_tests/statistics_tests/test_meanvar.py::TestProductZeroLength::test_external_mean_zero_len[_param_2_{func='mean', params=((0, 0), None)}] tests/cupy_tests/statistics_tests/test_meanvar.py::TestProductZeroLength::test_external_mean_zero_len[_param_4_{func='mean', params=((0, 0, 0), None)}] /home/kmaeh/.pyenv/versions/rocm-ci/lib/python3.8/site-packages/numpy/core/_methods.py:170: RuntimeWarning: invalid value encountered in double_scalars ret = ret.dtype.type(ret / rcount) tests/cupy_tests/statistics_tests/test_meanvar.py::TestProductZeroLength::test_external_mean_zero_len[_param_1_{func='mean', params=((0,), None)}] tests/cupy_tests/statistics_tests/test_meanvar.py::TestProductZeroLength::test_external_mean_zero_len[_param_2_{func='mean', params=((0, 0), None)}] tests/cupy_tests/statistics_tests/test_meanvar.py::TestProductZeroLength::test_external_mean_zero_len[_param_4_{func='mean', params=((0, 0, 0), None)}] /home/kmaeh/.pyenv/versions/rocm-ci/lib/python3.8/site-packages/numpy/core/_methods.py:170: RuntimeWarning: invalid value encountered in true_divide ret = ret.dtype.type(ret / rcount) tests/cupy_tests/statistics_tests/test_meanvar.py::TestProductZeroLength::test_external_mean_zero_len[_param_1_{func='mean', params=((0,), None)}] tests/cupy_tests/statistics_tests/test_meanvar.py::TestProductZeroLength::test_external_mean_zero_len[_param_2_{func='mean', params=((0, 0), None)}] tests/cupy_tests/statistics_tests/test_meanvar.py::TestProductZeroLength::test_external_mean_zero_len[_param_4_{func='mean', params=((0, 0, 0), None)}] /home/kmaeh/.pyenv/versions/rocm-ci/lib/python3.8/site-packages/numpy/core/_methods.py:168: RuntimeWarning: invalid value encountered in true_divide ret = arr.dtype.type(ret / rcount) tests/cupy_tests/statistics_tests/test_meanvar.py::TestProductZeroLength::test_external_mean_zero_len[_param_7_{func='std', params=((0,), None)}] tests/cupy_tests/statistics_tests/test_meanvar.py::TestProductZeroLength::test_external_mean_zero_len[_param_8_{func='std', params=((0, 0), None)}] tests/cupy_tests/statistics_tests/test_meanvar.py::TestProductZeroLength::test_external_mean_zero_len[_param_9_{func='std', params=((0, 0), 1)}] tests/cupy_tests/statistics_tests/test_meanvar.py::TestProductZeroLength::test_external_mean_zero_len[_param_10_{func='std', params=((0, 0, 0), None)}] tests/cupy_tests/statistics_tests/test_meanvar.py::TestProductZeroLength::test_external_mean_zero_len[_param_11_{func='std', params=((0, 0, 0), (0, 2))}] /home/kmaeh/.pyenv/versions/rocm-ci/lib/python3.8/site-packages/numpy/core/_methods.py:233: RuntimeWarning: Degrees of freedom <= 0 for slice ret = _var(a, axis=axis, dtype=dtype, out=out, ddof=ddof, tests/cupy_tests/statistics_tests/test_meanvar.py::TestProductZeroLength::test_external_mean_zero_len[_param_7_{func='std', params=((0,), None)}] tests/cupy_tests/statistics_tests/test_meanvar.py::TestProductZeroLength::test_external_mean_zero_len[_param_8_{func='std', params=((0, 0), None)}] tests/cupy_tests/statistics_tests/test_meanvar.py::TestProductZeroLength::test_external_mean_zero_len[_param_10_{func='std', params=((0, 0, 0), None)}] tests/cupy_tests/statistics_tests/test_meanvar.py::TestProductZeroLength::test_external_mean_zero_len[_param_13_{func='var', params=((0,), None)}] tests/cupy_tests/statistics_tests/test_meanvar.py::TestProductZeroLength::test_external_mean_zero_len[_param_14_{func='var', params=((0, 0), None)}] tests/cupy_tests/statistics_tests/test_meanvar.py::TestProductZeroLength::test_external_mean_zero_len[_param_16_{func='var', params=((0, 0, 0), None)}] /home/kmaeh/.pyenv/versions/rocm-ci/lib/python3.8/site-packages/numpy/core/_methods.py:194: RuntimeWarning: invalid value encountered in true_divide arrmean = um.true_divide( tests/cupy_tests/statistics_tests/test_meanvar.py::TestProductZeroLength::test_external_mean_zero_len[_param_7_{func='std', params=((0,), None)}] tests/cupy_tests/statistics_tests/test_meanvar.py::TestProductZeroLength::test_external_mean_zero_len[_param_8_{func='std', params=((0, 0), None)}] tests/cupy_tests/statistics_tests/test_meanvar.py::TestProductZeroLength::test_external_mean_zero_len[_param_10_{func='std', params=((0, 0, 0), None)}] tests/cupy_tests/statistics_tests/test_meanvar.py::TestProductZeroLength::test_external_mean_zero_len[_param_13_{func='var', params=((0,), None)}] tests/cupy_tests/statistics_tests/test_meanvar.py::TestProductZeroLength::test_external_mean_zero_len[_param_14_{func='var', params=((0, 0), None)}] tests/cupy_tests/statistics_tests/test_meanvar.py::TestProductZeroLength::test_external_mean_zero_len[_param_16_{func='var', params=((0, 0, 0), None)}] /home/kmaeh/.pyenv/versions/rocm-ci/lib/python3.8/site-packages/numpy/core/_methods.py:226: RuntimeWarning: invalid value encountered in double_scalars ret = ret.dtype.type(ret / rcount) tests/cupy_tests/statistics_tests/test_meanvar.py::TestProductZeroLength::test_external_mean_zero_len[_param_7_{func='std', params=((0,), None)}] tests/cupy_tests/statistics_tests/test_meanvar.py::TestProductZeroLength::test_external_mean_zero_len[_param_8_{func='std', params=((0, 0), None)}] tests/cupy_tests/statistics_tests/test_meanvar.py::TestProductZeroLength::test_external_mean_zero_len[_param_10_{func='std', params=((0, 0, 0), None)}] tests/cupy_tests/statistics_tests/test_meanvar.py::TestProductZeroLength::test_external_mean_zero_len[_param_13_{func='var', params=((0,), None)}] tests/cupy_tests/statistics_tests/test_meanvar.py::TestProductZeroLength::test_external_mean_zero_len[_param_14_{func='var', params=((0, 0), None)}] tests/cupy_tests/statistics_tests/test_meanvar.py::TestProductZeroLength::test_external_mean_zero_len[_param_16_{func='var', params=((0, 0, 0), None)}] /home/kmaeh/.pyenv/versions/rocm-ci/lib/python3.8/site-packages/numpy/core/_methods.py:226: RuntimeWarning: invalid value encountered in true_divide ret = ret.dtype.type(ret / rcount) tests/cupy_tests/statistics_tests/test_meanvar.py::TestProductZeroLength::test_external_mean_zero_len[_param_13_{func='var', params=((0,), None)}] tests/cupy_tests/statistics_tests/test_meanvar.py::TestProductZeroLength::test_external_mean_zero_len[_param_14_{func='var', params=((0, 0), None)}] tests/cupy_tests/statistics_tests/test_meanvar.py::TestProductZeroLength::test_external_mean_zero_len[_param_15_{func='var', params=((0, 0), 1)}] tests/cupy_tests/statistics_tests/test_meanvar.py::TestProductZeroLength::test_external_mean_zero_len[_param_16_{func='var', params=((0, 0, 0), None)}] tests/cupy_tests/statistics_tests/test_meanvar.py::TestProductZeroLength::test_external_mean_zero_len[_param_17_{func='var', params=((0, 0, 0), (0, 2))}] /home/kmaeh/.pyenv/versions/rocm-ci/lib/python3.8/site-packages/numpy/core/fromnumeric.py:3621: RuntimeWarning: Degrees of freedom <= 0 for slice return _methods._var(a, axis=axis, dtype=dtype, out=out, ddof=ddof, tests/cupyx_tests/jit_tests/test_raw.py: 21 warnings /global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupyx/jit/_builtin_funcs.py:334: RuntimeWarning: mask 4294967295 is ignored on HIP warnings.warn(f'mask {mask} is ignored on HIP', RuntimeWarning) tests/cupyx_tests/jit_tests/test_raw.py::TestRaw::test_syncwarp_mask /global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupyx/jit/_builtin_funcs.py:131: RuntimeWarning: mask is ignored on HIP warnings.warn(f'mask {mask} is ignored on HIP', RuntimeWarning) tests/cupyx_tests/scipy_tests/ndimage_tests/test_interpolation.py::TestShift::test_shift_int[_param_217_{cval=nan, mode='constant', order=0, output=float64, prefilter=True, shift=0.1}] tests/cupyx_tests/scipy_tests/ndimage_tests/test_interpolation.py::TestShift::test_shift_int[_param_220_{cval=nan, mode='constant', order=1, output=float64, prefilter=True, shift=0.1}] tests/cupyx_tests/scipy_tests/ndimage_tests/test_interpolation.py::TestShift::test_shift_int[_param_223_{cval=nan, mode='constant', order=3, output=float64, prefilter=True, shift=0.1}] tests/cupyx_tests/scipy_tests/ndimage_tests/test_interpolation.py::TestShift::test_shift_int[_param_226_{cval=inf, mode='constant', order=0, output=float64, prefilter=True, shift=0.1}] tests/cupyx_tests/scipy_tests/ndimage_tests/test_interpolation.py::TestShift::test_shift_int[_param_229_{cval=inf, mode='constant', order=1, output=float64, prefilter=True, shift=0.1}] tests/cupyx_tests/scipy_tests/ndimage_tests/test_interpolation.py::TestShift::test_shift_int[_param_232_{cval=inf, mode='constant', order=3, output=float64, prefilter=True, shift=0.1}] tests/cupyx_tests/scipy_tests/ndimage_tests/test_interpolation.py::TestShift::test_shift_int[_param_235_{cval=-inf, mode='constant', order=0, output=float64, prefilter=True, shift=0.1}] tests/cupyx_tests/scipy_tests/ndimage_tests/test_interpolation.py::TestShift::test_shift_int[_param_238_{cval=-inf, mode='constant', order=1, output=float64, prefilter=True, shift=0.1}] tests/cupyx_tests/scipy_tests/ndimage_tests/test_interpolation.py::TestShift::test_shift_int[_param_241_{cval=-inf, mode='constant', order=3, output=float64, prefilter=True, shift=0.1}] /global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/tests/cupyx_tests/scipy_tests/ndimage_tests/test_interpolation.py:670: RuntimeWarning: invalid value encountered in remainder float_out = self._shift(xp, scp, a.astype(xp.float64)) % 1 tests/cupyx_tests/scipy_tests/ndimage_tests/test_measurements.py: 14 warnings /home/kmaeh/.pyenv/versions/rocm-ci/lib/python3.8/site-packages/scipy/ndimage/measurements.py:684: RuntimeWarning: invalid value encountered in true_divide return sum / numpy.asanyarray(count).astype(numpy.float64) tests/cupyx_tests/scipy_tests/ndimage_tests/test_measurements.py: 28 warnings /home/kmaeh/.pyenv/versions/rocm-ci/lib/python3.8/site-packages/scipy/ndimage/measurements.py:736: RuntimeWarning: invalid value encountered in true_divide return sum_c_sq / np.asanyarray(count).astype(float) tests/cupyx_tests/scipy_tests/signal_tests/test_signaltools.py::TestWiener::test_wiener[_param_0_{im=(10,), mysize=3, noise=False}] tests/cupyx_tests/scipy_tests/signal_tests/test_signaltools.py::TestWiener::test_wiener[_param_1_{im=(10,), mysize=3, noise=True}] tests/cupyx_tests/scipy_tests/signal_tests/test_signaltools.py::TestWiener::test_wiener[_param_2_{im=(10,), mysize=4, noise=False}] tests/cupyx_tests/scipy_tests/signal_tests/test_signaltools.py::TestWiener::test_wiener[_param_3_{im=(10,), mysize=4, noise=True}] tests/cupyx_tests/scipy_tests/signal_tests/test_signaltools.py::TestWiener::test_wiener[_param_4_{im=(10,), mysize=(3, 4, 5), noise=False}] tests/cupyx_tests/scipy_tests/signal_tests/test_signaltools.py::TestWiener::test_wiener[_param_5_{im=(10,), mysize=(3, 4, 5), noise=True}] /home/kmaeh/.pyenv/versions/rocm-ci/lib/python3.8/site-packages/scipy/signal/signaltools.py:1598: RuntimeWarning: divide by zero encountered in true_divide res *= (1 - noise / lVar) tests/cupyx_tests/scipy_tests/signal_tests/test_signaltools.py::TestWiener::test_wiener[_param_0_{im=(10,), mysize=3, noise=False}] tests/cupyx_tests/scipy_tests/signal_tests/test_signaltools.py::TestWiener::test_wiener[_param_1_{im=(10,), mysize=3, noise=True}] tests/cupyx_tests/scipy_tests/signal_tests/test_signaltools.py::TestWiener::test_wiener[_param_2_{im=(10,), mysize=4, noise=False}] tests/cupyx_tests/scipy_tests/signal_tests/test_signaltools.py::TestWiener::test_wiener[_param_3_{im=(10,), mysize=4, noise=True}] tests/cupyx_tests/scipy_tests/signal_tests/test_signaltools.py::TestWiener::test_wiener[_param_4_{im=(10,), mysize=(3, 4, 5), noise=False}] tests/cupyx_tests/scipy_tests/signal_tests/test_signaltools.py::TestWiener::test_wiener[_param_5_{im=(10,), mysize=(3, 4, 5), noise=True}] /home/kmaeh/.pyenv/versions/rocm-ci/lib/python3.8/site-packages/scipy/signal/signaltools.py:1598: RuntimeWarning: invalid value encountered in multiply res *= (1 - noise / lVar) tests/cupyx_tests/scipy_tests/sparse_tests/test_csr.py: 32 warnings /home/kmaeh/.pyenv/versions/rocm-ci/lib/python3.8/site-packages/scipy/sparse/base.py:595: RuntimeWarning: invalid value encountered in true_divide return np.true_divide(self.todense(), other) tests/cupyx_tests/scipy_tests/sparse_tests/test_csr.py: 28 warnings /home/kmaeh/.pyenv/versions/rocm-ci/lib/python3.8/site-packages/scipy/sparse/base.py:595: RuntimeWarning: divide by zero encountered in true_divide return np.true_divide(self.todense(), other) tests/cupyx_tests/scipy_tests/sparse_tests/test_csr.py: 96 warnings /global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/tests/cupyx_tests/scipy_tests/sparse_tests/test_csr.py:1888: SparseEfficiencyWarning: Taking maximum (minimum) with > 0 (< 0) number results to a dense matrix. return getattr(a, self.opt)(0.5) tests/cupyx_tests/scipy_tests/sparse_tests/test_csr.py: 96 warnings /global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/tests/cupyx_tests/scipy_tests/sparse_tests/test_csr.py:1893: SparseEfficiencyWarning: Taking maximum (minimum) with > 0 (< 0) number results to a dense matrix. return getattr(a, self.opt)(-0.5) tests/cupyx_tests/scipy_tests/sparse_tests/test_csr.py: 16 warnings /home/kmaeh/.pyenv/versions/rocm-ci/lib/python3.8/site-packages/scipy/sparse/compressed.py:211: RuntimeWarning: invalid value encountered in less res = self._with_data(op(self.data, other), copy=True) tests/cupyx_tests/scipy_tests/sparse_tests/test_csr.py: 16 warnings /home/kmaeh/.pyenv/versions/rocm-ci/lib/python3.8/site-packages/scipy/sparse/compressed.py:211: RuntimeWarning: invalid value encountered in greater res = self._with_data(op(self.data, other), copy=True) tests/cupyx_tests/scipy_tests/sparse_tests/test_csr.py: 16 warnings /home/kmaeh/.pyenv/versions/rocm-ci/lib/python3.8/site-packages/scipy/sparse/compressed.py:211: RuntimeWarning: invalid value encountered in less_equal res = self._with_data(op(self.data, other), copy=True) tests/cupyx_tests/scipy_tests/sparse_tests/test_csr.py: 16 warnings /home/kmaeh/.pyenv/versions/rocm-ci/lib/python3.8/site-packages/scipy/sparse/compressed.py:211: RuntimeWarning: invalid value encountered in greater_equal res = self._with_data(op(self.data, other), copy=True) tests/cupyx_tests/scipy_tests/sparse_tests/test_csr.py: 6 warnings tests/cupyx_tests/scipy_tests/sparse_tests/test_index.py: 368 warnings /home/kmaeh/.pyenv/versions/rocm-ci/lib/python3.8/site-packages/scipy/sparse/_index.py:125: SparseEfficiencyWarning: Changing the sparsity structure of a csr_matrix is expensive. lil_matrix is more efficient. self._set_arrayXarray(i, j, x) tests/cupyx_tests/scipy_tests/sparse_tests/test_index.py: 16 warnings /home/kmaeh/.pyenv/versions/rocm-ci/lib/python3.8/site-packages/scipy/sparse/_index.py:116: SparseEfficiencyWarning: Changing the sparsity structure of a csr_matrix is expensive. lil_matrix is more efficient. self._set_arrayXarray_sparse(i, j, x) tests/cupyx_tests/scipy_tests/sparse_tests/test_index.py: 16 warnings /home/kmaeh/.pyenv/versions/rocm-ci/lib/python3.8/site-packages/scipy/sparse/_index.py:116: SparseEfficiencyWarning: Changing the sparsity structure of a csc_matrix is expensive. lil_matrix is more efficient. self._set_arrayXarray_sparse(i, j, x) tests/cupyx_tests/scipy_tests/sparse_tests/test_index.py: 368 warnings /home/kmaeh/.pyenv/versions/rocm-ci/lib/python3.8/site-packages/scipy/sparse/_index.py:125: SparseEfficiencyWarning: Changing the sparsity structure of a csc_matrix is expensive. lil_matrix is more efficient. self._set_arrayXarray(i, j, x) tests/cupyx_tests/scipy_tests/sparse_tests/test_index.py: 12 warnings /home/kmaeh/.pyenv/versions/rocm-ci/lib/python3.8/site-packages/scipy/sparse/_index.py:82: SparseEfficiencyWarning: Changing the sparsity structure of a csr_matrix is expensive. lil_matrix is more efficient. self._set_intXint(row, col, x.flat[0]) tests/cupyx_tests/scipy_tests/sparse_tests/test_index.py: 12 warnings /home/kmaeh/.pyenv/versions/rocm-ci/lib/python3.8/site-packages/scipy/sparse/_index.py:82: SparseEfficiencyWarning: Changing the sparsity structure of a csc_matrix is expensive. lil_matrix is more efficient. self._set_intXint(row, col, x.flat[0]) tests/cupyx_tests/scipy_tests/sparse_tests/test_linalg.py: 48 warnings /home/kmaeh/.pyenv/versions/rocm-ci/lib/python3.8/site-packages/scipy/sparse/linalg/dsolve/linsolve.py:553: SparseEfficiencyWarning: CSR matrix format is required. Converting to CSR matrix. warn('CSR matrix format is required. Converting to CSR matrix.', tests/cupyx_tests/scipy_tests/sparse_tests/test_linalg.py: 24 warnings /global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/cupyx/scipy/sparse/linalg/_solve.py:437: SparseEfficiencyWarning: CSR or CSC format is required. Converting to CSR format. warnings.warn('CSR or CSC format is required. Converting to CSR ' tests/cupyx_tests/scipy_tests/sparse_tests/test_linalg.py: 24 warnings /home/kmaeh/.pyenv/versions/rocm-ci/lib/python3.8/site-packages/scipy/sparse/linalg/dsolve/linsolve.py:318: SparseEfficiencyWarning: splu requires CSC matrix format warn('splu requires CSC matrix format', SparseEfficiencyWarning) tests/cupyx_tests/scipy_tests/sparse_tests/test_linalg.py: 24 warnings /home/kmaeh/.pyenv/versions/rocm-ci/lib/python3.8/site-packages/scipy/sparse/linalg/dsolve/linsolve.py:407: SparseEfficiencyWarning: splu requires CSC matrix format warn('splu requires CSC matrix format', SparseEfficiencyWarning) tests/cupyx_tests/scipy_tests/sparse_tests/test_linalg.py: 80 warnings /global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/tests/cupyx_tests/scipy_tests/sparse_tests/test_linalg.py:1572: RuntimeWarning: divide by zero encountered in true_divide M = xp.diag(1.0 / xp.diag(a)) tests/cupyx_tests/scipy_tests/sparse_tests/test_linalg.py: 80 warnings /home/kmaeh/.pyenv/versions/rocm-ci/lib/python3.8/site-packages/scipy/sparse/linalg/isolve/minres.py:190: RuntimeWarning: invalid value encountered in multiply v = s*y -- Docs: https://docs.pytest.org/en/stable/warnings.html - generated html file: file:///global/scratch/kmaeh/cupy-rocm-ci-work/tmp.oE8iI9EKKE/cupy/_output/report.html - =========================== short test summary info ============================ FAILED tests/cupy_tests/core_tests/test_ndarray.py::TestNdarrayCopy::test_copy_multi_device_with_stream FAILED tests/cupy_tests/core_tests/test_ndarray_copy_and_view.py::TestArrayCopyAndView::test_astype_boolean_view FAILED tests/cupy_tests/core_tests/test_ndarray_elementwise_op.py::TestArrayElementwiseOp::test_add_array_boolarray FAILED tests/cupy_tests/core_tests/test_ndarray_elementwise_op.py::TestArrayElementwiseOp::test_iadd_array_boolarray FAILED tests/cupy_tests/linalg_tests/test_product.py::TestDot_param_0_{shape=((2, 3, 4), (3, 4, 2)), trans_a=True, trans_b=True}::test_dot FAILED tests/cupy_tests/linalg_tests/test_product.py::TestDot_param_1_{shape=((2, 3, 4), (3, 4, 2)), trans_a=True, trans_b=False}::test_dot FAILED tests/cupy_tests/linalg_tests/test_product.py::TestDot_param_2_{shape=((2, 3, 4), (3, 4, 2)), trans_a=False, trans_b=True}::test_dot FAILED tests/cupy_tests/linalg_tests/test_product.py::TestDot_param_3_{shape=((2, 3, 4), (3, 4, 2)), trans_a=False, trans_b=False}::test_dot FAILED tests/cupy_tests/linalg_tests/test_product.py::TestDot_param_8_{shape=((1, 1), (1, 2)), trans_a=True, trans_b=True}::test_dot FAILED tests/cupy_tests/linalg_tests/test_product.py::TestDot_param_9_{shape=((1, 1), (1, 2)), trans_a=True, trans_b=False}::test_dot FAILED tests/cupy_tests/linalg_tests/test_product.py::TestDot_param_10_{shape=((1, 1), (1, 2)), trans_a=False, trans_b=True}::test_dot FAILED tests/cupy_tests/linalg_tests/test_product.py::TestDot_param_11_{shape=((1, 1), (1, 2)), trans_a=False, trans_b=False}::test_dot FAILED tests/cupy_tests/linalg_tests/test_product.py::TestDot_param_16_{shape=((2, 1), (1, 1)), trans_a=True, trans_b=True}::test_dot FAILED tests/cupy_tests/linalg_tests/test_product.py::TestDot_param_17_{shape=((2, 1), (1, 1)), trans_a=True, trans_b=False}::test_dot FAILED tests/cupy_tests/linalg_tests/test_product.py::TestDot_param_18_{shape=((2, 1), (1, 1)), trans_a=False, trans_b=True}::test_dot FAILED tests/cupy_tests/linalg_tests/test_product.py::TestDot_param_19_{shape=((2, 1), (1, 1)), trans_a=False, trans_b=False}::test_dot FAILED tests/cupy_tests/linalg_tests/test_product.py::TestDot_param_20_{shape=((1, 2), (2, 3)), trans_a=True, trans_b=True}::test_dot FAILED tests/cupy_tests/linalg_tests/test_product.py::TestDot_param_21_{shape=((1, 2), (2, 3)), trans_a=True, trans_b=False}::test_dot FAILED tests/cupy_tests/linalg_tests/test_product.py::TestDot_param_22_{shape=((1, 2), (2, 3)), trans_a=False, trans_b=True}::test_dot FAILED tests/cupy_tests/linalg_tests/test_product.py::TestDot_param_23_{shape=((1, 2), (2, 3)), trans_a=False, trans_b=False}::test_dot FAILED tests/cupy_tests/linalg_tests/test_product.py::TestDot_param_24_{shape=((2, 1), (1, 3)), trans_a=True, trans_b=True}::test_dot FAILED tests/cupy_tests/linalg_tests/test_product.py::TestDot_param_25_{shape=((2, 1), (1, 3)), trans_a=True, trans_b=False}::test_dot FAILED tests/cupy_tests/linalg_tests/test_product.py::TestDot_param_26_{shape=((2, 1), (1, 3)), trans_a=False, trans_b=True}::test_dot FAILED tests/cupy_tests/linalg_tests/test_product.py::TestDot_param_27_{shape=((2, 1), (1, 3)), trans_a=False, trans_b=False}::test_dot FAILED tests/cupy_tests/linalg_tests/test_product.py::TestDot_param_28_{shape=((2, 3), (3, 1)), trans_a=True, trans_b=True}::test_dot FAILED tests/cupy_tests/linalg_tests/test_product.py::TestDot_param_29_{shape=((2, 3), (3, 1)), trans_a=True, trans_b=False}::test_dot FAILED tests/cupy_tests/linalg_tests/test_product.py::TestDot_param_30_{shape=((2, 3), (3, 1)), trans_a=False, trans_b=True}::test_dot FAILED tests/cupy_tests/linalg_tests/test_product.py::TestDot_param_31_{shape=((2, 3), (3, 1)), trans_a=False, trans_b=False}::test_dot FAILED tests/cupy_tests/linalg_tests/test_product.py::TestDot_param_32_{shape=((2, 3), (3, 4)), trans_a=True, trans_b=True}::test_dot FAILED tests/cupy_tests/linalg_tests/test_product.py::TestDot_param_33_{shape=((2, 3), (3, 4)), trans_a=True, trans_b=False}::test_dot FAILED tests/cupy_tests/linalg_tests/test_product.py::TestDot_param_34_{shape=((2, 3), (3, 4)), trans_a=False, trans_b=True}::test_dot FAILED tests/cupy_tests/linalg_tests/test_product.py::TestDot_param_35_{shape=((2, 3), (3, 4)), trans_a=False, trans_b=False}::test_dot FAILED tests/cupy_tests/linalg_tests/test_product.py::TestDot_param_64_{shape=((2,), (2, 4)), trans_a=True, trans_b=True}::test_dot FAILED tests/cupy_tests/linalg_tests/test_product.py::TestDot_param_65_{shape=((2,), (2, 4)), trans_a=True, trans_b=False}::test_dot FAILED tests/cupy_tests/linalg_tests/test_product.py::TestDot_param_66_{shape=((2,), (2, 4)), trans_a=False, trans_b=True}::test_dot FAILED tests/cupy_tests/linalg_tests/test_product.py::TestDot_param_67_{shape=((2,), (2, 4)), trans_a=False, trans_b=False}::test_dot FAILED tests/cupy_tests/linalg_tests/test_product.py::TestDot_param_68_{shape=((4, 2), (2,)), trans_a=True, trans_b=True}::test_dot FAILED tests/cupy_tests/linalg_tests/test_product.py::TestDot_param_69_{shape=((4, 2), (2,)), trans_a=True, trans_b=False}::test_dot FAILED tests/cupy_tests/linalg_tests/test_product.py::TestDot_param_70_{shape=((4, 2), (2,)), trans_a=False, trans_b=True}::test_dot FAILED tests/cupy_tests/linalg_tests/test_product.py::TestDot_param_71_{shape=((4, 2), (2,)), trans_a=False, trans_b=False}::test_dot FAILED tests/cupy_tests/linalg_tests/test_product.py::TestProduct::test_dot_with_single_elem_array1 FAILED tests/cupy_tests/linalg_tests/test_product.py::TestProduct::test_dot_with_single_elem_array2 FAILED tests/cupy_tests/linalg_tests/test_product.py::TestProduct::test_kron FAILED tests/cupy_tests/linalg_tests/test_product.py::TestProduct::test_multidim_inner FAILED tests/cupy_tests/linalg_tests/test_product.py::TestProduct::test_multidim_kron FAILED tests/cupy_tests/linalg_tests/test_product.py::TestProduct::test_multidim_outer FAILED tests/cupy_tests/linalg_tests/test_product.py::TestProduct::test_outer FAILED tests/cupy_tests/linalg_tests/test_product.py::TestProduct::test_reversed_kron FAILED tests/cupy_tests/linalg_tests/test_product.py::TestProduct::test_reversed_outer FAILED tests/cupy_tests/linalg_tests/test_product.py::TestProduct::test_tensordot FAILED tests/cupy_tests/linalg_tests/test_product.py::TestProduct::test_tensordot_with_int_axes FAILED tests/cupy_tests/linalg_tests/test_product.py::TestProduct::test_tensordot_with_list_axes FAILED tests/cupy_tests/linalg_tests/test_product.py::TestProduct::test_transposed_dot FAILED tests/cupy_tests/linalg_tests/test_product.py::TestProduct::test_transposed_dot_with_out FAILED tests/cupy_tests/linalg_tests/test_product.py::TestProduct::test_transposed_higher_order_inner FAILED tests/cupy_tests/linalg_tests/test_product.py::TestProduct::test_transposed_tensordot FAILED tests/cupy_tests/linalg_tests/test_product.py::TestProduct::test_transposed_tensordot_with_int_axes FAILED tests/cupy_tests/linalg_tests/test_product.py::TestProduct::test_transposed_tensordot_with_list_axes FAILED tests/cupy_tests/linalg_tests/test_product.py::TestMatrixPower::test_matrix_power_2 FAILED tests/cupy_tests/linalg_tests/test_product.py::TestMatrixPower::test_matrix_power_3 FAILED tests/cupy_tests/linalg_tests/test_product.py::TestMatrixPower::test_matrix_power_large FAILED tests/cupy_tests/linalg_tests/test_product.py::TestMatrixPower::test_matrix_power_of_two FAILED tests/cupy_tests/linalg_tests/test_solve.py::TestLstsq::test_lstsq_solutions FAILED tests/cupy_tests/math_tests/test_matmul.py::TestMatmul_param_0_{shape_pair=((3, 2), (2, 4))}::test_cupy_matmul FAILED tests/cupy_tests/math_tests/test_matmul.py::TestMatmul_param_0_{shape_pair=((3, 2), (2, 4))}::test_operator_matmul FAILED tests/cupy_tests/math_tests/test_matmul.py::TestMatmul_param_4_{shape_pair=((2,), (2, 4))}::test_cupy_matmul FAILED tests/cupy_tests/math_tests/test_matmul.py::TestMatmul_param_4_{shape_pair=((2,), (2, 4))}::test_operator_matmul FAILED tests/cupy_tests/math_tests/test_matmul.py::TestMatmul_param_6_{shape_pair=((3, 2), (2,))}::test_cupy_matmul FAILED tests/cupy_tests/math_tests/test_matmul.py::TestMatmul_param_6_{shape_pair=((3, 2), (2,))}::test_operator_matmul FAILED tests/cupy_tests/math_tests/test_misc.py::TestConvolve_param_0_{mode='valid'}::test_convolve_diff_types FAILED tests/cupy_tests/math_tests/test_misc.py::TestConvolve_param_1_{mode='same'}::test_convolve_diff_types FAILED tests/cupy_tests/math_tests/test_misc.py::TestConvolve_param_2_{mode='full'}::test_convolve_diff_types FAILED tests/cupy_tests/statistics_tests/test_correlation.py::TestCorrelate_param_0_{mode='valid'}::test_correlate_diff_types FAILED tests/cupy_tests/statistics_tests/test_correlation.py::TestCorrelate_param_1_{mode='full'}::test_correlate_diff_types FAILED tests/cupy_tests/statistics_tests/test_correlation.py::TestCorrelate_param_2_{mode='same'}::test_correlate_diff_types FAILED tests/cupyx_tests/scipy_tests/ndimage_tests/test_interpolation.py::TestMapCoordinatesHalfInteger::test_map_coordinates_float[_param_46_{mode='grid-constant', order=4}] FAILED tests/cupyx_tests/scipy_tests/ndimage_tests/test_interpolation.py::TestMapCoordinatesHalfInteger::test_map_coordinates_float[_param_47_{mode='grid-constant', order=5}] FAILED tests/cupyx_tests/scipy_tests/ndimage_tests/test_interpolation.py::TestZoomOutputSize1::test_zoom_output_size1[_param_4_{grid_mode=False, mode='constant', order=4, shape=(5, 5, 2), zoom=(2, 2, 0.5)}] FAILED tests/cupyx_tests/scipy_tests/ndimage_tests/test_interpolation.py::TestZoomOutputSize1::test_zoom_output_size1[_param_5_{grid_mode=False, mode='constant', order=5, shape=(5, 5, 2), zoom=(2, 2, 0.5)}] FAILED tests/cupyx_tests/scipy_tests/ndimage_tests/test_interpolation.py::TestZoomOutputSize1::test_zoom_output_size1[_param_16_{grid_mode=False, mode='mirror', order=4, shape=(5, 5, 2), zoom=(2, 2, 0.5)}] FAILED tests/cupyx_tests/scipy_tests/ndimage_tests/test_interpolation.py::TestZoomOutputSize1::test_zoom_output_size1[_param_17_{grid_mode=False, mode='mirror', order=5, shape=(5, 5, 2), zoom=(2, 2, 0.5)}] FAILED tests/cupyx_tests/scipy_tests/ndimage_tests/test_interpolation.py::TestZoomOutputSize1::test_zoom_output_size1[_param_22_{grid_mode=False, mode='wrap', order=4, shape=(5, 5, 2), zoom=(2, 2, 0.5)}] FAILED tests/cupyx_tests/scipy_tests/ndimage_tests/test_interpolation.py::TestZoomOutputSize1::test_zoom_output_size1[_param_23_{grid_mode=False, mode='wrap', order=5, shape=(5, 5, 2), zoom=(2, 2, 0.5)}] FAILED tests/cupyx_tests/scipy_tests/ndimage_tests/test_interpolation.py::TestZoomOutputSize1::test_zoom_output_size1[_param_46_{grid_mode=False, mode='grid-constant', order=4, shape=(5, 5, 2), zoom=(2, 2, 0.5)}] FAILED tests/cupyx_tests/scipy_tests/ndimage_tests/test_interpolation.py::TestZoomOutputSize1::test_zoom_output_size1[_param_47_{grid_mode=False, mode='grid-constant', order=5, shape=(5, 5, 2), zoom=(2, 2, 0.5)}] FAILED tests/cupyx_tests/scipy_tests/ndimage_tests/test_interpolation.py::TestZoomOutputSize1::test_zoom_output_size1[_param_52_{grid_mode=True, mode='constant', order=4, shape=(5, 5, 2), zoom=(2, 2, 0.5)}] FAILED tests/cupyx_tests/scipy_tests/ndimage_tests/test_interpolation.py::TestZoomOutputSize1::test_zoom_output_size1[_param_53_{grid_mode=True, mode='constant', order=5, shape=(5, 5, 2), zoom=(2, 2, 0.5)}] FAILED tests/cupyx_tests/scipy_tests/ndimage_tests/test_interpolation.py::TestZoomOutputSize1::test_zoom_output_size1[_param_64_{grid_mode=True, mode='mirror', order=4, shape=(5, 5, 2), zoom=(2, 2, 0.5)}] FAILED tests/cupyx_tests/scipy_tests/ndimage_tests/test_interpolation.py::TestZoomOutputSize1::test_zoom_output_size1[_param_65_{grid_mode=True, mode='mirror', order=5, shape=(5, 5, 2), zoom=(2, 2, 0.5)}] FAILED tests/cupyx_tests/scipy_tests/ndimage_tests/test_interpolation.py::TestZoomOutputSize1::test_zoom_output_size1[_param_70_{grid_mode=True, mode='wrap', order=4, shape=(5, 5, 2), zoom=(2, 2, 0.5)}] FAILED tests/cupyx_tests/scipy_tests/ndimage_tests/test_interpolation.py::TestZoomOutputSize1::test_zoom_output_size1[_param_71_{grid_mode=True, mode='wrap', order=5, shape=(5, 5, 2), zoom=(2, 2, 0.5)}] FAILED tests/cupyx_tests/scipy_tests/ndimage_tests/test_interpolation.py::TestZoomOutputSize1::test_zoom_output_size1[_param_94_{grid_mode=True, mode='grid-constant', order=4, shape=(5, 5, 2), zoom=(2, 2, 0.5)}] FAILED tests/cupyx_tests/scipy_tests/ndimage_tests/test_interpolation.py::TestZoomOutputSize1::test_zoom_output_size1[_param_95_{grid_mode=True, mode='grid-constant', order=5, shape=(5, 5, 2), zoom=(2, 2, 0.5)}] FAILED tests/cupyx_tests/scipy_tests/signal_tests/test_signaltools.py::TestFFTConvolve::test_fftconvolve[_param_44_{mode='valid', size1=(3, 4, 10), size2=3}] FAILED tests/cupyx_tests/scipy_tests/sparse_tests/test_linalg.py::TestLOBPCG::test_small_generate_input_for_mikota_pair FAILED tests/cupyx_tests/scipy_tests/sparse_tests/test_linalg.py::TestLOBPCG::test_maxit_None = 96 failed, 84058 passed, 8582 skipped, 661 deselected, 4520 xfailed, 1692 warnings in 28639.36s (7:57:19) =