-#include "stdlib.h"
__kernel void
test_copy_buffer(__global float* src, __global float* dst)
{
+++ /dev/null
-; ModuleID = 'test_copy_buffer.cl.o'
-target datalayout = "e-p:32:32-i64:64:64-f64:64:64-n1:8:16:32:64"
-target triple = "ptx32--"
-
-define ptx_kernel void @test_copy_buffer(float addrspace(1)* nocapture %src, float addrspace(1)* nocapture %dst) nounwind noinline {
-get_global_id.exit:
- %call.i.i = tail call ptx_device i32 @__gen_ocl_get_local_id0() nounwind readnone
- %call.i3.i = tail call ptx_device i32 @__gen_ocl_get_local_size0() nounwind readnone
- %call.i10.i = tail call ptx_device i32 @__gen_ocl_get_group_id0() nounwind readnone
- %mul.i = mul i32 %call.i10.i, %call.i3.i
- %add.i = add i32 %mul.i, %call.i.i
- %arrayidx = getelementptr inbounds float addrspace(1)* %src, i32 %add.i
- %0 = load float addrspace(1)* %arrayidx, align 4, !tbaa !1
- %arrayidx1 = getelementptr inbounds float addrspace(1)* %dst, i32 %add.i
- store float %0, float addrspace(1)* %arrayidx1, align 4, !tbaa !1
- ret void
-}
-
-declare ptx_device i32 @__gen_ocl_get_group_id0() nounwind readnone
-
-declare ptx_device i32 @__gen_ocl_get_local_size0() nounwind readnone
-
-declare ptx_device i32 @__gen_ocl_get_local_id0() nounwind readnone
-
-!opencl.kernels = !{!0}
-
-!0 = metadata !{void (float addrspace(1)*, float addrspace(1)*)* @test_copy_buffer}
-!1 = metadata !{metadata !"float", metadata !2}
-!2 = metadata !{metadata !"omnipotent char", metadata !3}
-!3 = metadata !{metadata !"Simple C/C++ TBAA", null}
--- /dev/null
+__kernel void
+test_copy_buffer_row(__global int *src, __global int *dst, __global int *data)
+{
+ int row = data[0];
+ int size = data[1];
+ int id = (int) get_global_id(0);
+ for (; id < size; id += row) dst[id] = src[id];
+}
+
+++ /dev/null
-#include "stdlib.h"
-__kernel void
-test_write_only(__global float* dst )
-{
- int id = (int)get_global_id(0);
- dst[id] = 1;
-}
-
+++ /dev/null
-; ModuleID = 'test_write_only_2.cl.o'
-target datalayout = "e-p:32:32-i64:64:64-f64:64:64-n1:8:16:32:64"
-target triple = "ptx32--"
-
-define ptx_kernel void @test_write_only(float addrspace(1)* nocapture %dst) nounwind noinline {
-get_global_id.exit:
- %call.i.i = tail call ptx_device i32 @__gen_ocl_get_local_id0() nounwind readnone
- %call.i3.i = tail call ptx_device i32 @__gen_ocl_get_local_size0() nounwind readnone
- %call.i10.i = tail call ptx_device i32 @__gen_ocl_get_group_id0() nounwind readnone
- %mul.i = mul i32 %call.i10.i, %call.i3.i
- %add.i = add i32 %mul.i, %call.i.i
- %arrayidx = getelementptr inbounds float addrspace(1)* %dst, i32 %add.i
- store float 1.000000e+00, float addrspace(1)* %arrayidx, align 4, !tbaa !1
- ret void
-}
-
-declare ptx_device i32 @__gen_ocl_get_group_id0() nounwind readnone
-
-declare ptx_device i32 @__gen_ocl_get_local_size0() nounwind readnone
-
-declare ptx_device i32 @__gen_ocl_get_local_id0() nounwind readnone
-
-!opencl.kernels = !{!0}
-
-!0 = metadata !{void (float addrspace(1)*)* @test_write_only}
-!1 = metadata !{metadata !"float", metadata !2}
-!2 = metadata !{metadata !"omnipotent char", metadata !3}
-!3 = metadata !{metadata !"Simple C/C++ TBAA", null}
ADD_EXECUTABLE(test_write_only tests/test_write_only.c)
ADD_EXECUTABLE(test_copy_buffer tests/test_copy_buffer.c)
+ADD_EXECUTABLE(test_copy_buffer_row tests/test_copy_buffer_row.c)
ADD_EXECUTABLE(test_eot tests/test_eot.c)
TARGET_LINK_LIBRARIES(test_eot cl_test m)
TARGET_LINK_LIBRARIES(test_write_only cl_test m)
TARGET_LINK_LIBRARIES(test_copy_buffer cl_test m)
+TARGET_LINK_LIBRARIES(test_copy_buffer_row cl_test m)
#ADD_EXECUTABLE(test_copy_buffer tests/test_copy_buffer.c)
#ADD_EXECUTABLE(test_copy_image tests/test_copy_image.c)