[Tutorial] Use OpenCL on ARM Board

[tqchen]

While http://docs.tvmlang.org/tutorials/deployment/cross_compilation_and_rpc.html#sphx-glr-tutorials-deployment-cross-compilation-and-rpc-py covers how to cross compile and deploy CPU code on ARM board, it did not touch how to do so for the OpenCL code.

We have unittest examples to demonstrate that https://github.com/dmlc/tvm/blob/master/tests/python/unittest/test_runtime_rpc.py#L92 already. It would be nice to translate these lessons into the cross compilation tutorial

[ma-xiaodong]

Here is an example of using OpenCL on ARM Board.
https://github.com/ma-xiaodong/tvm_tuning/blob/master/02_conv_order.py
which corresponding the program
http://docs.tvmlang.org/tutorials/optimize/opt_conv_cuda.html#sphx-glr-tutorials-optimize-opt-conv-cuda-py
of optimizing conv on cuda.

[kaishijeng]

ma-xiaodong:
Are you to tune Mali-T860 opencl to pass test_topi_conv2d_nchw.py unittest in topi/tests/python directory?

I always encounter an error because of insufficient resources on a firefly3399 platform running Ubuntu 16.04
#530

Thanks,

[ma-xiaodong]

I also encountered the error CL_OUT_OF_RESOURCES.
In order to use opencl, you must split a dimension of an array and bind the results to threadIdx and blockIdx.
I guess this error occurs because three reasons:

1. You split a inner dimension instead of an outer dimension.
2. You split a dimension whose length is short.
3. Maybe the split factor is too big.

I suggest:

1. Split the out dimension first
2. Split dimension with big length (you can reorder the array if needet)
3. Use a small factor(4, or 8, no big than 16)

[ma-xiaodong]

#530

[kaishijeng]

In order to use opencl, you must split a dimension of an array and bind the results to threadIdx and >>blockIdx.

Which file do I need to change threadidx and blockidx? Is there a program which can tune to get optimal values of threadidx and blockidx?

[ma-xiaodong]

Which file do I need to change threadidx and blockidx? Is there a program which can tune to get optimal values of threadidx and blockidx?

Here is an example of using OpenCL on ARM Board.
https://github.com/ma-xiaodong/tvm_tuning/blob/master/02_conv_order.py

[kaishijeng]

@ma-xiaodong

Why are you not using the following conv2d modules

https://github.com/dmlc/tvm/blob/master/topi/python/topi/cuda/conv2d_nchw.py
https://github.com/dmlc/tvm/blob/master/topi/python/topi/cuda/conv2d_hwcn.py

conv2d_nchw.py looks more complex than your 02_conv_order.py.
My firefly 3399 has a problem to pass conv2d_nchw.py due to insufficient resources.

Thanks,

[ndcuong91]

@ma-xiaodong i couldn't access to your file: https://github.com/ma-xiaodong/tvm_tuning/blob/master/02_conv_order.py, can you update the link?

[ma-xiaodong]

@titikid My example is too old, so I remove it. TVM already has good example.
https://github.com/dmlc/tvm/blob/master/topi/python/topi/cuda/conv2d_hwcn.py

[ndcuong91]

@ma-xiaodong i see. Anyway, after use smaller split factor in conv2d_nchw.py as your hint i could solve my issue related to openCL. Thank you