kiritigowda commented 2 years ago

List of issues observed - List below to reproduce

caffe-resnet152 - Batch size 16 - HIP Backend on Ubuntu 20.04

caffe-resnet152 - Batch size 16
OK: loaded 39 kernels from libvx_nn.so
OK: OpenVX using GPU device#0 AMD Radeon VII (gfx906:sramecc+:xnack-) (with 60 CUs) on PCI bus 0b:00.0
clang-14: error: cannot specify -o when generating multiple output files
MIOpen Error: /home/kiriti/mivisionx-deps/MIOpen-2.14.0/src/tmp_dir.cpp:45: Can't execute cd /tmp/miopen-static_kernel_gridwise_convolution_implicit_gemm_v4r1_nchw_kcyx_nkhw_lds_double_buffer.cpp-50bb-d1f8-7c9e-84f8;  /opt/rocm/llvm/bin/clang++  -DCK_PARAM_PROBLEM_N=16 -DCK_PARAM_PROBLEM_K=128 -DCK_PARAM_PROBLEM_C=256 -DCK_PARAM_PROBLEM_HI=56 -DCK_PARAM_PROBLEM_WI=56 -DCK_PARAM_PROBLEM_HO=28 -DCK_PARAM_PROBLEM_WO=28 -DCK_PARAM_PROBLEM_Y=1 -DCK_PARAM_PROBLEM_X=1 -DCK_PARAM_PROBLEM_CONV_STRIDE_H=2 -DCK_PARAM_PROBLEM_CONV_STRIDE_W=2 -DCK_PARAM_PROBLEM_CONV_DILATION_H=1 -DCK_PARAM_PROBLEM_CONV_DILATION_W=1 -DCK_PARAM_PROBLEM_LEFT_PAD_H=0 -DCK_PARAM_PROBLEM_LEFT_PAD_W=0 -DCK_PARAM_PROBLEM_RIGHT_PAD_H=0 -DCK_PARAM_PROBLEM_RIGHT_PAD_W=0 -DCK_PARAM_PROBLEM_CONV_GROUP_COUNTS=1 -DCK_PARAM_PROBLEM_CONV_DIRECTION_FORWARD=1 -DCK_PARAM_PROBLEM_CONV_DIRECTION_BACKWARD_DATA=0 -DCK_PARAM_PROBLEM_CONV_DIRECTION_BACKWARD_WEIGHT=0 -DCK_PARAM_TUNABLE_BLOCK_SIZE=256 -DCK_PARAM_TUNABLE_B_PER_BLOCK=16 -DCK_PARAM_TUNABLE_K_PER_BLOCK=128 -DCK_PARAM_TUNABLE_E_PER_BLOCK=8 -DCK_PARAM_DEPENDENT_GRID_SIZE=98 -DCK_PARAM_GEMM_N_REPEAT=2 -DCK_PARAM_GEMM_M_PER_THREAD_SUB_C=4 -DCK_PARAM_GEMM_N_PER_THREAD_SUB_C=4 -DCK_PARAM_GEMM_M_LEVEL0_CLUSTER=4 -DCK_PARAM_GEMM_N_LEVEL0_CLUSTER=4 -DCK_PARAM_GEMM_M_LEVEL1_CLUSTER=4 -DCK_PARAM_GEMM_N_LEVEL1_CLUSTER=4 -DCK_PARAM_IN_BLOCK_COPY_CLUSTER_LENGTHS_E=8 -DCK_PARAM_IN_BLOCK_COPY_CLUSTER_LENGTHS_N1=2 -DCK_PARAM_IN_BLOCK_COPY_CLUSTER_LENGTHS_B=16 -DCK_PARAM_IN_BLOCK_COPY_CLUSTER_LENGTHS_N2=1 -DCK_PARAM_IN_BLOCK_COPY_SRC_DATA_PER_READ_B=1 -DCK_PARAM_WEI_BLOCK_COPY_CLUSTER_LENGTHS_E=2 -DCK_PARAM_WEI_BLOCK_COPY_CLUSTER_LENGTHS_K=128 -DCK_PARAM_WEI_BLOCK_COPY_SRC_DATA_PER_READ_E=1 -DCK_PARAM_EPACK_LENGTH=2 -DCK_THREADWISE_GEMM_USE_AMD_INLINE_ASM=1 -DCK_USE_AMD_INLINE_ASM=1 --std=c++14 -DCK_USE_AMD_BUFFER_ATOMIC_FADD=0 -DCK_BLOCK_SYNC_LDS_WITHOUT_SYNC_VMEM=1 -DCK_WORKAROUND_SWDEV_229564=1 -DCK_WORKAROUND_SWDEV_231101=1 -DCK_USE_AMD_BUFFER_ADDRESSING=1 -DCK_USE_AMD_V_FMAC_F32=0 -DMIOPEN_USE_FP16=1 -DMIOPEN_USE_FP32=0 -DMIOPEN_USE_INT8=0 -DMIOPEN_USE_INT8x4=0 -DMIOPEN_USE_BFP16=0 -DMIOPEN_USE_INT32=0 -DMIOPEN_USE_RNE_BFLOAT16=1 -DCK_PARAM_IN_BLOCK_COPY_DST_DATA_PER_WRITE_EPACK=2 -DCK_PARAM_WEI_BLOCK_COPY_DST_DATA_PER_WRITE_EPACK=2 -mcpu=gfx906 -Wno-everything --cuda-gpu-arch=gfx906:xnack- --cuda-device-only -c -O3  -Wno-unused-command-line-argument -I. -mllvm -amdgpu-early-inline-all=true -mllvm -amdgpu-function-calls=false -x hip -isystem /opt/rocm-5.0.0/hip/../include -isystem /opt/rocm/llvm/lib/clang/14.0.0/include/.. -D__HIP_PLATFORM_HCC__=1 -D__HIP_PLATFORM_AMD__=1 -isystem /opt/rocm-5.0.0/hip/include -isystem /opt/rocm/include /opt/rocm/llvm/lib/clang/14.0.0/lib/linux/libclang_rt.builtins-x86_64.a --hip-link /opt/rocm/llvm/lib/clang/14.0.0/lib/linux/libclang_rt.builtins-x86_64.a -mllvm --amdgpu-spill-vgpr-to-agpr=0 -DHIP_PACKAGE_VERSION_FLAT=5000022042 static_kernel_gridwise_convolution_implicit_gemm_v4r1_nchw_kcyx_nkhw_lds_double_buffer.cpp -o /tmp/miopen-static_kernel_gridwise_convolution_implicit_gemm_v4r1_nchw_kcyx_nkhw_lds_double_buffer.cpp-50bb-d1f8-7c9e-84f8/static_kernel_gridwise_convolution_implicit_gemm_v4r1_nchw_kcyx_nkhw_lds_double_buffer.cpp.o
ERROR: fatal error occured at /home/kiriti/develop/profile/MIVisionX/amd_openvx_extensions/amd_nn/src/convolution_layer.cpp#416

caffe-xyznet - Batch size 32 - HIP Backend on Ubuntu 20.04

caffe-xyznet - Batch size 32
OK: loaded 39 kernels from libvx_nn.so
OK: OpenVX using GPU device#0 AMD Radeon VII (gfx906:sramecc+:xnack-) (with 60 CUs) on PCI bus 0b:00.0
ERROR: hipMalloc Failed with status: 2
ERROR: agoGpuHipAllocBuffer(MEM_READ_WRITE,872415232,0,*) FAILED
ERROR: vxVerifyGraph(...) failed (-1)

caffe-yolov3 - Batch size 2 - HIP Backend on Ubuntu 20.04

caffe-yolov3 - Batch size 2
OK: loaded 39 kernels from libvx_nn.so
OK: OpenVX using GPU device#0 AMD Radeon VII (gfx906:sramecc+:xnack-) (with 60 CUs) on PCI bus 0b:00.0
clang-14: error: cannot specify -o when generating multiple output files
clang-14: error: cannot specify -o when generating multiple output files
terminate called after throwing an instance of 'miopen::Exception'
terminate called recursively
what():  /home/kiriti/mivisionx-deps/MIOpen-2.14.0/src/tmp_dir.cpp:45: Can't execute cd /tmp/miopen-convolution_forward_implicit_gemm_v6r1_dlops_nchw_kcyx_nkhw.cpp-6017-d909-909a-0c8f;  /opt/rocm/llvm/bin/clang++  -DCK_PARAM_ABDataTypeEnum=1 -DCK_PARAM_AccDataTypeEnum=1 -DCK_PARAM_CDataTypeEnum=1 -DCK_PARAM_BlockSize=128 -DCK_PARAM_GN0=1 -DCK_PARAM_GK1=1 -DCK_PARAM_GM1PerBlockGM11=64 -DCK_PARAM_GN1PerBlockGN11=128 -DCK_PARAM_GK0PerBlock=8 -DCK_PARAM_BM1PerThreadBM11=4 -DCK_PARAM_BN1PerThreadBN11=4 -DCK_PARAM_BK0PerThread=1 -DCK_PARAM_BM10BN10ThreadClusterBM10Xs=4,2 -DCK_PARAM_BM10BN10ThreadClusterBN10Xs=8,2 -DCK_PARAM_ABlockTransferThreadSliceLengths_GK0_GM0_GM10_GM11_GK1=4,1,1,1,1 -DCK_PARAM_ABlockTransferThreadClusterLengths_GK0_GM0_GM10_GM11_GK1=2,1,1,64,1 -DCK_PARAM_ABlockTransferSrcVectorTensorLengths_GK0_GM0_GM10_GM11_GK1=4,1,1,1,1 -DCK_PARAM_ABlockTransferDstVectorTensorLengths_GK0_GM0_GM10_GM11_GK1=1,1,1,1,1 -DCK_PARAM_BBlockTransferThreadSliceLengths_GK0_GN0_GN10_GN11_GK1=8,1,1,1,1 -DCK_PARAM_BBlockTransferThreadClusterLengths_GK0_GN0_GN10_GN11_GK1=1,1,1,128,1 -DCK_PARAM_BBlockTransferSrcVectorTensorLengths_GK0_GN0_GN10_GN11_GK1=1,1,1,1,1 -DCK_PARAM_BBlockTransferDstVectorTensorLengths_GK0_GN0_GN10_GN11_GK1=1,1,1,1,1 -DCK_PARAM_CThreadTransferDstScalarPerVector=4 -DCK_PARAM_HasMainKBlockLoop=1 -DCK_PARAM_HasDoubleTailKBlockLoop=1 --std=c++17 -DCK_AMD_GPU_GFX906 -DCK_USE_AMD_BUFFER_ATOMIC_FADD=0 -DCK_BLOCK_SYNC_LDS_WITHOUT_SYNC_VMEM=1 -DCK_USE_AMD_BUFFER_ADDRESSING=1 -mcpu=gfx906 -Wno-everything --cuda-gpu-arch=gfx906:xnack- --cuda-device-only -c -O3  -Wno-unused-command-line-argument -I. -mllvm -amdgpu-early-inline-all=true -mllvm -amdgpu-function-calls=false -x hip -isystem /opt/rocm-5.0.0/hip/../include -isystem /opt/rocm/llvm/lib/clang/14.0.0/include/.. -D__HIP_PLATFORM_HCC__=1 -D__HIP_PLATFORM_AMD__=1 -isystem /opt/rocm-5.0.0/hip/include -isystem /opt/rocm/include /opt/rocm/llvm/lib/clang/14.0.0/lib/linux/libclang_rt.builtins-x86_64.a --hip-link /opt/rocm/llvm/lib/clang/14.0.0/lib/linux/libclang_rt.builtins-xAborted (core dumped)

caffe-yolov3 - Batch size 4 - HIP Backend on Ubuntu 20.04

caffe-yolov3 - Batch size 4
OK: loaded 39 kernels from libvx_nn.so
OK: OpenVX using GPU device#0 AMD Radeon VII (gfx906:sramecc+:xnack-) (with 60 CUs) on PCI bus 0b:00.0
OK: graph initialization with annAddToGraph() took 761.874 msec
ERROR: kernel com.amd.nn_extension.upsample_nearest_layer exec failed (-2:(null))
ERROR: vxProcessGraph() failed (-1)

kiritigowda commented 2 years ago

@hansely - Please look into kernel com.amd.nn_extension.upsample_nearest_layer exec failed (-2:(null)) with YoloV3

kiritigowda commented 2 years ago

@paveltc -- need to track these nn issues.

LakshmiKumar23 commented 2 years ago

@kiritigowda I thought the clang14 issue was fixed on develop? I still see it with onnx squeezenet. Can you please check?

paveltc commented 2 years ago

@kiritigowda I can repro the issue with resnet152, will try xyznet next.

paveltc commented 2 years ago

@kiritigowda Can also reproduce with xyznet.

kiritigowda commented 2 years ago

@hansely can you verify if YoloV3 works for all batch sizes with TOT HIP/OpenCL

kiritigowda commented 2 years ago

Error Updates

caffe-resnet152 - Batch size 16 - HIP Backend on Ubuntu 20.04 -- FIX

caffe-resnet152 - Batch size 16
OK: loaded 39 kernels from libvx_nn.so
OK: OpenVX using GPU device#0 AMD Radeon PRO W6800 (gfx1030) (with 30 CUs) on PCI bus 0b:00.0
OK: graph initialization with annAddToGraph() took 430.691 msec
OK: vxProcessGraph() took 7756.354 msec (1st iteration)
OK: vxProcessGraph() took 44.986 msec (average over 100 iterations)
OK: HIP buffer usage: 430905504, 783/783
OK: successful

caffe-xyznet - Batch size 32 & above - HIP Backend on Ubuntu 20.04 -- maybe lack of memory on the board -- @asalmanp can you confirm this and reopen if this is something else.

caffe-xyznet - Batch size 32
OK: loaded 39 kernels from libvx_nn.so
OK: OpenVX using GPU device#0 AMD Radeon PRO W6800 (gfx1030) (with 30 CUs) on PCI bus 0b:00.0
ERROR: hipMalloc Failed with status: 2
ERROR: agoGpuHipAllocBuffer(MEM_READ_WRITE,805306368,0,*) FAILED
ERROR: vxVerifyGraph(...) failed (-1)

caffe-yolov3 - Batch size 2 - HIP Backend on Ubuntu 20.04 -- FIX

caffe-yolov3 - Batch size 2
OK: loaded 39 kernels from libvx_nn.so
OK: OpenVX using GPU device#0 AMD Radeon PRO W6800 (gfx1030) (with 30 CUs) on PCI bus 0b:00.0
OK: graph initialization with annAddToGraph() took 8286.610 msec
OK: vxProcessGraph() took 58.715 msec (1st iteration)
OK: vxProcessGraph() took 12.393 msec (average over 100 iterations)
OK: HIP buffer usage: 389538444, 376/376
OK: successful

caffe-yolov3 - Batch size 4 - HIP Backend on Ubuntu 20.04 -- FIX

caffe-yolov3 - Batch size 4
OK: loaded 39 kernels from libvx_nn.so
OK: OpenVX using GPU device#0 AMD Radeon PRO W6800 (gfx1030) (with 30 CUs) on PCI bus 0b:00.0
OK: graph initialization with annAddToGraph() took 8493.686 msec
OK: vxProcessGraph() took 66.456 msec (1st iteration)
OK: vxProcessGraph() took 20.852 msec (average over 100 iterations)
OK: HIP buffer usage: 531069860, 376/376
OK: successful

@hansely @LakshmiKumar23 @paveltc -- TOT master with latest deps is verified -- Need to add this to regular review.

ROCm / MIVisionX

VX_NN - Inference Issues #802

Error Updates