Page MenuHomePhabricator

Compute (SPIR-V 1.0)Milestone
ArchivedPublic

Watchers

  • This project does not have any watchers.
  • View All

Details

Description

Convert SPIR-V 1.0 to NV50 IR

Recent Activity

Aug 23 2020

pmoreau archived P105 SPIR-V -> NV50 IR: Texture R/W.
Aug 23 2020, 10:23 AM · Compute (SPIR-V 1.0)
pmoreau archived P111 Old patch for clCreateProgramWithSource() support.
Aug 23 2020, 10:23 AM · Compute (SPIR-V 1.0)
pmoreau archived Compute (SPIR-V 1.0).
Aug 23 2020, 10:21 AM
pmoreau closed T76: Fix FunctionCall as Invalid.

Translating from SPIR-V to NV50 IR has been superseded by translating from SPIR-V to NIR and from NIR to NV50 IR.

Aug 23 2020, 10:21 AM · Compute (SPIR-V 1.0)
pmoreau closed T86: Add an out-of-SSA pass as Invalid.

Translating from SPIR-V to NV50 IR has been superseded by translating from SPIR-V to NIR and from NIR to NV50 IR.

Aug 23 2020, 10:20 AM · Compute (SPIR-V 1.0)
pmoreau closed T80: Implement decorations as Invalid.

Translating from SPIR-V to NV50 IR has been superseded by translating from SPIR-V to NIR and from NIR to NV50 IR.

Aug 23 2020, 10:20 AM · Compute (SPIR-V 1.0)

Dec 2 2017

pmoreau edited the content of Current Status of OpenCL Through SPIR-V.
Dec 2 2017, 12:21 PM · Compute (SPIR-V 1.0)

Sep 20 2017

pmoreau changed the edit policy for P111 Old patch for clCreateProgramWithSource() support.
Sep 20 2017, 1:26 AM · Compute (SPIR-V 1.0)

Jun 25 2017

pmoreau edited the content of Current Status of OpenCL Through SPIR-V.
Jun 25 2017, 10:44 PM · Compute (SPIR-V 1.0)

May 26 2017

pmoreau edited the content of Current Status of OpenCL Through SPIR-V.
May 26 2017, 12:47 AM · Compute (SPIR-V 1.0)

Feb 18 2017

pmoreau edited the content of Current Status of OpenCL Through SPIR-V.
Feb 18 2017, 6:41 PM · Compute (SPIR-V 1.0)

Oct 29 2016

pmoreau edited the content of Current Status of OpenCL Through SPIR-V.
Oct 29 2016, 11:34 PM · Compute (SPIR-V 1.0)
pmoreau edited the content of Current Status of OpenCL Through SPIR-V.
Oct 29 2016, 11:33 PM · Compute (SPIR-V 1.0)

Sep 22 2016

pmoreau added a comment to P105 SPIR-V -> NV50 IR: Texture R/W.

int main(int argc, char ** argv)
{

cl_int error;
Sep 22 2016, 1:40 PM · Compute (SPIR-V 1.0)
pmoreau added a comment to P105 SPIR-V -> NV50 IR: Texture R/W.
__kernel void rw_image(__read_only image2d_t in, __write_only image2d_t out)
{
  int2 coords = (int2) (get_global_id(0), get_global_id(1));
  if (coords.x > 4 || coords.y > 4)
    return;
Sep 22 2016, 1:39 PM · Compute (SPIR-V 1.0)
pmoreau changed the edit policy for P105 SPIR-V -> NV50 IR: Texture R/W.
Sep 22 2016, 1:19 PM · Compute (SPIR-V 1.0)

Aug 15 2016

pmoreau closed T84: Use operation's signedness if available, rather than operands' signedness as Resolved.
Aug 15 2016, 12:26 AM · Compute (SPIR-V 1.0)
pmoreau closed T88: Implement `clCreateProgramWithIL()` as Resolved.
Aug 15 2016, 12:11 AM · Compute (SPIR-V 1.0)

Apr 23 2016

pmoreau added a comment to T88: Implement `clCreateProgramWithIL()`.

Done, need to clean up the code.

Apr 23 2016, 2:03 AM · Compute (SPIR-V 1.0)

Apr 22 2016

pmoreau closed T59: Fix Fermi+ support as Resolved.
Apr 22 2016, 7:21 PM · Compute (SPIR-V 1.0)
pmoreau closed T78: Undertand `NVC0LoweringPass::visit()` treatment of STORE as Resolved.

Fixed by Hans

Apr 22 2016, 7:21 PM · Compute (SPIR-V 1.0)
pmoreau closed T78: Undertand `NVC0LoweringPass::visit()` treatment of STORE, a subtask of T59: Fix Fermi+ support, as Resolved.
Apr 22 2016, 7:21 PM · Compute (SPIR-V 1.0)

Apr 13 2016

pmoreau archived P95 [GK106] Wrong CVT from U64 to U32, as src became U32.
Apr 13 2016, 11:46 PM · Compute (SPIR-V 1.0)
pmoreau archived P94 [GK106] Unknown opcode in CVT from U32 to U64.
Apr 13 2016, 11:46 PM · Compute (SPIR-V 1.0)
pmoreau archived P93 [GF119] hello_world.
Apr 13 2016, 11:46 PM · Compute (SPIR-V 1.0)

Apr 10 2016

pmoreau closed T89: `clGetDeviceInfo()` should return supported ILs as Resolved.

Hacked a quick patch which should do for now.

Apr 10 2016, 3:43 PM · Compute (SPIR-V 1.0)
pmoreau renamed T89: `clGetDeviceInfo()` should return supported ILs from `clGetDeviceInfoquery()` should return supported ILs to `clGetDeviceInfo()` should return supported ILs.
Apr 10 2016, 3:42 PM · Compute (SPIR-V 1.0)

Apr 9 2016

pmoreau updated the task description for T88: Implement `clCreateProgramWithIL()`.
Apr 9 2016, 12:25 PM · Compute (SPIR-V 1.0)
pmoreau changed the edit policy for T89: `clGetDeviceInfo()` should return supported ILs.
Apr 9 2016, 12:25 PM · Compute (SPIR-V 1.0)
pmoreau changed the edit policy for T88: Implement `clCreateProgramWithIL()`.
Apr 9 2016, 12:20 PM · Compute (SPIR-V 1.0)
pmoreau moved T86: Add an out-of-SSA pass from Doing to Initial Serie on the Compute (SPIR-V 1.0) board.
Apr 9 2016, 12:15 PM · Compute (SPIR-V 1.0)
pmoreau moved T59: Fix Fermi+ support from Backlog to Initial Serie on the Compute (SPIR-V 1.0) board.
Apr 9 2016, 12:14 PM · Compute (SPIR-V 1.0)
pmoreau moved T78: Undertand `NVC0LoweringPass::visit()` treatment of STORE from Backlog to Initial Serie on the Compute (SPIR-V 1.0) board.
Apr 9 2016, 12:14 PM · Compute (SPIR-V 1.0)
pmoreau moved T84: Use operation's signedness if available, rather than operands' signedness from Backlog to Initial Serie on the Compute (SPIR-V 1.0) board.
Apr 9 2016, 12:14 PM · Compute (SPIR-V 1.0)

Mar 6 2016

pmoreau changed the edit policy for P95 [GK106] Wrong CVT from U64 to U32, as src became U32.
Mar 6 2016, 3:35 AM · Compute (SPIR-V 1.0)
pmoreau changed the edit policy for P94 [GK106] Unknown opcode in CVT from U32 to U64.
Mar 6 2016, 1:54 AM · Compute (SPIR-V 1.0)
pmoreau closed T79: Kepler: NTID -> c0[0x1000] seems weird as Resolved.

Nope, it was filled in. :-)

Mar 6 2016, 1:27 AM · Compute (SPIR-V 1.0)
pmoreau closed T79: Kepler: NTID -> c0[0x1000] seems weird, a subtask of T59: Fix Fermi+ support, as Resolved.
Mar 6 2016, 1:26 AM · Compute (SPIR-V 1.0)

Mar 4 2016

pmoreau added a comment to T78: Undertand `NVC0LoweringPass::visit()` treatment of STORE.

imirkin pmoreau: yeah, i wouldn't expect the blob to do anything like that -- this is just the way that i expressed buffer objects... you can decree that fileIndex == -1 means that you shouldn't do any of that. or some other indication of the base address needing to be loaded... perhaps that should be done right at the from_tgsi level for buffers
[…]
imirkin i know that gmem on nv50 can have a legit fileindex though
imirkin so that might not be the best thing to do =/
imirkin but this needs to be communicated *somehow*
[…]
imirkin perhaps i should have created a new nv50/ir file
imirkin and then lowered it away

Mar 4 2016, 4:45 PM · Compute (SPIR-V 1.0)
pmoreau edited P93 [GF119] hello_world.
Mar 4 2016, 1:45 PM · Compute (SPIR-V 1.0)
pmoreau changed the edit policy for P93 [GF119] hello_world.
Mar 4 2016, 1:43 PM · Compute (SPIR-V 1.0)

Mar 3 2016

pmoreau claimed T78: Undertand `NVC0LoweringPass::visit()` treatment of STORE.
Mar 3 2016, 11:50 PM · Compute (SPIR-V 1.0)
pmoreau claimed T79: Kepler: NTID -> c0[0x1000] seems weird.
Mar 3 2016, 11:50 PM · Compute (SPIR-V 1.0)
pmoreau added a comment to T79: Kepler: NTID -> c0[0x1000] seems weird.

Comes from here. Basically NTID, NCTAID and GRIDID should be uploaded by the driver in c0[], which is not done currently (apparently).

Mar 3 2016, 11:50 PM · Compute (SPIR-V 1.0)
pmoreau renamed T78: Undertand `NVC0LoweringPass::visit()` treatment of STORE from Fix hello_world on Fermi to Undertand `NVC0LoweringPass::visit()` treatment of STORE.
Mar 3 2016, 11:32 PM · Compute (SPIR-V 1.0)

Feb 22 2016

pmoreau lowered the priority of T59: Fix Fermi+ support from High to Normal.
Feb 22 2016, 10:27 PM · Compute (SPIR-V 1.0)

Feb 21 2016

pmoreau moved T86: Add an out-of-SSA pass from Backlog to Doing on the Compute (SPIR-V 1.0) board.
Feb 21 2016, 8:13 PM · Compute (SPIR-V 1.0)
pmoreau changed the edit policy for T86: Add an out-of-SSA pass.
Feb 21 2016, 8:13 PM · Compute (SPIR-V 1.0)

Feb 19 2016

pmoreau closed T81: Clean load/store handling as Resolved.
Feb 19 2016, 11:32 PM · Compute (SPIR-V 1.0)
pmoreau closed T81: Clean load/store handling, a subtask of T61: `getIndirect()` returns `nullptr` in `CodeEmitterNV50::emitLoad()`, as Resolved.
Feb 19 2016, 11:32 PM · Compute (SPIR-V 1.0)