<commit_before>---
title: CUDA로 배워보는 병렬 프로그래밍(Parallel Programming) - 행렬(Matrix) 계산
categories: [Parallel]
tags: [parallel, cuda]
description: 병렬 프로그래밍의 기본적인 연산 방법 및 원리 및 행렬(Matrix)에 직접 적용
date : 2016-10-01 22:00:00 +0900
---

우선 행렬 계산을 하기 전에 CUDA에서 어떤 방식으로 여러 개의 쓰레드를 돌리는지를 알아야한다. CUDA에선 `BLOCK`과 `GRID`로 쓰레드 그룹을 관리한다.

`BLOCK`에서는 여러 개의 `THREAD`를 (x, y, z) 즉, 3차원 이하로 가지고 있을 수 있고, `GRID`는 여러 `BLOCK`을 2차원 이하(x, y) 로 가지고 있을 수 있다. 그림으로 나타내면 다음과 같다. 그리고 한 `BLOCK`에는 1024개가 넘는 `THREAD`를 가지고 있을 수 없다.

![image]("https://c2.staticflickr.com/9/8679/30000520041_e2bb9beeff_b.jpg)

그리고 이와 같은 `BLOCK`과 `GRID`를 정의한 다음에 `Kernel Function` 즉 Device에서 돌아갈 함수랑 같이 쓰여야한다.

{% highlight c++ %}
__global__ void kernelFunc();

dim3 DimGrid(10, 10); //100 thread blocks
dim3 DimBlock(4, 5, 6); // 120 thread per Blocks

kernelFunc<<< DimGrid, DimBlock>>>();
{% endhighlight %}

위와 같이 C++에서 템플릿을 넘기듯이 사용하면 된다.
`dim3` 자료형 같은 경우 3개 미만으로 parameter를 넘기는 것도 가능하다.

위와 같이 정의하면 각 `thread`는 고유한 `blockIdx`와 `threadIdx`를 가진다.

위 두 개의 고유한 값을 이용해서 각 `thread`는 자기의 고유한 위치를 알 수 있다. 그리고 이렇게 정의된 `Blocks of Threads`는 CUDA 자체 Queue를 통해서 `SMs(Streaming multiprocessors)`에 할당되어 CUDA가 계산을 해준다.

## 행렬 덧셈 구현

{% highlight c++ %}
#include <cstdio>

__global__ void addKernel(int *c, const int *a, const int *b) {
    int x = threadIdx.x;
    int y = threadIdx.y;
    int i = y * (blockDim.x) + x; // 1차원 인덱스로 변환

    c[i] = a[i] + b[i];<commit_msg>add image<commit_after>---
title: CUDA로 배워보는 병렬 프로그래밍(Parallel Programming) - 행렬(Matrix) 계산
categories: [Parallel]
tags: [parallel, cuda]
description: 병렬 프로그래밍의 기본적인 연산 방법 및 원리 및 행렬(Matrix)에 직접 적용
date : 2016-10-01 22:00:00 +0900
---

우선 행렬 계산을 하기 전에 CUDA에서 어떤 방식으로 여러 개의 쓰레드를 돌리는지를 알아야한다. CUDA에선 `BLOCK`과 `GRID`로 쓰레드 그룹을 관리한다.

`BLOCK`에서는 여러 개의 `THREAD`를 (x, y, z) 즉, 3차원 이하로 가지고 있을 수 있고, `GRID`는 여러 `BLOCK`을 2차원 이하(x, y) 로 가지고 있을 수 있다. 그림으로 나타내면 다음과 같다. 그리고 한 `BLOCK`에는 1024개가 넘는 `THREAD`를 가지고 있을 수 없다.

![image](https://c2.staticflickr.com/9/8679/30000520041_e2bb9beeff_b.jpg)

그리고 이와 같은 `BLOCK`과 `GRID`를 정의한 다음에 `Kernel Function` 즉 Device에서 돌아갈 함수랑 같이 쓰여야한다.

{% highlight c++ %}
__global__ void kernelFunc();

dim3 DimGrid(10, 10); //100 thread blocks
dim3 DimBlock(4, 5, 6); // 120 thread per Blocks

kernelFunc<<< DimGrid, DimBlock>>>();
{% endhighlight %}

위와 같이 C++에서 템플릿을 넘기듯이 사용하면 된다.
`dim3` 자료형 같은 경우 3개 미만으로 parameter를 넘기는 것도 가능하다.

위와 같이 정의하면 각 `thread`는 고유한 `blockIdx`와 `threadIdx`를 가진다.

위 두 개의 고유한 값을 이용해서 각 `thread`는 자기의 고유한 위치를 알 수 있다. 그리고 이렇게 정의된 `Blocks of Threads`는 CUDA 자체 Queue를 통해서 `SMs(Streaming multiprocessors)`에 할당되어 CUDA가 계산을 해준다.

## 행렬 덧셈 구현

{% highlight c++ %}
#include <cstdio>

__global__ void addKernel(int *c, const int *a, const int *b) {
    int x = threadIdx.x;
    int y = threadIdx.y;
    int i = y * (blockDim.x) + x; // 1차원 인덱스로 변환

    c[i] = a[i] + b[i];
{'fraction_non_alphanumeric': 0.13145707160669753, 'fraction_numerical': 0.04346277164232276, 'mean_word_length': 3.6184210526315788, 'pattern_counts': {'":': 0, '<': 11, '<?xml version=': 0, '>': 11, 'https://': 2, 'lorem ipsum': 0, 'www.': 0, 'xml': 0}, 'pii_count': 0, 'substrings_counts': 0, 'word_list_counts': {'cursed_substrings.json': 2, 'profanity_word_list.json': 0, 'sexual_word_list.json': 0, 'zh_pornsignals.json': 0}}
{'dir': 'git-commits-cleaned', 'id': '3131635', 'n_tokens_mistral': 1827, 'n_tokens_neox': 2010, 'n_words': 434}