test_kernel.cpp 11 KB

123456789101112131415161718192021222324252627282930313233343536373839404142434445464748495051525354555657585960616263646566676869707172737475767778798081828384858687888990919293949596979899100101102103104105106107108109110111112113114115116117118119120121122123124125126127128129130131132133134135136137138139140141142143144145146147148149150151152153154155156157158159160161162163164165166167168169170171172173174175176177178179180181182183184185186187188189190191192193194195196197198199200201202203204205206207208209210211212213214215216217218219220221222223224225226227228229230231232233234235236237238239240241242243244245246247248249250251252253254255256257258259260261262263264265266267268269270271272273274275276277278279280281282283284285286287288289290291292293294295296297298299300301302303304305306307308309310311312313314315316317318319320321322323324325326327328329330331332333334335
  1. //---------------------------------------------------------------------------//
  2. // Copyright (c) 2013 Kyle Lutz <kyle.r.lutz@gmail.com>
  3. //
  4. // Distributed under the Boost Software License, Version 1.0
  5. // See accompanying file LICENSE_1_0.txt or copy at
  6. // http://www.boost.org/LICENSE_1_0.txt
  7. //
  8. // See http://boostorg.github.com/compute for more information.
  9. //---------------------------------------------------------------------------//
  10. #define BOOST_TEST_MODULE TestKernel
  11. #include <boost/test/unit_test.hpp>
  12. #include <boost/compute/buffer.hpp>
  13. #include <boost/compute/kernel.hpp>
  14. #include <boost/compute/types.hpp>
  15. #include <boost/compute/system.hpp>
  16. #include <boost/compute/utility/source.hpp>
  17. #include "context_setup.hpp"
  18. #include "check_macros.hpp"
  19. namespace compute = boost::compute;
  20. BOOST_AUTO_TEST_CASE(name)
  21. {
  22. compute::kernel foo = compute::kernel::create_with_source(
  23. "__kernel void foo(int x) { }", "foo", context
  24. );
  25. BOOST_CHECK_EQUAL(foo.name(), "foo");
  26. compute::kernel bar = compute::kernel::create_with_source(
  27. "__kernel void bar(float x) { }", "bar", context
  28. );
  29. BOOST_CHECK_EQUAL(bar.name(), "bar");
  30. }
  31. BOOST_AUTO_TEST_CASE(arity)
  32. {
  33. compute::kernel foo = compute::kernel::create_with_source(
  34. "__kernel void foo(int x) { }", "foo", context
  35. );
  36. BOOST_CHECK_EQUAL(foo.arity(), size_t(1));
  37. compute::kernel bar = compute::kernel::create_with_source(
  38. "__kernel void bar(float x, float y) { }", "bar", context
  39. );
  40. BOOST_CHECK_EQUAL(bar.arity(), size_t(2));
  41. compute::kernel baz = compute::kernel::create_with_source(
  42. "__kernel void baz(char x, char y, char z) { }", "baz", context
  43. );
  44. BOOST_CHECK_EQUAL(baz.arity(), size_t(3));
  45. }
  46. BOOST_AUTO_TEST_CASE(set_buffer_arg)
  47. {
  48. const char source[] = BOOST_COMPUTE_STRINGIZE_SOURCE(
  49. __kernel void foo(__global int *x, __global int *y)
  50. {
  51. x[get_global_id(0)] = -y[get_global_id(0)];
  52. }
  53. );
  54. compute::kernel foo =
  55. compute::kernel::create_with_source(source, "foo", context);
  56. compute::buffer x(context, 16);
  57. compute::buffer y(context, 16);
  58. foo.set_arg(0, x);
  59. foo.set_arg(1, y.get());
  60. }
  61. BOOST_AUTO_TEST_CASE(get_work_group_info)
  62. {
  63. const char source[] =
  64. "__kernel void sum(__global const float *input,\n"
  65. " __global float *output)\n"
  66. "{\n"
  67. " __local float scratch[16];\n"
  68. " const uint gid = get_global_id(0);\n"
  69. " const uint lid = get_local_id(0);\n"
  70. " if(lid < 16)\n"
  71. " scratch[lid] = input[gid];\n"
  72. "}\n";
  73. compute::program program =
  74. compute::program::create_with_source(source, context);
  75. program.build();
  76. compute::kernel kernel = program.create_kernel("sum");
  77. using compute::ulong_;
  78. // get local memory size
  79. kernel.get_work_group_info<ulong_>(device, CL_KERNEL_LOCAL_MEM_SIZE);
  80. // check work group size
  81. size_t work_group_size =
  82. kernel.get_work_group_info<size_t>(device, CL_KERNEL_WORK_GROUP_SIZE);
  83. BOOST_CHECK(work_group_size >= 1);
  84. }
  85. #ifndef BOOST_COMPUTE_NO_VARIADIC_TEMPLATES
  86. BOOST_AUTO_TEST_CASE(kernel_set_args)
  87. {
  88. compute::kernel k = compute::kernel::create_with_source(
  89. "__kernel void test(int x, float y, char z) { }", "test", context
  90. );
  91. k.set_args(4, 2.4f, 'a');
  92. }
  93. #endif // BOOST_COMPUTE_NO_VARIADIC_TEMPLATES
  94. // Originally failed to compile on macOS (several types are resolved differently)
  95. BOOST_AUTO_TEST_CASE(kernel_set_args_mac)
  96. {
  97. compute::kernel k = compute::kernel::create_with_source(
  98. "__kernel void test(unsigned int a, unsigned long b) { }", "test", context
  99. );
  100. compute::uint_ a;
  101. compute::ulong_ b;
  102. k.set_arg(0, a);
  103. k.set_arg(1, b);
  104. }
  105. #ifdef BOOST_COMPUTE_CL_VERSION_1_2
  106. BOOST_AUTO_TEST_CASE(get_arg_info)
  107. {
  108. REQUIRES_OPENCL_VERSION(1, 2);
  109. const char source[] = BOOST_COMPUTE_STRINGIZE_SOURCE(
  110. __kernel void sum_kernel(__global const int *input,
  111. const uint size,
  112. __global int *result)
  113. {
  114. int sum = 0;
  115. for(uint i = 0; i < size; i++){
  116. sum += input[i];
  117. }
  118. *result = sum;
  119. }
  120. );
  121. compute::program program =
  122. compute::program::create_with_source(source, context);
  123. program.build("-cl-kernel-arg-info");
  124. compute::kernel kernel = program.create_kernel("sum_kernel");
  125. BOOST_CHECK_EQUAL(kernel.get_info<CL_KERNEL_NUM_ARGS>(), compute::uint_(3));
  126. BOOST_CHECK_EQUAL(kernel.get_arg_info<std::string>(0, CL_KERNEL_ARG_TYPE_NAME), "int*");
  127. BOOST_CHECK_EQUAL(kernel.get_arg_info<std::string>(0, CL_KERNEL_ARG_NAME), "input");
  128. BOOST_CHECK_EQUAL(kernel.get_arg_info<std::string>(1, CL_KERNEL_ARG_TYPE_NAME), "uint");
  129. BOOST_CHECK_EQUAL(kernel.get_arg_info<std::string>(1, CL_KERNEL_ARG_NAME), "size");
  130. BOOST_CHECK_EQUAL(kernel.get_arg_info<std::string>(2, CL_KERNEL_ARG_TYPE_NAME), "int*");
  131. BOOST_CHECK_EQUAL(kernel.get_arg_info<std::string>(2, CL_KERNEL_ARG_NAME), "result");
  132. BOOST_CHECK_EQUAL(kernel.get_arg_info<CL_KERNEL_ARG_TYPE_NAME>(0), "int*");
  133. BOOST_CHECK_EQUAL(kernel.get_arg_info<CL_KERNEL_ARG_NAME>(0), "input");
  134. BOOST_CHECK_EQUAL(kernel.get_arg_info<CL_KERNEL_ARG_TYPE_NAME>(1), "uint");
  135. BOOST_CHECK_EQUAL(kernel.get_arg_info<CL_KERNEL_ARG_NAME>(1), "size");
  136. BOOST_CHECK_EQUAL(kernel.get_arg_info<CL_KERNEL_ARG_TYPE_NAME>(2), "int*");
  137. BOOST_CHECK_EQUAL(kernel.get_arg_info<CL_KERNEL_ARG_NAME>(2), "result");
  138. }
  139. #endif // BOOST_COMPUTE_CL_VERSION_1_2
  140. #ifdef BOOST_COMPUTE_CL_VERSION_2_0
  141. #ifndef CL_KERNEL_MAX_SUB_GROUP_SIZE_FOR_NDRANGE_KHR
  142. #define CL_KERNEL_MAX_SUB_GROUP_SIZE_FOR_NDRANGE_KHR CL_KERNEL_MAX_SUB_GROUP_SIZE_FOR_NDRANGE
  143. #endif
  144. #ifndef CL_KERNEL_SUB_GROUP_COUNT_FOR_NDRANGE_KHR
  145. #define CL_KERNEL_SUB_GROUP_COUNT_FOR_NDRANGE_KHR CL_KERNEL_SUB_GROUP_COUNT_FOR_NDRANGE
  146. #endif
  147. BOOST_AUTO_TEST_CASE(get_sub_group_info_ext)
  148. {
  149. compute::kernel k = compute::kernel::create_with_source(
  150. "__kernel void test(float x) { }", "test", context
  151. );
  152. // get_sub_group_info(const device&, cl_kernel_sub_group_info, const std::vector<size_t>)
  153. std::vector<size_t> local_work_size(2, size_t(64));
  154. boost::optional<size_t> count = k.get_sub_group_info<size_t>(
  155. device,
  156. CL_KERNEL_SUB_GROUP_COUNT_FOR_NDRANGE_KHR,
  157. local_work_size
  158. );
  159. #ifdef BOOST_COMPUTE_CL_VERSION_2_1
  160. if(device.check_version(2, 1))
  161. {
  162. BOOST_CHECK(count);
  163. }
  164. else
  165. #endif // BOOST_COMPUTE_CL_VERSION_2_1
  166. if(device.check_version(2, 0) && device.supports_extension("cl_khr_subgroups"))
  167. {
  168. // for device with cl_khr_subgroups it should return some value
  169. BOOST_CHECK(count);
  170. }
  171. else
  172. {
  173. // for device without cl_khr_subgroups ext it should return null optional
  174. BOOST_CHECK(count == boost::none);
  175. }
  176. // get_sub_group_info(const device&, cl_kernel_sub_group_info, const size_t, const void *)
  177. count = k.get_sub_group_info<size_t>(
  178. device,
  179. CL_KERNEL_SUB_GROUP_COUNT_FOR_NDRANGE_KHR,
  180. 2 * sizeof(size_t),
  181. &local_work_size[0]
  182. );
  183. #ifdef BOOST_COMPUTE_CL_VERSION_2_1
  184. if(device.check_version(2, 1))
  185. {
  186. BOOST_CHECK(count);
  187. }
  188. else
  189. #endif // BOOST_COMPUTE_CL_VERSION_2_1
  190. if(device.check_version(2, 0) && device.supports_extension("cl_khr_subgroups"))
  191. {
  192. // for device with cl_khr_subgroups it should return some value
  193. BOOST_CHECK(count);
  194. }
  195. else
  196. {
  197. // for device without cl_khr_subgroups ext it should return null optional
  198. BOOST_CHECK(count == boost::none);
  199. }
  200. }
  201. #endif // BOOST_COMPUTE_CL_VERSION_2_0
  202. #ifdef BOOST_COMPUTE_CL_VERSION_2_1
  203. BOOST_AUTO_TEST_CASE(get_sub_group_info_core)
  204. {
  205. compute::kernel k = compute::kernel::create_with_source(
  206. "__kernel void test(float x) { }", "test", context
  207. );
  208. // get_sub_group_info(const device&, cl_kernel_sub_group_info, const size_t)
  209. boost::optional<std::vector<size_t>> local_size =
  210. k.get_sub_group_info<std::vector<size_t> >(
  211. device,
  212. CL_KERNEL_LOCAL_SIZE_FOR_SUB_GROUP_COUNT,
  213. size_t(1)
  214. );
  215. if(device.check_version(2, 1))
  216. {
  217. // for 2.1 devices it should return some value
  218. BOOST_CHECK(local_size);
  219. BOOST_CHECK(local_size.value().size() == 3);
  220. }
  221. else
  222. {
  223. // for 1.x and 2.0 devices it should return null optional,
  224. // because CL_KERNEL_LOCAL_SIZE_FOR_SUB_GROUP_COUNT is not
  225. // supported by cl_khr_subgroups (2.0 ext)
  226. BOOST_CHECK(local_size == boost::none);
  227. }
  228. // get_sub_group_info(const device&, cl_kernel_sub_group_info, const size_t)
  229. boost::optional<size_t> local_size_simple =
  230. k.get_sub_group_info<size_t>(
  231. device,
  232. CL_KERNEL_LOCAL_SIZE_FOR_SUB_GROUP_COUNT,
  233. size_t(1)
  234. );
  235. if(device.check_version(2, 1))
  236. {
  237. // for 2.1 devices it should return some value
  238. BOOST_CHECK(local_size_simple);
  239. }
  240. else
  241. {
  242. // for 1.x and 2.0 devices it should return null optional,
  243. // because CL_KERNEL_LOCAL_SIZE_FOR_SUB_GROUP_COUNT is not
  244. // supported by cl_khr_subgroups (2.0 ext)
  245. BOOST_CHECK(local_size_simple == boost::none);
  246. }
  247. // get_sub_group_info(const device&, cl_kernel_sub_group_info)
  248. boost::optional<size_t> max =
  249. k.get_sub_group_info<size_t>(
  250. device,
  251. CL_KERNEL_MAX_NUM_SUB_GROUPS
  252. );
  253. if(device.check_version(2, 1))
  254. {
  255. // for 2.1 devices it should return some value
  256. BOOST_CHECK(max);
  257. }
  258. else
  259. {
  260. // for 1.x and 2.0 devices it should return null optional,
  261. // because CL_KERNEL_MAX_NUM_SUB_GROUPS is not
  262. // supported by cl_khr_subgroups (2.0 ext)
  263. BOOST_CHECK(max == boost::none);
  264. }
  265. }
  266. #endif // BOOST_COMPUTE_CL_VERSION_2_1
  267. #ifdef BOOST_COMPUTE_CL_VERSION_2_1
  268. BOOST_AUTO_TEST_CASE(clone_kernel)
  269. {
  270. REQUIRES_OPENCL_PLATFORM_VERSION(2, 1);
  271. compute::kernel k1 = compute::kernel::create_with_source(
  272. "__kernel void test(__global int * x) { x[get_global_id(0)] = get_global_id(0); }",
  273. "test", context
  274. );
  275. compute::buffer x(context, 5 * sizeof(compute::int_));
  276. k1.set_arg(0, x);
  277. // Clone k1 kernel
  278. compute::kernel k2 = k1.clone();
  279. // After clone k2 0th argument (__global float * x) should be set,
  280. // so we should be able to enqueue k2 kernel without problems
  281. queue.enqueue_1d_range_kernel(k2, 0, x.size() / sizeof(compute::int_), 0).wait();
  282. }
  283. #endif // BOOST_COMPUTE_CL_VERSION_2_1
  284. BOOST_AUTO_TEST_SUITE_END()