Skip to content

Commit

Permalink
Add Christina Koutsou to contributors and add support for CUDA presen…
Browse files Browse the repository at this point in the history
…tation and blog post (#238)
  • Loading branch information
kchristin22 authored Aug 7, 2024
1 parent 8914a6f commit a4877f9
Show file tree
Hide file tree
Showing 6 changed files with 136 additions and 0 deletions.
25 changes: 25 additions & 0 deletions _data/contributors.yml
Original file line number Diff line number Diff line change
Expand Up @@ -374,6 +374,31 @@
mentors: Vassil Vassilev, Matheus Izvekov
proposal: /assets/docs/Sahil_Patidar_Proposal_2024_ClangRepl.pdf

- name: "Christina Koutsou"
photo: christina_koutsou.jpg
info: "Google Summer of Code 2024 Contributor"
email: christinakoutsou22@gmail.com
education: "Integrated Master's in Electrical and Computer Engineering, Aristotle University of Thessaloniki, Greece"
github: "https://github.com/kchristin22"
active: 1
linkedin: "https://www.linkedin.com/in/christina-koutsou-69416b28a/"
projects:
- title: "Reverse-mode automatic differentiation of GPU (CUDA) kernels using Clad"
status: Ongoing
description: |
Nowadays, the rise of AI has shed light into the power of GPUs. The notion of General Purpose GPU Programming is
becoming more and more popular and it seems that the scientific community is increasingly favoring it over CPU Programming.
Consequently, implementation of mathematics and operations needed for such projects are getting adjusted to GPU's architecture.
Automatic differentiation is a notable concept in this context, finding applications across diverse domains from ML to Finance to Physics.
Clad is a clang plugin for automatic differentiation that performs source-to-source transformation and produces a function capable of
computing the derivatives of a given function at compile time. This project aims to widen Clad’s use range and audience by enabling
the reverse-mode automatic differentiation of CUDA kernels. The total goal of the project is to support the differentiation of CUDA
kernels that may also include typical CUDA built-in objects (e.g. threadIdx, blockDim etc.), which are employed to prevent race conditions,
using Clad. These produced kernels will compute the derivative of an argument specified by the user as the output based on an input parameter
of their choosing. In addition, the user must be able to call these kernels with a custom grid configuration.
proposal: /assets/docs/Christina_Koutsou_GSoC_2024.pdf
mentors: Vassil Vassilev, Parth Arora, Alexander Penev

- name: "This could be you!"
photo: rock.jpg
info: See <a href="/careers">openings</a> for more info
Expand Down
4 changes: 4 additions & 0 deletions _data/standing_meetings.yml
Original file line number Diff line number Diff line change
Expand Up @@ -21,6 +21,10 @@
link: "[Slides](/assets/presentations/CaaS_Weekly_12_06_2024_Maksym_Andriichuk_Activity.pdf)"
- title: Updates
link: "[Meeting Notes](https://docs.google.com/document/d/1HMUtM8rRPeTYDqsWZVD-iQityjJj9_GKgu0FHHbiBr0)"
- title: "Enabling reverse-mode automatic differentiation of GPU (CUDA) kernels using Clad"
date: 2024-05-15 17:00:00 +0200
speaker: "Christina Koutsou"
link: "[Slides](/assets/presentations/CaaS_Weekly_15_05_2024_Christina_Enbale_reverse_mode_autodiff_for_kernels_first_presentation.pdf)"
- title: "Integrate a large language model with the Xeus-cpp jupyter kernel"
date: 2024-06-05 17:00:00 +0200
speaker: "Tharun Anandh"
Expand Down
107 changes: 107 additions & 0 deletions _posts/2024-05-18-reverse-mode-autodiff-of-cuda-kernels.md
Original file line number Diff line number Diff line change
@@ -0,0 +1,107 @@
---
title: "Enabling reverse-mode automatic differentiation of GPU kernels"
layout: post
excerpt: "Nowadays, the rise of General Purpose GPU programming has caused more and more tools used by the scientific community to adjust to GPU's architecture. This project aims to allow `Clad` to ride that tide and broaden its use-range by enabling reverse-mode automatic differentiation of CUDA kernels."
sitemap: false
author: Christina Koutsou
permalink: blogs/gsoc24_christina_koutsou_project_introductory_blog/
date: 2024-05-18

---
### Overview of the project

Nowadays, the rise of General Purpose GPU programming has caused more and more tools used by the scientific community to adjust to GPU's architecture. This project aims to allow `Clad` to ride that tide and broaden its use-range by enabling reverse-mode automatic differentiation of CUDA kernels.

Mentors: Vassil Vassilev, Parth Arora, Alexander Penev
[Proposal](/assets/docs/Christina_Koutsou_GSoC_2024.pdf)
[Slides](/assets/presentations/CaaS_Weekly_15_05_2024_Christina_Enbale_reverse_mode_autodiff_for_kernels_first_presentation.pdf)

---
### Current status and end goal

Currently, `Clad` supports differentiation of both host (functions executed by the CPU) and device functions (functions executed by the GPU). However, since the device function resides in the GPU's memory, there's a need for an overlap between the CPU and the GPU. Hence, a global function is used as an intermediary. The global function, also known as a kernel, is executed by the GPU and launched by the CPU. The following example illustrates that interaction:

```cpp
__device__ double fn(double *a) {
  return *a * *a;
}

__global__ void compute(double* d_a, double* d_result) {
  auto fn_grad = clad::gradient(fn, "a");
  fn_grad.execute(d_x, d_result);
}

int main(void) {
  (...) // memory allocations and initializations
  compute<<<1, 1>>>(d_a, d_result);
  cudaDeviceSynchronize();
  // copy back result to CPU
  cudaMemcpy(result.data(), d_result, N * sizeof(double),
                                      cudaMemcpyDeviceToHost);
  return 0;
}
```
It is evident that if the device function was a global one instead, the program would be much more simple and efficient. Kernels are also more widely used in general.
```cpp
__global__ void fn(double *a) {
  *a *= *a;
}
int main(void) {
  (...) // memory allocations and initializations
  auto fn_grad = clad::gradient(fn, "a");
  fn_grad.execute(d_x, d_result);
  cudaDeviceSynchronize();
  // copy back result to CPU
  cudaMemcpy(result.data(), d_result, N * sizeof(double),
                                      cudaMemcpyDeviceToHost);
  return 0;
}
```

Thus, the end goal of this project would be to compute the gradient of a basic kernel function using Clad.

### Approach

In order to define what a basic kernel looks like and solidify the project's objectives, it is helpful to research some problems that utilize GPU kernels for their solution. Furthermore, since reverse-mode automatic differentiation is particularly beneficial when the number of inputs exceeds the number of outputs in a function, identifying such problems will be especially advantageous. Gathering such examples provide a concrete framework for the project's goals.
However, it is possible to already identify some problems that will arise during the process. These issues can be broken down into two categories:
1) General Support: The ability to produce and execute a kernel’s derivative function without it necessarily being correct. Essentially, this refers to the handling of the kernel calls.
2) Support of kernel’s nature: Derive the kernel with respect to its characteristics. This refers to correct handling of the function's body and computation of its derivative.
###### General Support: Kernel’s derivative compilation
What differentiates CUDA kernels from device or host functions is that they cannot be called without a provided grid configuration to be used. Hence, any call to them should be accompanied by such an expression. When the function to be derived is identified as a kernel through its global attribute, a default configuration expression is to be created by calling `clang::Sema::ActOnCUDAExecConfigExpr()`. This is passed to the call expression of the derived function. The configuration is single-threaded to ensure that a single thread computes the gradient.
###### General Support: Kernel’s derivative execution
Similarly to the kernel's compilation, when the derived kernel is to be executed, it too needs to be provided with a grid configuration. For that purpose, the API can be modified to include an overloaded function of `execute()` and its nested calls, `execute_helper()` and `execute_with_default_args()`. The latter will eventually invoke the derived kernel like so:
```cpp
return f<<<grid_size, block_size, shared_mem_size, stream>>>
(static_cast<Args>(args)..., static_cast<Rest>(nullptr)...);
```
###### General Support: Derivation of a kernel call
To further enhance the support for kernel differentiation, we should ensure that deriving a function that includes a kernel invocation can successfully call the derived kernel. Thus, when the `CUDAKernelCallExpr` is visited during the top level of derivation, we should store its configuration to use for the pullback kernel call. Specifically, a Visit function for kernels calls will be written, that mimics the one used for typical CallExpr nodes. However, it will additionally include storing the specified configuration using `getConfig()` on the kernel node and then passing it to the creation of the pullback kernel call expression through `clang::Sema::ActOnCallExpr()`.
###### Support of kernel’s characteristics: Specify output argument to derive
Kernels are void functions, which means that derivation based on the return statement is not possible. As a result, we need to know the output of the function. This can be accomplished by the expansion of the API to include an overload of the `gradient()` function, where the user can specify the output argument of only void functions. This argument will be passed to the differentiation request of this function and stored to the variable list whose expressions are derived.
###### Support of kernel’s characteristics: Account for write race conditions in computation of the derivative value
To also account for the multithreaded environment of a kernel, the plus-assign operation to the derived function's output would result to incorrect results, as each thread would add the derivative value. Hence, this operation should be instead replaced with a plain assign operation. Furthermore, if array indexing is involved, the derived variable should also be considered as an array to avoid write race conditions in case the derivative array of the output is not initialized equally.
```cpp
__global__ void compute(double *in, double *out, double val)
{
  int index = threadIdx.x;
  out[index] = in[index] + val;
}

(derived function){
...
double _r_d0 = _d_out[index0];
* _d_val += _r_d0; —> * _d_val[index0] = _r_d0;
}
```
###### Support of kernel’s characteristics: Handling of CUDA built-in objects
Another common characteristic of CUDA kernels is the use of built-in objects, e.g. `threadIdx`, as depicted above, to ensure that each thread computes its own share of the final result. These nodes should only be cloned when visited in the global initializations and be treated as integers when differentiated, or in other words their derivatives should be 0. In addition,  support for the shared memory macro must also be included in the project’s scope, by not discarding it when visiting a variable or array declaration.
---
### About me
Hi, I'm Christina, an Electrical and Computer Engineering major in Aristotle University of Thessaloniki, Greece. My participation in Google Summer of Code was fueled by a deep appreciation for open-source projects, stemming from my two-year experience with an open-source student team that particpates in ESA Education's Fly Your Satellite program. Despite my passion for physics, it did not lead me down that career path, but I'm trying to channel my enthusiasm into developing tools for the scientific community. Clad offers a great opportunity to do just that and contribute to research advancements.
Binary file added assets/docs/Christina_Koutsou_GSoC_2024.pdf
Binary file not shown.
Binary file not shown.
Binary file added images/team/christina_koutsou.jpg
Loading
Sorry, something went wrong. Reload?
Sorry, we cannot display this file.
Sorry, this file is invalid so it cannot be displayed.

0 comments on commit a4877f9

Please sign in to comment.