[초보 자 는 매일 CUDAC]여러 개의 CUDA 흐름 을 사용 하여 프로그램의 실행 효율 을 높 인 다.

CUDA 흐름 은 GPU 작업 대기 열 을 표시 하 며,이 대기 열 에 추 가 된 작업 은 대기 열의 우선 순위 로 실 행 됩 니 다.CUDA 흐름 을 사용 하면 작업 수준의 병행 을 실현 할 수 있 습 니 다.예 를 들 어 GPU 가 핵 함 수 를 실행 하 는 동시에 호스트 와 장치 간 에 데 이 터 를 교환 할 수 있 습 니 다(전 제 는 GPU 가 중첩 을 지원 하고 property 의 deviceOverlay 는 true 입 니 다).
cudaMemcpy Async 함수 의 기능 은 GPU 와 호스트 사이 에서 데 이 터 를 복사 하 는 것 입 니 다.이것 은 비동기 함수 입 니 다.즉,함수 가 호출 된 후에 하나의 요청 만 설치 하여 스 트림 에서 메모리 복사 작업 을 수행 하 는 것 을 표시 합 니 다.함수 가 되 돌 아 왔 을 때 복사 작업 이 반드시 시작 되 거나 끝 날 필 요 는 없습니다.이 작업 은 실행 대기 열 에 넣 고 다음 스 트림 에 넣 기 전에 실 행 됩 니 다.
실험 은 한 조 의 데 이 터 를 GPU 로 나 누 어 실행 결 과 를 되 돌려 줌 으로 써 cuda 흐름 을 사용 하면 프로그램의 실행 효율 을 높 일 수 있다 는 것 을 설명 한다.원 리 는 주로 데이터 복제 작업 과 핵 함수 실행 작업 을 교차 집행 하 게 하 는 것 이다.첫 번 째 핵 함수 실행 이 끝 날 때 까지 기 다 렸 다가 두 번 째 데이터 복 제 를 시작 하지 않 아 도 순서 집행 에 따 른 지연 을 줄 일 수 있다(컴 파일 에서 파이프라인 을 사용 하여 충돌 을 해결 하 는 전제 에서 효율 을 높이 는 것 과 유사 하 다).
프로그램 코드 는 다음 과 같 습 니 다:

#include "cuda_runtime.h"
#include "cutil_inline.h"

#include <stdio.h>
#include <math.h>

static void HandleError( cudaError_t err,const char *file,int line ) 
{

	if (err != cudaSuccess)

	{

		printf( "%s in %s at line %d
", cudaGetErrorString( err ),
			file, line );
		exit( EXIT_FAILURE );
	}
}
#define HANDLE_ERROR( err ) (HandleError( err, __FILE__, __LINE__ ))
#define N (1024*1024)
#define FULL_DATA_SIZE N*20

__global__ void kernel(int* a, int *b, int*c)
{
	int idx = blockIdx.x * blockDim.x + threadIdx.x;
	int offset = gridDim.x * blockDim.x;

	if (idx < N)
	{
		int idx1 = (idx + 1) % 256;
		int idx2 = (idx + 2) % 256;
		float as = (a[idx] + a[idx1] + a[idx2]) / 3;
		float bs = (b[idx] + b[idx1] + b[idx2]) / 3;
		c[idx] = (as + bs) / 2;
	}
}

int main()
{
	cudaDeviceProp prop;
	int devID;

	HANDLE_ERROR(cudaGetDevice(&devID));
	HANDLE_ERROR(cudaGetDeviceProperties(&prop, devID));

	if (!prop.deviceOverlap)
	{
		printf("No device will handle overlaps. so no speed up from stream.
");
		return 0;
	}

	cudaEvent_t start, stop;
	float elapsedTime;
	HANDLE_ERROR(cudaEventCreate(&start));
	HANDLE_ERROR(cudaEventCreate(&stop));
	HANDLE_ERROR(cudaEventRecord(start, 0));

	cudaStream_t stream0;
	cudaStream_t stream1;
	HANDLE_ERROR(cudaStreamCreate(&stream0));
	HANDLE_ERROR(cudaStreamCreate(&stream1));

	int *host_a, *host_b, *host_c;
	int *dev_a0, *dev_b0, *dev_c0;
	int *dev_a1, *dev_b1, *dev_c1;

	HANDLE_ERROR(cudaMalloc((void**)&dev_a0, N*sizeof(int)));
	HANDLE_ERROR(cudaMalloc((void**)&dev_b0, N*sizeof(int)));
	HANDLE_ERROR(cudaMalloc((void**)&dev_c0, N*sizeof(int)));

	HANDLE_ERROR(cudaMalloc((void**)&dev_a1, N*sizeof(int)));
	HANDLE_ERROR(cudaMalloc((void**)&dev_b1, N*sizeof(int)));
	HANDLE_ERROR(cudaMalloc((void**)&dev_c1, N*sizeof(int)));

	HANDLE_ERROR(cudaHostAlloc((void**)&host_a, FULL_DATA_SIZE*sizeof(int), cudaHostAllocDefault));
	HANDLE_ERROR(cudaHostAlloc((void**)&host_b, FULL_DATA_SIZE*sizeof(int), cudaHostAllocDefault));
	HANDLE_ERROR(cudaHostAlloc((void**)&host_c, FULL_DATA_SIZE*sizeof(int), cudaHostAllocDefault));


	for (int i=0; i<FULL_DATA_SIZE; i++)
	{
		host_a[i] = rand();
		host_b[i] = rand();
	}


	// tasks are put into stack for gpu execution
	for (int i=0; i<FULL_DATA_SIZE; i+=2*N)
	{
		HANDLE_ERROR(cudaMemcpyAsync(dev_a0, host_a+i, N*sizeof(int), cudaMemcpyHostToDevice, stream0));
		HANDLE_ERROR(cudaMemcpyAsync(dev_a1, host_a+i+N, N*sizeof(int), cudaMemcpyHostToDevice, stream1));

		HANDLE_ERROR(cudaMemcpyAsync(dev_b0, host_b+i, N*sizeof(int), cudaMemcpyHostToDevice, stream0));
		HANDLE_ERROR(cudaMemcpyAsync(dev_a1, host_a+i+N, N*sizeof(int), cudaMemcpyHostToDevice, stream1));

		kernel<<<N/256, 256, 0, stream0>>>(dev_a0, dev_b0, dev_c0);
		kernel<<<N/256, 256, 0, stream1>>>(dev_a1, dev_b1, dev_c1);

		HANDLE_ERROR(cudaMemcpyAsync(host_c+i, dev_c0, N*sizeof(int), cudaMemcpyDeviceToHost, stream0));
		HANDLE_ERROR(cudaMemcpyAsync(host_c+i+N, dev_c1, N*sizeof(int), cudaMemcpyDeviceToHost, stream1));

// 		HANDLE_ERROR(cudaMemcpyAsync(dev_a0, host_a+i, N*sizeof(int), cudaMemcpyHostToDevice, stream0));
// 		HANDLE_ERROR(cudaMemcpyAsync(dev_b0, host_b+i, N*sizeof(int), cudaMemcpyHostToDevice, stream0));
// 		kernel<<<N/256, 256, 0, stream0>>>(dev_a0, dev_b0, dev_c0);
// 		HANDLE_ERROR(cudaMemcpyAsync(host_c+i, dev_c0, N*sizeof(int), cudaMemcpyDeviceToHost, stream0));
// 
// 		HANDLE_ERROR(cudaMemcpyAsync(dev_a1, host_a+i+N, N*sizeof(int), cudaMemcpyHostToDevice, stream1));
// 		HANDLE_ERROR(cudaMemcpyAsync(dev_a1, host_a+i+N, N*sizeof(int), cudaMemcpyHostToDevice, stream1));
// 		kernel<<<N/256, 256, 0, stream1>>>(dev_a1, dev_b1, dev_c1);
// 		HANDLE_ERROR(cudaMemcpyAsync(host_c+i+N, dev_c1, N*sizeof(int), cudaMemcpyDeviceToHost, stream1));
	}

	// wait until gpu execution finish
	HANDLE_ERROR(cudaStreamSynchronize(stream0));
	HANDLE_ERROR(cudaStreamSynchronize(stream1));

	HANDLE_ERROR(cudaEventRecord(stop, 0));
	HANDLE_ERROR(cudaEventSynchronize(stop));
	HANDLE_ERROR(cudaEventElapsedTime(&elapsedTime, start, stop));

	printf("Time taken: %3.1f ms
", elapsedTime);

	// free stream and mem
	HANDLE_ERROR(cudaFreeHost(host_a));
	HANDLE_ERROR(cudaFreeHost(host_b));
	HANDLE_ERROR(cudaFreeHost(host_c));
	HANDLE_ERROR(cudaFree(dev_a0));
	HANDLE_ERROR(cudaFree(dev_b0));
	HANDLE_ERROR(cudaFree(dev_c0));
	HANDLE_ERROR(cudaFree(dev_a1));
	HANDLE_ERROR(cudaFree(dev_b1));
	HANDLE_ERROR(cudaFree(dev_c1));
	HANDLE_ERROR(cudaStreamDestroy(stream0));
	HANDLE_ERROR(cudaStreamDestroy(stream1));
	return 0;
}

순서대로 실행 하 는 것 에 비해 두 개의 cuda 흐름 을 사용 하여 프로그램의 실행 시간 을 20ms 줄 였 습 니 다(데이터 양 이 많 지 않 기 때문에 흐름 을 사용 하 는 장점 이 뚜렷 하지 않 습 니 다).