Tuning a Monte Carlo Algo on GPU

http://www.pgroup.com/lit/articles/insider/v2n1a4.htm

CUDA shared memory array variable

http://stackoverflow.com/questions/9187899/cuda-shared-memory-array-variable

You have two choices for declaring shared memory inside a kernel - static or dynamic. I presume what you are doing at the moment looks something like this:

#define BLOCK_SIZE (16)

__global__ void sgemm0(const float *A, const float *B, float *C)
{
    __shared__ float As[BLOCK_SIZE][BLOCK_SIZE];

}

and you would like to be able to easily change BLOCK_SIZE.

One possibility is to continue to use static shared memory allocation, but make the allocation size a template parameter, like this:

template<int blocksize=16>
__global__ void sgemm1(const float *A, const float *B, float *C)
{
    __shared__ float As[blocksize][blocksize];

}
template void sgemm1<16>(const float *, const float *, float *C);

Then you can instantiate as many different block size variants at compile time as you need.

If you want to dynamically allocate the memory, define it like this:

__global__ void sgemm2(const float *A, const float *B, float *C)
{
    extern __shared__ float As[];

} 

and then add the size of the allocation as an argument to the kernel call:

size_t blocksize = BLOCK_SIZE * BLOCK_SIZE;
sgemm2<<< gridDim, blockDim, sizeof(float)*blocksize >>>(....);

If you have multiple statically declared arrays which you wish to replace with dynamically allocated shared memory, then be aware that there is only ever one dynamic shared memory allocation per kernel, so multiple items exits within (share) that memory segment. So if you had something like:

#define BLOCK_SIZE (16)

__global__ void sgemm0(const float *A, const float *B, float *C)
{
    __shared__ float As[BLOCK_SIZE][BLOCK_SIZE];
    __shared__ float Bs[BLOCK_SIZE][BLOCK_SIZE];

}

you could replace it with:

#define BLOCK_SIZE (16)

__global__ void sgemm3(const float *A, const float *B, float *C)
{
    extern __shared__ buffer[];

    float *As = &buffer[0];
    float *Bs = &buffer[BLOCK_SIZE*BLOCK_SIZE];

}

and launch the kernel like this:

size_t blocksize = 2 * BLOCK_SIZE * BLOCK_SIZE;
sgemm3<<< gridDim, blockDim, sizeof(float)*blocksize >>>(....);

All are equally valid, although I personally favour the template version because it can allow other compiler optimisation like automatic loop unrolling that the dynamic version cannot without extra work.

Rapid Problem Solving Using Thrust

http://on-demand.gputechconf.com/gtc-express/2011/presentations/Rapid-Problem-Solving-Using-Thrust.pdf

CUDA Libraries - Thrust
http://www.bu.edu/pasi/files/2011/07/Lecture6.pdf

Array of Vectors using Thrust

thrust::device_vector<float> vectors[3];
//thrust::device_vector<float> *vectors = new thrust::device_vector<float>[3];

vectors[0] = thrust::device_vector<float>(10);
vectors[1] = thrust::device_vector<float>(10);
vectors[2] = thrust::device_vector<float>(10);

http://stackoverflow.com/questions/12406948/array-of-vectors-using-thrust

Generating a random number vector between 0 and 1.0 using Thrust

http://stackoverflow.com/questions/12614164/generating-a-random-number-vector-between-0-and-1-0-using-thrust

#include <thrust/random.h>
#include <thrust/device_vector.h>
#include <thrust/transform.h>
#include <thrust/iterator/counting_iterator.h>
#include <iostream>

struct prg
{
    float a, b;

    __host__ __device__
    prg(float _a=0.f, float _b=1.f) : a(_a), b(_b) {};

    __host__ __device__
        float operator()(const unsigned int n) const
        {
            thrust::default_random_engine rng;
            thrust::uniform_real_distribution<float> dist(a, b);
            rng.discard(n);

            return dist(rng);
        }
};


int main(void)
{
    const int N = 20;

    thrust::device_vector<float> numbers(N);
    thrust::counting_iterator<unsigned int> index_sequence_begin(0);

    thrust::transform(index_sequence_begin,
            index_sequence_begin + N,
            numbers.begin(),
            prg(1.f,2.f));

    for(int i = 0; i < N; i++)
    {
        std::cout << numbers[i] << std::endl;
    }

    return 0;
}

How to measure time in NVIDIA CUDA?

Ivan's blog
http://ivanlife.wordpress.com/2011/05/09/time-cuda/

CUDA Developer Forum Discussion:
cudaEvent timer vs. Host timers
http://forums.developer.nvidia.com/devforum/discussion/7541/cudaevent-timers-vs-host-timers/p1
Parallel Nsight, NVIDIA Visual Profiler, CUDA profiler, and the CUPTI SDK provide the most accurate  methods to measure the execution time of a kernel. The measured time does not include the overhead to launch the kernel.

cudaEventRecord is the most accurate method to measure the setup and execution time of a kernel.

A high percision CPU timer can be used to measure the overhead of the launch, the execution of the kernel, and the completion notification. If you use this method I recommend that you call cudaDeviceSynchronize() before the first clock to make sure there is no outstanding work that might delay the launch of the kernel. This method will have the highest variance as OS context switching and other applications using the GPU will show up in this method.

Greg Simth

GPU: architecture and programming (NYU Course)

http://cs.nyu.edu/courses/spring12/CSCI-GA.3033-012/index.html

Contains some interesting links for GPU tools from webpage above
Multi2Sim Simulation Framework
http://www.multi2sim.org/

GPUocelot
http://code.google.com/p/gpuocelot/
Dynamic Compilation for PTX

Short CUDA tutorial of Colorado School of Mines
http://geco.mines.edu/tesla/cuda_tutorial_mio/index.html