![]() To statically allocate shared memory we use: _shared_ float sdata In our example, these parameters are the coefficient of the convolution kernel. The best practice is to use the shared memory for parameters that remain constant during the execution of the CUDA kernel and used in multiple calculations. Shared memory is allocated per thread block, so all threads in the block have access to the same shared memory. It is important to say that the kernel_size must be an odd number so that its center has an integer value.įor each block thread, the memory location of the corresponding pixel can be calculated by: int idx = iy*Nx ix We can define the center of the convolution kernel as it will be used in differnet calculations: //center of kernel in both dimensions The offset (kernel_size - 1)/2 is added to the iy, ix variables as the convolution will not be computed for the image pixels lying at the boundary layers of the original image (computations are performed only when the discrete filter kernel lies completely within the original image). Int ix = threadIdx.x (kernel_size - 1)/2 each thread is assigned to a pixel of a row, ix integer index of x Int iy = blockIdx.x (kernel_size - 1)/2 Thus, the index of a pixel in the image can be defined throught: //each block is assigned to a row of an image, iy integer index of y In this example, each block of threads will compute a row of the output image and each block thread will compute a single pixel value on this row. These variables are threadIdx.x and blockIdx.x respectively. Thus, a mapping mechanism is needed for each thread to compute a specific pixel of the output image and store the result to the corresponding memory location.ĬUDA kernels have access to device variables identifying both the thread index within the block and the block index. The CUDA kernel will be executed by each thread. Where *img is a pointer to the original image vector, *kernel is a pointer to the convolution kernel vector, *imgf is a pointer to the convoluted image, Nx and Ny are the dimensions of both the original and convoluted image, and kernel_size is the dimension of the convolution kernel. The CUDA kernel is already defined: void conv_img_cpu(float *img, float *kernel, float *imgf, int Nx, int Ny, int kernel_size) Open the source file LoG_gpu_exercise.cu with your favorite editor (e.g. Task 2: Following the steps 1 to 3 provided bellow write a CUDA kernel for the computation of the convolution operator. Salloc -reservation=hpcschool-gpu -p interactive -C gpu -ntasks-per-node 1 -c7 -G 1Ī CUDA kernel for the Convolution Operator or using the HPC School reservation 'hpcschool-gpu' # salloc -p interactive -qos debug -C gpu -c7 -G 1 -mem-per-cpu 27000 # /!\ warning: append -G 1 to really reserve a GPU Parameters: float *img, float *kernel, float *imgf, int Nx, int Ny, int kernel_size The convolution operator is calculated at each iteration for each image pixel using the double sum provided in the equation above. Where F is the original image, H is the convolution kernel and G is the resulted image.Ī serial code implementing the image convolution on a CPU employs two loops to compute the values of the pixels of the output image. Used for edge detection and noise detection You should have followed the Introduction to GPU programming with CUDA Laplacian of Gaussian (LoG): A convolution kernel for edge detectionĭerivative Filter used to find rapid changes in signals and especially images Now you'll need to pull the latest changes in your working copy of the ULHPC/tutorials you should have cloned in ~/git//ULHPC/tutorials (see "preliminaries" tutorial) (access)$> cd ~/git//ULHPC/tutorials ![]() ![]() (access)$> cp /etc/dotfiles.d/screen/.screenrc ~/ # always work within an GNU Screen session named with 'screen -S ' (Adapt accordingly) # /!\ Advanced (but recommended) best-practice: ![]() In particular, recall that the module command is not available on the access frontends. This tutorial will cover the following aspects of CUDA programming:Įnsure you are able to connect to the UL HPC clusters. Image Convolution with GPU and CUDA Copyright (c) 2020-2021 L.
0 Comments
Leave a Reply. |