DE19616199B4

DE19616199B4 - Rechenwerk für einen Computertomographen

Info

Publication number: DE19616199B4
Application number: DE19616199A
Authority: DE
Inventors: Jürgen Dipl.-Ing. Wohlrab (FH)
Original assignee: Siemens AG
Current assignee: Siemens AG
Priority date: 1996-04-23
Filing date: 1996-04-23
Publication date: 2005-02-10
Anticipated expiration: 2016-04-24
Also published as: DE19616199A1; JPH1066692A; US5802135A

Abstract

Computertomograph mit einem Rechner (6) zur Rekonstruktion eines Schnittbildes, bei dem ein spezielles Rechenwerk in Form von
a) einem oder mehreren Prozessoren PROC,
b) einem oder mehreren Medianfilter-ASICs MAS,
c) einem oder mehreren Winkelverschleifungs-ASICs WAS und
d) zwei Speichern DAT und ADR
zur Realisierung eines adaptiven Bildfilters zur Artefaktreduktion vorgesehen sind (2 und 3),
wobei die Realisierung des Medianfilters des adaptiven Bildfilters zur Artefaktreduktion mit nachfolgender Pipeline mit pixelweisen Durchlauf vorgenommen wird:
– Bestimmung der Inkremente zur Generierung der Stützstellenadressen durch einen oder mehr Prozessoren PROC
– Berechnung der Zugriffsadressen und Adressierung der Stützwerte vom Speicher DAT durch das Medianfilter-ASIC MAS
– Medianwertbildung durch das Medianfilter-ASIC MAS
– Differenzbilderzeugung durch das Medianfilter-ASIC MAS, und wobei die Realisierung der Winkelverschleifung des adaptiven Bildfilters zur Artefaktreduktion mit nachfolgender Pipeline mit pixelweisen Durchlauf vorgenommen wird:
– Bestimmung der Koordinaten des Pixels in...

Description

In der Computertomographie können bei Geräten mit umlaufendem Röntgenstrahler-Detektorsystem Ringartefakte in den rekonstruierten Bildern auftreten. Es kann ein Algorithmus angewendet werden, der in Form von einer Bildnachverarbeitung die Ringartefakte unterdrückt. Dieser Algorithmus wird nachfolgend als adaptives Filter bezeichnet.

Der Erfindung liegt die Aufgabe zugrunde, das adaptive Filter für einen Computertomographen zur Artefaktreduktion effizient zu implementieren.

Diese Aufgabe ist erfindungsgemäß gelöst durch die Merkmale des Patentanspruches. Bei dem erfindungsgemäßen Rechenwerk für einen Computertomographen ist eine Spezialhardware in Form von einem Prozessor PROC, einem Medianfilter-ASIC MAS, einem Winkelverschleifungs-ASIC WAS und zwei Speichern DAT und ADR vorgeschlagen.

In EP 0 544 507 ist zwar ein Computertomograph mit Mitteln zur Reduktion von Bildartefakten beschrieben, bei dem ein Spezialzweckcomputer vorgesehen ist, die spezielle Lösung mit Medianfilter-ASICs und Winkelverschleifungs-ASICs gemäß dem Patentanspruch ergibt sich daraus jedoch nicht. In der US-Patentschrift 4,513,440 ist ferner ein Medianfilter zur Verarbeitung von hochauflösenden Bilddaten beschrieben. Die Kombination eines solchen Medianfilters mit einem Winkelverschleifungs-ASIC ergibt sich daraus aber nicht.

Die Erfindung ist nachfolgend anhand von Zeichnungen näher erläutert. Es zeigen:
1 die wesentlichen Teile eines Computertomographen mit Fächerstrahl zur Erläuterung des Erfindungsgedankens, und
2 bis 6 graphische Darstellungen zur Erläuterung der Erfindung.
In der 1 ist ein Röntgenstrahler 1 und ein aus einer Matrix von Detektorelementen bestehender Detektor 2 dargestellt. Die Komponenten 1 und 2 rotieren zur Abtastung eines Objektes 3 um eine Systemachse 4, so daß das Objekt 3 unter verschiedenen Richtungen von dem vom Röntgenstrahler 1 ausgehenden fächer förmigen Röntgenstrahlbündel 5 durchstrahlt wird. Die von den Detektorelementen des Detektors 3 gelieferten Daten werden einem Rechner 6 zugeführt, der Schichtbilder innerhalb des durchstrahlten Volumens des Objektes 3 rekonstruiert, welche auf einem Monitor 7 wiedergegeben werden. Das Objekt 3 liegt dabei in einem durch das Röntgenstrahlbündel 5 erfaßten Meßfeld 8.
Die nachfolgend beschriebene vorschlagene Spezialhardware nutzt die maximale Zugriffsbandbreite zum Datenspeicher DAT voll aus und stellt somit ein optimales Rechenwerk für das adaptive Filter im Rechner 6 dar. Bei einem Multiprozessorsystem kann die maximale Zugriffsbandbreite zum Datenspeicher aufgrund von Arbitrierungsverlusten nicht voll genutzt werden.
Die grundsätzlichen Rechenschritte des adaptiven Filters sind

A) Abschneiden von Knochen und Luft: Begrenzung aller Pixelwerte eines Eingangsbildes INB auf eine vorgebbare obere und untere Schwelle. Dieser Rechenschritt ergibt eine neue Bildmatrix KLB.
B) eine erste Medianfilterung der gesamten Matrix KLB auf einer Geraden in Richtung vom jeweiligen Pixel zu einem vorgebbaren Drehzentrum Z (xc, yc). Das Raster der M1 Stützstellen auf dieser Geraden sei durch bildspezifische Parameter vorgebbar. Die somit erhaltenen Pixelwerte werden von den Pixeln der Matrix KLB abgezogen und die Differenz einer Schwellwertbewertung mit einer vorgebbaren Schwelle SART unterworfen. Somit erhält man eine neue Differenzbildmatrix UDB'.
C) eine zweite Medianfilterung einer drehzentrumsnahen Submatrix der Differenz der Matrizen KLB und UDB' auf einer Geraden in Richtung vom jeweiligen Pixel zum Drehzentrum Z (xc, yc). Das Raster der M2 Stützstellen auf dieser Geraden sei durch bildspezifische Parameter vorgebbar. Die somit erhaltenen Pixelwerte werden von den Pixeln der Matrix KLB abgezogen und die Differenz einer Schwellwertbewertung mit einer vorgebbaren Schwelle SART unterworfen. Somit erhält man eine neue Submatrix UDB''. Die Submatrix UDB'' wird in die Matrix UDB' kopiert. Dadurch erhält man die unverschliffene Differenzbildmatrix UDB.
D) Winkelverschleifung: Mittelung über W Stützstellen der Matrix UDB auf einem vom jeweiligen Pixel in beide Richtungen ausgehenden Kreisbogen um das Drehzentrum Z. Die Pixel werden dabei je nach Abstand zum Drehzentrum in Pixel im Innenbereich (Flag WIAU = 0) und Pixel im Außenbereich (Flag WIAU = 1) eingeteilt. Das Winkelinkrement zwischen den Stützstellen ist im Innen- und Außenbereich unterschiedlich.
E) Ergebnisbilderzeugung: Die winkelverschliffenen Differenzbildpixel, die sich aus dem Schritt D ergeben, werden von der Eingangsbildmatrix INB abgezogen und die erhaltenen Pixelwerte auf den Bereich [0,4095] begrenzt. Man erhält somit das Ergebnisbild OUTB.

In 2 ist der Minimalausbau der vorschlagenen Spezialhardware dargestellt. Sie besteht aus einem Prozessor PROC, einem Medianfilter-ASIC MAS, einem Winkelverschleifungs-ASIC WAS und zwei Speichern DAT und ADR.
Das Datenformat der Matrizen im Datenspeicher DAT ist 16 Bit Integer. Die Adreßrechnungen werden in Gleitkomma durchgeführt und die Ergebnisse auf Datenformat Integer gerundet.
Die Schritte A, B und C sind in ein Medianfilter-ASIC MAS integriert.
Die Schritte D und E sind in ein Winkelverschleifungs-ASIC WAS integriert.
Die ASICs MAS und WAS sollen von einem Prozessor PROC parametriert, angesteuert und überwacht werden. Der Prozessor PROC stellt außerdem eine intelligente, programmierbare Verbindung zum System-Bus dar, der eine Anbindung zum Computertomographen realisiert.
Der Prozessor PROC soll zur Unterstützung der beiden ASICs nachfolgende Berechnungen erledigen, auf dessen Ergebnisse die ASICs über den SP-Bus zugreifen können: Für jedes Pixel muß die Richtung vom Pixel zum Drehzentrum in Form einer x-Komponente DX und einer y-Komponente DY berechnet werden. Ebenso muß für jedes Pixel das Flag WIAU bestimmt werden, das angibt, ob das Pixel im Innen- oder im Außenbereich liegt. Für diese Berechnungen ist u.a. die Operation 1/SQURT(x) notwendig, die nicht die ASICs integiert werden soll. Diese Berechnungen erfolgen vom Prozessor PROC mit einem geringen zeitlichen Vorlauf parallel zur Abarbeitung des Medianfilters durch das Medianfilter-ASIC MAS. Die Verarbeitungsleistung des Medianfilter-ASICs MAS ist an die Verarbeitungsleisung der Prozessors PROC anzupassen, damit die vom Prozessor PROC berechneten Werte DX und DY eines Pixels bei der Bearbeitung des Pixels durch das Medianfilter-ASIC MAS bereits zur Verfügung stehen.
Die vom Prozessor PROC berechneten Werte DX werden in der Matrix MDX abgelegt, die Werte DY in der Matrix MDY und die Werte WIAU in der Matrix MWIAU.
Die ASICs MAS und WAS führen die Zugriffe auf die Datenfelder im Speicher DAT über einen 16-Bit-Bus D aus. Der Prozessor PROC soll über einen Wandler (siehe 2) über den Bus D auf den Speicher DAT per DMA zugreifen können, um folgende Aufgaben zu erledigen:

a) Schreiben des Eingangsbildes vor dem Start der Berechnungen,
b) Lesen des Ergebnisbildes nach dem Ende der Berechnungen,
c) Debugging im Falle eines während den Berechnungen aufgetretenen Fehlers.

Die Einführung eines zweiten Busses D neben dem Bus SP, an dem der Prozessor PROC hängt, bringt folgende Vorteile:

a) Zugriff des Prozessors PROC beim Schreiben der Matrizen MDX, MDY und MWIAU stört den Zugriff des Medianfilter-ASICs MAS auf die Datenfelder im Speicher DAT nicht.
b) Paralleler Zugriff auf die Matrizen MDX/MDY im Speicher ADR und KLB/UDB im Speicher DAT beim Medianfilter möglich.
c) Keine Arbitrierungsverluste beim Zugriff auf den Speicher DAT über den Bus D. Die ASICs MAS und WAS können vom Ablauf des Algorithmus her nicht konkurrierend auf den Speicher DAT zugreifen. An dem Bus SP, an dem der Prozessor PROC angeschlossen ist, ist wegen des nicht voraussehbaren Zugriffszeitpunktes des Prozessors PROC eine Arbitrierungslogik notwendig, die eine erhöhte Zugriffszeit auf dem BUS SP bedingt.

Für die Matrizen UDB' und UDB'' sind keine eigenen Datenfelder reserviert. Die Ergebnismatrix UDB' des 1. Medianfilters wird im Datenfeld UDB abgelegt. Während dem 2. Medianfilter werden die Matrizen KLB und UDB' (UDB' darf bis zum Rückkopieren am Ende des 2. Medianfilters nicht zerstört werden) als Eingangsdaten benötigt. Die Ausgangsdaten des 2. Medianfilters ist die Submatrix UDB'', die im Datenfeld OUTB angelegt wird. Am Ende des 2. Medianfilters wird dann die Submatrix UDB'' aus dem Datenfeld OUTB in die Matrix UDB' im Datenfeld UDB kopiert. Nach diesem Rückkopieren ist die Matrix UDB vollständig aufgebaut.
Dann kann die Winkelverschleifung beginnen. Die winkelverschliffenen Differenzbildpixel werden nicht als Datenfeld im Speicher DAT abgelegt, stattdessen wird gleich schrittweise die Ergebnisbilderzeugung von Schritt E vollzogen.
Es ist folgender Ablauf einzuhalten:
Schritt A wird vollständig abgearbeitet, bevor Schritt B beginnt. Der Prozessor PROC kann die Berechnungen der Datenfelder MDX, MDY und MWIAU schon während dem Ablauf von Schritt A beginnen.
Schritt B wird vollständig abgearbeitet, bevor Schritt C beginnt.
Schritt C wird vollständig abgearbeitet, bevor Schritt D beginnt.
Schritt D und E erfolgen für jedes Pixel schrittweise hintereinander.
Sofern durch 1 Medianfilter-ASIC MAS und 1 Winkelverschleifungs-ASIC WAS die Busse D und SP nicht voll ausgelastet sind, können mehrere Medianfilter-ASICs MAS1, MAS2, ..., MASm bzw. mehrere Winkelverschleifungs-ASICs WAS1, WAS2, ..., WASn parallel eingesetzt werden, um die Verarbeitung zu beschleunigen. Da die Verarbeitungsleistung des Medianfilter-ASICs an die Verarbeitungsleistung des Prozessors PROC angepaßt ist, sind genauso viele Prozessoren PROC1, PROC2, ..., PROCm wie Medianfilter-ASICs zu verwenden. Die Verarbeitung des Medianfilters wird in diesem Fall pixelweise auf mehrere Medianfilter-ASICs aufgeteilt, bzw. die Verarbeitung der Winkelverschleifung auf mehrere Winkelverschleifungs-ASICs. In 3 ist dies für den Fall m = 3 und n = 2 dargestellt.
Sofern mehrere ASICs MAS bzw. mehrere ASICs WAS parallel eingesetzt werden, greifen die ASICs untereinander konkurrierend auf den Speicher DAT zu. Hier ist eine vorausschauende Zugriffsarbitrierung D-Sync auf dem Bus D notwendig, die Arbitrierungsverluste beim Zugriff auf den Speicher DAT vermeidet. Die vorausschauende Zugriffslogik ist in die ASICs MAS und WAS eingebaut. Die Arbitrierung zwischen den ASICs MAS untereinander, bzw. unter den ASICs WAS untereinander findet vorausschauend bereits einige Takte vor dem Zugriffszeitpunkt statt. Da die ASICs MAS untereinander, bzw. die ASICs WAS untereinander die gleiche Verarbeitungsleistung haben, ist die Arbitrierung schon nach wenigen Zugriffen eingeschwungen, und es kommt zu keinerlei Arbitrierungsverlusten mehr. Die Verarbeitung auf den ASICs läuft dann entsprechend verzahnt ab, so daß sich die ASICs MAS untereinander, bzw. die ASICs WAS untereinander nicht mehr behindern.
Die Anzahl der ASICs MAS und WAS kann zum Erreichen der optimalen Performance des Rechenwerkes für das adaptive Filter bis zur vollständigen Auslastung der Busse D und SP erhöht werden. Ist einer der Busse D oder SP voll ausgelastet, bringt eine größere Anzahl von ASICs keine Performance-Steigerung des Rechenwerkes für das adaptive Filter mit sich.
Abbildung der Rechenschritte auf das Rechenwerk des adaptiven Filters und Abschätzung der Verarbeitungsleistung:
Die Abschätzung der Verarbeitungsleistung soll in Rechentakten vorgenommen werden. Beim Zugriff auf dem Bus D werden 0 Wait States angenommen (keine Arbitrierung), beim Zugriff auf dem Bus SP wird 1 Wait State angenommen (1 Takt Arbitrierung). Es wird von einer Bildmatrix mit MATDIM·MATDIM Pixeln ausgegangen. Der Wertebereich eines Pixels sei [0,4095].
Nachfolgend aufgeführte Rechenwerke sollen nach Möglichkeit in die ASICs integriert werden. Sofern die ASICs die Unterbringung dieser Rechenwerke aufgrund ihrer Komplexität nicht zulassen, sind die Rechenwerke von einem ASIC auf mehrere ASICs sinnvoll zu verteilen.
Medianfilter-ASIC:

a) 1 Integer-Comparator (16 Bit) mit nachgeschalteter Multiplexlogik (für Clipping gemäß Rechenschritt A bzw. der Schwellwertbewertung von Schritt B und C bzw. Begrenzung von Koordinaten auf Bilddimension)
b) 1 Integer-Comparator (16 Bit) mit nachgeschalteter Multiplexlogik (für Medianwertbildung von Schritt B und C)
c) 1 NINT-Operator (Gleitkomma → 16 Bit)
d) 1 Gleitkomma-Multiplizierer
e) 1 Integer-Addierer (32 Bit)
f) 1 Integer-Multiplizierer (16 Bit Eingang, 32 Bit Ausgang)

Winkelverschleifungs-ASIC:

a) 4 Gleitkomma-Multiplizierer: GMUL1, GMUL2, GMUL3, GMUL4
b) 3 Gleitkomma-Addierer: GADD1, GADD2, GADD3
c) 2 Integer-Comparatoren (16 Bit): COMP1, COMP2
d) 1 NINT-Operator: NINT
e) 1 FLOAT-Operator: FLOAT
f) 2 Integer-Addierer (32 Bit): IADD1, IADD2
g) 1 Integer-Multiplizierer (16 Bit Eingang, 32 Bit Ausgang): IMUL
h) 1 Integer-Comparator (16 Bit) mit nachgeschalteter Multiplexlogik (für Clipping der Pixelwerte auf den gültigen Pixel-Wertebereich)

Rechenschritt A: Abschneiden von Knochen und Luft
Hierzu wird die Comparator-Schaltung von 4 verwendet. Es sind für jedes Pixel folgende Schritte notwendig:

1.) Einlesen eines Pixels aus der Matrix INB vom Speicher DAT (1 Takt)
2.) 2 Vergleiche des eingelesenen Pixelwertes mit den Grenzen SWOMOD und SWUMOD, Durchführung mit Integer-Comparator (2 Takte)
3.) Vom Integer-Comparator gesteuerte Multiplexlogik zur Generierung des entsprechenden Output-Enable für das Schreiben des Ausgangspixels in die Matrix KLB (1 Takt)
4.) Inkrementierung der Pixeladresse (1 Takt)
5.) Schreiben des Pixels in die Matrix KLB im Speicher DAT (1 Takt).

Unter der Annahme, daß Schritt 3 und 4 und die beiden Zugriffe auf dem D-Bus hinter den beiden Vergleichen von Schritt 2 versteckt werden können, ergibt sich eine Verarbeitungsleistung von 2 Takten pro Pixel.
Der Rechenschritt A erfordert somit 2·MATDIM·MATDIM Rechentakte.
Rechenschritt B: 1. Medianfilter
Zuerst sind die Inkremente für die Generierung der Stützwertadressen durch den Prozessor PROC zu bestimmen. Die nachfolgende Abbildung der Rechenschritte auf Rechentakte stellt nur ein Beispiel für eine Realisierung auf einem Prozessor dar. Die Anzahl von Rechentakten ist prinzipiell von der Performance des verwendeten Prozessors PROC abhängig.
Hinweis für die folgenden Operationen: (u, v) bzw. (iu, iv) stellen Koordinatenpaare von Pixeln dar; mitte, rzoom und azoom sind bildspezifische Parameter.
Pro Pixel sind folgende Operationen notwendig:

1.) 1 Takt: u = iu – mitte
2.) 1 Takt: v = iv – mittex1 = u·rzoom
3.) 1 Takt: x = x1 + xcy1 = v·rzoom
4.) 1 Takt: y = y1 + ycx2 = x·x
5.) 1 Takt: y2 = y·y
6.) 1 Takt: r2 = x2 + y2XAZOOM = x·azooms
7.) 13 Takte: 1/R = 1/SQRT(r2) YAZOOM = y·azoom
8.) 1 Takt: DX = XAZOOM·1/R
9.) 2 Takte: DY = YAZOOM·1/R Ablegen von DX in Matrix MDX im Speicher ADR
10.) 2 Takte: COMPARE (1/R, 1/IAU) Ablegen von DY in Matrix MDY im Speicher ADR
11.) 1 Takt: if 1/R >= 1/IAU: WIAU = 0
12.) 1 Takt: if 1/R < 1/IAU: WIAU = 1
13.) 2 Takte: Inkrement von Pixeladresse, Ablegen von WIAU in Datenfeld MWIAU im Speicher ADR

Verarbeitungsleistung insgesamt: 28 Takte pro Pixel
Anschließend ist die Berechnung der Zugriffsadresse und die Adressierung der Stützwerte für die Medianwertbildung notwendig. Für jedes zu bearbeitende Pixel müssen M1 Stützwerte k = {–(M1-1)/2, –(M1-1)/2 + 1, ..., 0, (M1-1)/2 – 1, (M1-1)/2} adressiert werden.
Ein Stützwert ist dabei das zu bearbeitende Pixel selbst (Adreßrechnungen für dieses Pixel entfallen).
Die Operationen zur Adressierung eines Stützwertes sind:

1.) Einlesen von DX über SP-Bus
2.) Einlesen von DY über SP-Bus
3.) Floating-Point-Multiplikation: du1 = k·DX
4.) Floating-Point-Multiplikation: dv1 = k·DY
5.) NINT-Operation: idu1 = nint(du1)
6.) NINT-Operation: idv1 = nint(dv1)
7.) Integer-Addition: inu1 = iu + idu1
8.) Integer-Addition: inv1 = iv + idv1
9.) Clipping-Operation: inu = clipping von inu1 auf {0, MAT-DIM-1}
10.) Clipping-Operation: inv = clipping von inv1 auf {0, MAT-DIM-1}
11.) Integer-Multiplikation: rel_row_addr = inv·MATDIM
12.) 2 Integer-Additionen: tap_adr = base_KLB + rel_row_adr + inu
13.) Einlesen des Stützstellenwertes aus Matrix KLB über D-Bus

Operationen insgesamt für pro Pixel (M1 Stützwerte):

1.) 2·M1 Zugriffe auf SP-Bus: 4·M1 Takte
2.) M1 Zugriffe auf D-Bus: M1 Takte
3.) 2·(M1-1) Floating-Point-Multiplikationen: 2·(M1-1) Takte
4.) 2·(M1-1) NINT-Operationen: 2·(M1-1) Takte
5.) 2·(M1-1) Clipping-Operationen (mit Comparatorschaltung 1 von, 2 Takte pro Clipping-Operation): 4·(M1-1) Takte
6.) (M1-1) Integer-Multiplikation: (M1-1) Takte
7.) 4·(M1-1) Integer-Additionen: 4·(M1-1) Takte

Bei geeigneter Verzahnung und Pipelining dürfte sich eine Verarbeitungsleistung von 5·M1 Takten pro Pixel erreichen lassen.
Für jedes Pixel muß der Medianwert aus M1 Werten bestimmt werden. Dazu sind (M1-1) + (M1-2)+ ... + 2 = 0.5·M1·(M1-1) – 1 Vergleiche notwendig, die über eine Comparatorschaltung 2 gebildet werden. In 5 ist die Comparatorschaltung 2 für den Fall M1 = 5 dargestellt. Die Vorgehensweise beim Vergleich von 2 Werten A und B ist:
Wenn der Wert A größer ist als der Wert B, dann sollen die Werte A und B getauscht werden. Das Vertauschen ist in der Comparatorschaltung 2 durch 2 Multiplexer realisiert, die abhängig vom Vergleichsergebnis die Werte tauschen oder nicht. In 5 ist ein Vergleich zwischen den Werten von Reg. Wert 1 und Reg. Wert 3 dargestellt.
Der ermittelte Medianwert eines Pixels sei MED[Pixel].
Mit Hilfe der Comparatorschaltung 2 von 5 dürfte es möglich sein, innerhalb von 3 Takten einen Vergleich durchzuführen. Die Verarbeitungsleistung bei der Medianwertbildung dürfte damit 3·[0.5·M1·(M1-1) – 1] Takte betragen.
Bei der Differenzbilderzeugung hat jedes Pixel folgende Operationen zu durchlaufen:

1.) Integer-Subtraktion: IDIFF [Pixel] = KLB [Pixel] – MED [Pixel] → 1 Takt
2.) Clipping-Operation: UDB'(Pixel] = Clipping von IDIFF[Pixel] auf {–SART, SART} → 2 Takte (mit Comparatorschaltung 1 von)
3.) Schreiben von UDB'(Pixel] in Speicher DAT → 1 Takt

Bei der Differenzbilderzeugung dürfte eine Verarbeitungsleistung von 4 Takten pro Pixel möglich sein.
Jedes Pixel hat im Medianfilter-ASIC MAS die in 6 dargestellte Verarbeitungspipeline zu durchlaufen. Das langsamste Glied der Kette bestimmt die Verarbeitungsleistung der Pipeline. Für den Fall M1 = 5 hat die Generierung der Stützwertadressen durch den Prozessor PROC mit den abgeschätzten 28 Takten/Pixel ungefähr den gleichen Aufwand wie die Medianwertbildung mit 27 Takten/Pixel. Für M1 > 5 stellt die Medianwertbildung den größten Aufwand dar.
Der Rechenschritt B erfordert [0.5·M1·(M1-1) – 1]·3·MATDIM ·MATDIM Rechentakte, wenn die Medianwertbildung das langsamste Glied in der Verarbeitungspipeline ist.
Rechenschritt C: 2. Medianfilter
Das 2. Medianfilter wird mit M2 Stützstellen durchgeführt. Das Raster zwischen den Stützstellen ist gegenüber dem Raster des 1. Medianfilters halbiert. Die vom Prozessor PROC abgelegten Inkremente zur Generierung der Stützstellenadressen sind demnach zu halbieren.
Das 2. Medianfilter ist nur in einem zentrumsnahen Bereich zu durchlaufen. Dieser Bereich ist im Vergleich zur Bildmatrix sehr klein.
Die Eingangsbildmatrix ist die Differenz der Matrizen KLB[Pixel] und UDB'[Pixel]. Die Differenz wird aus Aufwandsgründen nur für die Stützstellenelemente des 2. Medianfilters berechnet.
Die Berechnungen des 2. Medianfilters können mit leichten Modifikationen mit den Recheneinheiten des 1. Medianfilters durchgeführt werden.
Selbst bei großem Zoom dürften beim 2. Medianfilter höchstens 10% des Rechenaufwandes des 1. Medianfilters anfallen. Die Verarbeitungsleistung wird beim 2. Medianfilter im Fall von M2 < 5 durch die Berechnung der Zugriffsadresse der Stützwerte und die Adressierung der Stützwerte für die Medianfilterung bestimmt. Der maximale Rechenaufwand für das 2. Medianfilter wird mit Hilfe der Darstellung des Rechenaufwandes für das 1. Medianfilter folgendermaßen abgeschätzt:
Der Rechenschritt B erfordert maximal
0.5·M2·MATDIM·MATDIM Rechentakte.
Rechenschritt D: Winkelverschleifung und Rechenschritt E: Ergebnisbilderzeugung
Zuerst sind die Koordinaten der Pixel in Bezug auf das Drehzentrum durch den Prozessor PROC zu bestimmen.
Diese Berechnungen erfolgen mit einem geringen zeitlichen Vorlauf gegenüber dem Start des Winkelverschleifungs-ASICs WAS, damit die x- und y-Koordinaten eines Pixels bei der Bearbeitung des Pixels durch das Winkelverschleifungs-ASIC bereits zur Verfügung stehen.
Die Datenfelder MDX und MDY der Medianfilter werden nicht mehr benötigt und können überschrieben werden. Der Prozessor PROC legt die x-Koordinate im Datenfeld MDX der Medianfilter an, die y-Koordinate im Datenfeld MDY.
Pro Pixel sind folgende Operationen notwendig:

1.) 1 Takt: u = iu – mitte
2.) 1 Takt: v = iv – mitte x1 = u·rzoom
3.) 1 Takt: x = x1 + xc y1 = v·rzoom
4.) 2 Takte: y = y1 + yc Ablegen von x in Matrix MDX im Speicher ADR
5.) 2 Takte: Inkrement von Pixeladresse Ablegen von y in Matrix MDY im Speicher ADR

Verarbeitungsleistung: 7 Takte pro Pixel
Die vom Prozessor PROC berechneten Daten MDX, MDY und MWIAU werden vom Winkelverschleifungs-ASIC WAS über den Bus SP vom Speicher ADR eingelesen.
Die Pipeline zur Generierung und Akkumulation von Stützstellen (im folgenden als Akkumulationspipeline WAKKUS bezeichnet) be nötigt folgende Rechenwerke, um einen Stützwert innerhalb von 2 Takten aufakkumulieren zu können:
4 Gleitkomma-Multiplizierer: GMUL1, GMUL2, GMUL3, GMUL4
3 Gleitkomma-Addierer: GADD1, GADD2, GADD3
2 Integer-Comparatoren (16 Bit): COMP1, COMP2
1 NINT-Operator: NINT
2 Integer-Addierer (32 Bit): IADD1, IADD2
1 Integer-Multiplizierer (16 Bit Eingang, 32 Bit Ausgang): IMUL
Die Pipeline WAKKUS wird bei der Generierung und Aufakkumulation jedes Stützwertes durchlaufen. Bei der Winkelverschleifung muß für jedes Pixel eine Anzahl von W Stützstellen auf akkumuliert werden.
In der Pipeline WAKKUS wird in jeweils 2 Takten ein neuer Stützwert generiert und akkumuliert. Es ergibt sich somit eine Verarbeitungsleistung von 2 Takten pro aufzuakkumulierendem Stützwert.
Ablauf der Operationen in der Akkumulationspipeline WAKKUS zur Adressierung und Aufakkumulation eines Stützwertes (jeder Unterpunkt stellt eine Verzahnung um 1 Takt bzw. einer Pipelinestufe dar)

1.) GMUL1: y1 = x·a21 GMUL2: y2 = y·a22 GMUL3: x1 = x·a11 GMUL4: x2 = y·a12 Der Zustand vom Flag WIAU gibt an, ob ein zu bearbeitendes Pixel im Innen- oder Außenbereich liegt. Anhand von WIAU wird unterschieden, ob für a11, a12, a21 und a22 die Koeffizienten für den Innen- oder Außenbereich benutzt werden.
2.) GADD1: yn = y1 + y2 GADD2: xn = x1 + x2 x: = xn y: = yn → Einspeisung in Schritt 1 für Generierung des nächsten Stützwertes
3.) GADD1: y3 = yn – yc GADD2: x3 = xn – xc
4.) GMUL1: vn = y3·1/rzoom (1/rzoom ist eine bildspezifische Konstante) GMUL2: un = x3·1/rzoom
5.) GADD3: finv = vn + mitte (mitte ist eine Konstante)
6.) GADD3: finu = un + mitte NINT : inv = nint (finv)
7.) COMP1: inv < 0? NINT: inu = nint(finu)
8.) COMP1: inv > MATDIM – 1? COMP2: inu < 0?
9.) IMUL: izei = inv·matdim COMP2: inu > MATDIM-1?
10.) IADD1: sadr = izei + inu
11.) Lese-Zugriff auf Datenfeld UDB mit relativer Adresse sadr ergibt Stützwert sval
12.) IADD2: Aufakkumulation der Stützwerte ak = ak + sval

Es ergibt sich folgende zeitliche Abfolge der Rechenschritte für aufeinanderfolgende Pixel:
Wie man sieht, laufen im eingeschwungenen Zustand abwechselnd die geradzahligen Rechenoperationen und die ungeradzahligen Rechenoperationen parallel ab. Diese parallele Abarbeitung ist mit den Rechenwerken der Akkumulationspipeline WAKKUs möglich. Es ergibt sich eine Verarbeitungsleistung von 2·W Takten pro Pixel.
Der aus W Stützstellen aufakkumulierte Wert ak wird für die Ergebnisbilderzeugung als Wert akend abgespeichert und der Wert ak für die Bearbeitung eines neuen Pixels rückgesetzt.
Die Erzeugung des Ergebnispixels aus dem Wert akend ergibt sich wie folgt:

1.) Lesen des zu bearbeitenden Pixels von der Matrix INB[Pixel] über den D-Bus
2.) Normierung des Wertes akend: diff = akend·1/W durch Wandlung von akend in Float durch Float-Operator FLOAT GMUL3: Multiplikation mit der Konstante 1/W NINT: Wandlung des Ergebnisses der Multiplikation in Integer
3.) Differenzbildung OUTB[Pixel] = INB[Pixel] – diff
4.) Begrenzung des Wertbereiches von OUTB[Pixel] auf den Bereich {0,4095} durch einen Integer-Comparator mit nachgeschalteter Multiplexlogik (siehe 4)
5.) Schreiben des Ergebnispixels OUTB[Pixel] in Ergebnismatrix über den D-Bus

Die Operationen zur Erzeugung des Ergebnispixels können auf jeden Fall hinter den benötigen 2·W Takten zur Stützstellenakkumulation versteckt werden: Während für das Pixel n die Ergebnisbilderzeugung vorgenommen wird, kann für das Pixel n + 1 schon mit der Stützstellenakkumulation begonnen werden. Das Einlesen von MDX, MDY und MWIAU über den Bus SP vom Speicher ADR kann ebenfalls hinter der Stützstellenakkumulation versteckt werden.
Damit dürfte sich für die Rechenschritte D (Winkelverschleifung) und E (Ergebnisbilderzeugung) eine Verarbeitungsleistung von 2·W Takten pro Pixel ergeben.
Die Rechenschritte D und E erfordern somit 2·W·MATDIM MATDIM Rechentakte.
Abschätzungen zur Verarbeitungsleistung des Gesamtsystems:
Bei dem in 2 angegebenen Minimalausbau mit 1 Prozessor PROC, 1 Medianfilter-ASIC MAS und 1 Winkelverschleifungs-ASIC WAS ergibt sich folgende Verarbeitungsleistung in Rechentakten:
Rechenschritt A: 2·MATDIM·MATDIM
Rechenschritt B: [0.5·M1·(M1-1) – 1]·3·MATDIM·MATDIM
Rechenschritt C: max. 0.5·M2·MATDIM·MATDIM
Rechenschritt D und E: 2·W·MATDIM·MATDIM
Insgesamt für Rechenschritte A bis E:
{2·(W + 1) + 3·[0.5·M1·(M1-1) – 1] + 0.5·M2}·MATDIM·MATDIM
Hinzu kommt die Zeit für die Datenversorgung von Eingangs- und Ergebnisbild.
Sofern durch 1 Medianfilter-ASIC MAS und 1 Winkelverschleifungs-ASIC WAS die Busse D und SP nicht voll ausgelastet sind, können mehrere Medianfilter-ASICs MAS1, MAS2, ..., MASm bzw. mehrere Winkelverschleifungs-ASICs WAS1, WAS2, ..., WASn parallel eingesetzt werden, um die Verarbeitung zu beschleunigen. Da die Verarbeitungsleistung des Medianfilter-ASICs an die Verarbeitungsleistung des Prozessors PROC angepaßt ist, sind genauso viele Prozessoren PROC1, PROC2, ..., PROCm wie Medianfilter-ASICs zu verwenden. Die maximale Verarbeitungsleistung wird dann durch die Zugriffe auf den Bussen D und SP bestimmt:
1.) Zugriffe pro Pixel beim Rechenschritt A:

a) 1 Zugriff auf dem D-Bus zum Einlesen des Pixels aus der Matrix INB (1 Takt)
b) 1 Zugriff auf dem D-Bus zum Schreiben des Pixels in die Matrix KLB (1 Takt)

Anzahl der Zugriffstakte beim Rechenschritt A auf dem Bus D: 2·MATDIM·MATDIM.
Anzahl der Zugriffstakte beim Rechenschritt A auf dem Bus SP: keiner
2.) Zugriffe pro Pixel beim Rechenschritt B:

a) M1 Zugriffe auf dem D-Bus zum Einlesen der Stützstellen aus der Matrix KLB (insgesamt M1 Takte)
b) 1 Zugriff auf dem D-Bus zum Schreiben des Pixels in die Matrix UDB (1 Takt)
c) 3 Zugriffe auf dem SP-Bus zum Ablegen von MDX[Pixel], MDY(Pixel] und MWIAU[Pixel] von den Prozessoren (insgesamt 6 Takte)
d) 2 Zugriffe auf dem SP-Bus zum Einlesen von MDX[Pixel] und MDY[Pixel] vom Medianfilter-ASIC (insgesamt 4 Takte.

Anzahl der Zugriffstakte beim Rechenschritt B auf dem Bus D: (M1 + 1)·MATDIM·MATDIM
Anzahl der Zugriffstakte beim Rechenschritt B auf dem Bus SP: 10·MATDIM·MATDIM
Anmerkung: 2·MATDIM·MATDIM Zugriffstakte können dabei schon während dem Rechenschritt A getätigt werden, wenn die Berechnung von MDX, MDY und MWIAU schon zu Beginn des Rechenschrittes A begonnen wird. Somit verbleiben nur noch 8·MATDIM·MATDIM Zugriffstakte auf dem Bus SP nach dem Start des Medianfilter-ASICs.
3.) Zugriffe pro zu betrachtendem Pixel beim Rechenschritt C:

a) 2·M2 Zugriffe auf dem D-Bus zum Einlesen der Stützstellen aus den Matrix KLB und UDB' (insgesamt 2·M2Takte)
b) 1 Zugriff auf dem D-Bus zum Schreiben des Pixels in die Matrix UDB'' (1 Takt)
c) 2 Zugriffe auf dem D-Bus zum Rückkopieren der Matrix UDB'' in die Matrix UDB (insgesamt 2 Takte)
d) 2 Zugriffe auf dem SP-Bus zum Einlesen von MDX[Pixel] und MDY[Pixel] vom Medianfilter-ASIC (insgesamt 4 Takte).

Anzahl der Zugriffstakte beim Rechenschritt C auf dem Bus D: 0.1·(2·M2 + 3)·MATDIM·MATDIM.
Anzahl der Zugriffstakte beim Rechenschritt C auf dem Bus SP: 0.4·MATDIM·MATDIM
4.) Zugriffe pro zu betrachtendem Pixel bei den Rechenschritten D und E:

a) W Zugriffe auf dem D-Bus zum Einlesen der Stützstellen aus der Matrix UDB (W Takte),
b) 1 Zugriff auf dem D-Bus zum Einlesen des Eingangsbildpixels aus der Matrix INB (1 Takt)
c) 1 Zugriff auf dem D-Bus zum Schreiben des Ergebnispixels in die Matrix OUTB (1 Takt)
d) 2 Zugriffe auf dem SP-Bus zum Ablegen von MDX(Pixel] und MDY[Pixel) vom Prozessor (insgesamt 4 Takte)
e) 3 Zugriffe auf dem SP-Bus zum Einlesen von MDX [Pixel], MDY[Pixel] und MWIAU[Pixel] vom Winkelverschleifungs-ASIC (insgesamt 6 Takte)

Anzahl der Zugriffstakte bei den Rechenschritten D und E auf dem Bus D: (W + 2)·MATDIM·MATDIM.
Anzahl der Zugriffstakte bei den Rechenschritten D und E auf dem Bus SP: 10·MATDIM·MATDIM
Die größtmögliche Verarbeitungsleistung ergibt sich dadurch durch die Zugriffstakte auf dem am stärksten belasteteten Bus (Angabe in Zugriffstakten):
Rechenschritt A: 2·MATDIM·MATDIM
Rechenschritt B: 8·MATDIM·MATDIM (für M1 < = 7)
Rechenschritt C: max. 0.1·(2·M2 + 3)·MATDIM·MATDIM
Rechenschritt D und E: (W + 2)·MATDIM·MATDIM (für W > = 8)
Insgesamt für Rechenschritte A bis E: [W + 12 + 0.1·(2·M2 + 3)]·MATDIM·MATDIM

Claims

Computertomograph mit einem Rechner (6) zur Rekonstruktion eines Schnittbildes, bei dem ein spezielles Rechenwerk in Form von a) einem oder mehreren Prozessoren PROC, b) einem oder mehreren Medianfilter-ASICs MAS, c) einem oder mehreren Winkelverschleifungs-ASICs WAS und d) zwei Speichern DAT und ADR zur Realisierung eines adaptiven Bildfilters zur Artefaktreduktion vorgesehen sind (2 und 3), wobei die Realisierung des Medianfilters des adaptiven Bildfilters zur Artefaktreduktion mit nachfolgender Pipeline mit pixelweisen Durchlauf vorgenommen wird: – Bestimmung der Inkremente zur Generierung der Stützstellenadressen durch einen oder mehr Prozessoren PROC – Berechnung der Zugriffsadressen und Adressierung der Stützwerte vom Speicher DAT durch das Medianfilter-ASIC MAS – Medianwertbildung durch das Medianfilter-ASIC MAS – Differenzbilderzeugung durch das Medianfilter-ASIC MAS, und wobei die Realisierung der Winkelverschleifung des adaptiven Bildfilters zur Artefaktreduktion mit nachfolgender Pipeline mit pixelweisen Durchlauf vorgenommen wird: – Bestimmung der Koordinaten des Pixels in Bezug auf das Drehzentrum durch einen Prozessor PROC – Akkumulationspipeline WAKKUS im Winkelverschleifungs-ASIC WAS zur Generierung und Akkumulation von Stützwerten aus dem von der Medianfilterung resultierenden Differenzbild – Erzeugung des Ergebnisbildes vom Winkelverschleifungs-ASIC WAS.