aaron/ispc
1
0
Fork 0
Code Issues Pull Requests Releases Wiki Activity

handwired CDP launch

Browse Source
This commit is contained in:
Evghenii
2013-11-12 21:20:10 +01:00
parent 76bfcc29c2
commit f587e0a459
2 changed files with 7 additions and 2 deletions
Download Patch File Download Diff File Expand all files Collapse all files

7
builtins/target-nvptx64.ll
View File

@@ -115,7 +115,8 @@ define i8* @ISPCAlloc(i8**, i64, i32) nounwind alwaysinline
{
ret i8* null
}
define void @ISPCLaunch(i8**, i8* %func_ptr, i8** %func_args, i32 %ntx, i32 %nty, i32 %ntz) nounwind alwaysinline
declare i64 @cudaGetParameterBuffer(i64, i64) nounwind
define void @ISPCLaunch(i8**, i8* %func_ptr, i8** %func_args, i32 %nargs, i32 %ntx, i32 %nty, i32 %ntz) nounwind alwaysinline
{
entry:
%func_i64 = ptrtoint i8* %func_ptr to i64
@@ -136,6 +137,10 @@ entry:
if.then: ; preds = %entry
%param = call i64 @cudaGetParameterBuffer(i64 8, i64 24);
%ptr = inttoptr i64 %param to i8*;
%res_tmp = call i32 asm sideeffect "{
.reg .s32 %r<8>;
.reg .s64 %rd<3>;

2
examples_cuda/aobench/ao.cu
View File

@@ -418,7 +418,7 @@ void ao_ispc_tasks(
const int nbz = 1;
const dim3 blocks (nbx, nby, nbz);
if (threadIdx.x == 0)
ao_task<<<blocks, 128>>>(w,h,nsubsamples,image);
ao_task<<<blocks, 128>>>(w,w,nsubsamples,image);
cudaDeviceSynchronize();
}
#endif