The Fashionable CUDA Toolbox in Follow: A Step-by-Step Optimization Walkthrough

NVIDIA CUDA stays the inspiration of GPU-accelerated computing, powering every thing from scientific simulations to large-scale AI coaching.

However writing appropriate, maintainable, and performant CUDA code may be difficult: reminiscence bugs cover in plain sight, efficiency bottlenecks may be invisible with out the best instrumentation, and hand-rolled GPU algorithms hardly ever match the effectivity of optimized libraries. Thankfully, the trendy CUDA toolchain has matured considerably, and lots of of those challenges now have simple options.

On this weblog submit, we’ll stroll via the instruments NVIDIA gives to debug, benchmark, and enhance your code. With solely small line modifications every time, we’re going to make the instance code safer, simpler to take care of, and sooner.

Throughout six incremental steps, this submit will cowl:

Methods to simply discover indexing bugs by adopting the trendy CCCL API and Compute Sanitizer

Methods to enhance Nsight Programs benchmarks with NVTX

Methods to use CUB’s optimized algorithms on the block and system degree

Methods to handle GPU reminiscence via pooled containers

Methods to pace up host-to-device transfers with pinned containers

Methods to parallelize GPU work by giving every thread its personal stream and asynchronous transfers

As a companion to this weblog submit, we offer the code and the choice to run on Google Colab.

Start line: A picture processing pipeline instance

From an enter stream of pink, inexperienced, and blue pictures, begin by transferring the info from the CPU to the GPU. Then convert these from RGB to grayscale.

Then, for every 32 by 32 pixel tile within the picture, compute the median by sorting the pixels and choosing the center worth. Lastly, copy the median of every tile again to the CPU.

The bottom code instance

Under is the complete beginning code. Every step on this submit improves on it.

#outline CUDA_CHECK_ERROR(name) do {
cudaError_t err = name;
if (err != cudaSuccess) {
std::cerr << “CUDA error in “ << __FILE__ << ” at line “ << __LINE__ << “: “
<< cudaGetErrorString(err) << std::endl;
std::exit(EXIT_FAILURE);
}
} whereas (0)
// Alias for a picture pixel
utilizing pixel_t = uint8_t;
// Kernel changing the pink, inexperienced and blue pictures right into a single grey picture
__global__ void computeRGBToGray(const pixel_t* d_image_r, const pixel_t* d_image_g, const pixel_t* d_image_b, pixel_t* d_image_gray, int width, int top) {
// Compute the thread international index within the grid
const int x = threadIdx.x + blockIdx.x * blockDim.x;
const int y = threadIdx.y + blockIdx.y * blockDim.y;
// Boundary test choosing solely threads inside the picture boundary
if (x < width && y < top) {
// Compute the thread index within the picture
const int i = x + y * width;
// Convert from rgb to grayscale and retailer the lead to international reminiscence
d_image_gray[i] = static_cast(0.299f * d_image_r[i] + 0.587f * d_image_g[i] + 0.114f * d_image_b[i]);
}
}
// Kernel computing the median of every tile within the grayscale picture
template <int TILE_WIDTH, int HISTO_SIZE>
__global__ void computeMedian(pixel_t *d_image_gray, pixel_t *d_median, int width, int top) {
// Compute the thread international index within the grid
const int x = threadIdx.x + blockIdx.x * blockDim.x;
const int y = threadIdx.y + blockIdx.y * blockDim.y;
// Boundary test choosing solely threads inside the picture boundary
if (!(x < width && y < top))
return;
// Allocate the shared reminiscence wherein we’ll retailer the tile
__shared__ pixel_t tile[TILE_WIDTH * TILE_WIDTH];
// Compute the thread index within the picture
const int index = x + y * width;
// Load the tile’s grayscale worth from international reminiscence into shared reminiscence
tile[index] = d_image_gray[index];
// Synchronize to verify all threads have loaded their knowledge
__syncthreads();
// Type the tile array utilizing a single threaded bubble type
if (threadIdx.x == 0 && threadIdx.y == 0) {
for (int i = 0; i < TILE_WIDTH * TILE_WIDTH; ++i)
for (int j = i + 1; j < TILE_WIDTH * TILE_WIDTH; ++j)
if (tile[i] > tile[j])
cuda::std::swap(tile[i], tile[j]);
// Every thread block shops the median, discovered within the center index after sorting, within the international median array
const int medianIndex = (TILE_WIDTH * TILE_WIDTH) / 2;
d_median[blockIdx.x + blockIdx.y * gridDim.x] = tile[medianIndex];
}
}
int essential() {
// Outline all the instance constants
constexpr auto TILE_WIDTH = 32;
constexpr auto HISTO_SIZE = 256;
constexpr auto NB_TILE_X = 250;
constexpr auto NB_TILE_Y = NB_TILE_X;
constexpr auto IMAGE_LENGTH = TILE_WIDTH * NB_TILE_X;
constexpr auto IMAGE_SIZE = IMAGE_LENGTH * IMAGE_LENGTH;
constexpr auto NB_IMAGES = 3;
constexpr auto INIT_VALUE = 4;
// Allocate the CPU reminiscence to retailer the photographs tiles medians and for the pink, inexperienced, blue and grayscale pictures
std::vector> h_images_r(NB_IMAGES, std::vector(IMAGE_SIZE, 4));
std::vector> h_images_g(NB_IMAGES, std::vector(IMAGE_SIZE, 4));
std::vector> h_images_b(NB_IMAGES, std::vector(IMAGE_SIZE, 4));
std::vector> h_images_gray(NB_IMAGES, std::vector(IMAGE_SIZE, 0));
std::vector> h_medians(NB_IMAGES, std::vector(NB_TILE_X * NB_TILE_Y));
// Run the picture processing pipeline for every picture, in parallel
#pragma omp parallel for
for (int i = 0; i < NB_IMAGES; ++i)
{
pixel_t *d_image_r, *d_image_g, *d_image_b, *d_image_gray, *d_median;
// Allocate the GPU reminiscence for every container
CUDA_CHECK_ERROR(cudaMalloc(&d_image_r, IMAGE_SIZE * sizeof(pixel_t)));
CUDA_CHECK_ERROR(cudaMalloc(&d_image_g, IMAGE_SIZE * sizeof(pixel_t)));
CUDA_CHECK_ERROR(cudaMalloc(&d_image_b, IMAGE_SIZE * sizeof(pixel_t)));
CUDA_CHECK_ERROR(cudaMalloc(&d_image_gray, IMAGE_SIZE * sizeof(pixel_t)));
CUDA_CHECK_ERROR(cudaMalloc(&d_median, (NB_TILE_X * NB_TILE_Y) * sizeof(pixel_t)));
// Copy the reminiscence of every container from CPU to GPU
CUDA_CHECK_ERROR(cudaMemcpy(d_image_r, h_images_r[i].knowledge(), IMAGE_SIZE * sizeof(pixel_t), cudaMemcpyHostToDevice));
CUDA_CHECK_ERROR(cudaMemcpy(d_image_g, h_images_g[i].knowledge(), IMAGE_SIZE * sizeof(pixel_t), cudaMemcpyHostToDevice));
CUDA_CHECK_ERROR(cudaMemcpy(d_image_b, h_images_b[i].knowledge(), IMAGE_SIZE * sizeof(pixel_t), cudaMemcpyHostToDevice));
// Launch a GPU kernel to transform the RGB pictures to grayscale
dim3 blockSize(TILE_WIDTH, TILE_WIDTH);
dim3 gridSize(cuda::ceil_div(IMAGE_LENGTH, blockSize.x), cuda::ceil_div(IMAGE_LENGTH, blockSize.y));
computeRGBToGray<<>>(d_image_r, d_image_g, d_image_b, d_image_gray, IMAGE_LENGTH, IMAGE_LENGTH);
CUDA_CHECK_ERROR(cudaGetLastError());
// Launch the GPU kernel to compute the median of each tile within the picture
computeMedian<<>>(d_image_gray, d_median, IMAGE_LENGTH, IMAGE_LENGTH);
CUDA_CHECK_ERROR(cudaGetLastError());
// Copy the GPU median reminiscence again to the CPU
CUDA_CHECK_ERROR(cudaMemcpy(h_medians[i].knowledge(), d_median, (NB_TILE_X * NB_TILE_Y) * sizeof(pixel_t), cudaMemcpyDeviceToHost));
// Free the GPU reminiscence
CUDA_CHECK_ERROR(cudaFree(d_image_r));
CUDA_CHECK_ERROR(cudaFree(d_image_g));
CUDA_CHECK_ERROR(cudaFree(d_image_b));
CUDA_CHECK_ERROR(cudaFree(d_image_gray));
CUDA_CHECK_ERROR(cudaFree(d_median));
}
return 0;
}

This code begins by defining two kernels:

computeRGBToGray masses the values of the pink, inexperienced, and blue enter pictures to transform and write them to the grayscale output picture.

computeMedian computes the median of every tile of an enter grayscale picture. Every thread block masses the tile from international reminiscence to shared reminiscence. Then, a single thread is used to type the array and write the worth discovered on the center index, akin to the median, within the international output median array.

In the principle, after defining the constants used for the instance, the CPU reminiscence is allotted for every picture and for the medians.

The picture processing pipeline is then run for every of the three pictures, in parallel, utilizing OpenMP. The pipeline begins by allocating the required reminiscence on the GPU earlier than transferring the info from the CPU to the GPU. The 2 kernels to transform from RGB to grayscale and to compute the median are launched afterward. Lastly, we copy again the median outcomes to the CPU earlier than releasing the reminiscence.

This code has a number of flaws that will likely be addressed, step-by-step.

1. Compute Sanitizer and the CCCL API: Discover bugs simply and write safer code

Let’s begin by working the code.

Whereas the code has some error checking in it, once you get an error message like this “unlawful reminiscence entry,” you must begin by utilizing compute-sanitizer to research additional.

Utilizing Compute Sanitizer, the NVIDIA useful correctness checking suite, we are able to immediately establish a bug that may be exhausting to identify:

Working the above immediately exhibits that the code suffers from an out-of-bound shared write at line 55 of 0_base_error_example.cu.

The road incorrectly masses knowledge in shared reminiscence, utilizing a world index. Since shared reminiscence is outlined on the thread block degree, we have to change the indexing. To keep away from indexing errors, a brand new API was launched in CCCL to differentiate international and block-level indexing. To make use of it, you first must launch your kernel utilizing the brand new cuda::launch API:

Then, utilizing the brand new indexing API contained in the kernel:

With out utilizing compute-sanitizer or the brand new API, this error might even have been immediately noticed by utilizing cuda::std::span or its n-dimensional variant cuda::std::mdspan as a substitute of uncooked pointers. cuda::std::span and cuda::std::mdspan are non-owning views over contiguous reminiscence and are helpful to summary the precise container away. Accessing knowledge via spans is safer than via uncooked pointers partly as a result of in debug mode, out-of-bounds accesses will set off an assertion.

The kernel must be up to date as:

If you happen to run the code with these modifications you’ll get one thing like the next:

Reminiscence accesses in shared reminiscence also needs to be protected by utilizing a cuda::shared_memory_mdspan as within the following snippet:

Now you’ll be able to run and every thing ought to execute correctly with out errors.

Utilizing the brand new launch API and its indexing mechanism, spans over uncooked pointers, and compute-sanitizer, out-of-bounds accesses both don’t occur or are caught straight away. For extra data on compute-sanitizer, see Environment friendly CUDA Debugging: Methods to Hunt Bugs with NVIDIA Compute Sanitizer.

2. Nsight Programs and NVTX: Benchmark your code correctly

Now the code is bug-free, it is able to be benchmarked utilizing NVIDIA Nsight Programs. It lets you visualize this system timeline: know when every operate is named and for a way lengthy.

To make the timeline visualization simpler, we wrap each attention-grabbing code part utilizing NVTX:

Which yields the next outcome:

Within the GPU {hardware} (CUDA HW) part of the profiler output in Determine 3, above, it’s reported that the GPU is especially busy with kernels (98.5% of the GPU time) whereas the reminiscence operations solely take 1.5% of the GPU time.

Of the 2 kernels, the one computing the medians is taking nearly all of the runtime with 2.1 seconds for every picture (see the yellow field on the best, the place the stats for computeMedian are proven, and the elapsed time is 2.142s).

From the CPU (thread) part, we see the picture computation takes 6.8 seconds in whole, with more often than not being spent on computing the medians for the three grayscale pictures.

We now know the primary operation to optimize with a view to have the best influence. For extra data on Nsight Programs, see Optimizing CUDA Reminiscence Transfers with NVIDIA Nsight Programs. For extra data on NVTX, see CUDA Professional Tip: Generate Customized Utility Profile Timelines with NVTX.

3. CUB: Specific algorithms immediately on the GPU

When coping with widespread algorithms, writing customized kernels is error-prone and can more than likely lead to an inefficient implementation. Each time doable, each for device-side patterns and for in-kernel primitives, it’s endorsed to make use of CUB.

CUB is the NVIDIA parallel algorithm library shipped via CCCL. It exposes extremely optimized routines at a number of granularities: device-wide (cub::System*), block-level (cub::Block*), and warp-level (cub::Warp*).

For the RGB to grayscale step, we are able to change the customized kernel with cub::DeviceTransform::Remodel. It applies a user-provided operate to a tuple of enter iterators and writes the outcome to an output iterator, executing on the GPU:

For the median, programming a parallel block-level type by hand is advanced and sluggish. As an alternative, we immediately leverage CUB’s block-level radix type contained in the kernel:

Following this modification, we benchmark once more utilizing Nsight Programs:

The time required to compute the median is now solely 773 microseconds (once more, take a look at the elapsed time within the computeMedian yellow pop-out picture), 2717x sooner. The general time to compute all three pictures is now 635 milliseconds, 10x sooner.

If we reassess the present bottleneck: the time spent on reminiscence allocations represents round 83% of the whole picture compute runtime.

This may be significantly improved.

4. Pooled Reminiscence Containers: Handy and sooner reminiscence administration

Allocating GPU reminiscence utilizing cudaMalloc can have surprising adverse results: leaks by forgetting to name cudaFree and dear reminiscence operations in crucial elements of your code.

As an alternative, we suggest utilizing CCCL’s asynchronous reminiscence containers, cuda::device_buffer. Like C++ std::vector, the reminiscence is robotically deallocated as soon as the container goes out of scope.

Moreover, a reminiscence pool backs the buffer, so repeated allocations and deallocations don’t pay the complete price of cudaMalloc / cudaFree each time.

To make use of the GPU reminiscence containers, we replace the code accordingly:

After this modification, we analyze the timeline once more:

The time spent on reminiscence allocation is now virtually nonexistent; the time it takes to compute a picture has improved by 2.6x.

The GPU time is now memory-dominated. Nearly on a regular basis to compute all pictures is spent on first copying the pink, inexperienced and blue, for the three pictures, from CPU to GPU.

It’s doable to significantly pace up the host-to-device reminiscence transfers.

5. Pinned reminiscence: Sooner host-to-device reminiscence transfers

CPU knowledge allocations are pageable by default, which the GPU can not entry immediately. The CUDA driver should first allocate a short lived page-locked, or pinned, host array, copy the host knowledge to the pinned array, after which switch the info from the pinned array to system reminiscence.

When it’s identified prematurely that CPU reminiscence will likely be copied to the GPU, it’s suggested to immediately allocate utilizing pinned reminiscence.

CCCL exposes a pinned-memory host container, cuda::host_buffer, that may be constructed via the cuda::make_pinned_buffer manufacturing facility:

Following these modifications, we are able to benchmark once more:

The time it takes to do the host-to-device reminiscence transfers has been considerably lowered; it now solely takes 25 ms to compute all pictures, 10x sooner.

For extra data on pinned reminiscence, see Methods to Optimize Information Transfers in CUDA C/C++.

One shocking conduct may need caught the attention of some readers because the starting:

Although we’re utilizing completely different CPU threads, all operations (reminiscence and kernels) are being executed sequentially on the GPU.

Let’s repair it.

6. Streams: Parallelize operations on the GPU

By default, all operations (kernels, reminiscence allocations, or transfers) are launched on what we name the default stream: it may be considered as a queue of duties the GPU must execute so as.

On this instance, we want a stream for every picture/thread. CCCL supplies cuda::stream, an proudly owning self-managed model of CUDA streams. It could possibly merely be constructed contained in the parallel for loop, so every OpenMP thread will get its personal queue of GPU work.

To effectively leverage streams, we additionally want to make use of the asynchronous API: every GPU operation launched by the CPU shouldn’t be waited upon till completion. To saturate the GPU, every CPU thread ought to launch as many operations as doable, as quick as doable, with out ready for them to first full. Kernels and CUB system calls are already asynchronous by default and are launched on the handed stream. To launch asynchronous copies between the host and the system, we use the brand new CCCL cuda::copy_bytes API.

It’s suggested in any fashionable CUDA code by no means to depend on the default stream and to all the time depend on streams.

We replace the code accordingly:

A devoted init_stream is used for the preliminary allocations of the pinned host buffers. Every iteration of the parallel for loop now owns its personal cuda::stream for the computation pipeline:

Following these modifications, we are able to take a remaining take a look at the timeline:

We now have an entire overlap between our kernels and reminiscence copies.

The ultimate length to compute all three pictures following all our enhancements is 23 milliseconds, ranging from 6.8 seconds.

Your flip

Utilizing the CUDA Developer’s Toolbox, we made the code safer, simpler to take care of and sooner. No low-level optimizations have been used, but the code is 300x sooner.

Check out this code your self, and run on Google Colab in the event you like.

We have now additionally constructed a full class to discover ways to use these instruments intimately. It’s freely obtainable on YouTube alongside hyperlinks to follow on Google Colab.