Codiary

Freitag, 25. September 2015

ArchLinux Manjaro Broadcom Wireless - Channel 12 + 13

Problem: On Manjaro (ArchLinux) there are only WLAN channels available up to 11, which is US standard, but not European.
I also could see this problem on Linux Mint and Xubuntu, but not on grml. Drivers like b43, broadcom-wl STA firmware could not give me success. What I did, after many tries failed (compiled own regulatory.bin, playing with crda, iw reg set, ...)

Uninstalled some packages:
pacman -Ru crda
pacman -Ru linux318-broadcom-wl
pacman -Ru zd1211-firmware
pacman -Ru wireless-regdb
pacman -Ru ipw2200-fw
pacman -Ru b43-fwcutter
pacman -Ru ipw2100-fw

Removed manjaro hardware config:
mhwd -r pci network-broadcom-wl

lspci -vnn -d 14e4:
... Broadcom Corporation BCM43225 802.11b/g/n [14e4:4357] (rev 01)
...
Kernel driver in use: bcma-pci-bridge
Kernel modules: bcma

iw reg get
global
country 00: DFS-UNSET
(2402 - 2472 @ 40), (6, 20), (N/A)
(2457 - 2482 @ 40), (6, 20), (N/A), PASSIVE-SCAN
(2474 - 2494 @ 20), (6, 20), (N/A), NO-OFDM, PASSIVE-SCAN
(5170 - 5250 @ 160), (6, 20), (N/A), PASSIVE-SCAN
(5250 - 5330 @ 160), (6, 20), (0 ms), DFS, PASSIVE-SCAN
(5490 - 5730 @ 160), (6, 20), (0 ms), DFS, PASSIVE-SCAN

After restart I could see Channel 13 (iwlist scan | grep Channel:)
iwlist chan also shows 13 channels.

Montag, 6. Juli 2015

E-Mailverschlüsselung mit Thunderbird+Enigmail

E-Mailverschlüsselung für Anfänger ohne viel Schnick-Schnack ;)

1] Thunderbird installieren
[E-Mailprogramm, welches deine E-Mailkonten verwalten wird]

3] *Enigmail*-Addon installieren
- Thunderbird -> Extras -> Add-ons
- nach Add-on 'Enigmail' suchen und installieren

4] Enigmail Einrichtungs-Assistenten durchklicken

6] Gpg4Win (eigentliches Verschlüsselungsprogramm) installieren, falls nicht schon geschehen

12] Widerrufszertifikat erzeugen [optional]. Mit diesem kann man seinen Schlüssel ungültig machen, falls man seine Passphrase verloren hat oder der eigene Schlüssel kompromittiert wurde. Wenn der öffentliche Schlüssel auf einem KeyServer aber nicht öffentlich zugänglich ist, braucht man auch kein Widerrufszertifikat (welches ohnehin auf einem separaten System hätte aufbewahrt werden müssen).

13] Falls man Widerrufszertifikat anlegt, dann sollte dieses auf einem anderen Rechner oder auf ein externes Speichermedium gespeichert werden. Denn falls der Rechner angegriffen wird, reicht das Widerrufszertifikat aus, um ohne Passphrase den Schlüssel ungültig zu machen.

16] Test-E-Mail zusammen mit eigenem öffentlichen Schlüssel an adele@gnupp.de schicken

17] Passphrase für eigenen Schlüssel eingeben zur Entschlüsselung der Mail

18] Mail entschlüsselt. Im Text findet sich der Public Key von adele. (Normalerweise kommen *.asc Dateien im Anhang - gleich mehr dazu)

19] Public Key Text markieren und in Zwischenablage kopieren (STRG+C):
-----BEGIN PGP PUBLIC KEY BLOCK----- ... -----END PGP PUBLIC KEY BLOCK-----

20] Schlüsselverwaltung: Thunderbird Menü->Enigmail->Schlüssel verwalten...
Dann: Bearbeiten->Aus Zwischenablage importieren

21] Public Key wurde erfolgreich importiert (Dialognachrichten manchmal falsch kodiert, einfach ignorieren)

22] Wenn Public Keys im Anhang verschickt werden, dann genügt ein Rechtsklick auf die *.asc-Datei zum Importieren des Schlüssels

25] Schlüsselechtheit verifizieren:
Fingerabdruck-Zahl (Fingerprint) des Schlüssels muss über einen zweiten, möglichst sicheren Kanal mit dem Eigentümer abgeglichen werden (Telefon, realer Kontakt)
Das ist notwendig, da ja der Mail-Übertragungsweg durch einen Angreifer kompromittiert sein kann.

26] Bei adele ist das nicht möglich, daher unterschreiben wir mit "Ich habe es nicht überprüft".

27] Die Einstellung zum Besitzervertrauen definiert die Vertrauenswürdigkeit des anderen Benutzers.

28] adele prüft als Bot die Fingerprints überhaupt nicht, daher nehmen wir in diesem Beispiel "Ich vertraue ihm NICHT".

29] Testkommunikation mit adele@gnupp.de Teil 2
- verschlüsselte und unterschriebene E-Mail an adele schreiben (siehe die gelben Buttons in der Enigmail-Symbolleiste)
-- E-Mailunterschrift sichert Integrität (Nachricht wurde nicht manipuliert) und Authentizität (Absender ist echt; unter Annahme Schlüssel ist gleich Absender), alles unter der Voraussetzung, dass man den Schlüssel des Absenders beim Importieren überprüft hat (Fingerabdruck/Fingerprint-Check über zweiten sicheren Kanal)

30] Sie erhalten daraufhin wieder eine verschlüsselte E-Mail mit der Klartext-Nachricht, die Sie eben verschlüsselt an adele geschickt haben.

31] Wenn Sie verschlüsselte E-Mails standardmäßig unterschreiben möchten:
- Extras->Konten-Einstellungen->OpenPGP-Sicherheit
- - Verschlüsselte Nachrichten unterschreiben

32] Bei verschlüsselten E-Mails mit Anhängen "Jeden Anhang einzeln verschlüsseln und unterschreiben ..." auswählen - dieser Auswahl-Dialog poppt beim Absenden auf

« »

Thunderbird installieren
[E-Mailprogramm, welches deine E-Mailkonten verwalten wird]
E-Mailkonto anlegen
Enigmail-Addon installieren
- Thunderbird -> Extras -> Add-ons
- nach Add-on 'Enigmail' suchen und installieren
- Enigmail Einrichtungs-Assistenten durchklicken und dabei
- - automatisch Gpg4Win (eigentliches Verschlüsselungsprogramm) installieren lassen
- - Schlüsselpaar erzeugen
Testkommunikation mit adele@gnupp.de Teil 1
- Testemail zusammen mit eigenem öffentlichen Schlüssel an adele@gnupp.de schicken
- nach ein paar Minuten erhält man verschlüsselte E-Mail, die automatisch entschlüsselt wird - dabei muss die Passphrase eingegeben werden
- Gesamten PGP PUBLIC KEY BLOCK markieren und kopieren
- eben kopierter öffentlicher Schlüssel in Enigmail-Schlüsselverwaltung einfügen per "Bearbeiten->Aus Zwischenablage importieren"
- - für gewöhnlich schickt man sich *.asc Dateien als Anhang, die dann ganz leicht per Rechtsklick importiert werden, ohne dass Text markiert werden muss
Schlüsselverwaltung
- in Schlüsselverwaltung auf importierten Schlüssel doppelklicken, um Schlüssel zu prüfen
- - Fingerabdruck-Zahl muss über einen zweiten Kanal mit dem Eigentümer abgeglichen werden (Telefon, realer Kontakt)
- - bei adele ist das nicht möglich, daher unterschreiben wir mit "Ich habe es nicht überprüft"
- - das Besitzervertrauen regelt die Vertrauenswürdigkeit des anderen Benutzers
- - adele prüft garantiert keine Fingerprints anderer Benutzer, daher in diesem Beispiel "Ich vertraue ihm NICHT" auswählen
Testkommunikation mit adele@gnupp.de Teil 2
- verschlüsselte und unterschriebene E-Mail an adele schreiben
- - Antwort mit verschlüsselter Nachricht von adele erhalten
Verschlüsselte E-Mails standardmäßig unterschreiben
- Extras->Konten-Einstellungen->OpenPGP-Sicherheit
- - Verschlüsselte Nachrichten unterschreiben
Bei verschlüsselten E-Mails mit Anhängen "Jeden Anhang einzeln verschlüsseln und unterschreiben ..." auswählen - dieser Auswahl-Dialog poppt beim Absenden auf

Nicht abgedeckt in diesem Tutorial werden: Keyserver, Experteneinstellungen, manuelle Key-Erstellung, Export von Unterschriften, Sicherheitsoptimierungen, ... Für tiefergehende Anleitungen und Erklärungen zu gpg verweise ich auf: http://www.hauke-laging.de/sicherheit/openpgp.html
Bei entsprechendem Bedarf folgen ggf. weitere Screenshot-Anleitungen.

Donnerstag, 3. Oktober 2013

Serializing Arrays and Alignment in CUDA

Perhaps you have to deal with message buffers in general or serialized arrays of mixed types, i.e., different arrays packed together into a generic char buffer array on host side and then unpacking on device side again. There you have to take care about aligning the data right, otherwise kernel launches will fail due to misaligned memory accesses. A 64-bit address is only allowed to start at multiples of 8 bytes.

If there is a char array with packed data in this order:

[3x int, 12x double]

... then misalignment takes place. The double starts not at a 64-bit aligned address. Launching a kernel which unpacks the buffer will fail. cuda-memcheck helps you out by this hint:

========= Invalid __global__ read of size 8
=========     at 0x00000248 in kernel(char*, unsigned int)
=========     by thread (0,0,0) in block (0,0,0)
=========     Address 0x500340004 is misaligned

You can place the data in descending order according to the type size:

[12x double, 3x int]

So remember about aligning next time (also pointing at me!).

If you just come here to see some old-fashioned serialization in action, here you get it:

//nvcc -O2 -m64 -gencode arch=compute_30,code=sm_30 mwe_char.cu
#include <stdio.h>
#include <cuda.h>
#include <cuda_runtime_api.h>
#include <device_launch_parameters.h>
#include <driver_types.h>


#define HANDLE_ERROR(err) __handleError(err, __FILE__, __LINE__)
#define HANDLE_LAST(msg)  __handleLastError(msg, __FILE__, __LINE__)

void __handleError( cudaError err, const char *file, const int line )
{
  if( cudaSuccess != err) {
    fprintf(stderr, "%s(%i) : Runtime API error %d: %s.\n",
                file, line, (int)err, cudaGetErrorString( err ) );
    exit(-1);
  }
}
void __handleLastError( const char *errorMessage, const char *file, const int line )
{
  cudaError_t err = cudaGetLastError();
  if( cudaSuccess != err) {
    fprintf(stderr, "%s(%i) : CUDA error : %s : (%d) %s.\n",
            file, line, errorMessage, (int)err, cudaGetErrorString( err ) );
  }
}

/*
 *
 */
__device__ __host__ int toBuffer(char* buffer, unsigned var, unsigned long long llvar)
{
  unsigned long long* ptr0 = reinterpret_cast<unsigned long long*>(buffer);
  unsigned* ptr1 = reinterpret_cast<unsigned*>(buffer+sizeof(unsigned long long));

  *ptr1 = var;
  *ptr0 = llvar;
  return 0;
} 
/*
 *
 */
__device__ __host__ int toVars(char* buffer, unsigned* var, unsigned long long* llvar)
{
  unsigned long long* ptr0 = reinterpret_cast<unsigned long long*>(buffer);
  unsigned* ptr1 = reinterpret_cast<unsigned*>(buffer+sizeof(unsigned long long));
  //printf("Addresses: %p %p\n", ptr0, ptr1);

  *var = *ptr1;
  *llvar = *ptr0;
  return 0;
} 
/*
 *
 */
__global__ void kernel(char* buffer, unsigned size)
{
  unsigned id = threadIdx.x + blockDim.x*blockIdx.x;
  unsigned var = 0;
  unsigned long long llvar = 0;
  if(id>0)
    return;

  for(unsigned k=0; k<size; ++k)
    printf("%u ", buffer[k]);
  printf("\n");

  int err = toVars(buffer, &var, &llvar);

  if(err)
   printf("Error %d.\n", err);
  else
  {
    printf("CUDA: %u %llu\n", var, llvar);
  }
}
/*
 *
 */
int main()
{
  unsigned var = 42;
  unsigned long long llvar = 123456789123456789;
  const size_t BUFFER_COUNT = sizeof(unsigned) + sizeof(unsigned long long);
  char* buffer = new char[BUFFER_COUNT];
  char* dbuffer = NULL;
  int err = 0;
  printf("Vars: %u %llu\n", var, llvar);
  // --- CPU ---
  err = toBuffer(buffer, var, llvar);
  if(err)
    return err;
  var = llvar = 0;
  err = toVars(buffer, &var, &llvar);
  if(err)
    return err;
  printf("CPU: toVars(): %u %llu\n", var, llvar);
  // --- CUDA ---
  HANDLE_ERROR( cudaSetDevice(0) );
  HANDLE_ERROR( cudaMalloc((void**)(&dbuffer), BUFFER_COUNT*sizeof(char)) );
  HANDLE_ERROR( cudaMemcpy(dbuffer, buffer, BUFFER_COUNT*sizeof(char), cudaMemcpyHostToDevice) );
  kernel<<<1,1>>>(dbuffer, BUFFER_COUNT);
  HANDLE_LAST("Kernel launch failed.");
  HANDLE_ERROR( cudaFree(dbuffer) );
  HANDLE_ERROR( cudaDeviceReset() );
  delete[] buffer;
  return 0;
}

Samstag, 18. Mai 2013

Planet Gravitation Map Simulator

Screenshot CUDA / OpenGL Demo of Gravitation Map Simulator

This is just another fun demo for CUDA and OpenGL. It is a slightly modified version of the so called "random oscillating magnetic pendulum" (ROMP). A pixel represents a start position of one particle. This object becomes attracted by all the planets around due to their gravity. After a while the object is going to hit a planet. Now that corresponding pixel gets the color of that planet. You see an evolving map of "gravity" structures in realtime (more or less *cough*).

Download: (tested on Linux and Windows (VS2010), needs CUDA, OpenGL, GLUT, GLEW)

Download Source Code v1.1
Download Source Code v1.0

Positions and masses of the planets can be changed by the user as well as the scale of the map.

The simulation uses Euler or Runge-Kutta integration method for solving the differential equation. It is also possible to use doubles instead of floats for better precision.

The idea for this demo comes from a friend, he also implemented the algorithm in processing. I ported this to CUDA, so the stuff is entirely computed and rendered on the GPU. You need a Nvidia CUDA capable GPU for that.


Screenshot 2.1, different coloring

Screenshot 2.2

Computation

Given a particle at position $\vec{y}(t)=\left(x(t), y(t)\right)$ at a time t. This particle has a (positive) mass m. In our space there are n planets with their fixed positions $\vec p_1,\ldots,\vec p_n$ and (positive) masses $m_1,\ldots,m_n$ .
The following equation represents the gravitational force of planet i acting on our particle. This force comes from Newton's law of universal gravitation:

$\displaystyle \vec F_i = G\frac{m_i\cdot m}{r_i^2}\vec e_{r_i}$

G is just a gravitation constant and can be neglected for our purposes. r is the distance between the particle and the planet, $\vec e_r$ is the force direction vector:

$\displaystyle \vec e_r = \frac{\vec p_i - \vec y}{\Vert \vec p_i - \vec y\Vert}$

Hence, there is the following force at a time t:

$\displaystyle \sum_{i=1}^n\vec F_i(t) = \sum_{i=1}^n m\cdot \frac{m_i\cdot\left(\vec p_i - \vec y(t)\right)}{\Vert\vec p_i-\vec y(t)\Vert^3}$

Newton's second law gives the equality:

$\displaystyle \vec F=m\cdot\vec a\Rightarrow \sum_{i=1}^n\vec F_i(t) = m\cdot \frac{\mathrm d^2 \vec y(t)}{\mathrm dt^2}$

The mass m of the particle can be eliminated in the equation. This ordinary differential equation system of order 2 need to be integrated two times to obtain the position function:

$\begin{displaymath}\begin{split}\vec v(t) =& \frac{\mathrm dy(t)}{\mathrm dt}=\i... ...c y(t) =& \int\limits_0^t \vec v(\tau)\mathrm d\tau \end{split}\end{displaymath}$

This integration can be solved numerically, e.g. by Euler's method:

$\begin{displaymath}\begin{split}\vec v_{k+1} =& \vec v_k + h\cdot\sum_{i=1}^n\fr... ...y_k\Vert^3}\\ \vec y_{k+1} =& \vec y_k + h\cdot v_k \end{split}\end{displaymath}$

with $y_k{=}y(t_k),\, v_k{=}(t_k),$ and $y(t_0){=}y_0,\, v(t_0){=}v_0$ as the initial values. h is the size of every step and v represents the velocity of the particle.

Here you see some videos, which are showing the simulation more or less in realtime.

More:
Article by Ingo Berg, 2006, on magnetic pendulum with implementation (using Beeman's algorithm):
http://www.codeproject.com/Articles/16166/The-magnetic-pendulum-fractal

Mathematica implementation:
http://nylander.wordpress.com/2007/10/27/magnetic-pendulum-strange-attractor/

Some more stuff, videos:
http://magnetmfa.wikispaces.com/pendula?responseToken=05ae0f7708a5c4c989455051dc970f570
http://www.youtube.com/watch?v=duy8s8C7-Uc
http://www.youtube.com/watch?v=QXf95_EKS6E

Freitag, 26. April 2013

Raycasted Spheres and Point Sprites vs. Geometry Instancing (OpenGL 3.3)

Source code of sphere_shader examples (raycasted spheres and mesh-based spheres) written in OpenGL 3.3. Code also includes handy camera, shader classes and (GPU) timer functions based on glQueryCounter. Tested on Linux and Windows.

Dependencies: OpenGL, GLEW, and FREEGLUT and glm.

Download source code:
https://github.com/tdd11235813/spheres_shader
(Update: 2016/03/12 - refactoring of code framework)
~~http://tubafun.bplaced.net/public/sphere_shader.zip~~
~~(Update: 2013/05/21 - added: common shaders with VBO and duplicated parameters)~~

Developing the QPTool I was interested in the performance of the different types of sphere visualizations.

Spheres with Geometry Instancing (mesh-based)
Sphere Impostors with hardware based Point Sprites (GL_POINT_SPRITE_ARB)
Sphere Impostors with Geometry Shaders
Sphere Impostors with common Shaders (Parameters via TBO)
Sphere Impostors with common Shaders (Parameters duplicated in VBO)

Mesh-based Spheres

The first implementation uses mesh-based spheres by the help of geometry instancing [2]. In the initialization a sphere mesh is sent down to the GPU. By glDrawElementsInstanced(...) the object geometry gets cloned many times. A shader places these objects into the scene. Usually this is a good improvement for visualizing a massive amount of same (complex) geometry. In our case there are two disadvantages:

Meshes are a visible inaccurate approximation of the spheres.
Polygon data increases with the resolution of the spheres, slows down the application regardless of the screen resolution.

In this picture you see in the bottom part the artifacts by the polygon approximation of the spheres. The upper two parts show raycasted sphere impostors (screenshot from QPTool v0.9.7).

Sphere Impostors by Raycast-Shaders

These two pictures are from the sphere_shader examples. First picture shows the mesh-based spheres and the second one the impostors.

If we want to use shaders to render the spheres in the fragment shader on our own, we need a billboard.

A billboard is a flat object, usually a quad (square), which faces the camera. [3]

This quad must be generated with texture coordinates before we can use it. A good explanation can be found here [4]. There are several ways to achieve this. It is possible to use hardware based point sprites, where for each vertex a billboard is created. In the vertex shader you just have to set the point size, which makes up the billboard size later. There are some issues with compatibility and correct visuals. If you take the particles demo from CUDA SDK, where this rendering method is utilized, you might notice wrong quad sizes, especially when the quads move to the borders of the screen. I tried to fix that, but there are still some glitches.

Another possibility uses geometry shader. Dummy vertices (one per sphere) are sent down to GPU and the geometry shader creates four vertices making the billboard for one sphere.

The last way would be sending four vertices per sphere with predefined texture coordinates down to the GPU. By this either a second data buffer is needed for the parameters OR all the parameters are duplicated (4x) in the vertex buffer object.

One disadvantage of the implemented sphere impostor shaders is, Early depth test is not possible, since z-buffer is computed in fragment shader.

Results

Setting: 800x600
Hardware: GPU Nvidia GTX 470 (Fermi), CPU Intel Quad Core i7 @3.07GHz, 12GB RAM
OS: Windows 7 64bit
Number of spheres: 100,000
Parameters per sphere: Position (4 floats), Color (4 floats), Radius (1 float)
Frames measured after warmup: 200

Benchmark Results of Rendering Methods for Spheres (in ms)
FOV = 60
Method	Avg1	Avg2	Min1	Min2	Max1	Max2
Sphere Geometry Instanced	121.523	0.013	121.173	0.006	122.134	0.023
Sphere Billboard Quads (TBO)	7.080	0.022	6.938	0.017	7.201	0.029
Sphere Billboard Quads (VBO)	1.879	0.002	1.874	0.001	1.886	0.003
Sphere Impostor Geometry Shader	1.915	0.001	1.900	0.001	1.933	0.002
Sphere Point Sprites, with Glitches :(	1.541	0.001	1.501	0.001	1.619	0.002
FOV = 20
Sphere Geometry Instanced	121.606	0.015	121.081	0.003	122.083	0.027
Sphere Billboard Quads (TBO)	8.320	0.023	8.291	0.019	8.359	0.031
Sphere Billboard Quads (VBO)	8.230	0.001	8.227	0.001	8.234	0.001
Sphere Impostor Geometry Shader	8.229	0.001	8.224	0.001	8.243	0.001
Sphere Point Sprites, with Glitches :(	7.756	0.001	7.752	0.001	7.763	0.001
FOV: Field of view. The lower the value, the more sphere-pixels have to be drawn. It is like zoom.
***1: Value for whole frame in ms.
***2: Value for binding and shader config in ms.

Sphere Impostor Geometry Shader: Takes 1 dummy vertex, creates 4 vertices for billboard on the fly.

Sphere Billboard Quads: Takes 4 vertices with preset texture coordinates for billboard. Uses either texture buffer object (TBO) for sphere parameters or parameters which are duplicated (4x) as vertex attributes in the vertex buffer object (VBO).

Sphere Geometry Instanced: Polygon representation of sphere. Using geometry instancing.

Sphere Point Sprites. Hardwarebased Billboard generation (point sprites). Alas, there are glitches and incompatibilities, so I do not recommend to use this.

My current implementation uses the geometry shader method, because it has no glitches so far and it is fast. Nevertheless I would recommend to use predefined billboards with duplicated data, since you will not find geometry shader support "everywhere".

On Linux with a Kepler GPU I obtained following interesting results:

Setting: 800x600
Hardware: GPU Nvidia GTX 670 (Kepler), CPU Intel Quad Core i7 @3.4GHz, 8GB RAM
OS: Linux Kubuntu 12.04
Number of spheres: 100,000
Parameters per sphere: Position (4 floats), Color (4 floats), Radius (1 float)
Frames measured after warmup: 200

Benchmark Results of Rendering Methods for Spheres (in ms)
FOV = 60
Method	Avg1	Avg2	Min1	Min2	Max1	Max2
Sphere Geometry Instanced	50.943	0.003	50.370	0.001	53.117	0.013
Sphere Billboard Quads (TBO)	1.087	0.003	1.069	0.001	1.159	0.010
Sphere Billboard Quads (VBO)	1.061	0.001	1.059	0.001	1.064	0.002
Sphere Impostor Geometry Shader	1.066	0.001	1.063	0.001	1.069	0.003
Sphere Point Sprites, with Glitches :(	0.897	0.001	0.895	0.001	0.902	0.003
FOV = 20
Sphere Geometry Instanced	50.798	0.004	50.356	0.001	53.487	0.015
Sphere Billboard Quads (TBO)	6.191	0.001	6.137	0.001	6.747	0.008
Sphere Billboard Quads (VBO)	6.133	0.001	6.130	0.001	6.137	0.003
Sphere Impostor Geometry Shader	6.121	0.001	6.116	0.001	6.127	0.002
Sphere Point Sprites, with Glitches :(	5.442	0.001	5.440	0.001	5.446	0.002
FOV: Field of view. The lower the value, the more sphere-pixels have to be drawn. It is like zoom.
***1: Value for whole frame in ms.
***2: Value for binding and shader config in ms.

Please leave a comment, if you encounter bugs or know more or if you have other ideas.

[1] - ~~http://www.arcsynthesis.org/gltut/Illumination/Tutorial%2013.html~~
https://web.archive.org/web/20150215085719/http://www.arcsynthesis.org/gltut/illumination/tutorial%2013.html
[2] - http://wiki.delphigl.com/index.php/shader_Instancing
[3] - http://nehe.gamedev.net/article/billboarding_how_to/18011/
[4] - http://www.sunsetlakesoftware.com/2011/05/08/enhancing-molecules-using-opengl-es-20