CUBLAS performance improved 50% to 300% on Fermi architecture
GPUs, for matrix multiplication of all datatypes and transpose
variations
CUFFT performance tuned for radix-3, -5, and -7 transform sizes on Fermi architecture GPUs, now 2x to 10x faster than MKL
New CUSPARSE library of GPU-accelerated sparse matrix routines for
sparse/sparse and dense/sparse operations delivers 5x to 30x faster
performance than MKL
New CURAND library of GPU-accelerated random number generation (RNG)
routines, supporting Sobol quasi-random and XORWOW pseudo-random
routines at 10x to 20x faster than similar routines in MKL
H.264 encode/decode libraries now included in the CUDA Toolkit
deviceQuery 예문은 말그대로 장치에 질의의 던져 어떤 스펙인지
몇개의 thread가 존재하는지를 파악하는 프로그램이다.
아무튼 NVIDIA_CUDA_C_ProgrammingGuide.pdf 문서를 보면 아래의 내용이 나오니 참고를 해서 보자면,
Geforce 8800GT(이하 8800GT)는 14개의 멀티 프로세서를 가졌고 112개의 코어를 지녔다.
Grid는 Block을 포함하고, Block은 Thread를 포함한다.
단순 계산으로는 하나의 프로세서별로 8개의 코어가 존재하며
멀티프로세서는 14개 코어는 총 112개가 존재한다.
그리드(Grid)는 블럭의 2차원 배열로 존재하고,
블럭(Block)은 쓰레드의 2차원 배열로 존재한다.
쓰레드(Thread)는 일을하는 최소단위이다.
한번에 묶이는 최소 쓰레드의 숫자(Warp size)는 32개 이며
하나의 블럭으로 묶을수 있는 최대 쓰레드는 512개 이다.
블럭의 최대 차원은 3차원 512x512x64 이며
그리드의 최대 차원은 2차원 65535x65535x1 이다.
라고 이해하면 되려나?
A multithreaded program is partitioned into blocks of threads that execute independently from each
other, so that a GPU with more cores will automatically execute the program in less time than a GPU
with fewer cores.
=> 멀티쓰레드화된 프로그램은 서로 독립적으로 실행되는 쓰레드의 블럭으로 나누어지며, 그러한 이유로 더욱 많은 코어를 포함하는 CPU는 적은 코어를 지닌 GPU보다 짧은 시간에 프로그램을 실행할 수 있다.
Host는 일반적인 CPU 환경이고, Device는 GPU 환경이다.
컴파일시에 Device 코드만 nvcc에서 처리하고, 나머지는 일반적인 컴파일러에서 처리하는 이원화된 구조이다.
아래는 커널이다. 디바이스 코드를 생성하는 부분이며
MatAdd 함수는 __global__ 선언을 앞에 붙여 device 코드임을 명시해야 한다.
블럭당 쓰레디의 크기를 16x16 thread로, 그리드를 N/16개로 분할하는 예제이다.
물론 2차원은 1차원과 같이 사용이 가능하므로 int형도 허용하는 듯?
B.15 Execution Configuration
Any call to a __global__ function must specify the execution configuration for that call. The execution configuration defines the dimension of the grid and blocks that will be used to execute the function on the device, as well as the associated stream (see Section 3.3.10.1 for a description of streams).
When using the driver API, the execution configuration is specified through a series of driver function calls as detailed in Section 3.3.3.
When using the runtime API (Section 3.2), the execution configuration is specified by inserting an expression of the form <<< Dg, Db, Ns, S >>> between the function name and the parenthesized argument list, where:
Dg is of type dim3 (see Section B.3.2) and specifies the dimension and size of the grid,
such that Dg.x * Dg.y equals the number of blocks being launched; Dg.z must be equal to 1;
Db is of type dim3 (see Section B.3.2) and specifies the dimension and size of each block,
such that Db.x * Db.y * Db.z equals the number of threads per block;
Ns is of type size_t and specifies the number of bytes in shared memory that is dynamically allocated per block for this call in addition to the statically allocated memory; this dynamically allocated memory is used by any of the variables declared as an external array as mentioned in Section B.2.3; Ns is an optional argument which defaults to 0;
S is of type cudaStream_t and specifies the associated stream; S is an optional argument which defaults to 0.
As an example, a function declared as
__global__ void Func(float* parameter);
must be called like this:
Func<<< Dg, Db, Ns >>>(parameter);
The arguments to the execution configuration are evaluated before the actual function arguments and like the function arguments, are currently passed via shared memory to the device. The function call will fail if Dg or Db are greater than the maximum sizes allowed for the device as specified in Appendix G, or if Ns is greater than the maximum
amount of shared memory available on the device, minus the amount of shared memory required for static allocation, functions arguments (for devices of compute capability 1.x), and execution configuration.
아무튼 원래는 아래의 예제 내용을 분석하기 위한 내용인데
점점 미궁으로 빠지는 느낌 -_-
CUDA Device Query (Runtime API) version (CUDART static linking)
There is 1 device supporting CUDA
Device 0: "GeForce 8800 GT" CUDA Driver Version: 3.20 CUDA Runtime Version: 3.10 CUDA Capability Major revision number: 1 CUDA Capability Minor revision number: 1 Total amount of global memory: 536543232 bytes Number of multiprocessors: 14 Number of cores: 112 Total amount of constant memory: 65536 bytes Total amount of shared memory per block: 16384 bytes Total number of registers available per block: 8192 Warp size: 32 Maximum number of threads per block: 512 Maximum sizes of each dimension of a block: 512 x 512 x 64 Maximum sizes of each dimension of a grid: 65535 x 65535 x 1 Maximum memory pitch: 2147483647 bytes Texture alignment: 256 bytes Clock rate: 1.50 GHz Concurrent copy and execution: Yes Run time limit on kernels: Yes Integrated: No Support host page-locked memory mapping: Yes Compute mode: Default (multiple host threads can use this device simultaneously) Concurrent kernel execution: No Device has ECC support enabled: No
deviceQuery, CUDA Driver = CUDART, CUDA Driver Version = 3.20, CUDA Runtime Vers ion = 3.10, NumDevs = 1, Device = GeForce 8800 GT
PASSED
Press <Enter> to Quit... -----------------------------------------------------------
CL_DEVICE_COMPUTE_CAPABILITY_NV: 1.1 NUMBER OF MULTIPROCESSORS: 14 NUMBER OF CUDA CORES: 112 CL_DEVICE_REGISTERS_PER_BLOCK_NV: 8192 CL_DEVICE_WARP_SIZE_NV: 32 CL_DEVICE_GPU_OVERLAP_NV: CL_TRUE CL_DEVICE_KERNEL_EXEC_TIMEOUT_NV: CL_TRUE CL_DEVICE_INTEGRATED_MEMORY_NV: CL_FALSE CL_DEVICE_PREFERRED_VECTOR_WIDTH_<t> CHAR 1, SHORT 1, INT 1, LONG 1, FLOAT 1, DOUBLE 0
--------------------------------- 2D Image Formats Supported (71) --------------------------------- # Channel Order Channel Type
경로설정이 잘못되었나 했는데, gcc 호환성으로 인해 구버전을 쓰라던 이야기가 떠오르게 하는 아래의 내용 -_-
NVIDIA Cuda ¶
Before running the Makefile, you will need to install gcc 4.3 and g++ 4.3. This is because the NVIDIA Cuda SDK 3.0 has not yet worked with gcc 4.0 and g++ 4.0. There should be no issue compiling cuda files with gcc 4.3 and g++ 4.3 on newer NVIDIA Cuda SDK versions. For a successful compilation, please follow these steps:
...
3) Create a directory and create symlinks to gcc-4.3/g++-4.3
$ gcc --version
gcc (Ubuntu 4.4.3-4ubuntu5) 4.4.3
Copyright (C) 2009 Free Software Foundation, Inc.
This is free software; see the source for copying conditions. There is NO
warranty; not even for MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.
$ g++ --version
g++ (Ubuntu 4.4.3-4ubuntu5) 4.4.3
Copyright (C) 2009 Free Software Foundation, Inc.
This is free software; see the source for copying conditions. There is NO
warranty; not even for MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.
일단.. 어거지로
~/NVIDIA_GPU_Computing_SDK/C 에서
make를 하니 어느정도 컴파일을 하는데 GLU 어쩌구 하면서 중단 OTL
libcudart.so.3: cannot open shared object file: No such file or directory
요런 에러가 발생하면
단순하게 LD_LIBRARY_PATH를 정해주고, sudo ldconfig 를 해서는 해결이 되지 않았다. /etc/ld.so.conf.d/ 에 libcuda.conf를 만들고 cuda 설치 경로인
/usr/local/cuda/lib
를 넣어주고 나서 sudo ldconfig를 해야 제대로 설정이 되었다.
2010.12.07 추가
아래의 경로에서 libcutil_i386.a 발견! (별 의미는 없음)
NVCC는 NV(Nvidia)CC(C Compiler) 인데, 구조적으로 아래와 같은 컴파일 과정을 거친다.
호스트 코드는 일반적인 C 컴파일러(예를 들면 비쥬얼 스튜디오 커맨드 라인이나 gcc)로 컴파일을 떠넘기고
nvcc는 머신코드(CUDA device용 PTX)를 생성한다.
즉, 어떠한 코드를 컴파일 하는데 있어 nvcc만으로는 독립적으로 컴파일이 진행될수 없다.
그런 이유로 윈도우에서는 Visual Studio에 빌붙고, 리눅스에서는 gcc에 빌붙는다.
nvcc의 목적에 나온 내용으로, CUDA가 아닌 내용은 범용 C 컴파일러로 투척(forward)한다고 되어있고
윈도우에서는 MS Visual Studio의 cl을 실행(instance)하여 사용한다고 되어있다.
Purpose of nvcc
This compilation trajectory involves several splitting, compilation, preprocessing,
and merging steps for each CUDA source file, and several of these steps are subtly
different for different modes of CUDA compilation (such as compilation for device
emulation, or the generation of device code repositories). It is the purpose of the
CUDA compiler driver nvcc to hide the intricate details of CUDA compilation from
developers. Additionally, instead of being a specific CUDA compilation driver,
nvcc mimics the behavior of the GNU compiler gcc: it accepts a range of
conventional compiler options, such as for defining macros and include/library
paths, and for steering the compilation process. All non-CUDA compilation steps are forwarded to a general purpose C compiler that is supported by nvcc, and on Windos platforms, where this compiler is an instance of the Microsoft Visual Studio compiler, nvcc will translate its options into appropriate ‘cl’ command syntax. This
extended behavior plus ‘cl’ option translation is intended for support of portable
application build and make scripts across Linux and Windows platforms.
그리고 내 컴퓨터에는 일단..
Visual Studio 6.0이 설치되어 있고, 개인적인 .net 거부반응으로 인해 2002나 2008 이런 녀석들은 설치되어 있지 않다.
아무튼, host compiler에서 Windows platform은
"Microsoft Visual Studio compiler, cl" 이라고 되어 있는디..
VS2002 부터 지원하는지는 모르겠지만 아무튼, cl은 command line이라고
clcc.exe 같은 녀석으로 지원하는 커맨드 라인 MSVS 컴파일러 이다.
혹시나 openCL인줄 알았더니 그것도 아니네 -_-
그리고 Supported build enviroment 에서는 Windows + MinGW shell이 존재한다.
gcc가 아니다 shell 이다 -_- 즉, 죽어도 컴파일러는 Visual Studio를 설치할 수 밖에 없다(윈도우에서는)
컴파일러가 ATI 쪽만 인식하도록 되어있는지, nVidia의 GPU를 제대로 활용하지는 못한다.
아니면 예제 프로그램들이 openCL을 이용하기는 하지만, nVidia의 openCL과는 달라서 그럴려나?
Number of platforms: 1 Platform Profile: FULL_PROFILE Platform Version: OpenCL 1.1 ATI-Stream-v2.2 (302) Platform Name: ATI Stream Platform Vendor: Advanced Micro Devices, Inc. Platform Extensions: cl_khr_icd cl_amd_event_callback cl_khr_d3d10_sharing
Platform Name: ATI Stream Number of devices: 2 Device Type: CL_DEVICE_TYPE_CPU Device ID: 4098 Max compute units: 4 Max work items dimensions: 3 Max work items[0]: 1024 Max work items[1]: 1024 Max work items[2]: 1024 Max work group size: 1024 Preferred vector width char: 16 Preferred vector width short: 8 Preferred vector width int: 4 Preferred vector width long: 2 Preferred vector width float: 4 Preferred vector width double: 0 Max clock frequency: 2393Mhz Address bits: 32 Max memory allocation: 536870912 Image support: No Max size of kernel argument: 4096 Alignment (bits) of base address: 1024 Minimum alignment (bytes) for any datatype: 128 Single precision floating point capability Denorms: Yes Quiet NaNs: Yes Round to nearest even: Yes Round to zero: Yes Round to +ve and infinity: Yes IEEE754-2008 fused multiply-add: No Cache type: Read/Write Cache line size: 64 Cache size: 32768 Global memory size: 1073741824 Constant buffer size: 65536 Max number of constant args: 8 Local memory type: Global Local memory size: 32768 Profiling timer resolution: 427 Device endianess: Little Available: Yes Compiler available: Yes Execution capabilities: Execute OpenCL kernels: Yes Execute native function: Yes Queue properties: Out-of-Order: No Profiling : Yes Platform ID: 00C3D40C Name: Intel(R) Core(TM) i5 CPU M 450 @ 2.40GHz Vendor: GenuineIntel Driver version: 2.0 Profile: FULL_PROFILE Version: OpenCL 1.1 ATI-Stream-v2.2 (302) Extensions: cl_amd_fp64 cl_khr_global_int32_base_atomics cl_khr_global_int32_extended_atomics cl_khr_local_int32_base_atomics cl_khr_local_int32_extended_atomics cl_khr_byte_addressable_store cl_khr_gl_sharing cl_ext_device_fission cl_amd_device_attribute_query cl_amd_printf cl_khr_d3d10_sharing Device Type: CL_DEVICE_TYPE_GPU Device ID: 4098 Max compute units: 2 Max work items dimensions: 3 Max work items[0]: 128 Max work items[1]: 128 Max work items[2]: 128 Max work group size: 128 Preferred vector width char: 16 Preferred vector width short: 8 Preferred vector width int: 4 Preferred vector width long: 2 Preferred vector width float: 4 Preferred vector width double: 0 Max clock frequency: 720Mhz Address bits: 32 Max memory allocation: 134217728 Image support: No Max size of kernel argument: 1024 Alignment (bits) of base address: 32768 Minimum alignment (bytes) for any datatype: 128 Single precision floating point capability Denorms: No Quiet NaNs: Yes Round to nearest even: Yes Round to zero: Yes Round to +ve and infinity: Yes IEEE754-2008 fused multiply-add: Yes Cache type: None Cache line size: 0 Cache size: 0 Global memory size: 268435456 Constant buffer size: 65536 Max number of constant args: 8 Local memory type: Global Local memory size: 16384 Profiling timer resolution: 1 Device endianess: Little Available: Yes Compiler available: Yes Execution capabilities: Execute OpenCL kernels: Yes Execute native function: No Queue properties: Out-of-Order: No Profiling : Yes Platform ID: 00C3D40C Name: ATI RV710 Vendor: Advanced Micro Devices, Inc. Driver version: CAL 1.4.838 Profile: FULL_PROFILE Version: OpenCL 1.0 ATI-Stream-v2.2 (302) Extensions: cl_khr_icd cl_khr_gl_sharing cl_amd_device_attribute_query cl_khr_d3d10_sharing
Passed!
Number of platforms: 2 Platform Profile: FULL_PROFILE Platform Version: OpenCL 1.0 CUDA 3.2.1 Platform Name: NVIDIA CUDA Platform Vendor: NVIDIA Corporation Platform Extensions: cl_khr_byte_addressable_store cl_khr_icd cl_khr_gl_sharing cl_nv_d3d9_sharing cl_nv_compiler_options cl_nv_device_attribute_query cl_nv_pragma_unroll Platform Profile: FULL_PROFILE Platform Version: OpenCL 1.1 ATI-Stream-v2.2 (302) Platform Name: ATI Stream Platform Vendor: Advanced Micro Devices, Inc. Platform Extensions: cl_khr_icd cl_amd_event_callback
Platform Name: NVIDIA CUDA Number of devices: 2 Device Type: CL_DEVICE_TYPE_GPU Device ID: 4318 Max compute units: 4 Max work items dimensions: 3 Max work items[0]: 512 Max work items[1]: 512 Max work items[2]: 64 Max work group size: 512 Preferred vector width char: 1 Preferred vector width short: 1 Preferred vector width int: 1 Preferred vector width long: 1 Preferred vector width float: 1 Preferred vector width double: 0 Max clock frequency: 1350Mhz Address bits: 5347096844566560 Max memory allocation: 134217728 Image support: Yes Max number of images read arguments: 128 Max number of images write arguments: 8 Max image 2D width: 4096 Max image 2D height: 32768 Max image 3D width: 2048 Max image 3D height: 2048 Max image 3D depth: 2048 Max samplers within kernel: 16 Max size of kernel argument: 4352 Alignment (bits) of base address: 2048 Minimum alignment (bytes) for any datatype: 128 Single precision floating point capability Denorms: No Quiet NaNs: Yes Round to nearest even: Yes Round to zero: Yes Round to +ve and infinity: Yes IEEE754-2008 fused multiply-add: Yes Cache type: None Cache line size: 0 Cache size: 0 Global memory size: 268107776 Constant buffer size: 65536 Max number of constant args: 9 Local memory type: Scratchpad Local memory size: 16384 Profiling timer resolution: 1000 Device endianess: Little Available: Yes Compiler available: Yes Execution capabilities: Execute OpenCL kernels: Yes Execute native function: No Queue properties: Out-of-Order: Yes Profiling : Yes Platform ID: 003E8750 Name: GeForce 8600 GT Vendor: NVIDIA Corporation Driver version: 260.99 Profile: FULL_PROFILE Version: OpenCL 1.0 CUDA Extensions: cl_khr_byte_addressable_store cl_khr_icd cl_khr_gl_sharing cl_nv_d3d9_sharing cl_nv_compiler_options cl_nv_device_attribute_query cl_nv_pragma_unroll cl_khr_global_int32_base_atomics cl_khr_global_int32_extended_atomics Device Type: CL_DEVICE_TYPE_GPU Device ID: 4318 Max compute units: 4 Max work items dimensions: 3 Max work items[0]: 512 Max work items[1]: 512 Max work items[2]: 64 Max work group size: 512 Preferred vector width char: 1 Preferred vector width short: 1 Preferred vector width int: 1 Preferred vector width long: 1 Preferred vector width float: 1 Preferred vector width double: 0 Max clock frequency: 1188Mhz Address bits: 5347096844566560 Max memory allocation: 134217728 Image support: Yes Max number of images read arguments: 128 Max number of images write arguments: 8 Max image 2D width: 4096 Max image 2D height: 32768 Max image 3D width: 2048 Max image 3D height: 2048 Max image 3D depth: 2048 Max samplers within kernel: 16 Max size of kernel argument: 4352 Alignment (bits) of base address: 2048 Minimum alignment (bytes) for any datatype: 128 Single precision floating point capability Denorms: No Quiet NaNs: Yes Round to nearest even: Yes Round to zero: Yes Round to +ve and infinity: Yes IEEE754-2008 fused multiply-add: Yes Cache type: None Cache line size: 0 Cache size: 0 Global memory size: 268107776 Constant buffer size: 65536 Max number of constant args: 9 Local memory type: Scratchpad Local memory size: 16384 Profiling timer resolution: 1000 Device endianess: Little Available: Yes Compiler available: Yes Execution capabilities: Execute OpenCL kernels: Yes Execute native function: No Queue properties: Out-of-Order: Yes Profiling : Yes Platform ID: 003E8750 Name: GeForce 8600 GT Vendor: NVIDIA Corporation Driver version: 260.99 Profile: FULL_PROFILE Version: OpenCL 1.0 CUDA Extensions: cl_khr_byte_addressable_store cl_khr_icd cl_khr_gl_sharing cl_nv_d3d9_sharing cl_nv_compiler_options cl_nv_device_attribute_query cl_nv_pragma_unroll cl_khr_global_int32_base_atomics cl_khr_global_int32_extended_atomics
Error : Bytes mismatch! Error : glSharing mismatch! Error : images mismatch! Error : printf mismatch! Error : deviceAttributeQuery mismatch! Failed! Platform Name: ATI Stream Number of devices: 1 Device Type: CL_DEVICE_TYPE_CPU Device ID: 4098 Max compute units: 2 Max work items dimensions: 3 Max work items[0]: 1024 Max work items[1]: 1024 Max work items[2]: 1024 Max work group size: 1024 Preferred vector width char: 16 Preferred vector width short: 8 Preferred vector width int: 4 Preferred vector width long: 2 Preferred vector width float: 4 Preferred vector width double: 0 Max clock frequency: 2211Mhz Address bits: 32 Max memory allocation: 536870912 Image support: No Max size of kernel argument: 4096 Alignment (bits) of base address: 1024 Minimum alignment (bytes) for any datatype: 128 Single precision floating point capability Denorms: Yes Quiet NaNs: Yes Round to nearest even: Yes Round to zero: Yes Round to +ve and infinity: Yes IEEE754-2008 fused multiply-add: No Cache type: Read/Write Cache line size: 64 Cache size: 65536 Global memory size: 1073741824 Constant buffer size: 65536 Max number of constant args: 8 Local memory type: Global Local memory size: 32768 Profiling timer resolution: 279 Device endianess: Little Available: Yes Compiler available: Yes Execution capabilities: Execute OpenCL kernels: Yes Execute native function: Yes Queue properties: Out-of-Order: No Profiling : Yes Platform ID: 01DFD40C Name: AMD Athlon(tm) 64 X2 Dual Core Processor 4200+ Vendor: AuthenticAMD Driver version: 2.0 Profile: FULL_PROFILE Version: OpenCL 1.1 ATI-Stream-v2.2 (302) Extensions: cl_amd_fp64 cl_khr_global_int32_base_atomics cl_khr_global_int32_extended_atomics cl_khr_local_int32_base_atomics cl_khr_local_int32_extended_atomics cl_khr_byte_addressable_store cl_khr_gl_sharing cl_ext_device_fission cl_amd_device_attribute_query cl_amd_printf
