CUDAカーネルですべての乗算を__umul24関数に置き換える価値はありますか?私は異なった反対の意見を読みました、そして私はそれを理解するためにまだ手がかりを作ることができません
2 に答える
Ashwin/fabrizioMとは少し違う意見でチャイムを鳴らしたかっただけです...
CUDAを自分で学ぼうとしているだけなら、おそらく彼らの答えは多かれ少なかれ受け入れられます。しかし、実際に本番環境向けのアプリを商用または研究環境にデプロイしようとしている場合、エンドユーザー(またはエンドユーザーの場合はあなた)が絶対に確信している場合を除いて、そのような態度は一般的に受け入れられません。ユーザー)はFermi以降です。
おそらく、コンピュートレベルの適切な機能を使用することでメリットを享受するレガシーマシンでCUDAを実行するユーザーがたくさんいます。そして、Ashwin/fabrizioMがそれを実現するほど難しくはありません。
たとえば、私が取り組んでいるコードでは、次のものを使用しています。
//For prior to Fermi use umul, for Fermi on, use
//native mult.
__device__ inline void MultiplyFermi(unsigned int a, unsigned int b)
{ a*b; }
__device__ inline void MultiplyAddFermi(unsigned int a, unsigned int b,
unsigned int c)
{ a*b+c; }
__device__ inline void MultiplyOld(unsigned int a, unsigned int b)
{ __umul24(a,b); }
__device__ inline void MultiplyAddOld(unsigned int a, unsigned int b,
unsigned int c)
{ __umul24(a,b)+c; }
//Maximum Occupancy =
//16384
void GetComputeCharacteristics(ComputeCapabilityLimits_t MyCapability)
{
cudaDeviceProp DeviceProperties;
cudaGetDeviceProperties(&DeviceProperties, 0 );
MyCapability.ComputeCapability =
double(DeviceProperties.major)+ double(DeviceProperties.minor)*0.1;
}
ここに欠点があります。それは何ですか?
乗算を使用するカーネルには、2つの異なるバージョンのカーネルが必要です。
その価値はありますか?
よく考えてみてください。これは簡単なコピー&ペーストの仕事であり、効率が上がっていると思います。結局のところ、CUDAは概念的に最も簡単なプログラミング形式ではありません(並列プログラミングでもありません)。パフォーマンスが重要でない場合は、自問してみてください。なぜCUDAを使用しているのですか?
パフォーマンスが重要な場合は、怠惰なコーディングを怠り、レガシーデバイスを放棄するか、最適とは言えない実行を投稿します。ただし、展開のレガシーサポートを放棄できる(最適な実行を可能にする)ことが絶対に確信できる場合を除きます。
ほとんどの場合、レガシーサポートを提供することは理にかなっています。それを行う方法を理解すれば、それほど難しくはないからです。これは、将来のアーキテクチャの変更に適応するために、コードも更新する必要があることを意味することに注意してください。
一般に、コードが作成されたときに、コードが対象とした最新バージョンに注意し、最新の実装が最適化されている以上の計算機能をユーザーが持っている場合は、ユーザーに何らかの警告を出力する必要があります。
fermi より前のアーキテクチャを持つデバイス、つまり整数演算ユニットが 24 ビットである 2.0 より前の cuda 機能を持つデバイスのみ。
機能 >= 2.0 の Cuda デバイスでは、アーキテクチャは 32 ビットであり、_umul24 は高速ではなく低速になります。その理由は、32 ビット アーキテクチャで 24 ビット操作をエミュレートする必要があるためです。
問題は次のとおりです。速度を上げるために努力する価値はありますか? おそらくそうではありません。