TORCH_CHECK(x.is_contiguous(), #x " must be contiguous") #define CHECK_CUDA(x) TORCH_CHECK(x.is_cuda(), #x " must be a CUDA tensor") It also resets CUDA’s default device when destructed (i.e. The at::DeviceGuard sets CUDA’s default device to the desired one (if you have multiple GPUs) such that subsequent calls to cudaGetDevice and at::cuda::getCurrentCUDAStream can return the correct values. Additional checks on tensors shapes / constraints on function arguments can be added but are omitted here. Typical CUDA checks simply verify that input tensors are on CUDA and contiguous.
DIM3 GRID CALCULATION CODE
Source code 3.1 CUDA Input Checks + Device Guard cpp is to check input validity, set default CUDA device and create Python binding. upfirdn2d.py: Python API and Pytorch autograd.Function definition.upfirdn2d_kernel.cu: Actual CUDA kernel.Having understood the purpose of the upfirdn2d kernel, let’s look at the relevant files: In the upfirdn2d case, upsampling was done before the convolution but since the upsampling only involes zero padding, the blurring kernel still sort of removes the resulting high-frequency noise. Iīelieve it sort of achieves the same purpose as applying a blurring kernel before bilinear upsampling, which is to avoid aliasing using a low-pass filter before upsampling see this. In StyleGAN2, up/downsampling operations use a separable 2D Gaussian kernel ( ). downsampling (by dropping zeros after each pixel).FIR filter (by convolution with a kernel).padding (by adding zeros on each side of the image).upsampling (by adding zeros after each pixel).image in StyleGAN2) goes through the following steps: The upfirdn2d is a very common signal processing operation where a signal (i.e. The former, on the other hand, is used for up/downsampling (by replacing the bilinear up/downsampling in the original StyleGAN paper). The latter is very simple as it is just a leaky ReLU that comes with a bias term. The StyleGAN2 team implemented two custom kernels: upfirdn2d and used_bias_act. Hence, I took the time and effort to investigate its inner workings, simplify it and summarize my learnings here. Sadly, the source code was not properly commented nor systematically explained anywhere. I wanted to understand the implementation of the CUDA kernels and the design considerations that went into the process. In configuration E at $1024^2$ resolution, our optimizations improved the overall training time by about 30% and memory footprint by about 20%. We implemented filtered up/downsampling as a single fused operation, and bias and activation as another one. This motivated us to optimize these operations using hand-written CUDA kernels. In the StyleGAN2 paper, the team mentioned that they implemented custom CUDA kernels to speed up the training time. I found the latter more stable during training and much simpler (such that I can incorporate it in my project repo more easily). I experimented with both NVIDIA’s official implementation and rosinality’s implementation. I would be clear where the configuration of the threads has been defined, and the 1D, 2D and 3D access pattern depends on how you are interpreting your data and also how you are accessing them by 1D, 2D and 3D blocks of threads.During my final year project in university, I needed to train a StyleGAN2 model for X-ray image synthesis. To sumup, it does it matter if you use a dim3 structure. Int y = blockIdx.y * blockDim.y + threadIdx.y īecause blockIdx.y and threadIdx.y will be zero. So, in both cases: dim3 blockDims(512) and myKernel>(.) you will always have access to threadIdx.y and threadIdx.z.Īs the thread ids start at zero, you can calculate a memory position as a row major order using also the ydimension: int x = blockIdx.x * blockDim.x + threadIdx.x The same happens for the blocks and the grid. When defining a variable of type dim3, any component left unspecified is initialized to 1. However, the access pattern depends on how you are interpreting your data and also how you are accessing them by 1D, 2D and 3D blocks of threads.ĭim3 is an integer vector type based on uint3 that is used to specify dimensions. The memory is always a 1D continuous space of bytes. The way you arrange the data in memory is independently on how you would configure the threads of your kernel.
0 kommentar(er)
