public static void Sobel(cudaTextureObject_t texObj, cudaSurfaceObject_t surfObj, int width, int height) { for (int i = threadIdx.y + blockIdx.y * blockDim.y; i < height; i += blockDim.y * gridDim.y) { for (int j = threadIdx.x + blockIdx.x * blockDim.x; j < width; j += blockDim.x * gridDim.x) { int pixelId = i * width + j; if (i != 0 && j != 0 && i != height - 1 && j != width - 1) { float tl = TextureHelpers.tex2D(texObj, j - 1.0F, i - 1.0F); float t = TextureHelpers.tex2D(texObj, j - 1.0F, i); float tr = TextureHelpers.tex2D(texObj, j - 1.0F, i + 1.0F); float l = TextureHelpers.tex2D(texObj, j, i - 1.0F); float r = TextureHelpers.tex2D(texObj, j, i + 1.0F); float br = TextureHelpers.tex2D(texObj, j + 1.0F, i + 1.0F); float bl = TextureHelpers.tex2D(texObj, j + 1.0F, i - 1.0F); float b = TextureHelpers.tex2D(texObj, j + 1.0F, i); float data = (Math.Abs((tl + 2.0F * l + bl - tr - 2.0F * r - br)) + Math.Abs((tl + 2.0F * t + tr - bl - 2.0F * b - br))); if (data > 255) { data = 255; } TextureHelpers.surf2Dwrite(data, surfObj, j * sizeof(float), i); } else { TextureHelpers.surf2Dwrite(0.0F, surfObj, j * sizeof(float), i); } } } }
public static void Sobel(cudaTextureObject_t texObj, float[] output, int width, int height) //make output to be a surface { for (int i = threadIdx.y + blockIdx.y * blockDim.y; i < height; i += blockDim.y * gridDim.y) { for (int j = threadIdx.x + blockIdx.x * blockDim.x; j < width; j += blockDim.x * gridDim.x) { int pixelId = i * width + j; if (i != 0 && j != 0 && i != height - 1 && j != width - 1) { float tl = TextureHelpers.tex2D(texObj, j - 1.0F, i - 1.0F); float t = TextureHelpers.tex2D(texObj, j - 1.0F, i); float tr = TextureHelpers.tex2D(texObj, j - 1.0F, i + 1.0F); float l = TextureHelpers.tex2D(texObj, j, i - 1.0F); float r = TextureHelpers.tex2D(texObj, j, i + 1.0F); float br = TextureHelpers.tex2D(texObj, j + 1.0F, i + 1.0F); float bl = TextureHelpers.tex2D(texObj, j + 1.0F, i - 1.0F); float b = TextureHelpers.tex2D(texObj, j + 1.0F, i); float data = (Math.Abs((tl + 2.0F * l + bl - tr - 2.0F * r - br)) + Math.Abs((tl + 2.0F * t + tr - bl - 2.0F * b - br))); if (data > 255) { data = 255; } output[pixelId] = data; } else { output[pixelId] = 0; } } } }
static void Main(string[] args) { HybRunner runner = HybRunner.Cuda().SetDistrib(32, 32, 16, 16, 1, 0); GrayBitmap image = GrayBitmap.Load("../../images/lena512.bmp"); uint height = image.Height, width = image.Width; ushort[] inputPixels = image.PixelsUShort; float[] imageFloat = new float[width * height]; float[] imageCompute = new float[width * height]; for (int i = 0; i < width * height; ++i) { imageFloat[i] = (float)inputPixels[i]; } IntPtr src = runner.Marshaller.MarshalManagedToNative(imageFloat); //bind texture cudaChannelFormatDesc channelDescTex = TextureHelpers.cudaCreateChannelDesc(32, 0, 0, 0, cudaChannelFormatKind.cudaChannelFormatKindFloat); cudaArray_t cuArrayTex = TextureHelpers.CreateCudaArray(channelDescTex, src, (int)width, (int)height); cudaResourceDesc resDescTex = TextureHelpers.CreateCudaResourceDesc(cuArrayTex); //create Texture descriptor cudaTextureDesc texDesc = TextureHelpers.CreateCudaTextureDesc(); //create Texture object cudaTextureObject_t texObj; cuda.CreateTextureObject(out texObj, ref resDescTex, ref texDesc); //create and bind surface dynamic wrapper = runner.Wrap(new Program()); wrapper.Sobel(texObj, imageCompute, (int)width, (int)height); ushort[] outputPixel = new ushort[width * height]; for (int i = 0; i < width * height; ++i) { outputPixel[i] = (ushort)imageCompute[i]; } GrayBitmap imageSobel = new GrayBitmap(width, height); imageSobel.PixelsUShort = outputPixel; imageSobel.Save("../../output-03-surface/sobel.bmp"); }
static void Main(string[] args) { HybRunner runner = HybRunner.Cuda().SetDistrib(32, 32, 16, 16, 1, 0); GrayBitmap image = GrayBitmap.Load("../../images/lena512.bmp"); uint height = image.Height, width = image.Width; ushort[] inputPixels = image.PixelsUShort; float[] imageFloat = new float[width * height]; for (int i = 0; i < width * height; ++i) { imageFloat[i] = (float)inputPixels[i]; } IntPtr src = runner.Marshaller.MarshalManagedToNative(imageFloat); //bind texture cudaChannelFormatDesc channelDescTex = TextureHelpers.cudaCreateChannelDesc(32, 0, 0, 0, cudaChannelFormatKind.cudaChannelFormatKindFloat); cudaArray_t cuArrayTex = TextureHelpers.CreateCudaArray(channelDescTex, src, (int)width, (int)height); cudaResourceDesc resDescTex = TextureHelpers.CreateCudaResourceDesc(cuArrayTex); //create Texture descriptor cudaTextureDesc texDesc = TextureHelpers.CreateCudaTextureDesc(); //create Texture object cudaTextureObject_t texObj; cuda.CreateTextureObject(out texObj, ref resDescTex, ref texDesc); //bind surface cudaChannelFormatDesc channelDescSurf = TextureHelpers.cudaCreateChannelDesc(32, 0, 0, 0, cudaChannelFormatKind.cudaChannelFormatKindFloat); cudaArray_t cuArraySurf; cuda.MallocArray(out cuArraySurf, ref channelDescSurf, width, height, cudaMallocArrayFlags.cudaArraySurfaceLoadStore); //create cudaResourceDesc for surface cudaResourceDesc resDescSurf = TextureHelpers.CreateCudaResourceDesc(cuArraySurf); //create surface object cudaSurfaceObject_t surfObj; cuda.CreateSurfaceObject(out surfObj, ref resDescSurf); dynamic wrapper = runner.Wrap(new Program()); wrapper.Sobel(texObj, surfObj, (int)width, (int)height); //pinned float array to allow the copy of the surface object on the host float[] imageCompute = new float[width * height]; GCHandle handle = GCHandle.Alloc(imageCompute, GCHandleType.Pinned); IntPtr dest = handle.AddrOfPinnedObject(); cuda.MemcpyFromArray(dest, cuArraySurf, 0, 0, width * height * sizeof(float), cudaMemcpyKind.cudaMemcpyDeviceToHost); ushort[] outputPixel = new ushort[width * height]; for (int i = 0; i < width * height; ++i) { outputPixel[i] = (ushort)imageCompute[i]; } GrayBitmap imageSobel = new GrayBitmap(width, height); imageSobel.PixelsUShort = outputPixel; imageSobel.Save("../../output-3-surface/sobel.bmp"); }