CUDA / OpenCL ready Solid Classes

README.md

The Saint on Porting C++ Classes to CUDA with Unified Memory | Parallel Forall

codesketch.cpp

// skeleton

// Porting Solid Classes to OpenCl

// Porting Solid Classes to CUDA

demo.cu

#include <cuda.h>
#include <cuda_runtime.h>
#include <stdio.h>
#include <assert.h>
#include "simpleSolid.cu"

#define dim 64
// interface for CPU
__global__ void foo(Coordinate singleton){
    // utilize shared memory etc ..
    __shared__ SimpleSolid a;
    if (threadIdx.x == 0) {
        // Coordinate singleton(1.0, 1.0, 1.0);
        a.setCoordinate(singleton);
    }
    __syncthreads();
    Coordinate atom;
    // make a good distribution
    atom.x = threadIdx.x - (float) dim/2;
    atom.y = threadIdx.x - (float) dim/2;
    atom.z = threadIdx.x - (float) dim/2;
    printf("hello from octant %d, with distance %g \
        \t(thread %d: %g %g %g)\n",
        a.getOctant(atom), a.getDistance(atom), 
        blockIdx.x*blockDim.x+threadIdx.x, 
        atom.x, atom.y, atom.z);
};

int main(int argc, char const *argv[])
{
    Coordinate singleton;
    singleton.x = 0.f;
    singleton.y = 0.f;
    singleton.z = 0.f;
    foo<<<dim,dim>>>(singleton);
    cudaDeviceSynchronize();
    // foo();
    return 0;
}

Makefile

all:
	nvcc -gencode arch=compute_20,code=sm_20 GPU_solid.cu

simpleSolidClass_GPU.cu

typedef struct
{
    float x;
    float y;
    float z;
} Coordinate;

enum octant {first = 1, second, third, fourth,
    fifth, sixth, seventh, eighth};

// SimpleSolid is a dot for demo purpose only
    class SimpleSolid
    {
    public:
        __device__ SimpleSolid()
        {
            coor.x = 0.f;
            coor.y = 0.f;
            coor.z = 0.f;
        }
        __device__ int setCoordinate(Coordinate a)
        {
            coor = a;
            return 0;
        }
        __device__ Coordinate getCoordinate(){return coor;}
        __device__ float getDistance(Coordinate b){
            float dx = powf(b.x-coor.x, 2.0f);
            float dy = powf(b.y-coor.y, 2.0f);
            float dz = powf(b.z-coor.z, 2.0f);
            return sqrtf (dx + dy + dz);
        }
        __device__ octant getOctant(Coordinate b){
            // I, IV, V, VIII
            if (b.x - coor.x > 0)
            {
                // I, V
                if (b.y - coor.y > 0)
                {
                    if (b.z - coor.z > 0)
                    {
                        return first;
                    }else{
                        return fifth;
                    }
                }else{
                    if (b.z - coor.z > 0)
                    {
                        return fourth;
                    }else{
                        return eighth;
                    }
                }
            }else{
            // II, III, VI, VII
                // II, VI
                if (b.y - coor.y > 0)
                {
                    if (b.z - coor.z > 0)
                    {
                        return second;
                    }else{
                        return sixth;
                    }
                }else{
                    if (b.z - coor.z > 0)
                    {
                        return third;
                    }else{
                        return seventh;
                    }
                }
            }
        }
    private:
    // float x, y, z;
        Coordinate coor;
    };