doe300
Posts: 45
Joined: Thu Dec 29, 2016 1:41 pm

Re: OpenCL on the VideoCore IV!

Sun Aug 19, 2018 2:37 pm

vottghern wrote:Result without sudo: [...]
I know that requiring root access to run an OpenCL program is not a good idea. But there are limitations, which are hard/impossible to go around. So for now, programs using VC4CL will always require admin rights.
vottghern wrote:Result WITH sudo: [...]
See here, also for more bugs/updates supporting OpenCV. I haven't had time yet to look into the warning, but it should not affect the program (except that it takes longer).
The error
vottghern wrote:OpenCL error CL_INVALID_WORK_GROUP_SIZE (-54) during call: clEnqueueNDRangeKernel [...]
comes from OpenCV (as sadly also so many other OpenCL programs) not heeding the implementation limit for the local work-group size. This can only be fixed inside the OpenCV code. VC4CL supports only a work-group size of 12 elements (and not the 16 required in this error).

Zenitur
Posts: 3
Joined: Fri Mar 16, 2018 7:27 pm

Re: OpenCL on the VideoCore IV!

Wed Aug 22, 2018 12:58 pm

Hello doe300! Thank you for updating of Wiki. I'd posted my GCC build because I'd used Raspbian 8 at that moment. Now I'm Using Raspbian 9, and VC4CL was successfully installed using the instruction from Wiki.

I'd tested VC4CL with a cgminer application. I had got 0 h/s. After few hours, an application stopped to work. Although, the GPU had some heat when cgminer was started. I'd tested with a nightly builds versions 0.4-2018-07-01 and 0.4-2018-08-21. This is my binary, and this is a command that I used:

Code: Select all

sudo ./cgminer -T --scrypt -o stratum+tcp://ltc-eu.give-me-coins.com:3333 -u Zenitur.rpi -p raspberry --no-submit-stale

doe300
Posts: 45
Joined: Thu Dec 29, 2016 1:41 pm

Re: OpenCL on the VideoCore IV!

Wed Aug 22, 2018 3:40 pm

What mining program do you use? I couldn't find any cgminer which uses OpenCL.
Zenitur wrote:Although, the GPU had some heat when cgminer was started
What do you mean? Actual heat or usage?

Zenitur
Posts: 3
Joined: Fri Mar 16, 2018 7:27 pm

Re: OpenCL on the VideoCore IV!

Wed Aug 22, 2018 5:34 pm

doe300, I'm sorry for my bad English. This is not my native language.

The last cgminer version that using OpenCL is 3.7.2. Starting from 3.8, Con Colivas dropped OpenCL support.

I'd attached my build in my previous message. It's a cloud storage, don't afraid of them, this is not a virus.
doe300 wrote:What do you mean? Actual heat or usage?
Is "heat" a temperature in celsius? I mean that my Raspberry Pi's GPU is 40° when idle, and 65° when cgminer is running. But it doesn't shows any results.

My target is not cgminer because of nobody mine Litecoin on GPU's today. I just making a tests. My target is sgminer-gm 5.5.6, a fork of cgminer. I want to know, how much hashrate can I get on Raspberry Pi's GPU in cryptonightv7 mining.

User avatar
ab1jx
Posts: 868
Joined: Thu Sep 26, 2013 1:54 pm
Location: Heath, MA USA
Contact: Website

Re: OpenCL on the VideoCore IV!

Fri Mar 22, 2019 9:32 pm

I installed on a ZeroW using the debs mentioned in this thread. I think it's mostly OK, I get this clinfo output:

Code: Select all

This is from my ZeroW 3/22/2019

Number of platforms                               1
  Platform Name                                   OpenCL for the Raspberry Pi VideoCore IV GPU
  Platform Vendor                                 doe300
  Platform Version                                OpenCL 1.2 VC4CL 0.4
  Platform Profile                                EMBEDDED_PROFILE
  Platform Extensions                             cl_khr_il_program cl_khr_spir cl_khr_create_command_queue cl_altera_device_temperature cl_altera_live_object_tracking cl_khr_icd cl_vc4cl_performance_counters
  Platform Extensions function suffix             VC4CL

  Platform Name                                   OpenCL for the Raspberry Pi VideoCore IV GPU
Number of devices                                 1
  Device Name                                     VideoCore IV GPU
  Device Vendor                                   Broadcom
  Device Vendor ID                                0xa5c
  Device Version                                  OpenCL 1.2 VC4CL 0.4
  Driver Version                                  0.4
  Device OpenCL C Version                         OpenCL C 1.2 
  Device Type                                     GPU
  Device Profile                                  EMBEDDED_PROFILE
  Device Available                                Yes
  Compiler Available                              Yes
  Linker Available                                Yes
  Max compute units                               1
  Available core IDs                              0, 64
  Max clock frequency                             300MHz
  Device Partition                                (core)
    Max number of sub-devices                     0
    Supported partition types                     None
    Supported affinity domains                    (n/a)
  Max work item dimensions                        3
  Max work item sizes                             12x12x12
  Max work group size                             12
  Preferred work group size multiple              <getWGsizes:1213: build program : error -15>
  Preferred / native vector sizes                 
    char                                                16 / 16      
    short                                               16 / 16      
    int                                                 16 / 16      
    long                                                 0 / 0       
    half                                                 0 / 0        (n/a)
    float                                               16 / 16      
    double                                               0 / 0        (n/a)
  Half-precision Floating-point support           (n/a)
  Single-precision Floating-point support         (core)
    Denormals                                     No
    Infinity and NANs                             No
    Round to nearest                              No
    Round to zero                                 Yes
    Round to infinity                             No
    IEEE754-2008 fused multiply-add               No
    Support is emulated in software               No
    Correctly-rounded divide and sqrt operations  No
  Double-precision Floating-point support         (n/a)
  Address bits                                    32, Little-Endian
  Global memory size                              67108864 (64MiB)
  Error Correction support                        No
  Max memory allocation                           67108864 (64MiB)
  Unified memory for Host and Device              Yes
  Minimum alignment for any data type             64 bytes
  Alignment of base address                       512 bits (64 bytes)
  Global Memory cache type                        Read/Write
  Global Memory cache size                        32768 (32KiB)
  Global Memory cache line size                   64 bytes
  Image support                                   No
  Local memory type                               Global
  Local memory size                               67108864 (64MiB)
  Max number of constant args                     64
  Max constant buffer size                        67108864 (64MiB)
  Max size of kernel argument                     256
  Queue properties                                
    Out-of-order execution                        No
    Profiling                                     Yes
  Prefer user sync for interop                    Yes
  Profiling timer resolution                      1ns
  Execution capabilities                          
    Run OpenCL kernels                            Yes
    Run native kernels                            No
  printf() buffer size                            0
  Built-in kernels                                (n/a)
  Device Extensions                               cl_khr_global_int32_base_atomics cl_khr_global_int32_extended_atomics cl_khr_local_int32_base_atomics cl_khr_local_int32_extended_atomics cl_khr_byte_addressable_store cl_nv_pragma_unroll cl_arm_core_id cl_ext_atomic_counters_32 cl_khr_initialize_memory cl_arm_integer_dot_product_int8 cl_arm_integer_dot_product_accumulate_int8 cl_arm_integer_dot_product_accumulate_int16

NULL platform behavior
  clGetPlatformInfo(NULL, CL_PLATFORM_NAME, ...)  OpenCL for the Raspberry Pi VideoCore IV GPU
  clGetDeviceIDs(NULL, CL_DEVICE_TYPE_ALL, ...)   Success [VC4CL]
  clCreateContext(NULL, ...) [default]            Success [VC4CL]
  clCreateContextFromType(NULL, CL_DEVICE_TYPE_DEFAULT)  Success (1)
    Platform Name                                 OpenCL for the Raspberry Pi VideoCore IV GPU
    Device Name                                   VideoCore IV GPU
  clCreateContextFromType(NULL, CL_DEVICE_TYPE_CPU)  No devices found in platform
  clCreateContextFromType(NULL, CL_DEVICE_TYPE_GPU)  Success (1)
    Platform Name                                 OpenCL for the Raspberry Pi VideoCore IV GPU
    Device Name                                   VideoCore IV GPU
  clCreateContextFromType(NULL, CL_DEVICE_TYPE_ACCELERATOR)  No devices found in platform
  clCreateContextFromType(NULL, CL_DEVICE_TYPE_CUSTOM)  No devices found in platform
  clCreateContextFromType(NULL, CL_DEVICE_TYPE_ALL)  Success (1)
    Platform Name                                 OpenCL for the Raspberry Pi VideoCore IV GPU
    Device Name                                   VideoCore IV GPU

ICD loader properties
  ICD loader Name                                 OpenCL ICD Loader
  ICD loader Vendor                               OCL Icd free software
  ICD loader Version                              2.2.11
  ICD loader Profile                              OpenCL 2.1
But something in the apt/dpkg, etc. isn't happy, it thinks vc4c needs some configuration step.

Code: Select all

zero3# apt-get --fix-broken install
Reading package lists... Done
Building dependency tree       
Reading state information... Done
0 upgraded, 0 newly installed, 0 to remove and 0 not upgraded.
2 not fully installed or removed.
After this operation, 0 B of additional disk space will be used.
Setting up vc4c (0.4-2019-03-09) ...
/var/lib/dpkg/info/vc4c.postinst: 3: /var/lib/dpkg/info/vc4c.postinst: /usr/local/bin/vc4c: not found
dpkg: error processing package vc4c (--configure):
 subprocess installed post-installation script returned error exit status 127
dpkg: dependency problems prevent configuration of vc4cl:
 vc4cl depends on vc4c; however:
  Package vc4c is not configured yet.

dpkg: error processing package vc4cl (--configure):
 dependency problems - leaving unconfigured
Errors were encountered while processing:
 vc4c
 vc4cl
E: Sub-process /usr/bin/dpkg returned an error code (1)
zero3# 
BTW I was hoping to GPU mine on a Zero with this sgminer https://github.com/sgminer-dev/sgminer but I seem to be having jansson problems, which I'll submit under Issues to Github, even though nobody else has mentioned it.
Last edited by ab1jx on Sun Mar 24, 2019 1:15 am, edited 1 time in total.

User avatar
ab1jx
Posts: 868
Joined: Thu Sep 26, 2013 1:54 pm
Location: Heath, MA USA
Contact: Website

Re: OpenCL on the VideoCore IV!

Sun Mar 24, 2019 12:45 am

Alright, maybe I got it fixed. From my notes:

Code: Select all



zero3# apt-get --fix-broken install
Reading package lists... Done
Building dependency tree       
Reading state information... Done
0 upgraded, 0 newly installed, 0 to remove and 0 not upgraded.
2 not fully installed or removed.
After this operation, 0 B of additional disk space will be used.
Setting up vc4c (0.4-2019-03-09) ...
/var/lib/dpkg/info/vc4c.postinst: 3: /var/lib/dpkg/info/vc4c.postinst: /usr/local/bin/vc4c: not found
dpkg: error processing package vc4c (--configure):
 subprocess installed post-installation script returned error exit status 127
dpkg: dependency problems prevent configuration of vc4cl:
 vc4cl depends on vc4c; however:
  Package vc4c is not configured yet.

dpkg: error processing package vc4cl (--configure):
 dependency problems - leaving unconfigured
Errors were encountered while processing:
 vc4c
 vc4cl
E: Sub-process /usr/bin/dpkg returned an error code (1)
zero3# 

------------

The vc4c postinst gets a file not found, when I look at
/var/lib/dpkg/info/vc4c.postinst

I see

#!/bin/sh                                                                       
set -e                                                                          
/usr/local/bin/vc4c --precompile-stdlib -o /usr/local/include/vc4cl-stdlib/ /usr
/local/include/vc4cl-stdlib/VC4CLStdLib.h                                       
echo "VC4CL standard library precompiled into /usr/local/include/vc4cl-stdlib/" 

but

zero3# ls -la /usr/local/bin/vc4c
ls: cannot access '/usr/local/bin/vc4c': No such file or directory
zero3# 

and

zero3# vc4c
bash: vc4c: command not found
zero3# 

It's not getting installed into the path.

I look in /usr/local/bin and I see VC4C (not vc4c).  Uh-oh.  So I symlinked
it.     

Then I did 
apt-get install --fix-broken again and eventually I see:

ib//VC4CLStdLib.h.pch /usr/local/include/vc4cl-stdlib/VC4CLStdLib.h
[W] Sat Mar 23 20:10:27 2019: Warnings in precompilation:
[W] Sat Mar 23 20:10:27 2019: In file included from /usr/local/include/vc4cl-stdlib/VC4CLStdLib.h:17:
In file included from /usr/local/include/vc4cl-stdlib/_config.h:12:
/usr/local/include/vc4cl-stdlib/opencl-c.h:12828:13: warning: unknown attribute 'convergent' ignored
void __ovld __conv barrier(cl_mem_fence_flags flags);
            ^
/usr/local/include/vc4cl-stdlib/opencl-c.h:28:31: note: expanded from macro '__conv'
#define __conv __attribute__((convergent))
                              ^
1 warning generated.

[I] Sat Mar 23 20:10:27 2019: Pre-compiling standard library with: /usr/bin/clang-3.9 -cc1 -triple spir-unknown-unknown -O3 -ffp-contract=off -cl-std=CL1.2 -cl-kernel-arg-info -cl-single-precision-constant -Wno-all -Wno-gcc-compat -Wdouble-promotion -Wno-undefined-inline -x cl -emit-llvm-bc -o /usr/local/include/vc4cl-stdlib//VC4CLStdLib.bc /usr/local/include/vc4cl-stdlib/VC4CLStdLib.h
[W] Sat Mar 23 20:12:34 2019: Warnings in precompilation:
[W] Sat Mar 23 20:12:34 2019: In file included from /usr/local/include/vc4cl-stdlib/VC4CLStdLib.h:17:
In file included from /usr/local/include/vc4cl-stdlib/_config.h:12:
/usr/local/include/vc4cl-stdlib/opencl-c.h:12828:13: warning: unknown attribute 'convergent' ignored
void __ovld __conv barrier(cl_mem_fence_flags flags);
            ^
/usr/local/include/vc4cl-stdlib/opencl-c.h:28:31: note: expanded from macro '__conv'
#define __conv __attribute__((convergent))
                              ^
1 warning generated.

VC4CL standard library precompiled into /usr/local/include/vc4cl-stdlib/
Setting up vc4cl (0.4-2019-02-09) ...
W: APT had planned for dpkg to do more than it reported back (3 vs 7).
   Affected packages: vc4c:armhf
zero3# 

Finally I did 
vc4c -v                                
and got

Running VC4C in version: 0.4 Build configuration: debug mode; multi-threaded
optimization; clang 3.9+ OpenCL features; clang in /usr/bin/clang-3.9; LLVM
library front-end with libLLVM 3.9; vc4asm verification
I don't know about the unknown attribute 'convergent' warning but the main problem was one of case, VC4C vs vc4c, which caused the config step to fail.

I should probably file this as an issue on Github too. Done, see https://github.com/doe300/VC4CL/issues/63

doe300
Posts: 45
Joined: Thu Dec 29, 2016 1:41 pm

Re: OpenCL on the VideoCore IV!

Sun Mar 24, 2019 7:26 am

ab1jx wrote:
Sun Mar 24, 2019 12:45 am
I don't know about the unknown attribute 'convergent' warning but the main problem was one of case, VC4C vs vc4c, which caused the config step to fail.

I should probably file this as an issue on Github too. Done, see https://github.com/doe300/VC4CL/issues/63
Thanks for investigating further and filing a ticket. The warning with the unknown attribute can be ignored, this is due to the clang version of Raspbian being to old and not supporting the attribute yet, which is not a problem. I should probably completely hide the warning to avoid confusion though.

User avatar
ab1jx
Posts: 868
Joined: Thu Sep 26, 2013 1:54 pm
Location: Heath, MA USA
Contact: Website

Re: OpenCL on the VideoCore IV!

Sun Mar 24, 2019 4:39 pm

Anybody get it to mine yet? I don't want it just for that but I'm not up to writing OpenCL code yet. I'm trying to use a cgminer from https://github.com/ozbenh/cgminer I've been mining Litecoin and Bitcoin as a hobby for a couple years, I've built and used several versions of cgminer, never tried with a GPU before.

My cgminer.conf:

Code: Select all

{
"pools" : [
	{
		"url" : "stratum+tcp://us.litecoinpool.org:3333",
		"user" : "ab1jx.3",
		"pass" : "3"
	}
]
,
"intensity" : "10",
"vectors" : "1",
"worksize" : "12",
"lookup-gap" : "0",
"thread-concurrency" : "1",
"shaders" : "0",
"api-mcast-port" : "4028",
"api-port" : "4028",
"expiry" : "120",
"gpu-dyninterval" : "7",
"gpu-platform" : "0",
"gpu-threads" : "1",
"log" : "5",
"no-pool-disable" : true,
"queue" : "1",
"scan-time" : "30",
"shares" : "0",
"kernel-path" : "/usr/local/bin"
}
The results:

Code: Select all

I left it running overnight but it never actually did any mining.

 [2019-03-24 04:18:48] Started cgminer 3.7.2
 [2019-03-24 04:18:48] Started cgminer 3.7.2
 [2019-03-24 04:18:48] Loaded configuration file /root/.cgminer/cgminer.conf
 [2019-03-24 04:18:49] Probing for an alive pool
 [2019-03-24 04:18:49] Pool 0 difficulty changed to 65536
 [2019-03-24 04:19:38] Pool 0 difficulty changed to 16384
 [2019-03-24 04:20:29] Pool 0 difficulty changed to 4096
 [2019-03-24 04:21:16] Pool 0 difficulty changed to 1024
 [2019-03-24 04:21:39] Network diff set to 9.42M
 [2019-03-24 04:21:41] Stratum from pool 0 requested work restart
 [2019-03-24 04:22:06] Pool 0 difficulty changed to 256
 [2019-03-24 04:22:25] Stratum from pool 0 requested work restart
 [2019-03-24 04:22:54] Stratum from pool 0 requested work restart
 [2019-03-24 04:23:14] Stratum from pool 0 detected new block
 [2019-03-24 04:23:18] Stratum from pool 0 requested work restart
 [2019-03-24 04:27:48] Stratum from pool 0 requested work restart
 [2019-03-24 04:28:26] Stratum from pool 0 requested work restart
 [2019-03-24 04:29:08] Error -11: Building Program (clBuildProgram)
 [2019-03-24 04:29:08] [W] Sun Mar 24 04:19:12 2019: 64-bit operations are not s
upported by the VideoCore IV architecture, further compilation may fail!
[W] Sun Mar 24 04:19:12 2019: 64-bit operations are not supported by the VideoCo
re IV architecture, further compilation may f
 [2019-03-24 04:29:08] Failed to init GPU thread 0, disabling device 0
 [2019-03-24 04:29:08] Restarting the GPU from the menu will not fix this.
 [2019-03-24 04:29:08] Try restarting cgminer.
Press enter to continue:

......

 cgminer version 3.7.2 - Started: [2019-03-24 04:18:50]
--------------------------------------------------------------------------------
 (5s):0.000 (avg):0.000h/s | A:0  R:0  HW:0  WU:0.0/m
 ST: 2  SS: 0  NB: 136  LW: 515  GF: 0  RF: 0
 Connected to us.litecoinpool.org diff 256 with stratum as user ab1jx.3
 Block: 766e651b...  Diff:9.42M  Started: [10:58:02]  Best share: 0
--------------------------------------------------------------------------------
 [P]ool management [S]ettings [D]isplay options [Q]uit
 GPU 0:                | OFF   / 0.000h/s | A:0 R:0 HW:0 WU:0.0/m I:10
--------------------------------------------------------------------------------

 [2019-03-24 10:36:32] Stratum from pool 0 requested work restart
 [2019-03-24 10:37:05] Stratum from pool 0 detected new block
 [2019-03-24 10:37:10] Stratum from pool 0 requested work restart
 [2019-03-24 10:39:56] Stratum from pool 0 requested work restart
 [2019-03-24 10:40:11] Stratum from pool 0 requested work restart
 [2019-03-24 10:42:56] Stratum from pool 0 requested work restart
 [2019-03-24 10:43:18] Stratum from pool 0 requested work restart
 [2019-03-24 10:43:19] Stratum from pool 0 detected new block
 [2019-03-24 10:45:30] Stratum from pool 0 detected new block
 [2019-03-24 10:45:35] Stratum from pool 0 requested work restart
 [2019-03-24 10:46:01] Stratum from pool 0 requested work restart
 [2019-03-24 10:46:57] Stratum from pool 0 detected new block
 [2019-03-24 10:47:02] Stratum from pool 0 requested work restart
 [2019-03-24 10:47:47] Stratum from pool 0 requested work restart
 [2019-03-24 10:47:56] Stratum from pool 0 requested work restart
 [2019-03-24 10:49:18] Stratum from pool 0 requested work restart
 [2019-03-24 10:50:37] Stratum from pool 0 requested work restart
 [2019-03-24 10:50:48] Stratum from pool 0 requested work restart
 [2019-03-24 10:52:24] Stratum from pool 0 requested work restart
 [2019-03-24 10:52:42] Stratum from pool 0 requested work restart
 [2019-03-24 10:53:12] Stratum from pool 0 detected new block
 [2019-03-24 10:53:18] Stratum from pool 0 requested work restart
 [2019-03-24 10:53:27] Stratum from pool 0 requested work restart
 [2019-03-24 10:56:31] Stratum from pool 0 requested work restart
 [2019-03-24 10:58:02] Stratum from pool 0 detected new block
When it's working you should see some "accepted" lines in there after a couple minutes, then just about every line.

It never actually did anything but I doubt I have it set up right. There is the message in there
64-bit operations are not supported by the VideoCore IV architecture
It may be possible to configure it to only use 32 bit, I never ran into this before. My first GPU mining attempt. cgminer's been around for years, I have a different fork of it running a Gridseed ASIC miner on another Zero. And another running a GeckoScience 2Pac ASIC Bitcoin plugminer, you seem to have to find the right fork made for what you're trying to do, with the right drivers. You get cgminer running once with a few command line parameters, then if those are close enough to run there's a menu option to write out a config file. After that you can just start cgminer and it picks up settings from the config file. There are many more settings in there than I started out with, usually the defaults are reasonable.

I'm suspicious of these lines in the cgminer.conf not being right, I don't know how to set them. These are the first settings I've tried that didn't kick out an error message.

Code: Select all

"intensity" : "10",
"vectors" : "1",
"worksize" : "12",
"lookup-gap" : "0",
"thread-concurrency" : "1",
"shaders" : "0",
"gpu-dyninterval" : "7",
"gpu-platform" : "0",
"gpu-threads" : "1",
This was on a ZeroW, same GPU. And cheap to buy but they need a handful of adapters to hook them up. I'm running headless, talking to it over ssh via wifi.
------------ later --------
I did some more looking through what little documentation on this cgminer there is, and grepped its source directory for "64 bit" and "64-bit". I find stuff about whether the host computer is 64 bit or not, but noting about whether the GPU is 64 bit. It was mostly written for ATI and Nvidia cards. I can't even find the phrase "64-bit operations" in there, could it be coming from vc4c? This stuff worries me:

Code: Select all

[2019-03-24 18:59:53] Error -11: Building Program (clBuildProgram) [2019-03-24 18:59:53] [W] Sun Mar 24 18:49:53 2019: 64-bit operations are not supported by the VideoCore IV architecture, further compilation may fail!
[W] Sun Mar 24 18:49:53 2019: 64-bit operations are not supported by the VideoCore IV architecture, further compilation may f

[12d[4M[20d [2019-03-24 18:59:53] [W] Sun Mar 24 18:49:53 2019: 64-bit operations are not s[21;1Hupported by the VideoCore IV architecture, further compilation may fail!
I guess I should look up clBuildProgram and see if there are any options there.

Return to “Graphics programming”