mirror of
https://github.com/KevinMidboe/linguist.git
synced 2025-10-29 17:50:22 +00:00
samples: add Terra samples
source: https://github.com/zdevito/terra/tree/master/tests
This commit is contained in:
47
samples/Terra/cudashared.t
Normal file
47
samples/Terra/cudashared.t
Normal file
@@ -0,0 +1,47 @@
|
||||
if not terralib.cudacompile then
|
||||
print("CUDA not enabled, not performing test...")
|
||||
return
|
||||
end
|
||||
|
||||
local tid = cudalib.nvvm_read_ptx_sreg_tid_x--terralib.intrinsic("llvm.nvvm.read.ptx.sreg.tid.x",{} -> int)
|
||||
|
||||
N = 1024
|
||||
|
||||
somedata = cudalib.sharedmemory(int,N)
|
||||
|
||||
terra bar(result : &int)
|
||||
var t = tid()
|
||||
somedata[t] = t
|
||||
cudalib.nvvm_barrier0()
|
||||
result[t] = somedata[N - 1 - t]
|
||||
end
|
||||
|
||||
terralib.includepath = terralib.includepath..";/usr/local/cuda/include"
|
||||
|
||||
local C = terralib.includecstring [[
|
||||
#include "cuda_runtime.h"
|
||||
#include <stdlib.h>
|
||||
#include <stdio.h>
|
||||
]]
|
||||
|
||||
local R = terralib.cudacompile({ bar = bar },true)
|
||||
|
||||
|
||||
terra doit(N : int)
|
||||
var data : &int
|
||||
C.cudaMalloc([&&opaque](&data),sizeof(int)*N)
|
||||
var launch = terralib.CUDAParams { 1,1,1, N,1,1, 0, nil }
|
||||
R.bar(&launch,data)
|
||||
var results : &int = [&int](C.malloc(sizeof(int)*N))
|
||||
C.cudaMemcpy(results,data,sizeof(int)*N,2)
|
||||
var result = 0
|
||||
for i = 0,N do
|
||||
--C.printf("result = %d\n",results[i])
|
||||
result = result + results[i]
|
||||
end
|
||||
return result
|
||||
end
|
||||
|
||||
local test = require("test")
|
||||
local expected = (N - 1)*N/2
|
||||
test.eq(doit(N),expected)
|
||||
Reference in New Issue
Block a user