CUDA(17)의 메모리 조작과 접근 빈도 문제

요약
본고는 주로 CUDA의 메모리 조작과 주파수 접근 문제를 다루고 있다.
 
1. 문제 설명
Register의 데이터가 글로벌 메모리의 데이터에 대한 조작에 소모되는 시간은 글로벌 메모리의 데이터가 글로벌 메모리의 데이터에 대한 조작보다 길다. 이 가설의 구상이 정확한지 더 많은 검증이 필요하다.
 
2. 유효성 검사 시나리오
데이터 초기화
Data = [0, 1, 2, 3 ……, 1024*192-1]; Tata = [1, 1, 1, 1, ……, 1];
Plan A: Data를 Register에 넣고, tata를 글로벌 메모리에 넣고, Data[i] = Data[i] + Tata[i]를 조작합니다.Plan B: Data를 글로벌 메모리에 넣고 tata를 글로벌 메모리에 넣고 Data[i] = Data[i] + Tata[i]를 조작합니다.Plan C: Data를 register에 넣고, tata를 shared memory에 넣고, Data[i] = Data[i] + Tata[i]를 수행합니다.PlanD: Data를 글로벌 메모리에 배치하고 tata를 shared 메모리에 배치하여 Data[i] = Data[i] + Tata[i]를 수행합니다.
 
3. 코드 구현
 

#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include 
#include 
#include 

using namespace std;

#define N 1024*19200 // N threads

// data->register 
// tata->global
__global__ void R2G(int *data_gpu, int *tata_gpu, int threads_number){	
	
	int tid  = blockIdx.x * blockDim.x + threadIdx.x;
	int temp;

	if(tid < threads_number){
		temp = data_gpu[tid];
		temp = temp + tata_gpu[tid];
		data_gpu[tid] = temp;
	}
}

// data->global 
// tata->global
__global__ void G2G(int *data_gpu, int *tata_gpu, int threads_number){	
	
	int tid  = blockIdx.x * blockDim.x + threadIdx.x;

	if(tid < threads_number){
		data_gpu[tid] = data_gpu[tid]+ tata_gpu[tid];
	}
}

// data->register 
// tata->shared
__global__ void R2S(int *data_gpu, int *tata_gpu, int threads_number){	
	
	int tid  = blockIdx.x * blockDim.x + threadIdx.x;
	int temp;
	__shared__ int t;
	
	if(blockIdx.x < 19200){
		t = tata_gpu[blockIdx.x];
	}

	if(tid < threads_number){
		temp = data_gpu[tid];
		temp = temp + t;
		data_gpu[tid] = temp;
	}
}

// data->global 
// tata->shared
__global__ void G2S(int *data_gpu, int *tata_gpu, int threads_number){	
	
	int tid  = blockIdx.x * blockDim.x + threadIdx.x;
	__shared__ int t;
	if(blockIdx.x < 19200){
		t = tata_gpu[blockIdx.x];
	}

	if(tid < threads_number){
		data_gpu[tid] = data_gpu[tid]+ t;
	}
}

// data->register 
// tata->register
__global__ void R2R(int *data_gpu, int *tata_gpu, int threads_number){	
	
	int tid  = blockIdx.x * blockDim.x + threadIdx.x;
	int temp;
	int t;

	if(tid < threads_number){
		temp = data_gpu[tid];
		t = tata_gpu[tid];
		temp = temp + t;
		data_gpu[tid] = temp;
	}
}



int main(){
	
	//--------------------------CPU version-------------------------------------------------------
	//  
	int *data;
	data = (int *)malloc(sizeof(int)*N);
	
	int *tata;
	tata = (int *)malloc(sizeof(int)*N);

	//  
	for(int i = 0; i < N; i++){
		data[i] = i;
		tata[i] = 1;
	}


	////  
	//for(int i = 0; i < N; i++){
	//	cout << data[i] << " ";
	//}
	//cout << endl;
	
	

	//--------------------------GPU version-------------------------------------------------------	
	cudaSetDevice(0);
	cudaError_t cuda_error;

	// load data to GPU
	int *data_gpu;
	cuda_error = cudaMalloc((void**)&data_gpu, sizeof(int)*N);
	cuda_error = cudaMemcpy(data_gpu, data, sizeof(int)*N, cudaMemcpyHostToDevice);
	
	int *tata_gpu;
	cuda_error = cudaMalloc((void**)&tata_gpu, sizeof(int)*N);
	cuda_error = cudaMemcpy(tata_gpu, tata, sizeof(int)*N, cudaMemcpyHostToDevice);

	// init gpu
	dim3 threads(1024,1,1);
	dim3 blocks(19200,1,1);

	// start timing
	float time_elapsed=0;
	cudaEvent_t start,stop;
	cudaEventCreate(&start);    
	cudaEventCreate(&stop);
	cudaEventRecord( start,0);

	// kernel
	//R2G<<>>(data_gpu,tata_gpu, N); // data->register tata->global
	//G2G<<>>(data_gpu,tata_gpu, N); // data->global tata->global
	//G2S<<>>(data_gpu,tata_gpu, N); // data->global tata->shared
	//R2S<<>>(data_gpu,tata_gpu, N); // data->register tata->shared
	R2R<<>>(data_gpu,tata_gpu, N); // data->register tata->register

	// finish timing
	cudaEventRecord(stop,0);    
	cudaEventSynchronize(start);    
	cudaEventSynchronize(stop);   
	cudaEventElapsedTime(&time_elapsed,start,stop);  

	// back data to CPU
	cuda_error = cudaMemcpy(data, data_gpu, sizeof(int)*N, cudaMemcpyDeviceToHost);

	// check result
	for(int i = 0; i < N; i++){
		//cout << data[i] << endl;
		if(data[i]!=i+1)
			cout << "failed:"<

 
 
 
4. 결과 분석
메모리 접근 주파수: 메모리 조작을 한 번 하면 글로벌 메모리에서 글로벌 메모리로 효율이 가장 높습니다.여러 번 조작하면 shared memory와 register의 속도 우위가 서서히 나타난다.