protected override void OnLoad(System.EventArgs e)
{
try
{
Console.WriteLine(ClientSize.Width + " " + ClientSize.Height);
Console.WriteLine(GL.GetString(StringName.Version));
Console.WriteLine(GL.GetString(StringName.ShadingLanguageVersion));
GL.Viewport(0, 0, ClientSize.Width, ClientSize.Height);
VSync = VSyncMode.On;
GL.Enable(EnableCap.Texture2D);
GL.Enable(EnableCap.DepthTest);
GL.DepthFunc(DepthFunction.Less);
// Setup quad VAO
quadVAO = GL.GenVertexArray();
int quadVBO = GL.GenBuffer();
GL.BindVertexArray(quadVAO);
GL.BindBuffer(BufferTarget.ArrayBuffer, quadVBO);
GL.BufferData(BufferTarget.ArrayBuffer, quadVertices.Length * sizeof(float), quadVertices, BufferUsageHint.StaticDraw);
GL.EnableVertexAttribArray(0);
GL.VertexAttribPointer(0, 3, VertexAttribPointerType.Float, false, 5 * sizeof(float), 0);
GL.EnableVertexAttribArray(1);
GL.VertexAttribPointer(1, 2, VertexAttribPointerType.Float, false, 5 * sizeof(float), 3 * sizeof(float));
uchar4[] data = new uchar4[Width * Height];
textureID = GL.GenTexture();
GL.BindTexture(TextureTarget.Texture2D, textureID);
GL.TexImage2D(TextureTarget.Texture2D, 0, PixelInternalFormat.Rgba, Width, Height, 0, PixelFormat.Bgra, PixelType.UnsignedByte, data);
GL.TexParameter(TextureTarget.Texture2D, TextureParameterName.TextureWrapS, Convert.ToInt32(TextureWrapMode.Repeat));
GL.TexParameter(TextureTarget.Texture2D, TextureParameterName.TextureWrapT, Convert.ToInt32(TextureWrapMode.Repeat));
GL.TexParameter(TextureTarget.Texture2D, TextureParameterName.TextureMinFilter, Convert.ToInt32(TextureMinFilter.Linear));
GL.TexParameter(TextureTarget.Texture2D, TextureParameterName.TextureMagFilter, Convert.ToInt32(TextureMagFilter.Linear));
GL.BindTexture(TextureTarget.Texture2D, 0);
string vsource = File.ReadAllText(@"vertex.glsl");
string fsource = File.ReadAllText(@"fragment.glsl");
shaderProgram = LoadShaderProgram(vsource, fsource);
IntPtr resource;
cuda.ERROR_CHECK(cuda.GraphicsGLRegisterImage(out resource, (uint)textureID, (uint)GL_TEXTURE_MODE.GL_TEXTURE_2D, (uint)cudaGraphicsRegisterFlags.SurfaceLoadStore));
cuda.ERROR_CHECK(cuda.GraphicsMapResources(1, new IntPtr[1] {
resource
}, cudaStream_t.NO_STREAM));
cudaArray_t array;
cuda.ERROR_CHECK(cuda.GraphicsSubResourceGetMappedArray(out array, resource, 0, 0));
cudaChannelFormatDesc channelDescSurf = TextureHelpers.cudaCreateChannelDesc(32, 0, 0, 0, cudaChannelFormatKind.cudaChannelFormatKindUnsigned);
cudaResourceDesc resDescSurf = TextureHelpers.CreateCudaResourceDesc(array);
cuda.CreateSurfaceObject(out surface, ref resDescSurf);
}
catch (Exception ex)
{
Console.WriteLine(ex.Message);
}
}
public void TypeConversion()
{
float f = float.MaxValue / 2;
byte b = (byte)CLTypeConverter.Convert(typeof(byte), f);
float4 f4 = new float4(f);
uchar4 i4 = (uchar4)CLTypeConverter.Convert(typeof(uchar4), f4);
Assert.True(b == 128);
for (int i = 0; i < 4; i++)
{
byte s = i4[i];
Assert.True(s == 128);
}
}
public static void Save(string path, uchar4[] colors, int width, int height)
{
Bitmap bitmap = new Bitmap(width, height);
int index = 0;
for (int j = 0; j < height; ++j)
{
for (int i = 0; i < width; ++i, ++index)
{
uchar4 pixel = colors[j * width + i];
bitmap.SetPixel(i, j, Color.FromArgb(pixel.x, pixel.y, pixel.z));
}
}
bitmap.Save(path, ImageFormat.Png);
}
public static extern CUResult cuMemcpyDtoH_v2(ref uchar4 dstHost, CUdeviceptr srcDevice, SizeT ByteCount);
public static void surf2Dwrite(uchar4 data, cudaSurfaceObject_t surfObj, int x, int y)
{
}
static void Main(string[] args)
{
// init CUDA
IntPtr d;
cuda.Malloc(out d, sizeof(int));
cuda.Free(d);
HybRunner runner = HybRunner.Cuda();
cudaDeviceProp prop;
cuda.GetDeviceProperties(out prop, 0);
dynamic wrapped = runner.Wrap(new Program());
runner.saveAssembly();
cudaStream_t stream;
cuda.StreamCreate(out stream);
NppStreamContext context = new NppStreamContext
{
hStream = stream,
nCudaDevAttrComputeCapabilityMajor = prop.major,
nCudaDevAttrComputeCapabilityMinor = prop.minor,
nCudaDeviceId = 0,
nMaxThreadsPerBlock = prop.maxThreadsPerBlock,
nMaxThreadsPerMultiProcessor = prop.maxThreadsPerMultiProcessor,
nMultiProcessorCount = prop.multiProcessorCount,
nSharedMemPerBlock = 0
};
Random rand = new Random();
using (NPPImage input = NPPImage.Load(inputFileName, stream))
{
uchar4[] output = new uchar4[input.width * input.height];
IntPtr d_output;
cuda.Malloc(out d_output, input.width * input.height * 4 * sizeof(byte));
// working area
IntPtr oDeviceDst32u;
size_t oDeviceDst32uPitch;
cuda.ERROR_CHECK(cuda.MallocPitch(out oDeviceDst32u, out oDeviceDst32uPitch, input.width * sizeof(int), input.height));
IntPtr segments;
size_t segmentsPitch;
cuda.ERROR_CHECK(cuda.MallocPitch(out segments, out segmentsPitch, input.width * sizeof(ushort), input.height));
NppiSize oSizeROI = new NppiSize {
width = input.width, height = input.height
};
int nBufferSize = 0;
IntPtr pScratchBufferNPP1, pScratchBufferNPP2;
// compute maximum label
NPPI.ERROR_CHECK(NPPI.LabelMarkersGetBufferSize_16u_C1R(oSizeROI, out nBufferSize));
cuda.ERROR_CHECK(cuda.Malloc(out pScratchBufferNPP1, nBufferSize));
int maxLabel;
NPPI.ERROR_CHECK(NPPI.LabelMarkers_16u_C1IR_Ctx(input.deviceData, input.pitch, oSizeROI, 165, NppiNorm.nppiNormInf, out maxLabel, pScratchBufferNPP1, context));
// compress labels
NPPI.ERROR_CHECK(NPPI.CompressMarkerLabelsGetBufferSize_16u_C1R(maxLabel, out nBufferSize));
cuda.ERROR_CHECK(cuda.Malloc(out pScratchBufferNPP2, nBufferSize));
NPPI.ERROR_CHECK(NPPI.CompressMarkerLabels_16u_C1IR_Ctx(input.deviceData, input.pitch, oSizeROI, maxLabel, out maxLabel, pScratchBufferNPP2, context));
uchar4[] colormap = new uchar4[maxLabel + 1];
for (int i = 0; i <= maxLabel; ++i)
{
colormap[i] = new uchar4 {
x = (byte)(rand.Next() % 256), y = (byte)(rand.Next() % 256), z = (byte)(rand.Next() % 256), w = 0
};
}
IntPtr d_colormap;
cuda.Malloc(out d_colormap, (maxLabel + 1) * 4 * sizeof(byte));
var handle = GCHandle.Alloc(colormap, GCHandleType.Pinned);
cuda.Memcpy(d_colormap, handle.AddrOfPinnedObject(), (maxLabel + 1) * 4 * sizeof(byte), cudaMemcpyKind.cudaMemcpyHostToDevice);
handle.Free();
NPP_ImageSegmentationx46Programx46ColorizeLabels_ExternCWrapperStream_CUDA(
8 * prop.multiProcessorCount, 1, 256, 1, 1, 0, stream, // cuda configuration
input.deviceData, d_output, d_colormap, maxLabel + 1, input.pitch * input.height / sizeof(ushort), input.width, input.pitch / sizeof(ushort));
handle = GCHandle.Alloc(output, GCHandleType.Pinned);
cuda.Memcpy(handle.AddrOfPinnedObject(), d_output, input.width * input.height * sizeof(byte) * 4, cudaMemcpyKind.cudaMemcpyDeviceToHost);
handle.Free();
NPPImage.Save(segmentsFileName, output, input.width, input.height);
Process.Start(segmentsFileName);
cuda.ERROR_CHECK(cuda.Free(oDeviceDst32u));
cuda.ERROR_CHECK(cuda.Free(segments));
cuda.ERROR_CHECK(cuda.Free(pScratchBufferNPP1));
cuda.ERROR_CHECK(cuda.Free(pScratchBufferNPP2));
}
}
public void Render(cudaSurfaceObject_t surface, int cWidth, int cHeight, int iterations, float3 viewDirection, float3 nearFieldLocation, float3 eyeLocation, float3 lightDirection)
{
int cHalfWidth = cWidth / 2;
float pixel = ((float)DEPTH_OF_FIELD) / (((cHeight + cWidth) / 2.0F));
float halfPixel = pixel / 2.0F;
float smallStep = 0.01F;
float bigStep = 0.02F;
for (int y = threadIdx.y + blockIdx.y * blockDim.y; y < cHeight; y += blockDim.y * gridDim.y)
{
var ny = y - cHeight / 2;
float3 tempViewDirectionY = viewDirection;
float3 tempViewDirectionX1 = viewDirection;
MathFunctions.turnOrthogonal(ref tempViewDirectionY);
MathFunctions.crossProduct(ref tempViewDirectionY, viewDirection);
tempViewDirectionY = tempViewDirectionY * (ny * pixel);
MathFunctions.turnOrthogonal(ref tempViewDirectionX1);
for (int x = threadIdx.x + blockDim.x * blockIdx.x; x < cWidth; x += blockDim.x * gridDim.x)
{
int nx = x - cHalfWidth;
float3 pixelLocation = nearFieldLocation;
float3 tempViewDirectionX2 = tempViewDirectionX1 * (nx * pixel);
pixelLocation = pixelLocation + tempViewDirectionX2 + tempViewDirectionY;
float3 rayLocation = pixelLocation;
float3 rayDirection = rayLocation - eyeLocation;
MathFunctions.normalize(ref rayDirection);
var distanceFromCamera = 0.0;
var d = Distance(rayLocation, iterations);
int i = 0;
for (; i < MAX_ITER; i++)
{
if (d < halfPixel)
{
break;
}
//Increase rayLocation with direction and d:
rayDirection = rayDirection * d;
rayLocation = rayLocation + rayDirection;
//And reset ray direction:
MathFunctions.normalize(ref rayDirection);
//Move the pixel location:
distanceFromCamera = MathFunctions.length(nearFieldLocation - rayLocation);
if (distanceFromCamera > DEPTH_OF_FIELD)
{
break;
}
d = Distance(rayLocation, iterations);
}
if (distanceFromCamera < DEPTH_OF_FIELD && distanceFromCamera > 0)
{
rayLocation.x -= smallStep;
var locationMinX = Distance(rayLocation, iterations);
rayLocation.x += bigStep;
var locationPlusX = Distance(rayLocation, iterations);
rayLocation.x -= smallStep;
rayLocation.y -= smallStep;
var locationMinY = Distance(rayLocation, iterations);
rayLocation.y += bigStep;
var locationPlusY = Distance(rayLocation, iterations);
rayLocation.y -= smallStep;
rayLocation.z -= smallStep;
var locationMinZ = Distance(rayLocation, iterations);
rayLocation.z += bigStep;
var locationPlusZ = Distance(rayLocation, iterations);
rayLocation.z -= smallStep;
float3 normal = new float3();
//Calculate the normal:
normal.x = (locationMinX - locationPlusX);
normal.y = (locationMinY - locationPlusY);
normal.z = (locationMinZ - locationPlusZ);
MathFunctions.normalize(ref normal);
//Calculate the ambient light:
var dotNL = MathFunctions.dotProduct(lightDirection, normal);
var diff = MathFunctions.saturate(dotNL);
//Calculate specular light:
float3 halfway = rayDirection + lightDirection;
MathFunctions.normalize(ref halfway);
var dotNH = MathFunctions.dotProduct(halfway, normal);
float s = MathFunctions.saturate(dotNH);
float s2 = s * s;
float s4 = s2 * s2;
float s8 = s4 * s4;
float s16 = s8 * s8;
float s32 = s16 * s16;
float spec = s * s2 * s32; // s^35
var shad = shadow(lightDirection, rayLocation, iterations, 1.0F, DEPTH_OF_FIELD, 16.0F) + 0.25F;
var brightness = (10.0F + (200.0F + spec * 45.0F) * shad * diff) / 270.0F;
var red = 10 + (380 * brightness);
var green = 10 + (280 * brightness);
var blue = (180 * brightness);
red = MathFunctions.clamp(red, 0, 255.0F);
green = MathFunctions.clamp(green, 0, 255.0F);
blue = MathFunctions.clamp(blue, 0, 255.0F);
uchar4 color = new uchar4();
color.x = (byte)red;
color.y = (byte)green;
color.z = (byte)blue;
TextureHelpers.surf2Dwrite(color, surface, x * sizeof(int), y);
}
else
{
uchar4 color = new uchar4();
color.x = (byte)(155.0F + MathFunctions.clamp(i * 1.5F, 0.0F, 100.0F));
color.y = ((byte)(205.0F + MathFunctions.clamp(i * 1.5F, 0.0F, 50.0F)));
color.z = 255;
TextureHelpers.surf2Dwrite(color, surface, x * sizeof(int), y);
}
}
}
}
private void Image_MouseMove(object sender, MouseEventArgs e)
{
System.Windows.Point position = GetPositionWithDpi(e);
int pX = (int)position.X;
int pY = (int)position.Y;
Point pixel = GetImagePixelFromMouseCoordinate(position);
uchar4[] p = new uchar4[1];
if (!double.IsInfinity(pixel.X) && !double.IsInfinity(pixel.Y))
{
d3dimage.Lock();
_graphicsres.MapAllResources();
CudaArray2D arr = _graphicsres[0].GetMappedArray2D(0, 0);
CUDAMemCpy2D copy = new CUDAMemCpy2D();
GCHandle handle = GCHandle.Alloc(p, GCHandleType.Pinned);
copy.dstHost = handle.AddrOfPinnedObject();
copy.srcArray = arr.CUArray;
copy.srcMemoryType = CUMemoryType.Array;
copy.dstMemoryType = CUMemoryType.Host;
copy.Height = 1;
copy.WidthInBytes = 4;
copy.srcXInBytes = (int)pixel.X * 4;
copy.srcY = (int)pixel.Y;
arr.CopyData(copy);
_graphicsres.UnmapAllResources();
arr.Dispose();
handle.Free();
d3dimage.Unlock();
}
SetValue(ColorOfPixelProperty, Color.FromArgb(p[0].w, p[0].z, p[0].y, p[0].x));
SetValue(PixelCoordinateProperty, pixel);
if (_clicked)
{
_viewShiftX += (-_lastX + pX) / _scaleFac * _projFac;
_viewShiftY += (-_lastY + pY) / _scaleFac * _projFac;
_lastX = pX;
_lastY = pY;
SlimDX.Matrix matScale = SlimDX.Matrix.Scaling(_scaleFac, _scaleFac, 1);
float shiftScale = Math.Min((float)ActualWidthDpi, (float)ActualHeightDpi);
SlimDX.Matrix matTrans = SlimDX.Matrix.Translation(_viewShiftX / shiftScale * _scaleFac, _viewShiftY / shiftScale * _scaleFac, 0);
SlimDX.Matrix mat = matScale * matTrans;
SlimDX.Matrix rotMat = new SlimDX.Matrix();
switch (_rotation)
{
case Rotation._0:
rotMat = SlimDX.Matrix.RotationZ(0);
break;
case Rotation._90:
rotMat = SlimDX.Matrix.RotationZ((float)(90.0 / 180.0 * Math.PI));
break;
case Rotation._180:
rotMat = SlimDX.Matrix.RotationZ((float)(180.0 / 180.0 * Math.PI));
break;
case Rotation._270:
rotMat = SlimDX.Matrix.RotationZ((float)(270.0 / 180.0 * Math.PI));
break;
}
_device.SetTransform(TransformState.View, rotMat * mat);
updateFrame();
}
}
