In topi schedules, for OpenCL target, num_thread == 512 is too big for AMD GPU

[masahi]

Hi, I noticed that some schedules in topi.cuda use thread block of size 512 in one dimension.
For example, here and here.

512 threads in one dimension is too big for AMD OpenCL, because their driver restricts the maximum blockDim.x to be 256. According to the discussion here, this restriction is due to their implementation of driver, not their hardware itself.

I understand that schedules in topi/cuda are optimized for CUDA, but if I change the 512 threads to 256 threads, all tests in https://github.com/dmlc/tvm/tree/master/topi/tests/python pass for my AMD gpu as well. I was also able to run tutorials and resnet benchmark from the latest nnvm, on AMD gpu.

So, unless there is a strong reason to use thread block of size 512, is it possible to change 512 to 256?
Of course, if you are planning to implement dedicated schedules for opencl, that's better :)

[tqchen]

I agree that we should have dedicated schedules for devices(e.g. AMD). We will do a upgrade of the topi interface to make plugin of device-specific schedule easy

[kaishijeng]

On firefly3399, I can pass all unittests except test_topi_conv2d_nchw.py which complains running out of resources. I probably need to reduce # of threads to a smaller number.

Just curious, are you able to run NNVM benchmark with OpenCL on amd platform?
NNVM benchmark has 2 benchmarks, one is for rasp and the other one is for cuda. There is no Opencl benchmark

Thanks,

[masahi]

@kaishijeng yes, I modified cuda bench to use opencl target and context. I also modified nnvm internal to enable opencl target.
Proof:
$python cuda_imagenet_bench.py --model resnet
[09:38:49] src/runtime/opencl/opencl_device_api.cc:195: Initialize OpenCL platform 'AMD Accelerated Parallel Processing'
[09:38:49] src/runtime/opencl/opencl_device_api.cc:215: opencl(0)='gfx803' cl_device_id=0x7f85d4dc46f0
benchmark args: Namespace(model='resnet', num_iter=100, opt_level=3)
ProfileResult(mean=0.01073681044)
ProfileResult(mean=0.01073411568)
ProfileResult(mean=0.01073584366)

I'm not sure if performance is good or bad. My card is R9 Nano, with 8 TFLOPs of peak perf.

For test_topi_conv2d_nchw, you can use my modified schedule here

The fact that I can load PyTorch or MXNet - trained model and run inference with it on AMD gpu is pretty exciting. Plus, my opencl driver and tool chains are fully open source with rocm.

[masahi]

To enable opencl target in nnvm, I changed

if target == "cuda":
to
if target == "cuda" or target == "opencl":

in
python/nnvm/top/nn.py
python/nnvm/top/reduction.py
python/nnvm/top/tensor.py

[kaishijeng]

With your modified schedule, I still got out of resource error below. I guess firefly3399 has smaller GPU than your amd. Do you mind to share your modifed cuda_imagenet_bench.py and how you modify nnvm internal to enable opencl target? TVM: Initializing cython mode... Use memoize topi.tests.test_topi_conv2d.verify_con2d_nchw.get_ref_data.pkl(5, (1, 3, 224, 224), 'float32', 2, 3, (64, 3, 7, 7)) Skip because cuda is not enabled [17:55:55] src/runtime/opencl/opencl_device_api.cc:195: Initialize OpenCL platform 'ARM Platform' [17:55:55] src/runtime/opencl/opencl_device_api.cc:215: opencl(0)='Mali-T860' cl_device_id=0x7f93fa8dd8 [17:55:55] /home/firefly/2TB/src/firefly/tvm/dmlc-core/include/dmlc/logging.h:308: [17:55:55] src/runtime/opencl/opencl_module.cc:223: Check failed: e == CL_SUCCESS OpenCL Error, code=-5: CL_OUT_OF_RESOURCES Stack trace returned 10 entries: [bt] (0) /usr/local/lib/python2.7/dist-packages/tvm-0.1.0-py2.7-linux-aarch64.egg/tvm/libtvm.so(_ZN4dmlc15LogMessageFatalD1Ev+0x44) [0x7f944bc6e4] [bt] (1) /usr/local/lib/python2.7/dist-packages/tvm-0.1.0-py2.7-linux-aarch64.egg/tvm/libtvm.so(_ZNK3tvm7runtime17OpenCLWrappedFuncclENS0_7TVMArgsEPNS0_11TVMRetValueEPPv+0xbac) [0x7f94854b9c] [bt] (2) /usr/local/lib/python2.7/dist-packages/tvm-0.1.0-py2.7-linux-aarch64.egg/tvm/libtvm.so(_ZNSt17_Function_handlerIFvN3tvm7runtime7TVMArgsEPNS1_11TVMRetValueEEZNS1_6detail17PackFuncVoidAddr_ILi4ENS1_17OpenCLWrappedFuncEEENS1_10PackedFuncET0_RKSt6vectorINS6_14ArgConvertCodeESaISC_EEEUlS2_S4_E_E9_M_invokeERKSt9_Any_dataOS2_OS4_+0xd0) [0x7f94854f10] [bt] (3) /usr/local/lib/python2.7/dist-packages/tvm-0.1.0-py2.7-linux-aarch64.egg/tvm/libtvm.so(TVMFuncCall+0x74) [0x7f9482951c] [bt] (4) [0x7f9676a9dc] [bt] (5) [0x7f9676a2ac] [bt] (6) /usr/local/lib/python2.7/dist-packages/tvm-0.1.0-py2.7-linux-aarch64.egg/tvm/libtvm.so(_ZZN3tvm7codegen14LLVMModuleNode11GetFunctionERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEERKSt10shared_ptrINS_7runtime10ModuleNodeEEENKUlNSB_7TVMArgsEPNSB_11TVMRetValueEE0_clESG_SI_+0x3c) [0x7f946f2a3c] [bt] (7) /usr/local/lib/python2.7/dist-packages/tvm-0.1.0-py2.7-linux-aarch64.egg/tvm/libtvm.so(_ZNSt17_Function_handlerIFvN3tvm7runtime7TVMArgsEPNS1_11TVMRetValueEEZNS0_7codegen14LLVMModuleNode11GetFunctionERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEERKSt10shared_ptrINS1_10ModuleNodeEEEUlS2_S4_E0_E9_M_invokeERKSt9_Any_dataOS2_OS4_+0x2c) [0x7f946f2cfc] [bt] (8) /usr/local/lib/python2.7/dist-packages/tvm-0.1.0-py2.7-linux-aarch64.egg/tvm/libtvm.so(TVMFuncCall+0x74) [0x7f9482951c] [bt] (9) /usr/local/lib/python2.7/dist-packages/tvm-0.1.0-py2.7-linux-aarch64.egg/tvm/_ffi/_cy2/core.so(+0x180f8) [0x7f8a5280f8] [17:55:55] /home/firefly/2TB/src/firefly/tvm/dmlc-core/include/dmlc/logging.h:308: [17:55:55] src/codegen/llvm/llvm_module.cc:59: Check failed: ret == 0 (-1 vs. 0) [17:55:55] src/runtime/opencl/opencl_module.cc:223: Check failed: e == CL_SUCCESS OpenCL Error, code=-5: CL_OUT_OF_RESOURCES Stack trace returned 10 entries: [bt] (0) /usr/local/lib/python2.7/dist-packages/tvm-0.1.0-py2.7-linux-aarch64.egg/tvm/libtvm.so(_ZN4dmlc15LogMessageFatalD1Ev+0x44) [0x7f944bc6e4] [bt] (1) /usr/local/lib/python2.7/dist-packages/tvm-0.1.0-py2.7-linux-aarch64.egg/tvm/libtvm.so(_ZNK3tvm7runtime17OpenCLWrappedFuncclENS0_7TVMArgsEPNS0_11TVMRetValueEPPv+0xbac) [0x7f94854b9c] [bt] (2) /usr/local/lib/python2.7/dist-packages/tvm-0.1.0-py2.7-linux-aarch64.egg/tvm/libtvm.so(_ZNSt17_Function_handlerIFvN3tvm7runtime7TVMArgsEPNS1_11TVMRetValueEEZNS1_6detail17PackFuncVoidAddr_ILi4ENS1_17OpenCLWrappedFuncEEENS1_10PackedFuncET0_RKSt6vectorINS6_14ArgConvertCodeESaISC_EEEUlS2_S4_E_E9_M_invokeERKSt9_Any_dataOS2_OS4_+0xd0) [0x7f94854f10] [bt] (3) /usr/local/lib/python2.7/dist-packages/tvm-0.1.0-py2.7-linux-aarch64.egg/tvm/libtvm.so(TVMFuncCall+0x74) [0x7f9482951c] [bt] (4) [0x7f9676a9dc] [bt] (5) [0x7f9676a2ac] [bt] (6) /usr/local/lib/python2.7/dist-packages/tvm-0.1.0-py2.7-linux-aarch64.egg/tvm/libtvm.so(_ZZN3tvm7codegen14LLVMModuleNode11GetFunctionERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEERKSt10shared_ptrINS_7runtime10ModuleNodeEEENKUlNSB_7TVMArgsEPNSB_11TVMRetValueEE0_clESG_SI_+0x3c) [0x7f946f2a3c] [bt] (7) /usr/local/lib/python2.7/dist-packages/tvm-0.1.0-py2.7-linux-aarch64.egg/tvm/libtvm.so(_ZNSt17_Function_handlerIFvN3tvm7runtime7TVMArgsEPNS1_11TVMRetValueEEZNS0_7codegen14LLVMModuleNode11GetFunctionERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEERKSt10shared_ptrINS1_10ModuleNodeEEEUlS2_S4_E0_E9_M_invokeERKSt9_Any_dataOS2_OS4_+0x2c) [0x7f946f2cfc] [bt] (8) /usr/local/lib/python2.7/dist-packages/tvm-0.1.0-py2.7-linux-aarch64.egg/tvm/libtvm.so(TVMFuncCall+0x74) [0x7f9482951c] [bt] (9) /usr/local/lib/python2.7/dist-packages/tvm-0.1.0-py2.7-linux-aarch64.egg/tvm/_ffi/_cy2/core.so(+0x180f8) [0x7f8a5280f8] Stack trace returned 6 entries: [bt] (0) /usr/local/lib/python2.7/dist-packages/tvm-0.1.0-py2.7-linux-aarch64.egg/tvm/libtvm.so(_ZN4dmlc15LogMessageFatalD1Ev+0x44) [0x7f944bc6e4] [bt] (1) /usr/local/lib/python2.7/dist-packages/tvm-0.1.0-py2.7-linux-aarch64.egg/tvm/libtvm.so(_ZZN3tvm7codegen14LLVMModuleNode11GetFunctionERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEERKSt10shared_ptrINS_7runtime10ModuleNodeEEENKUlNSB_7TVMArgsEPNSB_11TVMRetValueEE0_clESG_SI_+0x1f4) [0x7f946f2bf4] [bt] (2) /usr/local/lib/python2.7/dist-packages/tvm-0.1.0-py2.7-linux-aarch64.egg/tvm/libtvm.so(_ZNSt17_Function_handlerIFvN3tvm7runtime7TVMArgsEPNS1_11TVMRetValueEEZNS0_7codegen14LLVMModuleNode11GetFunctionERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEERKSt10shared_ptrINS1_10ModuleNodeEEEUlS2_S4_E0_E9_M_invokeERKSt9_Any_dataOS2_OS4_+0x2c) [0x7f946f2cfc] [bt] (3) /usr/local/lib/python2.7/dist-packages/tvm-0.1.0-py2.7-linux-aarch64.egg/tvm/libtvm.so(TVMFuncCall+0x74) [0x7f9482951c] [bt] (4) /usr/local/lib/python2.7/dist-packages/tvm-0.1.0-py2.7-linux-aarch64.egg/tvm/_ffi/_cy2/core.so(+0x180f8) [0x7f8a5280f8] [bt] (5) python2(PyObject_Call+0x50) [0x458ee8] Traceback (most recent call last): File "./test_topi_conv2d_nchw.py", line 72, in <module> test_conv2d_nchw() File "./test_topi_conv2d_nchw.py", line 58, in test_conv2d_nchw verify_conv2d_nchw(1, 3, 224, 64, 7, 3, 2) File "./test_topi_conv2d_nchw.py", line 54, in verify_conv2d_nchw check_device(device) File "./test_topi_conv2d_nchw.py", line 48, in check_device func1(a, w, b) File "/usr/local/lib/python2.7/dist-packages/tvm-0.1.0-py2.7-linux-aarch64.egg/tvm/_ffi/function.py", line 128, in __call__ return f(*args) File "tvm/_ffi/_cython/./function.pxi", line 264, in core.FunctionBase.__call__ File "tvm/_ffi/_cython/./function.pxi", line 213, in core.FuncCall File "tvm/_ffi/_cython/./function.pxi", line 205, in core.FuncCall3 File "tvm/_ffi/_cython/./base.pxi", line 131, in core.CALL tvm._ffi.base.TVMError: [17:55:55] src/codegen/llvm/llvm_module.cc:59: Check failed: ret == 0 (-1 vs. 0) [17:55:55] src/runtime/opencl/opencl_module.cc:223: Check failed: e == CL_SUCCESS OpenCL Error, code=-5: CL_OUT_OF_RESOURCES Stack trace returned 10 entries: [bt] (0) /usr/local/lib/python2.7/dist-packages/tvm-0.1.0-py2.7-linux-aarch64.egg/tvm/libtvm.so(_ZN4dmlc15LogMessageFatalD1Ev+0x44) [0x7f944bc6e4] [bt] (1) /usr/local/lib/python2.7/dist-packages/tvm-0.1.0-py2.7-linux-aarch64.egg/tvm/libtvm.so(_ZNK3tvm7runtime17OpenCLWrappedFuncclENS0_7TVMArgsEPNS0_11TVMRetValueEPPv+0xbac) [0x7f94854b9c] [bt] (2) /usr/local/lib/python2.7/dist-packages/tvm-0.1.0-py2.7-linux-aarch64.egg/tvm/libtvm.so(_ZNSt17_Function_handlerIFvN3tvm7runtime7TVMArgsEPNS1_11TVMRetValueEEZNS1_6detail17PackFuncVoidAddr_ILi4ENS1_17OpenCLWrappedFuncEEENS1_10PackedFuncET0_RKSt6vectorINS6_14ArgConvertCodeESaISC_EEEUlS2_S4_E_E9_M_invokeERKSt9_Any_dataOS2_OS4_+0xd0) [0x7f94854f10] [bt] (3) /usr/local/lib/python2.7/dist-packages/tvm-0.1.0-py2.7-linux-aarch64.egg/tvm/libtvm.so(TVMFuncCall+0x74) [0x7f9482951c] [bt] (4) [0x7f9676a9dc] [bt] (5) [0x7f9676a2ac] [bt] (6) /usr/local/lib/python2.7/dist-packages/tvm-0.1.0-py2.7-linux-aarch64.egg/tvm/libtvm.so(_ZZN3tvm7codegen14LLVMModuleNode11GetFunctionERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEERKSt10shared_ptrINS_7runtime10ModuleNodeEEENKUlNSB_7TVMArgsEPNSB_11TVMRetValueEE0_clESG_SI_+0x3c) [0x7f946f2a3c] [bt] (7) /usr/local/lib/python2.7/dist-packages/tvm-0.1.0-py2.7-linux-aarch64.egg/tvm/libtvm.so(_ZNSt17_Function_handlerIFvN3tvm7runtime7TVMArgsEPNS1_11TVMRetValueEEZNS0_7codegen14LLVMModuleNode11GetFunctionERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEERKSt10shared_ptrINS1_10ModuleNodeEEEUlS2_S4_E0_E9_M_invokeERKSt9_Any_dataOS2_OS4_+0x2c) [0x7f946f2cfc] [bt] (8) /usr/local/lib/python2.7/dist-packages/tvm-0.1.0-py2.7-linux-aarch64.egg/tvm/libtvm.so(TVMFuncCall+0x74) [0x7f9482951c] [bt] (9) /usr/local/lib/python2.7/dist-packages/tvm-0.1.0-py2.7-linux-aarch64.egg/tvm/_ffi/_cy2/core.so(+0x180f8) [0x7f8a5280f8] Stack trace returned 6 entries: [bt] (0) /usr/local/lib/python2.7/dist-packages/tvm-0.1.0-py2.7-linux-aarch64.egg/tvm/libtvm.so(_ZN4dmlc15LogMessageFatalD1Ev+0x44) [0x7f944bc6e4] [bt] (1) /usr/local/lib/python2.7/dist-packages/tvm-0.1.0-py2.7-linux-aarch64.egg/tvm/libtvm.so(_ZZN3tvm7codegen14LLVMModuleNode11GetFunctionERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEERKSt10shared_ptrINS_7runtime10ModuleNodeEEENKUlNSB_7TVMArgsEPNSB_11TVMRetValueEE0_clESG_SI_+0x1f4) [0x7f946f2bf4] [bt] (2) /usr/local/lib/python2.7/dist-packages/tvm-0.1.0-py2.7-linux-aarch64.egg/tvm/libtvm.so(_ZNSt17_Function_handlerIFvN3tvm7runtime7TVMArgsEPNS1_11TVMRetValueEEZNS0_7codegen14LLVMModuleNode11GetFunctionERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEERKSt10shared_ptrINS1_10ModuleNodeEEEUlS2_S4_E0_E9_M_invokeERKSt9_Any_dataOS2_OS4_+0x2c) [0x7f946f2cfc] [bt] (3) /usr/local/lib/python2.7/dist-packages/tvm-0.1.0-py2.7-linux-aarch64.egg/tvm/libtvm.so(TVMFuncCall+0x74) [0x7f9482951c] [bt] (4) /usr/local/lib/python2.7/dist-packages/tvm-0.1.0-py2.7-linux-aarch64.egg/tvm/_ffi/_cy2/core.so(+0x180f8) [0x7f8a5280f8] [bt] (5) python2(PyObject_Call+0x50) [0x458ee8] [17:55:55] /home/firefly/2TB/src/firefly/tvm/dmlc-core/include/dmlc/logging.h:308: [17:55:55] src/runtime/workspace_pool.cc:92: Check failed: allocated_.size() == 1 (3 vs. 1) Stack trace returned 4 entries: [bt] (0) /usr/local/lib/python2.7/dist-packages/tvm-0.1.0-py2.7-linux-aarch64.egg/tvm/libtvm.so(_ZN3tvm7runtime13WorkspacePool4Pool7ReleaseE9DLContextPNS0_9DeviceAPIE+0x488) [0x7f94838fc0] [bt] (1) /usr/local/lib/python2.7/dist-packages/tvm-0.1.0-py2.7-linux-aarch64.egg/tvm/libtvm.so(_ZN3tvm7runtime13WorkspacePoolD1Ev+0x48) [0x7f94837710] [bt] (2) /usr/local/lib/python2.7/dist-packages/tvm-0.1.0-py2.7-linux-aarch64.egg/tvm/libtvm.so(_ZN3tvm7runtime2cl17OpenCLThreadEntryD1Ev+0x18) [0x7f9484c818] [bt] (3) /lib/aarch64-linux-gnu/libc.so.6(__call_tls_dtors+0x44) [0x7f96985474] terminate called after throwing an instance of 'dmlc::Error' what(): [17:55:55] src/runtime/workspace_pool.cc:92: Check failed: allocated_.size() == 1 (3 vs. 1) Stack trace returned 4 entries: [bt] (0) /usr/local/lib/python2.7/dist-packages/tvm-0.1.0-py2.7-linux-aarch64.egg/tvm/libtvm.so(_ZN3tvm7runtime13WorkspacePool4Pool7ReleaseE9DLContextPNS0_9DeviceAPIE+0x488) [0x7f94838fc0] [bt] (1) /usr/local/lib/python2.7/dist-packages/tvm-0.1.0-py2.7-linux-aarch64.egg/tvm/libtvm.so(_ZN3tvm7runtime13WorkspacePoolD1Ev+0x48) [0x7f94837710] [bt] (2) /usr/local/lib/python2.7/dist-packages/tvm-0.1.0-py2.7-linux-aarch64.egg/tvm/libtvm.so(_ZN3tvm7runtime2cl17OpenCLThreadEntryD1Ev+0x18) [0x7f9484c818] [bt] (3) /lib/aarch64-linux-gnu/libc.so.6(__call_tls_dtors+0x44) [0x7f96985474] Aborted (core dumped)

…

On Tue, Oct 10, 2017 at 5:46 PM, masahi ***@***.***> wrote: @kaishijeng <https://github.com/kaishijeng> yes, I modified cuda bench to use opencl target and context. I also modified nnvm internal to enable opencl target. Proof: $python cuda_imagenet_bench.py --model resnet *[09:38:49] src/runtime/opencl/opencl_device_api.cc:195: Initialize OpenCL platform 'AMD Accelerated Parallel Processing' [09:38:49] src/runtime/opencl/opencl_device_api.cc:215: opencl(0)='gfx803' cl_device_id=0x7f85d4dc46f0 benchmark args: Namespace(model='resnet', num_iter=100, opt_level=3) ProfileResult(mean=0.01073681044) ProfileResult(mean=0.01073411568) ProfileResult(mean=0.01073584366)* For test_topi_conv2d_nchw, you can use my modified schedule here <https://gist.github.com/masahi/18b4ed1fb8fa1a049ffdba35ed807509> — You are receiving this because you were mentioned. Reply to this email directly, view it on GitHub <#530 (comment)>, or mute the thread <https://github.com/notifications/unsubscribe-auth/AMGg3kpr7vIftAYXZ7FlNEcyTzrgMd_Qks5srA_9gaJpZM4PzdCV> .

[tqchen]

This is mainly because AMD card and ARM GPU have different resources and requires different kind of schedules

[kaishijeng]

I change num_thread from your 256 to 32, still got out of resources which looks strange to me. I also see num_thread_x = 64 in the code, do I need to chage this variable too? Thanks,

…

On Tue, Oct 10, 2017 at 6:01 PM, Tianqi Chen ***@***.***> wrote: This is mainly because AMD card and ARM GPU have different resources and requires different kind of schedules — You are receiving this because you were mentioned. Reply to this email directly, view it on GitHub <#530 (comment)>, or mute the thread <https://github.com/notifications/unsubscribe-auth/AMGg3v0WfHr-F5Hx3fDgxcFWPist5RDSks5srBN4gaJpZM4PzdCV> .

[masahi]

@kaishijeng Sorry I don't know anything about ARM gpus. Maybe you can try smaller inputs?
Say, verify_conv2d_nchw(1, 3, 64, 16, 3, 1, 0) ?
Run clinfo (not sure if it exists in ARM platform) to see the maximum work group size.

I can push my repo, after I clean it up.

[kaishijeng]

Yes, it works OK with verify_conv2d_nchw(1, 3, 64, 64, 3, 1, 0) Thanks,

…

On Tue, Oct 10, 2017 at 6:37 PM, masahi ***@***.***> wrote: @kaishijeng <https://github.com/kaishijeng> Sorry I don't know anything about ARM gpus. Maybe you can try smaller inputs? Say, verify_conv2d_nchw(1, 3, 64, 64, 3, 1, 0) ? I can push my repo, after I clean it up. — You are receiving this because you were mentioned. Reply to this email directly, view it on GitHub <#530 (comment)>, or mute the thread <https://github.com/notifications/unsubscribe-auth/AMGg3rzJGs8Uq8nmRtLKl9f_D8_uTtiwks5srBvDgaJpZM4PzdCV> .

[masahi]

Glad to hear that.
So your CL_OUT_OF_RESOURCES is due to insufficient memory, not the amount of threads.

[kaishijeng]

Can you share your modified cuda_imagenet_bench.py to enable OpenCL? Thanks,

…

On Tue, Oct 10, 2017 at 7:04 PM, kaishi Jeng ***@***.***> wrote: Yes, it works OK with verify_conv2d_nchw(1, 3, 64, 64, 3, 1, 0) Thanks, On Tue, Oct 10, 2017 at 6:37 PM, masahi ***@***.***> wrote: > @kaishijeng <https://github.com/kaishijeng> Sorry I don't know anything > about ARM gpus. Maybe you can try smaller inputs? > Say, verify_conv2d_nchw(1, 3, 64, 64, 3, 1, 0) ? > > I can push my repo, after I clean it up. > > — > You are receiving this because you were mentioned. > Reply to this email directly, view it on GitHub > <#530 (comment)>, or mute > the thread > <https://github.com/notifications/unsubscribe-auth/AMGg3rzJGs8Uq8nmRtLKl9f_D8_uTtiwks5srBvDgaJpZM4PzdCV> > . >

[masahi]

Sure it's here

Note that I all changed is just
target = "opencl"
ctx = tvm.cl(0)

With the same change you can also run tutorial/from_onnx.py and tutorial/from_mxnet.py etc.

You can also use
python/nnvm/top/nn.py
python/nnvm/top/reduction.py
python/nnvm/top/tensor.py

This should be enough to run resnet bench and tutorials.

[kaishijeng]

I encountered CL_OUT_RESOURCES again with your changes. tvm._ffi.base.TVMError: [19:44:47] src/codegen/llvm/llvm_module.cc:59: Check failed: ret == 0 (-1 vs. 0) [19:44:47] src/runtime/opencl/opencl_module.cc:223: Check failed: e == CL_SUCCESS OpenCL Error, code=-5: CL_OUT_OF_RESOURCES

…

On Tue, Oct 10, 2017 at 7:26 PM, masahi ***@***.***> wrote: Sure it's here <https://gist.github.com/masahi/de658ec98b1e275de6331a4cf4c873b3> Note that I all changed is just *target = "opencl" ctx = tvm.cl <http://tvm.cl>(0)* With the same change you can also run tutorial/from_onnx.py and tutorial/from_mxnet.py etc. You can also use python/nnvm/top/nn.py <https://gist.github.com/masahi/a9b67fe27b24b783334ef8fc6f877ef9> python/nnvm/top/reduction.py <https://gist.github.com/masahi/706fd19cb4c589bc7ea59d68cde99a8a> python/nnvm/top/tensor.py <https://gist.github.com/masahi/1d503f9e7839636dc8086d3737523d64> This should be enough to run resnet bench and tutorials. — You are receiving this because you were mentioned. Reply to this email directly, view it on GitHub <#530 (comment)>, or mute the thread <https://github.com/notifications/unsubscribe-auth/AMGg3rLt4GKrsFxLsDkZ7bU_HeNNS6abks5srCdZgaJpZM4PzdCV> .

[masahi]

Are you running resnet bench? If so, try mobilenet bench first. It requires less memory I think. 2017/10/11 午前11:47 "kaishijeng" <notifications@github.com>:

…

I encountered CL_OUT_RESOURCES again with your changes. tvm._ffi.base.TVMError: [19:44:47] src/codegen/llvm/llvm_module.cc:59: Check failed: ret == 0 (-1 vs. 0) [19:44:47] src/runtime/opencl/opencl_module.cc:223: Check failed: e == CL_SUCCESS OpenCL Error, code=-5: CL_OUT_OF_RESOURCES On Tue, Oct 10, 2017 at 7:26 PM, masahi ***@***.***> wrote: > Sure it's here > <https://gist.github.com/masahi/de658ec98b1e275de6331a4cf4c873b3> > > Note that I all changed is just > > *target = "opencl" ctx = tvm.cl <http://tvm.cl>(0)* > > With the same change you can also run tutorial/from_onnx.py and > tutorial/from_mxnet.py etc. > > You can also use > python/nnvm/top/nn.py > <https://gist.github.com/masahi/a9b67fe27b24b783334ef8fc6f877ef9> > python/nnvm/top/reduction.py > <https://gist.github.com/masahi/706fd19cb4c589bc7ea59d68cde99a8a> > python/nnvm/top/tensor.py > <https://gist.github.com/masahi/1d503f9e7839636dc8086d3737523d64> > > This should be enough to run resnet bench and tutorials. > > — > You are receiving this because you were mentioned. > Reply to this email directly, view it on GitHub > <#530 (comment)>, or mute > the thread > <https://github.com/notifications/unsubscribe- auth/AMGg3rLt4GKrsFxLsDkZ7bU_HeNNS6abks5srCdZgaJpZM4PzdCV> > . > — You are receiving this because you authored the thread. Reply to this email directly, view it on GitHub <#530 (comment)>, or mute the thread <https://github.com/notifications/unsubscribe-auth/ABsbE1POGJTHEVVfNe-2fQ8nVIAyfVBVks5srCwwgaJpZM4PzdCV> .

[kaishijeng]

No luck with mobilenet. It is probably related to how tvm/nnvm allocates GPU resources in various layers because OpenCL functions in ARM ComputeLibrary is running ok on Firefly3399 Thanks,

…

On Tue, Oct 10, 2017 at 7:50 PM, masahi ***@***.***> wrote: Are you running resnet bench? If so, try mobilenet bench first. It requires less memory I think. 2017/10/11 午前11:47 "kaishijeng" ***@***.***>: > I encountered CL_OUT_RESOURCES again with your changes. > > tvm._ffi.base.TVMError: [19:44:47] src/codegen/llvm/llvm_module.cc:59: > Check failed: ret == 0 (-1 vs. 0) [19:44:47] > src/runtime/opencl/opencl_module.cc:223: Check failed: e == CL_SUCCESS > OpenCL Error, code=-5: CL_OUT_OF_RESOURCES > > On Tue, Oct 10, 2017 at 7:26 PM, masahi ***@***.***> wrote: > > > Sure it's here > > <https://gist.github.com/masahi/de658ec98b1e275de6331a4cf4c873b3> > > > > Note that I all changed is just > > > > *target = "opencl" ctx = tvm.cl <http://tvm.cl>(0)* > > > > With the same change you can also run tutorial/from_onnx.py and > > tutorial/from_mxnet.py etc. > > > > You can also use > > python/nnvm/top/nn.py > > <https://gist.github.com/masahi/a9b67fe27b24b783334ef8fc6f877ef9> > > python/nnvm/top/reduction.py > > <https://gist.github.com/masahi/706fd19cb4c589bc7ea59d68cde99a8a> > > python/nnvm/top/tensor.py > > <https://gist.github.com/masahi/1d503f9e7839636dc8086d3737523d64> > > > > This should be enough to run resnet bench and tutorials. > > > > — > > You are receiving this because you were mentioned. > > Reply to this email directly, view it on GitHub > > <#530 (comment)>, or mute > > the thread > > <https://github.com/notifications/unsubscribe- > auth/AMGg3rLt4GKrsFxLsDkZ7bU_HeNNS6abks5srCdZgaJpZM4PzdCV> > > . > > > > — > You are receiving this because you authored the thread. > Reply to this email directly, view it on GitHub > <#530 (comment)>, or mute > the thread > <https://github.com/notifications/unsubscribe-auth/ABsbE1POGJTHEVVfNe- 2fQ8nVIAyfVBVks5srCwwgaJpZM4PzdCV> > . > — You are receiving this because you were mentioned. Reply to this email directly, view it on GitHub <#530 (comment)>, or mute the thread <https://github.com/notifications/unsubscribe-auth/AMGg3psKYLJ6VQRek8IY7MJQ47HLedFmks5srCzZgaJpZM4PzdCV> .

[masahi]

@tqchen Halide has auto scheduling work merged in master (here is an example of conv layer with auto scheduling). The results in the paper are impressive. Do you think auto scheduling in tvm is possible? That will be pretty cool.

Note that it is different from auto tuning. They automatically generate a image processing schedule that is on par or better than hand written ones by Halide experts, without actually running it. I wonder how auto scheduling can apply to conv and gemm.

[tqchen]

yes, we are looking into it. The auto schedule pattern for GPU and multi-core CPU could go beyond what pattern we can use for image processing works and we are gaining some experience on what new changes need to be done here

[masahi]

That's great, looking forward to learning the techniques.

[kaishijeng]

Why is resource required in conv2d_nchw.py much biigger than conv2d_hwnc? I still try to figure out why there is not enough resource in conv2d_nchw.py, but conv2d_hwnc.py is OK on unittest Thanks

…

On Wed, Oct 11, 2017 at 4:17 PM, masahi ***@***.***> wrote: That's great, looking forward to learning the techniques. — You are receiving this because you were mentioned. Reply to this email directly, view it on GitHub <#530 (comment)>, or mute the thread <https://github.com/notifications/unsubscribe-auth/AMGg3jUrLUfkqemOnJeLuCYeapJ79M9mks5srUykgaJpZM4PzdCV> .

[tqchen]

The conv2d_nchw is quite specially tuned for K80 card, which will launch quite a lot more threads with shared memory resources. They are not as good schedule for the mobile GPUs.

[kaishijeng]

1) Can I always use conv2d_hwnc even though it is not optimized for a particular processor/platform? 2) How do I know whch one , ie conv2d_nchw or conv2d_hwnc is used in the cuda_imagenet_bench.py <https://github.com/dmlc/nnvm/blob/master/examples/benchmark/cuda_imagenet_bench.py> ? Thanks,

…

On Thu, Oct 12, 2017 at 10:23 AM, Tianqi Chen ***@***.***> wrote: The conv2d_nchw is quite specially tuned for K80 card, which will launch quite a lot more threads with shared memory resources. They are not as good schedule for the mobile GPUs. — You are receiving this because you were mentioned. Reply to this email directly, view it on GitHub <#530 (comment)>, or mute the thread <https://github.com/notifications/unsubscribe-auth/AMGg3n-DLJxIH3Iw6_6jt6xK8UD19JdFks5srksRgaJpZM4PzdCV> .

[tqchen]

The imagenet bench always use the nchw one, the hnwc one is designed for the batch kernel in cuda. I think we can try to make the schedule more adjustable so it can adapt to more platforms as a future goal

[kaishijeng]

Will wait for your new design. I think if NNVM/TVM can take advantage of mobile GPU to speed up deep learning, it will be enable more deep learning applications on mobile devices. Thanks,

…

On Thu, Oct 12, 2017 at 3:52 PM, Tianqi Chen ***@***.***> wrote: The imagenet bench always use the nchw one, the hnwc one is designed for the batch kernel in cuda. I think we can try to make the schedule more adjustable so it can adapt to more platforms as a future goal — You are receiving this because you were mentioned. Reply to this email directly, view it on GitHub <#530 (comment)>, or mute the thread <https://github.com/notifications/unsubscribe-auth/AMGg3jY291ZSSjLacwyrpD4r2vrdcZr1ks5srpgUgaJpZM4PzdCV> .

[masahi]

Solved by #556

[kaishijeng]

Masahi, How do #556 solve this issue? With the latest code, I got different error of test_topi_conv2d_nchw.py below. The error of previous code is an insufficient resource of CL allocation. TVM: Initializing cython mode... Use memoize topi.tests.test_topi_conv2d.verify_con2d_nchw.get_ref_data.pkl(5, (1, 3, 224, 224), 'float32', 2, 3, (64, 3, 7, 7)) Skip because cuda is not enabled [10:56:06] src/runtime/opencl/opencl_device_api.cc:195: Initialize OpenCL platform 'ARM Platform' [10:56:06] src/runtime/opencl/opencl_device_api.cc:215: opencl(0)='Mali-T860' cl_device_id=0x7f75be3dd8 [10:56:06] /home/firefly/2TB/src/firefly/tvm/dmlc-core/include/dmlc/logging.h:308: [10:56:06] src/runtime/opencl/opencl_module.cc:223: Check failed: e == CL_SUCCESS OpenCL Error, code=-55: CL_INVALID_WORK_ITEM_SIZE

…

On Sat, Oct 14, 2017 at 4:55 PM, masahi ***@***.***> wrote: Closed #530 <#530>. — You are receiving this because you were mentioned. Reply to this email directly, view it on GitHub <#530 (comment)>, or mute the thread <https://github.com/notifications/unsubscribe-auth/AMGg3j_Ea_PtPZxSj4n-PtcyY_EdIEp9ks5ssUnfgaJpZM4PzdCV> .

[tqchen]

To enable Mali GPUs, we need to further mobile GPU options here https://github.com/dmlc/tvm/blob/master/python/tvm/target.py#L63

Note that most of the current cuda schedules need tweaks for the mali target. So additional registered function from topi is needed

[masahi]

Thanks, it looks like not all schedules under topi.cuda have been updated to use max_num_threads. I think we shouldn't simply replace every occurrence of 512 with max_num_threads. I need to discuss this with @tqchen

I am at work now, I will check the code later.

[kaishijeng]

Tiangi/masahi Just curious, why is there no opencl directory in tvm/topi/python/topi? There are rasp and cuda., but no opencl.

…

On Mon, Oct 16, 2017 at 2:39 AM, masahi ***@***.***> wrote: Thanks, it looks like not all schedules under topi.cuda have been updated to use max_num_threads. I need to discuss this with @tqchen <https://github.com/tqchen> — You are receiving this because you were mentioned. Reply to this email directly, view it on GitHub <#530 (comment)>, or mute the thread <https://github.com/notifications/unsubscribe-auth/AMGg3mxaHZH9RHjd6GQtU_0x78wfI-Gyks5ssvf3gaJpZM4PzdCV> .

[tqchen]

This is mainly due to fact that there is not yet effort to optimize schedule for opencl devices. We do expect that there should be topi/rocm(for amd gpus) and topi/mali (for mobile gpus)

[kaishijeng]

masahi

Does the latest code have proper  max_num_threads enforced?

If yes, what changes do I need to do to reduce max_num_threads for firefly3399?

Thanks,

[masahi]

@kaishijeng yes, at least it solves my AMD opencl problem. Not sure if it will work for you, though.

If max_num_threads = 256 is still too big for your card, you can set a smaller number here.

[kaishijeng]

Masahi, WIll try different values if 256 is too high. Just curious, I saw the change of max_num_threads is under target_name in ("rocm", "opencl"). Will firefly 3399 use the same target name? Thanks,

…

On Sun, Oct 22, 2017 at 7:28 PM, masahi ***@***.***> wrote: @kaishijeng <https://github.com/kaishijeng> yes, at least it solves my AMD opencl problem. Not sure if it will work for you, though. If max_num_threads = 256 is still too big for your card, you can set a smaller number here <https://github.com/dmlc/tvm/blob/122183583663ae11a5389e46d3ef6c92e2b82d05/python/tvm/target.py#L113> . — You are receiving this because you were mentioned. Reply to this email directly, view it on GitHub <#530 (comment)>, or mute the thread <https://github.com/notifications/unsubscribe-auth/AMGg3tQ8a_q9SQmmAZVjZdCS1FIvXKYLks5su_mtgaJpZM4PzdCV> .

[masahi]

Yes, what matters is the keyword "gpu" in self.keys += ("rocm", "gpu")
This will cause schedules registered under "gpu" to be called by opencl backend.

[kaishijeng]

I have tried different values of self.max_num_threads in target.py and even down to 8,still got Check failed: e == CL_SUCCESS OpenCL Error, code=-5: CL_OUT_OF_RESOURCES when doing test_conv2d_nchw Thanks,

…

On Sun, Oct 22, 2017 at 11:12 PM, masahi ***@***.***> wrote: Yes, what matters is the keyword "gpu" in self.keys += ("rocm", "gpu") This will cause schedules registered under "gpu" to be called by opencl backend. — You are receiving this because you were mentioned. Reply to this email directly, view it on GitHub <#530 (comment)>, or mute the thread <https://github.com/notifications/unsubscribe-auth/AMGg3lpV7hB1Y7P3_gwLQTBC6LB08OWrks5svC4zgaJpZM4PzdCV> .

[masahi]

Seems your card doesn't have enough memory (global or local) to run test_conv2d_nchw
You should try smaller inputs, as I told you earlier.

[kaishijeng]

But it is not useful to have a smaller input because most of nets have input size of (224,224, 3). Thanks,

…

On Sun, Oct 22, 2017 at 11:24 PM, masahi ***@***.***> wrote: Seems your card doesn't have enough memory to run test_conv2d_nchw You should try smaller inputs, as I told you earlier. — You are receiving this because you were mentioned. Reply to this email directly, view it on GitHub <#530 (comment)>, or mute the thread <https://github.com/notifications/unsubscribe-auth/AMGg3hn17G0KiMLMJEIJ45njMDbQCmXQks5svDEWgaJpZM4PzdCV> .

[masahi]

you can try reducing output channels or smaller filter size.

[tqchen]

The existing schedule for AMD/CUDA GPUs do not work well on the ARM GPUs, and we will need independent schedules for them. For example, it is usually not helpful to take use of shared memory