Programming/openCL & CUDA2014.02.15 22:10

cuda 6.0 rc

아직 함정인가?
링크는 로그인 할 수 없다고 하네?
그나저나 이제야 조금 맛을 보았는데 금세 6.x 대라니 ㅠㅠ


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

cuda 6.0 rc  (0) 2014.02.15
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
Posted by 구차니
Trackback 0 Comment 0

댓글을 달아 주세요

Programming/openCL & CUDA2014.01.17 19:41

cuda + openmp 적용 예제

어제 했던 소스에서 대충 openmp 적용해봄
2014/01/16 - [Programming/openCL / CUDA] - cuda 1차원 데이터, 2차원 처리 예제

CUDA Runtime API - Host - Extra C++ Options 에서 /openmp 를 추가하면 된다.

[링크 : https://devtalk.nvidia.com/.../vs2010-cuda4rc2-howto-compile-openmp-host-code...]

물론 #inclued <omp.h>를 추가하지 않으면
컴파일 시에는 문제가 없으나 실행시에 이런 문제가 발생한다.

[링크 : http://blog.naver.com/changfull7/70110120004]

openmp 적용
cpu Time : 0.048000
gpu Time : 0.001000 

openmp 미적용
cpu Time : 0.131000
gpu Time : 0.001000  		 


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

#include < omp.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;
	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));

#pragma omp parallel for
	for(idx = 0;idx < ARRAY_SIZE ; idx++)
	{
		a[idx] = rand() & 0xFFFF;
		b[idx] = rand() & 0xFFFF;
		c[idx] = 0;
	}

	start_time = clock();

#pragma omp parallel for
	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 6.0 rc  (0) 2014.02.15
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
Posted by 구차니
Trackback 0 Comment 0

댓글을 달아 주세요

Programming/openCL & CUDA2014.01.17 11:44

cuda + openmp 프로젝트 생성은....

사용자 지정 빌드규칙을 추가하면 된다는데 빌드는 아직 안해봤고


문제는 cuda 프로젝트로 생성시


프로젝트에 디버깅 아래 C/C++이 존재하지 않는다.


아무튼.. C/C++ 의 언어에 OpenMP가 있으므로
CUDA와 OpenMP를 동시에 사용하려면 일반프로젝트 생성 + CUDA 사용자 룰 추가 이런식으로 해야 할 듯 하다.


결론 : 일단은.. openMP + CUDA는 생각하지 말자.. 귀차나 ㅠㅠ

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

cuda 6.0 rc  (0) 2014.02.15
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
Posted by 구차니
Trackback 0 Comment 0

댓글을 달아 주세요

Programming/openCL & CUDA2014.01.16 00:03

cuda 1차원 데이터, 2차원 처리 예제

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 구차니
Trackback 0 Comment 4

댓글을 달아 주세요

  1. BlogIcon Q의 성공

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

    2014.01.16 09:49 [ ADDR : EDIT/ DEL : REPLY ]

  2. 비밀댓글입니다

    2015.01.28 20:05 [ ADDR : EDIT/ DEL : REPLY ]
    • BlogIcon 구차니

      저도 여전히 헷갈리는 부분이 있지만...
      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 ]

Programming/openCL & CUDA2014.01.15 21:58

CUDA 쓰레드 계산

머리를 데굴데굴 굴려봐도 깔끔한 계산식이 안나온다..
걍 2차원 배열로 하는수 밖에 없으려나?

일단 1차원으로 할당한 녀석을 2차원 block/thread 로 구성시에 대한 수식이다.

threadIdx.x +blockIdx.x * blockDim.x
+ (gridDim.x * blockDim.x) * (blockIdx.y * blockDim.y + threadIdx.y);

근데.. 걍.. 데이터를 2차원 으로 해서 하는게 속 편할지도 모르겠네...


뻘짓 끝에 디버거로 돌려보니..
blockIdx 에는 thread block의 dimension이 들어간다....

다시 글들을 보니..
block에 대한 카운트와 인덱스는 gridDim 과 blockIdx에서
thread에 대한 카운트와 인덱스는 blockDim 과 threadIdx 에서 내장변수로 제공한다.

---
끄아아.. 벌써 세번째 글이었구나 ㅠㅠ

2011/01/16 - [Programming/openCL / CUDA] - CUDA 내장변수 - built in variable
2012/04/30 - [Programming/openCL / CUDA] - cuda 내장변수 

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

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
cuda 1 block 1차원 & 2차원 thread 예제  (0) 2014.01.14
Posted by 구차니
Trackback 0 Comment 0

댓글을 달아 주세요

Programming/openCL & CUDA2014.01.15 11:29

visual studio 2008 nsight 실행..된거 맞나?

Visual Studio 2008 SP1설치후
CUDA 5.5 재설치 - Nsight Visual Studio Edition
음.. 그런데 이녀석.. CUDA SDK에 포함되어서 굳이 Nvidia에서 가입/등록하고 안해도 되네?
물론.. SDK를 통채로 다운받아서 갈아 엎어야 하는 귀차니즘이 있긴 하겠지만..


Installed에 똭! 추가!


기본 값은 Launcher debugging enabled False이므로


True로 바꾸어 주고(위에것도 해야하려나?)


VS2008의 디버거를 사용하지 말고 그냥
Start CUDA Debugging을 누르면 지가 알아서 실행되고 프로파일링 한다.


머하다가 이게 떴더라..?


막 누르다 보니.. 정신없는데..
Application Control에 Launch를 누르면 프로그램이 실행되면서
Nsight에서 녹화하고 그 데이터를 기반으로 디버깅 정보를 출력한다.


Summary Report 아직 보는 법은 모르겠고


nsight를 설치해서인지 VS2008 자체 디버거로 실행시 kernel에서 내장 변수들이 출력된다.
근데.. 막 건너뛰는 기분..(마스터 쓰레드만 하나 출력되려나?)


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

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
cuda 1 block 1차원 & 2차원 thread 예제  (0) 2014.01.14
cuda 5.5 조금은 더 빠르게 컴파일 하기  (0) 2014.01.14
Posted by 구차니
Trackback 0 Comment 0

댓글을 달아 주세요

Programming/openCL & CUDA2014.01.15 10:41

nsight 설치... -_-

Nsight는 VS2008일 경우 SP1 이상을 요구한다

Requirement not met: Service Pack 1 for Microsoft Visual Studio 2008 was not found.


하지만 문제는 Nsight 샘플 프로젝트가 없다는 것!!!


분위기로 봐서는.. CUDA 5.5 설치시에
VS2008 SP1이 없어서 Nsight for Visual Studio 2008이 설치 되지 않으면서 생긴 문제로 생각되니..
이래저래 CUDA 5.5도 전부 다시 재설치.. 끄아아 ㅠㅠ


[링크 : https://developer.nvidia.com/install-nsight-visual-studio-edition]

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

CUDA 쓰레드 계산  (0) 2014.01.15
visual studio 2008 nsight 실행..된거 맞나?  (0) 2014.01.15
nsight 설치... -_-  (0) 2014.01.15
cuda 1 block 1차원 & 2차원 thread 예제  (0) 2014.01.14
cuda 5.5 조금은 더 빠르게 컴파일 하기  (0) 2014.01.14
GTX650 / ion devicequery  (4) 2014.01.13
Posted by 구차니
Trackback 0 Comment 0

댓글을 달아 주세요

Programming/openCL & CUDA2014.01.14 19:53

cuda 1 block 1차원 & 2차원 thread 예제

visual studio 2008 / cuda 5.5 / GTX650 기준 작성

    0 a:26326 b:14567 c:40893 == r:40893
    1 a:20769 b:29469 c:50238 == r:50238
    2 a:19293 b:19828 c:39121 == r:39121
    3 a: 5720 b:16164 c:21884 == r:21884
    4 a:10116 b:16010 c:26126 == r:26126
    5 a:24503 b: 1380 c:25883 == r:25883
    6 a: 1261 b:20500 c:21761 == r:21761
    7 a:32527 b:14265 c:46792 == r:46792
    8 a: 6165 b: 1639 c: 7804 == r: 7804
    9 a:16881 b: 7619 c:24500 == r:24500
   10 a:14636 b: 3016 c:17652 == r:17652
   11 a:20766 b: 1675 c:22441 == r:22441
   12 a:24356 b: 3886 c:28242 == r:28242
   13 a: 9279 b:15721 c:25000 == r:25000
   14 a:20744 b:   74 c:20818 == r:20818
   15 a:24023 b:17957 c:41980 == r:41980
   16 a:  399 b:19653 c:20052 == r:20052
   17 a: 9077 b: 9308 c:18385 == r:18385
   18 a:18673 b:  713 c:19386 == r:19386
   19 a:17966 b:12837 c:30803 == r:30803
   20 a:28921 b:31938 c:60859 == r:60859
   21 a:20298 b:18933 c:39231 == r:39231
   22 a:18267 b:31334 c:49601 == r:49601
   23 a:17726 b:18368 c:36094 == r:36094
   24 a:10825 b:19187 c:30012 == r:30012
   25 a:15579 b: 9569 c:25148 == r:25148
   26 a:17217 b:27831 c:45048 == r:45048
   27 a: 2756 b:13884 c:16640 == r:16640
   28 a:25641 b:17878 c:43519 == r:43519
   29 a:10533 b:17954 c:28487 == r:28487
   30 a:15005 b:23112 c:38117 == r:38117
   31 a: 9634 b: 8053 c:17687 == r:17687 		 


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

#include "cuda_runtime.h"

#define ARRAY_SIZE	32

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

void main()
{
	int a[ARRAY_SIZE],b[ARRAY_SIZE],c[ARRAY_SIZE],res[ARRAY_SIZE];
	int *dev_a,*dev_b,*dev_c;
	int idx = 0;
	dim3 block(1);
	dim3 thread(ARRAY_SIZE);

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

	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);

	kernel_test<<<block,thread>>>(dev_a,dev_b,dev_c);

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

	for(idx = 0;idx < ARRAY_SIZE ; 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]);
	}

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


#define ARRAY_SIZE	64

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

void main()
{
	dim3 block(1);
	dim3 thread(8,8);
}

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

visual studio 2008 nsight 실행..된거 맞나?  (0) 2014.01.15
nsight 설치... -_-  (0) 2014.01.15
cuda 1 block 1차원 & 2차원 thread 예제  (0) 2014.01.14
cuda 5.5 조금은 더 빠르게 컴파일 하기  (0) 2014.01.14
GTX650 / ion devicequery  (4) 2014.01.13
cuda dim3 변수 초기화  (0) 2014.01.13
Posted by 구차니
Trackback 0 Comment 0

댓글을 달아 주세요

Programming/openCL & CUDA2014.01.14 17:14

cuda 5.5 조금은 더 빠르게 컴파일 하기

vs2008의 문제인지 cuda5.5의 변경점인지 모르겠지만
컴파일시 하나의 파일에 대해서 경고가 여러번 뜨는 것 봐서는
여러번의 컴파일을 시도하는 것으로 생각되어 프로젝트 옵션을 뒤져보고 대충 테스트 해본결과

기본값으로 GPU Architecture가 sm_10과 sm_20 두개가 설정이 되어 있어 2번의 컴파일을 수행하고 있었다.
그래픽 카드의 아키텍쳐를 확실히 안하면 제대로 설정해서 한번만 컴파일 하는 것도
컴파일 시간 단축에 도움이 많이 될 듯 하다.

1,2에 sm_10 / sm_20을

 
GTX650을 쓰고 있으니 sm_30으로 설정!

 

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

nsight 설치... -_-  (0) 2014.01.15
cuda 1 block 1차원 & 2차원 thread 예제  (0) 2014.01.14
cuda 5.5 조금은 더 빠르게 컴파일 하기  (0) 2014.01.14
GTX650 / ion devicequery  (4) 2014.01.13
cuda dim3 변수 초기화  (0) 2014.01.13
vs2008 cuda syntax highlight  (0) 2014.01.13
Posted by 구차니
Trackback 0 Comment 0

댓글을 달아 주세요

Programming/openCL & CUDA2014.01.13 22:12

GTX650 / ion devicequery

가장 중요한 차이점 아래 스샷 참조
왼쪽이 ion 오른쪽이 GTX650 이다.


전체 비교


사이즈 문제로 텍스트 버전은 아래에


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

cuda 1 block 1차원 & 2차원 thread 예제  (0) 2014.01.14
cuda 5.5 조금은 더 빠르게 컴파일 하기  (0) 2014.01.14
GTX650 / ion devicequery  (4) 2014.01.13
cuda dim3 변수 초기화  (0) 2014.01.13
vs2008 cuda syntax highlight  (0) 2014.01.13
VS2008 + cuda5.5 프로젝트 생성  (0) 2014.01.13
Posted by 구차니
Trackback 0 Comment 4

댓글을 달아 주세요

  1. BlogIcon S매니저

    덕분에 잘 보고 갑니다~
    즐거운 하루 보내세요^^

    2014.01.14 09:55 [ ADDR : EDIT/ DEL : REPLY ]
  2. BlogIcon Q의 성공

    다녀갑니다 ^^
    의미있는 오늘이 되세요~

    2014.01.14 10:05 [ ADDR : EDIT/ DEL : REPLY ]
  3. BlogIcon 씨디맨

    780 ti로도 테스트 해보고 파요. 물론 게임요 ㅋ;

    2014.01.14 10:55 신고 [ ADDR : EDIT/ DEL : REPLY ]