C# language, C++, CodeProject, GPGPU, Machine Learning, Technical

Large Scale Machine Learning using NVIDIA CUDA


You may have heard about the Stanford University’s machine learning on-line course given by Prof. Andrew Ng. in 2011; it was a great course with lots of real-world examples. During the course, I’ve realized that GPUs are the perfect solution for large-scale machine learning problems. In fact, there are many examples of supervised and unsupervised learning all around the internet. Being a fan of both GPGPU and Machine Learning technologies, I came up with my own perspective to run machine learning algorithms with a huge amount of data on the GPUs.

I’ve already presented the solution recently at the South Florida Code Camp 2012. Everybody was interested in these two subjects a lot; therefore I’ve decided to share it on my blog. The example in this post is neither the only solution nor the best solution. I hope it will help you one day solve your own machine learning problem.

There is a lot of concepts to machine learning but in this post, I’m only scratching the surface. If you already know about GPGPU and Machine Learning you can just go to the source code at this link, download the Visual Studio 2010 projects and try it out.

I’ve also prepared the same example using CUBLAS with vectorized implementation of the polynomial regression algorithm, but the CUBLAS example would require more in-depth explanations. Therefore I’m posting this example first which is a simplified implementation. If you are interested in CUBLAS implementation please let me know and I can send you that copy.




Machine Learning

If you are already familiar with machine learning you can skip the brief introduction and jump directly to the Large Scale Machine Learning section. Or if you want to learn more about machine learning please follow the links or check out the Stanford course I’ve mentioned at the beginning.

Machine learning algorithms allow computers recognize complex patterns. It focuses on the prediction, based on known properties learned from the training data. We are using machine learning algorithms every day dozens of times maybe unknowingly: every time we get a book or movie recommendation or every time we do a web search. In 1959 Arthur Samuel described Machine learning as Field of study that gives computers the ability to learn without being explicitly programmed. It has been a while machine learning was first introduced, and it is gaining popularity again with the rise of Big Data.

Figure 1 shows how some of the machine learning processes work. On phase 1, given a data set a machine learning algorithm can recognize complex patterns and come up with a model. In most cases, this phase is the bulk of the computation. In the second phase, any given data can run through the model to make a prediction. For example if you have a data set of house prices by size, you could let the machine learn from the data set and let it predict house price of any given size.

Figure 1

It does this by recognizing the function which defines the relation between the different features of the problem. A linear problem with two dimensions, like house price (the house size is the feature and the house price is the label data), can be expressed with the f(x) = ax + b model. Figure 2 shows how one feature can be used on a linear regression problem to predict new house prices. The term “hypothesis” was used in the Stanford course to describe the model.

Figure 2

Depending to the data set, more complex functions can be used. On Figure 3 you can see how the complexity can grow easily from 2 dimensions linear to hundreds of dimensions polynomial. In a spam filtering problem the different features could be words in the email or in a face recognition problem the features could be the pixels of the image. In the house price prediction example, features are properties of the house which are affecting the price. e.g. size, room count, floors, neighborhood, crime rate etc.

Figure 3

There are many machine learning algorithms for different problem types. The most common groups of these algorithms are Supervised Learning, Unsupervised Learning. Supervised learning is used on problems where we can provide the output values for the learning algorithm. For example: house prices for some house features is the output value, therefore house price prediction is a supervised learning problem. Data with these output values is named as “labeled data”. On the other hand unsupervised learning does not require output values, patterns or hidden structures can be recognized just with feature data. For example: clustering social data to determine groups of people by interest would not require to define any output value, therefore it is an unsupervised learning problem.

Gradient Descent

In supervised learning problems, the machine can learn the model and come up with a hypothesis by running a hypothesis with different variables and testing if the result is close to the provided labels (calculating the error). Figure 4 shows how a training data is plot and the error is calculated. An optimization algorithm named Gradient Descent (Figure 5) can be used to find the optimum hypothesis. In this simple two dimensional problem, the algorithm would run for every different value of “a” and “b”, and would try to find the minimum total error.

Figure 4

The pseudo code below shows how the gradient descent algorithm in Figure 5 works :

for every a and b loop until converge
errors = 0
for i = 1 to data.length
    fx = a * data[i] + b
    errors += (fx - labelData[i]) * data[i]
end for
gradient = gradient - learningRate * 1/data.length * errors  
end for

Figure 5 (Fromml-class.org)

Large Scale Machine Learning

Machine learning problems become computationally expensive when the complexity (dimensions and polynomial degree) increases and/or when the amount of data increases. Especially on big data sources with hundreds of millions of samples, the time to run optimization algorithms increases dramaticaly. That’s why we are looking for parallelization opportunities in the algorithms. The error summation of gradient descent algorithm is a perfect candidate for parallelization. We could split the data into multiple parts and run gradient descent on these parts in parallel. In Figure 6 you can see how the data is split into four parts and fed into four different processors. On the next step the result is gathered together to run the rest of the algorithm.

Figure 6

Clearly, this approach would speed up the machine learning computation by almost four times. But what if we would have more cores and split the data further? That is where GPUs step into the solution. With GPUs we can parallelize in two layers: multiple GPUs and multiple cores in every GPU. Assuming a configuration with 4 GPUs and 512 cores each, we could split down the data into 512 more pieces. Figure 7 shows this configuration along with the parallelized part on the GPU cores.

Figure 7



Utilizing GPUs to enable dramatic increases in computing performance of general purpose scientific and engineering computing is named GPGPU. NVIDIA is providing a parallel computing platform and programming model named CUDA to develop GPGPU software on C, C++ or Fortran which can run on any NVIDIA GPU. NVIDIA CUDA comes with many high level APIs and libraries like basic linear algebra, FFT, imaging etc. to allow you concentrate on the business logic rather than re-writing well known algorithms.

You can visit my previous blog posts where I’ve explained how to use NVIDIA CUDA capable GPUs to perform massively parallel computations. The examples include Monte Carlo simulation, random number generators and sorting algorithms.


House Price Prediction Example

On this post I’ll show you how you can implement house price prediction on NVIDIA CUDA. Given a house price data set based on bedrooms, square feet and year built, it is possible to let the machine learn from this data set and provide us with a model for future predictions. Because the error calculation part of the Gradient Descent algorithm is highly parallelizable, we can offload it to the GPUs.

The machine learning algorithm in this example is Polynomial Regression, a form of the well known Linear Regression algorithm. In Polynomial Regression the model is fit on a high order polynomial function. In our case we will be using bedrooms, square feet, year built, square root of bedrooms, square root of square feet, square root of year built and the product of bedrooms and square feet. The reason we added the four polynomial terms to the function is because of the nature of our data. Fitting the curve correctly is the main idea behind building a model for our machine leaning problem. Logically house prices increase by these features not in a linear or exponential way and they don’t drop after a certain peek. Therefore the graph is more like a square root function, where house prices increase less and less compared to increasing any other feature.

Finding the right polynomial terms is very important for the success of the machine learning algorithm: having a very complex, tightly fitting function would generate too specific model and end up with overfitting, having a very simple function, like a straight line would generate too general model and end up with under fitting. Therefore we are using additional methods like adding regularization terms to provide a better fit to your data. Figure 8 shows the gradient descent algorithm including with the regularization term lambda.

Figure 8 (Fromml-class.org)

Application Architecture

The sample application consist of a C++ native DLL named LR_GPULib, for the machine learning implementation on the GPU and a C# Windows application named TestLRApp for the user interface. The DLL implements Data Normalization and Polynomial Regression using the high level parallel algorithm library Thrust on NVIDIA CUDA. I’ve mentioned on my previous blog posts about Thrust more in detail, therefore I’m not going into much detail on this post. Figure 9 shows the application architecture and also the program flow from loading the training data all the way down to making a prediction.

Figure 9

The application provides the UI shown below on Figure 10 to load the data, train and make a prediction with new data set. The UI also shows the hypothesis on the bottom of the dialog with all constants and features.

Figure 10


The LR_GPU_Functors.cu file in the DLL contains the functors used as kernels on Thrust methods. The LR_GPU.cu file in the DLL contains the normalization, learning and prediction methods. The Learn method accepts the training data and the label data, which are all the features and all prices in two float arrays. The first thing the Learn method does is to allocate memory, add bias term and normalize the features. The reason we added the bias term is to simplify the gradient loop and the reason we normalize the features is because the data ranges are too different. E.g. square feet is four digits and bedrooms is single digit. By normalizing the features we bring them into the same range, between zero and one. Normalization gets also executed on the GPU using the NormalizeFeatures. But the normalization requires the mean and standard deviation (std), therefore mean and std are calculated first and provided to the NormalizeFeaturesByMeanAndStd method to calculate the mean normalization.

void NormalizeFeaturesByMeanAndStd(unsigned int trainingDataCount, float * d_trainingData, 
thrust::device_vector<float> dv_mean, thrust::device_vector<float> dv_std)
	//Calculate mean norm: (x - mean) / std
	unsigned int featureCount = dv_mean.size();
	float * dvp_Mean = thrust::raw_pointer_cast( &dv_mean[0] );
	float * dvp_Std = thrust::raw_pointer_cast( &dv_std[0] );
	FeatureNormalizationgFunctor featureNormalizationgFunctor(dvp_Mean, dvp_Std, featureCount); 
		thrust::device_ptr<float> dvp_trainingData(d_trainingData); 
	thrust::transform(thrust::counting_iterator<int>(0), thrust::counting_iterator<int>
		(trainingDataCount * featureCount), dvp_trainingData, dvp_trainingData, featureNormalizationgFunctor);

The Normalization code running on the GPU is implemented in the FeatureNormalizationgFunctor functor, which is simply calculating data - mean / std in parallel for every element of the data, as seen below:

  __host__ __device__
  float operator()(int tid, float trainingData)
	  int columnIdx = tid % featureCount;
	  float fnorm = trainingData - meanValue[columnIdx];
	  if (stdValue[columnIdx] > 0.0)
		fnorm /= stdValue[columnIdx];
	  return fnorm;

On the next step in the Learn method, the gradient descent is calculated with the for(int i = 0; i < gdIterationCount; i++) loop. As I mentioned before, the error calculation part of the gradient descent is executed in parallel but the rest is calculated sequentialy. The thrust::transform is used with the TrainFunctor to calculate f(x)-y in parallel for every sample. f(x) is simply the A*x1 + Bx2 + Cx3 + Dx4 + Ex5 + Fx6 + Gx7 + H hypothesis where x1 through x7 are the features (x1=bedrooms, x2=square feet, x3=year built, x4=square root of bedrooms, x5=square root of square feet, x6=square root of year built and x7=the product of bedrooms and square feet) and A through H are the constants which gradient descent will find out. This is shown with the Green Square on Figure 11. The TrainFunctor code snippet and the usage code snippet are shown below:

Figure 11 (Fromml-class.org)

__host__ __device__
  float operator()(int tid, float labelData)
		float h = 0;
		for (int f = 0; f < featureCount; f++)
			h += hypothesis[f] * trainingData[tid * featureCount + f];
		return h - labelData;
	dv_labelData.begin(), dv_costData.begin(), tf);

The thrust::transform_reduce is used with the TrainFunctor2 to apply the features to the error result and sum up all of them. This is shown with the code snippet below and the Red Square on Figure 11. Rest of the Learn method calculates gradient descent part marked with Blue Square on Figure 11.

float totalCost = thrust::transform_reduce(thrust::counting_iterator<int>(0), 
	thrust::counting_iterator<int>(trainingDataCount),  tf2, 0.0f, thrust::plus<float>());

Once gradient descent converges, the constants A though H of the hypothesis is returned back to the TestLRApp with the result array.

As you may guess the prediction works by using the constants with new sample data on the hypothesis. This is done using the Predict method in the LR_GPULib library. As seen below the Predict method normalizes the given features set and calculates the hypothesis using the constants and the normalized data with the help of the PredictFunctor. The result is a the predicted house price for the given features.

	NormalizeFeaturesByMeanAndStd(testDataCount, pdv_testData, dv_mean, dv_std);

	PredictFunctor predictFunctor(pdv_testData, pdv_hypothesis, featureCount);
		thrust::counting_iterator(testDataCount), dv_result.begin(), predictFunctor);
struct PredictFunctor : public thrust::unary_function
	float * testData;
	float * hypothesis;
	unsigned int featureCount;

	PredictFunctor(float * _testData, float * _hypothesis, unsigned int _featureCount) 
		: testData(_testData), hypothesis(_hypothesis), featureCount(_featureCount)

  __host__ __device__
  float operator()(int tid)
	  float sum = 0;
	  for(int i = 0; i < featureCount; i++)
		sum += testData[tid * featureCount + i] * hypothesis[i];
	  return sum;


GPGPU, Machine Learning and Big Data are three rising fields in the IT industry. There is so much more about these fields than what I’m providing in this post. As much as I get deeper into these fields I figure out how well they fit together. I hope this sample gave you some basic idea and maybe just one perspective how you can use NVIDIA CUDA easily on machine learning problems. As in any other software solution this example is not the only way to do polynomial regression on house price prediction with GPUs. In fact an enhancement would be supporting multiple GPUs and splitting down the data set into more parts.

AWS, C++, GPGPU, Technical

How to set up Amazon EC2 Windows GPU instance for NVIDIA CUDA development


Amazon Elastic Compute Cloud web service provides a very useful platform on the cloud. Especially for software developers who don’t have access to expensive hardware. Some time ago as I was looking for a better CUDA enabled GPU solution than my Mac Book Pro, I’ve realized that it is time to switch from laptop to a desktop. But luckily, Amazon introduced couple months ago the GPU instances, running on Windows Server 2008 OS. I’ve been using the scalable and cost efficient Amazon EC2’s since couple years without any problem and now that they are providing a platform with two Tesla M2050s to test my CUDA apps, I just want to say Thank You Amazon.

On this post I want to share with you my experience how to set up a full NVIDIA CUDA development environment on a Windows EC2 GPU instance. And I’ll also walk you through couple CUDA examples.

If you were following my previous blog posts and were not able to try them out because of not having a CUDA capable hardware, you will have a chance to do it after reading this blog.

One of the reasons I’m providing this blog post is also to use this information in our HPC & GPU Supercomputing group of South Florida hands-on lab meetups. If you are from the group, you’ve most probably received already the AMI. Therefore you can skip the set up part.


About Amazon EC2 GPU Instances

Amazon Elastic Compute Cloud (Amazon EC2) is a web service that provides resizable compute capacity in the cloud. It is designed to make web-scale computing easier for developers.

The GPU instances provide general-purpose graphics processing units (GPUs) with proportionally high CPU and increased network performance for applications benefitting from highly parallelized processing, including HPC, rendering and media processing applications. The Windows GPU instance is named Cluster GPU Quadruple Extra Large instance and has

22 GB memory, 33.5 EC2 Compute Units, 2 x NVIDIA Tesla “Fermi” M2050 GPUs, 1690 GB of local instance storage, 64-bit platform, 10 Gigabit Ethernet.


Utilizing GPUs to do general purpose scientific and engineering computing is named GPGPU. You can visit my previous blog posts where I’ve explained how to use NVIDIA CUDA capable GPUs to perform massively parallel computations.


Browse to http://aws.amazon.com/ec2/ and click the link on the top of the page saying “Sign in to the AWS Management Console”.

Please be aware that you will get charged by Amazon for the usage of their services. Therefore, check for running objects before leaving the AWS Management console. Please check out the Amazon pricing web page for more information.

The next couple paragraphs explain how to create your AWS account and set up your environment. You can skip this section if you already have and account and familiar with Amazon EC2.

Registering for Amazon AWS

If you already have an Amazon account you can use it to log in, otherwise you can create a new account from the same screen. Once you logged into the AWS console, it may ask you to sign up for a Amazon S3 account. In that case just follow the links to finish the sign up. Once it is done, you should receive the confirmation email. Now, login to you account to finish the registration and to go through a phone verification.

Setting up your AWS environment

Login to your Amazon AWS account, the AWS management console will show up. Select the EC2 tab from the top to see your EC2 dashboard. We will create a security group and a key pair for later use.

First, click the Key Pairs link on the right and after that click the Create Key Pair button. Enter a name for your private key file, like My_KeyPair and then after save the .pem file somewhere to use it later. You will also see the new key pair on the screen.

Go back to the EC2 dashboard and click the Security Group link on the right. This will open the security group console. Click the Create Security Group button and create a group named GPGPU_SecurityGroup. Select the Inbound tab for the new group and the rule editor will open. Add an RDP group by selecting RDP from the rules drop down and clicking the Add Rule button. Now click the Apply Rule Changes button to save the changes.

Creating the GPU EC2 Instance

  1. Go to the EC2 dashboard and click the Launch Instance button.
  2. Select the Launch Classic Wizard and click Continue.
  3. Find the Microsoft Windows 2008 R2 64-bit for Cluster Instances (AMI Id: ami-c7d81aae) in the list and click the Select button right next to it.
  4. Select Cluster GPU(cg1.4xlarge, 22GB) from the Instance Type drop down and click continue. If you have other instances and you are planning to transfer data between your instances, I’m suggesting selecting the same region for all of them to prevent in cloud data transfer charges.
  5. Select Continue on the Advanced Instance Options page.
  6. Give a name to your instance. e.g. GPGPU.
  7. Select the Key Pair you have created and click the continue button.
  8. Select the Security Group you have created and click the continue button.
  9. Click the launch button to finish the wizard.

Running the GPU EC2 Instance

You can click the instances link on the left hand Navigation menu to see the instance you’ve just created. The instance will be in pending state for a while until it will boot up completely.
Right click on the newly created instance and select Get Windows Password. You may have to come back after couple minutes if the password generation is pending.
Paste the content of the .pem file you’ve received while creating the key pair, to the Private Key field on the password retrieval dialog and click the Decrypt Password button.
Copy the Decrypted Password to use it later to log into the instance.

Connecting to the Instance using RDP

In order to connect to the newly created instance :

  1. Right click on it and select Connect.
  2. Click “Download shortcut file” link and save the RDP shortcut to your local machine.
  3. Open the saved RDP shortcut and logon to the instance by enter the retrieved password.
  4. Change your random generated password from the Control Panel / User Accounts section.

Installing GPGPU Developer Tools

Go to the CUDA Downloads website to see available downloads. At this time we will download the 4.1 RC2 version from CUDA Toolkit 4.1 web site.
Download and install the following items in the same order :

  1. Visual Studio C++ 2010 Express.
  2. CUDA Toolkit.
  3. GPU Computing SDK.
  4. Developer Drivers for WinVista and Win7 (285.86). The default drivers coming
  5. (Optional) Parallel Nsight 2.1RC2. In order to download this you have to sign up for the Parallel Nsight Registered Developer Program.

Backup the GPU EC2 instance

You will get charged for any instance which is not terminated, even for those in stopped state. Therefore, it is a good practice to backup to S3 and terminate your instance once you are done with testing to prevent any charges in downtime. You can do this in two ways: you can detach the EBS volume (storage) and terminate the instance or you can take a snapshot and delete the instance and volume. As of today the EBS volume costs $0.10 per GB-month and the snapshot costs $0.14 per GB-month. You can visit the Amazon EC2 pricing web site for a more up to date pricing.

Please follow the steps below for a snapshot backup:

  1. Click the volumes link on the navigation bar on the left hand side. You will see the volume ( storage ) attached to your EC2 instance.
  2. Right click on the volume and select Create Snapshot.
  3. Provide a name for the new snapshot and click the Yes, Create button.
  4. Go to the Snapshots section from the navigation menu and click refresh. You should see the new snapshot in pending mode. It will take a while to create the snapshot.

Running CUDA Samples

Now you are ready to compile and run a CUDA sample from the GPU Computing SDK. Please follow these steps :

  1. Login to the instance using the RDP shortcut.
  2. The samples require cutil32d.lib in order to function, therefore you need to compile the cutil project first. For that browse to the C:\ProgramData\NVIDIA Corporation\NVIDIA GPU Computing SDK 4.1\C\common folder and open the cutil_vs2010.sln visual studio solution file. Compile the solution.
  3. It is convenient to have syntax highlighting on .cu files. Therefore go to C:\ProgramData\NVIDIA Corporation\NVIDIA GPU Computing SDK 4.1\C\doc\syntax_highlighting\visual_studio_8 folder and follow the instructions in the readme.txt file.
  4. Our first example is the deviceQuery, which shows the properties of your GPU. Browse to the C:\ProgramData\NVIDIA Corporation\NVIDIA GPU Computing SDK 4.1\C\src\deviceQuery folder and open the deviceQuery_vs2010.sln. Compile the solution.
  5. The output executable will be placed into the C:\ProgramData\NVIDIA Corporation\NVIDIA GPU Computing SDK 4.1\C\bin\win32\Debug folder. Open a administrative command prompt and run the deviceQuery.exe.
  6. You should see two Tesla M2050 devices each with device capability 2.0, 448 CUDA cores, 3GB memory, 515 GFlops, 148 GB/sec memory bandwidth. This feels like 400hp under the hood!

Let’s run one more sample to see the performance difference of our GPUs. The sample we are going to run is matrixMul, located under the same C:\ProgramData\NVIDIA Corporation\NVIDIA GPU Computing SDK 4.1\C\src root folder. On Tesla M2050 this sample will multiply a 640 x 640 matrix with a 640 x 960 matrix to generate a 640 x 960 matrix.

Open the solution, go to the project properties and add the C:\ProgramData\NVIDIA Corporation\NVIDIA GPU Computing SDK 4.1\shared\inc path to the Include Directories under the VC++ Directories configuration properties. ( I’ve noticed that the path can not be found. )

Compile and run the project in a command window. You should see 0.001 sec for CUBLAS kernel execution and 0.021 sec for CUDA execution. CUBLAS is CUDA’s Basic Linear Algebra Library with optimized algorithm.

Let’s compare the GPU with the Intel Xeon 2.93Ghz CPU of the current instance. In order to do this we need to modify the code a little :

  1. Open the matrixMul.cu file.
  2. Add the #include <time.h> at line 41, under the kernel include.
  3. Find the line with computeGold(reference, h_A, h_B, uiHA, uiWA, uiWB); ( around line 417) and replace it with the following code.

    clock_t startTime,endTime;
    startTime = clock() * CLK_TCK;
    computeGold(reference, h_A, h_B, uiHA, uiWA, uiWB);
    endTime = clock() * CLK_TCK;
    shrLogEx(LOGBOTH | MASTER, 0, "> Host matrixMul Time = %.5f s\n", 
    				(double)(endTime - startTime) / 1000000.0 );
  4. Compile the code and execute it. You should see something around 3.463 sec. This means that the CUBLAS GPU version is about 3500x faster than the single core CPU version. A fair comparison with all cores utilized can be found on the CUBLAS web site, which is about 6-17x.


GPGPU is rising since the last couple years and now that Amazon provides a Windows GPU instance, it is much easier to jump onto the massively parallel software track as a Windows developer.

C++, CodeProject, GPGPU, Technical

Massively Parallel Monte Carlo Simulation using GPU


In my previous blog posts I’ve explained how you can utilize the GPU on your computer to perform massively parallel computation with the help of NVIDIA CUDA and Thrust technologies. On this blog post I’m diving deeper into Thrust usage scenarios with a simple implementation of Monte Carlo simulation.

My influence was the PI prediction sample on Thrust web site. The sample is running Monte Carlo simulation with 10K samples on a unit circle to estimate the PI number. You can visit this Wikipedia page if you are interested into how Monte Carlo Simulation can be used to approximate PI. Actually it is a solution to the famous Buffon’s Needle problem.

I’m taking the original example one step further to show you how to send device variables to functors in Thrust methods, and also using a slightly different problem. Perhaps there are many other methods to do the same logic, but on this blog post I’m just concentrating on this specific implementation.


About Monte Carlo Simulation

Monte Carlo simulation is an approach to solve deterministic problems with probabilistic analog. That is exactly what we are accomplishing in our example: estimating the area of intersecting disks. Monte Carlo methods are especially useful for simulating systems with many coupled degrees of freedom, such as fluids, disordered materials, strongly coupled solids, and cellular structures.

Our simulation is about predicting the intersect area of four overlapping unit disks as seen on the image below. (intersection of A,B,C and D disks) Actually, the problem can also be solved easily with the help of Geometry as explained here. I’ve calculated the area as 0.31515. On the other hand the simulation estimated 0.3149.

About Thrust

Writing code using CUDA API is very powerful in terms of controlling the hardware, but there are high level libraries like Thrust C++ template library, which provides many fundamental programming logic like sorting, prefix-sums, reductions, transformations etc.. The best part is that Thrust consists only of header files and is distributed with CUDA 4.0 installation.

If you are not familiar with terms like GPGPU and Thrust, I’m suggesting you to check out the background information on my previous posts.


The example is a console application written in C++. But you can easily transform it to a DLL to use it from your C# application. (previous posts)

I’ve used Visual Studio 2010 create the C++ console application. If you already did not, you need to install the NVIDIA CUDA Toolkit 4.0 and a supported graphics device driver from the same link. The new CUDA Toolkit 4.1 RC1 is also available at CUDA zone, but the project files are built on 4.0. Also do not forget to install the Build Customization BUG FIX Update from the same link for CUDA Toolkit 4.0.

Once the CUDA Toolkit is installed, creating CUDA enabled projects is really simple. For those who are not familiar using native C++ CUDA enabled projects, please follow the steps below to create one:

  • Create a Visual C++ console project in Visual Studio 2010 by selecting Empty project on the wizard,
  • Open Build Customization from the C++ project context menu, and check the CUDA 4.0(.targets, .props) checkbox,
  • Open the project properties, expand the Configuration Properties, Linker and select Input. Edit the additional dependencies and add cudart.lib.
  • Add a new empty source file ending with .cu.

You can also skip the steps above and download the example solution and project files directly from here.


The main application consists of calling thrust::transform_reduce 50 times to run the intersection estimation simulation. transform_reduce performs a reduction on the transformation of the sequence [first, last) according to unary_op. The unary_op is applied to each element of the sequence and then the result is reduced to a single value with binary_op.

The main code is as follows:

int main(void)
  // use 50 independent seeds
  int M = 50;
  //Create some circles in the device
  thrust::host_vector dCircles;
  dCircles.push_back(CIRCLE(0.0f, 0.0f));
  dCircles.push_back(CIRCLE(1.0f, 0.0f));
  dCircles.push_back(CIRCLE(1.0f, 1.0f));
  dCircles.push_back(CIRCLE(0.0f, 1.0f));

 //The kernel can not access host or device vector directly,
 //therefore get the device pointer to the circles to pass to the kernel
  thrust::device_vector circles = dCircles;
  CIRCLE * circleArray = thrust::raw_pointer_cast( &circles[0] );
  float estimate = thrust::transform_reduce(thrust::counting_iterator(0),
           estimate_intersection(circleArray, circles.size()),
  estimate /= M;
  std::cout << std::setprecision(6);
  //calculate area with gemometry : (pi + 3 - 3*sqrt(3)) / 3 = 0.31515s
  std::cout << "the area is estimated as " << estimate
            << std::endl << ". It should be 0.31515." ;
  return 0;

The unary_op has the Monte Carlo simulation logic implemented in the estimate_intersection functor. The estimate_intersection is a method derived from the thrust::unary_function class and returning the estimated intersect area as float. Using estimate_intersection in tranform_reduce means estimating the intersect area for every data element provided to tranform_reduce. For the data elements we are using two thrust::counting_iterators. This creates a range filled with a sequence from 1 to 50, without explicitly storing anything in the memory. Using a sequence of numbers helps us to assign different thread id for every estimate_intersection call. This is important to generate distinct seed for the random number generator of the simulation. (I’ve mentioned about random number generator seeds in my previous posts.)

For the reduction part of the tranform_reduce we are using the thrust::plus() binary functor, which sums all results into one number. At last we divide the result into 50 to find the average intersect area value.

Our goal with this code is to run the simulation on the device (GPU) and retrieve the result back to the host. Therefore any data we are going to use on the simulation must be placed into the device memory. That is exactly what is happening before we call thrust::transform_reduce. We are preparing properties of all circles we will try to intersect using the CIRCLE object defined below.

struct CIRCLE{
   float x,y;
   CIRCLE(float _x, float _y) : x(_x), y(_y){}
} ;

With thrust::host_vector dCircles; in the main code, we are defining a vector object in the host memory. Using a Thrust host vector object over a custom memory simplifies transferring data directly to device with the thrust::device_vector circles = dCircles; call. As you may know, transferring data between device and host memory in CUDA C is handled with cudaMemcpy. But Thrust has the equal operator overload, which allows you to copy memory easily.

On the next line we access the raw pointer of the circles object with the help of the thrust::raw_pointer_cast method. We do this because the estimate_intersection method can only accept a device pointer to the CIRCLE object array.

Simulation Method

The estimate_intersection unary function implements the simulation logic. A unary function is a function which takes one argument, has a () operator overload and returns one value. In our case the function takes the thrust::counting_iterator generated unique index number and returns the area of the intersection as float. Another important part of the method is the constructor (seen below) which takes in the device pointer to the CIRCLE array and the length of the allocated memory.

struct estimate_intersection : public thrust::unary_function
CIRCLE * Circles;
int CircleCount;

estimate_intersection(CIRCLE * circles, int circleCount) :
   Circles(circles), CircleCount(circleCount)

  __host__ __device__
  float operator()(unsigned int thread_id)
    float sum = 0;
    unsigned int N = 30000; // samples per thread
    unsigned int seed = hash(thread_id);
    // seed a random number generator
    thrust::default_random_engine rng(seed);
    // create a mapping from random numbers to [0,1)
    thrust::uniform_real_distribution u01(0,1);

    // take N samples
    for(unsigned int i = 0; i < N; ++i)
      // draw a sample from the unit square
      double x = u01(rng);
      double y = u01(rng);
      bool inside = false;

      //check if the point is inside all circles
      for(unsigned int k = 0; k < CircleCount; ++k)
       double dy,dx;
       //check if the point is further from
       //the center of the circle than the radius
       dx = Circles[k].x - x;
       dy = Circles[k].y - y;
       if ((dx*dx + dy*dy) <= 1)
        inside = true;
        inside = false;
      if (inside)
       sum += 1.0f;
    // divide by N
   return sum / N;

In order to run the code on the device and call it from the host, the () operator overload has to be defined as __host__ __device__. The rest of the code is the Monte Carlo simulation logic as follows:

1) Initiate the thrust default random number generator

2) Generate 30K random x and y values

3) Loop through all circles and check if the x and y value is inside the circle by calculating the hypotenuse

4) If all circles are inside the x and y coordinates then increase the found points count

5) return the average found points count

That’s it! I hope you enjoy it.

In addition to the code I included here, there are header includes and a hashing algorithm. You can download the code from here.

About the Implementations

The Monte Carlo simulation I provided on this post is an example and therefore I’m not guaranteeing that it will perform good enough in your particular solution. Also, for clarity there is almost no exception handling and logging implemented. This is not an API; my goal is to give you a high level idea how you can use utilize the GPU for simulations. Therefore, it is important that you re-factor the code for your own use.

Some of the code is taken from original sample and is under Apache License V 2, the rest is my code which is free to use without any restriction or obligation.


Thrust is a powerful library providing you with simple ways to accomplish complicated parallel computation tasks. There are many libraries like Thrust which are built on CUDA C. These libraries will save you many engineering hours on parallel algorithm implementation and allow you to concentrate on your real business problem. You can check out GPU Computing Webinars for presentations on this area.

C# language, C++, CodeProject, GPGPU, Technical

Massively Parallel RNG using CUDA C, Thrust and C#


On this post I’ll give you some simple examples how to use the massively parallel GPU on your computer to generate uniformly distributed psuedo-random numbers. Why GPU? Because it is orders of magnitude faster than CPU, does not occupy your CPU time, it is already on all computers and many other reasons I mentioned in my previous post. While there are maybe hundreds of ways to generate pseudorandom numbers I only covered four ways to do it on NVIDIA cards using CUDA related APIs:

1) A Basic Linear Congruential Generator (LCG) implementation using CUDA C
2) A Thrust C++ template library implementation
3) An NVIDIA CURAND implementation
4) A Mersenne Twister implementation using CUDA C

In order to demonstrate how to utilize the GPU, the implementations are provided as DLL’s and used within C# sample application. Perhaps there are many other APIs and ways worth to talk about to utilize your GPU within your C# application, but this post’s scope is limited only to the subject I mentioned above. I’m suggesting you to visit http://gpgpu.org/, if you already did not, to see the endless possibilities in this area.

While I was preparing these samples I saw that visualizing the data is very important to understand the algorithms. Therefore, I used Microsoft WPF on C# to visualize the generated random numbers. You can use your own application and copy the classes under the RNGLibs folder.

All code can be downloaded from this link: RNGTests09182011.zip.


About Random Number Generators (RNG)

The generation of random numbers is important in many applications like simulations, cryptography, sampling and mostly in statistics. A sequence of numbers is random when it does not have a recognizable pattern in it or in other words if it is non-deterministic. Although non-deterministic random numbers are ideal, the computer generated, deterministic random numbers can be statistically “random enough”. These random numbers are named as pseudo-random numbers and can have easily identifiable patterns if the algorithm is not chosen wisely. ( http://en.wikipedia.org/wiki/Pseudorandom_number_generator )

There are many pseudo-random number generators and also many different implementations of them in sequential and parallel environments. ( http://en.wikipedia.org/wiki/List_of_pseudorandom_number_generators ) On this post I used only Linear Congruential Generators, Xorshift and Mersenne Twister. Therefore, I explained only these three algorithm, but you can use CUDA to develop also other RNGs.


As I mentioned in my previous post, writing code using CUDA API is very powerful in terms of controlling the hardware, but there are high level libraries like Thrust C++ template library, which provides many fundamental programming logic like sorting, prefix-sums, reductions, transformations etc.. The best part is that Thrust consists only of header files and is distributed with CUDA 4.0 installation.


I’ve used Visual Studio 2010 to host one C# Windows Application and native C++ dlls for RNG implementations as seen in the solution structure below:

  • RNGVisual (Main C# application)
  • CURACRNGLib (CUDA C RNG implementation)
  • CURANDRNGLib (CURAND RNG implementation)
  • ThrustRNGLib (Thrust RNG Implementation)
  • MersenneTwisterRNGLib (Mersenne Twister RNG implementation)

The only additional API is the NVIDIA CUDA Toolkit 4.0, which you will need to install along with a supported graphics device driver from the same link. Also do not forget to install the Build Customization BUG FIX Update from the same link or from here.

Once the CUDA Toolkit is installed creating CUDA enabled projects is really simple. For those who are not familiar using native C++ CUDA enabled projects, please follow the steps below to create one:

  • Create a Visual C++ console project in Visual Studio 2010 by selecting DLL and Empty project on the wizard,
  • Open Build Customization from the C++ project context menu, and check the CUDA 4.0(.targets, .props) checkbox,
  • Open the project properties, expand the Configuration Properties, Linker and select Input. Edit the additional dependencies and add cudart.lib.
  • Add a new empty source file ending with .cu.


WPF Application

The RNGVisual C# WPF Application provides visualization of the random numbers in 2D and 3D. It allows you to select an RNG algorithm (.Net, CUDA C, CURAND, Thrust or Merseene Twister) and allows you to set some display parameters and processor parameters. With any number count below 10K, all RNGs calculate in about one millisecond and most of the time is spent drawing the squares to the screen. Therefore, the time should not confuse you in terms of performance comparison. You can run the algorithms with 100K numbers without the visualization and see the difference on your hardware. But please be aware that it is better to use CUDA events with cudaEventRecord to time GPU execution more precisely.



RNGVisual implements various proxy classes, which uses platform invoke to call RNG methods exported from the native C++ dlls. I used the same export and import technique in my previous post. The RNG libraries have the following two exports, one for CPU implementation and one for GPU implementation:

extern "C" __declspec(dllexport) void __cdecl
    GPU_RNG(float*, unsigned int, unsigned int);

extern "C" __declspec(dllexport) void __cdecl
    CPU_RNG(float*, unsigned int, unsigned int);

The first argument is a pointer to the memory location to hold the random numbers. The second argument is the size of the array and the last argument is the initial seed.

An important point of random number generation is selecting the seed value, because the same seed will give the same result. While there are many different techniques studied, I used my own method of combining current time, CPU load and available physical memory with the help of the Windows Management Instrumentation (WMI); it still does not perform well in multi-threaded solutions, but it gives at least a better random start. The implementation is in the CPUHelper class of the RNGVisual application.

A Linear Congruential Generator (LCG) implementation using CUDA C

The first RNG project is using native CUDA Runtime API to implement the oldest and best-known pseudorandom number generator algorithms named LCG. LCG is fast and requires minimal memory to retain state. Therefore, LCG is very efficient for simulating multiple independent streams. But LCGs have some disadvantages and should not be used for applications where high-quality randomness is critical. In fact the simple example I implemented repeats numbers in a very short period and should be enhanced with methods like explained in GPUGems 3 (37-4).

LCG is as simple as seen in the formula below; starting with a seed (Xn), the next random number is determined with (a * seed + c) mod m.
Xn+1 = (a Xn + c) (mod m)
where 0 < m, 0 < a < m, 0 <= x0 < m and 0 <= c < m

Below is a sequential implementation of the LCG algorithm, which generates 100 pseudo-random numbers:

random[0] = 123; //some initial seed
for(int i = 1; i < 100; i++)
 random[i] = ( a * random[i-1] + c) % m;

The CUDACRNGLib project has a very basic implementation of LCG by distributing the work onto 256 threads. Because the same seed will result in the same random number, first we generate different random seeds for every thread. When the kernel below is executed, every thread generates one section of the random number sequence:

__global__ void RNGKernel(float * randomNumbers, unsigned int numberCount,
    unsigned int * seeds, unsigned int c, unsigned int a, unsigned int M) 
    int startIdx = threadIdx.x * numberCount; 
    unsigned int x = seeds[threadIdx.x]; 
    for(int i=0; i < numberCount; i++) { 
        x = (a * x + c) % M; //M is shown for purpose of example 
        randomNumbers[startIdx + i]= (float)x / (float)M; //normalize  

As I mentioned before, this implementation is simplified to give you an idea how you can start using CUDA C. It even has static block count of one and thread count of 256. If you plan to go for production code, it is good to start many blocks of threads. You may want to check out a better implementation on GPU Gems 3 (37-7) or check out Arnold and Meel’s implementation, which provides also better randomness.

A Thrust C++ template library implementation

The Thrust library default random engine ( default_random_engine ) is a Linear Congruential Generator ( may change in the future ) with a = 48271, c = 0 and m = 2^31. Because c equals to zero, the algorithm is also named multiplicative congruential method or Lehmer RNG.

The ThrustRNGLib has a very basic implementation of Thrust default random engine by running the following functor to generate one random number. A functor is a type of class in C++ that overloads the operator() and therefore allows to be called like an ordinary function. Thrust provides unary_function and binary_function functors. Below I used the unary_function because my functor requires on argument to passed into the function:

struct RandomNumberFunctor : 
    public thrust::unary_function<unsigned int, float>
    unsigned int mainSeed;
    RandomNumberFunctor(unsigned int _mainSeed) : 
        mainSeed(_mainSeed) {}

    __host__ __device__
        float operator()(unsigned int threadIdx) 
        unsigned int seed = hash(threadIdx) * mainSeed;

        // seed a random number generator
        thrust::default_random_engine rng(seed);

        // create a mapping from random numbers to [0,1)
        thrust::uniform_real_distribution<float> u01(0,1);

        return u01(rng);

Using thrust to utilize the GPU is the simplest way to go. You can see the difference by comparing the GPU_RNG from below to the CUDACRNGLib GPU_RNG implementation. Using CUDA C gives you full control of the toolkit but it comes with the price of writing more code.

extern void GPU_RNG(float * h_randomData, unsigned int dataCount, unsigned int mainSeed)
    //Allocate device vector
    thrust::device_vector<float> d_rngBuffer(dataCount);

    //generate random numbers
        d_rngBuffer.begin(), RandomNumberFunctor(mainSeed));

    //copy the random mask back to host
    thrust::copy(d_rngBuffer.begin(), d_rngBuffer.end(), h_randomData);

Another good part of Thrust is that every implementation (except copy) exist for GPU as well as for CPU. The CPU implementation is also another three lines of code, this time by using the thrust::generate to generate the random numbers by using the C++ standard template library rand method and then after using thrust::transform to normalize the integer result into float with the help of the [](float n) {return n / (RAND_MAX + 1);} lambda expression. I used the lambda expression instead of a functor to show you also this possibility. Especially on the upcoming
Microsoft C++ AMP, lambda expressions play a big role. Lambda expression are handy in C++ as well as in C#, but it comes with a price of giving up unit testing of the inline expression.

An NVIDIA CURAND implementation

The NVIDIA CURAND library provides an API for simple and efficient generation of high-quality pseudorandom and quasirandom numbers. The CURAND library default pseudorandom engine is a XORWOW implementation of the Xorshift RNG (page 5) and it produces higher quality random numbers than LCG.
In order to start using CURAND, you only need to include the curand.h header and add the curand.lib to the Linker additional dependencies on the Linker settings.

Like ThrustRNGLib Thrust implementation, the CURANDRNGLib has a very basic implementation by running the following main code to generate a series of random numbers:

    //Create a new generator
    curandCreateGenerator(&m_prng, CURAND_RNG_PSEUDO_DEFAULT);
    //Set the generator options
    curandSetPseudoRandomGeneratorSeed(m_prng, (unsigned long) mainSeed);
    //Generate random numbers
    curandGenerateUniform(m_prng, d_randomData, dataCount);

CURAND provides the curandCreateGeneratorHost method besides the curandCreateGenerator method, to generate random numbers on the CPU instead of the GPU. Therefore the CPU part is as simple as the GPU part.

A Mersenne Twister implementation using CUDA C

Mersenne Twister ( MT ) is an algorithm developed by Makoto Matsumoto and provides very fast generation of high-quality random numbers. ( MT Home Page ) A common Mersenne twister implementation, uses an LCG to generate seed data.
Originally there are two MT algorithms suitable to use with CUDA: TinyMT and Mersenne Twister for Graphics Processors (MTGP). But I implemented part of the code from the NVIDIA CUDA Toolkit 4.0 MersenneTwister sample, which uses the original code from Makoto Matsumoto anyways.

The Mersenne Twister RNG is maybe the most complicated implementation out of the other three RNGs I provided, but with that you can look into the algorithm, unlike CURAND. The MersenneTwisterRNG.cpp file from the MersenneTwisterRNGLib project is the entry point to the library and exports the same GPU_RNG and CPU_RNG methods as the other libraries. I simplified the host code as much as possible and placed all GPU logic into the GPURNG.cu file. The remaining simple host code can be seen below:

extern void GPU_RNG(float * h_randomData, unsigned int dataCount, 
    unsigned int mainSeed)
	float * d_randomData = 0;

	//load GPU twisters configuration

	//find the rounded up data count 
    //because the generator generates in multiples of 4096
	int numbersPerRNG = iAlignUp(iDivUp(dataCount, MT_RNG_COUNT), 2);
	int randomDataCount = MT_RNG_COUNT * numbersPerRNG;

	//allocate device memory
	size_t randomDataSize = randomDataCount * sizeof(float);
	cudaMalloc((void**)&d_randomData, randomDataSize);

	//Call the generator
	RNGOnGPU(32, 128, d_randomData, numbersPerRNG);

	//Make sure all GPU work is done

	//Copy memory back to the device
	cudaMemcpy(h_randomData, d_randomData, dataCount * sizeof(float), 

	//free device memory

About the Implementations

The Pseudo-random number generators I provided on this post are widely used algorithms, but still I’m not guaranteeing that any of them will perform good enough in your particular solution. In fact, I left some generators poor by purpose to point out the core algorithm and provide variations in randomness. Also, for sake of clarity there is almost no exception handling and logging implemented. This is not an API; my goal is to give you a high level idea how you can use Thrust, CUDA C, CURAND to generate pseudo-random number. Therefore, it is important that you research the algorithms on-line and re-factor the code for your own use.

Some of the code is taken from original NVIDIA samples and have copyright notice on them, the rest is my code which is free to use without any restriction or obligation.


As in every field of computer science, there are many ways to solve a problem and the possibilities expand exponentially. I just scratched the surface of the possibility of using CUDA to add pseudorandom number generation to your C# application. I hope this post will help you in your project.

C# language, C++, CodeProject, GPGPU, Technical

Faster Sorting in C# by Utilizing GPU with NVIDIA CUDA


On this post I would like to give an entry level example how you can use NVIDIA CUDA technology to achieve better performance within C# with minimum possible amount of code. This is a very specific scenario where the main application is in C#, and the sorting algorithm is provided in C++ as an API and you have an NVIDIA CUDA capable graphics card. I hope this basic sample will give you some insight and encourage you to invest into parallel technologies.


Parallel Computing
For more than two decades CPUs drove rapid performance increase and price reductions. With that, we may get used to having speed boost of our code running on a new piece of hardware. As you may recall, this relentless performance improvement is described with two words: Moore’s law. While the continuation of Moore’s law is debated on many platforms, because of the slow down starting at 2003, many of us already jumped to the parallel bandwagon using systems based on SIMD or MIMD models. With parallel computing, application logic suitable for parallelization will most likely run faster without waiting for tomorrow’s hardware. (It is “most likely” because there are theoretical limits explained by Amdahl’s law.)

A GPU is a massively multithreaded processor which supports thousands of concurrent threads. Using GPU to do general purpose scientific and engineering computing is named GPGPU. NVIDIA is one of the leading GPU manufacturers and they are providing fully programmable GPU technology for high-level languages like C, C++ and Fortran with the architecture named CUDA. The NVIDIA hardware uses the SIMT (single-instruction multiple-thread) execution model rather than the popular SIMD model. I gave a SIMD example in one my previous posts, how to utilize an SSE4 CPU instruction within C#. And in this post we will be using SIMT execution model through NVIDIA CUDA. (Under the cover, on SIMD multiple data elements for a single instruction are collected and packed into a single register. On the other hand on SIMT, all threads process data in their own registers. )

Writing code using CUDA API is very powerful in terms of controlling the hardware, but it requires to handle memory and execution which is outside the scope of this post. Therefore I’ll use a high-level programming interface named Thrust. As describe at http://code.google.com/p/thrust/, Thrust is a CUDA library of parallel algorithms with an interface resembling the C++ Standard Template Library (STL). Thrust library provides many fundamental programming logic like sorting, prefix-sums, reductions, transformations etc..

Thrust consists of header files and is distributed with CUDA 4.0 installation. Therefore, you don’t have to install it additionally.

Radix Sort
The reason I choose sorting is that it is a widely used fundamental computational building block to many problems. Radix sort is one of the oldest algorithms and is very efficient for sorting small keys sequentially with the algorithmic complexity of O(n). One of the two sorting algorithms Thrust provides is Radix Sort, therefore thrust fits our example very well. There is a paper about Designing Efficient Sorting Algorithms for Manycore GPUs, which explains how the algorithm works on the GPU.


We will be using Visual Studio 2010 to host one C# application and one C++ library. The only additional API is the NVIDIA CUDA Toolkit 4.0. You will need to install the CUDA Toolkit and a supported graphics device driver from the same link.
Also do not forget to install the Build Customization BUG FIX Update from the same link or from here.


The solution consist of one C# console application and a C++ native dll. The C++ native dll exports a method which has the code to run the sort algorithm on the GPU. The C# code uses platform invoke to call the exported method of the C++ dll.

For comparison purpose I’m also providing a C# implementation of Radix Sort. The C# implementation takes an array of 32 bit integers and process it in three steps for every 8 bit part:
1) Creating histogram of data,
2) Running prefix sum(scan),
3) Adding the data to bin.

On some implementations the histogram is named as count or difference, but here I used the name histogram, because there is very good CUDA implementation I may write about in a future post. It is basically finding out how many of every byte is in the array. Because the max byte value is 0xFF, therefore the size of the histogram is 256, reflecting the maximum number of different values.

The second step is to run prefix sum on the histogram, which is simply adding one value to the right neighbor in the array. In that way every element in the array will contain it’s left neighbor. Actually, there is a very efficient way to to prefix sum on parallel environment explained by the paper Efficient Parallel Scan Algorithms for GPUs. Which you may find usefull for understanding the logic behind the GPU implementation.

The last step is to place the value into the correct bin by using the prefix sum result at the value index to determine the temporary array index.

The code is as follows:

public static void CPURadixSort(int[] sourceData)
 int size = sourceData.Length;
 //temporary array for every byte iteration
 int[] tempData = new int[size]; 
 //histogram of the last byte 
 int[] histogram = new int[256]; 
 //The prefix sum of the histogram
 int[] prefixSum = new int[256]; 
  fixed (int* pTempData = tempData)
   fixed (int* pSourceData = sourceData)
    int* pTemp = pTempData;
    int* pBck;
    int* pSource = pSourceData;
    //Loop through every byte of 4 byte integer
    for (int byteIdx = 0; byteIdx < 4; byteIdx++)
     	int shift = byteIdx * 8; //total bits to shift the numbers
     	//Calculate histogram of the last byte of the data
     	for (int i = 0; i < size; i++)
     	 histogram[(pSource[i] >> shift) & 0xFF]++;
     	//Calculate prefix-sum of the histogram
     	prefixSum[0] = 0;
     	for (int i = 1; i < 256; i++)
     	 prefixSum[i] = prefixSum[i - 1] + histogram[i - 1];
     	//Get the prefix-sum array index of the last byte, increase 
        //it by one. That gives us the the index we want to place 
        //the data
     	for (int i = 0; i < size; i++)
     	 pTemp[prefixSum[(pSource[i] >> shift) & 0xFF]++] = pSource[i];
     	//Swap the pointers
     	pBck = pSource;
     	pSource = pTemp;
     	pTemp = pBck;
     	//reset the histogram
     	for (int i = 0; i < 256; i++)
     	 histogram[i] = 0;

One more detail about the code is that I used unsafe code in order to swap the pointers for the temporary buffer and source data to gain another half second. But, it is possible to remove that part and use the tempData.CopyTo(sourceData, 0) instead.

When I run the code above on an Intel Xeon 2.40 Ghz CPU for 2^25 (33,554,432) random integers, it executed in 3.038 seconds.

Next step is to create the sorting code on the GPU and call it within our C# code. For this we need to create a DLL in C++ by using the default .Net project wizard. Please follow the steps to acomplish this step:
Step 1) Create a Visual C++ console project named “RadixSortGPU.dll” in Visual Studio 2010 by selecting DLL and Empty project on the wizard,

Step 2) Open Build Customization from the C++ project context menu, and ceck the CUDA 4.0(.targets, .props) checkbox,

Step 3) Open the project properties, expand the Configuration Properties, Linker and select Input. Edit the additional dependencies and add cudart.lib.

Step 4) Add a new empty source file ending with .cu and paste the following code into it:

#include <thrust/device_vector.h>
#include <thrust/sort.h>
#include <thrust/copy.h>
#include <thrust/detail/type_traits.h>

extern "C" __declspec(dllexport) void __cdecl GPURadixSort(int*, unsigned int);

extern void GPURadixSort(int* data, unsigned int numElements)
	thrust::device_vector<int> d_data(data, data + numElements);

	thrust::stable_sort(d_data.begin(), d_data.end());

	thrust::copy(d_data.begin(), d_data.end(), data);

This code is the only program logic you need to sort a given integer array on the GPU. For the sake of this example I did not include any exception handling and validation. But for any production code you most probably want to add exception handling and validation logic.

As you can see in the code, starting using Thrust is very simple by just including the required headers.

The dllexport definition makes the GPURadixSort usable from within another library. In this way we will be able to call this method from our C# library.

The GPURadixSort method itself is simply copying the data array into a container named thrust::device_vector, resulting the data to be placed into the GPU global memory, sorting the data on the GPU and copying back the data from the GPU memory to our input array. and
the begin() and end() methods are iterators pointing to the data at the beginning or end of the array. For more information how to use vectors and iterators please visit the Thrust Quick Start Guide.

Step 5) I’m assuming you already have a C# project you can paste the following code into your class,

[DllImport("RadixSortGPU.dll", CallingConvention = CallingConvention.Cdecl)]
public static extern void GPURadixSort(
 [MarshalAsAttribute(UnmanagedType.LPArray, ArraySubType = UnmanagedType.I4)]
 int[] data, uint numElements);

This code provides the information needed to call the GPURadixSort function exported from our native DLL.

Step 6) That’s it. Now you can call the GPURadixSort within your C# code to sort your array.

Based on my environment with a NVIDIA Quadro 600 GPU, the sort takes about 1.211 seconds, which is almost 3x faster than the CPU version. The NVIDIA Quadro 600 GPU I used is a low profile graphics processor which cost about $150. 96 cores, 1GB memory and 25.6 GB/s memory bandwidth maybe sound powerful, but it is low profile compared to GPUs like Quadro 6000, with 448 cores, 6 GB memory and 144 GB/s memory bandwidth. There are also the Tesla family GPUs for high performance computing needs you may want to check out at NVIDIA.

Below I provided some code to check if the GPU memory is enough to hold the given array, which is one of the important validations you want to perform if you target heterogeneous environments.

using namespace std;
extern void GPURadixSort(int* data, unsigned int numElements)
 int deviceID = -1;
 if (cudaSuccess == cudaGetDevice(&deviceID))
     cudaDeviceProp devprop;
     cudaGetDeviceProperties(&devprop, deviceID);
     unsigned int totalMem = numElements * sizeof(int);
     if (devprop.totalGlobalMem < totalMem)
        //Handle error

As I mentioned before, under the high-level Thrust library the powerful CUDA API resides. If Thrust does not have the functionality you are looking for, you can always use CUDA directly. It is not as easy as using Thrust, but if you will follow the resources below, it will not take you much to start using this great technology.


We are using parallelism mostly to gain performance when a program runs slow or we need to handle huge amount of calculation. But, I believe that parallelism is a very important software aspect allowing us to use the underlying hardware wisely and gain advantage in this competitive market, where two functionality wise similar applications can only compete on software qualities like performance. There are many parallel technologies which will allow you to achieve higher performance with the same hardware. One of them is definitely NVIDIA CUDA, and some others you want to look into are TPL, PPL, C++ AMP, CUDA, OpenCL, MPI.

C# language, C++, CodeProject, Technical

How to use CPU instructions in C# to gain performance


Today, .NetFramework and C# became very common for developing even the most complex application very easily. I remember that before getting our hands-on on C# in 2002 we were using all kind of programming languages for different purposes ranging from Assembly, C++ to PowerBuilder. But, I also remember the power of using assembly or C++ to use every small drop of power of your hardware resources. Every once in a while I’m getting to a project where using regular framework functionality puts my computer to grill for couple days to calculate something. In those cases I go back to the good old C++ or Assembly routines to use the power of my computer. On this blog I will be showing you the simplest way to take advantage of your hardware without introducing any code complexity.


I believe that samples are the best teacher; therefore, I’ll be using a sample CPU instruction from Streaming SIMD Extensions (SSE). SSE is just one of the many instruction set extension to X86 architecture. I’ll be using an instruction named PTEST from the SSE4 instructions, which is almost in all Intel and AMD CPUs today. You can visit the links above for supported CPUs. PTEST helps us to perform bitwise comparison between two 128-bit parameters. I picked this because it was a good sample using also data structures. You can easily lookup online for any other instruction set for your project requirements.

We will be using unmanaged C++, wrap it with managed C++ and call it from C#. Don’t worry, it is easier than it sounds.
Thanks to MSDN we will have all necessary information at Alphabetical Listing of Intrinsic Functions. The PTEST _mm_testc_si128 is documented at http://msdn.microsoft.com/en-us/library/bb513983.aspx. If we would use C++ we would end up having the code from the MSDN link:

#include <stdio.h>
#include <smmintrin.h>

int main ()
    __m128i a, b;

    a.m128i_u64[0] = 0xAAAA55551111FFFF;
    b.m128i_u64[0] = 0xAAAA55551111FFFF;

    a.m128i_u64[1] = 0xFEDCBA9876543210;
    b.m128i_u64[1] = 0xFEDCBA9876543210;

    int res1 = _mm_testc_si128(a, b);

    a.m128i_u64[0] = 0xAAAA55551011FFFF;

    int res2 = _mm_testc_si128(a, b);

    printf_s("First result should be 1: %d\nSecond result should be 0: %d\n",
                res1, res2);

    return 0;

I would like to point out that there are many ways to develop a software and the one I’m providing here is maybe not the best solution for your requirement. I’m just providing one way to accomplish a task, it depends to you to fit into your solution.

Ok, let’s start with a fresh new solution: ( I’m assumig you have Visual Studio 2010 with C++ language installed on it.)

1 ) Add a new C# console application to your solution, for the testing purpose
2 ) Add a new Visual C++ > CLR > Class Library project to your solution, named TestCPUInt
3 ) Add a reference to your console application and select TestCPUInt from the projects

Now we are ready to code.
4 ) Open the TestCPUInt.h file, if it is already not opened
5 ) Insert the following code right on top of the document

#include <smmintrin.h>
#include <memory.h>

#pragma unmanaged

class SSE4_CPP
//Code here

#pragma managed

This is the infrastructure for our unmanaged C++ code which will do the SSE4 call. As you can see I placed our unmanaged code between #pragma unmanaged and #pragme managed. I think it is a great feature to be able to write unmanaged and managed code together. You may prefer to place the unmanaged code to another file, but I used one file for clarity. We used here two include header files, smmintrin.h and memory.h: the first one is for the SSE4 instructions and the other one is for a method I used to copy memory.
6 ) Now paste the following code at the //Code here location:

	int PTEST( __int16* bufferA, __int16* bufferB)
		__m128i a, b;

		//transfer the buffers to the _m128i data type, because we do not want to handle with that in managed code
		memcpy(a.m128i_i16, bufferA, sizeof(a.m128i_i16));
		memcpy(b.m128i_i16, bufferB, sizeof(b.m128i_i16));

		//Call the SSE4 PTEST instructions 
		return _mm_testc_si128(a, b);

This _mm_testc_si128 will emit the SSE4 PTEST instruction. We have a little memory operation right before it to fill out the __m128i data structure on the C++ code. I used memcpy to trasfer the data from the bufferA and bufferB arguments to the __m128i data structure to push it to the PTEST. I prefered to do this here to separate tthe whole SSE4 specific implementation. I could also send the __m128i to the PTEST method, but that would be more complex.

As I mentioned before, in this example I used the PTEST sample with a data structure, you may run into other instructions which require only a pointer, in that case you don’t need to do the memcpy operation. There will maybe some challenges if you are not familiar with C++, especially when the IntelliSense is removed in Visual Studio 2010 VC++, but you can always check out for online answers. For example the __m128i data structure is visible in the emmintrin.h file, which is located somewhere [Program Files]\Microsoft Visual Studio 10.0\VC\include. Or you can check all fundamental data types if you are not sure what to use instead of __int16*.

7 ) Now paste the following code on top of your managed C++ code. Which is on the bottom of your TestCPUInt.h file in the namespace section.

namespace TestCPUInt {

	public ref class SSE4
		int PTestWPointer(__int16* pBufferA, __int16* pBufferB)
			SSE4_CPP * sse4_cpp = new SSE4_CPP();
			return sse4_cpp->PTEST(pBufferA, pBufferB);


What we did here is to pass forward the pointers pBufferA and pBufferB, which we are going to call from C#, into the unmanaged code. For those who are not familiar with pointers, the * sign defines a pointer: __int16* means a pointer to a 16 bit integer. In our case that is the address of the first element of an array.
There are also ways without using managed C++ to call a dynamic library, but as I mentioned before, I’m showing only the simplest way for a C# developer.

8 ) Let’s go to our C# code in the console application to use this functionality.
9 ) First we have to switch the application to allow unsafe code. For that goto the project properties and check the “Allow unsafe code” under the build tab.
10 ) Add the following code to your Program.cs file:

static int TestCPUWithPointer(short[] bufferA, short[] bufferB)
	SSE4 sse4 = new SSE4();
		//fix the buffer variables in memory to prevent from getting moved by the garbage collector
		fixed (short* pBufferA = bufferA)
			fixed (short* pBufferB = bufferB)
				return sse4.PTestWPointer(pBufferA, pBufferB);

If you never used unsafe code before, you can check out unsafe (C# Reference). Actually, it is fairly simple logic; the PTestWPointer required a pointer to an array and the only way to get the pointer to an array is to use the fixed statement. The fixed statement pins my buffer array in memory in order to prevent the garbage collector to move it around. But that comes with a cost: in one of my projects the system was slowing down because of too many fixed objects in memory. Anyways, you may have to experiment for your own project.
That’s it!

But we will not stop here, for comparison I did the same operation in C#, as seen below:

static int TestCLR(short[] bufferA, short[] bufferB)
	//We want to test if all bits set in bufferB are also set in bufferA
	for (int i = 0; i < bufferA.Length; i++)
		if ((bufferA[i] & bufferB[i]) != bufferB[i])
			return 0;
	return 1;

Here I simply calculate if every bit of bufferB is in bufferA; PTEST does the same.

On the rest of the application I compared the performance of these two methods. Below is a code which does the comparison for sake of testing:

static void Main(string[] args)
	int testCount = 10000000;
	short[] buffer1 = new short[8];
	short[] buffer2 = new short[8];

	for (int i = 0; i < 8; i++)
		buffer1[i] = 32100;
		buffer2[i] = 32100;
	Stopwatch sw = new Stopwatch();
	int testResult = 0;
	for (int i = 0; i < testCount; i++)
		testResult = TestCPUWithPointer(buffer1, buffer2);	
	Console.WriteLine("SSE4 PTEST took {0:G} and returned {1}", sw.Elapsed, testResult);

	for (int i = 0; i < testCount; i++)
		testResult = TestCLR(buffer1, buffer2);
	Console.WriteLine("C# Test took {0:G} and returned {1}", sw.Elapsed, testResult);


On my environment I gained %20 performance. On some of my projects I gained up to 20 fold performance.
A last thing I would like to do is to show you how to move the fixed usage from C# to managed C++. That makes you code little cleaner like the one below:

static int TestCPU(short[] bufferA, short[] bufferB)
	SSE4 sse4 = new SSE4();
	return sse4.PTest(bufferA, bufferB);

As you can see, it is only a method call. In order to do this we have to add the following code to the TestCPUInt class in the TestCPUInt.h file:

int PTest(array<__int16>^ bufferA, array<_int16>^ bufferB)
	pin_ptr<__int16> pinnedBufferA = &bufferA[0]; // pin pointer to first element in arr
	__int16* pBufferA = (__int16*)pinnedBufferA; // pointer to the first element in arr
	pin_ptr<__int16> pinnedBufferB = &bufferB[0]; // pin pointer to first element in arr
	__int16* pBufferB = (__int16*)pinnedBufferB; // pointer to the first element in arr

	SSE4_CPP * sse4_cpp = new SSE4_CPP();
	return sse4_cpp->PTEST(pBufferA, pBufferB);

This time our method takes a managed array object instead of a __int16 pointer and pins it in the memory like we did using fixed in C#.


I believe that as much as higher level frameworks we are using there will always be situations where we have to use our hardware resources more wisely. Sometimes these performance improvements save us big amount of hardware investment.
Please go ahead look into CPU instructions, Parallel computing, GPGPU etc. and understand the hardware; knowing your tools better will make you a better software architect.

C# language, CodeProject, Technical

Careful with Optional Arguments in C#4.0


Some time passed since optional arguments are introduced with Visual C# 2010 and we all got used to the convenience of not having to define method overloads for every different method signature. But, recently I came across of a limitation using optional arguments on enterprise solutions and now I’m using it with care.

The limitation is that if you use the optional arguments across libraries, the compiler will hard code the default value to the consumer and prevent you from re-deploying the provider library separately. It is very common scenario on enterprise applications where we have many libraries with different versions married to each other and this limitation makes it impossible to re-deploy only one dll without re-deploying all related libraries.

On this post I will explain you the details of this limitation. As I mentioned in my previous post Under the hood of anonymous methods in c#, it is very important to know the underlying architecture of a functionality you are using. This post is similar to that by explaining mechanics of the optional arguments.


I will not explain what optional arguments are, because it is already out since a while. But I will just give a quick reference to the description for those who are new to C#:

Optional arguments enable you to omit arguments for some parameters. … . Each optional parameter has a default value as part of its definition. If no argument is sent for that parameter, the default value is used. Default values must be constants.” (Named and Optional Arguments (C# Programming Guide) )

I run into this limitation while I was using some functionality from a secondary library, injected through an IoC container. The secondary library was accessed through an interface where some methods had optional arguments. I had everything deployed and working until I had to make some changes to the secondary library and alter the optional argument. After re-deploying the secondary library I figured out that the changes did not take effect and went I went into the IL code I figured out that the main library had the constant hard-coded into it.

In my situation I had interfaces, injections and a lot of complexity; to better picture the situation I will be using the simplest form of the limitation as in the following sample:

  • ProjectB has a method named Test with the following signature:

public void Test(string arg1 = "none")

  • ProjectA is referencing ProjectB and using the Test method with its default argument by making the following call:

static void Main(string[] args)
Class1 class1 = new Class1();

This works very well, because ProjectA is just using arg1 value as “none”. Now let’s look what is happening behing the scene:

If we compile ProjectA along with ProjectB and analyze the IL code, we will see that the optional argument is nothing more than a compiler feature. Because calling Test() is no different from calling Test(“none”), the compiler decides to compile our code as Test(“none”). That can be seen in the IL code and disassembled C# code below; the string constant “none” is hard-coded into ProjectA.

.method private hidebysig static void Main(string[] args) cil managed
L_0008: ldstr "none"
L_000d: callvirt instance void [ProjectB]ProjectB.Class1::Test(string)

private static void Main(string[] args)
new Class1().Test("none");

For libraries tightly coupled or in-library usage it is good that the compiler helps us eliminating some code and makes our life easier. But this comes with a price:

Let’s say we had to modify the Test method in ProjectB as Test( string arg1 = “something” ) and re-deploy it without re-deploying ProjectA. In this case ProjectA would still be calling the Test method with “none”.


Knowing this, it is good to use the optional arguments with caution across libraries when you have to support deployment scenarios with partial solutions.