Raymond Tay

Scala's implicit need to be handled with care

2011-12-29T11:54:00.002+08:00

Quick post when using Scala's implicits. These constructs are really useful in the manner that they allow the user to create rather complex expressions in type design & bridging frameworks in the Java world. However, i did find some oddity when i was playing with it and found something interesting, to me.

Let's begin with a simple example.

scala> object holder {
| trait Foo {
| implicit val x = new Foo { override def toString = "trait foo" } }
| object Foo {
| implicit val y = new Foo { override def toString = "object foo" } }
| }
defined module holder

scala> import holder.Foo
import holder.Foo

scala> def method(implicit foo: Foo) = println(foo)
method: (implicit foo: holder.Foo)Unit

scala> method
java.lang.StackOverflowError
at holder$Foo$$anon$1.(:18)
at holder$Foo$class.$init$(:18)
at holder$Foo$$anon$1.(:18)
at holder$Foo$class.$init$(:18)
at holder$Foo$$anon$1.(:18)
at holder$Foo$class.$init$(:18)
at holder$Foo$$anon$1.(:18)
at holder$Foo$class.$init$(:18)
at holder$Foo$$anon$1.(:18)
at holder$Foo$class.$init$(:18)

The interesting thing is that the Scala runtime can't seem to make up its mind whether to invoke the "x" or the "y" when its being looked up in the call to "method".

Another example would be related to Scala's package objects. To see how the problem can manifest itself, here's an example with a file named package.scala in Scala's convention and here are its contents:

package object foo {
implicit def foo = new Foo
implicit def foo2 = new Foo
}
package foo {
class Foo {
override def toString = "FOO!"
}
}

when you run the compilation via scalac, it generates a directory foo and dumps the class files there. Nothing to it. Now run scala and use the implicitly function and you'll see what i mean.

Welcome to Scala version 2.9.1.final (Java HotSpot(TM) 64-Bit Server VM, Java 1.6.0_29).
Type in expressions to have them evaluated.
Type :help for more information.

scala> implicitly[foo.Foo]
:8: error: ambiguous implicit values:
both method foo in package foo of type => foo.Foo
and method foo2 in package foo of type => foo.Foo
match expected type foo.Foo
implicitly[foo.Foo]

This isn't Scala's fault, really; because the compiler can't possibly know what you plan to do with it but it certainly is something you should be aware of and i do think this is a good thing (rather than having the compiler just forget about it)

The idea is to really limit the scope in which the implicits exist so as to reduce the amount of pain you have to go through when debugging

Scala's type erasure in writing functions

2011-12-18T11:28:00.000+08:00

Wanted to write a quick post to anyone whose out there interested. The topic of this post has to do with type erasure in Scala. I won't dwell on history but to get an idea what it is (other than googling or ref the wikipedia) you should read this well written article here.

I like to use an example to illustrate something here. Whilst playing with Scala, i find myself writing code like this

class MyTest {

def printSomething = println("blah blah blah")

def compute1Value = 41 * 41

def compute2Value = (41*41*2, 1)

}

object SimpleTest extends App {

def test1(a : {def printSomething: Unit ; def compute1Value: Int; def compute2Value: Tuple2[Int, Int]} ) {

.... // do something

}

I thought nothing special here. Compiled it and it works. But i thought i should be able to write something more generic than that because the structural subtyping technique here is REALLY useful in real world applications but there're caveats because of type erasure (this post isn't about that)

So i proceeded to modify the type of 'a' in 'test1' to become like this

...

def test1(a : {def print1Line: Unit; def fn1:Function0[Int]; def fntuple2: Function0[Tuple2[Int,Int]]}) {

.. // do something

}

But this didn't compile. Hmm...(honestly, WTF was the word that came to my mind) what happened. Afaik, it should work because i have a valid function definition on both LHS & RHS of the equation. However, after figuring out what happened by examining the java bytecode output i realized it and the following definition works

...

def test1(a : {def print1Line: Unit; def compute1Value:Int; def compute1Value2: Tuple2[Int,Int]} ) {

... // do something

}

Mystery solved. Don't get me wrong here but you should probably be able to find the answer by examining the Scala Language Specification & still i chose to de-assemble the code and from the output below you should be able to guess that the return type was determined and instead embedded into the signature of the function.

...

public int compute1Value();

Code:

0: sipush 1681

3: ireturn

public scala.Tuple2 compute1Value2();

Code:

0: new #24; //class scala/Tuple2$mcII$sp

3: dup

...

Alternatively, the following code works too

class MyAnonFnClass {

def print1Line = println("blah blah blah")

def fn1 = new Function0[Int] {

def apply() = 41 * 41}

def fntuple2 = new Function0[Tuple2[Int,Int]] {

def apply() = (41 * 41 * 2,1) }

}

object SimpleTest extends App {

def test1(a : {def print1Line: Unit; def fn1:Function0[Int]; def fntuple2: Function0[Tuple2[Int,Int]]}) {

... // do something

}

This definition of "test1" is what i like better but the definitions of "fn1" and "fntuple2" is somewhat awkward. Still, some say it doesn't look right ... there're better ways of writing this and an example is that of James Iry at stackoverflow

GPU development needs a push in Singapore

2011-10-12T09:29:00.002+08:00

Recently, i got to know of a event where geeks would gather for 1 day every year. Turns out that Geekcamp was held once every year and some companies would provide sponsorship either providing a venue, meals, refreshments, collaterals etc and i thought this year, CUDA should be given some highlight. So i went and gave a introductory talk on the subject and the slides are here. Glad to see many enthusiasts and heartened by their questions.

Lastly, i noticed that one of the organizers took photos of the event though its on flickr.

OpenCL Programming Guide - A book review

2011-09-18T10:05:00.000+08:00

Hi everyone,

This post is about a book i helped reviewed during this last month. This book is called OpenCL Programming Guide by Addison Wesley where one of the co-authors is the main guy behind OpenCL, Dr. Tim G. Mattson.

First, let me begin by what i feel about the book and then i'll proceed to give a more detailed review. When i first saw the book, i thought that it was a rather large book (approx 600 pages) for a relatively young technology as OpenCL and the reason for this was that the book was split into 3 major sections: Section 1: All about OpenCL, Section 2: Case studies, Section 3: OpenCL API guide. What i next felt was that there was a lack lustre when it came to the overall cover design of the book but the layout of the contents is good, clear and concise. Fits well into the human hand and i reckon can be carried around the bus, library, shopping mall whenever you feel like reading. I thoroughly enjoyed the case studies since those are real world applications of OpenCL and where details are desired, the appropriate reference are provided.

Overall, i think this is a good guide for programmers in OpenCL, CUDA and the general audience would benefit, too. Thanks

This book covers the OpenCL 1.1 specification in its entirety which is a good thing. I don't think i'll cover all the chapters but i'll point out which chapters sort of stood out, for me.

Chapter 1:

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

What i liked about this chapter was the cautionary note to readers that implementation of work-groups (including execution concurrency) in OpenCL is largely vendor dependent. This marks a great difference between the equivalent in the NVIDIA's CUDA model. I also liked the fact that the authors reminded the readers that they need to be aware of the IEEE 754 floating-point support issues - they did make known that it would suffice for current heterogenous platforms. The authors took a painstaking task, i believe, to lay the ground for OpenCL by going through the core concepts through clear explanation. I thought it was a pleasure to read :)

What i didn't like about this chapter was that there was too much write up about the API, perhaps they could illustrate a few and point the readers to the specs for details?

Chapter 2:

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

This chapter begins with a sample OpenCL in the classic "Hello World" application and HOWTOs configure 3 IDEs to use OpenCL. I think this was helpful for folks who come to OpenCL having used Eclipse, CodeBlocks and Microsoft Visual Studio. However, i didn't quite appreciate the lengthy explanation of that example but still i think it has value because the book assumes you don't have experience in using OpenCL so i think it still has tremendous value.

Chapter 3:

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

This section also suffered, imo, from elaborate explanation of the APIs ... in another word, kind of boring but still nonetheless important. It would serve you well when you hit those bugs because you didn't read the API docs well. However, this OpenCL 1.1 reference card is handy (http://www.khronos.org/files/opencl-1-1-quick-reference-card.pdf)

Chapter 4, 5, 6:

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

This is the chapter i liked. They go into great detail on the data types in OpenCL like the "vector", "half" and illustrate how to manipulate them through code in snippet form. They talked about implicit type conversions among types, IEEE 754 rounding modes and the usual arithmetic operations you can apply to the new data types in OpenCL. Something that caught my eye was the fact that there isn't short circuit evaluation for "vector" data types. This chapter also introduced kernel attribute qualifiers necessary for compiler optimization. They talked about the kinds of memory spaces allowed in OpenCL and i liked the fact that they pointed out that pointer types cannot be re-interpreted across different memory address spaces. Cool.

Coming back to vector data types, the authors didn't show whether binary operations applied but they did show what'll happen when scalars are applied to vectors.

Atomic/synchronization functions are introduced in chapter 5 and they more or less look similar to CUDA's synchronization primitives.

In Chapter 6, i liked that the authors mentioned that you can load NVIDIA's PTX into OpenCL. Interoperability but i suspect that's only pertinent to the OpenCL driver provided by NVIDIA.

Case studies

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

Totally love this section of the book where many examples of real-world, working OpenCL are applied. This is where the money's at, i think.

Unwind A Program's Stack

2011-09-04T16:21:00.000+08:00

Memories

Recalling my fond days in taking CS101 more than a decade ago, one of the critical concepts i had to learn was to understand how to debug a program. As time went on, i became interested in compilers and the technologies used and took up Compiler Writing course. What a journey, i'll do it all over again and this time i'll do better ;)

In this post, i wanted to share with you how you can really unwind the program stack using a C library, libunwind. Really cool stuff - wish i knew this earlier. Anyways, here goes...

Question Why do i ever need to unwind the program stack programmatically?

Answer This shouldn't come as a surprise since in many programming languages, you would have added a catch exception clause that basically is unwinding the stack (when it encounters an error) but what this library does for you is that you can unwind the stack under more favourable circumstances other than receiving an exception ;)

libunwind - is exactly the library you need to build this sort of support into your programs. Imagine how cool it'll be if your friends can debug your program other than GDB - they'll love you!

Building libunwind & issues

Download libunwind or clone it thru Git.

When you first build/install the software, you'll likely encounter two problems

/usr/include/bits/setjmp2.h:26:13: error: ‘longjmp’ aliased to undefined symbol ‘_longjmp’

libtool: install: (cd /home/tayboonl/libunwind/src; /bin/bash /home/tayboonl/libunwind/libtool --tag CC --mode=relink gcc -U_FORTIFY_SOURCE -fexceptions -Wall -Wsign-compare -XCClinker -nostartfiles -version-info 0:0:0 -o libunwind-setjmp.la -rpath /home/tayboonl/ray-opt/lib setjmp/longjmp.lo setjmp/siglongjmp.lo x86_64/longjmp.lo x86_64/siglongjmp.lo libunwind-elf64.la libunwind-x86_64.la libunwind.la -lc )
mv: cannot stat `libunwind-setjmp.so.0.0.0': No such file or directory
libtool: install: error: relink `libunwind-setjmp.la' with the above command before installing it
make[3]: *** [install-libLTLIBRARIES] Error 1

At first i was stumped, then after browsing through the net, i've discovered that the main cause was because the default optimization level applied (-O2) actually triggers the GCC on the Ubuntu 11 to conduct security checks and it thinks that libunwind has flouted the rules. Admittedly, i'm not too concern at this point in time over security since what i wanted to do is implement this library ASAP and use it. Using the Ubuntu's wiki on compiler flags, circumventing it was to apply the -U_FORTIFY_SOURCE (which stands for unknown fortify source) to the build.

One thing you'll immediately notice is that the optimization level (-O2) is no longer applied. Just be aware of this.

Second thing i've discovered is that a critical library cannot be linked and the hack i've provided below on LDFLAGS will mitigate the shared library creation error.

./configure CFLAGS=-U_FORTIFY_SOURCE LDFLAGS=-L<full pathname>/libunwind/src/.libs --prefix=<prefix directory>

Last thing is that building the program requires that config.h header file be present. Then i realize that this file has been relieved from the latest Ubuntu 11 kernel and then checking through a couple of forum posts, blogs i realized that it could be hacked by simply creating an empty file i.e. touch /usr/include/config.h - reason for double checking this is because config.h has been used in various linux releases to define variables (i guess that's why its called a configuration header file) and building any simple program that uses the library is something like this

gcc (prog name.c) -I(install dir of libunwind) -L(location of libraries of libunwind) -lunwind -o (prog name)

A note of caution here is that you really shouldn't leave the config.h file you just created lying around like that.

Here's a simple program that shows you how to create a show backtrace command in C.

#define UNW_LOCAL_ONLY
#include <libunwind.h>
#include <stdio.h>

void show_backtrace(void) {
	unw_cursor_t cursor; unw_context_t uc;
	unw_word_t ip, sp;

	unw_getcontext(&uc);
	unw_init_local(&cursor, &uc);
	while( unw_step(&cursor) > 0) {
		unw_get_reg(&cursor, UNW_REG_IP, &ip);
		unw_get_reg(&cursor, UNW_REG_SP, &sp);
		printf("ip = %lx, sp = %lx\n", (long) ip, (long) sp);
	}
}

void simple_func_call_no_args() {
	int i, j;
	print("value of i=%d, j=%d\n", i, j);
	show_backtrace();
}

int main(int argc, char** argv) {
	simple_func_call_no_args();
	return 0;
}

Cilk

2011-08-23T14:37:00.000+08:00

Cilk. Not Silk.

This is another parallel programming paradigm from MIT where the design goal was to allow the programmer to concentrate on expressing parallel constructs while the Cilk runtime takes care of the scheduling etc. Cilk has a commercial edition and was acquired by Intel in 2009 but recently, Intel open sourced Cilk Plus and available in the "cilkplus" branch in GCC 4.7 so i'm expecting a lot of activity in the next 12 months.

Issues getting Cilk to build on your operating system

Ubuntu 11
This OS comes with the GNU compiler suite version at 4.5 but there are issues while building the Cilk compiler, otherwise known as cilkc. The issue is that the Cilk-C compiler doesn't like unnamed bit fields though this is perfectly legal with typical C compilers on Mac OS X and Ubuntu 11/10.

Solving this unnamed bit fields is actually easy. Just give it a name and make sure there's no name collision. There is 1 file you need to take care and you can just give a unique name to the unnamed bit field and re-build your project.

/usr/include/bits/waitstatus.h

Mac OS X

This OS comes with a pre-installed GNU compiler suite customized for Apple at 4.2.1 Again it faces the similar problems as encountered in Ubuntu 11. This time round, the files are different and that's owed to the fact that although they are both Linux-es, their kernel files are arranged in another manner.

Fix the following files by providing a dummy name that's unique

/usr/include/stdlib.h
/usr/include/mac/i386/_structs.h

Another issue concerns the concept of "restrict" pointers and how Cilk-c compiler ignores it on the GCC 4.2.1 on Mac OS X. You can go read up about it on Wikipedia or your favourite C authors/books but the observation here is that Cilk doesn't understand restricted pointers in the C API memcpy(...) when its attempting to compile its example cilk programs. Looks to me like a self-test where a compiler This obviously can be a TODO for an enthusiast but the simplest thing i did was to ignore the build of those example programs but turns out skipping them is a bad idea as the final binary "cilkc" cannot be generated.

Drats.

The simplest solution is to replace the "offending" functions with their equivalents in "standard" C. e.g. coding a for-loop to replace memcpy(...) & strcpy(...)

Having that done, cilkc was generated and its static/dynamic libraries are generated.

Hopefully this helps and the next couple of posts, i'll write something up about using Cilk

MongoDB and JOINs

2011-07-18T21:27:00.005+08:00

MongoDB's Great!

MongoDB's a popular NoSQL database that i've heard of for the last 4 months or so but never had any time to consider doing something until last week. I guess the lack of enthusiasm has a lot to do with the fact that 'traditional' databases like Oracle can be quite a RPITA to work with but they've withstood the test of time - though i admit things are changing pretty fast these days. MongoDB is really a good database tool but it's got a catch!

MongoDB's Great BUT ...

This catch was a direct result of my fictitious assumption that MongoDB supports the concept of multiple table queries a.k.a SQL JOIN e.g. a SQL expression below illustrates

'SELECT CAR.plate, CAR.owner, Fines.owed from CAR, Fines where CAR.owner = Fines.doneBy'

which would retrieve all records done by a particular owner when he/she ran a traffic fine. Contrived really but i quickly realized it couldn't be done with MongoDB, at least not now anyway. I was kind of annoyed but forgiving for reasons that i've come to expect the same lot with Oracle Databases but realized MongoDB != OracleDB.

MongoDB is different. Accept that. MongoDB is simpler. Accept that too.

For someone whose played with MongoDB would realize that its operations evolved from the concept of collection & document; w.r.t SQL parlance its table and row respectively. You thought this is the only difference it has? You're wrong. There're others i've discovered which includes Stored Javascript (akin to Stored Procedures), Indexing etc which are freaking cool and the data structures that represent all forms of data you can store or retrieved from it is based in BSON. For web folks, you know what this means don't ya? ;) for everyone else, it simply means you can get your apps into market at a faster rate.

Back to what you were saying about JOIN and MongoDB ...

JOINs are pretty important in a lot of applications since everything we know is relational, in SQL parlance. SQL databases have taught us that all data is relational and we need to think in terms of tables and rows. If a relationship needs to established, then we need to select data from multiple tables and derive it.

Question is whether we can get around it? Spending 2 hrs mucking around MongoDB's documentation revealed two things (1) MongoDB's MapReduce (2) write one yourself.

Obviously i attempted the easier solution and a member on stackoverflow seem to think its possible. But i ran aground. To understand why, the MapReducer's reducer function has a primary purpose and it is to reduce the size of key value pairs the system has to handle at any point in time and hence it prohibits code like

// aggregating values to the array
function(k,v) {
  var ret = v;
  return v;
}

// a.k.a bypass
function(k, v) {
}

MongoDB will complain that it doesn't like the fact your reducer does nothing, and also aggregating values to an array since it does absolutely nothing to reduce the key value pairs. There might be other ways to do this but at the time of this writing, i'm not aware of any (Please drop me a mail if you do). Turns out somebody has filed a feature support ticket to sort out this JOIN. Yay!

Parser for MongoDB multi-collection query

I had no choice but to attempt option (2) which was fun! cos i get to built a simple SQL parser and translate the typical SQL syntax to a multi-collection/table query. The approach i took is something about the line that i like to use a SQL-like syntax that supports multi-table/collection (depending on your parlance) and allows me/user to write something like

select a.name, a.data, b.name, b.data, c.name, c.data from a, b, c where a.name = b.name and b.name = c.name

while underneath the machinery will conduct the query against MongoDB. The implementation approach is pretty simple where i parse the expression and isolate all expressions belonging to each table and keep it in a hashtable like data structure. What happens next is that each table query will be sent against Mongo and retrieved records are stored in RAM.

As Mongo can't do multi-joins yet, the approach is to conduct a O(N^N) where N's the number of tables operation where the JOIN query is run against each known bunch of records, in turn i.e. for a for b for c, and finally apply the SQL LIMIT expression to return the final collection of records.

There are obvious short comings of this, i don't claim to have the answers to them. One shortcoming is that it doesn't scale very well as its limited by the RAM/DRAM available and its a compute intensive job. And if you're interested please visit https://github.com/raygit/m2query All feed back is welcome!

Announcing the launch of HPC SIG Group in Singapore

2011-06-17T09:23:00.001+08:00

Hello everyone,

Birth of a Special Interest Group in Singapore

I just wanted to take this post to announce that myself and another fellow CUDA enthusiast have started a Special Interest Group (SIG) in Singapore that focuses on High Performance Computing which we affectionately call HPCSIG Singapore.

Group's Focus

This group's focuses on gathering administrators, researchers, developers or anyone whose interested in HPC to come together to share and learn new knowledge. This includes technology topics in Fortran, C/C++, OpenMP, MPI, CUDA, OpenCL, Python, Ruby etcc

Meetup

Our first meetup was held on 8 June 2011 at the Institute of High Performance Computing at
Resonance Room, level 13, Connexis North Tower 1 Fusionopolis Way Singapore 138632 and we had about 20 attendees and 2 talks:

Group's Information/Forum

https://groups.google.com/group/hpcsig-singapore?hl=en (Google Group)
http://www.facebook.com/home.php?sk=group_210520412314382 (Facebook)

Speakers for First Meetup

Simulations for Maritime Industry in SimPlus and the Need for HPC
Mr. Ye Rong
CTO
SimPlus Pte Ltd

Introduction to CUDA Programming
Mr. Raymond Tay
Software Engineer
HP Labs, Singapore

Participants

And we had the pleasure of Dr. Simon See (Chief Solution Architect of Nvidia), Dr. Rich Goh (Director in Institute of High Performance Computing) graced our first meetup. In the lineup were a HPC engineer who helped engineered the special effects for the Lord of the Rings Trilogy, other HPC engineers whom are currently residing in Singapore helping to bring forth this movement.

What's next?

The meetup for July is likely to be canceled as many of the participants are busy with their respective commitments so it would not be in the interest of the group to meet but August is on!

Founder Co-Founder

Mr. Chung Shin Yee (Creator of SGC Ruby CUDA)

Mr. Raymond Tay (Cloud & GPU Developer)

Exploring CUDA through Google's PageRank (Part 1)

2011-06-11T10:35:00.002+08:00

Page Ranking

PageRank is new to me. In fact, i've only known about it some time in mid of 2010; talk about coming to a scene pretty late. For those of you whom don't know what PageRank is about, i'll let Wikipedia tell you all about it or as Google tells it "... a technology that determined the 'importance' of a webpage by looking at what other pages link to it, as well as other data." Talk about 'peer' appraisal on the exascale. I do apologize for overloading the terms to the point of abuse.

What i do like about the PageRank is that it takes into account two aspects of human beings when it comes to sharing info:

1) People just love to share information through the internet; thus enabling others to receive that information rapidly;

2) People will eventually tire of sharing probably at the end of a day...before lunch, before a nap etc

I just love the sort of things people do with math, here's another one where mathematicians can describe a city using equations. Ok, that should be enough digression.

Why i did this?

Reason is because i asked myself : Can i model the PageRank algorithm into the GPU so that i can use the super computer under my desk to compute how popular a web page really is?

First of all, take some time to understand the mathematics. Frankly, i'm no math major but it seems pretty alright since the equation is

Trusting the wikipedia *makes the sign of the cross* and confirming that equation with my favourite algorithm texts, i was assured the math is correct. Whew. From the documentation, i understood that this equation needs to be iterated a number of times so that the results can be more reliable. Those folks in HPC shouldn't have much problem I imagine but how much a novice like myself go about doing this?

Modeling the problem ... some ideas ...

Few considerations i can think of the top of my mind:

Mine the Internet (that's crawling the web, really and for ease of programming i've used Python)
Represent the mined data into a graph (i use the adjacency matrix for starters)
Decide a damping factor (that's the d in the equation and i found that 0.85 seems to be the accepted value from here and what 0.85 says is that the random surfer has 85% chance of clicking on ANY link and you are free to test different models by adjusting the damping factor)
Decide a number of iterations you are going to use to compute the probability distribution. From this fella's description in his/her post it appears that 100 iterations should suffice (Did i mention there would be hundreds of millions of pages?? That'll give your CPU a good run for the money you've paid)

Those are just some of the stuff i've thought about, definitely not exhaustive.

I'm assuming you know how to do a web crawling in multi-threading/processing form (note that the liberal use of 'thread' and 'process' differs from 1 programming language to another)

Now, to the implementation and benchmarks!!! Btw, the simulations i have is in the order of hundreds of millions of web links (some cyclic, some not) and running it on my mac book pro didn't work well as it crashed and burned as i didn't have enough RAM and for best results i've used by GTX 480 FERMI 2.0

Results of performance tests and optimizations applied

There are a couple of scenarios in the tests i've ran. To give you an idea what the following sub-sections is going to be about, i'll like to show you the benchmark results off the GT 330M Nvidia processor and GTX 480. One key thing i wanted to point out is that the two devices of different compute capabilities DO make a difference.

Results of test runs on GT 330M Nvidia processor (Compute Capability 1.2)

Having to cut back on the size of my data as it crashed much to my horror.

Moving forward, the tests are modelled based on 4 million web links, the average run times is taken from 10 test runs discarding 1 or 2 runs depending whether the deviation is too wide. The CPU code was compiled to level 2 optimization and GPU code was compiled using the production release of CUDA 4.0 SDK.

Version 1 (straight forward):
The average GPU run time (over ten runs discarding 1) = 7850.5 ms
The average CPU run time (over ten runs discarding 1) = 22187.89 ms
The speed up over the CPU version is = 2.83 times.

Version 2 (using shared memory, loop optimization etc):
The average GPU run time (over ten runs discarding 1) = 3033.8 ms
The average CPU run time (over ten runs discarding 1) = 20412.2 ms

The speed up over the CPU version is = 6.7 times

The speed up over the two GPU versions is approximately 2.4 times. The key improvement was to make use of the device's shared memory so that it can be as close as possible to the GPU considering the latency of global memory access is between 400 - 600 cycles.

The discrepancy in the CPU run times is probably due to a Mac OS X system process running in the background and also i was listening to coding music.

Here's a glimpse of the device's configuration i used calculated by CUDA's occupancy calculator (an invaluable tool in your arsenal); see below for the diagram for GTX 480 (CC 2.0)

Results of test runs on GTX 480 Nvidia GPU (Compute Capability 2.0)

Consume all of shared memory on device (no L1), compiled for sm_20, use fast math

The same code used in Version 2 previously is tested here compiled for -Xptxas -dlcm=cg -use_fast_math, -arch=sm_20 and the run times shown here was fine tuned with the help of the visual profiler.

The average GPU run time (over ten runs discarding 1) = 1176.9 ms
The average CPU run time (over ten runs discarding 1) = 20362.6 ms
The speed up over the CPU version is = 17 times

Consume all of shared memory on device (no L1), compiled for sm_20, use fast math, texture memory

One last ditch effort came in the form of using texture memory, specifically speaking linear texture memories, on the GTX 480. Reason for poor run is because i've only managed to achieve a 3% hit rate in the texture memory. That's really because of scatter reads and writes; hence the REAL problem is actually the kernel has significant scatter read / write - so it doesn't really matter if i used texture memory or not because the problem's still there.

The average GPU run time (over ten runs discarding 1) = 1142.1 ms
The average CPU run time (over ten runs discarding 1) = 20092.3 ms
The speed up over the CPU version is = 17 times

Conclusion

I can't really conclude this at this point in time since the kernel hasn't reached its maximum potential but a couple of things are useful to take note of and possibly worth sharing:

Choosing a graphics card that supports Fermi or later should be a good bet i.e. compute capability 2.x
Attempt to reduce global memory bandwidth by placing data as close as possible to the processing units e.g. shared memory, L1/L2 caches, Texture memories, constant memories.
Remember to compile your CUDA kernel code for architectures that best match your graphics card e.g. -arch=sm_12 -arch=sm_13 -arch=sm_20 since the default would be sm_10
Experiment with different ways to implement your kernel. Always think about improving memory access.
If you are determined to get the best performing code, then i'd recommend you get really familiar with the visual profiler and understand the metrics that's being measured and how each strategy you applied worked best.

Roadmap ahead

This kernel isn't optimal definitely (see diagrams above which implies low occupancy). I've plans to rewrite the kernel (after hours) and run some more tests - i'll update it in a separate post when i can.

Exploring CUDA through computing Coulomb's Law

2011-04-27T09:48:00.000+08:00

This is more like a physics post than anything else. I loved physics when i was doing my secondary school (that's high school in some countries) til pre-university days and at one point in time - i decided i was going to be a physicist! but i chose Computer Science. Well, that's a tad short bit of history for you anyway.

So, this post isn't how i DIDN'T make it to be a physicist but rather an interesting application (admittedly, contrived) of CUDA into physics, namely the computation of force vectors between sets of charged particles.

Some basics in Physics is in order ...

Coulomb's Law
Atoms and molecules have the same number of protons as electrons and are neutral (without overall charge). Electrons can be transferred from one object to another. When this happens, there is an excess of electrons in one place and a deficiency of electrons in another. Charge is the result of an excess or deficiency of electrons. Where an object has excess electrons, the object is negatively charged. Where there is a deficiency of electrons, the object is positively charged. Electric charge is usually represented in equations by the letter q or Q.

So, the next question is to ask whether 2 charged particles next to one another would "feel" a force of attraction or repulsion? (FYI: 2 charged particles would attract one another if they had opposite charges and repel one another if they had same charge.)

In fact, they do and this felt force is computed via Coulomb's Law. Two charges, Q and q, separated by a distance, r, each experience a force of magnitude

F = k|Qq|/r² where k = 9 x 10⁹ and |Qq| is the positive value of the product of Q and q

The above formulation is what is called the scalar form. The k is known as the Coulomb Constant. Here's an image i dragged off Wikipedia to illustrate how charged particles can be thought of in their scalar & vector forms:

Force Vectors

The takeaway from this picture is that (1) like charges repel (top row) (2) opposite charges attract (bottom row) and (3) each repulsion/attraction can be expressed as a vector in 2D (You can extend this to 3D if you like). This is something i can imagine quite well but when i attempt to extrapolate it to say 6 x 10 pow 23 or some other big number; then it becomes a bit mind boggling. Question that i have is whether i can find some way to compute the electrostatic force from this imaginary large number of charged particles?

Turns out you can by applying the Law of Superposition to Coulomb's Law over large numbers of such particles and the net effect is saying that the net force on any particle, p , is the vector sum of the other charged particles acting on this particle, p.

Modelling the problem

To model it, i need to create 2 or more sets of charged particles and have CUDA compute their net force by applying both laws. Each particle can be represented by a float4 type (basically, a element that has 4 ordinates x,y,z,w) ; each particle would be given a set of coordinates in 3D populated into their respective x, y, z and w will become the charge.

The parallel portion is shown below and in case you are wondering about force_calc; don't worry about it because its just the computation of Coulomb's Law which you can find samples of anywhere.

Side Note:

Another important thing you should remember (or probably know) is that we should always create the exact algorithm that'll run on a CPU (single core preferably). This serves the purpose of creating a reference implementation which you can debug your parallel stuff against and also to determine what kind of optimizations you need when you compare their run times.

done with host force eval kernel
host force eval took 157.8 ms
done with gpu force eval kernel
gpu force eval took 33.3 ms

Summary:

Think you get the idea now :) also from the output, you can see that the run times is a 4.7 fold drop :) and that's w/o any kind of optimization (you can tell right?) and a couple of optimizations you can apply and that's one simple application of CUDA to high performance computing.

Automation - My thoughts ...

2011-03-16T22:29:00.001+08:00

I like to take this post to share my thoughts on Software Automation.

Why the need for Automation?

This topic has been explored by many authors in books and blogs. Automation exists in many industries so don't think that its a "IT" thing. The cookies you are enjoying right now, the milk you're drinking to the foods you are eating are highly automated. Automation, in my opinion, arose from a human trait - being more effective and productive.

I'm not going to write about industrial automation since its not something i do from day to day back then though i suspect they share similar traits; so i'm going to focus on automating software or software automation and there's an area called "automated software testing" which is a subset of software automation,imo.

One reality of Automation,i realized, is that it'll take time to setup the necessary processes, applications and would draw some of your best QA talent from their daily work but the benefits you will reap will far outweigh its initial drawbacks for the simple reason that you would have freed up the resources you previously held and hence you have more QA resources to work on the really big problems that need attention.

What made me decide to do Automation from QA?

Quite some time back, a company i used to work for recognize that over a long run QA couldn't effectively run its business when you consider a server solution that supported tens of thousands of 3D-clients connected to a back-bone of 10,000 servers.

It was an exciting opportunity for me, personally, since white-box testing was a largely manual process but there were many other QA tasks that could be automated and hence raise overall efficiency and i wanted to explore that.

What do you automate?

This is tricky stuff. Some folks would like to automate everything but that is not practical. Focus on the problem you're trying to solve that is not is not easily solvable but if you pull it off, you can solve all other problems.

Many of the stuff you can list as potential candidates often manifest themselves as problems you've been solving repeatedly so a couple of examples could be:

Application startup, shutdown
Application deployment (I know some orgs actually still rely on human efforts)
Application build and release
Application build, test and release
...many others

Here's what we did on determining that:

Back then, QA implemented many processes w.r.t testing, coding standards, code reviews and the like and it was a rigorous effort. But still bugs are reported by the users and that means that our test coverage was inadequate. The answer to this question is not simply to add more QA resources but to solve the problem in a smarter manner...

Don't get me wrong here, there was automated testing implemented but neither QA nor the Devs had any visibility into how the tests relate to one another, how they vary from release to release and its a tough thing to keep track of considering that you have tens of thousands of tests cases. QA couldn't effectively exhaust all possible testing scenarios.

In effect, what we/QA/Dev wanted was something like this
1) Build the app and run automated tests on a platform type. e.g. Win32, Win64, Linux32 etc; report success/failures
2) Pre-release (automated) testing to a controlled environment; report success/failures
3) Collect & archive test metrics for trending purposes (afaik, it wasn't implemented by the time i left that company)

The end result of a series of brain storming sessions revealed that we needed a new way of doing things! Continuous Integration and SCM tools help realized 1) and 2) :)

Source Control Management - How does that come into the picture?

A big part of CI is to set up several gateways where the quality of the code can be checked before the actual build happens. As devs, we all know that when your code base expands rapidly, the time it takes to build an application becomes longer (not to mention source code locking) so it's best not to check-in code that doesn't compile or work else you'll incur the wrath of your co-workers.

Hence, the first gateway you should consider is scrutinizing the code that's being checked in. Modern SCM tools like Mercurial have a simple regression test framework in python & shell-script driven testing that runs automated tests for your code. Another example would be SVN, read about it here.

Check-in code -> SCM runs first-line defense -> rejects code if tests fail, accepts code otherwise.

SCM is not enough ...

Of course, this scrutinizing is somewhat limited in this effectiveness - so you need to ensure that devs check in tests that accompany the codes. This is TDD or Test Driven Development and back then, the product was a mix of C++, Python, Perl and so you can imagine that we used lightweight regression test frameworks to incorporate tests.

Alright, you say, now you've seem to have all the tests incorporated but how does that solve the problem? Who will run the tests and report the results? Before that, how does anyone know that code's been checked-in and builds need to happen?

You need Continuous Integration.

Continuous Integration - How does that help?

Martin Fowler said it best. Click this to read about it. Key thing to remember about CI is that frequent code check-ins trigger automated builds and that flashes errors faster than the usual build-test-release cycle you are used to.

A typical CI tool would allow have the capability to pull sources codes out onto its base, compile and run the automated tests you've built in - of course, tests results would be published to the tool's dashboard and alert the QA/Dev that a build/test failed and require attention.

Another thing that a CI tool would do is to allow you to upload modules that measure and informs you the coverage of your tests!

That is a really cool thing to have :) because a tool like this allows both QA and Dev to see the same problem in the same language.

CI Tool detects code checked-in -> runs build -> runs automated tests -> runs coverage tests -> reports build pass/fail & publishes test artifacts to dashboard.

But is that all it can do? Fortunately, no. Read on.

Extended Automation Framework via Python

Previously, i mentioned that we had a CI tool that builds, tests the app. In addition, the tool archived/published test results so that Devs/QA can determine which code check-ins caused a particular test/build to fail.

Now, i want to share with you something else we did. Part of the problem we had was that black box QA were quite involved in manual testing of the 3D client because there was no automated way to interact with such an application, even less report errors in a concise manner.

In a nutshell, we build a Python framework that allows any Dev/QA to control the 3D client in Linux, Windows & Mac OS X. The idea was to allow our QA/Devs to write automated test scripts in python so that the manual testing could be relieved. The framework came with a small set of APIs back then, but i'm sure it has grown much bigger these days.

Can you imagine how our test coverage blew sky high?!!!!

That was a tough problem to crack but we did it :) and that's because we had a bunch of dedicated and passionate engineers working on it.

But, that's not the end of the story ;)

With such a framework in place, QA now had the capability to run automated tests with a freshly build app (from the CI tool) in an environment of their choice. So, it resembles something like this:

CI Tool -> builds app -> test the (native) app through automated test scripts in Python -> publishes test results back to CI Tool's dashboard.

So that's it.

Installing CUDA 4.0 on Ubuntu 10.10 (watch out for a minor detail)

2011-03-10T15:31:00.000+08:00

Just a quick gotcha when installing CUDA 4.0 onto your Ubuntu 10.10.

If you recently encountered a situation (like myself) whereby you've downloaded the latest CUDA 4.0 Toolkit from Nvidia together with an Ubuntu image upgrade, rebooted your machine and discovered that your lovely X server doesn't start and complains of API mismatch then this is for you.

Re-install the downloaded device driver (i.e devdriver_XXXX.run) and it should work again (w/o rebooting) by issuing "sudo service gdm restart" or "sudo service gdm start"

Technical reason is simple: I installed the device drivers onto an existing image (while the newly upgraded image hasn't been loaded) and when it restarts, you would naturally boot into the upgrade image and the X server's like WTH ? and you're like WTH ?

Exploring CUDA via the Euclidean Distance

2011-02-02T10:25:00.002+08:00

So this post is really inspired after reading Chapter 1 of Toby Segaran's book "Programming Collective Intelligence". In that chapter, Toby talked about making recommendations based on the ideas behind the Euclidean Distance and Pearson Coefficient.

I thought of doing a simple exercise and see stuff works out in CUDA. Here's a sample implementation of the Euclidean Distance and conducted benchmarking both on a Nvidia GTX 480 and GT 330M.

The implementation is pretty straight forward since the block that is suitable for massively parallelism is easily identifiable and here's how. Using the wikipedia's article on Euclidean distance, we focus the attention on this equation:

Reading the equation from left-to-right or from right-to-left, you can tell that its basically performing computations on two 1-dimensional arrays, namely P and Q aka Euclidean 2-space. From there, you can easily tell that CUDA's technology is made for these forms of computations and best of all - you'll probably never encounter the problem of divergent computation and you can tap into the raw computing power of your GPU since the GPU's constantly busy executing as occupancy's close to 1! that's pretty cool in my dictionary :)

One implementation that's pretty straightforward to implement is to model these 2 spaces into 2 different input arrays i've named P and Q, coincidentally. If you have read the documentation or books in CUDA, you would know the next few steps is to allocate memory on the device and design your GPU kernel to run these computations. The pseudo algorithm for the GPU kernel is something like this:

1. Load input arrays, 'P' & 'Q', into shared memory
2. For each input element in 'P' & 'Q' in the shared memory, load them and  
   perform the computation of raising the result (i.e. taking the difference
   of 'P' & 'Q') to the power of 2
3. Store the previous result to the device memory of the GPU

. You'll notice that this GPU kernel's not taking the square-root of the result in this kernel - it'll be done in the typical C setting i.e. computing in the host and not on the GPU and the reason is due to the mathematical nature of performing square-roots since sqrt( a + b ) != sqrt(a) + sqrt(b)

One gotcha you need to be aware of is that on devices such as those with compute capability 1.2 and lower (e.g. yours truly's mac book pro running GT 330M) the double-precision FP computation is downgraded to single-precision FP. You'll notice a warning message like "warning : Double is not supported. Demoting to float"

Next question: How do i/you do all this with Thrust?

Thrust is the library you should probably use when it comes to developing stuff with CUDA. It just makes things easier. Here's one implementation with the CUDA Thrust that i've made - one thing you need to take note is that if you conduct this test on a machine that has compute capability of 1.x then you'll get the warning message i've highlighted above; otherwise it should work well on device's with capability of 2.x

1 #include <thrust/host_vector.h>
  2 #if __CUDA_ARCH__ == 200 
  3  #include <thrust/device_vector.h>
  4 #elif __CUDA_ARCH__ == 100
  5  #include <thrust/device_ptr.h>
  6 #endif
  7 #include <thrust/generate.h>
  8 #include <thrust/reduce.h>
  9 #include <thrust/transform.h>
 10 #include <thrust/functional.h>
 11 #include <cstdlib>
 12 #include <iostream>
 13 
 14 struct powfunctor
 15 {
 16     __host__ __device__
 17     float operator()(const float& p, const float& q) const {
 18         return pow( p - q, 2);
 19     }
 20 };
 21 
 22 template <typename inputiterator1,typename inputiterator2>
 23 double computeGold(InputIterator1 first1, InputIterator1 last1, InputIterator2 first2, InputIterator2 last2)
 24 {
 25     float sum = 0.0;
 26     for(; (first1 != last1) && (first2 != last2); ++first1, ++first2)
 27         sum += pow(*first1 - *first2, 2);
 28 
 29     float s1 = sqrt(sum);
 30     std::cout << "Gold=" << s1 << std::endl;
 31 
 32     return s1;
 33 }
 34 
 35 int main(void)
 36 {
 37 #if __CUDA_ARCH__ == 200
 38     thrust::device_vector<float> p_vec(1 << 20);
 39     thrust::device_vector<float> q_vec(1 << 20);
 40     thrust::device_vector<float> r_vec(1 << 20);
 41     thrust::generate(p_vec.begin(), p_vec.end(), rand);
 42     thrust::generate(q_vec.begin(), q_vec.end(), rand);
 43     // Current Thrust's transformations supports 2 input vectors, so we use it
 44     thrust::transform(p_vec.begin(), p_vec.end(), q_vec.begin(), r_vec.begin(), powfunctor());
 45 
 46     int sum = thrust::reduce(r_vec.begin(), r_vec.end(), (int)0, thrust::plus<float>());
 47     std::cout << "sqrt(" << sum  << ")=" << sqrt(sum) << std::endl;
 48 #else
 49
50 
 51     unsigned int  N = 1 << 20;
 52     thrust::host_vector<float> p_vec(N);
 53     thrust::host_vector<float> q_vec(N);
 54     thrust::host_vector<float> r_vec(N);
 55     srand(0);

56     thrust::generate(p_vec.begin(), p_vec.end(), rand);
 57     thrust::generate(q_vec.begin(), q_vec.end(), rand);
 58 
 59     double referenceSoln = computeGold(p_vec.begin(), p_vec.end(), q_vec.begin(), q_vec.end());
 60 
 61     // device memory 'raw' pointers
 62     float* raw_ptr_P;
 63     float* raw_ptr_Q;
 64     float* raw_ptr_R;
 65 
 66     cudaMalloc( (void**)&raw_ptr_P, (N)*sizeof(float));
 67     cudaMalloc( (void**)&raw_ptr_Q, (N)*sizeof(float));
 68     cudaMalloc( (void**)&raw_ptr_R, (N)*sizeof(float));
 69 
 70     thrust::device_ptr<float> dev_ptr_P(raw_ptr_P);
 71     thrust::device_ptr<float> dev_ptr_Q(raw_ptr_Q);
 72     thrust::device_ptr<float> dev_ptr_R(raw_ptr_R);
 73 
 74     thrust::copy(p_vec.begin(), p_vec.end(), dev_ptr_P);
 75     thrust::copy(q_vec.begin(), q_vec.end(), dev_ptr_Q);
 76 
 77     // uncommenting the following will produce errors for 1.x devices 
 78     // complaining that CUDA doesn't support function pointers and function 
 79     // templates. reason is because a host function like 'rand' cannot be 
 80     // executed in the device i.e. GPU
 81     //thrust::generate(dev_ptr_P, dev_ptr_Q + N, rand);
 82     //thrust::generate(dev_ptr_Q, dev_ptr_Q + N, rand);
 83 
 84     thrust::transform(dev_ptr_P, dev_ptr_P + N, dev_ptr_Q, dev_ptr_R, powfunctor());
 85 
 86     float sum = thrust::reduce(dev_ptr_R, dev_ptr_R + N, (float)0, thrust::plus<float>());
 87     std::cout << "1. GPU " << sqrt(sum) << std::endl;
 88     std::cout << "2. CPU " << referenceSoln << std::endl;
 89 #endif
 90 
 91     std::cout << "END" << std::endl;
 92     return 0;
 93 }
 94

Before i go off, i leave with you Rich Hickey's latest talk on Hammock-drive Development. A really cool video and i strongly recommend that you catch all 40 mins of it. This is a great video to start thinking about improving self-thinking in 2011
Would this make computing say recommendation algorithms faster? By the pure fact that you're leveraging the GPU the answer is a quick "Yes" but the actual thing i wanted to show in this post is that there are always at least more than 1 way to get something done which is simply "freedom".

Have fun!

My humble google code project

2010-12-29T09:52:00.001+08:00

Just wanted to inform everyone that i've released some sample implementations of exact string matching algorithms in CUDA into my first google code project exactstrmatchgpu. As a first release ever, i've included the Bruteforce, Horspool and QuickSearch equivalents in CUDA - i surely like your comments and feedback :).

I like to leave you with this screen shot from Rich Hickey's recent talk in October 2010 which sums up pretty much the motivation behind this stuff i've just put out.

The code is messy but i'm working on cleaning it up. Thank you for your understanding and patience - meanwhile send me feedback!!!

I've planned to make more implementations available in the coming weeks before my classes start.

CUDA - An demonstration w.r.t Exact String Matching

2010-11-30T09:27:00.001+08:00

Haven't been posting about CUDA stuff for a while and i thought why not do something?

The stuff i'm going to show you is what i submitted for the GPU Computing Gems by Elsevier - though my work was eventually turned down by the editorial committee recently at the final presentation. Did a post-mortem and realized that i put in too little hours (typically you should spend 3 months per kernel for implementation, testing & optimization and i had worked out 3 CUDA algorithms). Personally disappointed but for some one whose a virgin writer in submitting papers, i'm thankful for being given the opportunity - i like to thank Prof Hwu for his generosity in his comments and review.

So...

Motivation
The idea is really simple:

How can i turn algorithms that were designed to run on CPUs to run on the graphics card(s) with possibly running hundreds of cores e.g. GTX2XX, GTX480, GTX580?

CUDA is possibly the answer to this (possibly in the next 10 years or so).

CUDA is a extension of the C language (makes sense doesn't it when it comes to controlling the hardware?) and its not that hard to learn it.

Getting to it

After reading David Kirk & Prof Hwu's book, i felt that i just gotta try it out! Next, i spent some time thinking about the sort of problems i can apply the techniques i've learnt and one subject stood out: Exact String Matching. They seemed perfect since the problem was well understood for decades and serial implementations made available via online texts or classic textbooks.

How do i parallelize a serial algorithm?

This is the first hurdle you've got to deal with ...

IMO, it helps to start with something simple (Think about the technology for implementation at a later time).

Take a pen/pencil and paper and start drawing what you think the algorithm should look like. I use squigglies to represent threads, drew blocks to house the swigglies etc. Plot a data-flow graph on how your implementation would look like. Check your algorithm for potential embarrassingly data parallel code (for-loops are good candidates), identify them and exploit their existence. Depending on your experiences and exposure to this stuff, it could be easy or extremely excruciating.

Which algorithms did i apply parallelism to?

Embarrassingly one of them is the brute-force, the other two are Horspool and QuickSearch. I've picked the latter two because by solving one of them, i get the solution to the other for free :)

The implementation i have is not optimized, that i know unfortunately since i was rushing to use it to submit an initial draft and thought i should be able to fix it in the inbetween hours of off-work and bedtime (i'm SOooooo wrong). Anyway that's history.

I don't have a multi-GPU solution at the moment since 1 GTX 480 set me back by 1000 singapore dollars and i wasn't ready to replace my rig which even though had 2 PCIe slots but were physically TOO DAMN CLOSE to house 2 GTX 480s.

I won't present the code here since i'm also in the preparing to host those codes into a repository - i'll let you know where it is if you are interested. Drop me a mail.

Initial results

Benchmarking is very important, so starting with a relatively small data set and string i went about testing the algorithm with and without shared memory. Through Nvidia's documentation, i realized that shared memory offers a aggregated bandwidth of 1 TB/s. Woot! To get an idea of how the algorihtm performed, Nvidia provided a Visual Profiler that's very useful in helping me to understand where the latencies are.

Initially, you will notice 2 long bars which respectively represent the actual kernel function strstr2 and the memory alignment function, memset32_aligned1D, used internally by CUDA.

That's not good at all, how to improve it?
Couple of things came to mind, one is removing dependencies on the global memory and improving memory coalescing. Register spilling doesn't appear to be occurring, whew.

Use Constant Memory
This sort of memory is very useful in other applications (10% reduction in run times) but not in my implementation. Reason why its effective is because you need all threads (e.g. in a block) to read the variable (allocated via the __const__) conversely why its not effective in my implementation is because the search against the memory access pattern in block-1 differs from say block-99 but if you have a situation where all blocks are executing an instruction which accesses the same memory location, then its good.

Respect data structure alignment
This is important because in the current architecture, global memory supports reads/writes of 1,2,4,8 & 16 bytes. If they are not, then it means the compiler needs to generate multiple instructions to support the non-aligned access.

Use Shared Memory
The kernel function, strstr2, is the actual workhorse conducting the search for the string and you can see that by just using shared memory on the device the run times have been reduced by 997%! woot! almost a thousand times faster!

Improvement Memory Coalescing
Coalescing is a concept that's applicable to a warp of CUDA threads but originated from compute science literature. It all comes down to optimal memory access patterns. A poor or random memory access pattern is sure to hit your application's effective bandwidth while a good memory access pattern (that is, one that respects the hardware) will surely send your application's effective bandwidth through the roof. This concept is applicable to threads in a warp.

Warp divergence
To understand this right away, immediately examine your kernel code and if there are conditional statements like "if"; if so, then there's a good chance you'll encounter warp divergence. The net effect is that throughput will be lowered since every time a warp divergence occurs, another warp cannot be executed since time's spent in deciphering and executing the threads in the if portion and the others spend time waiting on the else portion.

In my limited experiences, i've found it tough to eliminate warp divergence and appears to me that re-designing the kernels is necessary... *puts on his thinking hat*

Gotchas

If you are developing on a mac book pro, especially the newer models that i got where the laptop will drop the performance of your integrated graphics chip when you're unplugged from AC; just remember to plug in when running CUDA else you'll suffer a performance hit. Not entirely sure what's the cause: my suspects are CPU down-stepping, GT330 graphics card switching on my mac book pro etc.

I've observed that when developing CUDA apps on Mac OS X its always a good practice to use cudaMemset(...) to initialize your data structures to a default state.

Double-precision operations is supported for devices that's 1.3 and above so that means my GT 330M cannot perform such operations without suffering from lack of accuracies.

So....

That's the end of my sharing session and i hope i've been of some help to you, at least got you excited about CUDA. I know i still am :)

Clojure - I'm excited about functional programming !

2010-11-04T09:34:00.000+08:00

Oh man, these past few days have been quite a experience. I've just picked up a book Programming Clojure by Stuart Halloway (He wrote Practical Lisp too) and started to inject the functional blood into my system again.

Its very exciting i must say! There are new concepts to be learn, old-concepts to be re-mapped again. I'm planning to write a couple of stuff on Clojure when i finish the book. There are so many ideas inside Clojure and one of those things i'm interested in are Memoization & Software Transaction Memory.

Wish i could camp out at a corner of the coffeeshop for 8 hrs a day doing nothing but Clojure ; that reminds me of the time when i took time off just to learn Erlang ... these days i can't do that though i strongly believe that it helps anyone whose in the business of software development.

Stay tuned!

Ruby conference in Singapore

2010-10-29T09:28:00.000+08:00

Hi again everyone,

I like to take the chance to conduct a broadcast. Back a couple of weeks before today, i raised a question in the Singapore Ruby Brigade about the possibility of bringing a Ruby Conference to Singapore (This follows the synergy in the recent PyCON Asia Conf 2010) ; i have to admit that i have no prior experience in this and i'm not the organizer of this event but just trying to lend a helping hand.

So having said that,

I'm calling out to ALL Ruby enthusiasts out there to pitch in ideas at the Singapore Ruby Brigade group w.r.t the following:

1) What tracks would like the conference to have? e.g. education, research, application of ruby to real-world applications etc.

2) Speakers list. Currently, a fellow member of the SRB has contact several Ruby heavy weights and thus far a lot of many are interested. But if you know of someone that we don't and is doing GREAT Work, do not hesitate to submit their names

3) We are also trying to gather information about some practical issues like costs of running this conference so if anyone has expertise in this area, volunteering your expertise would be greatly appreciated.

4) We are also trying to get an idea of how much people like yourselves are willing to pay to attend the conference? Please send a note to @sausheong on twitter.

Alright, that's all i have for now. Would be great if we can come together to promote Ruby, share best practices and BE INSPIRED!

Thank you for your attention!

Shocked, Sad about DLR's IronPython IronRuby

2010-10-22T09:28:00.001+08:00

Just a day after my previous post on IronPython and IronRuby and i've read the news that Jim Hugunin of Jython, IronPython and DLR fame has left Microsoft for Google. Read his blog here.

This news was right after when Microsoft announced that it'll stop investing in IronPython. Read it here.

I'm shocked to say the least ...

I've got questions ranging from the whys to what's going to happen to IronPython; yesterday i've talked about the countless possibilities that DLR allowed developers like myself but now ...

Life goes on i suppose ... back to the drawing board ...

Automation

2010-10-21T21:11:00.001+08:00

What's been happening
Has been quite a while since i last posted anything as i've been busy learning new stuff, which includes new programming languages like C#, VB, Ruby and also building automation into a product that specializes in 3D modeling and visualization.

Laying down my perceived flaws
Before we come to exploring Automation per say, one of the challenges i had to face was to actually learn the tools i needed to hit the ground running. Wasn't exactly tough but it required me to put down my previous reservations about learning and using the Microsoft Technology suite "stack" if i can call it that. Once that was done, it was kind of cool and i believe people whom have extensive coding background in open source technologies would make the leap of faith to Microsoft easier than the other way round. That's what i saw at least.

The key thing that helped me cross that barrier was recognizing that i should looking at other technologies with an objective mind; focus on the problem and not the technology.

Automation
Having said that, next thing i want to talk about is Automation. After spending time working on it, i'm happy to say that we've shipped version 1.5 to our customers with the company. Now we're waiting for feedback. Yep, feedback's important to the whole process - no point building something that people don't like to use. Building automation is kind of a tough subject to begin with since it doesn't really fall into the standard SDLC; but that's what they said about Test Driven Development. TDD should be part of the SDLC period

The usual questions are "How does automation help me?", "How much would it cost to build this automation?", "What should we automate?", "Can this be automated?". There are many more. The typical use cases of automation seem to lie more in complementing QA i.e Automated Testing.

We've built the usual stuff that comes along with Automated Testing Frameworks using C#, VB - Test Management Tools, Test Execution Tools (local, distributed), Test Reporting/Archival/Trending Tools...what did i miss? Hmm...whatever

But i see a problem. We have to get someone whose skilled in the Microsoft Technology stack to be able to write tests (Well, there are a lot of Microsoft professionals aren't there?...yeah sure...but what if you can garner the people whom practice Python, Ruby, Scala and wouldn't that open you to a greater pool of talents?). Those professionals may/may not need to re-write libraries not found in the C# arsenal etc

Where is this leading to?
Think IronPython, IronRuby. These two ports originated from Python and Ruby respectively and they both have a great number of followers but most important of all, the integration of Microsoft suite with open source is a very powerful notion. I've been working on a DLR integration with the Automation framework lately and progress has been promising. This is exciting because of a number of reasons: I can still program in Python/Ruby, i can still re-use libraries that was previously developed (re-invention of the wheel aint gonna happen), integration with Rails/Django to pull/publish data, you get the idea? The possibilities are endless.

Book cover for GPU GEMs Volume 2

2010-09-18T05:23:00.000+08:00

Here's a quick peek at one possible book cover for GPU GEMs Volume 2 circa 2011 :)

How to catch unhandled exceptions and generate diagnostic data

2010-08-31T22:08:00.000+08:00

Recently, i got involved in the C# world and came across an interesting problem where one of our applications would crash at times and nobody would know why and how. The idea was to find a way to detect the crash, break into and obtain its dump file & a nice stack trace. That always help developers.

There are a number of solutions but i needed something that was relatively easy for a C# newbie (like yours truly), it had to be a runtime artifact where the code would reference (think dynamic loading) among other requirements. Voila, found a nice little library at Microsoft's Open Source website - yeah i know, for microsoft newbies like myself i had NO IDEA microsoft was INTO open source.

And so the story goes, there is this neat library called Application Management. Download it and check it out yourself. Might save you the time you needed. I know ;)

Let me give you a quick rundown of what its about. Basically, its a DLL and supports the major microsoft programming languages like C#, VB etc and so that's what made it nice because i created a sample application in C# and since it works, i know the VB.NET code i'm working on is likely to work. Cool.

So what you do?

Create a C# project in Visual Studio 2010 (You can try VS2008, i didn't try it)
Reference the DLL by adding it to your project (its ApplicationManagement.dll)
You need to initialize it in your application. In this case, its the Program.cs

using ApplicationManagement;
namespace MyTestApp
{
    static class Program
    {
     ///
     /// some other stuff
     ///

     public static ExceptionHandler exceptionHandler; 
        static void Main()
        {
            exceptionHandler = new ExceptionHandler(); 
            Application.EnableVisualStyles();
            Application.SetCompatibleTextRenderingDefault(false);
            AppDomain.CurrentDomain.UnhandledException += new UnhandledExceptionEventHandler(CurrentDomain_UnhandledException);
            Application.Run(new Form1());
        }

        static void CurrentDomain_UnhandledException(object sender, UnhandledExceptionEventArgs e)
        {
                exceptionHandler.UnhandledException(sender, e);
        }

So in your application code, it were to throw some form of exception then a user dump + stack trace would be created in your project's working directory. On my setup, its found at <project>\bin\Debug\Logs\AppErrors\

Well that's it. Have fun!

The World of Windows (C#, F#, VB, .NET)

2010-08-24T22:35:00.000+08:00

I haven't been blogging much lately because i've found a new job that requires me to learn new stuff - and this time its not exactly the LAMP stack ;)

So what is it exactly? well its the mysterious world of windows. Like all things, its mysterious because quite simply i have not worked with it previously. But these past few weeks opened my eyes to another platform that i have never considered before. My new employer is receptive to the idea that i lack experience in Windows development but that didn't deter me so i guess its probably why it doesn't deter them :) yay

One of the questions i asked myself is how do i get to speed with C# and the likes? Kachink! I read the ebook Head First C# (i was used to their layout and presentation format, so it works much better for me) and that took me where i wanted to pretty quickly enough. And i thought a very well written book on that subject was Essential C# 4.0 That was done. How about Visual Basic? I looked to Microsoft's MSDN (which is a great resource btw)

For a guy whose done Java,J2EE - you'll probably find C# pretty easy to follow since it was created *supposedly* to be Java's rival :) though the Visual Studio 2010 was packed with lots of goodies that the developer can take advantage of. One cool thing from VS2010 was the fact that it reduce the amount of boilerplate code for accessor/mutator methods so much easier.

Here's an example in VB
Public Class SimpleTest
    ' --- replace the string "newPropertyValue" consistently
    Private newPropertyValue As String
    Public Property NewProperty() As String
        Get
            Return newPropertyValue
        End Get
        Set(ByVal value As String)
            newPropertyValue = value
        End Set
    End Property

End Class

Another example in C#
namespace SimpleHelloWorld
{
    public partial class Form1 : Form
    {
        public Form1()
        {
            InitializeComponent();
        }

        private void button1_Click(object sender, EventArgs e)
        {

        }
    }
    //
    // Accessor/Mutator methods to hide 'backing' value
    class SimpleTest
    {
        private int state;
        public int GetState
        {
            get { return this.state; }
            set { this.state = value; }
        }
    }
}

From a learner's point of view, the language's syntax does help make a lot of concepts stick much better and faster since i can quickly identify the commonalities these 2 languages share. Btw, i thought the idea of a 'partial' class was quite nice :) and the idea of namespaces pretty collide with what i understood from C++'s concept of 'namespace'. Of course this is not exhaustive.

From a guy whose done Erlang,Python,Ruby i wonder what F# holds for me? I guess i'll find out soon...

I cannot possibly end this post without mention two Scotts: Scott Gu's Blog and Scott Cate's blog.

My bias is possibly gone as i embrace the new technologies (at least to me)

Author for 4th Edition of GPU GEMs

2010-07-16T16:26:00.000+08:00

Let me start by saying: Yayyy!!! It was a whole lot of fun! I'll do it again!

Reason for my joy is that my article has been accepted for GPU GEMs 4th edition and now finalizing my article before its released and all that stuff. That happened rather recently when i started working on the article around April 2010 on a feverish pace and made it before the deadline; now i'm cleaning up the article and i've agreed to release the source code (hope it won't be too contrived for you experts out there) - details will be finalized over the course of time.

I'm not trying to be secretive on the subject of my work but the topic relates to the application of CUDA (Nvidia's Compute Unified Device Architecture) and so far, i've had tons of fun with it. As part of the finalization process, i've got new hardware namely GTX480 (this hardware is OSSM!) running benchmarks, looking at possibility of embedding PTX (its assembly language for CUDA), learning illustrations etc

Btw, if you are interested in CUDA - you should check NVIDIA's website on CUDA and go download it from here and get yourself started. What's interesting with the latest CUDA is the support for C/C++ printf() and function pointers in the CUDA kernel which opens up a HUGE treasure of possibilities for application.

Debunking the 100X GPU vs. CPU myth: An evaluation of throughput computing on CPU and GPU

2010-07-08T09:58:00.000+08:00

Just read a very interesting article on the continuing debate between GPU and CPU and their relative performance. Its between Intel and Nvidia; not surprising its not AMD vs. ATI or Intel vs. ATI- that would REALLY be interesting. Here's the abstract

Abstract:

Recent advances in computing have led to an explosion in the amount of data being generated. Processing the ever-growing data in a timely manner has made throughput computing an important aspect for emerging applications. Our analysis of a set of important throughput computing kernels shows that there is an ample amount of parallelism in these kernels which makes them suitable for today’s multi-core CPUs and GPUs. In the past few years there have been many studies claiming GPUs deliver substantial speedups (between 10X and 1000X) over multi-core CPUs on these kernels. To understand where such large performance difference comes from, we perform a rigorous performance analysis and find that after applying optimizations appropriate for both CPUs and GPUs the performance gap between an NVIDIA GTX280 processor and the Intel Core i7-960 processor narrows to only 2.5x on average. In this paper, we discuss optimization techniques for both CPU and GPU, analyze what architecture features contributed to performance differences between the two architectures, and recommend a set of architectural features which provide significant improvement in architectural efficiency for throughput kernels.

I'm a big fan of GPUs and recognize the value of CPUs but one thing i'll like to keep in mind is that GPUs though not a young technology (in 2010) but they are young as compared to the CPU as we know it today. Go read it and form your own ideas about it.

My thoughts...

The design and purpose of the CPU and GPUs were entirely different and for decades the GPU have always played a complementary role to the CPU, but things are shifting as we speak from massively-parallel antivirus engines etc. Like any IT product in the market today, it will go through the usual adoption and evolution phases before i can even see it attempt to supplant the CPU. How far away are we? I like to know too.

So...

If i were a bystander reading the article, first reaction would be - Has the GPU finally caught up with the CPU to prompt an intensive study of this nature by the largest CPU producer in the world ? Perhaps i should REALLY pay attention to it.

Another adventure begins!

2010-06-29T22:58:00.000+08:00

Hello everyone,

To start, i'd like to say that i'm leaving Linden Lab. If you caught the news from the internet, then you'll understand why but if you didn't here are two i found recently by googling about it. http://www.vizworld.com/2010/06/linden-lab-laying-staff-closing-singapore-office/ and http://digitalmedia.strategyeye.com/article/70432b60e0/2010/06/10/Linden_Lab_cuts_30_of_Second_Life_staff/

To be clear about the motivation about this post, i'm only sharing my thoughts on what i think about the entire thing. If there is one takeaway from this article i like you to have, that would be " when one door closes, another door opens "

Obviously, being laid off is not a pleasant experience but its something i have to deal with. Having support from family and co-workers during this period in time really made it easier to cope. Support from co-workers in the U.S. also helped made things easier, collectively speaking it wasn't too bad an experience. Lindens are just a warm and fuzzy bunch ;)

Our lab in Singapore is "tight" - yeah i mean it. Having smart and enthusiastic people come to office every day and hearing their ideas and how they work feverishly to accomplish is what its all about. period. work hard, play hard. You should see some of the stuff we did every quarter - we did cycling at pulau ubin (offshore island at singapore), had great lunches, bowling etc. Our farewell lunch was quite a different experience, at least for me, as i quickly realized that i would no longer be working, having lunches, discussions/debates with these people. The bantering, ohhhh how i would miss that.

Someone once said " what makes a person want to stay with a company is always the people ". Was it Mitch Kapor? Well i don't know

Moving on, i realized that i'll soon be out of a job. Jobless in Singapore! Hehe, i'll admit these are uncertain times and was wondering what i should do next and thanks to the great support rendered by co-workers and family, i quickly found myself doing interviews and i'm really surprised how many people are doing amazing work in startups and established companies. I had a blast doing these interviews! These companies are recognizing the power of open source and field those leadership roles with people with the right mindset :D

I'm still looking for my next adventure :D

Dennis Hopper in "Crash" said " it's a beautiful thing man ... "

Touring C++ Coroutines

2010-06-05T00:10:00.000+08:00

In the last couple of months, i have been exposed to C++ Co-routines and i wanted to write about my learning. The library i'm referring to is the Boost.coroutine where it has a family of class templates that wrap function objects in coroutines and these coroutines can be reentered and returned more than once without causing the destruction of automatic objects. Basically, it allow states to be store/retrieve across a function call. I've had the immense pleasure of figuring this stuff out because another favourite scripting language of mine, Python, uses this concept extensively in its 'yield' statement.

There are two important concepts you'll need to comprehend before you can use coroutines effectively in C++. Number 1, the subroutines we're so used to programming (e.g. functions, procedures) lose their state when they return (You'll probably wonder that function objects also save state but there's a difference and that is: function objects stored their state in class member variables while coroutines store the state in the stack as automatic objects - these automatic objects are not destroyed btw which is COOL!). Number 2, coroutines preserved the point of execution and is able to resume execution from where it left off which is unlike function objects

Simple use of Coroutines
One example is the generator function (yep, python has it too) and what it is is really a function that returns a sequence of values instead of returning all of them; this programming concept is very useful in practice in the form of a dynamic prime number generator.

I love the function objects because it allows me to indulge in the world of functional programming to an extent though i know i'm still working in C++. You can pretty much recognize a function object in C++ by looking for the operator()() in the source code and possibly a couple of operators that are overloaded that suits the logic.

Simple example is the generator that generates all numbers (integers) between lower and upper inclusive. e.g. the lower bound is 1 and 4.

class generator {
public:
  generator(int lower, int upper): lb(lower), ub(upper) {}
  int operator()() { return lb++; }
  operator bool() const { return lb < ub; }
private:
  int lb, ub;
};
int main(int argc, char** argv) {
  generator gen(1,4);
 // invokes the operator bool() 
  while(gen)
    // invokes the operator()() - function object! 
    std::cout << gen() << "\n";
}

There's another way and that is using input iterators as generators but IMO, its not that straightforward and simple as opposed to using function objects which i do consider more elegance but then again, i'll iterate a philosophy of mine "whatever that works simply" when it comes to choosing paradigms

So i know i can write generators using function objects and can i do that using coroutines? Sure i can and here's what i found out

// In a .h file, there's the contents 
#ifndef GEN_TEMPLATE_H
#define GEN_TEMPLATE_H

#include 
namespace coro = boost::coroutines;
typedef coro::generator gen_type;
/*
Notice that range_generator body is entered for the first itme when the generator is constructed
(from the main entry point) then at every iteration range_iterator is reentered from 
yield(). In this case range_iterator is reentered when generator::operator++ is invoked.
*/
int range_generator(gen_type::self& self, int min, int max) {
    while ( min < max - 1)
        self.yield(min++);
    self.exit(); // This is used instead
}
#endif
...
...
// In another .cpp file
#ifdef GEN_TEMPLATE_H
    gen_type generator(boost::bind(range_generator, _1, 100, 200) );
    while( generator != gen_type() )
        std::cout << *generator++ << "\n";
#endif

Now this example is the input iterator fashion of creating a generator function and the creation of the generator is in the statement boost::bind(range_generator, _1, 100, 200) where the boost::coroutines::generator::self is bound to range_generator; another thing is that you can see the generator entry function "yield" and exit function "exit" invoked through the "self" object. The "self" object is used to identify the different instances of a generator.

So i asked myself: "Great, it works on built-ins, how about user defined types?" Yep can be done too and here's what i did, its not great but it'll do for this example

// in some .h file

#ifndef GEN_TEMPLATE_H
#define GEN_TEMPLATE_H
class test {
public:
    typedef int value_type;
    test(int min) : m_current(min) {}
    test() {}
    test& operator++() {
        m_current++;
        if ( m_current > max ) m_current = -1;
        return *this;
    }
    test& operator++(int) {
        test t(*this);
        ++*this;
        return t;
    }
    friend bool operator<(test& rhs, const test& lhs) {
        if ( rhs.max != lhs.max ) rhs.max = lhs.max; // only invoked the first time
        return rhs.m_current < lhs.m_current;
    }
    friend bool operator==(const test& rhs, const test& lhs) {
        return rhs.m_current == lhs.m_current;
    }
    friend bool operator!=(const test& rhs, const test& lhs) {
        return !(rhs == lhs);
    }
    operator int() { return m_current;}
    int operator* () { return m_current; }
private:
    int m_current;
    int max;
};
#include 
namespace coro = boost::coroutines;

typedef coro::generator gen_type;
/*
Notice that range_generator body is entered for the first itme when the generator is constructed
(from the main entry point) then at every iteration range_iterator is reentered from 
yield(). In this case range_iterator is reentered when generator::operator++ is invoked.
*/
int range_generator(gen_type::self& self, test min, test max) {
    while ( min < max)
        self.yield(min++);
    self.exit(); // This is used instead
}
#endif

...

...

// in some .c file

#ifdef GEN_TEMPLATE_H
    gen_type generator(boost::bind(range_generator, _1, test(100), test(200)) );
    while( generator != gen_type() )
        std::cout << *generator++ << "\n";
#endif

Pretty basic stuff there, so it can be done.

Another concept: the producer/consumer pattern
I love this concept because its basically a messaging system and i have Nat Goodspeed & Brad Kittenbrink to thank ;) thanks guys!

Here's how it works (its not perfect code but it works to illustrate my example)

// In some .h file

#ifndef MESSAGE_H
#define MESSAGE_H
#include 
#include 


class Message {
public:
    Message(int msg) : m(msg) { std::cout << "Created:" << m << std::endl; }
    int operator *() { return m; }
    friend std::ostream& operator<<(std::ostream& os, const Message& m);
private:
    int m;
};

std::ostream& operator<<(std::ostream& os, const Message& m) {
    os << m.m << std::endl;
    return os;
}

template 
class MessageQ {
public:
    MessageQ(int capacity = 10) : capacity(capacity) {
        for( int i = 0; i < capacity; ++i) q.push_back(Message(i));
    }
    operator bool() { return q.size() > 0; }
    Message operator *() { Message lastEle = q.back(); q.pop_back(); return lastEle; }
    MessageQ& operator++() {
        if ( q.size() <= 0 ) return false;
        q.pop_back();
        return *this;
    }
private:
    std::vector q;
    int capacity;
};

#endif

#ifdef MESSAGE_H
#include 
namespace coro = boost::coroutines;

typedef coro::generator gen;

const Message& producer(gen::self& self, MessageQ msgq) {
    while (msgq) {
        self.yield(*msgq);
    }
    self.exit();
}

template
void consumer(Producer prod) {
    do {
        std::cout << *prod << std::endl;
    } while( ++prod );
}

#endif
...
...
// and in some other .c file
#ifdef MESSAGE_H
    consumer(gen(boost::bind(producer, _1, MessageQ())));
#endif

So what i've basically done is to create a message queue of sorts and populate it with 10 simple messages and use the generator, in this case the consumer, to iterate through the messages.

JUnit Cookbook

2010-05-15T10:41:00.002+08:00

This is simply excellent writing :)

http://junit.sourceforge.net/doc/cookbook/cookbook.htm

Mudflap - Instrumentation of pointers

2010-05-10T23:23:00.000+08:00

Heard of mudflap? This technology is exciting in the sense that it allows one to instrument C/C++ pointers that it finds to be error prone to cause an application to crash and its got potential. Read about the wiki and download the original paper.

It is enabled by passing -fmudflap to the compiler. For front-ends that support it (C and very simple C++ programs), it instruments all risky pointer/array dereferencing operations, some standard library string/heap functions, and some other associated constructs with range/validity tests. Modules so instrumented should be immune to buffer overflows, invalid heap use, and some other classes of C/C++ programming errors. The instrumentation relies on a separate runtime library (libmudflap), which will be linked into a program if -fmudflap -lmudflap is given at link time. Run-time behavior of the instrumented program is controlled by the MUDFLAP_OPTIONS environment variable.

Customizing TUT framework for TeamCity 5.x Reporting Tests

2010-02-17T16:48:00.003+08:00

Recently, i've had the pleasure of figuring out how to get CPP unit tests "report" results to a Distributed Build System like TeamCity 5.x Btw, i'm not accepting any sort of monies from Jetbrains nor was i invited to do this post; its purely to share information.

The first thing to realize about the TUT framework (lightweight C++ unit testing framework - i like it :) ) is basically how it works and how TeamCity 5.x service messages work. There isn't a shortcut to understanding TUT's mechanism except to read the code so download it and understand it. On TeamCity 5.x Service Messages, refer to this article here.

Once you've done that, you'll realize that TeamCity 5.x needs you to output messages of the form

##[teamcity testSuiteStarted name='suite.name']
...
##[teamcity testStarted name='XXX']
...
##[teamcity testIgnored name='ccc']
...
##[teamcity testFinished name='XXX']
...
##[teamcity testIgnored name='ddd']
...
##[teamcity testSuiteFinished name='suite.name']

during a build process and you'll realize that testStarted and testFinished must begin and finished in pairs, so are the testSuiteStarted and testSuiteFinished messages.

Next thing you should figure out is how and where to insert print statements in your test code. Well the short answer is to study the C++ class callback and all classes that inherit from it. Currently, tut_console_reporter.hpp and tut_cppunit_reporter.hpp are the ones i know about.

What you need to do is to realize that the TUT runner as the framework calls it, is actually invoking whatever class that inherits tut::callback (i.e. class console_reporter) and you need to implement/override the virtual C++ member functions in tut::callback so that you can place the appropriate TeamCity 5.x service messages. I won't place the code here cos its pretty easy :) and there are more than one way to do it - depends on how your unit tests are crafted.

Once you get it to work, you should see statistics on your TeamCity dashboard via the Projects dashboard and you can drill down into how long each test took, details etc.

Quick update on 8 May 2010: I've since then released this modifications to sourceforge and here's a URL to JetBrains list of available plugins.

PyCUDA on Mac OS X

2009-12-23T21:47:00.003+08:00

PyCUDA came to my attention sometime back and i've got a chance to install it and starting to play with it - this package is essentially a python binding to CUDA (Compute Unified Device Architecture) by Nvidia; and this package (download it here) provides an easier way to program your CUDA graphics card via Python.

Btw, i've installed PyCuda 0.93, boost 1.41.0 and CUDA 2.3 on Mac OS X 10.5.8

So in this post of mine, i'd like to share some of the problems i've faced but before i do that its important that you read through the following wiki pages:

And try out the installation of PyCUDA on those pages but if you found a problem, i might have a solution for you. So i'll run a couple of Q&As here and hope it helps

Q: After i installed PyCUDA as instructed, i ran the test_driver.py and an error message was thrown out Fatal Python error: Interpreter not initialized (version mismatch?). Why is this happening?
A: The version of the Python used is incorrect. You need to rebuild boost and PyCUDA using a consistent version of Python and you need to check the following files

siteconf.py (; this file is produced after launching configure.py)
project-config.jam (; this file is produced after launching bootstrap.sh)

I changed the following lines in project-config.jam to look like this:

# Python configuration
# using python : 2.5 : /Library/Frameworks/Python.framework/Versions/2.5 ;
using python : 2.5 : /System/Library/Frameworks/Python.framework/Versions/2.5/ ;

Noticed that i changed the python interpreter to use another location (This same location must be used to build PyCuda subsequently (i.e. /System/Library/Frameworks/Python.framework/Versions/2.5/bin/python ./setup.py build etc)

Should be alright now...

Q: When i ran the test_driver.py, i found an error message like this

Traceback (most recent call last):
File "test_driver.py", line 15, in
gen = pycuda.autoinit.device.get_attributes()
File "/Library/Python/2.5/site-packages/pycuda-0.93-py2.5-macosx-10.5-i386.egg/pycuda/driver.py", line 50, in device_get_attributes
print getattr(device_attribute,i), dev.get_attribute(getattr(device_attribute, i))
pycuda._driver.LogicError: cuDeviceGetAttribute failed: not found

Why is this happening?
A: Frankly, i don't know at the moment but the other tests in the package i've run and tested is alright but something odd is happening (I'm sure its got something to do with my lack of understanding of how libraries are loaded/unloaded)

What i know at this point is that when i re-worked the package a little to provide me debug output, i've found an interesting scenario and you can see this below

cuDeviceGetAttribute
CAN_MAP_HOST_MEMORY 1
cuDeviceGetAttribute
CLOCK_RATE
Traceback (most recent call last):
File "test_driver.py", line 15, in
gen = pycuda.autoinit.device.get_attributes()
File "/Library/Python/2.5/site-packages/pycuda-0.93-py2.5-macosx-10.5-i386.egg/pycuda/driver.py", line 50, in device_get_attributes
print getattr(device_attribute,i), dev.get_attribute(getattr(device_attribute, i))
pycuda._driver.LogicError: cuDeviceGetAttribute failed: not found

You will see a output that resembles this:
cuDeviceGetAttribute <--- this was enabled by setting CUDA_TRACE=True (False is default)
CAN_MAP_HOST_MEMORY, 1 <--- CUDA device attribute mapped to PyCUDA and the number 1 is the end result of calling the actual API in CUDA
and from some simple deduction you can see that for some reason the program reported that it could not locate the API by the stack trace. .... Will investigate further when i have the bandwidth again but meanwhile feel free to email me if you have an idea why its failing

---- update on 25 dec 2009 -----
A small and quick update is that after removing the attribute CLOCK_RATE in driver.py, it ran fine and the verification passed.
---- end of update ------

---- update on 30 dec 2009 ----
Finally, i've got spare time to look at this and turns out to be a problem with GeForce graphics card present on my machine (there are 2 9400m and 9600m GT)
In a nutshell, i've used the low powered version i.e. 9400m and when i changed to 9600m GT (high performance chip) the problem was resolved.
This is odd because this problem was not present in CUDA 2.0 - Apple / Nvidia bug? Hrmm...
---- end of update ----

But other than that, so far so good :) Have fun !

PythonWin and Python Ctypes

2009-12-17T17:55:00.002+08:00

Ran into a situation when using PythonWin by Mark Hammond - btw its a really nice IDE with syntax highlighting and debugging capabilities :) kudos!. So what happened is this:
When i ran this simple piece of code against the PythonWin and a normal Python session from the Windows command prompt, i get two results.

>>> from ctypes import *
>>> cdll.msvcrt.printf(c_char_p('hello world'))

When i ran it against the PythonWin, i was only seeing the string '11' returned and in the other case i was seeing 'hello world11' which is what i expected. For some reason unknown to me at this point in time, PythonWin is capturing the result of the call to printf() whilst ignoring/hid the I/O conducted by printf(). I think it is a bug and i've filed it <- in case you're interested - but in a nutshell its basically the group of I/O functions sys.stdout etc being overridden

Of course i'm not going to dismiss PythonWin (i think its a great job at that, thanks Mark) but i just thought that you might like to know incase you bumped into this.

Tsung - Load Testing Tool in Erlang

2009-11-02T16:17:00.005+08:00

You've must have heard of Tsung - a load testing tool. Well i've got the opportunity to learn and use it in my current employment and its really versatile in conducting load tests for web-based applications. It has many features that i've come to appreciate after spending a couple of years with Mercury Interactive (its defunct now after being bought over by HP in 2006); admittedly Mercury's solutions were more versatile than what Tsung can offer but considering that web is the common platform where many applications are being hosted on; i think its a safe bet :)

So, IMO i think

Tsung is good for conducting load testing scenarios and executing them in a local / distributed manner
Load Testing can be relatively light weight (Erlang processes are hitting the target app) and hence the cost of using relatively heavy weight machines is likely to reduce since there is lesser need to use those machines since more concurrent users can be simulated on 1 machine in Tsung. (I'll see whether this assumption is correct)

Good starting points (URLs of interest):

A load testing tool would be useless if it didn't know how to capture server response data (e.g. checking whether an expected string is returned) and reuse it (e.g. session ids returned from servers which you can use subsequently in a HTTP POST/GET in an URL-rewrite type of string) subsequently, generate dynamic data, read data from an external file (commonly used in storing <username, password> pairs)
So my example would illustrate the logging in to a website (e.g. http://projecteuler.net) and logging out from it - simple enough to illustrate my point.
Next what i did was to record the series of events that mimick my user logging and logging out of the website and this is captured in the tsung_recorder_timestamp.xml
and i used that XML file and included some other stuff so that it looks like what i have for you below (this is a basic load test scenario)

<?xml version="1.0"?>
<!DOCTYPE tsung SYSTEM "/usr/local/share/tsung/tsung-1.0.dtd">

<tsung loglevel="info" dumptraffic="false" version="1.0">

<clients>
<client host="localhost" use_controller_vm="true"/>
</clients>
<servers>
<server host="78.110.165.8" port="80" type="tcp"></server>
</servers>

<load>
<arrivalphase phase="1" duration="2" unit="minute">
<users interarrival="1" unit="minute"></users>
</arrivalphase>
<user session="rec20091102-01:30" start_time="0" unit="second"></user>
</load>

<options>
<option name="file_server" value="/tmp/userlist.csv"></option>
</options>

<sessions>
<session name='rec20091102-01:30' probability='100'  type='ts_http'>
<request><http url='http://projecteuler.net/' version='1.1' method='GET'></http></request>
<request><http url='/style_main.css' version='1.1' if_modified_since='Sat, 29 Nov 2008 20:04:00 GMT' method='GET'></http></request>
<request><http url='/images/logo.jpg' version='1.1' if_modified_since='Thu, 28 Dec 2006 14:12:43 GMT' method='GET'></http></request>
<request><http url='/images/icon_register.png' version='1.1' if_modified_since='Fri, 31 Dec 2004 12:48:16 GMT' method='GET'></http></request>
<request><http url='/images/icon_about.png' version='1.1' if_modified_since='Fri, 31 Dec 2004 12:48:28 GMT' method='GET'></http></request>
<request><http url='/images/icon_problems.png' version='1.1' if_modified_since='Fri, 31 Dec 2004 12:48:26 GMT' method='GET'></http></request>
<request><http url='/images/icon_login.png' version='1.1' if_modified_since='Fri, 31 Dec 2004 12:48:32 GMT' method='GET'></http></request>
<request><http url='http://projecteuler.net/images/corner_tr.gif' version='1.1' if_modified_since='Thu, 10 Apr 2008 19:35:02 GMT' method='GET'></http></request>
<request><http url='/images/corner_tl.gif' version='1.1' if_modified_since='Thu, 10 Apr 2008 19:34:41 GMT' method='GET'></http></request>
<request><http url='/images/corner_br.gif' version='1.1' if_modified_since='Thu, 10 Apr 2008 19:35:10 GMT' method='GET'></http></request>
<request><http url='/images/corner_bl.gif' version='1.1' if_modified_since='Thu, 10 Apr 2008 19:34:55 GMT' method='GET'></http></request>
<request><http url='/images/euler_main.jpg' version='1.1' if_modified_since='Mon, 21 Jan 2002 19:18:20 GMT' method='GET'></http></request>

<thinktime random='true' value='2'/>

<request><http url='http://projecteuler.net/index.php?section=login' version='1.1' method='GET'></http></request>

<thinktime random='true' value='6'/>

<request subst="true">
<match do="continue" when="match">Logged in as %%readcsv:getUsername%%</match>
<http url='/index.php' version='1.1'  contents='%%readcsv:getUserString%%' content_type='application/x-www-form-urlencoded' method='POST'></http>
</request>

<thinktime random='true' value='4'/>

<request><http url='/images/icon_tick.png' version='1.1' method='GET'></http></request>

<thinktime random='true' value='4'/>

<thinktime random='true' value='3'/>

<request><http url='http://projecteuler.net/index.php?section=logout' version='1.1' method='GET'></http></request>
</session>
</sessions>

</tsung>

Hence, the main thing you should note is the use of dynamic substitution (e.g. %%readcsv:getUsername%%) where i wrote a simple erlang program to read my username and password from a file (see the XML tag option above) and replacing each simulated user with a valid user id and password.
Next, i checked that the server response contains a string Logged in as XXX where XXX would be dynamically generated by the function (Check out the erlang code for the function, simple stuff).
The erlang program is shown below.

1 -module(readcsv).
2 -export([getUserString/1, getUsername/1]).
3
4 getUserString({Pid, DynVar}) ->
5  {ok, Line} = ts_file_server:get_next_line(),
6  [Uid,Pwd] = string:tokens(Line, ","),
7  "username=" ++ Uid ++ "&password=" ++ Pwd ++ "&login=Login".
8
9 getUsername({Pid, DynVar}) ->
10  {ok, Line} = ts_file_server:get_next_line(),
11  [Uid,_] = string:tokens(Line, ","),
12  Uid.

The above erlang code must be compiled via erlc and you place the .beam file into the directory via a command like this

sudo mv readcsv.beam /usr/local/lib/erlang/lib/tsung-1.3.1/ebin/

Now, running the load test should be alright.
Note: In this load test, i defined a duration of 2 minutes with 2 users since load testing using 800 gazillion users is considered a chargeable offense so DON'T DO IT.

Have fun!

Simple problem solving using Scala

2009-10-24T10:21:00.008+08:00

If you like solving mathematical problems using Scala, i would suggest that you sign up at Project Euler (Thanks to my good friend Chi Hung for getting me hooked!) and use your favourite programming language or languages to solve it (There is more than 1 way to accomplish a goal)

So here's my take on it using Scala and just to provide an example, i'm using problem 4

scala> var largest = 1
largest: Int = 1

scala> for(i <- 1 until 1000) for( j <- i until 1000) {
|  val v = i*j                                            
|  if ((v.toString.reverse.mkString) == (v.toString)) {     
|   if (v > largest) largest = v               
|  }
| }

scala> largest
res5: Int = 906609

This is not efficient since i'm looking for the largest palindrome made from 2 3-digit numbers so i should reverse this and possibly limit the data ranges to perhaps half. Anyway, i'd like to show you how this simple function can be factored into a style more Scala-like.

scala> var largest = 1
largest:Int = 1

scala> def isPalindrome(s:Any):(Boolean,Any) = {
|  if (s.toString.reverse.mkString == s.toString) (true,s)
|  else (false,s)
| }
isPalindrome: (Any)(Boolean, Any)
scala> (1 to 100).map(i => (1 to 100).map(j => if ( isPalindrome(i*j)._1 && (i*j) > largest) largest = (i*j)))
scala> largest
res12: Int = 9009
scala> (1 to 1000).map(i => (1 to 1000).map(j => if ( isPalindrome(i*j)._1 && (i
*j) > largest) largest = (i*j)))
res13: RandomAccessSeq.Projection[RandomAccessSeq.Projection[Unit]] = RangeM(Ran
geM((), (), (), (), (), (), (), (), (), (), (), (), (), (), (), (), (), (), (),
(), (), (), (), (), (), (), (), (), (), (), (), (), (), (), (), (), (), (), (),
(), (), (), (), (), (), (), (), (), (), (), (), (), (), (), (), (), (), (), (),
(), (), (), (), (), (), (), (), (), (), (), (), (), (), (), (), (),...
scala> largest
res14: Int = 906609

Ignore the part on the variable res13 and this version ran pretty fast for me too and i like this version much better than the former. It can be faster but i'd like to show you how to solve problems for now.

Feedback is appreciated!

Scala - I like it!

2009-09-17T12:54:00.013+08:00

Have you heard of Scala? Its a new programming language designed by Martin Odersky and i'm exploring it as i write this.

The motivation is this:

I'm working in an environment whereby Java is not the mainstream language being used and i don't want to lose the accumulated programming experiences using Java and i want to find a new way (think functional + imperative) to use it and create apps that scale on multi-core efficiently.

For those whom have followed my other blog on Erlang (You can pretty much tell i like functional programming) but here's a new language that combines both OOP + FP. Its a relatively young language when compared to Ruby, Python, Erlang ... but i think its got great potential. Here are two organizations that are using it *drum rolls*

Here's a class hierarchy diagram in Scala:

Job openings at Linden Lab

2009-08-03T18:11:00.000+08:00

The company i'm working for (Linden Lab) has a great list of available positions and i would like to invite interested and passionate individuals about Virtual Worlds and Second Life to go check it out :D ---> http://lindenlab.com/employment

If you are good with software development (that means programming) and/or QA-ing experience, feel free to drop me a mailto:tay_boon_leong@yahoo.com.sg

Should free software depend on Mono or C#?

2009-07-15T11:09:00.003+08:00

Here's a fantastic read on why free software should never depend on Mono or C# authored by one of the favourite authors of OSes Mr. Richard Stallman. Click here for the full story.

Debian's decision to include Mono in its principal way of installing GNOME, for the sake of Tomboy which is an application written in C#, leads the community in a risky direction. It is dangerous to depend on C#, so we need to discourage its use.
...
We should systematically arrange to depend on the free C# implementations as little as possible. In other words, we should discourage people from writing programs in C#. Therefore, we should not include C# implementations in the default installation of GNU/Linux distributions or in their principal ways of installing GNOME, and we should distribute and recommend non-C# applications rather than comparable C# applications whenever possible.

Make your own conclusions on this but personally, i've never been a fan of C# myself....not now and to be quite frank possibly never.

OpenMP - 'firstprivate' and 'lastprivate' caveat

2009-07-12T10:20:00.000+08:00

I've just started experimenting OpenMP and its another super cool multicore development paradigm as compared to NVIDIA's CUDA where the former concentrates on utilizing the actual CPU cores while the latter is on using the vast numbers of GPUs in the Nvidia's graphics cards.

To begin, let me do a demonstration of what 'firstprivate' or 'lastprivate' do.


extern int n; // external storage linked to some variable i've defined somewhere else.
void demo_firstprivate(void) {
 int i, indx, TID;
 int a[n];
 for(i = 0; i < n; i++ )
  a[i] = -i-1;
 indx = 4;
 int n1 = 1;
#pragma omp parallel default(none) firstprivate(indx) private(i, TID) shared(n1,a)
 {
  TID = omp_get_thread_num();
  indx += n1*TID;
  for( i = indx; i < indx + n1; i++)
   a[i] = TID + 1;
 }// end of parallel region
 
 printf("After the parallel region:\n");
 for( i = 0; i < n; i++ )
  printf("a[%d] = %d\n", i, a[i]);
}

The output is


Demo 'firstprivate' clause begin..
After the parallel region:
a[0] = -1
a[1] = -2
a[2] = -3
a[3] = -4
a[4] = 1
a[5] = 2
a[6] = -7
a[7] = -8
a[8] = -9
a[9] = -10
Demo 'firstprivate' clause end..

There are two things you need to realize when using firstprivate
(1) The firstprivate variable is initialized once per thread
(2) In C++, the firstprivate object is constructed by calling its copy constructor with the master thread's copy of the variable as its argument.

Demo of lastprivate is the following code


extern int n;
void demo_lastprivate(void) {
 int a = 0;
 int i;
#pragma omp parallel for private(i) lastprivate(a)
  for(i = 0; i < n; i++) {
   a = i + 1;
   printf("Thread:%d got value: %d, iteration: %d\n", omp_get_thread_num(), a, i);
  }

 // End of parallel region
 printf("value of 'lastprivate' variable 'a' is %d\n", a);
}

The output is


Demo 'lastprivate' clause begin..
Thread:1 got value: 6, iteration: 5
Thread:0 got value: 1, iteration: 0
Thread:1 got value: 7, iteration: 6
Thread:0 got value: 2, iteration: 1
Thread:1 got value: 8, iteration: 7
Thread:0 got value: 3, iteration: 2
Thread:1 got value: 9, iteration: 8
Thread:1 got value: 10, iteration: 9
Thread:0 got value: 4, iteration: 3
Thread:0 got value: 5, iteration: 4
value of 'lastprivate' variable 'a' is 10
Demo 'lastprivate' clause end..

The caveat in using the lastprivate clause is
(1) If the lastprivate variable is some sort of an array or structure and only some elements or fields are assigned in the last iteration; then after the parallel execution, the elements or fields that were not assigned in the final iteration are undefined.
(2) In C++, this variable/object needs to have its copy assignment operator invoked with the master thread's copy with the sequentially last value of the variable as the argument.

That is, both copy assignment operator and copy constructor must be publicly available otherwise you'll find yourself in quite a fix - basically a complier error will be reflected and depending on your compiler e.g. Xcode, VS, gcc you might be able to figure out why the error is there in the first place.

The last thing i wanted to share is that under OpenMP, variables pointing to heap storage are shared by all threads in the program. So you need to be careful while dealing with memory allocation otherwise you will get runtime errors like

Non-aligned pointer being freed ... or double free

Netflix has a winner! Well not if someone beats them to it in 30 days

2009-06-28T21:37:00.003+08:00

Caught this on New York Times while trying to catchup with the rest of the world. See the link here for more details. I've got a snapshot of that story here in case you want to just get the gees of it.

Turns out that they're not revealing their secret success yet (since they hope there are no challengers within the next 30 days) which is understandable but i sure would like to find out how that team engineered the solution.

CUDA and libcurl - a simple demo

2009-05-31T19:23:00.008+08:00

Recently, i got to know cURL and libcurl and wondered whether i could get sample apps to run under CUDA. However i could only get it to run under device emulation mode under CUDA for the simple reason that libcurl is not CUDA enabled if i may use this term.

So here's the sample run with runtimes and the source code.

The one below is run under simple multi-threading implementation

ray:src ray$ time ./multiDownload 1>/dev/null 2>&1

real 0m3.261s
user 0m0.014s
sys 0m0.028s

This one below is run under CUDA-device emulation mode

ray:src ray$ time ./cudaDownload 1>/dev/null 2>&1

real 0m31.893s
user 0m0.022s
sys 0m0.038s

The performance was disappointing while running under device emulation mode under CUDA and I looked into the runtimes to find an explanation as to why. Here are 2 notes i've made ...

1. Under device emulation mode, the CUDA still launched as many threads as implied by the source code(s) which is cool stuff. See gdb output of runtime below. From the output below, you can see that the threads are in the semaphore_wait_trap() which probably explains why the runtimes suck that much.


(gdb) info threads
33 process 28855 thread 0x5303  0x945b32c2 in semaphore_wait_trap ()
32 process 28855 thread 0x5103  0x945b32c2 in semaphore_wait_trap ()
31 process 28855 thread 0x4f03  0x945b32c2 in semaphore_wait_trap ()
30 process 28855 thread 0x4d03  0x945b32c2 in semaphore_wait_trap ()
29 process 28855 thread 0x4b03  0x945b32c2 in semaphore_wait_trap ()
28 process 28855 thread 0x4903  0x945b32c2 in semaphore_wait_trap ()
27 process 28855 thread 0x4703  0x945b32c2 in semaphore_wait_trap ()
26 process 28855 thread 0x4503  0x945b32c2 in semaphore_wait_trap ()
25 process 28855 thread 0x4303  0x945b32c2 in semaphore_wait_trap ()
24 process 28855 thread 0x4103  0x945b32c2 in semaphore_wait_trap ()
23 process 28855 thread 0x3f03  0x945b32c2 in semaphore_wait_trap ()
22 process 28855 thread 0x3d03  0x945b32c2 in semaphore_wait_trap ()
21 process 28855 thread 0x3b03  0x945b32c2 in semaphore_wait_trap ()
20 process 28855 thread 0x3903  0x945b32c2 in semaphore_wait_trap ()
19 process 28855 thread 0x3703  0x945b32c2 in semaphore_wait_trap ()
18 process 28855 thread 0x3503  0x945b32c2 in semaphore_wait_trap ()
17 process 28855 thread 0x3303  0x945b32c2 in semaphore_wait_trap ()
16 process 28855 thread 0x3103  0x945b32c2 in semaphore_wait_trap ()
15 process 28855 thread 0x2f03  0x945b32c2 in semaphore_wait_trap ()
14 process 28855 thread 0x2d03  0x945b32c2 in semaphore_wait_trap ()
13 process 28855 thread 0x2b03  0x945b32c2 in semaphore_wait_trap ()
12 process 28855 thread 0x2903  0x945b32c2 in semaphore_wait_trap ()
11 process 28855 thread 0x2703  0x945b32c2 in semaphore_wait_trap ()
10 process 28855 thread 0x2503  0x945b32c2 in semaphore_wait_trap ()
9 process 28855 thread 0x2303  0x945b32c2 in semaphore_wait_trap ()
8 process 28855 thread 0x2103  0x945b32c2 in semaphore_wait_trap ()
7 process 28855 thread 0x1f03  0x945b32c2 in semaphore_wait_trap ()
6 process 28855 thread 0x1d03  0x945b32c2 in semaphore_wait_trap ()
5 process 28855 thread 0x1b03  0x945b32c2 in semaphore_wait_trap ()
4 process 28855 thread 0x1903  0x945b32c2 in semaphore_wait_trap ()
3 process 28855 thread 0x1703  0x946026fa in select$DARWIN_EXTSN ()
2 process 28855 thread 0x1503  0x945b32c2 in semaphore_wait_trap ()
* 1 process 28855 local thread 0x2d03  0x945ba46e in __semwait_signal ()
(gdb) disassemble semaphore_wait_trap
Dump of assembler code for function semaphore_wait_trap:
0x945b32b8 : mov    $0xffffffdc,%eax
0x945b32bd : call   0x945b3ad4 <_sysenter_trap>
0x945b32c2 : ret
0x945b32c3 : nop
End of assembler dump.

2. The semaphore_wait_trap() resolves into launching a system call into the Mac OS X kernel which is not surprising as that's how CUDA implement kernels in device emulation mode and that causes most of the latencies since under this mode of execution, threads are executed on the CPUs and not on the GPUs.

Its expected that i could not cuda my sample application since its not possible to call host function from within kernel function (to borrow CUDA's terminology) so i had to compile and build it under device emulation but what this experiment demonstrated was that CUDA's device emulation mode may not be the answer that i was looking for but it raises my question "Wouldn't it be great if Nvidia could provide the software library to allow kernel functions to call host functions in the CUDA manner? " Perhaps its a work in progress.

A likely candidate for this sort of computing could be in OpenCL (Open Computing Language) and it'll be in the next Mac OS (Snow Leopard) Yay! Read the press release here.

Here are the sources codes i used (this multi-threaded program was lifted from the libcurl website's example code and i merely modified some stuff to fit my experiment)

#include <stdio.h>
#include <pthread.h>
#include <curl/curl.h>

#define NUMT 32

const char* const urls[NUMT] = {
"http://www.yahoo.com",
"http://www.cnn.com",
"http://www.hotmail.com",
"http://www.gmail.com",
"http://www.hp.com",
"http://www.microsoft.com",
"http://www.sun.com",
"http://blogs.sun.com/",
"http://www.acm.org",
"http://blogs.sun.com/d/",
"http://blogs.sun.com/jonathan",
"http://blogs.sun.com/jimgris",
"http://blogs.sun.com/theaquarium",
"http://blogs.sun.com/arungupta",
"http://blogs.sun.com/katakai",
"http://blogs.sun.com/webmink",
"http://blogs.sun.com/startups",
"http://blogs.sun.com/geertjan",
"http://blogs.sun.com/eclectic",
"http://blogs.sun.com/theplanetarium",
"http://blogs.sun.com/SDNProgramNews",
"http://blogs.sun.com/GullFOSS",
"http://blogs.sun.com/richb",
"http://blogs.sun.com/chrisg",
"http://blogs.sun.com/ontherecord",
"http://blogs.sun.com/HPC",
"http://blogs.sun.com/bblfish",
"http://blogs.sun.com/enterprisetechtips",
"http://blogs.sun.com/ahl",
"http://blogs.sun.com/jag",
"http://blogs.sun.com/bigadmin",
"http://blogs.sun.com/brendan"
};

static void *pull_one_url(void* url) {
 CURL* curl;
 curl = curl_easy_init();
 curl_easy_setopt(curl, CURLOPT_URL, url);
 curl_easy_perform(curl);
 curl_easy_cleanup(curl);

return NULL;
}

int main(int argc, char** argv) {
 pthread_t tid[NUMT];
 int i;
 int error;
 
 curl_global_init(CURL_GLOBAL_ALL);
 for( i = 0; i < NUMT; i++)
  pthread_create(&tid[i], NULL, pull_one_url, (void*)urls[i]);

 for( i = 0; i < NUMT; i++)
  pthread_join(tid[i], NULL);

return 0;
}

Here's the portion in CUDA-style (I've shown only the portion where its different)


__global__ void pull_one_url(char** url) {
 int tid = threadIdx.x;
 CURL* curl;
 curl = curl_easy_init();
 curl_easy_setopt(curl, CURLOPT_URL, url[tid]);
 curl_easy_perform(curl);
 curl_easy_cleanup(curl);
 printf("%d finished@%s\n", tid, url[tid]);
return ;
}

int main(int argc, char** argv) {
 char** d_a;
 int memSize=0;
 for( int i = 0; i < NUMT; i++)
  memSize += strlen(urls[i]);
 printf("size=%d\n", memSize);
 cudaMalloc((void**)&d_a, memSize);
 cudaMemcpy( d_a, urls, memSize, cudaMemcpyHostToDevice );
 for( int i = 0; i < NUMT; i++)
  printf("%s\n", d_a[i]);

 pull_one_url<<<1, NUMT>>>(d_a);

 cudaFree(d_a);
return 0;
}

RunSnakeRun - View performance profile in Python

2009-05-09T16:33:00.011+08:00

In my previous post on Greenlet/Eventlet, i went about testing and profiling the apps i wrote just to get a feel but admittedly, i do find reading the statistical data gathered thru profiling laborious and at times, it could get quite tiring..

So, i've found a nice application that can read python application's profiling data. Its called RunSnakeRun which is pretty cool as it helps to visualize the data in a much more interesting and readable manner and the caveat is that you will need to download the appropriate wxPython binaries or source code to build for the platform you're using. In this case, i'm using Ubuntu Linux to run the profiling.

Besides the proper wxPython, you will need to follow the instructions on the site to begin profiling and viewing your profiled data.

On a side note, if you are curious about profiling applications using python but have little idea what its all about, here's a link that gives you an idea why profiling is important in python code.

The screen shot i have for you is a caller-callee graph with runtimes (accumulated and exclusive) which allows a python developer to quickly isolate poor performing code and also provides you an idea how code is being executed. E.g. this screenshot illustrated the concurrent execution of the Greenlets

Obviously, this application allows a developer to quickly isolate problematic code using a square map since it highlights the largest consumer of time by the largest square and also provides color highlighting when hovering your mouse over the different squares. Neat!

Python's Greenlet/Linden Lab's Eventlet

2009-05-09T09:52:00.019+08:00

Have you guys used Python's greenlet or eventlet before? Its pretty cool and i thought i write something about it. So here's a little experiment i did using a couple of Python's module i.e. optparse, greenlet, eventlet. I have also did some simple measuring to see what was causing latencies etc.

What i did was basically, collected statistics on running a computation against the standard matrix multiplication which is highly parallel. There are 3 programs i created using the normal iterative-approach i.e. for-loop, greenlet and lastly eventlet implementation.

Overall, i find that greenlet is much more suitable to eventlet (w.r.t time) and the iterative-approach (w.r.t flexibility and elegance). I include my scripts below at the end of this post.

Here's how i executed and profiled my script:


ray:~ ray$ python -m cProfile -o eventletprof_10 ./myEventletDemo.py
Sum of matrix multiplication: 285

ray:~ ray$ python -m cProfile -o eventletprof_100 ./myEventletDemo.py --i 100
Sum of matrix multiplication: 328350

ray:~ ray$ python -m cProfile -o eventletprof_1000 ./myEventletDemo.py --i 1000
Sum of matrix multiplication: 332833500

ray:~ ray$ python -m cProfile -o eventletprof_10000 ./myEventletDemo.py --i 10000
Sum of matrix multiplication: 333283335000

ray:~ ray$ python -m cProfile -o eventletprof_100000 ./myEventletDemo.py --i 100000
Sum of matrix multiplication: 333328333350000
...
...
ray:~ ray$ python -m cProfile -o noneventletprof_10 ./myNoneventletDemo.py
iteration version of matrix add: 285

ray:~ ray$ python -m cProfile -o noneventletprof_100 ./myNoneventletDemo.py --i 100
iteration version of matrix add: 328350

ray:~ ray$ python -m cProfile -o noneventletprof_1000 ./myNoneventletDemo.py --i 1000
iteration version of matrix add: 332833500

ray:~ ray$ python -m cProfile -o noneventletprof_10000 ./myNoneventletDemo.py --i 10000
iteration version of matrix add: 333283335000

ray:~ ray$ python -m cProfile -o noneventletprof_100000 ./myNoneventletDemo.py --i 100000
iteration version of matrix add: 333328333350000
...
...
ray:~ ray$ python -m cProfile -o greenletprof_10 ./myGreenletDemo.py
iteration version of matrix add: 285

ray:~ ray$ python -m cProfile -o greenletprof_100 ./myGreenletDemo.py --i 100
iteration version of matrix add: 328350

ray:~ ray$ python -m cProfile -o greenletprof_1000 ./myGreenletDemo.py --i 1000
iteration version of matrix add: 332833500

ray:~ ray$ python -m cProfile -o greenletprof_10000 ./myGreenletDemo.py --i 10000
iteration version of matrix add: 333283335000

ray:~ ray$ python -m cProfile -o greenletprof_100000 ./myGreenletDemo.py --i 100000
iteration version of matrix add: 333328333350000

What profiling does is to output the results into a file (e.g. greenletprof_100000) and you can view the statistics in the Python interpreter using the following:


Python 2.5.1 (r251:54863, Jan 13 2009, 10:26:13)
[GCC 4.0.1 (Apple Inc. build 5465)] on darwin
Type "help", "copyright", "credits" or "license" for more information.
>>> import pstats
>>> profiledData = pstats.Stats('greenletprof_100000')
>>> profiledData.print_stats()

And you'll see the statistics printed in much detail. However, this is rather tedious so i would suggest a graphical UI to view the profiled data instead and one such option is RunSnakeRun.

Comparing the statistics in detail between the different runs reveals that the iterative approach is the most efficient followed by Greenlet and last eventlet. Here are the statistics in detail:


>>> import pstats
>>> stats10 = pstats.Stats('nongreenleteventlet100000.profile')
>>> stats10.print_stats()
Sun May 10 22:30:16 2009    nongreenleteventlet100000.profile

      6 function calls in 0.069 CPU seconds

Random listing order was used

ncalls  tottime  percall  cumtime  percall filename:lineno(function)
     1    0.060    0.060    0.069    0.069 /home/tayboonl/Desktop/Greenlet_Eventlet/nongreenleteventletdemo.py:8(execMatrixAdd_Iter)
     1    0.000    0.000    0.000    0.000 {method 'disable' of '_lsprof.Profiler' objects}
     1    0.000    0.000    0.069    0.069 :1()
     3    0.009    0.003    0.009    0.003 {range}

...
>>> stats10 = pstats.Stats('greenlet100000.profile')
>>> stats10.print_stats()
Sun May 10 22:16:45 2009    greenlet100000.profile

      500011 function calls in 1.511 CPU seconds

....
>>> stats10 = pstats.Stats('eventlet100000.profile')
>>> stats10.print_stats()
Sun May 10 22:20:17 2009    eventlet100000.profile

      6500027 function calls (5800031 primitive calls) in 34.901 CPU seconds

You might have noticed the remarkable number of functions calls made in the scenario using Eventlet. It suffice at this point for me to say is that Eventlet was not designed for high performance computing w.r.t Greenlet since it was designed to be a asynchronous networking library and to me, it makes alot of sense to always profile the application or framework before using it on a wide basis.

In case you are interested, here are the scripts i used

All together, there are 3 in total first being that implemented using Greenlet (Cool stuff!)


#!/usr/bin/python

import greenlet
import optparse

tasks = []
accum = list()
sum = 0
vecA = vecB = None

def mulNAccum(idx, val1, val2) :
global accum
accum.insert(idx, val1*val2)
#print idx,val1*val2,accum

def createTasks(HowMany):
global vecA, vecB
vecA = range(0, HowMany)
vecB = range(0, HowMany)
for i in range(0,HowMany):
       tasks.append(greenlet.greenlet(run=mulNAccum,parent=greenlet.getcurrent()))

def executeTasks(HowMany):
global tasks
global accum
global sum

for i in range(0,HowMany):
       tasks[i].switch(i, vecA[i],vecB[i])
#print "tasks finished executing ...\n"
# print accum
for i in range(0, len(accum)):
       sum = sum + accum[i]

print "Sum of matrix multiplication: %d\n" % sum

if __name__ == "__main__":
parser = optparse.OptionParser()
parser.add_option("--items", "-i", default=10, action="store", type="int", dest="HowMany", help="number of elements in matrix array to process")
(options, args) = parser.parse_args()

createTasks(options.HowMany)
executeTasks(options.HowMany)

Similar program using the for-loop a.k.a iterative approach


#!/usr/bin/python

import optparse

sum = 0
vecA = vecB = None

def execMatrixAdd_Iter(HowMany):
global vecA, vecB, sum
vecA = range(0, HowMany)
vecB = range(0, HowMany)
for i in range(0, HowMany):
       temp = vecA[i] * vecB[i]
       sum = temp + sum
print "iteration version of matrix add: %d\n" % sum


if __name__ == "__main__":
parser = optparse.OptionParser()
parser.add_option("--items", "-i", default=10, action="store", type="int", dest="HowMany", help="number of elements in matrix array to process")
(options, args) = parser.parse_args()
execMatrixAdd_Iter(options.HowMany)

Last script using Linden Lab's eventlet


#!/usr/bin/python

import eventlet
import eventlet.api
import optparse

tasks = []
accum = list()
sum = 0
vecA = vecB = None

def mulNAccum(idx, val1, val2) :
global accum
accum.insert(idx, val1*val2)
#print idx,val1*val2,accum

def createTasks(HowMany):
global vecA, vecB
vecA = range(0, HowMany)
vecB = range(0, HowMany)
for i in range(0,HowMany):
       tasks.append(eventlet.api.spawn(mulNAccum, i, vecA[i-1], vecB[i-1]))

def executeTasks(HowMany):
global tasks
global accum
global sum

for i in range(0,HowMany):
       tasks[i].switch()
#print "tasks finished executing ...\n"
# print accum
for i in range(0, len(accum)):
       sum = sum + accum[i]

print "Sum of matrix multiplication: %d\n" % sum

if __name__ == "__main__":
parser = optparse.OptionParser()
parser.add_option("--items", "-i", default=10, action="store", type="int", dest="HowMany", help="number of elements in matrix array to process")
(options, args) = parser.parse_args()

createTasks(options.HowMany)
executeTasks(options.HowMany)

Designing efficient sorting algorithms for manycore GPUs

2009-04-19T11:08:00.004+08:00

There's a link "Design efficient sorting algorithms for many core GPUs" which i recommend people to read to understand and think about how CUDA can help implement fast sorting algorithms which has benefitted me which i thought to share with everyone

------------ Updated 26 Apr'09 ----------

On a similar note, here's 2 illustrations from a recent report i read on Nvidia's CUDA which you can download by clicking here. I believe these two pictures would incite curiosity in you :)

CUDA 2.1 Nbody linkage issue under Mac OS X

2009-03-26T11:25:00.010+08:00

Just hours ago, i downloaded CUDA 2.1 from NVIDIA's website and the reason i did that was because a fellow CUDA enthusiast ran into some problem while compiling the samples and in particular N-body.

Well the problem was that there's a duplicate symbol/definition of the name beginWinCoords where a particular static library had 2 definitions and this is disallowed by the standards. So the resolution was to remove the directories under /Developer/CUDA and reinstall the CUDA 2.1 SDK entirely.

In detail, here's what happened.

When i first installed CUDA 2.1 ontop of CUDA 2.0 and next attempt to rebuild the sample projects, in particular Nbody i hit into the following problem:


ray:nbody ray$ make
ld: duplicate symbol beginWinCoords()     in ../../lib/libparamgl.a(paramgl.cpp.o) and ../../lib/libparamgl.a(paramgl.cpp_o)
collect2: ld returned 1 exit status
make: *** [../../bin/darwin/release/nbody] Error 1

What's happened is that there are two definitions of the same function i.e. beginWinCoords() in the static archived file libparamgl.a (.a is the file extension for static archive libraries instead of .dylib, .so, .dll)

Next thing i did was to compare the contents of the two static libraries and found that the 2.1 version had 1 definition (great) while the 2.0 version had two definitions. See below for the comparison (If you are unfamiliar with nm, consult your man page for your platform)


ray:Developer ray$ nm CUDA/lib/libparamgl.a | grep beginWinCoords
000016b0 s __GLOBAL__I__Z14beginWinCoordsv
00000000 a __GLOBAL__I__Z14beginWinCoordsv.eh
00000000 T __Z14beginWinCoordsv
000019c4 S __Z14beginWinCoordsv.eh


ray:Developer ray$ nm CUDA2.0/lib/libparamgl.a | grep beginWinCoords

000017e0 s __GLOBAL__I__Z14beginWinCoordsv
00000000 a __GLOBAL__I__Z14beginWinCoordsv.eh
00000000 T __Z14beginWinCoordsv
00001af8 S __Z14beginWinCoordsv.eh
000017e0 s __GLOBAL__I__Z14beginWinCoordsv
00000000 a __GLOBAL__I__Z14beginWinCoordsv.eh
00000000 T __Z14beginWinCoordsv
00001af8 S __Z14beginWinCoordsv.eh
000017e0 s __GLOBAL__I__Z14beginWinCoordsv
00000000 a __GLOBAL__I__Z14beginWinCoordsv.eh
00000000 T __Z14beginWinCoordsv
00001af8 S __Z14beginWinCoordsv.eh
000016b0 s __GLOBAL__I__Z14beginWinCoordsv
00000000 a __GLOBAL__I__Z14beginWinCoordsv.eh
00000000 T __Z14beginWinCoordsv
000019c4 S __Z14beginWinCoordsv.eh
000016b0 s __GLOBAL__I__Z14beginWinCoordsv
00000000 a __GLOBAL__I__Z14beginWinCoordsv.eh
00000000 T __Z14beginWinCoordsv
000019c4 S __Z14beginWinCoordsv.eh
000016b0 s __GLOBAL__I__Z14beginWinCoordsv
00000000 a __GLOBAL__I__Z14beginWinCoordsv.eh
00000000 T __Z14beginWinCoordsv
000019c4 S __Z14beginWinCoordsv.eh
ray:Developer ray$

The latter output indicated the two object files within libparamgl.a that caused the duplicate symbol/ definition issue. Okie Dokie, have fun!

Second Life featured in Channel New Asia

2009-03-17T08:24:00.004+08:00

Guess what?! Second Life is featured in Channel News Asia :). Read about the entire article here.

Here's an excerpt of that article from Linden Lab CEO Mark Kingdon

More than 1.3 million US dollars worth of transactions reportedly take place daily in Second Life, where the currency is the Linden dollar.   There are more than 15,000 merchants in Second Life selling snippets of computer code that become clothing, hair, art work or other items for avatars.   People spent 360 million dollars (US) in Second Life last year, according to Linden.   Schools continue to use Second Life for online classrooms and bands perform on in-world stages, albeit to sometimes meager audiences.   Linden is making avatar tools easier and "reworking the user experience," according to Kingdon.   "We have hired a world-class-team to lead the changes," Kingdon said. "You ain't seen nothing yet. A lot of work is going to be done in the next 9 to 12 months."

Its very very exciting to be working for this revolutionary company and with so many great talented people :D makes me wanna jump up from my bed everyday and say "What shall we achieve today!"

Maximizing application performance with Sun Studio 12

2009-02-26T10:24:00.013+08:00

Darryl Gove's blog entry has good notes on improving application performance on x86/SPARC platforms. I've pasted some stuff from his webex session but i do encourage you to listen to it, click here. Otherwise, i've taken some screen shots from the webex session and paste them here for your convenience.

I do apologize that i've not sorted them out, i'll get to them soon.

Notes on CUBLAS & CUFFT

2009-02-22T10:40:00.019+08:00

Some stuff you need to be aware of while using CUBLAS

CUBLAS uses column-major storage and 1-based indexing (because to maintain compatibility with Fortran and btw, i'm not expert in Fortran but i know at least this much about it); contrast this to C and C++'s of using row-major storage and 0-based indexing.

To overcome this while using CUBLAS in your C and C++ code, you can use macros or

inline functions to implement matrices on top of 1-dimensional arrays. As i code in C

and C++ primarily since i don't understand Fortran, the macro i used as recommended

in Nvidia CUDA's documentation is

#define IDX2C(i, j, ld) (((j)*(ld))+(i))

So here's a situation how i use the above macro to compute the index of an 1-D array

element using what looks like a 2-D matrix:

for( int j = 0; j < N; ++j) {

for( int i = 0; i < M; ++i) {

a[ IDX2C(i, j, M) ] = value;

}

...

// call a CUBLAS function (e.g. cublasScal) and using IDX2C to compute array index

cublasScal(...,..., &a[IDX2C(p,q,ldm)], ldm);

Include the header file cublas.h into your programs in case you forget. Normally, i don't use the syntax #include <cublas.h> but instead #include "cublas.h" for obvious reasons.
Remember to link your apps with the dynamic library provided cublas.so (Linux), cublash.dll (Windows) & cublas.dylib (Mac OS X) and they have the dynamic libraries for emulation purposes with naming conventions like cublasemu.so etc.

Notes of using CUFFT

The NVIDIA CUDA implemented their version of FFT using FFTW and follows what's known as a plan - which specifies the optimal or minimal number of flops for execution.
Depending on your graphics cards shared memory configuration, its best if you can fit your computation entirely in the CUDA's shared memory to minimize use of global memory.
Refer to the CUFFT documentation for more details as its continually evolving as hardware and software continue to mature over the next few months.

CUDA resources

2009-02-22T10:07:00.005+08:00

Here are links that illustrate how CUDA works and how it can help in maximizing the parallelism inherently in your Nvidia graphics cards.

Part1 - Introduction to CUDA

Part2 - Explanation of CUDA programming constructs

Part3 - Error handling and CUDA global memory limitations

Part4 - Understanding CUDA shared memory

Part5 - How to use CUDA shared memory effectively

Part6 - Profiling CUDA code

Part7 - Current and future implementations

Part8 - Using CUBLAS and CUFFT (CUDA versions of Fast Fourier Transformation & Basic Linear Algebra Subprograms)

Part9 - Extending CUDA to support other programming languages

Part10 - CUDPP

Part11 - Revisiting CUDA memory spaces

Personally, i like to thank Rob Farber for this series of REALLY awesome articles published on Dr. Dobb's Journal which is another great platform for developers world-wide. Yay!

While using CUDA to implement your test or commercial applications, its likely you'll create your own arrays/matrices or vectors which is fine but i would recommend that you re-implement that part of your program that uses these constructs using the CUBLAS library. Remember to try out the CUFFT library as well.

Second Life on Linux is out

2009-02-16T09:47:00.010+08:00

Second Life has been able to run on the Linux for some months now but there were glitches on the voice-chats but it looks good now. Have fun and try it out.

If anyone knows whether i've infringe any copyrights, do let me know and i'll take this down.

Getting started with CUDA 2.0

2009-02-01T10:29:00.023+08:00

I'm going to run a couple of stuff from this article onwards using Nvidia's CUDA technology - really cool stuff. But there're some glitches in the installation w.r.t installing CUDA 2.0 straight away and when you run the example code(s), you'll notice that you are running in emulation mode - i.e. its not running on the GPU (Graphics Processing Unit) which your machine has.

In my case, i have two Nvidia graphics cards loaded on my MBP 10.5.6 and when you first installed CUDA 2.0, you will not be able to run the examples or any of your code in the GPU itself because the CUDA driver was never installed !!! Nividia's website doesn't tell you straight that you need to install CUDA 1.1 Toolkit prior to installing CUDA 2.0 Toolkit so go ahead and do it.

Make sure the CUDA driver is installed and the way to check it is to go to /System/Library/Extensions/ and look for the file "CUDA.kext". Should be there by now and updated to CUDA 2.0 driver. Yay!

Once you are done with that, build the example code(s) given like this:

Go to /Developer/CUDA
Build the example codes using "make" or "make dbg=1, "make emu=1" and "make emu=1 dbg=1" which will create the directories of "release", "debug", "emu" etc.
Run the example code "deviceQuery" either in the directory "release" or "debug" etc but that's beside the point because they should all give the same result once the CUDA driver is installed properly

Here's the sample output after all these steps are done:

ray:release ray$ ./deviceQuery
There are 2 devices supporting CUDA

Device 0: "GeForce 9600M GT"
Major revision number:                         1
Minor revision number:                         337500
Total amount of global memory:                 536543232 bytes
Number of multiprocessors:                     1
Number of cores:                               8
Total amount of constant memory:               1 bytes
Total amount of shared memory per block:       16384 bytes
Total number of registers available per block: 8192
Warp size:                                     32
Maximum number of threads per block:           512
Maximum sizes of each dimension of a block:    512 x 512 x 64
Maximum sizes of each dimension of a grid:     65535 x 65535 x 1
Maximum memory pitch:                          262144 bytes
Texture alignment:                             256 bytes
Clock rate:                                    0.07 GHz
Concurrent copy and execution:                 No

Device 1: "GeForce 9400M"
Major revision number:                         3
Minor revision number:                         250000
Total amount of global memory:                 266010624 bytes
Number of multiprocessors:                     1
Number of cores:                               8
Total amount of constant memory:               1 bytes
Total amount of shared memory per block:       16384 bytes
Total number of registers available per block: 8192
Warp size:                                     32
Maximum number of threads per block:           512
Maximum sizes of each dimension of a block:    512 x 512 x 64
Maximum sizes of each dimension of a grid:     65535 x 65535 x 1
Maximum memory pitch:                          262144 bytes
Texture alignment:                             256 bytes
Clock rate:                                    0.07 GHz
Concurrent copy and execution:                 No

Test PASSED

Press ENTER to exit...

If the driver was not present, then all you'll get is an message "There is no device supporting CUDA"

Also make sure that the shell variables PATH and DYLD_LIBRARY_PATH are populated properly to something of the effect

export PATH=/usr/local/cuda/bin:$PATH
export DYLD_LIBRARY_PATH=/usr/local/cuda/lib:$DYLD_LIBRARY_PATH

Note: The above is for the Bash shell.

If you don't do the above, you're NOT going to be able to run any of the example codes with a error message like "dyld: Library not loaded: @rpath/libcudart.dylib" which refers to the runtime library of CUDA.

After all that is done, i ran the "oceanFFT" and if everything runs ay-ok on your machine, you would see something like what i have below:

If you check the output of the terminal screen, you will notice that CUDA is using my slightly more powerful graphics card i.e. GeForce 9600M GT instead of the GeForce 9400M and the reason is because i've configured my MBP's to do that via the System Preferences -> Energy Saver and select "Higher Performance->Power Adapter" and set "Computer Sleep" and "Display Sleep" to "Never" so that Mac will automatically pick out your more powerful graphics card :) neat trick i learnt from my colleagues at work! Go Mac!

Connect to Internet using Nokia E71 and Macbook Pro

2009-01-22T21:28:00.008+08:00

Basically, this weblog has been done by a previous writer, click here to read the original article but i've adapted this article for people living in Singapore and hold a Singnet account and wish to use their Nokia E71 to connect to the internet via the mobile phone.

Step-by-step:

============

1) Read Mathew Rothenberg's blog and get the script for Nokia E71. Follow the instructions.

2) Read the blog at here to get an general understanding. For techies, this should be easy but for non-techies it's pretty easy going too :D

3) For me, i'm using the USB cable to connect mobile phone and MacBook Pro OS X 10.5.6

4) For Singnet E-Ideas, you would need the following configuration

Telephone Number: e-ideas
User name: 65IDEAS
Password :

5) Click on the Advanced tab and make sure that you select HSDPA and CID=1 (Note that the CID could be different)

6) Click on "Connect" and if all goes well, then you should be connected.

Hope this helps

Mac OS X "Instruments"

2009-01-18T10:45:00.008+08:00

Been starting to use Mac OS for 1 month plus and i would say that it really offers developers a plethora of tools and for hardcore Mac fans, i don't think its a surprise but being a newbie user myself there are stuff which i am beginning to appreciate.

So the one thing i found was "Instruments" which is a cool tool to use as it allows me to profile the Second life viewer when it starts up etc. Using the user friendly UI, you can select the stuff you want to see/profile via the "Library" and simply click on "Record" to start and of course the same button allows you to "Stop" the profiling as well.

For starters, i opened up the Library as in the screen shot above to select what i would like to profile and next thing is to simply select the application and begin :) really easy and after that, voila! the stuff comes out like below

but the caveat of this is that my CPU is running close to 100% throughout a 5 minute session and i was using the default profiling...you can imagine what would happen if i select some other stuff

In this example of mine i was interested in object allocation and so i selected "ObjectAlloc" and one nice thing i discovered is that you can see how much memory is allocated by "drilling" down into any object on the column "Object Address" and the profiler will show you the object's memory allocations and de-allocations. Cool!

The age old question of whether PermGen is part of JVM memory

2008-12-22T20:13:00.004+08:00

This article of mine is lifted from the OpenJDK user group discussion on this topic

This is Jon Masamitsu's reply to the question (This has not been edited in any way and represents him 100%)

In the young generation there are 3 spaces - eden plus 2 survivor
spaces. This organization is to enable the collection of the young generation. See

http://java.sun.com/javase/technologies/hotspot/gc/gc_tuning_6.html

for more details. In calculating the total size of the heap, only 1 on the survivor spaces is being counted. That's because only
eden + 1 survivor space is available to the application for allocations.
That may be where the rest of the space is. If you add -XX:+PrintHeapAtGC
you will see something like

Heap PSYoungGen total 10752K, used 368K [0xf1000000,0xf1c00000, 0xfbc00000)
eden space 9216K, 4% used [0xf1000000,0xf105c308,0xf1900000)
from space 1536K, 0% used [0xf1a80000,0xf1a80000,0xf1c00000)
to space 1536K, 0% used [0xf1900000,0xf1900000,0xf1a80000)
PSOldGen total 24576K, used 0K [0xdb800000, 0xdd000000, 0xf1000000)
object space 24576K, 0% used [0xdb800000,0xdb800000,0xdd000000)
PSPermGen total 16384K, used 1489K [0xd7800000, 0xd8800000,
0xdb800000)
object space 16384K, 9% used [0xd7800000,0xd7974540,0xd8800000)

Which will tell you the sizes of the survivor spaces. By the way, the
perm generation is separate from the Java heap.

Notice the last statement he says "the perm generation is separate from the Java heap". I hope this statement more or less settles the issue i have been hearing for 10,000 years (You can verify this by downloading and browsing the code for it ... it will take you some time but the benefits are tremendous :D )

Homegrown store with Second Life

2008-12-18T08:16:00.008+08:00

Today in the newspapers, i saw a piece of news which pleased me. The Singapore store "Tangs" is going to set up a 3-D replica in Second Life. Read about the news by clicking here. Its warming to see that malls and stores in Singapore are beginning to appreciate 3-D virtual reality. Havingworked in Linden Lab for a couple of weeks, this is a place that really excites me because here's a place where people make it happen. The variety of technologies and the "lindens" (that's what we are called within the company, special isn't it? I thought so too ... has a familial bonding effect to that ring) here are special and dedicated to their craft. Awesome :)

Having said that, its time to go back to my Second Life (yeah, its corny i know)

Python Performance and the GIL

2008-11-20T10:30:00.002+08:00

Recently developed a simply multi-threaded Python program (CPU-bound) and tried its performance on my AMD Phenom X4 Quad-Core system. The results were not what i expected. After some investigation and R&D, one cause could be the way i write my application (i always suspect myself first) and the second would be the infamous GIL (Global Interpreter Lock).

Based on Amdahl's law, i should obtain roughly a four-fold improvement in runtime since the system-under-test (SUT) has 4 processing cores but because of Python's semantics and requirements i.e. GIL, there isn't a way to really push the performance of the application's runtime.

Why do i say that? To my understanding, the GIL is being acquired/release for almost every operation you do in Python and based on the specification of the language, it is possible to control the frequency of the lock/release cycle in GIL by changing it via sys.setcheckinterval(). The side effect of this is that it would increase the frequency of thread context switches and this means python checks events at a higher/lower rate (depending on the passed-in value) which might spell a good or bad thing depending on the nature of your application i.e. I/O-bound or CPU-bound. For those Solaris buffs, its something like "set hires_tick ..." not exactly but thinking along that line doesn't hurt.

There's a book called "Python Cookbook" that illustrates the issue and i draw a paragraph from there.

To make life easier for programmers, the interpreter releases and reacquires the lock every 100 bytecode instructions ( a value that can be changed using sys.setcheckinterval).
...
However, effective performance-boosting exploitation of multiple processors from multiple Python threads of the same process is just not in the cards.
...
you will not observe substantial performance increases by moving your multithreaded application to a multiprocessor machine

If i read this book earlier, i probably would not have done this exercise but its interesting to understand why it happens. As part of this "validation" exercise, i browse into the Python 2.6 sources, and here's what i think the author meant.


file: ------ ceval.c ------
/* for manipulating the thread switch and periodic "stuff" - used to be
per thread, now just a pair o' globals */
int _Py_CheckInterval = 100;
volatile int _Py_Ticker = 100;

These internal C variables defined the "interval" which i think the author meant and there's a C API that serves as the implementation of sys.setcheckinterval() in Python which happens to be


static PyObject *
sys_setcheckinterval(PyObject *self, PyObject *args)
{
      if (!PyArg_ParseTuple(args, "i:setcheckinterval", &_Py_CheckInterval))
              return NULL;
      Py_INCREF(Py_None);
      return Py_None;
}

And there are a couple of C-signal functions which uses the values of _Py_CheckInterval & _Py_Ticker to perform what the author has described but i omit them here.

However, there is a caveat to increase the frequency of this cycle and that is it will likely increase the occurences of race condition within the signal handlers i.e. your program's correctness. You can find plenty of such issues by google-ing

Building Python for Performance with Sun Studio

2008-11-12T12:36:00.007+08:00

This article is lifted from Sun's Wiki page on Python performance and you can view the entire article by clicking here. Thanks to Chi Hung for sending me this link.

In a nutshell, the link reports performance improvements on using Sun Studio to compile Python from source on 3 of Sun's processors namely the UltraSPARC IV+, Woodcrest & Niagara2 using compiler flags like "-XO4" and "-xipo" which contains optimizations like inlining, cross source file optimization among others.

Check out the link for the benchmarks. One interesting thing was that Sun did not publish any benchmarks for Sun Studio on Linux-based systems; second interesting thing was that there's a piece of software called "parrot benchmark" which contains a iterable option (in the link, the author suggests to change the profiling to occur 20 times instead of the original 2 times) which is used during the compilation process to allow the compiler to profile Python (i assume to run the binary against some benchmarking software, collect the results in a temp place) and optimize portions of it.

Haven't used much of this parrot benchmarking software so i am not sure what is the simulation procedure used inside it to measure which parts of Python needs to be optimized. Interesting.

CMT Presentation in SecondLife

2008-10-29T01:01:00.007+08:00

Today at 9am PST (1AM Singapore time), Darryl Gove from Sun Microsystems presented on Chip Multi Threading in Second Life, and it was great to hear him live (even though there were some issues with the sound) speak about the latest compiler (Sun Studio 12) / software technologies like POSIX, OpenMP, Autoparallelization are helping developers write fuss free, if i may, multi threaded code. Here's the link to the presentation slides.

Above are snapshots captured during the live session with him and a bunch other people from all around the world.

Psss...in case you didn't know, SecondLife is capable of much more ... download it (it's free) .. run it .. enjoy!

Saturday Night Live

2008-10-24T21:25:00.005+08:00

You got to see the video, Saturday Night Live and its a full episode :). Depending on the political sensitivity of your geography, you might or might not have access to it. But this episode has Will Ferrell, Tina Fey and others playing out the ongoings of the US Presidential Elections. Fyi, i am not for any US party but i do find this video extremely entertaining.

Here's the url http://www.nbc.com/Saturday_Night_Live/video/episodes/#vid=784241

J2EE/WebLogic Tuning project ... continued (Part 3)

2008-10-12T11:57:00.018+08:00

Finally, the tuning project came to a close (at least for now) and there were many takeaways from this experience with my customer which i like to share. Overall, i felt that the root cause of these issues were typical of many software development projects in J2EE and that was not understanding the technology well enough (not to mention project costing and time schedules which compromised software quality...after all these years, you would think that people got it...but apparently they don't...)

Understanding EJBs, Transactions and SQL
That application had more than 3,000 bean instances in active transacting but the average transaction rate was less than 5 per second, which highlighted an important fact that its heavily clogged. The real clog was really poor performing SQLs, injudicious use of CMT and BMT i.e. Container/Bean Managed Transaction, mixing CMT & BMT in the same bean, "chatty" EJBs. The EJB specifications clearly say that its best not to mix CMT & BMT as it makes transaction management complicated also never EVER place global variables of the kind that represent database connections, user preferences etc in stateless beans except of course, stateful beans. Since the concurrency now is pretty low, those kinds of issues related to concurrency won't surface but you can be sure they will surface once performance improves and you don't want these kinds of bugs are the HARDEST to locate.

It was so slow to the point that WebLogic actual emitted an "STUCK THREAD" message and these occur when execution exceeds 10 minutes.

The next thing to do is to re-look at the design of those beans and remove those CMT & BMT mixtures and to convert globals into locals.

Understanding Clustering Technology
Clustering can be so easily abused by developers and lack of design control could worsen the entire situation. This project's implementation resulted in creating a 1 cluster group of 12 machines and its observed that Minor + Major GCs were happening on such a frequency and pause times were worsening that it totally wrecked the concurrency of the application; after which they changed it to 3 cluster groups which significantly reduced GC (however it still wasn't good enough) ... i suggested to them to reduced to several cluster groups of 2 servers which should/will significantly reduced overall heap size usage, frequency of GCs (and GC pause times) since there's less pressure on the Java memory system now that there should see less replication activity.

Developers placed huge objects (> 500KB) into the session via attributes combined with a fact that session timeouts ranged in hours (can't do anything since that's a requirement of the business) implied that memory can't be freed during GCs. But i suspect its probably an application design issue since these objects tend to reflect database data displayed on the GUI and all those nifty controls that gives users flexbility in the business functions.

Understanding Garbage Collection
GC are quite a killer of typical J2EE applications and its important to keep the pause times as low as it is possible. One major improvement was in using the Concurrent Low Pause Collector i.e. -XX:+UseConcMarkSweepGC which allowed GC to happen in a concurrent fashion in both the Young and Tenured Generation. Before this change, the Major GC time took 10 - 17 secs on average occurring once every 3 hours on a typical day; and 15 - 20 secs on average occuring once every 30 - 45 minutes on a busy day. However, i am skeptical whether this will work because the average object creation rate is 800 MB per minute (since Minor GC occurs twice per minute and each time it ocurred freed approximately 400 MB of memory)...there's still a chance that the Stop-The-World GC will happen; to understand why do read Jon Masamitsu's blog here.

Class Loading
The developers of the application did something which is very commonly seen in development and that is dumping their classes into the application server's classpath which not only prevented the class from unloading and what's worse is that a type of memory leak a.k.a Class Loader Leaks occur and GC doesn't collect this memory space. I admit i used to do this but i don't now. The effect was OutOfMemoryError: PermGen Space and of course increasing the Perm Gen size would effectively stamp out this error.

Of course, now my customer is pursuing the Queue OverFlow exception with BEA Systems and let's see what happens next ... perhaps a patch or BEA/Oracle will say " Upgrade your WebLogic 10, it solves that problem !" Psss tell you a secret the problem is not resolved there either ...

J2EE/WebLogic Performance Tuning Project...continued

2008-10-02T16:09:00.007+08:00

In my previous article, my tuning project with a customer ran into some trouble with WebLogic's Work Manager and in particular, on the Java exception weblogic.utils.UnsyncCircularQueue$FullQueueException where the WebLogic server indicated that the queue where the server works on submitted requests. Checked the WebLogic docs and accordingly, the server will automatically resize the queue to fit the requests but what the docs didn't mention was that the resize will fail to work if the size of the queue equals the capacity which happens to be 65536 and that's the reason why it threw the error message "Queue exceeded the maximum capacity of '65536' elements".

However, checking the code reveals something quite peculiar and that is the constructor suggests that only queues of sizes exceeding 1 GB will throw this error but the default capacity is always 256 and reaching a maximum of 65536 elements and btw, its WebLogic 9.2 ; so my guess is that source codes need to be cleaned up ? If anybody has any idea why, do drop me a comment, thanks in advanced.

J2EE/WebLogic Performance Tuning Project

2008-09-30T09:25:00.025+08:00

This article contains a series of investigations for a customer of mine where the environment is running a WebLogic cluster of 20 machines in round-robin on HP-UX to service a global J2EE application and it performed slowly during peak periods and occasional hangs. The application was a typical 3-tier architecture whereby web relegates requests to the middle-tier (EJBs, MQs, MDBs) and this middleware goes to the Database-tier (SQL inserts, updates, deletes, stored procedures etc). The application was found to be experiencing heavy load during peak periods everyday.

There were a couple of issues related to poor performing SQLs, poorly designed middleware apps, WebLogic cluster design and runtime issues, JVM memory consumption and frequent garbage collections. Let me try to detail them a bit without giving away too much customer information. Hopefully, it can help you in your investigations in your environment.

During the peak period, the major contributing factors of the apps slow-ness were:

The heap size was 1.5GB (min,max), 512MB for Eden and the PermGen was 192MB. The minor GC kicked in frequently releasing approximately 60MB on average; the major GC kicked in twice every minute (avg. 3-5s on average, 40s on max) releasing 400 - 500 MB each time and reverse engineering the figures reveals that the object creation rate was roughly 800 M - 1.0 GB per minute. As GC is primarily a CPU-intensive operation (with saving state, freeing memory, compacting the heap etc). The large object creation rate combined with the relatively long pauses GCs occurences suggests that the application are creating objects in an in-efficient manner and that created problems with the cluster's session replication mechanism as the users of the system would see stale data - due to long pauses in GC, the data in the session was not replicated *properly* across to the other servers.

Applications were attached to the WebLogic system classpath which meant that the Java classes were never unloaded from memory and combined with the fact that there are ALWAYS classloader leaks meant that whenever the operation team redeploy a.k.a "hot"-redeployment the apps, it worsens the memory footprint since the previous memory was never release due to this leakage. If you keep hot-deploying these stuff you will almost certainly get an OutOfMemory Error: PermGen out of space.

EJBs (4,000+ EJBs deployed, in my opinion too many) were utilizing remote interfaces when there was no need as those apps were not doing cross-vm operations and based on my previous experimentation, you would get a 3-fold runtime improvement when you convert the EJBs to local interfaces. This improvement is because there is less object marshalling/unmarshalling via RMI since everything is on the same JVM heap and consumes less system resources like file descriptors/socket & memory since local interfaces implies a local/normal Java call.

As i mentioned previously, the apps were deployed in the cluster and that meant that all persistent objects (e.g. session data, user preferences etc) must be Serializable (i.e. persistent objects need to implement java.lang.Serializable) since there would be session replication across the servers in the cluster which further degraded the performance as the cluster needs to maintain state across all 20 machines. Source code analysis found that user's were keeping results of database fetches in session data! You can imagine the pressure faced by the JVM memory subsystem + WebLogic cluster replication.

WebLogic cluster was also malfunctioning during peak periods throwing an exception message like <WorkManager> <BEA-002911> <WorkManager weblogic.kernel.System failed to schedule a request due to weblogic.utils.UnsyncCircularQueue$FullQueueException: Queue exceed maximum capacity of: '65536' elements and this is an critical error thrown from the Work Manager which replaced the BEA traditional thread pool. What this meant was that the WebLogic cluster could no longer handle user's requests and hanged. *I plan to unravel this mystery in a while to understand why this is happening*

The hardware loadbalancer was in "sticky" mode even though the WebLogic cluster was in round-robin mode which negated this round-robin-ness and resulted in certain servers encountering more stress than others and this was made worse by the long session timeout of 20+ hrs. That's the cost of doing business....

After tracing the SQL statements execution times, it was found that they were causing alot of problems from missing indexes, lack of functional indexes, improper SQL statements which causes large database table joins and many "select count(*)..." from large table joins statements contributed to this object creation rate.

When i looked at these issues, the first couple of items i advised my customer was to do the following:
(1) Convert the EJBs to use local interfaces i.e. call-by-reference
(2) Tune the SQL statements via SQL reordering, indexes etc
(3) Tune the JVM heap to use more aggressive + parallel heap collectors via -XX:+UseParallelGC -XX:+UseConcMarkSweepGC (We are still experimenting this portion)
(4) Do not use system classpath to load application classes
(5) Review source codes to remove known classloader leaks

The customer and myself are still in the process of implementing/reviewing this so i hope to have an update for you in a couple of weeks time. Meanwhile, do visit Jon Masamitsu's blog for an understanding of some JVM tuning parameters.

"man" crashes on my Lenovo T61 OpenSolaris snv_95

2008-09-30T09:24:00.000+08:00

Didn't expect "man" to crash on me today. But it's an interesting problem because when i tried the same command that crashed on my OpenSolaris snv_95 on another machine running Solaris (SunOS raymond 5.10 Generic_118855-36 i86pc i386 i86pc) it didn't crash but gave me instead a error message "No manual entry for make" which is good since its more friendly to the user and i found similar behavior on Ubuntu Linux.

After investigating for a while, i believe its a bug. The command i attempted was


man -M /usr/bin/man make

You will noticed that i made a mistake and its intentional. "-M" is suppose to mean the path but in this case i gave the absolute path to a file (which is executable btw) and when i ran on my OpenSolaris it generated a coredump.

After examining the code + coredump, it appears to me that the reason of the crash is because the program attempted to release memory via "free" but wasn't assigned previously. Filed a bug report with OpenSolaris. Keep you updated on this when i have news.

This is an update on this post. Here is the bug id issued, click here for more details.

Building OpenJDK 7

2008-09-12T11:20:00.012+08:00

I guess alot of people already knew that Sun Microsystems has released the source code for Java. Well i tried my hand at compiling the OpenJDK 7 and its important you read the installation manual before you attempt to compile it. Reason i did it was because i wanted to try out the new features of the JVM and also to understand the build-release process. Of course, i am still learning it and this post is about my experiences of compiling it and hope that it can provide you useful information on building your own, customizing it, fixing bugs if you find it etc

What i did first was to read the instructions found here. Read it a couple of times to understand exactly what you need to do. For my experience, here's what i did

(1) Alter my .bashrc file to include the environment variables needed


ALT_BINARY_PLUGS_PATH=/export/home/tayboonl/build_jdk/openjdk-binary-plugs
ANT_HOME=/export/home/tayboonl/apache-ant-1.7.1
ALT_COMPILER_PATH=/opt/SunStudioExpress/bin
ALT_GCC_COMPILER_PATH=/usr/bin/
ALT_CUPS_HEADERS_PATH=/usr/include
ALT_JDK_IMPORT_PATH=/usr/jdk/jdk1.6.0_06
LANG=C

export ALT_BINARY_PLUGS_PATH
export ALT_CUPS_HEADERS_PATH
export ALT_COMPILER_PATH
export ALT_GCC_COMPILER_PATH
export ALT_JDK_IMPORT_PATH
export ANT_HOME
export LANG

(2) Invoke the sanity check as defined in the build instructions


gmake sanity ARCH_DATA_MODEL=32

Note: Its important to fix all the warnings and errors before proceeding. If in doubt, check the openjdk forums. Remember to set the ALT_* variables, they are pretty critical to the success of the build. Also remember to install findbugs

(3) Re-run the sanity checks till everything is fixed, then start the build (From the output below, you can tell i am building the JDK for 32-bit instead of 64-bit)


gmake ARCH_DATA_MODEL=32

(4) Wait a while...grab a coffee...grab a bagel...sleep a little
If you have reached this stage, you will probably find that your terminal screen is scrolling away building , compiling etc and you might have hit this error

"Error: ia_nice is not a member of iaparms."

This is a bug 6712505 and you can comment out the offending line. I wonder why didn't the folks at Sun took this away...

Next thing i encountered was a build error where it complains that it cannot find the file "sys/audio.h", "sys/audioio.h", "sys/mixer.h" and i subsequently found the 3 files from src.opensolaris.org and placed them into their proper directories.

After that, re-running the build took another hit in the form of attempting to locate the file "X11/Intrinsic.h" and some other header files from the X11 package and on OpenSolaris its known as SUNWxwinc. Well, here's the irritating part whereby the OpenSolaris package manager tells you that it has the files but when you visit the folder(s), i simply couldn't find it so i have to re-install the package SUNWxwinc and the header files are there.

(5) Finally, the entire build process kicked off without any further disappointments and i see IT, i saw the message


gmake[2]: Leaving directory `/export/home/tayboonl/build_jdk/openjdk/jdk/make'
gmake[1]: Leaving directory `/export/home/tayboonl/build_jdk/openjdk'
Control solaris i586 1.7.0-internal build_product_image build finished:
Control solaris i586 1.7.0-internal all_product_build build finished:
Control solaris i586 1.7.0-internal all build finished:
tayboonl@opensolaris:~/build_jdk/openjdk$

Next is to invoke the commandline and when you see something like


tayboonl@opensolaris:~/build_jdk/openjdk/build/solaris-i586$ ./j2sdk-image/bin/java -version
openjdk version "1.7.0-internal"
OpenJDK Runtime Environment (build 1.7.0-internal-tayboonl_2008_09_12_11_16-b00)
OpenJDK Server VM (build 14.0-b04, mixed mode)

You know you are DONE. Woo Hoo! What a relief, now i hope the next build is not going to make me weep even further but its best that you test out the build with the demo apps found in your latest build. Here's a snapshot of the ArcTest demo found after completing the build

Last thing i did was to dump the JVM and from the output it appears that its fine


2008-09-12 14:00:35
Full thread dump OpenJDK Server VM (14.0-b04 mixed mode):

"TimerQueue" daemon prio=3 tid=0x08320c00 nid=0x11 waiting on condition [0xb6d7e000..0xb6d7ebf0]
 java.lang.Thread.State: WAITING (parking)
 at sun.misc.Unsafe.park(Native Method)
 - parking to wait for  <0xf3d65360> (a java.util.concurrent.locks.AbstractQueuedSynchronizer$ConditionObject)
 at java.util.concurrent.locks.LockSupport.park(LockSupport.java:186)
 at java.util.concurrent.locks.AbstractQueuedSynchronizer$ConditionObject.await(AbstractQueuedSynchronizer.java:1974)
 at java.util.concurrent.DelayQueue.take(DelayQueue.java:209)
 at javax.swing.TimerQueue.run(TimerQueue.java:170)
 at java.lang.Thread.run(Thread.java:674)

...

Hopefully, what i have done would go to show that you can do it too. Now it's time for lunch.

Rules for Writing 64-bit Clean Code

2008-09-02T21:23:00.005+08:00

I lifted this from the book "Solaris Systems Programming" by Rich Teer. This is good advice. So no surprises there and if you know Rich Teer will feel uncomfortable about this, let me know and i will gladly take this off.

The following rules should be observed to ensure the 64-bit cleanliness of our code. Following these rules will also make it easier to port 32-bit code to the 64-bit model.

Do not assume a pointer and an "int" are the same size. Unfortunately, a lot of code relies on this assumption, because it is true in the ILP32 model. Pointers are sometimes cast to "int"s or "unsigned int"s to perform address arithmetic. Instead, they should be cast to "long" (or "unsigned long") because pointers and "long"s are the same size in both the ILP32 and LP64 data type models. Even better is to cast pointers to "uintptr_ts", because it expresses the intent more clearly, and make the code more portable.
Do not make assumptions about the relative sizes of variable types. A classic example of this is to assume that the size of an "int" is the same as the size of a "long" and use them indiscriminately while implicitly or explicitly assuming they are interchangeable. Although this is typically true for 32-bit processes, it is not true for 64-bit ones.
Be wary of sign extension problems. This is quite a common problem when converting code to 64-bits, and is hard to detect before it actually occurs because "lint" doesn't warn us about it. Also, the type conversion and promotion rules are quite obscure. Hence, we should use explicit casting to fix sign extension problems.
Use pointer arithmetic rather than address arithmetic. As well as leading to cleaner code, pointer arithmetic is independent of the data model...
By default, external variables and functions are assumed to be or return an "int". by the compiler unless we declare them otherwise....in the ILP32 model its OK but on the LP64 then we will loose information which will cause the program to crash because of an illegal memory reference.

There's a lot more where that came from ... which i suggest you buy a copy (new or used doesn't matter)

Update OpenSolaris using Boot Environment Administration tool

2008-08-30T20:03:00.005+08:00

Was updating my OpenSolaris snv_93 to snv_95 when all of a sudden my wireless connection was dropped and the update was incomplete, so the next thing to do is to remove it. So i removed it by first unmounting the volume and invoke beadm to remove it like as follows unless it's going to be pretty irritating to see it when i reboot my laptop.


-bash-3.2# umount rpool/ROOT/opensolaris-5/opt
-bash-3.2# umount rpool/ROOT/opensolaris-5
....
-bash-3.2# beadm destroy opensolaris-5
Are you sure you want to destroy opensolaris-5? This action cannot be undone (y/[n]):
y
The BE that was just destroyed was the 'active on boot' BE. opensolaris-4 is now the 'active on boot' BE. Use 'beadm activate' to change it.
...
-bash-3.2# beadm list

BE            Active Active on Mountpoint Space
Name                 reboot               Used
----          ------ --------- ---------- -----
opensolaris   no     no        -          66.89M
opensolaris-4 yes    yes       /          12.54G
...

If you don't un-mount it, then you are going to get error messages like "Unable to destroy XXX"

Quitting my day job

2008-08-22T20:12:00.005+08:00

I was hanging around a Starbucks caffe today doing my usual research, browsing the internet, reading my "Learn Python" oreilly book today when i was approach by a caucasian male and something really funny happen

Caucasian: I am sorry to bother you but i need to ask you a favour
Me: Errr...sure, what can i do for you?
Caucasian: I bought a USB hard drive from "funan centre" (here's a picture of the IT super mall a.k.a "Funan Centre")

and wanted to test it out
Me: Errr...*puzzled look* (which translates to didn't the store tested out for you?)
...
subsequently, we tested it out and chatted his travels, what he liked/disliked about Singapore (my country)
...
... 5 mins later (yeah no doubt! all that happened in 5 mins)
Caucasian: Ahhh...yes i can see the drive and there's a song
Me: (I looked at the song and recognized it to be Rihana's umbrella)and immediately volunteered to hymm a couple of lines
Caucasian: *giggle* and finally broke out into a ROAR!
Caucasian: Please...Please...DONT QUIT YOUR DAY JOB
...

And we both broke out into a good laugh!

OpenSolaris Pydoc browser cannot be launched

2008-08-16T09:59:00.005+08:00

Got inspired by my friend to learn Python and i am beginning to appreciate its power and flexibility and it has elements from another favourite of mine Erlang. So following his example, i got a book called "Learning Python - 3rd Edition" and begun exploring the language and when i tried to launch the pydoc (fyi, its a tool for displaying available APIs) the script was telling me that it couldn't find the browser.

Here was the stack trace and like any other language, the "root" exception/cause is always found at the last line (look at the line in bold)


tayboonl@opensolaris:~/Desktop$ Exception in Tkinter callback
Traceback (most recent call last):
File "/usr/lib/python2.4/lib-tk/Tkinter.py", line 1345, in __call__
  return self.func(*args)
File "/usr/lib/python2.4/pydoc.py", line 2086, in open
  webbrowser.open(url)
File "/usr/lib/python2.4/webbrowser.py", line 43, in open
  get().open(url, new, autoraise)
File "/usr/lib/python2.4/webbrowser.py", line 38, in get
   raise Error("could not locate runnable browser")
Error: could not locate runnable browser

So, i looked at the file /usr/lib/python2.4/webbrowser.py and discovered that it was attempting to execute a list of known browser commands and in my case i simply added the string "firefox" to the list and it worked. *A quick hack*. Here the code snippet that i edited on the file.


   # X browsers have more in the way of options
   if os.environ.get("DISPLAY")$
       _tryorder = ["galeon", "skipstone", "firefox",
                    "mozilla-firefox", "mozilla-firebird", "mozilla", "netscape",
                    "kfm", "grail"] + _tryorder

DTrace and Tail-Call Optimization

2008-08-08T10:53:00.004+08:00

In many of the C compilers today, there's an optimization known as tail-call optimization and you need to pay careful attention to this when using DTrace's Function Boundary Tracing a.k.a fbt provider. A common application would be to tail-recursion.

Having said that, you probably need to realize that it occurs most often in SPARC systems than compared to x86 Solaris systems. How can you tell that you ran against a tail-call optimized routine or function ?

In DTrace, the variable arg0 contains the assembly return instruction with its offset in the DTrace return probe. So in your scripts you need to put in a extra statement that traces this variable and check whether this offset is a pure return instruction or does it point to something else otherwise you might get a nasty shock while trying to evaluate a result thinking it was the result of a function and in some cases, produce misleading information.


# dtrace -n fbt::squeue*:return'{printf("%s, 0x%x\n", probefunc, arg0);}'

and you should see something like


squeue_enter_chain,0x3a5
squeue_enter,0x47a
squeue_enter,0x47a
squeue_fire,0x85
^C

and next you need to use mdb or your favourite disassembler and look at the function at that offset. In my case, i didn't find any tail-call optimized functions since i saw the ret-instruction but in your case you might. Be aware to take care on this when you are tracing the Tcl/Python code since there are providers for them now.

Solaris 10 Network Stack

2008-08-06T14:00:00.003+08:00

The new implementation in the latest Solaris 10 and OpenSolaris is different and faster. Click here for details. Alternatively, download a pdf version of it here.

Here's a picture of the latest architecture

The new implementation is built on re-using the Solaris STREAMS framework found in the pre-Solaris 10 OS. Using the picture below, you can see how its evolved through the versions and after reading the PDF, i believe you would have a better idea of this network implementation.

And DTrace currently supports the tracing of this and here's a sample output of tracing it on my OpenSolaris 2-CPU machine


tayboonl@opensolaris:~/SunStudioProjects/TimeServer# dtrace -s ./trace.d -c ./dist/Debug/Sun12-Solaris-x86/timeserver
dtrace: script './trace.d' matched 2 probes
CPU FUNCTION                              
1  -> squeue_getprivate                     Bound to CPU(1),Squeue Name:ip_squeue_cpu_1/1/0, Thread ID:380

           ip`tcp_get_conn+0x22
           ip`tcp_open+0x1c6
           ip`tcp_openv4+0x24
           genunix`qattach+0x160
           genunix`stropen+0x490
           sockfs`socktpi_open+0xac
           sockfs`sotpi_create+0x11e
           sockfs`so_socket+0x146
           unix`_sys_sysenter_post_swapgs+0x14b
OS Thread ID:63619

0  -> squeue_getprivate                     Bound to CPU(0),Squeue Name:ip_squeue_cpu_0/0/0, Thread ID:290

           ip`tcp_time_wait_collector+0x20
           genunix`callout_execute+0xbf
           genunix`taskq_thread+0x1a7
           unix`thread_start+0x8
OS Thread ID:34

1  -> squeue_getprivate                     Bound to CPU(1),Squeue Name:ip_squeue_cpu_1/1/0, Thread ID:380

           ip`tcp_time_wait_collector+0x20
           genunix`callout_execute+0xbf
           genunix`taskq_thread+0x1a7
           unix`thread_start+0x8
OS Thread ID:35
...

What you are seeing is that there is a 1-1-1 ratio to thread:cpu:squeue and what you see is something like
"squeue name: ip_squeue_cpu 1/1/0 is bound to CPU-1 where the bounded thread id is 380" and next you can see that the threads id 34,35, 63619 is interacting with the 2 squeues.

Here's another method you can try using mdb and attaching to the process in question


> ::squeue
          ADDR STATE CPU            FIRST             LAST           WORKER
ffffff01ca044dc0 02060   1 0000000000000000 0000000000000000 ffffff0007ec2c80
ffffff01ca044e80 02060   0 0000000000000000 0000000000000000 ffffff0007c25c80
> ffffff0007ec2c80::thread
          ADDR    STATE  FLG PFLG SFLG   PRI  EPRI PIL             INTR
ffffff0007ec2c80 sleep       8    0    3    60     0   0              n/a
> ffffff0007c25c80::thread
          ADDR    STATE  FLG PFLG SFLG   PRI  EPRI PIL             INTR
ffffff0007c25c80 sleep       8    0    3    60     0   0              n/a

Visit the opensolaris website and browse through its source code, but most importantly have fun!

Java Generics - exploration ....

2008-07-13T12:50:00.012+08:00

Coding in Java generics (download this tutorial to gain an understanding) needs some time and experimentation in order to get it right and i discovered one of the pitfalls is how generic methods and type parameters are easily mixed up. Here's what i did and it took a java byte code viewer like jclasslib to understand what is happening.



----------- WeirdBox.java --------------------------

package generics.box;

public class WeirdBox<T, X, Y , E> extends PaperBox<T, X, Y > implements WeirdBoxProp {

public E weird;

WeirdBox(T id, X name, Y manu, E weird) {
    super.changeId(id);
    super.setManufacturer(manu);
    super.setName(name);
    this.weird = weird;
}

// isGlob(E t) & isTwistable(E t) is declared in the interface WeirdBoxProp
public <E extends Boolean> E isGlob(E t) {
// this.weird = t;
// return this.weird;
    return t;
}

public <E extends Boolean> E isTwistable(E t) {
// this.weird = t;
// return this.weird;
    return t;
}
}

----------- WeirdBoxProp.java --------------------------
package generics.box;

public interface WeirdBoxProp {
<E extends Boolean> E isGlob(E t); // suppose to be a glob so subclass must return true
<E extends Boolean> E isTwistable(E t); // suppose to be twistable so subclass must return true
}

The thing to realize from this experiment was that the type parameter declaration statement "public E weird" in WeirdBox.java is not related to that declared in the overridden methods isGlob(...) nor isTwistable(...). This is evident when you de-compile the java class(s) and here's what i saw

public E weird is translated to a object of type "java.lang.Object"
isGlob(...) is translated to a method of signature "<(Ljava/lang/Boolean;)Ljava/lang/Boolean;>" and similarly for isTwistable(...) and therefore it explains why i had to comment out the statement "this.weird = t; return this.weird;" because it wasn't type-compatible.

One of the ways i found (please drop me a email if you found another way) was to do the following changes in WeirdBox.java


package generics.box;

public class WeirdBox<T, X, Y, E1 extends Boolean> extends PaperBox<T, X, Y > implements WeirdBoxProp {

 public E1 weird;

 WeirdBox(T id, X name, Y manu, E1 weird) {
     super.changeId(id);
     super.setManufacturer(manu);
     super.setName(name);
     this.weird = weird;
 }

 public <E extends Boolean> E isGlob(E t) {
     this.weird = (E1)t;
     return (E)this.weird;
 }

 public <E extends Boolean> E isTwistable(E t) {
     this.weird = (E1)t;
     return (E)this.weird;
 }
}

You have to provide unique type parameters inside the code otherwise it won't work. You can see that i had to perform some type casts to get it going and if you compiled it using "-Xlint" you would find that the compiler has the following warning messages:

found : E
required: E1
this.weird = (E1)t;
found : E1
required: E
return (E)this.weird;

which is again weird because i assumed that E and E1 are now of type java.lang.Boolean but the warning of an unchecked cast exception makes me a little jittery but perhaps i used it in the wrong way.

--- An update on this post ---

I found an alternative site on Java Generics (here's the link to the PDF version) and its done by Angelika Langer

Detecting Endian-ness in Java

2008-07-08T09:50:00.001+08:00

Got a question on this from someone and the exact question was "how to check Endianness in Java coding"

The fact that Java is so prevalent in the industry and IT world is because you don't have to worry about such stuff. Java takes care of that for you in the Java Virtual Machine implementation and there isn't any API (correct me if i am wrong) that you can invoke like "isEndian()" , "BigIndian()" etc.

In general, the only area i can think at the moment is when you are running Java on 2 different Endian machines like Windows <--> Unix. And in these cases, Java takes care of the byte-order/endian conversion to/fro in the JVM implementation. One good resource i can think of is the OpenJDK project where you can browse through their source codes to discover the mechanism.

Hope this helps clear the air.

BTrace glitches ?

2008-07-08T08:54:00.003+08:00

Some glitches here on tracing Kind.NEW. Fyi, it's an feature that is used to track object creation.


// UDO == User Defined Object
class UDO {
UDO() {
    System.out.println("UDO");
}
}

public class ObjAlloc {
    Object id;

    ObjAlloc() {
     // Does not appear to work for Kind.NEW
     //id = new String("ME");
     //id = new String();
     id = new Object();
     //id = new UDO();

     // Works for Kind.NEWARRAY
     //id = new int[10];
    }

    private void A(String... str) {
              // Works for Kind.NEW & Kind.NEWARRAY
              //id = new int[10];
              //id = new Object();
    }

    public static void main(String[] args) throws Exception {
            System.out.println("Here");
            ObjAlloc oa = new ObjAlloc();
            while(true) {
              new ObjAlloc().A();
              Thread.currentThread().sleep(500);
              System.out.println(".");
            }
    }

}

Didn't work for Kind.NEW but worked for Kind.NEWARRAY and filed an issue with BTrace's website. Read about it by clicking here and i'll update you when i have updates. This is not bashing BTrace because i still think its a cool tool and the potential to challenge commercial implementations out there and like all software there will be glitches here and there, give them a break right ?

--- This is an update on this article ---

It turns out that this wasn't a glitch at all but due to my lack of understanding of the subject matter and i decided not remove this post so as a reminder. Having said that, i should proceed to show you how to do it correctly.

Below is the correct way to do it on a BTrace script


@BTrace
public class test {
@OnMethod(clazz="Main", method="/.+/",

location=@Location(value=Kind.NEW,clazz="/.+/"))
public static void p() {
 println("Object Created");
}
}


import com.sun.btrace.annotations.*;
import static com.sun.btrace.BTraceUtils.*;

@BTrace
public class test {
@OnMethod(clazz="Main", method="/.+/",

location=@Location(value=Kind.NEW,clazz="java.lang.Object"))
public static void p() {
 println("Object Created");
}
}

what i failed to do previously was to investigate the possibility of adding the "clazz" attributed to the "Location" which was why i wasn't able to instrument and detect my code from the perspective of object creation.

Using BTrace in Java 6

2008-06-29T22:10:00.003+08:00

Got up to speed with BTrace recently and thanks to Sundararajan's clarifications on syntax. This tool is simply TOO cool to give it up since in my opinion, someone whom has developed in Java and/or J2EE would quickly pick this technology up in no time :-)

However, beware of some caveats as i found out while attempting to monitor certain Java classes in Java 6 - coincidentally i attempt to monitor the EJB lifecycle etc.


btrace DEBUG: java.lang.RuntimeException: JSR/RET are not supported with computeFrames option
java.lang.RuntimeException: JSR/RET are not supported with computeFrames option
       at org.objectweb.asm.Frame.a(Unknown Source)
       at org.objectweb.asm.MethodWriter.visitJumpInsn(Unknown Source)
       at org.objectweb.asm.MethodAdapter.visitJumpInsn(Unknown Source)
       at org.objectweb.asm.ClassReader.accept(Unknown Source)
       at org.objectweb.asm.ClassReader.accept(Unknown Source)
       at com.sun.btrace.runtime.InstrumentUtils.accept(InstrumentUtils.java:66)
       at com.sun.btrace.runtime.InstrumentUtils.accept(InstrumentUtils.java:62)
       at com.sun.btrace.agent.Client.instrument(Client.java:261)
       at com.sun.btrace.agent.Client.transform(Client.java:101)
       at sun.instrument.TransformerManager.transform(TransformerManager.java:169)
       at sun.instrument.InstrumentationImpl.transform(InstrumentationImpl.java:365)
       at sun.instrument.InstrumentationImpl.retransformClasses0(Native Method)
       at sun.instrument.InstrumentationImpl.retransformClasses(InstrumentationImpl.java:124)
       at com.sun.btrace.agent.Main.handleNewClient(Main.java:278)
       at com.sun.btrace.agent.Main.startServer(Main.java:245)
       at com.sun.btrace.agent.Main.access$000(Main.java:53)
       at com.sun.btrace.agent.Main$1.run(Main.java:127)
       at java.lang.Thread.run(Thread.java:619)

If you attempted to instrument certain code in your EJBs like ejbXXX() methods you might or would find something similar to the above and that would certainly kill your interest but i would advise you not to be too hasty in this as it still works in the J2SE 1.5.x platform.

But this certainly sparked my interest in the algorithm that is causing this problem and i went browsing through the fundamental technology in which BTrace was built - ASM. According to the ASM 3.0 documentation & the Java Virtual Machine specifications 2nd Edition, it turns out ASM is having a hard time trying to run its algorithm "execute()" which essentially will simulate the java bytecode instruction on the output stack frame ; so the current version of this software will actually throw a Java unchecked exception i.e. java.lang.RuntimeException to inform the user of this tool (in this case, myself) that its not supported .... yet.

BTrace - BCI "cool" tool

2008-06-22T14:10:00.002+08:00

Recently, i got to know about BTrace which is really another cool technology by A. Sundararajan on monitoring JVM behavior and it does this via instrumenting the Java bytecode.

For an actual & real world demonstration, you can read a good article by Binod which illustrates his adventure in using this technology to solve a memory leak problem on the sailfin which is in a nutshell a SIP container based largely on the JSR-289.

Detecting deadlocks in JVM using DTrace

2008-06-15T14:32:00.001+08:00

In this article, i am going to demonstrate using DTrace to detect deadlocks in the application code. Before i do that, here's how you can do it using jconsole.

Simply start up jconsole and connect to the JVM either locally or remotely using the jconsole GUI and once you are able to log into the java virtual machine, navigate to the "Thread Tab" and look for the "Detect Deadlock" and activate it by clicking it. If your code has a deadlock situation, jconsole would pick it up. See the screen shots below.

From the screen shots, you can see that deadlock detection is pretty easy. You can achieve this using DTrace as well. Here's what i did with a sample d script and the output and of course it needs further refining (JDK SE 6 currently supports limited functionality on this).


/*
Detect deadlocks in the JVM

@author: Raymond Tay
@date: 15 June 2008
@version: 1.0

*/
profile-100msec
{
}

monitor-contended-enter
{
this->threadid = arg0;
this->monitorid = arg1;
this->monitorclass = (string)copyin(arg2, arg3+1);
printf("Attempt to enter! tid=%d, monitor id=%d, monitor class=%s\n", this->threadid, this->monitorid, this->monitorclass);
}

monitor-contended-entered
{
this->threadid = arg0;
this->monitorid = arg1;
this->monitorclass = (string)copyin(arg2, arg3+1);
printf("Entered ! tid=%d, monitor id=%d, monitor class=%s\n", this->threadid, this->monitorid, this->monitorclass);
}

And the output from a typical run


bash-3.2# dtrace -q -s ./detectdl.d -c "java -XX:+ExtendedDTraceProbes -jar /export/home/tayboonl/NetBeansProjects/DeadlockApp/dist/DeadlockApp.jar "
MyThread1 got mon1
MyThread1 trying to get monitor2...
MyThread2 got mon2
MyThread2 trying to get monitor1...
Attempt to enter! tid=8, monitor id=135675668, monitor class=sun/misc/Launcher$AppClassLoader
Entered ! tid=8, monitor id=135675668, monitor class=sun/misc/Launcher$AppClassLoader
Attempt to enter! tid=8, monitor id=135675868, monitor class=java/lang/Object
Attempt to enter! tid=9, monitor id=135675968, monitor class=java/lang/Object

From my code below, i know that my threads are contending for a monitor of type "java.lang.Object" and the two threads contending for them are tid=8 and tid=9 but of course the information offered by jconsole is more detailed than what i offer here and its partly due to the support provided by SE6 at the time of this writing. Probably will refine it later to include the same kind of information as provided by jconsole.

Last thing to note about the output above is that you need to input the JVM parameter "-XX:+ExtendedDTraceProbes" to the JVM inorder to enable the probes firing where by default, its turned off since there's a performance penalty. As far as the documentation goes, the JDK team is currently fixing it and probably will enable it by default in the next major release of the JDK.

Here's the sample java code


/*
 * To change this template, choose Tools | Templates
 * and open the template in the editor.
 */
package deadlockapp;

/**
 *
 * @author tayboonl
 */
public class Main {

    Object monitor1 = new Object();
    Object monitor2 = new Object();

    class MyThread1 extends Thread {

        public void run() {
            synchronized (monitor1) {
                System.out.println("MyThread1 got mon1");
                System.out.println("MyThread1 trying to get monitor2...");
                synchronized (monitor2) {
                    System.out.println("MyThread1 got mon2");
                }
            }
        }
    }

    class MyThread2 extends Thread {

        public void run() {
            synchronized (monitor2) {
                System.out.println("MyThread2 got mon2");
                System.out.println("MyThread2 trying to get monitor1...");
                synchronized (monitor1) {
                    System.out.println("MyThread2 got mon1");
                }
            }
        }
    }

    /**
     * @param args the command line arguments
     */
    public static void main(String[] args) {
        Main m = new Main();
        m.new MyThread1().start();
        m.new MyThread2().start();
    // TODO code application logic here
    }
}

Have fun!

Monitoring JVM (JavaSE6) activity using DTrace

2008-06-15T12:45:00.003+08:00

DTrace has been around for what ... 3 years or so i think and there are plenty of articles whereby bloggers like myself used DTrace in one way or another to monitor the internals of the JVM and possibly use JVM-DTrace to solve some real-world problems. So in this article, i like to bring to your attention the fact that DTrace probes are shipped with the latest version of the Java Development Kit version 6.x

Here's a simple D script (that is written in the D programming language) which monitors the JVM activity and in particular garbage collection.


/*

Tracks the JVM via the "hotspot" provider
@author: Raymond Tay
@date:  15 June 2008
@version: 1.0

*/

struct mem_pool {
string name;
long long init_size;
long long mem_inuse;
long long max_size;
long long gc_start_time;
};

/*
Declare pool as a array of JVM memory pools
*/
struct mem_pool pool[string];
this string poolname;

BEGIN
{
printf("Tracing has begun\n");
}
mem-pool-gc-begin
/execname == "java" && pid == $1/
{
this->poolname = (string)copyin(arg2, arg3+1);
pool[this->poolname].name = this->poolname;
pool[this->poolname].init_size = arg4;
pool[this->poolname].mem_inuse = arg5;
pool[this->poolname].max_size = arg7;
pool[this->poolname].gc_start_time = vtimestamp;
printf("\tGC about to begin for name=%s->(init=%d,inuse=%d,max=%d)\n",this->poolname,arg4,arg5,arg7);
}
mem-pool-gc-end
/execname == "java" && pid == $1/
{
this->poolname = copyin(arg2, arg3+1);
this->FreedMem = pool[this->poolname].mem_inuse - arg5;
this->FreeMem = pool[this->poolname].max_size - arg5;
printf("\tGC ended for name=%s->(init=%d,inuse=%d,freed=%d,max=%d)\n",this->poolname,arg4,arg5,this->FreedMem,arg7);
pool[this-poolname].mem_inuse = arg5;
}

END
{
printf("Tracing has ended, thanks\n");
}

If you were to run this on your Solaris 10, you would find the output as below


...
    GC about to begin for name=Code Cache->(init=163840,inuse=4853312,max=33554432)
    GC about to begin for name=Eden Space->(init=16580608,inuse=16646144,max=33030144)
    GC about to begin for name=Survivor Space->(init=2031616,inuse=307648,max=4128768)
    GC about to begin for name=Tenured Gen->(init=247791616,inuse=99175912,max=495583232)
    GC about to begin for name=Perm Gen->(init=50331648,inuse=93384752,max=134217728)
    GC about to begin for name=Perm Gen [shared-ro]->(init=8388608,inuse=6633584,max=8388608)
    GC about to begin for name=Perm Gen [shared-rw]->(init=12582912,inuse=7532648,max=12582912)
    GC ended for name=Code Cache->(init=163840,inuse=4853312,freed=0,max=33554432)
    GC ended for name=Eden Space->(init=16580608,inuse=0,freed=16646144,max=33030144)
    GC ended for name=Survivor Space->(init=2031616,inuse=240352,freed=67296,max=4128768)
    GC ended for name=Tenured Gen->(init=247791616,inuse=99175912,freed=0,max=495583232)
    GC ended for name=Perm Gen->(init=50331648,inuse=93384752,freed=0,max=134217728)
    GC ended for name=Perm Gen [shared-ro]->(init=8388608,inuse=6633584,freed=0,max=8388608)
    GC ended for name=Perm Gen [shared-rw]->(init=12582912,inuse=7532648,freed=0,max=12582912)
...

Obviously, you can customize the script to include calculations for gc times, number of gc invocations, etc and all this made possible by the HotSpot Dtrace Providers in the latest JDK6. Below is the sample script which would have the capability to detect thread starts or stops, java classes loading and unloading, gc times etc


/*

Tracks the JVM via the "hotspot" provider
@author: Raymond Tay
@date:  15 June 2008
@version: 1.0

*/

struct mem_pool {
string name;
long long init_size;
long long mem_inuse;
long long max_size;
long long gc_start_time;
};

/*
Declare pool as a array of JVM memory pools
*/
struct mem_pool pool[string];
this string poolname;

BEGIN
{
printf("Tracing has begun\n");
}
/* ---------------------- */
vm-init-begin
/execname == "java" && pid == $1/
{
printf("JVM is initializing ...\n");
}
vm-init-end
/execname == "java" && pid == $1/
{
printf("JVM is initialized completely and going to run app. code\n");
}

vm-shutdown
/execname == "java" && pid == $1/
{
printf("JVM is shutting down now\n");
}

thread-start
/execname == "java" && pid == $1/
{
printf("Thread=(%s,%d),OS tid=(%d) (daemon:Y/N? %s) is starting\n", stringof(copyin(arg0,arg1+1)), arg2, arg3, arg4 == 0?"N":"Y");
}

thread-stop
/execname == "java" && pid == $1/
{
printf("Thread=(%s,%d),OS tid=(%d) (daemon:Y/N? %s) is stopping\n", stringof(copyin(arg0,arg1+1)), arg2, arg3, arg4 == 0?"N":"Y");
}

class-loaded
/execname == "java" && pid == $1/
{
printf("\t\t%s loaded\n", stringof(copyin(arg0,arg1+1)));
}
class-unloaded
/execname == "java" && pid == $1/
{
printf("\t\t%s un-loaded\n", stringof(copyin(arg0,arg1+1)));
}

gc-begin
/execname == "java" && pid == $1/
{
printf("\t%s about is starting..\n", arg0 == 0? "Full-GC":"GC");
}

gc-end
/execname == "java" && pid == $1/
{
printf("\tGC has completed\n");
}

mem-pool-gc-begin
/execname == "java" && pid == $1/
{
this->poolname = (string)copyin(arg2, arg3+1);
pool[this->poolname].name = this->poolname;
pool[this->poolname].init_size = arg4;
pool[this->poolname].mem_inuse = arg5;
pool[this->poolname].max_size = arg7;
pool[this->poolname].gc_start_time = vtimestamp;
printf("\tGC about to begin for name=%s->(init=%d,inuse=%d,max=%d)\n",this->poolname,arg4,arg5,arg7);
}
mem-pool-gc-end
/execname == "java" && pid == $1/
{
this->poolname = copyin(arg2, arg3+1);
this->FreedMem = pool[this->poolname].mem_inuse - arg5;
this->FreeMem = pool[this->poolname].max_size - arg5;
printf("\tGC ended for name=%s->(init=%d,inuse=%d,freed=%d,max=%d) took %d msec.\n",this->poolname,arg4,arg5,this->FreedMem,arg7, (vtimestamp - pool[this->poolname].gc_start_time)/1000000);
pool[this-poolname].mem_inuse = arg5;
}

/* ---------------------- */

END
{
printf("Tracing has ended, thanks\n");
}

Depending on your preference, you might find the information in jconsole equally appealing. Mandy Chung from Sun Microsystems Inc. wrote a good article on this some years back and you can read it by clicking it here.

Last thing here is that, if you ran your d script on a jvm and closely monitor the garbage collection data, you should notice something peculiar on the portion related to "Tenured Generation" where you will find that with every gc, the size increases and this is due to the object generational promotion. For more details on this, you can click on this article: Java SE6 HotSpot VM Tuning.

Oracle 10g installation on Solaris 10 x86-64 is annoying ...

2008-06-03T22:32:00.001+08:00

This is irritating as far as i can tell. I was trying to install Oracle 10g on Solaris 10 x86 and according to the Oracle website, it's supported and so happily downloaded everything and triggered the installer....so far so good but hit into a wall with the following error message:


bash-3.2$ ./runInstaller
Starting Oracle Universal Installer...

Checking installer requirements...

Checking operating system version: must be 5.10.    Actual 5.11
                                    Failed <<<<

Now that's obviously annoying since my OS is a solaris 10 so what's going on ? Well it turns out that this irritating thing is widely known and you can search for it on Google and see tons of complaints about this ...

What happens is that the Oracle 10 g installer invokes a particular program and its a hidden file (name: .oui) and i guess the name is Oracle Universal Installer and hence the name oui and so happens the program invokes "uname -r" and on my system returns 5.11 !

For some reason unknown to me, my native version of uname contains the hard-coded string 5.11 which of course is used to determine the platform on which i am using. Btw, use the Solaris/Unix utility strings to discover the embedded strings in the program and this works in general for compiled programs. But there's a caveat and that is you will see that the same hard-coded string is found in /lib/libc.so.1 yes that's right a shared library.... and that's the end of the road for me.

My advise to you is not to touch it the shared file as it might affect your other applications running on your machine ... unless you know EXACTLY 100% what you are doing, you might as well wait for a patch or the next release and perhaps Sun might fix this

Solaris utility "cputrack" glitches

2008-05-26T19:03:00.005+08:00

In my previous article whereby i demonstrated that you can use "cputrack" to monitor CPU hardware counters but it turns out that there was a glitch and subsequently i filed a bug report with Sun and they assigned an id for it: 6705264

IT salary survey 2008

2008-05-24T13:00:00.009+08:00

Found an interesting comparison survey conducted by ZDNet Asia on the latest IT salaries in the island state of Singapore. Here are the snapshots from the report. You can draw your own conclusions :-)