JP2002073579A - 分散共有メモリ向けコンパイル方法 - Google Patents
分散共有メモリ向けコンパイル方法Info
- Publication number
- JP2002073579A JP2002073579A JP2000264491A JP2000264491A JP2002073579A JP 2002073579 A JP2002073579 A JP 2002073579A JP 2000264491 A JP2000264491 A JP 2000264491A JP 2000264491 A JP2000264491 A JP 2000264491A JP 2002073579 A JP2002073579 A JP 2002073579A
- Authority
- JP
- Japan
- Prior art keywords
- array
- space
- processor
- program
- shared
- Prior art date
- Legal status (The legal status is an assumption and is not a legal conclusion. Google has not performed a legal analysis and makes no representation as to the accuracy of the status listed.)
- Pending
Links
Landscapes
- Devices For Executing Special Programs (AREA)
Abstract
(57)【要約】
【課題】分散メモリ型マルチプロセッサのプログラム開
発手法として,従来は,逐次プログラムから出発して,
プログラム中の主要部分を少しずつ修正し,効果を確か
めながらチューニングを進め最終的に最高性能を引出す
という,インクリメンタルな開発手法が困難であった。 【解決手段】分散共有メモリ機構を備える並列システム
向けに,配列を共有空間またはプライベート空間に配置
し,プライベート配列についてプロセッサ間通信コード
を生成し,共有配列については通常のメモリアクセス命
令を生成し,また共有空間とプライベート空間の間での
配列コピーコードを生成するコンパイラを提供する。 【効果】ユーザは,まず配列を共有空間に配置して分散
共有メモリ機構を用いることによってプログラムを動作
させ,次にカーネル部分に対してのみ,配列をプライベ
ート空間に配置して通信を明示的に制御するという,イ
ンクリメンタルなプログラム開発が可能になる。
発手法として,従来は,逐次プログラムから出発して,
プログラム中の主要部分を少しずつ修正し,効果を確か
めながらチューニングを進め最終的に最高性能を引出す
という,インクリメンタルな開発手法が困難であった。 【解決手段】分散共有メモリ機構を備える並列システム
向けに,配列を共有空間またはプライベート空間に配置
し,プライベート配列についてプロセッサ間通信コード
を生成し,共有配列については通常のメモリアクセス命
令を生成し,また共有空間とプライベート空間の間での
配列コピーコードを生成するコンパイラを提供する。 【効果】ユーザは,まず配列を共有空間に配置して分散
共有メモリ機構を用いることによってプログラムを動作
させ,次にカーネル部分に対してのみ,配列をプライベ
ート空間に配置して通信を明示的に制御するという,イ
ンクリメンタルなプログラム開発が可能になる。
Description
【0001】
【発明の属する技術分野】本発明は,分散共有メモリ機
構を備える並列システム向けのコンパイラに関する。
構を備える並列システム向けのコンパイラに関する。
【0002】
【従来の技術】並列計算機システムの一形態として,各
プロセッサに分散したメモリを備えるものがある。この
ようなシステムを分散メモリ型マルチプロセッサ(DMP)
と呼ぶ。例えば,プロセッサ数が100を超えるような高
並列計算機や,パーソナルコンピュータ(PC)をネットワ
ークで接続したPCクラスタなどが,この部類に属する。
DMPの中には,物理的に分散されたメモリを共有メモリ
としてアクセスできる分散共有メモリ(DSM: Distrbuted
Shared Memory)機構を備えるものがある。すなわち,
任意のプロセッサから,通常のメモリアクセス命令によ
って,自プロセッサに付随するメモリ(ローカルメモリ)
も他のプロセッサに付随するメモリ(リモートメモリ)も
区別することなくアクセスできる。DSMには,ハードウ
ェアによって実現されるものや,OSによってソフトウェ
ア的に実現されるものがあるが,通常のメモリアクセス
命令によってアクセスできるという点は共通である。DS
Mの利点は,ユーザがプログラミングを行うときに,通
常の共有メモリのインタフェースを用いることができる
点である。すなわち,処理の並列実行を記述するだけで
よく,プロセッサ間のデータ転送(通信)を気にする必要
がない。このようなプログラミングインタフェースの代
表的なものとしてOpenMPが知られている。OpenMPでのプ
ログラミング例を以下に示す。
プロセッサに分散したメモリを備えるものがある。この
ようなシステムを分散メモリ型マルチプロセッサ(DMP)
と呼ぶ。例えば,プロセッサ数が100を超えるような高
並列計算機や,パーソナルコンピュータ(PC)をネットワ
ークで接続したPCクラスタなどが,この部類に属する。
DMPの中には,物理的に分散されたメモリを共有メモリ
としてアクセスできる分散共有メモリ(DSM: Distrbuted
Shared Memory)機構を備えるものがある。すなわち,
任意のプロセッサから,通常のメモリアクセス命令によ
って,自プロセッサに付随するメモリ(ローカルメモリ)
も他のプロセッサに付随するメモリ(リモートメモリ)も
区別することなくアクセスできる。DSMには,ハードウ
ェアによって実現されるものや,OSによってソフトウェ
ア的に実現されるものがあるが,通常のメモリアクセス
命令によってアクセスできるという点は共通である。DS
Mの利点は,ユーザがプログラミングを行うときに,通
常の共有メモリのインタフェースを用いることができる
点である。すなわち,処理の並列実行を記述するだけで
よく,プロセッサ間のデータ転送(通信)を気にする必要
がない。このようなプログラミングインタフェースの代
表的なものとしてOpenMPが知られている。OpenMPでのプ
ログラミング例を以下に示す。
【0003】 第2行目がOpenMPの指示文であり,直後のdoループの繰
返しを各プロセッサに分割して,並列に実行することを
指示する。ユーザはa(i)やb(i+1)でアクセスされる配列
要素がどのプロセッサのメモリ上に存在するかを気にし
なくて良い。
返しを各プロセッサに分割して,並列に実行することを
指示する。ユーザはa(i)やb(i+1)でアクセスされる配列
要素がどのプロセッサのメモリ上に存在するかを気にし
なくて良い。
【0004】一方,DMP向けの別のプログラミングモデ
ルとして,メッセージ通信方式が知られている。これ
は,プロセッサ間でのデータの転送を,ライブラリ呼び
出しの形で明示的に記述するものである。通常のメモリ
アクセス命令でアクセスできるのは自プロセッサのロー
カルメモリだけであり,リモートメモリ上のデータが必
要な場合は,前もってメッセージ通信ライブラリ呼び出
しによってローカルメモリに転送しておく必要がある。
メッセージ通信方式の利点は,ユーザがプロセッサ間通
信を明示的に制御できるため,きめ細かい性能チューニ
ングが可能な点である。メッセージ通信方式で上記のOp
enMPプログラムと同じ計算を行うためのプログラミング
例を以下に示す。
ルとして,メッセージ通信方式が知られている。これ
は,プロセッサ間でのデータの転送を,ライブラリ呼び
出しの形で明示的に記述するものである。通常のメモリ
アクセス命令でアクセスできるのは自プロセッサのロー
カルメモリだけであり,リモートメモリ上のデータが必
要な場合は,前もってメッセージ通信ライブラリ呼び出
しによってローカルメモリに転送しておく必要がある。
メッセージ通信方式の利点は,ユーザがプロセッサ間通
信を明示的に制御できるため,きめ細かい性能チューニ
ングが可能な点である。メッセージ通信方式で上記のOp
enMPプログラムと同じ計算を行うためのプログラミング
例を以下に示す。
【0005】 real a(50),b(51) if(mypid==1) call MPI_Send(b(1),1,mypid-1) ! データ送信 if(mypid==0) call MPI_Recv(b(51),1,mypid+1) ! データ受信 if(mypid==1) ub=49 if(mypid==0) ub=50 do i = 1,ub a(i) = b(i+1)+... enddo この例ではプロセッサ数は2台と仮定している。1行目
の宣言文では,各プロセッサに分割された配列のサイズ
を宣言する必要があり,そのため配列aのサイズは50と
なっている。また配列bについてはさらにデータ受信用
に1要素分の領域を確保する必要があり,サイズは51と
なっている。2,3行目では通信ライブラリ呼出しによっ
て,後続のループでb(i+1)として参照される配列bの要
素を転送している。ここで,第1引数は転送する先頭要
素,第2引数は転送要素数,第3引数は転送相手のプロセ
ッサ番号である。4,5行目では,各プロセッサが担当す
べきループ範囲の上限を求めている。これらの処理の後
に初めて元のループの処理が実行できる。この例に示さ
れるように,メッセージ通信方式におけるプログラミン
グはきわめて煩雑になる。
の宣言文では,各プロセッサに分割された配列のサイズ
を宣言する必要があり,そのため配列aのサイズは50と
なっている。また配列bについてはさらにデータ受信用
に1要素分の領域を確保する必要があり,サイズは51と
なっている。2,3行目では通信ライブラリ呼出しによっ
て,後続のループでb(i+1)として参照される配列bの要
素を転送している。ここで,第1引数は転送する先頭要
素,第2引数は転送要素数,第3引数は転送相手のプロセ
ッサ番号である。4,5行目では,各プロセッサが担当す
べきループ範囲の上限を求めている。これらの処理の後
に初めて元のループの処理が実行できる。この例に示さ
れるように,メッセージ通信方式におけるプログラミン
グはきわめて煩雑になる。
【0006】メッセージ通信方式のプログラミングの負
担を軽減するための言語として,VPP Fortranが知られ
ている。これは逐次実行用のプログラムにユーザが指示
文を挿入することによって,メッセージ通信プログラム
を簡潔に記述できるようにするものである。VPP Fortra
nで上記のメッセージ通信プログラムと同じ計算を行う
例を以下に示す。
担を軽減するための言語として,VPP Fortranが知られ
ている。これは逐次実行用のプログラムにユーザが指示
文を挿入することによって,メッセージ通信プログラム
を簡潔に記述できるようにするものである。VPP Fortra
nで上記のメッセージ通信プログラムと同じ計算を行う
例を以下に示す。
【0007】 1行目の宣言文では,OpenMPプログラムと同様に,配列
a,bのサイズは100のままでよい。2行目以降が指示文で
ある。2行目はプロセッサが2台であることを宣言し,3
行目,4行目は,配列を2台のプロセッサ上に分散配置す
ることを指示する。さらに配列bについては,4行目のov
erlap(0:1)という指示により,各プロセッサでデータ受
信用の領域を1要素分確保することを指示している。5行
目のoverlapfix指示文は,この受信用領域に対して隣接
プロセッサ間で配列要素を転送することを指示してい
る。6行目のspread do指示文は,後続のループをプロセ
ッサ2台で分割して実行することを指示している。この
とき,b(i+1)として参照される要素は,上記のoverlapf
ix指示文によって既に転送完了している。この例からも
分かるように,VPP Fortranを使うと,通常の逐次プロ
グラムに指示文を追加することによって通信が記述でき
るので,メッセージ通信プログラムの作成が容易になる
という利点がある。
a,bのサイズは100のままでよい。2行目以降が指示文で
ある。2行目はプロセッサが2台であることを宣言し,3
行目,4行目は,配列を2台のプロセッサ上に分散配置す
ることを指示する。さらに配列bについては,4行目のov
erlap(0:1)という指示により,各プロセッサでデータ受
信用の領域を1要素分確保することを指示している。5行
目のoverlapfix指示文は,この受信用領域に対して隣接
プロセッサ間で配列要素を転送することを指示してい
る。6行目のspread do指示文は,後続のループをプロセ
ッサ2台で分割して実行することを指示している。この
とき,b(i+1)として参照される要素は,上記のoverlapf
ix指示文によって既に転送完了している。この例からも
分かるように,VPP Fortranを使うと,通常の逐次プロ
グラムに指示文を追加することによって通信が記述でき
るので,メッセージ通信プログラムの作成が容易になる
という利点がある。
【0008】VPP Fortranについては,岩下英俊, 進藤
達也, 岡田信,"VPP Fortran: 分散メモリ型並列計算機
向けプログラミング言語",情報処理学会論文誌, Vol.3
6, No. 7, pp. 1542-1550 (1995) に述べられている。
達也, 岡田信,"VPP Fortran: 分散メモリ型並列計算機
向けプログラミング言語",情報処理学会論文誌, Vol.3
6, No. 7, pp. 1542-1550 (1995) に述べられている。
【0009】
【発明が解決しようとする課題】従来の方法では,それ
ぞれ以下のような問題があった。
ぞれ以下のような問題があった。
【0010】DSMを用いたOpenMPプログラムでは,プロ
セッサ間通信はハードウェアやOSによって暗黙に行われ
るので,ユーザは通信を明示的に制御できなかった。ま
た,データの各プロセッサの分散配置は,固定サイズの
ページを単位として行われるので,ユーザの要求するデ
ータ分散が正確に実現できない場合があった。これらの
制限のため,ユーザがプログラムをチューニングして最
高性能を引出すのには限界があった。
セッサ間通信はハードウェアやOSによって暗黙に行われ
るので,ユーザは通信を明示的に制御できなかった。ま
た,データの各プロセッサの分散配置は,固定サイズの
ページを単位として行われるので,ユーザの要求するデ
ータ分散が正確に実現できない場合があった。これらの
制限のため,ユーザがプログラムをチューニングして最
高性能を引出すのには限界があった。
【0011】一方メッセージ通信方式では,逐次プログ
ラムから出発して並列プログラムを作成する場合に,効
果を確かめながら少しずつチューニングするという「イ
ンクリメンタル」な開発手法を取ることが困難であっ
た。例えば,ある配列を分散した場合,その配列を参照
するすべてのサブルーチンにおいて,配列サイズの変更
とそれに伴う通信の挿入やループ繰返し範囲の変更を行
わなければならなかった。これは事実上,プログラム全
体の修正を意味する。VPP Fortranを用いた場合でも同
様に,ある配列を分散した場合,関連するサブルーチン
をすべて修正する必要があった。
ラムから出発して並列プログラムを作成する場合に,効
果を確かめながら少しずつチューニングするという「イ
ンクリメンタル」な開発手法を取ることが困難であっ
た。例えば,ある配列を分散した場合,その配列を参照
するすべてのサブルーチンにおいて,配列サイズの変更
とそれに伴う通信の挿入やループ繰返し範囲の変更を行
わなければならなかった。これは事実上,プログラム全
体の修正を意味する。VPP Fortranを用いた場合でも同
様に,ある配列を分散した場合,関連するサブルーチン
をすべて修正する必要があった。
【0012】なおVPP Fortranでは配列に対して「グロ
ーバル属性」を指定することができる。この属性を持つ
配列は,たとえ分散されていたとしても,コンパイラが
自動的に実行時ライブラリを挿入してリモートメモリ参
照かどうかを実行時に判定するようなコードを生成す
る。これを利用すれば,ある配列を分散したときに,ユ
ーザがプログラム全体について通信を挿入する負担は軽
減される。しかし,この場合,配列要素参照のたびにラ
イブラリ呼出しのオーバヘッドが発生してしまい,非常
に遅くなってしまうという問題があった。
ーバル属性」を指定することができる。この属性を持つ
配列は,たとえ分散されていたとしても,コンパイラが
自動的に実行時ライブラリを挿入してリモートメモリ参
照かどうかを実行時に判定するようなコードを生成す
る。これを利用すれば,ある配列を分散したときに,ユ
ーザがプログラム全体について通信を挿入する負担は軽
減される。しかし,この場合,配列要素参照のたびにラ
イブラリ呼出しのオーバヘッドが発生してしまい,非常
に遅くなってしまうという問題があった。
【0013】以上をまとめると,従来は,逐次プログラ
ムから出発してインクリメンタルにDMPプログラムを作
成することが困難であった。すなわち,逐次プログラム
をまずDMP上でそのまま走らせ,その後,プログラム中
の主要部分を少しずつDMP向けに修正していき,効果を
確かめながら並列化およびチューニングを進め,最終的
に最高性能を引出すという開発手法を行うことができな
かった。これは,大規模ステップ数のプログラムをDMP
向けに移植する際に大きな障壁となっていた。
ムから出発してインクリメンタルにDMPプログラムを作
成することが困難であった。すなわち,逐次プログラム
をまずDMP上でそのまま走らせ,その後,プログラム中
の主要部分を少しずつDMP向けに修正していき,効果を
確かめながら並列化およびチューニングを進め,最終的
に最高性能を引出すという開発手法を行うことができな
かった。これは,大規模ステップ数のプログラムをDMP
向けに移植する際に大きな障壁となっていた。
【0014】本発明の目的は,DMP向けのインクリメン
タルなプログラミングモデルを提供することにある。
タルなプログラミングモデルを提供することにある。
【0015】
【課題を解決するための手段】上記課題の解決のため,
本発明では,DSM機構を備える並列システム向けのコン
パイラにおいて,配列を共有空間または各プロセッサの
プライベート空間のいずれかに配置し,プライベート空
間に配置した配列については明示的なプロセッサ間通信
コードを生成し,共有空間に配置した配列については通
常のメモリアクセス命令を生成する。さらに,一つの配
列をプログラム中のある個所では共有空間に置き,別の
個所ではプライベート空間に置いて,共有空間とプライ
ベート空間の間での配列コピーコードを生成する。
本発明では,DSM機構を備える並列システム向けのコン
パイラにおいて,配列を共有空間または各プロセッサの
プライベート空間のいずれかに配置し,プライベート空
間に配置した配列については明示的なプロセッサ間通信
コードを生成し,共有空間に配置した配列については通
常のメモリアクセス命令を生成する。さらに,一つの配
列をプログラム中のある個所では共有空間に置き,別の
個所ではプライベート空間に置いて,共有空間とプライ
ベート空間の間での配列コピーコードを生成する。
【0016】このコンパイラを利用することにより,ユ
ーザは,まず配列を共有空間に配置してDSM機構を用い
ることによってプログラムを動作させ,次にプログラム
中のカーネル部分に対してのみ,配列をプライベート空
間にコピーすることによって,通信を明示的に制御しか
つ正確なデータ分散を実現するという,インクリメンタ
ルなプログラム開発が可能になる。
ーザは,まず配列を共有空間に配置してDSM機構を用い
ることによってプログラムを動作させ,次にプログラム
中のカーネル部分に対してのみ,配列をプライベート空
間にコピーすることによって,通信を明示的に制御しか
つ正確なデータ分散を実現するという,インクリメンタ
ルなプログラム開発が可能になる。
【0017】本発明ではさらに,共有空間からプライベ
ート空間への配列コピーにおいて,仮想共有空間上のペ
ージが既にコピー先プロセッサの物理メモリに割り付け
られているかどうかを判定し,既に割り付けられている
場合はページの内容をコピーする代わりに,コピー先の
仮想プライベート空間上のページを当該物理ページにマ
ッピングする。これにより,データ移動量を最小限に抑
え,またメモリ消費量を削減できる。
ート空間への配列コピーにおいて,仮想共有空間上のペ
ージが既にコピー先プロセッサの物理メモリに割り付け
られているかどうかを判定し,既に割り付けられている
場合はページの内容をコピーする代わりに,コピー先の
仮想プライベート空間上のページを当該物理ページにマ
ッピングする。これにより,データ移動量を最小限に抑
え,またメモリ消費量を削減できる。
【0018】
【発明の実施の形態】図1は,本発明のコンパイル方法
の一実施形態におけるコンパイラの処理を示すフローチ
ャートである。図2は本コンパイラへの入力プログラム
の例である。図1の各ステップを説明する前に,図2のプ
ログラムの概要を説明する。
の一実施形態におけるコンパイラの処理を示すフローチ
ャートである。図2は本コンパイラへの入力プログラム
の例である。図1の各ステップを説明する前に,図2のプ
ログラムの概要を説明する。
【0019】図2のプログラムはmain(1行目〜5行目),d
sm_sub(11行目〜18行目),mp_sub(21行目〜33行目)の三
つのルーチンから構成される。mainの2行目は配列の宣
言文であり,各々100個の要素からなる二つの配列aおよ
ぴbを宣言している。これらの配列は,mainの3行目およ
び4行目で,dsm_subやmp_subに引数として渡され,各ル
ーチンの中でこれらの配列を用いた計算が行われる。ds
m_subは計算量の比較的少ないルーチンであり,ユーザ
にとっては最高性能を追求する必要は少ない。一方mp_s
ubは27行目に繰返し数が1000のループを含み,計算量が
大きい。したがって,ユーザはmp_subに関しては最高性
能を得るために十分なチューニングを行いたい。!$omd
で始まる行は,チューニング用の指示文であり,本発明
のコンパイラによって認識される。これらの指示文の意
味は後で説明する。
sm_sub(11行目〜18行目),mp_sub(21行目〜33行目)の三
つのルーチンから構成される。mainの2行目は配列の宣
言文であり,各々100個の要素からなる二つの配列aおよ
ぴbを宣言している。これらの配列は,mainの3行目およ
び4行目で,dsm_subやmp_subに引数として渡され,各ル
ーチンの中でこれらの配列を用いた計算が行われる。ds
m_subは計算量の比較的少ないルーチンであり,ユーザ
にとっては最高性能を追求する必要は少ない。一方mp_s
ubは27行目に繰返し数が1000のループを含み,計算量が
大きい。したがって,ユーザはmp_subに関しては最高性
能を得るために十分なチューニングを行いたい。!$omd
で始まる行は,チューニング用の指示文であり,本発明
のコンパイラによって認識される。これらの指示文の意
味は後で説明する。
【0020】図1に戻り,図2のプログラムへの適用を例
として,各ステップを説明する。プロセッサ数は2台と
する。
として,各ステップを説明する。プロセッサ数は2台と
する。
【0021】ステップ100では,入力プログラム中の配
列宣言文およびdist指示文を読み込み,情報を図3に示
す配列表30に登録する。配列宣言文は,mainの2行目,d
sm_subの12行目,mp_subの22行目である。これらの宣言
文を読み込んで,配列表30の配列名フィールド300,形
状フィールド301に情報を登録する。dist指示文は,mp_
subの23行目である。この指示文は,配列を均等なブロ
ックに分割して,各プロセッサに割り当てること指示す
るものである。すなわち,2プロセッサの場合は,配列
a'およびb'を50要素ずつに分割して,各プロセッサに割
り当てる。コンパイラはdist指示文を読み込んで,配列
表30のdist情報フィールド302に情報を登録する。な
お,mainやdsm_sub内の配列aおよびbにはdist指示文が
無いので,それらに対するdist情報フィールド302には
何も登録しない。
列宣言文およびdist指示文を読み込み,情報を図3に示
す配列表30に登録する。配列宣言文は,mainの2行目,d
sm_subの12行目,mp_subの22行目である。これらの宣言
文を読み込んで,配列表30の配列名フィールド300,形
状フィールド301に情報を登録する。dist指示文は,mp_
subの23行目である。この指示文は,配列を均等なブロ
ックに分割して,各プロセッサに割り当てること指示す
るものである。すなわち,2プロセッサの場合は,配列
a'およびb'を50要素ずつに分割して,各プロセッサに割
り当てる。コンパイラはdist指示文を読み込んで,配列
表30のdist情報フィールド302に情報を登録する。な
お,mainやdsm_sub内の配列aおよびbにはdist指示文が
無いので,それらに対するdist情報フィールド302には
何も登録しない。
【0022】ステップ101から104では,配列表30を参照
して各配列を共有空間に配置するかプライベート空間に
配置するかを決定する。ここで共有空間およびプライベ
ート空間とは,共に仮想メモリ空間上の領域であり,共
有空間はDSM機構を利用して各プロセッサからアクセス
できる領域,プライベート空間は各プロセッサに固有の
領域であり他プロセッサからはアクセスできないもので
ある。
して各配列を共有空間に配置するかプライベート空間に
配置するかを決定する。ここで共有空間およびプライベ
ート空間とは,共に仮想メモリ空間上の領域であり,共
有空間はDSM機構を利用して各プロセッサからアクセス
できる領域,プライベート空間は各プロセッサに固有の
領域であり他プロセッサからはアクセスできないもので
ある。
【0023】まずステップ101では配列表30の中に配置
先が未定の配列があるかどうかを判定し,あればステッ
プ102に進む。ステップ102では,dist情報フィールド30
2を参照して,当該配列がdist指示文で指定されている
かどうかを判定する。もし指定されていればステップ10
3に進んで,当該配列の配置先をプライベート空間と決
定する。すなわち,配置先フィールド304に「private」
を登録する。図3の例では,配列a'およびb'について,
本ステップが適用されている。一方dist指示文で指定さ
れていない場合にはステップ104に進んで,当該配列の
配置先を共有空間と決定する。すなわち,配置先フィー
ルド304に「shared」を登録する。図3の例では,mainお
よびdsm_sub内の配列aおよびbについて,本ステップが
適用されている。
先が未定の配列があるかどうかを判定し,あればステッ
プ102に進む。ステップ102では,dist情報フィールド30
2を参照して,当該配列がdist指示文で指定されている
かどうかを判定する。もし指定されていればステップ10
3に進んで,当該配列の配置先をプライベート空間と決
定する。すなわち,配置先フィールド304に「private」
を登録する。図3の例では,配列a'およびb'について,
本ステップが適用されている。一方dist指示文で指定さ
れていない場合にはステップ104に進んで,当該配列の
配置先を共有空間と決定する。すなわち,配置先フィー
ルド304に「shared」を登録する。図3の例では,mainお
よびdsm_sub内の配列aおよびbについて,本ステップが
適用されている。
【0024】以上の手順によってすべての配列に対して
配置先が決定したら,ステップ105に進む。
配置先が決定したら,ステップ105に進む。
【0025】ステップ105では,shadow指示文およびref
lect指示文を読み込む。図2では,shadow指示文はmp_su
bの24行目,reflect指示文は同じくmp_subの25行目であ
る。このうちshadow指示文は,各プロセッサに分散され
た配列の端に受信用のバッファを指定された要素数だけ
設けることを指示する。例えば図2のshadow指示文は,
分散された配列b'の両側に,それぞれ0要素分および1要
素分のバッファを設けることを意味する。コンパイラは
shadow指示文を読み込んで,配列表30のshadow情報フィ
ールド303に情報を登録する。
lect指示文を読み込む。図2では,shadow指示文はmp_su
bの24行目,reflect指示文は同じくmp_subの25行目であ
る。このうちshadow指示文は,各プロセッサに分散され
た配列の端に受信用のバッファを指定された要素数だけ
設けることを指示する。例えば図2のshadow指示文は,
分散された配列b'の両側に,それぞれ0要素分および1要
素分のバッファを設けることを意味する。コンパイラは
shadow指示文を読み込んで,配列表30のshadow情報フィ
ールド303に情報を登録する。
【0026】一方, reflect指示文はshadow指示文で指
定された受信バッファに対して,実際にデータ転送を実
行することを指示する。図2のプログラムでは,配列b'
は1要素分のshadow領域を持つので,分散された配列の
端の1要素を隣接プロセッサに転送して,該隣接プロセ
ッサのshadow領域に格納する。
定された受信バッファに対して,実際にデータ転送を実
行することを指示する。図2のプログラムでは,配列b'
は1要素分のshadow領域を持つので,分散された配列の
端の1要素を隣接プロセッサに転送して,該隣接プロセ
ッサのshadow領域に格納する。
【0027】ステップ106では,前ステップで読み込ん
だreflect指示文に基づいて,実際に通信コードを生成
する。この処理は実際にはコンパイラ中間語に対して行
われるが,ここでは図4に示すようなソースプログラム
の形式で説明する。
だreflect指示文に基づいて,実際に通信コードを生成
する。この処理は実際にはコンパイラ中間語に対して行
われるが,ここでは図4に示すようなソースプログラム
の形式で説明する。
【0028】図4は,図2の入力プログラムに対する出力
プログラムを,ソースプログラム形式で表現したもので
ある。図1のステップ106で生成する通信コードは,83行
目および84行目である。ここでOMD_Sendは,第1引数お
よび第2引数で指定された自プロセッサ上のデータを,
第3引数で指定されたプロセッサに送信するライブラリ
ルーチンとする。第1引数は転送する先頭要素,第2引数
は転送要素数を表す。またOMD_Recvは,第3引数で指定
されたプロセッサからデータを受信して,第1引数およ
び第2引数で指定された自プロセッサ上の領域に格納す
るライブラリルーチンとする。mypidは自プロセッサ番
号を表す。したがって,83行目は,プロセッサ1(以下P
1)がプロセッサ0(以下P0)に対して自身のプライベート
空間内の配列要素b'(1)を送信する処理を表し,84行目
は,P0がP1から受信したデータを自身のプライベート空
間内のb'(51)に格納する処理を表す。
プログラムを,ソースプログラム形式で表現したもので
ある。図1のステップ106で生成する通信コードは,83行
目および84行目である。ここでOMD_Sendは,第1引数お
よび第2引数で指定された自プロセッサ上のデータを,
第3引数で指定されたプロセッサに送信するライブラリ
ルーチンとする。第1引数は転送する先頭要素,第2引数
は転送要素数を表す。またOMD_Recvは,第3引数で指定
されたプロセッサからデータを受信して,第1引数およ
び第2引数で指定された自プロセッサ上の領域に格納す
るライブラリルーチンとする。mypidは自プロセッサ番
号を表す。したがって,83行目は,プロセッサ1(以下P
1)がプロセッサ0(以下P0)に対して自身のプライベート
空間内の配列要素b'(1)を送信する処理を表し,84行目
は,P0がP1から受信したデータを自身のプライベート空
間内のb'(51)に格納する処理を表す。
【0029】図1に戻って,ステップ107では,入力プロ
グラム中の手続き呼出しを検出して,各実引数につい
て,対応する仮引数の配置先空間が実引数と異なるかど
うかを判定する。もし異なる場合,すなわち,一方が共
有空間でもう一方がプライベート空間の場合には,ステ
ップ108に進む。例えば,図2の入力プログラムでは,ma
inの3行目でdsm_subを呼出しているが,実引数も仮引数
も共に共有空間に配置されているから条件に該当しな
い。一方,mainの4行目のmp_subの呼出しでは,実引数a
およびbは共有空間に配置されているが,対応する仮引
数a'およびb'は共にプライベート空間に配置されてい
る。したがって,この手続き呼出しについては,ステッ
プ108に進むことになる。
グラム中の手続き呼出しを検出して,各実引数につい
て,対応する仮引数の配置先空間が実引数と異なるかど
うかを判定する。もし異なる場合,すなわち,一方が共
有空間でもう一方がプライベート空間の場合には,ステ
ップ108に進む。例えば,図2の入力プログラムでは,ma
inの3行目でdsm_subを呼出しているが,実引数も仮引数
も共に共有空間に配置されているから条件に該当しな
い。一方,mainの4行目のmp_subの呼出しでは,実引数a
およびbは共有空間に配置されているが,対応する仮引
数a'およびb'は共にプライベート空間に配置されてい
る。したがって,この手続き呼出しについては,ステッ
プ108に進むことになる。
【0030】ステップ108では,当該引数に対して,共
有空間とプライベート空間の間でコピーを行うコード
(以下,配列コピーコード)を生成する。例えばmp_subの
呼出しに対しては,図4の出力プログラムの55行目から5
9行目に示されるように,mp_sub呼出しの前後に,OMD_a
rray_copyというライブラリ呼出しを生成する。ここでO
MD_array_copyは第1引数で指定された配列を第2引数で
指定された配列にコピーするライブラリルーチンであ
る。実際にはコピーする範囲などを表す引数が続くが,
ここでは省略してある。55行目および56行目によって,
共有空間の配列aおよびbからプライベート空間の配列a'
およびb'へのコピーが実行される。またmp_sub呼出しか
ら戻った後に58行目および59行目によって,逆にプライ
ベート空間から共有空間への配列コピーが実行される。
有空間とプライベート空間の間でコピーを行うコード
(以下,配列コピーコード)を生成する。例えばmp_subの
呼出しに対しては,図4の出力プログラムの55行目から5
9行目に示されるように,mp_sub呼出しの前後に,OMD_a
rray_copyというライブラリ呼出しを生成する。ここでO
MD_array_copyは第1引数で指定された配列を第2引数で
指定された配列にコピーするライブラリルーチンであ
る。実際にはコピーする範囲などを表す引数が続くが,
ここでは省略してある。55行目および56行目によって,
共有空間の配列aおよびbからプライベート空間の配列a'
およびb'へのコピーが実行される。またmp_sub呼出しか
ら戻った後に58行目および59行目によって,逆にプライ
ベート空間から共有空間への配列コピーが実行される。
【0031】ステップ109では,共有空間に配置された
配列の要素の参照に対して,当該配列要素のアドレスへ
のメモリアクセス命令を生成する。図4に示す出力プロ
グラムは実際には機械語コードの形で生成されるが,そ
こでは配列の参照はメモリアクセス命令に置換されてい
る。この命令を生成するのが本ステップである。例えば
dsm_sub内での配列要素b(i+1)の参照に対して,図5に示
すような命令列を生成する。図5において,r0,r1,な
どはレジスタを表す。レジスタr0にはbの先頭アドレ
ス,r1にはiの値が格納されているものとする。これら
の値を元に3行目のadd命令から6行目のadd命令まででb
(i+1)のアドレスを計算し,レジスタr2に格納してい
る。最後のld命令がメモリアクセス命令であり,r2で指
定されるアドレスの内容,すなわちb(i+1)の値をr3にロ
ードしている。
配列の要素の参照に対して,当該配列要素のアドレスへ
のメモリアクセス命令を生成する。図4に示す出力プロ
グラムは実際には機械語コードの形で生成されるが,そ
こでは配列の参照はメモリアクセス命令に置換されてい
る。この命令を生成するのが本ステップである。例えば
dsm_sub内での配列要素b(i+1)の参照に対して,図5に示
すような命令列を生成する。図5において,r0,r1,な
どはレジスタを表す。レジスタr0にはbの先頭アドレ
ス,r1にはiの値が格納されているものとする。これら
の値を元に3行目のadd命令から6行目のadd命令まででb
(i+1)のアドレスを計算し,レジスタr2に格納してい
る。最後のld命令がメモリアクセス命令であり,r2で指
定されるアドレスの内容,すなわちb(i+1)の値をr3にロ
ードしている。
【0032】以上で,図1のコンパイル方法の説明を終
わり,これ以降は,ステップ108で生成した配列コピー
コードの処理内容を,より詳細に説明する。
わり,これ以降は,ステップ108で生成した配列コピー
コードの処理内容を,より詳細に説明する。
【0033】図4の出力プログラムは,各々のプロセッ
サによって実行されるものである。52行目で宣言されて
いる配列aおよびbは共有空間に配置されているので,両
方のプロセッサから同じものが見えるが,53行目で宣言
されている配列a'およびb'はプライベート空間に配置さ
れているので,ぞれぞれのプロセッサで異なるものが見
える。この様子を図6を用いて説明する。
サによって実行されるものである。52行目で宣言されて
いる配列aおよびbは共有空間に配置されているので,両
方のプロセッサから同じものが見えるが,53行目で宣言
されている配列a'およびb'はプライベート空間に配置さ
れているので,ぞれぞれのプロセッサで異なるものが見
える。この様子を図6を用いて説明する。
【0034】図6は,図4のプログラムを実行するとき
の,配列のメモリ配置を示す。プログラム内で宣言され
ている配列は,図の中心に示されるように仮想空間上に
配置される。仮想空間は共有空間とプライベート空間に
分かれており,プライベート空間はさらに各プロセッサ
に固有の空間に分かれている。また,仮想空間は,図の
両側に示される各プロセッサの物理空間にマッピングさ
れている。このマッピングはページと呼ばれる固定サイ
ズを最小単位として行われる。ここでは,ページのサイ
ズは配列要素20個分であると仮定する。
の,配列のメモリ配置を示す。プログラム内で宣言され
ている配列は,図の中心に示されるように仮想空間上に
配置される。仮想空間は共有空間とプライベート空間に
分かれており,プライベート空間はさらに各プロセッサ
に固有の空間に分かれている。また,仮想空間は,図の
両側に示される各プロセッサの物理空間にマッピングさ
れている。このマッピングはページと呼ばれる固定サイ
ズを最小単位として行われる。ここでは,ページのサイ
ズは配列要素20個分であると仮定する。
【0035】配列aおよびbは全体(100個要素)が共有空
間に配置される。これは,ページを単位として物理空間
にマッピングされている。どのページがどのプロセッサ
にマッピングされるかは任意であるが,いずれにして
も,100要素は5ページに相当しプロセッサ数で割り切れ
ないので,マッピングされる要素数はプロセッサによっ
て偏りがあることになる。ここでは,前の3ページはP0
に,後の2ページはP1にマッピングされているものとす
る。
間に配置される。これは,ページを単位として物理空間
にマッピングされている。どのページがどのプロセッサ
にマッピングされるかは任意であるが,いずれにして
も,100要素は5ページに相当しプロセッサ数で割り切れ
ないので,マッピングされる要素数はプロセッサによっ
て偏りがあることになる。ここでは,前の3ページはP0
に,後の2ページはP1にマッピングされているものとす
る。
【0036】また,配列a'およびb'は各プロセッサのプ
ライベート空間に,それぞれ50要素または51要素ずつ配
置される。配列b'の最後の(51番目)の要素は,shadow領
域である。これらは,各プロセッサの物理空間にマッピ
ングされなければならない。マッピングはやはりページ
単位で行われるが,各プロセッサのプライベート空間は
仮想空間上で元々独立しているので,共有空間のように
偏りが生じることなく,そのままの大きさで物理空間に
マッピングされる。
ライベート空間に,それぞれ50要素または51要素ずつ配
置される。配列b'の最後の(51番目)の要素は,shadow領
域である。これらは,各プロセッサの物理空間にマッピ
ングされなければならない。マッピングはやはりページ
単位で行われるが,各プロセッサのプライベート空間は
仮想空間上で元々独立しているので,共有空間のように
偏りが生じることなく,そのままの大きさで物理空間に
マッピングされる。
【0037】前述のライブラリルーチンOMD_array_copy
の中では,自プロセッサのプライベート空間と対応する
共有空間との間で配列をコピーする。例えば,図4の55
行目の call OMD_array_copy(a,a',...) では,P0は共有空間のa(1:50)を自身のプライベート空
間のa'(1:50)にコピーし,P1は共有空間のa(51:100)を
自身のプライベート空間のa'(1:50)にコピーする。ま
た,図4の56行目の call OMD_array_copy(b,b',...) では,配列bおよびb'について同様の処理を行う。58行
目および59行目では逆方向のコピーを実行する。
の中では,自プロセッサのプライベート空間と対応する
共有空間との間で配列をコピーする。例えば,図4の55
行目の call OMD_array_copy(a,a',...) では,P0は共有空間のa(1:50)を自身のプライベート空
間のa'(1:50)にコピーし,P1は共有空間のa(51:100)を
自身のプライベート空間のa'(1:50)にコピーする。ま
た,図4の56行目の call OMD_array_copy(b,b',...) では,配列bおよびb'について同様の処理を行う。58行
目および59行目では逆方向のコピーを実行する。
【0038】以上が基本的な配列コピー方法であるが,
これをさらに改良した別の方法もある。この別方法を以
下に説明する。
これをさらに改良した別の方法もある。この別方法を以
下に説明する。
【0039】図7に,別方法における配列のメモリ配置
を示す。本方法では,仮想プライベート空間を予め物理
空間へマッピングしておくのではなく,配列コピーコー
ドの中でマッピングを行う。このとき,仮想共有空間上
のコピーすべきページが既に自プロセッサの物理空間に
マッピングされているならば,内容をコピーする代わり
にコピー先仮想プライベート空間のページを当該物理空
間のページにマッピングする。例えば,図7において,
共有空間の配列bの最初の2ページは,コピー先がP0のプ
ライベート空間であり,かつ,既にP0の物理空間にマッ
ピングされているから,P0のプライベート空間の配列b'
の最初の2ページを当該物理ページにマッピングする。
b'の残りの部分,すなわち最後の10要素および1要素分
のshadow領域については,新たな物理ページを確保して
そこにマッピングする。P1のプライベート空間の配列b'
については,後の2ページが既にP1の物理空間にマッピ
ングされているので,これらについては当該物理ページ
にマッピングする。残りの部分,すなわち最初の10要素
および最後のshadow領域については,それぞれ新たな物
理ページを確保してそこにマッピングする。最初の10要
素とshadow領域とはメモリ上連続でないので,別々のペ
ージにマッピングしなければならない。
を示す。本方法では,仮想プライベート空間を予め物理
空間へマッピングしておくのではなく,配列コピーコー
ドの中でマッピングを行う。このとき,仮想共有空間上
のコピーすべきページが既に自プロセッサの物理空間に
マッピングされているならば,内容をコピーする代わり
にコピー先仮想プライベート空間のページを当該物理空
間のページにマッピングする。例えば,図7において,
共有空間の配列bの最初の2ページは,コピー先がP0のプ
ライベート空間であり,かつ,既にP0の物理空間にマッ
ピングされているから,P0のプライベート空間の配列b'
の最初の2ページを当該物理ページにマッピングする。
b'の残りの部分,すなわち最後の10要素および1要素分
のshadow領域については,新たな物理ページを確保して
そこにマッピングする。P1のプライベート空間の配列b'
については,後の2ページが既にP1の物理空間にマッピ
ングされているので,これらについては当該物理ページ
にマッピングする。残りの部分,すなわち最初の10要素
および最後のshadow領域については,それぞれ新たな物
理ページを確保してそこにマッピングする。最初の10要
素とshadow領域とはメモリ上連続でないので,別々のペ
ージにマッピングしなければならない。
【0040】図8は,図7のメモリ配置を用いる場合の配
列コピー方法を表すフローチャートである。これはライ
ブラリルーチンOMD_array_copyの中で共有空間からプラ
イベート空間へ配列をコピーする場合に実行される処理
である。
列コピー方法を表すフローチャートである。これはライ
ブラリルーチンOMD_array_copyの中で共有空間からプラ
イベート空間へ配列をコピーする場合に実行される処理
である。
【0041】図8のステップ120では,共有空間上のコピ
ーすべき配列を構成する各ページにつき,コピー処理が
済んでいるかどうかを判定し,未処理であれば当該ペー
ジに対してステップ121以降を実行する。
ーすべき配列を構成する各ページにつき,コピー処理が
済んでいるかどうかを判定し,未処理であれば当該ペー
ジに対してステップ121以降を実行する。
【0042】ステップ121では,当該ページに含まれる
全要素についてコピー先プロセッサが同一かどうかを判
定する。これは必ずしも個々の要素について一つずつ判
定する必要はなく,より簡単な判定,例えば,プロセッ
サへの分割境界のインデックス(図7では50)がページ内
に含まれるかどうかの判定を行えば十分である。もしコ
ピー先プロセッサが同一だと判明したらステップ122に
進み,そうでない場合はステップ124に進む。
全要素についてコピー先プロセッサが同一かどうかを判
定する。これは必ずしも個々の要素について一つずつ判
定する必要はなく,より簡単な判定,例えば,プロセッ
サへの分割境界のインデックス(図7では50)がページ内
に含まれるかどうかの判定を行えば十分である。もしコ
ピー先プロセッサが同一だと判明したらステップ122に
進み,そうでない場合はステップ124に進む。
【0043】ステップ122では,当該ページが既にコピ
ー先プロセッサの物理空間にマッピングされているかど
うかを判定する。もしそうならばステップ123に進み,
さもなくばステップ124に進む。
ー先プロセッサの物理空間にマッピングされているかど
うかを判定する。もしそうならばステップ123に進み,
さもなくばステップ124に進む。
【0044】ステップ123では,コピー先のプライベー
ト空間の当該仮想ページを,ステップ122で求められた
物理ページにマッピングする。
ト空間の当該仮想ページを,ステップ122で求められた
物理ページにマッピングする。
【0045】一方,ステップ124に進んだ場合は,コピ
ー先プロセッサの物理空間に新たなページを確保する。
そしてステップ125で,コピー先プライベート空間の仮
想ページを当該物理ページにマッピングする。さらにス
テップ126で,共有空間のページに含まれる配列要素を
確保した当該物理ページに実際にコピーする。
ー先プロセッサの物理空間に新たなページを確保する。
そしてステップ125で,コピー先プライベート空間の仮
想ページを当該物理ページにマッピングする。さらにス
テップ126で,共有空間のページに含まれる配列要素を
確保した当該物理ページに実際にコピーする。
【0046】以上で配列コピー処理の別方法の説明を終
わる。
わる。
【0047】最後に,図1のコンパイル方法を実行する
コンパイラの構成を示す。図9はそのようなコンパイラ
の構成の一例である。コンパイラ4には,構文解析部4
0,配列配置先空間決定部41,中間語変換部42,および
出力プログラム生成部43が含まれる。
コンパイラの構成を示す。図9はそのようなコンパイラ
の構成の一例である。コンパイラ4には,構文解析部4
0,配列配置先空間決定部41,中間語変換部42,および
出力プログラム生成部43が含まれる。
【0048】構文解析部40は,入力プログラム20を読み
込んで,中間語50を生成する。入力プログラムの内容は
例えば図2に示したようなものである。中間語50はコン
パイラ内部でのプログラムの表現形式であり,その形式
は通常のコンパイラの場合と特に変わらないので,ここ
では特に詳細は述べない。構文解析部の中に含まれる指
示文解析部400は,指示文を読み込む処理を実行する。
すなわち,図1のステップ100,およびステップ105の処
理を実行し,配列表30に情報を登録する。
込んで,中間語50を生成する。入力プログラムの内容は
例えば図2に示したようなものである。中間語50はコン
パイラ内部でのプログラムの表現形式であり,その形式
は通常のコンパイラの場合と特に変わらないので,ここ
では特に詳細は述べない。構文解析部の中に含まれる指
示文解析部400は,指示文を読み込む処理を実行する。
すなわち,図1のステップ100,およびステップ105の処
理を実行し,配列表30に情報を登録する。
【0049】配列配置空間決定部41は,配列表30の中の
指示文情報を参照して,配列を共有空間に配置するかプ
ライベート空間に配置するかを決定する。これは図1の
ステップ102からステップ104の処理である。
指示文情報を参照して,配列を共有空間に配置するかプ
ライベート空間に配置するかを決定する。これは図1の
ステップ102からステップ104の処理である。
【0050】中間語変換部42には,通信コード生成部42
0および配列コピーコード生成部421およびメモリアクセ
ス命令生成部422が含まれる。通信コード生成部420は,
中間語内のreflect指示文にしたがって,通信コードの
中間語を生成する。これは図1のステップ106の処理であ
る。配列コピーコード生成部421は,手続き呼出しの実
引数と仮引数の配置先空間に応じて,配列コピーコード
の中間語を生成する。これは図1のステップ108の処理で
ある。メモリアクセス命令生成部422は,共有空間に配
置された配列の参照に対してメモリアクセス生成する。
これは図1のステップ109の処理である。
0および配列コピーコード生成部421およびメモリアクセ
ス命令生成部422が含まれる。通信コード生成部420は,
中間語内のreflect指示文にしたがって,通信コードの
中間語を生成する。これは図1のステップ106の処理であ
る。配列コピーコード生成部421は,手続き呼出しの実
引数と仮引数の配置先空間に応じて,配列コピーコード
の中間語を生成する。これは図1のステップ108の処理で
ある。メモリアクセス命令生成部422は,共有空間に配
置された配列の参照に対してメモリアクセス生成する。
これは図1のステップ109の処理である。
【0051】出力プログラム生成部43は,変換された中
間語50を読み込んで,出力プログラム21を生成する。出
力プログラム21の内容は例えば図4に示したようなもの
である。本実施形態の説明では図4のように高級言語ソ
ースプログラムの形式を用いて説明したが,代わりに機
械語形式を用いることもできる。出力プログラム生成部
43の内容は通常の特に変わらないので,これ以上の詳細
は述べない。
間語50を読み込んで,出力プログラム21を生成する。出
力プログラム21の内容は例えば図4に示したようなもの
である。本実施形態の説明では図4のように高級言語ソ
ースプログラムの形式を用いて説明したが,代わりに機
械語形式を用いることもできる。出力プログラム生成部
43の内容は通常の特に変わらないので,これ以上の詳細
は述べない。
【0052】以上で,本発明の実施形態の説明を終わ
る。
る。
【0053】
【発明の効果】本発明によれば,ユーザは,まず配列を
共有空間に配置してDSM機構を用いることによってプロ
グラムを動作させ,次に,プログラム中のカーネル部分
に対してのみ,最高性能を引出すためのチューニングを
行うという,インクリメンタルなプログラム開発が可能
になる。また,本発明によればさらに,上記チューニン
グによって発生する配列コピーにおいて,データ移動量
やメモリ消費量を最小限に抑え,プログラムの実行性能
を向上させることができる。
共有空間に配置してDSM機構を用いることによってプロ
グラムを動作させ,次に,プログラム中のカーネル部分
に対してのみ,最高性能を引出すためのチューニングを
行うという,インクリメンタルなプログラム開発が可能
になる。また,本発明によればさらに,上記チューニン
グによって発生する配列コピーにおいて,データ移動量
やメモリ消費量を最小限に抑え,プログラムの実行性能
を向上させることができる。
【図1】本発明のコンパイル方法の一例を示すフローチ
ャートである。
ャートである。
【図2】本発明のコンパイラへの入力プログラムの例で
ある。
ある。
【図3】コンパイラ内部で利用される配列表である。
【図4】図2の入力プログラムに本発明を適用したこと
によって得られる出力プログラムである。
によって得られる出力プログラムである。
【図5】図4の出力プログラムにおける配列要素の参照
に対する機械語の命令列である。
に対する機械語の命令列である。
【図6】本発明における配列のメモリ配置を示す図であ
る。
る。
【図7】本発明における配列のメモリ配置の別方法を示
す図である。
す図である。
【図8】図7のメモリ配置を用いだ場合の,共有空間か
らプライベート空間への配列コピー処理を示すフローチ
ャートである。
らプライベート空間への配列コピー処理を示すフローチ
ャートである。
【図9】本発明のコンパイル方法を実施するコンパイラ
の構成例である。
の構成例である。
4…コンパイラ,20…入力プログラム,21…出力プログ
ラム,30…配列表,400…指示文解析部,41…配列配置
空間決定部,420…通信コード生成部,421…配列コピー
コード生成部,422…メモリアクセス命令生成部。
ラム,30…配列表,400…指示文解析部,41…配列配置
空間決定部,420…通信コード生成部,421…配列コピー
コード生成部,422…メモリアクセス命令生成部。
───────────────────────────────────────────────────── フロントページの続き (72)発明者 佐藤 真琴 神奈川県川崎市麻生区王禅寺1099番地 株 式会社日立製作所システム開発研究所内 Fターム(参考) 5B045 DD04 GG11 5B081 BB08 CC28 CC29 CC32
Claims (4)
- 【請求項1】分散共有メモリ機構を備える並列システム
向けのコンパイル方法であって,配列を共有空間に配置
するか各プロセッサのプライベート空間に配置するかを
決定するステップと,共有空間に配置した配列の参照に
対して当該配列のアドレスへのメモリアクセス命令を生
成するステップと,プライベート空間に配置した配列に
対してプロセッサ間通信コードを生成するステップと,
共有空間とプライベート空間の間での配列コピーコード
を生成するステップとを含むことを特徴とする分散共有
メモリ向けコンパイル方法。 - 【請求項2】請求項1のコンパイル方法であって,該配
列コピーコードによって実行される処理は,共有空間上
のページが既にコピー先プロセッサの物理メモリに既に
割り付けられているかどうかを判定するステップと,既
に割り付けられている場合はコピー先の仮想プライベー
ト空間上のページを当該物理ページにマッピングするス
テップとを含むことを特徴とする分散共有メモリ向けコ
ンパイル方法。 - 【請求項3】分散共有メモリ機構を備える並列システム
向けのコンパイル方法であって,配列の配置空間を指定
する第1のユーザ指示および通信コード生成を指定する
第2のユーザ指示を読み込むステップと,該第1のユー
ザ指示に従って配列を共有空間に配置するか各プロセッ
サのプライベート空間に配置するかを決定するステップ
と,共有空間に配置した配列の参照に対して当該配列の
アドレスへのメモリアクセス命令を生成するステップ
と,該第2のユーザ指示に従ってプライベート空間に配
置した配列に対してプロセッサ間通信コードを生成する
ステップと,共有空間とプライベート空間の間で配列を
コピーするコードを生成するステップとを含むことを特
徴する分散共有メモリ向けコンパイル方法。 - 【請求項4】分散共有メモリ機構を備える並列システム
向けのコンパイル方法をコンピュータに実行させるため
の処理手順プログラムを格納するコンピュータ読み込み
可能な記録媒体であって,前記方法は、 配列を共有空間に配置するか各プロセッサのプライベー
ト空間に配置するかを決定する手順と,共有空間に配置
した配列の参照に対して当該配列のアドレスへのメモリ
アクセス命令を生成する手順と,プライベート空間に配
置した配列に対してプロセッサ間通信コードを生成する
手順と,共有空間とプライベート空間の間での配列コピ
ーコードを生成する手順とを有することを特徴とする記
録媒体。
Priority Applications (1)
Application Number | Priority Date | Filing Date | Title |
---|---|---|---|
JP2000264491A JP2002073579A (ja) | 2000-08-29 | 2000-08-29 | 分散共有メモリ向けコンパイル方法 |
Applications Claiming Priority (1)
Application Number | Priority Date | Filing Date | Title |
---|---|---|---|
JP2000264491A JP2002073579A (ja) | 2000-08-29 | 2000-08-29 | 分散共有メモリ向けコンパイル方法 |
Publications (1)
Publication Number | Publication Date |
---|---|
JP2002073579A true JP2002073579A (ja) | 2002-03-12 |
Family
ID=18751896
Family Applications (1)
Application Number | Title | Priority Date | Filing Date |
---|---|---|---|
JP2000264491A Pending JP2002073579A (ja) | 2000-08-29 | 2000-08-29 | 分散共有メモリ向けコンパイル方法 |
Country Status (1)
Country | Link |
---|---|
JP (1) | JP2002073579A (ja) |
Cited By (4)
Publication number | Priority date | Publication date | Assignee | Title |
---|---|---|---|---|
JP2009211458A (ja) * | 2008-03-05 | 2009-09-17 | Nec Corp | コンパイラ、変数最適化装置、方法、及び、プログラム |
US7698696B2 (en) | 2002-07-03 | 2010-04-13 | Panasonic Corporation | Compiler apparatus with flexible optimization |
JP2012530995A (ja) * | 2009-06-26 | 2012-12-06 | コードプレイ、ソフトウェア、リミテッド | 処理方法 |
CN110308908A (zh) * | 2018-03-20 | 2019-10-08 | 北京小米移动软件有限公司 | 应用的配置文件的生成和应用页面的展示方法、装置及存储介质 |
-
2000
- 2000-08-29 JP JP2000264491A patent/JP2002073579A/ja active Pending
Cited By (6)
Publication number | Priority date | Publication date | Assignee | Title |
---|---|---|---|---|
US7698696B2 (en) | 2002-07-03 | 2010-04-13 | Panasonic Corporation | Compiler apparatus with flexible optimization |
US8418157B2 (en) | 2002-07-03 | 2013-04-09 | Panasonic Corporation | Compiler apparatus with flexible optimization |
JP2009211458A (ja) * | 2008-03-05 | 2009-09-17 | Nec Corp | コンパイラ、変数最適化装置、方法、及び、プログラム |
JP2012530995A (ja) * | 2009-06-26 | 2012-12-06 | コードプレイ、ソフトウェア、リミテッド | 処理方法 |
CN110308908A (zh) * | 2018-03-20 | 2019-10-08 | 北京小米移动软件有限公司 | 应用的配置文件的生成和应用页面的展示方法、装置及存储介质 |
CN110308908B (zh) * | 2018-03-20 | 2023-07-18 | 北京小米移动软件有限公司 | 应用的配置文件的生成和应用页面的展示方法、装置及存储介质 |
Similar Documents
Publication | Publication Date | Title |
---|---|---|
Reinders et al. | Data parallel C++: mastering DPC++ for programming of heterogeneous systems using C++ and SYCL | |
US7926046B2 (en) | Compiler method for extracting and accelerator template program | |
JP5851396B2 (ja) | 処理方法 | |
Yang et al. | High-performance computing: paradigm and infrastructure | |
JP2669603B2 (ja) | コンパイラにおけるコード生成方法及びコンパイラ | |
KR20110093965A (ko) | 고수준 언어 코드를 hdl 코드로 변환하는 방법 및 시스템 | |
Hormati et al. | Macross: Macro-simdization of streaming applications | |
KR20160003485A (ko) | 원격 함수 호출을 위한 자동 코드 생성 | |
Ebcioglu et al. | Optimizations and oracle parallelism with dynamic translation | |
Moses et al. | Scalable automatic differentiation of multiple parallel paradigms through compiler augmentation | |
Theobald et al. | Overview of the Threaded-C language | |
US11599478B2 (en) | Reduced instructions to generate global variable addresses | |
JP4830108B2 (ja) | プログラム処理装置、プログラム処理方法、並列処理プログラム用コンパイラおよび並列処理プログラム用コンパイラを格納した記録媒体 | |
GB2394085A (en) | Generating code for a configurable microprocessor | |
JP2002073579A (ja) | 分散共有メモリ向けコンパイル方法 | |
Min et al. | Portable compilers for OpenMP | |
Matz et al. | Automated partitioning of data-parallel kernels using polyhedral compilation | |
Brent | Using program structure to achieve prefetching for cache memories | |
Chandramohan et al. | A compiler framework for automatically mapping data parallel programs to heterogeneous MPSoCs | |
Sakaguchi et al. | Programmability and Performance of New Global-View Programming API for Multi-Node and Multi-Core Processing | |
Kalra | Design and evaluation of register allocation on gpus | |
Lossing et al. | Automatic code generation of distributed parallel tasks | |
van der Wijst | An Accelerator based on the ρ-VEX Processor: an Exploration using OpenCL | |
Shabanov et al. | Features of Dataflow Processor Emulator Implementing | |
Ku | The design of an efficient and portable interface between a parallelizing compiler and its target machine |