성태의 닷넷 이야기
홈 주인
모아 놓은 자료
프로그래밍
질문/답변
사용자 관리
사용자
메뉴
아티클
외부 아티클
유용한 코드
온라인 기능
MathJax 입력기
최근 덧글
[정성태] Roll A Lisp In C - Reading ; https...
[정성태] Java - How to use the Foreign Funct...
[정성태] 제가 큰 실수를 했군요. ^^; Delegate를 통한 Bein...
[정성태] Working with Rust Libraries from C#...
[정성태] Detecting blocking calls using asyn...
[정성태] 아쉽게도, 커뮤니티는 아니고 개인 블로그입니다. ^^
[정성태] 질문이 잘 이해가 안 됩니다. 우선, 해당 소스코드에서 ILis...
[양승조
] var대신 dinamic으로 선언해서 해결은 했습니다. 맞는 해...
[양승조
] 또 막혔습니다. ㅠㅠ var list = props[i].Ge...
[양승조
] 아. 감사합니다. 어제는 안됐던것 같은데....정신을 차려야겠네...
글쓰기
제목
이름
암호
전자우편
HTML
홈페이지
유형
제니퍼 .NET
닷넷
COM 개체 관련
스크립트
VC++
VS.NET IDE
Windows
Team Foundation Server
디버깅 기술
오류 유형
개발 환경 구성
웹
기타
Linux
Java
DDK
Math
Phone
Graphics
사물인터넷
부모글 보이기/감추기
내용
<div style='display: inline'> <h1 style='font-family: Malgun Gothic, Consolas; font-size: 20pt; color: #006699; text-align: center; font-weight: bold'>CUDA의 인덱싱 관련 용어 - blockIdx, threadIdx, blockDim, gridDim</h1> <p> blockDim과 gridDim의 차이점을 이해하기에 앞서 우선 block이 뭔지에 대해 알아야 합니다.<br /> <br /> CUDA 런타임은 N 개의 커널 복사본을 만들고 그것을 병렬로 실행할 수 있는데 여기서의 병렬로 실행되는 복사본 각각을 block이라고 합니다. 그리고 kernel 함수를 개시할 때 첫 번째 인자로 그 크기를 명시합니다.<br /> <br /> <pre style='margin: 10px 0px 10px 10px; padding: 10px 0px 10px 10px; background-color: #fbedbb; overflow: auto; font-family: Consolas, Verdana;' > kernel<<<<span style='color: blue; font-weight: bold'>2</span>,1>>>(); </pre> <br /> 즉, 위의 커널 함수 호출은 2개의 복사본(block)이 생성됩니다.<br /> <br /> 이때 복사본 각각은 자신의 Task id를 blockId.x로 구분할 수 있습니다.<br /> <br /> <pre style='margin: 10px 0px 10px 10px; padding: 10px 0px 10px 10px; background-color: #fbedbb; overflow: auto; font-family: Consolas, Verdana;' > __global__ void kernel() { int taskId = blockIdx.x; // 2개의 block이므로 [0, 1]의 범위 } </pre> <br /> 커널 함수의 첫 번째 인자로 전달하는 block의 수는 3차원까지 지정할 수 있습니다.<br /> <br /> <pre style='margin: 10px 0px 10px 10px; padding: 10px 0px 10px 10px; background-color: #fbedbb; overflow: auto; font-family: Consolas, Verdana;' > dim3 grid(2, 3); kernel<<<grid, 1>>>(); </pre> <br /> 위의 경우 총 블록의 수는 2 * 3 = 6입니다. 아울러 블록의 각 차원에 대한 최댓값은 cudaGetDeviceProperties로 구할 수 있습니다. (cudaDeviceProp.maxGridSize[3])<br /> <br /> <pre style='margin: 10px 0px 10px 10px; padding: 10px 0px 10px 10px; background-color: #fbedbb; overflow: auto; font-family: Consolas, Verdana;' > GTX 1070, GTX 960, GT 640M의 cudaGetDeviceProperties 출력 결과 ; <a target='tab' href='http://www.sysnet.pe.kr/2/0/11472'>http://www.sysnet.pe.kr/2/0/11472</a> </pre> <br /> <pre style='margin: 10px 0px 10px 10px; padding: 10px 0px 10px 10px; background-color: #fbedbb; overflow: auto; font-family: Consolas, Verdana;' > Max dimension size of a grid size (x,y,z): (2147483647, 65535, 65535) </pre> <br /> 특이하게 1차원인 경우 2147483647까지 가능하지만, y, z에 해당하는 2와 3차원은 인자 값이 65535로 제한되어 있습니다. 비록 x, y, z 차원의 곱이 상상을 초월할 정도로 커지지만 최대 생성할 수 있는 블록의 수는 (테스트해보면) 1차원에 지정된 수의 값과 같습니다. (그런데.... 정말로 블록을 21억 개나 보유가 가능할까요? 아니면 부분적으로 GPU에 의해 적당하게 나눠서 관리할까요?)<br /> <br /> 즉,<br /> <br /> <pre style='margin: 10px 0px 10px 10px; padding: 10px 0px 10px 10px; background-color: #fbedbb; overflow: auto; font-family: Consolas, Verdana;' > kernel<<<span style='color: blue; font-weight: bold'>2147483648</span>, 1>>>(); </pre> <br /> 로 실행하면 cudaErrorInvalidConfiguration(9) 오류가 발생하는 반면, 범위 내에 속하는 경우,<br /> <br /> <pre style='margin: 10px 0px 10px 10px; padding: 10px 0px 10px 10px; background-color: #fbedbb; overflow: auto; font-family: Consolas, Verdana;' > kernel<<<span style='color: blue; font-weight: bold'>2147483647</span>, 1>>>(); </pre> <br /> 정상적으로 구동이 됩니다.<br /> <br /> <hr style='width: 50%' /><br /> <br /> 지정된 모든 블록은 Grid라고 불립니다. 즉 다음의 경우,<br /> <br /> <pre style='margin: 10px 0px 10px 10px; padding: 10px 0px 10px 10px; background-color: #fbedbb; overflow: auto; font-family: Consolas, Verdana;' > dim3 grid(2, 3); kernel<<<grid, 1>>>(); </pre> <br /> 총 6개의 블록을 갖는 grid가 운영되는 것입니다. 그리고 kernel 함수 내에서 grid에 속한 블록의 최댓값을 알고 싶을 수 있는데, 바로 이때 사용하는 내부 변수가 gridDim입니다.<br /> <br /> 즉, dim3 grid(2,3)으로 구동한 경우 kernel 함수 내에서 각각의 block 차원 수를 다음과 같이 구할 수 있습니다.<br /> <br /> <pre style='margin: 10px 0px 10px 10px; padding: 10px 0px 10px 10px; background-color: #fbedbb; overflow: auto; font-family: Consolas, Verdana;' > __global__ void kernel() { int maxBlockX = <span style='color: blue; font-weight: bold'>gridDim.x</span>; // 2 int maxBlockY = <span style='color: blue; font-weight: bold'>gridDim.y</span>; // 3 } </pre> <br /> 대개의 경우 gridDim이 사용되는 사례는, 지정된 블록의 수보다 많은 요소의 데이터를 다룰 때입니다. 가령, 100개의 byte 배열에 있는 값을 2개의 블록을 이용해 전체 배열의 값을 +1 하고 싶은 경우 다음과 같이 gridDim을 사용할 수 있습니다.<br /> <br /> <pre style='margin: 10px 0px 10px 10px; padding: 10px 0px 10px 10px; background-color: #fbedbb; overflow: auto; font-family: Consolas, Verdana;' > kernel<<2, 1>>>(buf); __global__ void kernel(BYTE *srcPtr) { int taskId = blockIdx.x; while (taskId < 100) { srcPtr[taskId] = srcPtr[taskId] + 1; <span style='color: blue; font-weight: bold'>taskId = taskId + gridDim.x;</span> // gridDim.x == 2 } } </pre> <br /> 위의 경우, blockIdx.x == 0번인 block이 구동되면 while 문을 통해 0, 2, 4, 6, 8, ...에 해당하는 버퍼의 값을 변경하게 됩니다. 반면 blockIdx.x == 1번인 경우 1, 3, 5, 7, 9, ...에 해당하는 버퍼의 값을 바꿔 [0, 100) 범위의 모든 BYTE 요소를 접근하게 됩니다.<br /> <br /> <hr style='width: 50%' /><br /> <br /> block에 대해 다뤘으니, 이제 thread에 대해 알아볼 차례입니다. kernel 함수 구동에서 두 번째 인자로 전달하는 수가 바로 block 당 구동되는 스레드의 수입니다. 가령, 다음의 kernel 구동은,<br /> <br /> <pre style='margin: 10px 0px 10px 10px; padding: 10px 0px 10px 10px; background-color: #fbedbb; overflow: auto; font-family: Consolas, Verdana;' > kernel<<2, <span style='color: blue; font-weight: bold'>3</span>>>>(); </pre> <br /> 2개의 커널 복사본(block)이 할당되고, 각각의 block에 대해 3개의 스레드가 실행하는 것으로 이렇게 되면 작업이 총 6개가 됩니다. <br /> <br /> block에 대해 몇 번째 스레드가 구동되는지는 내부 변수인 threadIdx를 통해 가능합니다. 따라서 6개의 요소를 가진 배열의 값을 +1 하고 싶다면 다음과 같이 커널 함수를 만들 수 있습니다.<br /> <br /> <pre style='margin: 10px 0px 10px 10px; padding: 10px 0px 10px 10px; background-color: #fbedbb; overflow: auto; font-family: Consolas, Verdana;' > kernel<<<span style='color: blue; font-weight: bold'>1, 6</span>>>>(buf); __global__ void kernel(BYTE *srcPtr) { int taskId = <span style='color: blue; font-weight: bold'>threadIdx.x</span>; srcPtr[taskId] = srcPtr[taskId] + 1; } </pre> <br /> block의 수를 gridDim 내부 변수로 구할 수 있었던 것처럼, block 당 스레드 수 역시 내부 변수가 제공되는데 그것이 바로 blockDim입니다. blockDim이 사용되는 사례도 gridDim과 유사합니다. 즉, 생성된 스레드보다 많은 요소의 데이터를 접근할 때인데 가령 위의 예제 코드에서 전달되는 buffer의 요소가 총 12개라면 이에 대해 +1을 하기 위해 다음과 같이 blockDim을 활용할 수 있습니다.<br /> <br /> <pre style='margin: 10px 0px 10px 10px; padding: 10px 0px 10px 10px; background-color: #fbedbb; overflow: auto; font-family: Consolas, Verdana;' > kernel<<1, 6>>>(buf); __global__ void kernel(BYTE *srcPtr) { int taskId = threadIdx.x; while (taskId < 12) { srcPtr[taskId] = srcPtr[taskId] + 1; taskId = taskId + <span style='color: blue; font-weight: bold'>blockDim.x</span>; // blockDim.x == 6 } } </pre> <br /> 위의 커널 코드는 1개의 블록으로 6개의 스레드가 실행되고 있기 때문에 각각의 스레드는 다음과 같이 taskId가 while 루프를 통해 증가하면서 작업을 처리합니다.<br /> <br /> <pre style='margin: 10px 0px 10px 10px; padding: 10px 0px 10px 10px; background-color: #fbedbb; overflow: auto; font-family: Consolas, Verdana;' > thread 0: 0, 6 thread 1: 1, 7 thread 2: 2, 8 thread 3: 3, 9 thread 4: 4, 10 thread 5: 5, 11 </pre> <br /> 참고로 블록 당 지정 가능한 스레드의 최대 수는 cudaGetDeviceProperties로 구할 수 있습니다. (cudaDeviceProp.maxThreadsPerBlock, cudaDeviceProp.maxThreadsDim[3])<br /> <br /> <pre style='margin: 10px 0px 10px 10px; padding: 10px 0px 10px 10px; background-color: #fbedbb; overflow: auto; font-family: Consolas, Verdana;' > Maximum number of threads per block: 1024 Max dimension size of a thread block (x,y,z): (1024, 1024, 64) </pre> <br /> 따라서 위의 상황에서 다음과 같은 커널은 정상적으로 구동되지만,<br /> <br /> <pre style='margin: 10px 0px 10px 10px; padding: 10px 0px 10px 10px; background-color: #fbedbb; overflow: auto; font-family: Consolas, Verdana;' > dim3 threads(1, 512, 2); // 1 * 512 * 2 = 1024 kernel <<<1, threads>>>(); </pre> <br /> 1024를 넘기면 cudaErrorInvalidConfiguration 오류가 발생합니다.<br /> <br /> <pre style='margin: 10px 0px 10px 10px; padding: 10px 0px 10px 10px; background-color: #fbedbb; overflow: auto; font-family: Consolas, Verdana;' > dim3 threads(1, 512, 3); // 1 * 512 * 3 = 1536 kernel <<<1, threads>>>(); cudaError_t cudaStatus = cudaGetLastError(); // cudaErrorInvalidConfiguration </pre> <br /> <hr style='width: 50%' /><br /> <br /> 지금까지의 예제에서는 block 또는 thread의 한 쪽을 1로 설정했었는데, 그렇지 않은 경우를 보겠습니다.<br /> <br /> <pre style='margin: 10px 0px 10px 10px; padding: 10px 0px 10px 10px; background-color: #fbedbb; overflow: auto; font-family: Consolas, Verdana;' > kernel <<<2, 2>>>(); </pre> <br /> 위의 커널은 2개의 블록에 대해 각각 2개의 스레드로 실행됩니다. 이런 경우 kernel 함수 내에서 작업 ID는 다음과 같이 구할 수 있습니다.<br /> <br /> <pre style='margin: 10px 0px 10px 10px; padding: 10px 0px 10px 10px; background-color: #fbedbb; overflow: auto; font-family: Consolas, Verdana;' > __global__ void kernel() { int taskId = <span style='color: blue; font-weight: bold'>threadIdx.x + blockIdx.x * blockDim.x</span>; } </pre> <br /> 계산해 보면 다음과 같이 [0, 3]까지의 작업 id가 구해지는 것을 확인할 수 있습니다.<br /> <br /> <pre style='margin: 10px 0px 10px 10px; padding: 10px 0px 10px 10px; background-color: #fbedbb; overflow: auto; font-family: Consolas, Verdana;' > thread\block 0 1 0 0 + 0 * 2 = <span style='color: blue; font-weight: bold'>0</span> 0 + 1 * 2 = <span style='color: blue; font-weight: bold'>2</span> 1 1 + 0 * 2 = <span style='color: blue; font-weight: bold'>1</span> 1 + 1 * 2 = <span style='color: blue; font-weight: bold'>3</span> </pre> <br /> 따라서 위의 작업 id와 같은 인덱싱 방식으로 4개의 요소를 가진 데이터를 병렬 처리할 수 있습니다. 그 이상의 데이터, 예를 들어 10개의 요소를 가진 배열을 처리하고 싶다면 다음과 같이 gridDim 내부 변수를 함께 사용하면 됩니다.<br /> <br /> <pre style='margin: 10px 0px 10px 10px; padding: 10px 0px 10px 10px; background-color: #fbedbb; overflow: auto; font-family: Consolas, Verdana;' > __global__ void kernel(BYTE *srcPtr) { int taskId = threadIdx.x + blockIdx.x * blockDim.x; while (taskId < 10) { srcPtr[taskId] = srcPtr[taskId] + 1; taskId = taskId + <span style='color: blue; font-weight: bold'>(blockDim.x * gridDim.x)</span>; // (blockDim.x = 2, gridDim.x = 2) == 4 } } </pre> <br /> 그럼, 블록 별 스레드에 대한 인덱싱은 다음과 같이 바뀌면서 작업을 처리하게 됩니다.<br /> <br /> <pre style='margin: 10px 0px 10px 10px; padding: 10px 0px 10px 10px; background-color: #fbedbb; overflow: auto; font-family: Consolas, Verdana;' > // 첫 번째 루프 thread\block 0 1 0 0 + 0 * 2 = 0 0 + 1 * 2 = 2 1 1 + 0 * 2 = 1 1 + 1 * 2 = 3 // 두 번째 루프 thread\block 0 1 0 4 6 1 5 7 // 세 번째 루프 thread\block 0 1 0 8 (루프탈출) 1 9 (루프탈출) </pre> <br /> 점점 더 복잡해지죠? ^^ 게다가 block과 thread에 대해 2차원 이상으로 늘릴 수도 있으니... 연산이 정확한지 신중을 기해 살펴봐야 합니다.<br /> <br /> <hr style='width: 50%' /><br /> <br /> 참고로, <a target='tab' href='http://www.yes24.com/24/goods/5782880?scode=032'>예제로 배우는 CUDA 프로그래밍</a>에 언급된 것에 의하면, 어떤 성능 측정의 경우 블록의 수가 멀티 프로세서 수의 정확히 2배가 되었을 때 최고의 성능을 발휘한다고 합니다. 그럼, cudaGetDeviceProperties로 MP 값(cudaDeviceProp.multiProcessorCount)을 구한 후,<br /> <br /> <pre style='margin: 10px 0px 10px 10px; padding: 10px 0px 10px 10px; background-color: #fbedbb; overflow: auto; font-family: Consolas, Verdana;' > (15) Multiprocessors, (128) CUDA Cores/MP: 1920 CUDA Cores </pre> <br /> 15 * 2 = 30 정도의 값이 블록의 수로 적당하다는 이야기가 됩니다. 그런데 아쉽게도 블록 당 스레드의 적절한 값에 대해서는 언급이 없는데요. 일단 GPU는,<br /> <br /> <pre style='margin: 10px 0px 10px 10px; padding: 10px 0px 10px 10px; background-color: #fbedbb; overflow: auto; font-family: Consolas, Verdana;' > Maximum number of threads per multiprocessor: 2048 </pre> <br /> 라고 cudaGetDeviceProperties를 통해 MP 당 지원가능한 스레드의 수(cudaDeviceProp.maxThreadsPerMultiProcessor)를 구할 수 있습니다. 따라서 위의 경우에는, 15 * 2048 = 30,720개의 동시 스레드를 제공할 수 있다는 것이 됩니다. 물론, 블록 당 최대치는 cudaDeviceProp.maxThreadsPerBlock으로 정해집니다. 아마도 스레드의 적절한 수는 로직에 따라 성능 테스트를 하고 결정하는 것이 좋을 듯 합니다.<br /> <br /> <hr style='width: 50%' /><br /> <br /> 자, 그럼 정리해 볼까요? ^^<br /> <br /> 우선, 여러분은 CUDA에게 grid 하나를 정의해야 합니다. grid는 1 ~ 3차원 block들의 묶음입니다. 그리고 개별 block들에서는 자신의 인덱스를 blockIdx 내부 변수로 알아낼 수 있고 grid의 크기를 gridDim 내부 변수를 통해 접근할 수 있습니다.<br /> <br /> 또한 하나의 block들은 1 ~ 3차원 thread들의 컨테이너 역할을 합니다. 개별 thread들에서는 자신의 인덱스를 threadIdx 내부 변수로 알아낼 수 있고 block이 소유한 스레드의 크기를 blockDim 내부 변수를 통해 접근할 수 있습니다.<br /> <br /> 이들을 한데 엮어서 그림으로 표현하면 아래와 같습니다.<br /> <br /> [Grid of Thread Blocks - 출처: <a target='tab' href='http://docs.nvidia.com/cuda/cuda-c-programming-guide/'>http://docs.nvidia.com/cuda/cuda-c-programming-guide/</a>]<br /> <img alt='grid-of-thread-blocks.png' src='/SysWebRes/bbs/grid-of-thread-blocks.png' /><br /> <br /> 사실 제가 혼란스러웠던 것이 gridDim, blockDim 변수명입니다. 차라리 다음과 같이 이름 지었으면,<br /> <br /> <pre style='margin: 10px 0px 10px 10px; padding: 10px 0px 10px 10px; background-color: #fbedbb; overflow: auto; font-family: Consolas, Verdana;' > gridDim => blockDim blockDim => threadDim </pre> <br /> 더 잘 이해가 되었을지도 모르겠습니다<br /> <br /> <hr style='width: 50%' /><br /> <br /> 참고로, CUDA 메모리에 관련된 값도 좀 알아볼까요? 우선 cudaGetDeviceProperties를 통해 알아낼 수 있는 정보는 다음과 같습니다. (cudaDeviceProp.totalGlobalMem)<br /> <br /> <pre style='margin: 10px 0px 10px 10px; padding: 10px 0px 10px 10px; background-color: #fbedbb; overflow: auto; font-family: Consolas, Verdana;' > Total amount of global memory: 4096 MBytes (4294967296 bytes) Maximum memory pitch: 2147483647 bytes </pre> <br /> cudaMalloc에 4GB(4294967295)를 전달하면 cudaErrorMemoryAllocation(2) 오류가 반환됩니다. 그럴 수 있습니다, 왜냐하면 다른 프로그램에서도 사용하고 있을 것이기 때문에 4GB가 온전히 남아 있을 리는 없기 때문입니다. CUDA의 가용 메모리는 cudaMemGetInfo 함수를 통해 알아낼 수 있습니다.<br /> <br /> <pre style='margin: 10px 0px 10px 10px; padding: 10px 0px 10px 10px; background-color: #fbedbb; overflow: auto; font-family: Consolas, Verdana;' > cudaMemGetInfo ; <a target='tab' href='https://www.cs.cmu.edu/afs/cs/academic/class/15668-s11/www/cuda-doc/html/group__CUDART__MEMORY_gd5d6772f4b2f3355078ecd6059e6aa74.html'>https://www.cs.cmu.edu/afs/cs/academic/class/15668-s11/www/cuda-doc/html/group__CUDART__MEMORY_gd5d6772f4b2f3355078ecd6059e6aa74.html</a> </pre> <br /> 재미있는 것은, 가용 메모리 내에서 1 ~ 3GB를 cudaMalloc으로 메모리를 할당받는 것과 cudaMemcpy 시에 cudaMemcpyHostToDevice 옵션으로 이동하는 것은 잘 되지만, cudaMemcpyDeviceToHost 단계에서는 cudaErrorLaunchFailure(4)가 떨어진다는 점입니다. (원인은 잘 모르겠습니다. 혹시 아시는 분 덧글 부탁드립니다.)<br /> <br /> 따라서 (4GB 정도의 GPU 메모리 환경에서) GB 단위의 메모리 할당은 현실적으로 피하는 것이 좋겠습니다.<br /> <br /> 할당에 성공만 한다면 256MB (268,435,456) 정도의 버퍼를 다룬다고 했을 때, GPU block의 수가 2,147,483,647까지 가능하기 때문에 kernel 함수 내부에서 별도의 루프 없이 연산할 수 있습니다. 따라서 메모리에 대한 인덱스를 blockIdx.x만으로 설정하는 것이 가능합니다.<br /> <br /> <pre style='margin: 10px 0px 10px 10px; padding: 10px 0px 10px 10px; background-color: #fbedbb; overflow: auto; font-family: Consolas, Verdana;' > __global__ void incrementArray(BYTE *srcPtr) // BYTE srcPtr[1024 * 1024 * 256]; { int taskId = blockIdx.x; srcPtr[taskId] = srcPtr[taskId] + 1; } </pre> <br /> 물론 위에서도 언급했지만 이상적인 블록의 수는 MP의 두 배이기 때문에 2147483647로 지정하는 경우가 많지는 않을 것입니다. 단지, 커널 프로그램을 직관적으로 쉽게 다루고 싶다면 저것도 하나의 방법이 될 수 있는 정도입니다.<br /> <br /> <hr style='width: 50%' /><br /> <br /> CUDA의 block과 thread를 이해했다면 이제 __shared__ 예약어로 지정한 공유 메모리도 이해할 수 있습니다.<br /> <br /> __shared__ 메모리는 block 단위로 할당된 메모리이므로, block 당 생성되는 thread들에 의해 공유가 되고 서로 다른 block 간에는 공유되지 않는 메모리입니다. 참고로 cudaGetDeviceProperties로 블록 당 공유 메모리의 최댓값(cudaDeviceProp.sharedMemPerBlock)을 구할 수 있습니다.<br /> <br /> <pre style='margin: 10px 0px 10px 10px; padding: 10px 0px 10px 10px; background-color: #fbedbb; overflow: auto; font-family: Consolas, Verdana;' > Total amount of shared memory per block: 49152 bytes </pre> <br /> 또한 상수 메모리(cudaDeviceProp.totalConstMem)가 있는데,<br /> <br /> <pre style='margin: 10px 0px 10px 10px; padding: 10px 0px 10px 10px; background-color: #fbedbb; overflow: auto; font-family: Consolas, Verdana;' > Total amount of constant memory: 65536 bytes </pre> <br /> CUDA에서는 Warp라는 단위로 32개의 스레드를 관리하는데, 그것의 절반인 half-warp, 즉 16개로 관리되는 스레드 중의 하나가 상수 메모리에 읽기 요청을 하는 경우 캐시가 되어 같은 half-warp에 속하는 또 다른 스레드들이 연이어 읽기 요청을 하는 경우 일기 속도가 향상이 됩니다. 이론적으로는 전역 메모리를 사용하는 경우보다 16분의 1까지 읽기 성능이 나아질 수 있습니다.<br /> </p><br /> <br /><hr /><span style='color: Maroon'>[이 글에 대해서 여러분들과 의견을 공유하고 싶습니다. 틀리거나 미흡한 부분 또는 의문 사항이 있으시면 언제든 댓글 남겨주십시오.]</span> </div>
첨부파일
스팸 방지용 인증 번호
9945
(왼쪽의 숫자를 입력해야 합니다.)