主要内容

coder.gpu.nokernel

ループの作成を无效プラグマプラグマ

说明

coder.gpu.nokernel()はループプラグマであり,,,の直前にすると,コードジェネレータージェネレーターでそのそのループループ内のステートメントステートメント®カーネルがれなくます。このはパラメーターを必要としません。

この关数コード生成ですです。matlab®では效果ありません。

すべて折りたたむ

この例は,关数にNokernelプラグマをてコードジェネレーターでそのののステートメントのののカーネルカーネルカーネルカーネルさされないないようように

1つのファイルに,サイズ32x512の2つのつの入力a,bを受け入れるエントリ关数Nestedloopを记述。关数には反复のささが异なる异なる异なる异なるつのつの入れ子にさたた为了1つは1つは1つは1つは1岁もうもう,,,,つはつは行にににに沿っ沿っ沿ったたたた演算演算演算のののループですですですです。。。。。。。。。。番目1の入れ子れループででそのその合计ををを倍倍スケーリングしていい。

功能[C] = Nestedloop(a,b)g =零(32,512);C =零(32,512);coder.gpu.kernelfun();%此嵌套环将融合为了i = 1:32为了j = 1:512 g(i,j)= a(1,j) + b(1,j);结尾结尾coder.gpu.nokernel();为了i = 1:32为了j = 1:512 c(i,j)= g(i,j) * 3;结尾结尾结尾

关数代码根Cuda Mex关数关数关数を生成生成し。。。

cfg = coder.gpuconfig('Mex');cfg.generatereport = true;代码根-configCFG-args{一个(1,512,'double'),一个(1,512,'double')}}Nestedloop

GPU编码器は2つのつの作成し。。。Nestedloop_kernel11番目番目入れ子されたループ计算计算g(i,j)= a(1,j) + b(1,j);を実行し,Nestedloop_kernel22番目番目番目入れ子されたの计算计算计算c(i,j)= g(i,j) * 3;を実行ます。番目番目カーネルはは番目番目入れ子されたループループの内部ループ用ますますししNokernelプラグマはステートメント直后のループのみされます生成生成されたカーネルカーネルののますますます

静态__global ____launch_bounds __(512,1)void nestedloop_kernel1(const real_t b [512],const real_t a [512],real_t g [16384]){uint32_t threadID;... if(i <32){g [i +(j << 5)] = a [j] + b [j];}} static __global ____launch_bounds __(512,1)void nestedloop_kernel2(real_t g [16384],int32_t i,real_t c [16384])...;if(j <512){c [i +(j << 5)] = g [i +(j << 5)] * 3.0;}

main关数のは,ジェネレーターによってによってによって番目番目入れ子にされループがが,,カーネルカーネルカーネルの起动起动でで示さ示されるれるとおりにに融合融合さされれれれててているががますますますますますますますます外侧のカーネルマッピングされませんため,コードジェネレーター,,循环ステートメントをを番目のののNestedloop_kernel2の呼び出し直前に配置。。

void nestedloop(const real_t a [512],const Real_t b [512],real_t c [16384]){int32_t i;... //这两个循环将被融合为cudamemcpy(gpu_b,(void *)&b [0],4096ul,cudamemcpyHostTodeVice);cudamemcpy(gpu_a,(void *)和a [0],4096ul,cudamemcpyhosttodevice);nestedloop_kernel1 <<< dim3(32u,1u,1u),dim3(512u,1u,1u,1u)>>>( *gpu_b, *gpu_a, *gpu_a, *gpu_g);for(i = 0; i <32; i ++){nestedloop_kernel2 <<< dim3(1u,1u,1u,1u),dim3(512u,1u,1u,1u)>>>( *gpu_g,i,i, *gpu_c);c_dirtyongpu = true;} ... cudafree(*gpu_c);}

バージョン履歴

R2019aで导入