Programming/openCL & CUDA2014. 1. 16. 00:03
Phenum2 945(3.0Ghz) / GTX650

cpu Time : 0.131000
gpu Time : 0.001000

일단.. 싱글 코어로 돌려서 130msec 정도 
GPU로 돌려서 동일 연산이 1msec 미만으로 걸린다.. ㄷㄷㄷ

일단 32*32*128*128*4 로
총 1024개의 쓰레드를 동시에 돌리고
16,384 개의 블럭으로 연산을 한다.
67,108,864 byte 대략 64MB + 64MB = 64MB 용량의 연산을 하는 소스이다.

#include < stdio.h >
#include < stdlib.h >
#include < time.h >

#include "cuda_runtime.h"

#define BLOCK_WID	128
#define THREAD_WID	32
#define ARRAY_SIZE	 (THREAD_WID * THREAD_WID * BLOCK_WID * BLOCK_WID)

__global__ void kernel_test(int *a, int *b, int *c)
{
	int idx = threadIdx.x +blockIdx.x * blockDim.x + (gridDim.x * blockDim.x) * (blockIdx.y * blockDim.y + threadIdx.y);
	c[idx] = a[idx] + b[idx];
}

void main()
{
	clock_t start_time, end_time;
	int *a, *b, *c, *res;
	int *dev_a,*dev_b,*dev_c;
	unsigned int idx = 0;
	dim3 block(BLOCK_WID,BLOCK_WID);
	dim3 thread(THREAD_WID,THREAD_WID);

	a = (int *)malloc(ARRAY_SIZE * sizeof(int));
	b = (int *)malloc(ARRAY_SIZE * sizeof(int));
	c = (int *)malloc(ARRAY_SIZE * sizeof(int));
	res = (int *)malloc(ARRAY_SIZE * sizeof(int));

	// initialize
	srand (time(NULL));
	for(idx = 0;idx < ARRAY_SIZE ; idx++)
	{
		a[idx] = rand() & 0xFFFF;
		b[idx] = rand() & 0xFFFF;
		c[idx] = 0;
	}

	start_time = clock();
	for(idx = 0;idx < ARRAY_SIZE ; idx++)
	{
		res[idx] = a[idx] + b[idx];
	}
	end_time = clock();
	printf("cpu Time : %f\n", ((double)(end_time-start_time)) / CLOCKS_PER_SEC); 

	cudaMalloc(&dev_a, ARRAY_SIZE * sizeof(int));
	cudaMalloc(&dev_b, ARRAY_SIZE * sizeof(int));
	cudaMalloc(&dev_c, ARRAY_SIZE * sizeof(int));

	cudaMemcpy(dev_a, a, ARRAY_SIZE * sizeof(int), cudaMemcpyHostToDevice);
	cudaMemcpy(dev_b, b, ARRAY_SIZE * sizeof(int), cudaMemcpyHostToDevice);

	start_time = clock();
	kernel_test<<<block,thread>>>(dev_a,dev_b,dev_c);
	end_time = clock();
	printf("gpu Time : %f\n", ((double)(end_time-start_time)) / CLOCKS_PER_SEC); 

	cudaMemcpy(c, dev_c, ARRAY_SIZE * sizeof(int), cudaMemcpyDeviceToHost);

	for(idx = 0;idx < ARRAY_SIZE ; idx++)
	{
		if(res[idx] != c[idx])
		{
			printf("%5d a:%5d b:%5d c:%5d", idx, a[idx], b[idx], c[idx]);
			if(res[idx] != c[idx])
					printf(" != ");
			else	printf(" == ");
			printf("r:%5d\n",res[idx]);
			break;
		}
	}

	cudaFree(dev_a);
	cudaFree(dev_b);
	cudaFree(dev_c);

	free(a);
	free(b);
	free(c);
	free(res);
}

'Programming > openCL & CUDA' 카테고리의 다른 글

cuda + openmp 적용 예제  (0) 2014.01.17
cuda + openmp 프로젝트 생성은....  (0) 2014.01.17
cuda 1차원 데이터, 2차원 처리 예제  (4) 2014.01.16
CUDA 쓰레드 계산  (0) 2014.01.15
visual studio 2008 nsight 실행..된거 맞나?  (0) 2014.01.15
nsight 설치... -_-  (0) 2014.01.15
Posted by 구차니

댓글을 달아 주세요

  1. 다녀갑니다^^
    의미있는 하루가 되세요~

    2014.01.16 09:49 [ ADDR : EDIT/ DEL : REPLY ]
  2. 비밀댓글입니다

    2015.01.28 20:05 [ ADDR : EDIT/ DEL : REPLY ]
    • 저도 여전히 헷갈리는 부분이 있지만...
      http://minimonk.net/2291
      http://minimonk.net/2287 을 일단 참고 하시고..

      하드웨어 적으로는 한번에 생성이 가능한 쓰레드(cuda core)는 갯수가 제한되어 있으나
      논리적으로 여러블럭으로 나누어 하드웨어를 반복 계산을 하는 식으로
      32*32 의 쓰레드 블럭을 생성하여 (1024 쓰레드) 4*4의 그리드를 형성하여
      총 128*128 개의 쓰레드 처럼 구현도 가능하기 때문에
      효율을 제외하면 어떤식으로든 상상하는 이상의 크기로 배열을 계산이 가능합니다. 물론 결합전송이라던가 이런 제한이 8800GT 에는 있었지만 650GTX만 해도 상당히 완화되고 780이나 Titan 에서는 거의 없다 싶을 정도로 좋아져서 일단은 돌리는데 집중을 하시고
      천천히 최적화를 해보시는게 좋지 않을까 합니다.

      2015.01.28 22:02 신고 [ ADDR : EDIT/ DEL ]