2013年09月19日

Linux icpcとWin Intel composer

ご存じだと思いますが,Intel c++ compilerのLinux版は,アカデミック用途に限り無料なんですよね.
GCCとICPC(-fast?)を比較したら,2倍以上速くなったような記憶があります.-fast? -o2?オプション付けたからかもしれませんが.(最適化オプションをどうしたかはちょっと記憶無いです.)

開発はWin 7 32Bit,Visual studioで行ってます.

Winのインテルコンパイラ(コンポーザ?)では,exit()が通るのですが,
Linuxでは「error: identifier "exit" is undefined」だったようなメッセージで,通りません.

何ででしょう?

そんなわけで,
#ifndef _win32
#define exit(n) _exit(n)
#endif
みたいにしたら,どっちでも通るようになりました.わーい
posted by にゃんこ at 21:44| Comment(0) | TrackBack(0) | C++

2013年09月09日

カーネル関数内で,何処まで進んだかPrint

カーネル関数内で,
my_cudaPass_inKernel(threadの一次元化ID, 任意の数字);
とすれば,そこまで進んだというのを表示できる関数.
ただ,カーネル関数のPrintの記述の順番と,表示の順番が何故か異なるケースもあるようで.


#define my_cudaPass_inKernel(idx, n) __my_cudaPass_inKernel(idx, n, __FILE__, __LINE__)

inline __device__ void __my_cudaPass_inKernel(const int idx, const int n, const char *file, const int line)
{
if(idx == n){
//compute Capabirity 2.0 or higher
printf("%s (%i) : Pass, in Kernel.\n", file, line);
}
return;
}
posted by にゃんこ at 21:26| Comment(0) | TrackBack(0) | CUDA C

参照渡しとデバイス関数内でのPrintf

以下のように,C++と同様に参照渡しが出来ます.

また,デバイス関数内でもPrintf()出来るのを,今日知ったのでメモ.
Printfは,ComputeCapability2.0以上でないと使用できないようです.
つまり,sm20以上にコンパイルオプションを指定しなければなりません.
他にも,C++とは違った制限・制約・機能があるようですが…
とりあえず動けば良いや的な



__device__ void D_func2(float3 &b)
{
b.x++;
return;
}
__global__ void D_func()
{
float3 a;
a.x= threadIdx.x;
a.y= threadIdx.x;
a.z= -1;
printf("Hello thread %d\n", threadIdx.x);
printf("a.x=%f, a.y=%f, a.z=%f\n", a.x, a.y, a.z);
D_func2(a);
printf("a.x=%f, a.y=%f, a.z=%f\n", a.x, a.y, a.z);
return;
}
posted by にゃんこ at 17:25| Comment(0) | TrackBack(0) | CUDA C

2013年09月05日

__device__変数間のコピー

デバイス間でメモリのコピーが上手く出来なかった.
D_aをD_bにコピーしたかったのですが,cudaMemcpyFromSymbol(D_b, D_a, sizeof(H_a), 0, cudaMemcpyDeviceToDevice);とすると引数エラーに.
cudaMemcpyToSymbolFromSymbol()的な扱いを想定していたんだけど.


いろいろ悩んだりおちんこだりしたけれど,とりあえず,Symbolのアドレス取得して,そのアドレスを使ってCudaMemcpyさせました.
CudaでのSymbolの取り扱いを理解できていないから,よく間違えるんだろうなぁ.


double H_a[110][110][110];
__device__ double D_a[110][110][110];
__device__ double D_b[110][110][110];

main(){
my_cudaSafeCall(cudaMemcpyToSymbol( D_a, H_a, sizeof(H_a)));
double *symbol_from, *symbol_to;
my_cudaSafeCall(cudaGetSymbolAddress((void **)&symbol_from, D_a));
my_cudaSafeCall(cudaGetSymbolAddress((void **)&symbol_to, D_b));
// my_cudaSafeCall(cudaMemcpyFromSymbol(D_b, D_a, sizeof(H_a), 0, cudaMemcpyDeviceToDevice));//ここで11 Invalid argument
my_cudaSafeCall(cudaMemcpy(symbol_to, symbol_from, sizeof(H_a), cudaMemcpyDeviceToDevice));
my_cudaSafeCall(cudaMemcpyFromSymbol(H_a, D_b, sizeof(H_a)));
}
posted by にゃんこ at 23:17| Comment(0) | TrackBack(0) | CUDA C

2013年09月04日

uses too much shared data

ptxas : error : Entry function '関数名' uses too much shared data (0x4e40 bytes, 0x4000 max)

と出て,コンパイル(リンク?)が通らないので調べてみた.

https://devtalk.nvidia.com/default/topic/465799/cuda-programming-and-performance/how-to-use-48k-shared-memory-in-gtx480-/
に,
とりあえず「-arch=sm_20」にしてみろよって書いてあったので,sm13だったのを「compute_20,sm_20」にプロジェクトのプロパティを変更した.

ComputeCapability1.3だと,GPUの性能に関わらず,シェアードメモリは16KB制限になってるっぽいね.
posted by にゃんこ at 16:33| Comment(0) | TrackBack(0) | CUDA C

2013年08月30日

RegisterとLmem,PTX

以下の最大値を見つけるだけのデバイス関数をデバッグ実行しようとしたら,
lmemが192Byteも使われていて悩んでいたので,メモ.

まずは,関数で
7 too many resources requested for launch
30 Unknown Error
などが発生していた.
各GPUで使用できるRegister数などのメモリ量が異なる.
(デバイスqueryで確認できる http://blog.tsp.me.uk/article/72126659.html )

で,各関数で使用されているRegister数などを確認しないといけない.
VC++では,通常の設定だとPTXが見つからなかったので,
プロパティ->構成->CUDA C/C++->Device Verbose PTXAS Output をYes
(もし最大Register数を指定するのであれば,Max Used Register に値を指定. 例;1,024threads/block使用して,許容が65,536Register/Blockならば,64を指定.)
これで,コンパイル時の出力に,各グローバルデバイス関数での各メモリ使用量が分かる(/thread)

ついでに,CUDA C/C++->Host Additional Compiler Options に-keepを付けておくと,PTXが(消されずに?)作られています.


「出力」に出力されるPTXASを確認すると,以下の関数で,
10 registers, 8216Byte shared mem, 24B cmem, 192B lmem(Local mem)
使用していました.
GTX 780は65,536Registers/Blockなので,1024threadsでも64個使えます.ので,Local使用するのはおかしいなと.Registerなら1クロックなのに,ローカルは100かかえるし.
Local memは-maxrregcount=64での使用量を超えないと使用されないと思い込んでいたので.

で,仕方ないので()プログラミングガイドを検索.しかし,原因となりそうな記述は見つからず.
・許容Register数を超えていない
・数学関数を使用していない
ので.

ここで,デバッグ情報に気付いて,リリースビルドしてみると,
13regs, 8216B shared, 12B cmemで,ローカルの記述が消えた.


あと,関数の実行後にGetLastErrorして問題なくても,その後変数をHostにコピーしようとしてUnknown Errorする場合は,
変数を確保し忘れているなどのミスの可能性もある.
変数がMallocされていなくても動作する(Pointerで動作するから?)




__global__ void max_search(double *D_e)
{
//extern
__shared__ double error[1024]; //収束誤差の計算用
int id= RTN_threadIdx_inBlock(threadIdx);
int3 ijk= RTN_index_ijk(blockIdx, threadIdx, false);
int i= ijk.x;
int j= ijk.y;
int k= ijk.z;
if(i
posted by にゃんこ at 16:21| Comment(0) | CUDA C

2013年08月21日

.cuのハイライト化・インテリセンス

そういえばblockDimとかインテリセンスの補間に出てこなかったり波線が出たりで微妙に不便だったので,ハイライトとインテリセンスを有効化しようとgoogle先生に聞いてみた.

Visual C++ 2010 でCUDAソースファイル(.cu)のシンタックスハイライトとインテリセンス有効化
http://feather.cocolog-nifty.com/weblog/2011/07/visual-c-2010-c.html

勝手に紹介.


要点は,
1. NvidiaがSDKで配布してるusertype.datをVC\Common7\IDEに入れて
2. VSのオプションのプロジェクトおよびソリューション->VC++プロジェクトの設定 内の「含める拡張子」に「.cu」追加.(;.cuを末尾に)
3. テキストエディタ->ファイル拡張子 拡張子に「cu」があるか確認
のようです.

文字色は,C/C++ユーザーキーワード扱いの様子.
posted by にゃんこ at 16:19| Comment(0) | CUDA C

2013年08月09日

GPU Computing SDKとincludeパス

帰省中にCUDAコード書こうと思って,ノートPCにインストしたのでメモ.


CUDA4.2では,CUDAはToolkitやSDKなどが分割されて提供されていたと思う.
けど,CUDA5.5は一つにまとめられて提供されている.
1個落として,そのままインストするだけで簡単にCUDA環境が整う.

問題は,cutil_inline.hのあるGPU Computing SDK.
いろいろ探したけど,見つからない.どこに行ってしまったのか.

で,仕方ないので4.2をDLした.
https://developer.nvidia.com/cuda-toolkit-42-archive

GPU Computing SDK 5.5はどうなってるかご存じの方,教えて下さい!



で,includeパスをVCに通さなきゃいけない.
CUDA_PATH=C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v5.5\
CUDA_BIN_PATH=%CUDA_PATH%\bin
CUDA_INC_PATH=%CUDA_PATH%\include
CUDA_LIB_PATH=%CUDA_PATH%\lib\Win32
っぽいので,
$(CUDA_BIN_PATH)
$(CUDA_INC_PATH)
$(CUDA_LIB_PATH)
を,
C:\Users\<ユーザー名>\AppData\Local\Microsoft\MSBuild\v4.0\Microsoft.Cpp.Win32.user.props
の該当箇所に挿入.

これがデフォルト的な扱いらしく,継承されてOK.


SDKのパスはよく分からなかったので,絶対パスにした.
posted by にゃんこ at 20:48| Comment(0) | CUDA C

配列とポインタ

よく間違える配列のポインタをメモ


int array[4][5]
「要素5個の配列」の要素4個の配列
1.int型要素20個が連続して確保されている
(2.Fortranとかでは配列要素外にアクセスできないはずだけど,C++ではポインタ的にアクセスできる.)

で,
&array:配列全体の先頭を示すポインタ
&array[0]:配列の先頭要素を示すポインタ.単にarrayと書いても良い(&array[0]に置き換えてくれる?)
array:&array[0]を示す場合もあれば,配列全体を示す場合も.
->sizeof(array) == sizeof(int)*4*5 != sizeof(&array[0])
この場合では,左2個は20B,右は4B.

なので,
int array[N][M];
...
sizeof(array);
としていたのを,

int **a= new int *[4];
for(i= 0; i < 4; i++){
a[i]= new int [5];
}
for(i= 0; i < 4; i++){
cout << a[i] << endl;//連続していないないよ
}

とすると,sizeof(a)はポインタの大きさを返す.

配列のアドレスが連続している必要性がある場合には,以下.

int (*b)[5]= new int [4][5];
for(i= 0; i < 4; i++){
cout << b[i] << endl;
for(j= 0; j < 5; j++){
b[i][j]= array0[i][j];
}
}
for( i= 0; i < 4; i++){
for(j= 0; j < 5; j++){
cout << b[i][j] << endl;
}
}
cout << sizeof(b) << endl;




使い道は分からないけど,
配列要素がアドレスに対して連続して確保されており,配列要素が範囲外でもポインタ的に動作するので,

int array0[4][5];
int i= 2, j= 3;
array0[i][j]= 23;
cout << array0[0][i*5+ j] << endl;

も一応できる.



for(i= 0; i < 4; i++){
cout << array0[i] << endl;
}

とすれば,sizeof(int)[B]*5=20B=x14づつアドレスが増えている.


配列を引数としても渡せる.が,ポインタを渡しているだけなので注意.
Privateメンバ変数なんかは無理かも?

for(i= 0; i < 4*5; i++){
array0[0][i]= i;
}
func(array0);
cout << array0[0][0] << endl;

void func(int (*a)[5])
{
cout << a[1] << endl;
for( int i= 0; i < 4; i++){
for(int j= 0; j < 5; j++){
cout << a[i][j] << endl;
a[i][j]++;
}
}
return;
}

ここで,aは「int型5個配列」の先頭を示すポインタ.int x[][5]としてもOK
posted by にゃんこ at 16:19| Comment(0) | C++

DeviceQueryとCurand


//cf. cuda toolkit curand guide

// NOTE
// compute capability 1.3 over; USING double (-floating decimal)


#include

#include
#include //using cutilsafecall
#include
#pragma comment (lib, "curand.lib")

using namespace std;

#define is_110 110

struct integer3d{
int x, y, z;
};
struct struct_cuda_device_info{
string devName;
int MaxThreadsNum_perBlock;
integer3d MaxThreads3D_perBlock;
integer3d MaxBlocks3D_perGrid;
int computeCapability_Major;
int computeCapability_Minor;
int sharedMem_perBlock;
int register_perBlock;
int warpSize;
};

bool success;
int devSlect;
struct_cuda_device_info *cuda_dev_info;

double H_a[is_110][is_110][is_110];
__device__ double D_a[is_110][is_110][is_110];
__device__ int D_count_0, D_count_1;

__device__ integer3d RTN_index_ijk(dim3, dim3, bool);
__global__ void d_add(double*, double*);
void cuda_initial_chk();


int main()
{
cout << "main start." << endl;
cuda_initial_chk();
if(success == false){
return 0;
}

cout << "curand start." << endl;
int size_a= sizeof(H_a);
double *D_b, *D_c;
cutilSafeCall(cudaMalloc(&D_b, size_a));
cutilSafeCall(cudaMalloc(&D_c, size_a));

for( int i= 0; i < is_110; i++){
for( int j= 0; j < is_110; j++){
for(int k= 0; k < is_110; k++){
H_a[i][j][k]= -2.0;
}
}
}

int H_0= -1; //0であると,バグなのか,ヒットしていないのか分からない
cutilSafeCall(cudaMemcpyToSymbol(D_count_0, &H_0, sizeof(int)));
cutilSafeCall(cudaMemcpyToSymbol(D_count_1, &H_0, sizeof(int)));

//curand失敗の場合,分かりやすい
cutilSafeCall(cudaMemcpy(D_b, H_a, size_a, cudaMemcpyHostToDevice));
cutilSafeCall(cudaMemcpy(D_c, H_a, size_a, cudaMemcpyHostToDevice));

// rand generate INITIALIZE
curandGenerator_t generator;
curandCreateGenerator(&generator, CURAND_RNG_PSEUDO_XORWOW); //create generator
curandSetPseudoRandomGeneratorSeed(generator, 1234ULL); //set seed

//** 0 to 1 **
int count= is_110* is_110* is_110;
curandGenerateUniformDouble(generator, D_b, count); //generate 0~1 random num.
cutilSafeCall(cudaGetLastError());
curandGenerateUniformDouble(generator, D_c, count);//generate 0~1 random num.
cutilSafeCall(cudaGetLastError());
//失敗していても,Error返さないときがある? //CUDA Memcpyして,30 Unknown errorが返る時もある.

cout << "curand generated." << endl;

dim3 blockNum, threadNum; //即ち,GridDim,BlockDim
threadNum.x= threadNum.y= 16;
int MaxThreadsNum= cuda_dev_info[devSlect].MaxThreadsNum_perBlock;
threadNum.z= (int)(floor(MaxThreadsNum/ (double)(threadNum.x* threadNum.y))+ 1.0e-5); //丸め誤差対策
//floor:切り捨て. ex. 1024/(16*15)= 4.2 -> 16*15*5=1,200

//check threads size
integer3d MaxThreads= cuda_dev_info[devSlect].MaxThreads3D_perBlock;
if(threadNum.x > MaxThreads.x){threadNum.x= MaxThreads.x;}
if(threadNum.y > MaxThreads.y){threadNum.y= MaxThreads.y;}
if(threadNum.z > MaxThreads.z){threadNum.z= MaxThreads.z;}
cout << "Threads SIZE;" << endl;
cout << "\tx:" << threadNum.x << "\ty:" << threadNum.y << "\tz:" << threadNum.z << endl;

blockNum.x= (int)(ceil(is_110/(float)threadNum.x) + 5.0e-7); //丸め誤差対策
blockNum.y= (int)(ceil(is_110/(float)threadNum.y) + 5.0e-7); //ex. ceil(9.0/3.0) -> 2.999999とか
blockNum.z= (int)(ceil(is_110/(float)threadNum.z) + 5.0e-7);
cout << "Blocks SIZE;" << endl;
cout << "\tx:" << blockNum.x << "\ty:" << blockNum.y << "\tz:" << blockNum.z << endl;

integer3d MaxBlocks= cuda_dev_info[devSlect].MaxBlocks3D_perGrid;
if(blockNum.x>MaxBlocks.x || blockNum.y>MaxBlocks.y || blockNum.z>MaxBlocks.z){
cout << "Block size ERROR! over the maxblock size of current GPU-device." << endl;
cout << "Current device name:\t" << cuda_dev_info[devSlect].devName << endl;
cout << "Max Block SIZE;" << endl;
cout << "\tx:" << MaxBlocks.x << "\ty:" << MaxBlocks.y << "\tz:" << MaxBlocks.z << endl;
return 0;
}

d_add<<>>(D_b, D_c);
cutilSafeCall(cudaGetLastError());
cudaThreadSynchronize();
cutilSafeCall(cudaMemcpyFromSymbol(H_a, D_a, size_a));

cout << "curand value;" << endl;
cutilSafeCall(cudaMemcpyFromSymbol(&H_0, D_count_0, sizeof(int)));
cout << "\t0.0 hits:\t" << H_0+1 << endl; //H_0を-1で初期化したので
cutilSafeCall(cudaMemcpyFromSymbol(&H_0, D_count_1, sizeof(int)));
cout << "\t1.0 hits:\t" << H_0+1 << endl;
//0,0Hitなら,乱数の範囲は0超過1未満であることが確認できる

curandDestroyGenerator(generator);
cutilSafeCall(cudaFree(D_b));
cutilSafeCall(cudaFree(D_c));
delete [] cuda_dev_info;

return 0;
}


__global__ void d_add(double *D_b, double *D_c)
{
integer3d ijk= RTN_index_ijk(blockIdx, threadIdx, false);
int i, j, k;
i= ijk.x, j= ijk.y, k= ijk.z;
if(is_110<=i || is_110<=j || is_110<=k){
return;
}

int id_array= (is_110* is_110* i)+ (is_110* j)+ k;
double b= D_b[id_array];
double c= D_c[id_array];
D_a[i][j][k]= b+ c;

//グローバルにアクセスしてるけど,基本的に書き込みしてないので衝突しにくい
if(b == 0.0){ D_count_0++; }
if(c == 0.0){ D_count_0++; }
if(b == 1.0){ D_count_1++; }
if(c == 1.0){ D_count_1++; }
return;
}

__device__ integer3d RTN_index_ijk(dim3 B_Idx, dim3 T_Idx, bool Reverce)
{
//blockDimは,1,2... //blockDimはブロック内スレッド数
//B_Idx,T_Idxは0,1,2…
integer3d ijk;
int i= (blockDim.x* B_Idx.x)+ T_Idx.x;
int j= (blockDim.y* B_Idx.y)+ T_Idx.y;
int k= (blockDim.z* B_Idx.z)+ T_Idx.z;

if(Reverce == false){
ijk.x= i;
ijk.y= j;
ijk.z= k;
}else{
ijk.z= i;
ijk.y= j;
ijk.x= k;
}
return ijk;
}

void cuda_initial_chk()
{
success= true;

int devCount; //何個?
cutilSafeCall(cudaGetDeviceCount(&devCount));
// n個のdevice情報を保存するために,動的確保
cuda_dev_info= new struct_cuda_device_info[devCount];

if(devCount == 0){
cout << "Your Computer has not the CUDA device (GPU-card). halt." << endl;
success= false;
return;
}else{
cout << "Your computer has CUDA devices, count is:\t" << devCount << endl;
}

int devNo;
cudaDeviceProp deviceProp;

for(devNo= 0; devNo < devCount; devNo++){ //n個繰り返し
cout << endl;
cout << "**** Device No.\t" << devNo << endl;
cutilSafeCall(cudaGetDeviceProperties(&deviceProp, devNo)); //型番取得
cuda_dev_info[devNo].devName= deviceProp.name;
cout << "Device name:\t" << deviceProp.name << endl;

//compute Capability
cuda_dev_info[devNo].computeCapability_Major= deviceProp.major;
cuda_dev_info[devNo].computeCapability_Minor= deviceProp.minor;
cout << "Compute Capability:\t" << deviceProp.major << "." << deviceProp.minor << endl;
if( (deviceProp.major<1) || (deviceProp.major==1 && deviceProp.minor<3) ){
//0. or 1.3<
cout << "**CATION** This GPU is disabled \"Double-precision floating\"." << endl;
}

cuda_dev_info[devNo].sharedMem_perBlock= deviceProp.sharedMemPerBlock;
cout << "Shared memory per block:\t"<< deviceProp.sharedMemPerBlock/1024 << " KB" << endl;

cuda_dev_info[devNo].register_perBlock= deviceProp.regsPerBlock;
cout << "Register size per block;\t" << deviceProp.regsPerBlock << endl;

cuda_dev_info[devNo].warpSize= deviceProp.warpSize;
cout << "Warp size:\t" << deviceProp.warpSize << endl;

cuda_dev_info[devNo].MaxThreadsNum_perBlock= deviceProp.maxThreadsPerBlock;
// ****perBlock と perMultiProcessorとは違うよ!!****
cout << "Max threads per block:\t" << deviceProp.maxThreadsPerBlock << endl;

cuda_dev_info[devNo].MaxThreads3D_perBlock.x= deviceProp.maxThreadsDim[0];
cuda_dev_info[devNo].MaxThreads3D_perBlock.y= deviceProp.maxThreadsDim[1];
cuda_dev_info[devNo].MaxThreads3D_perBlock.z= deviceProp.maxThreadsDim[2];
cout << "Max Threads per block, SIZE;" << endl;
cout << "\tx:\t" << deviceProp.maxThreadsDim[0];
cout << "\ty:\t" << deviceProp.maxThreadsDim[1];
cout << "\tz:\t" << deviceProp.maxThreadsDim[2] << endl;

cuda_dev_info[devNo].MaxBlocks3D_perGrid.x= deviceProp.maxGridSize[0];
cuda_dev_info[devNo].MaxBlocks3D_perGrid.y= deviceProp.maxGridSize[1];
cuda_dev_info[devNo].MaxBlocks3D_perGrid.z= deviceProp.maxGridSize[2];
cout << "Max Blocks per grid, SIZE;" << endl;
cout << "\tx:\t" << deviceProp.maxGridSize[0];
cout << "\ty:\t" << deviceProp.maxGridSize[1];
cout << "\tz:\t" << deviceProp.maxGridSize[2] << endl;
}

for( ; ; ){
cout << "For lunch CUDA device No, Please Input..." << endl;
cin >> devSlect;
if( cin.fail() ){
cout << "ConsoleIn has the fail state." << endl;
cin.clear(); //error clear
cin.ignore(INT_MAX, '\n'); //buffer clear
}else if(0<=devSlect && devSlect
posted by にゃんこ at 11:03| Comment(0) | CUDA C