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

150 Commits

Author SHA1 Message Date
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
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
f0b49995e5 fix offset compuations 2014-01-07 19:13:10 +01:00
Evghenii
21313e52b4 added local ptr correction of store instruction. change compilation to llvm ptx for tests 2014-01-07 18:54:23 +01:00
Evghenii
1303b07b72 added correction for local pointer computations 2014-01-07 18:45:22 +01:00
evghenii
7d37f7b634 added separate function that deal with local pointers 2014-01-07 18:29:44 +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
3972d740a6 added mask for tasking function 2014-01-06 16:18:28 +01:00
Evghenii
7fbe2eba59 fixed sync 2014-01-06 15:51:14 +01:00
Evghenii
91d4ae46f6 sort --fails 2014-01-06 15:38:30 +01:00
Evghenii
546f9cb409 MAJOR CHANGE--- STOP WITH THIS BRANCH-- 2014-01-06 13:51:02 +01:00
Evghenii
ddfe782151 merged 2013-12-13 11:56:43 +01:00
Dmitry Babokin
2d2d14744b Fixing --opt=force-aligned-memory for LLVM 3.3+ 2013-12-04 19:00:02 +04:00
evghenii
bb46b561fd Merged with upstream/master 2013-11-22 08:13:16 +01:00
Ilia Filippov
3fd9d5a025 support of LLVM 3.5 2013-11-21 19:09:43 +04:00
Evghenii
e162d5a99d programIndex still not working, found where change is needed... 2013-11-14 19:46:08 +01:00
Evghenii
918ca339b6 now programIndex returns laneIdx = %tid.x & (%warpsize-1) & programCount returns 32 2013-11-14 19:27:52 +01:00
Evghenii
268be7f0b5 fixed ISPCSync functionality 2013-11-13 11:19:10 +01:00
Evghenii
55bf0d23c2 resotred non-ptx functionality 2013-11-13 11:08:58 +01:00
Evghenii
f433aa3ad5 CDP works now 2013-11-13 10:43:52 +01:00
Evghenii
cf679187b1 added CDP calls into IR, next step ... check :) 2013-11-12 16:39:22 +01:00
Evghenii
fd17ad236a export functions are now also generated... next add proper CDP calls.. 2013-11-12 14:05:12 +01:00
Evghenii
dbb96c1885 need to fix launch code 2013-11-12 13:41:03 +01:00
Evghenii
4cd7e10ad3 reversed to original changes. Here is the plan to use CDP and genarate only device code with host wrapper.. 2013-11-12 12:51:56 +01:00
Evghenii
2cef101022 now emits host object file with ptx in it... next step .. testing 2013-10-31 18:05:04 +01:00
Evghenii
0a069f7de2 added comment 2013-10-31 16:06:44 +01:00
Evghenii
a2fd124997 forced module name & ptx string to be generaetd nly once 2013-10-31 16:04:30 +01:00
Evghenii
63917f8cc2 now generates CUDALaunch call. Few tweaks are still necesary 2013-10-31 16:01:34 +01:00
Evghenii
e7ddb9e642 now adds function&module name. next step adding pointer to parameter list 2013-10-30 22:41:01 +01:00
Evghenii
f9ec1a0097 .. work in programm to embed PTX into host code .. 2013-10-30 16:47:30 +01:00
Evghenii
8baef6daa3 +1 2013-10-29 14:01:53 +01:00
Evghenii
ac700d4860 checkpoint 2013-10-29 13:36:31 +01:00
egaburov
f89bad1e94 launch now passes the right info into tasking 2013-10-23 12:51:06 +02:00
Matt Pharr
1276ea9844 Revert "Remove support for building with LLVM 3.1" 
This reverts commit d3c567503b.

Conflicts:
	opt.cpp
 2013-08-06 17:00:35 -07:00
Matt Pharr
d3c567503b Remove support for building with LLVM 3.1 2013-07-31 06:46:45 -07:00
Matt Pharr
e7abf3f2ea Add support for mask vectors of 8 and 16-bit element types. 
There were a number of places throughout the system that assumed that the
execution mask would only have either 32-bit or 1-bit elements.  This
commit makes it possible to have a target with an 8- or 16-bit mask.
 2013-07-23 16:50:11 -07:00
Matt Pharr
946c39a5df Fix build with LLVM top-of-tree. 
The DIBuilder::getCU() method has been removed; we now just store the
compilation unit returned when we call DIBuilder::createCompileUnit.
 2013-07-22 15:42:52 -07:00
Dmitry Babokin
c85439e7bb Fix for the bug introduced by --intrumentation fix 2013-07-04 21:41:57 +04:00
Ilia Filippov
9fb981e9a0 correction of --instrument option support 2013-06-25 12:33:23 +04:00
Ilia Filippov
560acd5017 changes to support createFunction() with DICompositeType argument in LLVM_3_4 2013-06-04 15:48:39 +04: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
james.brodman
60c5bef90f Simplify ~mask codegen to emit single XOR like other places in the code. 2013-05-15 16:57:41 -04:00
Dmitry Babokin
b6b9daa3c5 Enabling llvm 3.4 2013-05-13 19:25:31 +04:00
james.brodman
658dd3486b Add check for enum type in Assert. 2013-04-30 16:10:57 -04:00
Dmitry Babokin
d36ab4cc3c Adding noalias attribute to malloc return 2013-04-25 20:39:01 +04:00
Dmitry Babokin
eb2e5f378c Comment fixes 2013-04-18 15:36:35 +04: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