Setting kernel constants using CUDA driver API

CUDA uses special qualifiers (namely, __constant__) to tell to the compiler that some variables should be placed direcly in a cached memory banks. The content of such variables (e.g., numeric values, constant arrays) can be set by the host before a kernel is executed. This is achieved, with CUDA runtime API, by using the following statement:

cudaMemcpyToSymbol(symbol, &variable, sizeof(variable));

Piace of cake. But how does it work with the CUDA driver API? Well, it is not clearly explained in the programming guide (at least, I couldn't find any examples). It turns out that the process requires two steps:

1. get a reference to the __constant__ variable in the device;
2. initialize it using a memcpy "Host to Device".

In practice, it works as follows:

CUdeviceptr constDevPtr;
cuModuleGetGlobal(&constDevPtr, NULL, module, "symbol");
cuMemcpyHtoD(constDevPtr, &variable, sizeof(variable));

I tested the code above and it works seemlessly. I hope this will save some headaches to CUDA developers.

On Golden Rules

Western society is permeated, on some extent, by a latent christian mentality. Even though a small part of western population is really christian, in a strict sense, "we" all agree on a passage of Matthew's gospel:

Whatever you want men to do to you, do also to them.
- Matthew, 7:12

This fundamental principle is generally referred to as Jesus Christ's golden rule, and it is considered the basis of christian god's will. Interestingly, this idea is way older than the New Testament. A symmetric version of this sentence was indeed pronounced by Confucius more than 4 centuries before Jesus Christ:

What you do not want done to yourself, do not do to others.
- The Analects of Confucius, XVː24

It is worth noting that similar concepts were proposed hundreds of years before. For instance, an egyptian papyrus of the late period reads:

That which you hate to be done to you, do not do to another.

Thus, it seems that christianity is just a statement of a "common-sense" on which everybody was already agreeing since centuries.  Nevertheless, I think that Christ did mess up with the proposition, leading to a deformed meaning of the principle which lead to proselytism and abuses.

Christianity assumes that something that is good to me is NECESSARILY good for everyone else, but this is far from true. We are all different, hence the assumption that we all want the same things from others is both naïve and a bit arrogant. The christian god wants us to overcome the personal space of the people surrounding us, in order to force the delivery of our "good" intentions. Naturally, there is here a second implicit assumption that what is "good" for christians must be good in any other culture or society, in any time and situation. Naturally, it is wrong.

On the contrary, the rest of golden rules (e.g., confucianism) respects the coexistence and promotes diversity (behavioral and cultural) by not doing anything that may cause frictions or arguments and, in particular, by avoiding any action that may cause suffering to me in the first place, because it is likely that it will make the others suffer too.

The paradigm shift is clear and consequences are straightforward: when a christian person feels that something is good for him, he will try to make the others feel the same. This misplaced urge, which completely ignores individual needs, generates absurd debates about mutual respect and morality. Moreover, christians have weird expectations from humanity, which is supposed to behave with a strict christian perspective (e.g., doing christian good actions) because of a wrong (or malicious) re-statement of a millenary concept produced by their fundamental prophet.

So long, LucasArts

Sia maledetta Disney per questo.

Semplici ipocriti o autentici fessi allo sbaraglio?

I grillini, aka movimentisti a cinque stelle, sono riusciti a diventare il primo secondo partito alla Camera* sulla base di un atteggiamento intransigente di rottura con la "vecchia" politica, intesa quantomeno come nomi e sigle. Grazie all'approccio "non faremo alleanze", hanno drenato voti di protesta e sono esplosi come partito.

Al di là del successo elettorale, trovo demenziale questo approccio per un partito sedicente rivoluzionario, perché a meno di non portare a casa il 50%+1 dei voti ("facile" alla Camera, impossibile al Senato) significa puntare a non voler mai governare, dove mai si intende in eterno. Altro che rivoluzione, si passa solo per gli scassapalle che assistono al gioco di altri senza proporre nulla di proprio, atteggiamento tipico di una nazione con 50milioni di commissarî tecnici ma  diametralmente opposto a quello che occorre in Italia oggi.

Questo è l'atteggiamento di Grillo. Rispecchia effettivamente quello che i suoi deputati stanno mantenendo? Alla prima verifica concreta - l'elezione del Presidente del Senato - alcuni grillini dissidenti hanno mancato di seguire  la linea oltranzista del non-capo e hanno deciso che, pur di evitare l'elezione dell'impresentabile Schifani, era meglio regalare voti al candidato di csx.

Atteggiamento saggio, rispettabile, intelligente. Purtuttavia ipocrita, perché si è  sempre professato il contrario che, puntualmente, avrebbe conseguenze nefaste, vedasi Schifani. Ci siete già arrivati?

Ora, tutto potrebbe risolversi in maniera serena, se non fosse per alcune  giustificazioni arrivate in queste ore. Partirei dal preambolo, ovvero dalle infelici (eufemismo) dichiarazioni della capogruppo 5 stelle alla Camera, Roberta Lombardi:

"Le elezioni dei rappresentanti sono troppo lente, sono uno spreco di denaro pubblico. Non si può perdere così una giornata".

Molto bene, ci sta. Le procedure sono troppo farraginose, arcaiche; meglio procedere più spediti, tanto si hanno le idee chiare e non c'è bisogno di ragionare, ci insegna la Lombardi.

Il Senatore Vito Crimi

Tutto vero? Forse no, visto che dopo il voto che ha portato all'elezione di Grasso, Grillo ha tuonato contro i "ribelli" che hanno appoggiato un candidato esterno, portando Vito Crimi (capogruppo al Senato) a ribattere così:

"Ci siamo subito riuniti, ci siamo fatti portare panini e acqua nella stanza, e confrontati a lungo (dalle 14.00 alle 16.30, NdMe) senza il tempo per poter approntare uno streaming ne poter interpellare in modo attendibile la rete… Alla fine si è deciso di votare per alzata di mano tre opzioni: 1) Votare Grasso 2) Lasciare libertà di voto 3) Non votare ne Grasso ne Schifani (scheda bianca, nulla, con scritto Orellana, non votare erano tutte opzioni equivalenti ai fini della terza opzione). Su ogni opzione si è votato liberamente, senza il vincolo che chi votava per una non potesse votare per l’altra, e per ciascuna si è cercata una maggioranza, e la maggioranza assoluta ha votato per la terza opzione".

Ecco alcune mie osservazioni:

• evidentemente, tutti i buoni propositi iniziali (nessuna alleanza, niente compromessi) erano pura demagogia, dato che quando si passa alla politica vera e propria l'intransigenza programmatica viene agilmente buttata alle ortiche. Buono a sapersi, ma faccio notare a Crimi e ai 5 Stelle che non c'era niente da dover discutere;
• è interessante anche il fatto che senza una connessione ad internet, i 5 Stelle non sappiano che pesci pigliare. Si deduce che un sotto-campionamento dell'intelligenza collettiva non ne conserva le caratteristiche. In altre parole, non è scale-free. Forse (forse?) serve più rappresentanza in Parlamento, piuttosto che invocare demagogicamente il taglio dei parlamentari;
• trovo stupendo il fatto che, a dispetto di una assemblea in cui si è deliberato di non votare nè Grasso, nè Schifani (cosa ovvia, ripeto, per i 5 stelle), poi grazie al voto segreto ognuno faccia il cazzo che gli pare;
• trovo infine magnifico accostare le parole di Climi a quelle della Lombardi: anziché bofonchiare stupidate riguardo sprechi e tempo eccessivo, perché non lo si impiega per RAGIONARE e DISCUTERE? 

Cari grillini, fate il vostro lavoro, fatelo bene, fatelo con intelligenza, piantatela di dar retta a Grillo che vi manda solo in confusione e vi obbliga ad una linea demenziale. Accendete il cervello individuale, solo questo serve all'Italia.

EDIT: niente, la realtà supera sempre la fantasia e spazza via ogni barlume di logica. Beppe Grillo ha unilateralmente deciso di introdurre due "coordinatori della comunicazione", uno per ogni ramo del Parlamento, affinché "si interfaccino con i capi-gruppo". Insomma, che fondamentalmente gli dicano che cosa dire. Naturalmente, controllati da Grillo e pagati con soldi pubblici, perché la democrazia dal basso funziona così.

EDIT: la domanda retorica è stata parzialmente fugata dalla Lombardi stessa.

* La mia amica Lara mi fa presente che, contando i voti degli italiani all'estero, il M5S è il secondo partito alla Camera.

Ubuntu, launch problem with Nvidia GeForce (GTX 590?)

I experienced serious difficulties in launching my Ubuntu live distribution. In particular, both the 11.10 and the 12.04 versions did not recognize any keyboard (PS/2 or USB, it was the same), and halted during the startup phase.

The interesting part is that on a completely identical computer, except the video card, everything worked correctly.

Googling around, I realized that it could be due to my dual-GPU Nvidia GTX 590, that may be not correctly supported by the open source driver, Noveau. The workaround is to use the "nomodeset" flag under the "other options" menu.

Once installed, you can set Nvidia's proprietary driver, which works in a seemless way.

Dear Hollywood, are you trying to tell me something?

IMDB's home page. Do you notice anything?

Today I bumped upon IMDB as I came to know that a fourth chapter of Toy Story is in production: a sad commercial move. Anyway, when I loaded the home page, something has flashed before my eyes, something that should be self-evident by my screenshot above.

Hollywood's approach with violence - and guns in a particular way - clashes with many recent crime news and even with recent President Obama's words.

How do you expect that people stop killing each other, even without plausible reasons, if movies and culture promote the use of guns? As long as the coolness of characters remains proportional to the size of their weapons, it is no surprise that we won't do any progress.

However, this probably is not a problem for the lobbyists and I'm pretty sure that weapon producers are rubbing their hands for this kind of brain-shaping.

Giornalismo italiano e copia/incolla

Qual è la probabilità che due redattori di due testate indipendenti e teoricamente concorrenti scrivano due articoli completamente identici, parola per parola? Diciamo zero?

Al di là dell'assoluta inutilità della notizia, è interessante che RaiNews24 e Repubblica.it abbiano pubblicato lo stesso identico testo, la cui unica differenza è nel titolo vagamente differente. Nella migliore delle ipotesi, l' articolo è stato offerto ad entrambe le testate da un freelancer ed è stato riversato online senza troppi complimenti.

Non conosco la produzione di RaiNews24, ma so per certo che la qualità degli articoli di Repubblica è totalmente in picchiata — con punte di autentico orrore quando vengono trattati argomenti scientifici — il che mi porta a pensare che questo tipo di approccio al risparmio usato per gli articoli di contorno venga applicato anche a quelli principali. Tristezza.

Meta-programming in Python

Python gives the programmer a plethora of facilities for the meta-programming, for instance the introspection of objects, modules, and such. Moreover, like many other interpreted languages, it allows the execution of generated code.

Which is cool, obviously, if you like algorithms that manipulate other algorithmic-generated algorithms.

http://xkcd.com/917/

This is definitively my case, for I need to generate, at run-time, some data structures and methods that shall be given as an input to an ODE integrator. As both things must be visible from the calling module I need to put them into the global namespace. Piece of cake:

exec open(PATH).read() in globals()

where PATH points to the program-generated file that contains the data structures and the functions I need. Again, the elegance and conciseness of Python are astonishing: exec command checks the syntax and interprets the content of the file, after having it read with read() method, and all the variables are kept in globals namespace.

One-liner power!

Python, my first time

When it comes to learning a new technology I need some starting project that pushes me and gives me "momentum". In particular, it was years that I wanted to learn Python. I touched this language two years ago, at CINECA's SciViz course, and I toyed around with Django, but they have been momentary encounters and I was waiting for some small-sized, not too trivial problem to solve with it.

Recently, a friend of mine was experiencing problems with some visual representations of temporal series, stochastic in nature and very irregular as number of points. The dataset was huge (in the order of gigabytes) therefore an automatic solution was needed.

Probably PERL would have been a better choice, but I have taken the chance and, this morning, in a couple of hours I've completed my first Python project. It took a while to learn the syntax for dictionaries, lambda functions and some useful I/O classes, but the rest was really piece of cake. My final implementation works as follows:

• a brief preprocessing, in order to determine the length of each temporal series, done with regular expressions and a dictionary;
• some filtering on the dictionary, in order to remove spurious results;
• reading the file again while skipping, in every time series, as many rows as are needed to make the time series uniform.

The final code is about 50 rows long (comments included); an equivalent C++ implementation, probably, would have been much longer and complex. For sure, files I/O is much more verbose than the Python's with open(...) as descriptor syntax. 

I'm very far from a serious knowledge of Python's potential, but I'm impressed: the learning curve is astonishing, as is the clarity of the source code. I'm pretty sure that my future implementations that will not require some low-level functionality (eg., CUDA) or extreme performances will be Python-based.

From Protein Data Bank to... Blender!

Visual Molecular Dynamics (VMD) is a stunning piece of software. Not only it is a powerful tool for molecules visualization and rendering, but it is very useful tool for converting PDB structures into other formats that 3D modeling programs can read. In particular, VMD exports into OBJ, which probably is the most common interchange format.

Moreover, it is usable! A+

DNA and fractals

DNA is a very interesting molecule, not only for its peculiar role in life. It has a solid structure and the Watson-Crick complementarity allows the construction of complex structures, by exploiting the so called sticky ends, that is, non-paired single strands that urges to be coupled with their complementary ones. 

So, if we manipulate some DNA and give it a regular shape, we can "connect" multiple pieces like a nano-puzzle. Actually, we don't do anything as the DNA makes a self-annealing and spontaneously builds nano-structures, according to the shapes we defined and a starting seed, that is, an initial DNA sequence. More precisely, the process is driven by the encoding of the sticky ends, that decides which shapes can link to each other and how. 

It is very easy to develop a simulator, so tonight I did my own. As a test, I used the "pieces" and the seed that make the DNA construct a Sierpinsky fractal.  

My simulation of DNA self-assembly, using Sierpinsky tile set

I implemented it in a couple of hours, so my simulation algorithm is largely inefficient. For instance, it makes a "scanline" testing of all cells at each iteration, instead of keeping track of the free sticky ends in the assembly.   This means that the complexity grows with lattice's size and with the number of different pieces, too much for testing serious scenarios. Furthermore, it's a strategy that makes the whole process strictly deterministic, introducing a "bias" in the annealing process which proceeds in a very regular way (from bottom to top, from left to right). In order to reproduce all the various mistakes that can happen in the annealing process, a stochastic algorithm should be employed.

By the way, the result is interesting and fascinating. The overall technique resembles the cellular automata, in particular the fact that local rules lead to emergent phenomena, complex and unpredictable.

Using Nsight with a GTX 590

Update: Nsight Development Platform, that is, Parallel Nsight v2.2 finally introduces the single-GPU debugging capability. This is good and allows you to skip the following paragraph.

The only way to make serious CUDA debugging is by using Parallel Nsight by NVIDIA, but it has a drawback: it takes two GPUs. So, you can remotely connect to a GPU-powered machine or you can do it locally, by installing two NVIDIA video cards on the same host. I chose the latter, and bought a GTX 590, a dual-GPU video card that NVIDIA itself considered "not profitable" and consequently stepped down the deliveries. Two GPUs are needed because one of them actually runs the kernels, whilst the other one "inspects" the memory banks of the former. Because of this, you must explicitly disable the SLI functionality, which makes the two GPU a single one:

Disable SLI by clicking "Disable multi-GPU mode"

So, one GPU is going to run the kernels, and the other one is used to "look inside" the video memory. Because of this architecture, the OS can't use the same GPU that run the kernels, so you must connect your screen on the other one. If you want, you can explicitly choose which GPU to use with the method cudaSetDevice().

Parallel Nsight has another, implicit drawback: it is a Windows-only application. In particular, you are going to need the full version Microsoft Visual Studio (VS Express is not supported). Windows has a functionality named timeout detection and recovery (TDR) that automatically determines if something is wrong with the video card driver and resets it. As the debug activity basically halts the execution of the kernels, TDR must be disabled:

Disable TDR by setting WDDM TDR enabled to "false"

Microsoft created some time ago a framework for the creation of nice GUIs in the Windows environment: the Windows Presentation Foundation (WPF). Naturally, they mess up with Nsight, so you have to turn WPF off. It can be done by looking for the DisableWpfHardwareAcceleration.reg file in the "common" directory of Parallel Nsight ("C:\Program Files\NVIDIA Parallel Nsight 1.51\Common" on 32-bit architectures, and "C:\Program Files (x86)\NVIDIA Parallel Nsight 1.51\Common" on 64-bit architectures). Double click on the file to install it.

WPF is not the only thing that breaks Nsight's work: the Aero interface does not allow debugging too. You can disable Aero it by setting a non-Aero theme:

 

Are we done? Not yet! Remember to open the Parallel Nsight Monitor before debugging your kernels. Moreover, if you try compiling a kernel in Visual Studio and debugging it, you could receive a message that basically says that Nsight can't find your executable. You have to explicitly tell where your compiled files are: 

Fill the "Working directory" field with the absolute path to your compiled file.

Now you can enjoy the CUDA debugging features of Parallel Nsight! Yikes!