Skip to content

Commit a4f2906

Browse files
cirosantillicirosantilli
committed
Flexible vector io
1 parent d9573ed commit a4f2906

File tree

10 files changed

+120
-28
lines changed

10 files changed

+120
-28
lines changed

‎opencl/clinfo.c

+2-2
Original file line numberDiff line numberDiff line change
@@ -52,12 +52,12 @@ int main(void) {
5252
* But yeah, likely the same for all dimensions,
5353
* and equal to DEVICE_MAX_WORK_GROUP_SIZE. */
5454
PRINT_SIZE_T(DEVICE_MAX_WORK_ITEM_SIZES);
55-
PRINT_CL_ULONG(DEVICE_LOCAL_MEM_SIZE);
55+
PRINT_CL_ULONG(DEVICE_LOCAL_MEM_SIZE);
5656
PRINT_CL_UINT(DEVICE_MAX_COMPUTE_UNITS);
5757

5858
/* Cleanup. */
5959
#ifdef CL_1_2
60-
clReleaseDevice(device);
60+
clReleaseDevice(device);
6161
#endif
6262
return EXIT_SUCCESS;
6363
}

‎opencl/common.h

+2-1
Original file line numberDiff line numberDiff line change
@@ -71,7 +71,8 @@ void common_create_kernel(
7171
) {
7272
if (NULL != source) {
7373
common_create_program(common, source, options, &common->program);
74-
common->kernel = clCreateKernel(common->program, "mymain", NULL);
74+
common->kernel = clCreateKernel(common->program, "kmain", NULL);
75+
assert(NULL != common->kernel);
7576
} else {
7677
common->kernel = NULL;
7778
common->program = NULL;

‎opencl/inc_vector.c

+21-21
Original file line numberDiff line numberDiff line change
@@ -21,22 +21,22 @@ int main(int argc, char **argv) {
2121
Common common;
2222
size_t i, n, io_sizeof;
2323

24-
if (argc > 1) {
25-
n = strtoul(argv[1], NULL, 10);
26-
} else {
27-
n = 2;
28-
}
29-
30-
/* Initialize data. */
31-
io_sizeof = n * sizeof(*io);
32-
io = malloc(io_sizeof);
33-
expected_output = malloc(n * sizeof(*expected_output));
34-
for (i = 0; i < n; ++i) {
35-
io[i] = i;
36-
expected_output[i] = i + 1;
37-
}
38-
39-
/* Run kernel. */
24+
if (argc > 1) {
25+
n = strtoul(argv[1], NULL, 10);
26+
} else {
27+
n = 2;
28+
}
29+
30+
/* Initialize data. */
31+
io_sizeof = n * sizeof(*io);
32+
io = malloc(io_sizeof);
33+
expected_output = malloc(n * sizeof(*expected_output));
34+
for (i = 0; i < n; ++i) {
35+
io[i] = i;
36+
expected_output[i] = i + 1;
37+
}
38+
39+
/* Run kernel. */
4040
common_init(&common, source);
4141
buffer = clCreateBuffer(common.context, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR, io_sizeof, io, NULL);
4242
clSetKernelArg(common.kernel, 0, sizeof(buffer), &buffer);
@@ -45,13 +45,13 @@ int main(int argc, char **argv) {
4545
clFinish(common.command_queue);
4646
clEnqueueReadBuffer(common.command_queue, buffer, CL_TRUE, 0, io_sizeof, io, 0, NULL, NULL);
4747

48-
/* Assertions. */
49-
common_vec_assert_eq_i(io, expected_output, n);
48+
/* Assertions. */
49+
common_vec_assert_eq_i(io, expected_output, n);
5050

51-
/* Cleanup. */
51+
/* Cleanup. */
5252
clReleaseMemObject(buffer);
5353
common_deinit(&common);
54-
free(io);
55-
free(expected_output);
54+
free(io);
55+
free(expected_output);
5656
return EXIT_SUCCESS;
5757
}

‎opencl/interactive/Makefile

+1
Original file line numberDiff line numberDiff line change
@@ -0,0 +1 @@
1+
../Makefile

‎opencl/interactive/Makefile_params

+1
Original file line numberDiff line numberDiff line change
@@ -0,0 +1 @@
1+
../Makefile_params

‎opencl/interactive/common.h

+1
Original file line numberDiff line numberDiff line change
@@ -0,0 +1 @@
1+
../common.h

‎opencl/interactive/vec_io.c

+82
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,82 @@
1+
/*
2+
Sample usage:
3+
4+
echo '1 2 3' | tr ' ' '\n' >vec_io.vec
5+
./prog vec_io.cl vec_io.vec
6+
7+
Or you can use the default kernel and stdin input:
8+
9+
echo '1 2 3' | tr ' ' '\n' | ./prog
10+
11+
Generic boilerplate that:
12+
13+
- takes a vector as input either from stdin or from a file, one per line
14+
15+
- processes it with a kernel read from a file, one vector item per work item
16+
17+
- produces as output a vector of the same size to stdout
18+
19+
This allows you to quickly play with different kernels without recompiling the C code.
20+
21+
But is unsuitable for real applications, which require querying the CL implementation
22+
for limits, specially work group and memory maximum sizes.
23+
*/
24+
25+
#include "common.h"
26+
27+
int main(int argc, char **argv) {
28+
char *cl_source_path;
29+
cl_float *io;
30+
cl_mem buffer;
31+
Common common;
32+
FILE *input_vector_file;
33+
float f;
34+
size_t i, n, nmax, io_sizeof;
35+
36+
/* Treat CLI arguments. */
37+
if (argc > 1) {
38+
cl_source_path = argv[1];
39+
} else {
40+
cl_source_path = "vec_io.cl";
41+
}
42+
if (argc > 2) {
43+
input_vector_file = fopen(argv[2], "r");
44+
} else {
45+
input_vector_file = stdin;
46+
}
47+
48+
/* Initialize data. */
49+
n = 0;
50+
nmax = n + 1;
51+
io = malloc(nmax * sizeof(*io));
52+
while(fscanf(input_vector_file, "%f", &f) != EOF) {
53+
io[n] = f;
54+
n++;
55+
if (n == nmax) {
56+
nmax *= 2;
57+
io = realloc(io, nmax * sizeof(*io));
58+
}
59+
}
60+
io_sizeof = n * sizeof(*io);
61+
62+
/* Run kernel. */
63+
common_init_file(&common, cl_source_path);
64+
buffer = clCreateBuffer(common.context, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR, io_sizeof, io, NULL);
65+
clSetKernelArg(common.kernel, 0, sizeof(buffer), &buffer);
66+
clEnqueueNDRangeKernel(common.command_queue, common.kernel, 1, NULL, &n, NULL, 0, NULL, NULL);
67+
clFlush(common.command_queue);
68+
clFinish(common.command_queue);
69+
clEnqueueReadBuffer(common.command_queue, buffer, CL_TRUE, 0, io_sizeof, io, 0, NULL, NULL);
70+
71+
/* Print result. */
72+
for (i = 0; i < n; ++i) {
73+
printf("%f\n", io[i]);
74+
}
75+
76+
/* Cleanup. */
77+
clReleaseMemObject(buffer);
78+
common_deinit(&common);
79+
free(io);
80+
fclose(input_vector_file);
81+
return EXIT_SUCCESS;
82+
}

‎opencl/interactive/vec_io.cl

+3
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,3 @@
1+
__kernel void kmain(__global float *io) {
2+
io[get_global_id(0)] += 1.0;
3+
}

‎opencl/vec_io.vec

+3
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,3 @@
1+
1
2+
2
3+
3

‎opencl/work_item_builtin.c

+4-4
Original file line numberDiff line numberDiff line change
@@ -17,7 +17,7 @@ int main(void) {
1717
cl_uint output[NUM_FUNCTIONS];
1818
Common common;
1919

20-
/* Run kernel. */
20+
/* Run kernel. */
2121
common_init_file(&common, "work_item_builtin.cl");
2222
buffer = clCreateBuffer(common.context, CL_MEM_WRITE_ONLY, sizeof(output), NULL, NULL);
2323
clSetKernelArg(common.kernel, 0, sizeof(cl_mem), &buffer);
@@ -26,8 +26,8 @@ int main(void) {
2626
clFinish(common.command_queue);
2727
clEnqueueReadBuffer(common.command_queue, buffer, CL_TRUE, 0, sizeof(output), output, 0, NULL, NULL);
2828

29-
/* Check the values. */
30-
puts("#work_item_builtin");
29+
/* Check the values. */
30+
puts("#work_item_builtin");
3131
printf("work_dim = %d\n", output[0]);
3232
printf("global_size = %d\n", output[1]);
3333
printf("global_id = %d\n", output[2]);
@@ -37,7 +37,7 @@ int main(void) {
3737
printf("group_id = %d\n", output[6]);
3838
printf("global_offset = %d\n", output[7]);
3939

40-
/* Cleanup. */
40+
/* Cleanup. */
4141
clReleaseMemObject(buffer);
4242
common_deinit(&common);
4343
return EXIT_SUCCESS;

0 commit comments

Comments
 (0)