4

nSight プロファイラーは、次のカーネルがスレッドごとに 52 個のレジスターを使用していることを示しています。

 //Just the first lines of the kernel.
__global__ void voles_kernel(float *params, int *ctrl_params, 
                             float dt, float currTime,
                             float *dev_voles, float *dev_weasels, 
                             curandStateMtgp32 *state) 
{

 __shared__ float dev_params[9];

 __shared__ int BuYeSimStep[4];

 if(threadIdx.x < 4)
 {
   BuYeSimStep[threadIdx.x] = ctrl_params[threadIdx.x];
 }

 if(threadIdx.x < 9){
     dev_params[threadIdx.x] = params[threadIdx.x];
 }

 __syncthreads();

float currVole = curand_uniform(&state[blockIdx.x]) + 3.0;
float currWeas = curand_uniform(&state[blockIdx.x]) + 0.1;
float oldVole = currVole;
float oldWeas = currWeas;

int jj;

if (blockIdx.x * blockDim.x + threadIdx.x < BuYeSimStep[2])
{
int dayIndex = 0;

/* Not declaring any new variable from here on, just doing arithmetics.
   ....... */

各レジスタに 4 バイトがある場合、配列 params[9] と ctrl_params[4] がレジスタで終わると仮定しても、どうやって 52 個のレジスタに到達するのか理解できません (この場合、私が行ったように共有メモリを使用しても作成されません)。検出）。占有率を上げたいのですが、なぜそんなに多くのレジスターを使用しているのかわかりません。何か案は？

4

3 に答える 3

9

一般に、C コードを見てレジスタの使用状況を予測することは困難です。コンパイラーは、レジスターの使用量を増やしてコードを積極的に最適化する場合があります。これは、おそらく命令をあちこちに保存するためです。Cコードの変数割り当てからレジスタの使用を予測できると仮定しているようですが、2つの間に何らかの関係がありますが、Cコードの変数の割り当てからレジスタの使用を直接計算できるとは想定できません。

コードを提供していないため、実際にレジスタの使用を手伝ってくれる人は誰もいません。レジスタの使用法をよりよく理解したい場合は、PTX コードを直接見る必要があります。これを行うにはnvcc、スイッチを使用してコードをコンパイル-ptxし、結果の .ptx ファイルを直接調べます。これを行うには、PTX のドキュメントとnvcc のドキュメントを参照して、さまざまなコンパイラ オプションを確認することをお勧めします。

コードを提供していないため、直接的な提案を行うことは実際には不可能ですが、定数の使用を減らしたり、算術使用を減らしたりリファクタリングしたり、からdoubleに切り替えたりすることで、レジスタの使用を減らすことができるかもしれません。float他にもたくさんの提案。-Gスイッチをコンパイラに渡す場合、レジスタの使用も影響を受けます。

適切なパラメーターを指定して-maxrregcountスイッチを渡すことにより、スレッドごとのレジスターのコンパイラーの使用を制限できます。たとえば、スレッドごとに 20 レジスターに制限するようにコンパイラーに指示します。ただし、この方法では良い結果が得られないか、パフォーマンスがあまり犠牲にならない値にパラメーターを調整する必要がある場合があります。ただし、基本性能をあまり犠牲にせず、居住性を向上できる最適な選択肢を見つけることができます。コンパイラに過度の制約を加えると、必要なレジスタの使用量がローカル メモリに流出し始め、一般にパフォーマンスが低下します。nvcc-maxrregcount 20

また、コンパイル時にコンパイラのレジスタの使用状況やその他の関連データ (スピルなど) に関する有用な出力を提供-Xptxas -vするwhich を渡すことができることにも注意してください。nvcc

于 2013-07-28T22:56:03.570 に答える
1

コードを Eclipse Nsight でデバッグすることをお勧めします。カーネルの最初の行にブレークポイントを作成し、そこに進みます。デバッグ パースペクティブの CUDA スレッド内に、現在のスタック トレースがあります。スタックを右クリックし、[命令ステッピング モード] をクリックします。[逆アセンブリ] ウィンドウで、カーネル PTX アセンブリが開きます。カーネルのステップ インを続行して、ソース コードとアセンブリの相関関係を追跡できます。したがって、どのレジスターが使用されているかを知ることができます。

于 2014-08-21T21:19:17.440 に答える