1D grid of 1D blocks

__device__ int getGlobalIdx_1D_1D()

{

return blockIdx.x *blockDim.x + threadIdx.x;

}

1D grid of 2D blocks

__device__ int getGlobalIdx_1D_2D()

{

return blockIdx.x * blockDim.x * blockDim.y + threadIdx.y * blockDim.x + threadIdx.x;

}

1D grid of 3D blocks

__device__ int getGlobalIdx_1D_3D()

{

return blockIdx.x * blockDim.x * blockDim.y * blockDim.z + threadIdx.z * blockDim.y * blockDim.x + threadIdx.y * blockDim.x + threadIdx.x;

} 

2D grid of 1D blocks 

__device__ int getGlobalIdx_2D_1D()

{

int blockId   = blockIdx.y * gridDim.x + blockIdx.x;

int threadId = blockId * blockDim.x + threadIdx.x; 

return threadId;

}

2D grid of 2D blocks  

 __device__ int getGlobalIdx_2D_2D()

{

int blockId = blockIdx.x + blockIdx.y * gridDim.x; 

int threadId = blockId * (blockDim.x * blockDim.y) + (threadIdx.y * blockDim.x) + threadIdx.x;

return threadId;

}

2D grid of 3D blocks

__device__ int getGlobalIdx_2D_3D()

{

int blockId = blockIdx.x 

+ blockIdx.y * gridDim.x; 

int threadId = blockId * (blockDim.x * blockDim.y * blockDim.z)

  + (threadIdx.z * (blockDim.x * blockDim.y))

  + (threadIdx.y * blockDim.x)

  + threadIdx.x;

return threadId;

} 

3D grid of 1D blocks

__device__ int getGlobalIdx_3D_1D()

{

int blockId = blockIdx.x 

+ blockIdx.y * gridDim.x 

+ gridDim.x * gridDim.y * blockIdx.z; 

int threadId = blockId * blockDim.x + threadIdx.x;

return threadId;

} 

3D grid of 2D blocks

__device__ int getGlobalIdx_3D_2D()

{

int blockId = blockIdx.x 

        + blockIdx.y * gridDim.x 

+ gridDim.x * gridDim.y * blockIdx.z; 

int threadId = blockId * (blockDim.x * blockDim.y)

 + (threadIdx.y * blockDim.x)

 + threadIdx.x;

return threadId;

}

3D grid of 3D blocks

__device__ int getGlobalIdx_3D_3D()

{

int blockId = blockIdx.x 

+ blockIdx.y * gridDim.x 

+ gridDim.x * gridDim.y * blockIdx.z; 

int threadId = blockId * (blockDim.x * blockDim.y * blockDim.z)

 + (threadIdx.z * (blockDim.x * blockDim.y))

 + (threadIdx.y * blockDim.x)

 + threadIdx.x;

return threadId;

}

1. visual Studio NSIGHT 메뉴에 Start Performance Analysis 항목을 선택. 

2. Activity Type 에서 Profile CUDA Application 항목을 선택한다.

3. Application Control 항목에서 Launch 버튼을 클릭한다.

4. Launch 버튼을 클릭하면, kill 버튼을 클릭할 때까지 반복해서 프로그램을 실행한다.

5. Kill 을 클릭하면, 리포팅을 확인 할 수 있다.

CUDA  디버깅 환경 구성.

- 프로그램 설치. [Nsight Visual Studio Edition 5.2 버전 기준]

- Nsight Monitor 

  4. Disable D3D acceleration for WPF

    - DisableWpfHardwareAcceleration.reg 항목을 검색하여, 등록한다. [C:\Program Files (x86)\NVIDIA Corporation\Nsight Visual Studio Edition 5.2\Host\Common ] 

 5. 컴퓨터 재부팅.

- Visual Studio 2013 

1. Nsight User Properties 설정

 1. 솔류션 검색창->프로젝트 선택 후 팝업메뉴에서 Nsight User Properties 항목을 선택한다.      

2. Launch 항목을 선택한다.

2. Run the Memory Checker 체크.

Native C++ 에서, WinRT 를 사용할 수 있게 됐다.

바로가기

https://docs.microsoft.com/ko-kr/cpp/cpp/support-for-cpp11-14-17-features-modern-cpp

#include <chrono>

using namespace std;

typedef std::chrono::duration<int, std::micro> millisecs_t;   // 시간단위 자료형  재정의.

std::chrono::steady_clock::time_point   nStart, nEnd; //  체크 시간을 저장할 변수 선언.

nStart = std::chrono::steady_clock::now(); //시간측정

~~

~~

~~

~~

nEnd = std::chrono::steady_clock::now(); // 시간측정.

//두 시간 간의 차를 duration_cast 캐스팅 연산자 변환 템플릿 클래스를 이용해서 변환한다.

long long nElapsedTime = std::chrono::duration_cast<millisecs_t>(nEnd - nStart)).count();  

TimeMeasure.zip

:: CUDA Visual Profiler 

  - 소 개

  - 실 행

    1. console

#include <stdio.h>
#include <stdlib.h>
#include <cutil_inline.h>
#include <cuda.h>

//벡터합 계산 커널
__global__ void vectorAdd(int *a, int *b, int *c)
{
	int idx = blockIdx.x*blockDim.x + threadIdx.x;
	c[idx] = a[idx] + b[idx];
}

int main()
{
	int i = 0;
	int nBlocks = 1024;
	int nThreads = 512;
	int Size = nBlocks* nThreads;
	size_t BufferSize = Size*sizeof(int);
	int *h_a, * h_b, * h_c;              // CPU 고정된 메모리 할당을 위한 포인터
	int *d_a, *d_b, *d_c;                // 메모리 맵핑을 위한 디바이스 포인터

	/* Allocate mapped CPU memory. */

	cudaHostAlloc((void **)& h_a, BufferSize, cudaHostAllocMapped);
	cudaHostAlloc((void **)& h_b, BufferSize, cudaHostAllocMapped);
	cudaHostAlloc((void **)& h_c, BufferSize, cudaHostAllocMapped);

	/* Initialize the vectors. */

	for(i = 0; i < Size; i++)
	{
		h_a[i] = i;
		h_b[i] = i;
	}

	//CPU 메모리를 맵핑하여 GPU 메모리 포인터로 가리키도록 한다.
	cudaHostGetDevicePointer((void **)&d_a, (void *)h_a, 0);
	cudaHostGetDevicePointer((void **)&d_b, (void *)h_b, 0);
	cudaHostGetDevicePointer((void **)&d_c, (void *)h_c, 0);

	////맵핑된 메모리 포인터를 이용하여 커널 호출

	vectorAdd<<< nBlocks, nThreads >>>(d_a, d_b, d_c);  
	cudaThreadSynchronize();

	cudaFreeHost(h_a);
	cudaFreeHost(h_b);
	cudaFreeHost(h_c);

	return 0;
}

bank conflict

SIMT를 실행시킬때 문제중의 하나가 memory access이다. GPU에서는 동시에 여러개의 데이터를 처리해야하기 때문에, 동시에 여러개의 데이터에 access를 허용한다. 이것을 하기 위해서 GPU는 shared memory를 각 warp마다 일정 갯수의 memory bank로 나누어 두었는데, 각각의 bank는 bank단위로 동시에 접근할 수 있다. 이때 bank conflict란 프로그래밍 잘못으로 동시에 서로 다른 thread가 특정 bank를 access할때 발생하는 문제이다.

서로 다른 thread가 하나의 특정 bank에 access하게 되면, 각각의 thread는 해당 bank에 접근하기 위해서 순차적으로 변하게 되고, 이는 병렬적으로 처리하려고 했던 의도를 벗어나게된다. 즉, 의도하지 않은 행동이 된다는 것이다.

이것은 GPU를 이용한 병렬처리에서 속도를 낮추게 하는 문제를 일으킨다. 여기서는 해당 문제 상황에 대해서 정리를 하고자 한다.

1.x대에서는 warp의 크기가 32이고, bank의 갯수가 16개이다. 이것은 최대 shared memory에 접근할 수 있는 thread의 갯수는 16개뿐이라는 이야기다. 이 문제를 해결하기 위해서 CUDA는 warp를 두개의 half-warp로 나눈다. 두개의 half-warp를 순차적으로 실행하므로써 bank conflict가 사라진다.(단, shared memory access를 하는데, 2 cycle이 걸리겠지?)

bank conflict가 발생하는 다른 경우를 살펴보자.

1. 다음과 같은 코드가 있다고 하자.

__shared__ struct type shared[32]; 

struct type data = shared[BaseIndex + tid];

이때

1) type 구조체가 다음과 같이 정의되었다면,

struct type { 
float x, y, z; 
};

위의 구조체에 의해서 실행되는 thread들은 bank conflict를 발생시키지 않는다. 
여기서 bank conflict가 발생하지 않는 이유를 설명하기 위해서, bank가 2개이고, thread도 2개라고 가정을 하자. 또한 데이터는 순서대로 bank 0과 bank 1을 왔다갔다하면서 쌓인다고 하자. 그럼 다음과 같이 데이터들이 저장되어 있을 것이다. (bank의 access단위는 32bits라고 하자.)

이때 위의 구조체대로라면, Thread 0은 [0]을 Thread 1은 [3]에 접근한다. [0]은 Bank 0에 있고, [3]은 Bank 1에 있으므로, 다른 Thread가 동시에 한 Bank에 접근하지 않는다. 

그러므로 Bank Conflict가 발생되지 않는다.

2) 만약 구조체가 다음과 같다면,

struct type { 
float x, y; 
};

위의 코드는 bank conflict를 발생시킨다. 이유는 위의 1)번과 같은 가정을 했을때, thread 0은 [0], thread 1은 [2]에 접근하는데, 이때 [0]과 [2]는 table에서 보듯이 같은 bank에 있으므로, 두개의 thread가 동시에 같은 bank를 access하게 된다. 그러므로 bank conflict가 발생한다.

예제 1

__shared__ char shared[32]; 
char data = shared[BaseIndex + tid];

1. 결과: 위의 코드는 Bank conflict를 일으킨다. 

2. 원인: thread 0은 shared[0]을 (BaseIndex는 0이라 하자.), thread 1은 shared[1]에 접근하는데, 이때 char는 1byte이고, bank의 한 데이터는 4bytes이므로, bank0에서 [0]및 [1]을 읽어야 한다. 즉, 같은 bank를 access하므로 bank conflict가 발생한다.

3.해결:

__shared__ char shared[32]; 
char data = shared[BaseIndex + 4 * tid];

와 같이 하면된다.

예제 2

__shared__ double shared[32]; 
double data = shared[BaseIndex + tid];

1. 결과: 위의 코드는 Bank conflict를 일으킨다.

2. 원인: thread 0은 shared[0]을 (BaseIndex는 0이라 하자.), thread 1은 shared[2]에 접근한다. 이유는 double은 8bytes이므로 한 bank에서만 읽어서는 안된다. 이와 같은 이유로 같은 bank를 access하므로 bank conflict가 발생한다.

3. 결과: 쪼개서 읽어야겠다.  어떻게 쪼개지?

예제 3 s는 stride이다.

__shared__ float shared[32]; 
float data = shared[BaseIndex + s * tid];

위와 같은 코드가 있을 때 각 thread가 s가 어떤 조건을 가져야지만, bank conflict가 나지 않을지를 찾아보자. 
bank갯수가 m이라고 하고, 위의 요청 갯수가 n개라고 하자.

이때 bank conflict가 일어나지 않을 조건은 s*n이 m의 배수이거나 m과 같을 경우(예, 위의 1)에 조건을 기준으로 설명을 하면, ), m과 s의 최대공약수가 d라고 하면, n이 m/d의 배수일 경우이다.

출처: https://dyanos.wordpress.com/2009/10/21/cuda-%ED%94%84%EB%A1%9C%EA%B7%B8%EB%9E%98%EB%B0%8D-%EC%A0%95%EB%A6%AC/  에서 정리 함. [임시보관]

2차원 배열일 때.

툴버전: Visual Studio 2013 

1.  솔류션 탐색창에서, 프로젝트를 마우스 우클릭-> Build Dependencies ->Build Customizations -> 쿠다버전 선택.

  - 설정하지 않을 경우, cu 파일 속성창에서, Item Type 항목을 CUDA C/C++로 선택할 수 없게 된다.

2. .cu 파일 속성 수정.

  -   cu 파일에서 마우스 우클릭 -> Properties 항목 클릭,

    Configuration Properties

        1.General

            Item Type 항목->  Drop-down 리스트에서   CUDA C/C++ 항목을 선택한다.

        2. CUDA C/C++

          -  Common 

              - Additional Include Directories : ./;../../common/inc 

              - Target Marhine Platform : 64-bit [--machine 64 ]

          - Device 

              Code Generation :