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