Orthogonal projection of some primitives (Sphere, Torus) using Pycuda based ray casting. Ported from - http://code.finkjensen.be/2011/11/ray-marching-on-implicit-surfaces-with-cuda-cc-and-visual-c/

pycuda_ray_cast.py

'''
Created on 09/10/2013

PyCUDA api usage for ray-casting

@author: Tisham
'''

import pycuda.driver as drv
import pycuda.tools
import pycuda.autoinit
import numpy
from pycuda.compiler import SourceModule
import pycuda.gpuarray as gpuarray
import pycuda.cumath
from pycuda.elementwise import ElementwiseKernel
from pylab import imshow,show

(free,total)=drv.mem_get_info()
print("Global memory occupancy:%f%% free"%(free*100/total))

for devicenum in range(drv.Device.count()):
    device=drv.Device(devicenum)
    attrs=device.get_attributes()

    #Beyond this point is just pretty printing
    print("\n===Attributes for device %d"%devicenum)
    for (key,value) in attrs.iteritems():
        print("%s:%s"%(str(key),str(value)))

mod = SourceModule("""
#define XRES 1024
#define YRES 1024

__device__ float radians(float a)
{
  return 0.017453292 * a;
}

__device__ float vecLen(float2 a){
    return sqrt(a.x*a.x+a.y*a.y);
}

__device__ float3 vecAdd(float3 a, float3 b){
    return make_float3((b.x+a.x),(b.y+a.y),(b.z+a.z));
}

__device__ float3 scalProd(float s,float3 v)
{
    return make_float3(s*v.x,s*v.y,s*v.z);
}


__device__ float3 rotXAxis(float3 p,float angle)
{
    float rad_a = radians(angle);
    return make_float3(p.x,p.y*cos(rad_a) - p.z*sin(rad_a),p.z*cos(rad_a) + p.y*sin(rad_a));
}


__device__ float3 rotZAxis(float3 p,float angle)
{
    float rad_a = radians(angle);
    return make_float3(p.x*cos(rad_a) - p.y*sin(rad_a),p.y*cos(rad_a) + p.x*sin(rad_a),p.z);
}

__device__ float torus(float3 p, float2 t)

{

    float2 q = make_float2(vecLen(make_float2(p.x, p.z))-t.x, p.y);
    return vecLen(q) - t.y;

}

// the point p is entered as its (x,y,z) coordinates...
__device__ float sphere(float x,float y,float z,float size)

{

    float length = sqrt(x * x + y * y + z * z);

    return length - size;

}

__device__ float distanceEstimator(float3 p)

{
    float d1 = torus(rotZAxis(rotXAxis(p, 45),45), 

                                   make_float2(200,80));

    float3 p2 = vecAdd(make_float3(50,0,200), p);

    float d2 = torus(rotZAxis(rotXAxis(p2, 75),75),

                                   make_float2(200, 80));

    float3 p3 = vecAdd(make_float3(200, 300, 0), p);
    
    float3 p4 = vecAdd(make_float3(-300,-200, 0), p);
    
    float3 p5 = vecAdd(make_float3(350,-150, 0), p);

    float d3 = sphere(p3.x, p3.y, p3.z, 50);

    float d4 = sphere(p4.x, p4.y, p4.z, 80);

    float d5 = sphere(p5.x, p5.y, p5.z, 65);

    return min(d5, min(d4, min(d3 ,min(d1, d2))));

}

__device__ int trace(float3 from, float3 direction)

{

    int MaxRaySteps = 100;

    float minDistance = 0.0007;

    float totalDistance = 0.0;

    int steps = 0;

    for (steps = 0; steps < MaxRaySteps; steps++)

    {

        float3 p = vecAdd(from, 

                     scalProd(totalDistance, direction));

        float distance = distanceEstimator(p);

        totalDistance += distance;

        if (distance < minDistance)

            break;

    }

    float v = 1.0 - float(steps)/float(MaxRaySteps);

    return (int)(v*255);
}

__global__ void kernel( unsigned int *ptr )

{

    int x = threadIdx.x + blockIdx.x * blockDim.x;

    int y = threadIdx.y + blockIdx.y * blockDim.y;

    int offset = x + y * blockDim.x * gridDim.x;

    float ox = (float)(x - XRES/2);

    float oy = (float)(y - YRES/2);

 

    float3 from = make_float3(ox, oy, -1500.0);

    float3 direction = make_float3(0.0, 0.0, 1.0);

 

    int colr = trace(from, direction);

    ptr[offset] = colr;

}
""")


kernel = mod.get_function("kernel")

a = numpy.zeros((1024,1024),dtype=numpy.uint32)
a_gpu = drv.mem_alloc(a.nbytes)

kernel(a_gpu,block=(32,32, 1),grid=(32,32))

a_trace = numpy.empty_like(a)
drv.memcpy_dtoh(a_trace,a_gpu)
 
print numpy.max(a_trace)
 
imshow(a_trace)
show()

Can you offer the VS2010 CUDA code for this code on:http://code.finkjensen.be/2011/11/ray-marching-on-implicit-surfaces-with-cuda-cc-and-visual-c/ ? the web site is invalid for me.... my email is : dongliang1005@163.com Thankyou !