Evghenii
4641a15287
Merge branch 'master' into nvptx
2014-03-19 10:53:07 +01:00
Dmitry Babokin
31b95b665b
Copyright update
2014-03-12 20:19:16 +04:00
Evghenii
739e4e73ef
make static illegal
2014-02-21 12:08:09 +01:00
Evghenii
4196c723eb
merged with nvptx
2014-02-20 11:01:58 +01:00
Evghenii
8ffa84f875
added some #ifdef .. #endif for control flow tests
2014-02-06 10:12:31 +01:00
Evghenii
686c1d676d
improvements
2014-02-05 12:04:36 +01:00
jbrodman
98cfc17843
Fix bug with printing due to uneven handling of bool types
2014-02-04 08:12:02 -08:00
Evghenii
98c82242c5
allowed static and disable memcpy/memmove/memset operations
2014-02-03 08:02:50 +01:00
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
