aaron/ispc
1
0
Fork 0
Code Issues Pull Requests Releases Wiki Activity
2,161 Commits 6 Branches 36 Tags
e6a6df105206b9e9acef40ffbbe2118565d5b6ba
Commit Graph

130 Commits

Author SHA1 Message Date
Evghenii
ddb9b2fc47 added basic printing from ptx 2014-01-24 13:44:38 +01:00
Evghenii
1a07aed6aa foreach_unique will work now on atomic data types, not pointers yet. enum is not tested. All tests/foreach-unique-*.ispc pass 2014-01-24 08:30:50 +01:00
Evghenii
1cf1dab649 fixed foreach_unique and local_atomics 2014-01-23 21:57:20 +01:00
Evghenii
da7a2c0c7f added emulation of "soa" data types via shared-memory 2014-01-23 16:17:06 +01:00
Evghenii
9389b6e3ef basic optimization path fails 2014-01-10 06:34:44 +01:00
evghenii
9053eed4b4 added basic optimization pass that promotes uniform into varying variables (not array) for nvptx target 2014-01-10 06:32:57 +01:00
Evghenii
b6b8855728 +1 2014-01-09 14:45:52 +01:00
Evghenii
2c043f67d0 +some uniform related improvements 2014-01-09 14:37:27 +01:00
Evghenii
1ed438dcdb cleaned up a bit code for treatment of non-static uniform variables. all stored in shared memory 2014-01-09 13:02:50 +01:00
Evghenii
db6f526b78 added experimental support for uniform variables, not only arrays. makes applications slower 2014-01-09 12:18:10 +01:00
Evghenii
5f859e4885 added addrspace(3,4)->addrspace(0) convertion to ctx->GetElementPtrInst. Appears to work now. 2014-01-08 19:31:28 +01:00
Evghenii
cc53fa4c14 fixed addrspace(3) pointer offset computation during conversion. Now it is done in GetElementPtr 2014-01-08 18:43:33 +01:00
Evghenii
69c5e0aae7 convert pointers in function arguments to addrspace(3). still there is poroblem with shared memory. need to figure out which one .. 2014-01-08 15:12:32 +01:00
Evghenii
d3dc5e0df1 patched examples to work with uniform for nvptx. function calls with non-generic pointers fail. need fix 2014-01-08 14:25:21 +01:00
Evghenii
de4d66c56f added addrspace(4)/constant memory for const uniform declarations 2014-01-08 13:27:24 +01:00
Evghenii
0a66f17897 experimental support for non-constant [non-static] uniform arrays mapped to addrspace(3) 2014-01-08 11:06:14 +01:00
Evghenii
9b74e60185 added conversion from addrspace(3)/__local/__shared__ to addspace(0)/generic when PtrToInt is called 2014-01-07 14:29:55 +01:00
Evghenii
96597d3716 identified where to add shared memory. other problems showed up.. 2014-01-07 11:07:54 +01:00
Evghenii
546f9cb409 MAJOR CHANGE--- STOP WITH THIS BRANCH-- 2014-01-06 13:51:02 +01:00
evghenii
298b4b8b5b +1 2014-01-02 13:16:14 +01:00
Evghenii
e192e8ea9e +change static naming in IR to make it compatible with NVVM 2013-11-22 14:43:14 +01:00
Evghenii
af75afeb7a foreach[_tiled] seems to work now 2013-11-14 16:29:40 +01:00
Evghenii
48644813d4 stmt.cpp forking on foreach 2013-11-14 11:30:22 +01:00
james.brodman
403d9e1059 Update break/continue test to use contribution of function mask. 2013-05-21 10:52:38 -04:00
james.brodman
4ea02c59d8 Disable break optimization and change return check to use full mask. 2013-05-21 10:00:22 -04:00
Dmitry Babokin
7497e86902 Adding Windows support for aligned memory allocation on Windows 2013-04-26 22:07:30 +02:00
Dmitry Babokin
4c35d9456a Additional cleanup to enable more broadcasts 2013-04-10 15:34:21 +04:00
Dmitry Babokin
5898532605 Broadcast implementation as InsertElement+Shuffle and related improvements 2013-04-10 02:18:24 +04:00
Dmitry Babokin
0f86255279 Target class redesign: data moved to private. Also empty target-feature attribute is not added anymore (generic targets). 2013-03-23 14:28:05 +04:00
Dmitry Babokin
3f8a678c5a Editorial change: fixing trailing white spaces and tabs 2013-03-18 16:17:55 +04:00
Matt Pharr
0bf1320a32 Remove support for building with LLVM 3.0 2013-01-06 12:27:53 -08:00
Matt Pharr
81dbd504aa Small fixes to eliminate compiler warnings when using clang 2013-01-06 12:10:54 -08:00
Matt Pharr
63dd7d9859 Fix build to work with LLVM top-of-tree again 2013-01-06 12:02:08 -08:00
Peng Tu
b80867d473 Move the call to RestoreContinuedLanes from bbBody to the correct place of bbCheckForMore for foreach_unique and foreach_active. 2012-10-29 17:27:11 -07:00
Matt Pharr
ddcd0a49ec Fix bugs with handling of 'continue' statements in foreach_* loops. 2012-09-05 10:16:58 -07:00
Matt Pharr
2a19cc1758 Fix cases where we were trying to type cast instead of type convert. 
Also, removed erroneous checks about the type of the test expression
in DoStmt and ForStmt.

These together were preventing conversion of pointer types to boolean
values, so things like "while (ptr)" would improperly not compile.

Issue #346.
 2012-08-03 12:47:53 -07:00
Matt Pharr
765a0d8896 Use puts() rather than printf() for printing assertion failure strings. 
This way, we don't lose '%'s in the assertion strings.

Issue #342.
 2012-08-03 11:31:38 -07:00
Matt Pharr
ae89a65dad Fix bug that caused unterminated basic blocks. 
Issue #339.
 2012-07-23 08:24:18 -07:00
Jean-Luc Duprat
aecd6e0878 All the smear(), setzero() and undef() APIs are now templated on the return type. 
Modified ISPC's internal mangling to pass these through unchanged.
Tried hard to make sure this is not going to introduce an ABI change.
 2012-07-17 17:06:36 -07:00
Matt Pharr
926b3b9ee3 Fix bugs with mask-handling for switch/do/for/while statements. 
All of these pass the current mask to FunctionEmitContext::SetBlockEntryMask()
so that when a break/continue/return is encountered, it can test to see if all
lanes have followed that path and then return; this in turn ensures that we never
run statements with an all-off execution mask.

These functions were passing the function internal mask, not the full mask, and
thus could end up executing code with the mask all off if some lanes were
disabled by an outer function.  (The new tests test this case.)
 2012-07-09 15:13:30 -07:00
Matt Pharr
2d8026625b Always check the execution mask after break/continue/return. 
When "break", "continue", or "return" is used under varying control flow,
we now always check the execution mask to see if all of the program
instances are executing it.  (Previously, this was only done with "cbreak",
"ccontinue", and "creturn", which are now deprecated.)

An important effect of this change is that it fixes a family of cases
where we could end up running with an "all off" execution mask, which isn't
supposed to happen, as it leads to all sorts of invalid behavior.

This change does cause the volume rendering example to run 9% slower, but
doesn't affect the other examples.

Issue #257.
 2012-07-06 11:09:11 -07:00
Matt Pharr
73afab464f Provide mask at block entry for switch statements. 
This fixes a crash if 'cbreak' was used in a 'switch'.  Renamed
FunctionEmitContext::SetLoopMask() to SetBlockEntryMask(), and
similarly the loopMask member variable.
 2012-07-06 11:08:05 -07:00
Matt Pharr
4186ef204d Fix build with LLVM top of tree. 2012-07-05 13:35:01 -07:00
Matt Pharr
ceb8ca680c Fix crash in codegen for assert() with malformed program. 
Issue #302.
 2012-06-26 11:54:55 -07:00
Matt Pharr
79ebcbec4b Fix crash in SwitchStmt::TypeCheck() with malformed programs. 2012-06-26 11:21:33 -07:00
Matt Pharr
54459255d4 Add unmasked { } statement. 
This reestablishes an "all on" execution mask for the gang, which can
be useful for nested parallelism..
 2012-06-22 14:30:58 -07:00
Matt Pharr
b4a078e2f6 Add foreach_active iteration statement. 
Issue #298.
 2012-06-22 10:35:43 -07:00
Matt Pharr
ed13dd066b Distinguish between 'regular' foreach and foreach_unique in FunctionEmitContext 
We need to do this since it's illegal to have nested foreach statements, but
nested foreach_unique, or foreach_unique inside foreach, etc., are all fine.
 2012-06-22 06:04:00 -07:00
Matt Pharr
3bc66136b2 Add foreach_unique iteration construct. 
Idea via Ingo Wald / IVL compiler.
 2012-06-20 10:04:24 -07:00
Matt Pharr
79e0a9f32a Fix codegen bug with foreach_tiled. 
When the outermost dimension(s) were partially active, but the innermost
dimension was all on, we'd inadvertently use an incorrect "all on"
execution mask.

Fixes issues #177 and #200.
 2012-06-08 14:56:18 -07:00