1 /******************************************************************
2 Copyright (c) 2016 The Khronos Group Inc. All Rights Reserved.
3
4 This code is protected by copyright laws and contains material proprietary to the Khronos Group, Inc.
5 This is UNPUBLISHED PROPRIETARY SOURCE CODE that may not be disclosed in whole or in part to
6 third parties, and may not be reproduced, republished, distributed, transmitted, displayed,
7 broadcast or otherwise exploited in any manner without the express prior written permission
8 of Khronos Group. The receipt or possession of this code does not convey any rights to reproduce,
9 disclose, or distribute its contents, or to manufacture, use, or sell anything that it may describe,
10 in whole or in part other than under the terms of the Khronos Adopters Agreement
11 or Khronos Conformance Test Source License Agreement as executed between Khronos and the recipient.
12 ******************************************************************/
13
14 #include "testBase.h"
15 #include "types.hpp"
16
17 #include <sstream>
18 #include <string>
19
20 template<typename Tv, typename Ts>
test_vector_times_scalar(cl_device_id deviceID,cl_context context,cl_command_queue queue,const char * Tname,std::vector<Tv> & h_lhs,std::vector<Ts> & h_rhs)21 int test_vector_times_scalar(cl_device_id deviceID,
22 cl_context context,
23 cl_command_queue queue,
24 const char *Tname,
25 std::vector<Tv> &h_lhs,
26 std::vector<Ts> &h_rhs)
27 {
28 if(std::string(Tname).find("double") != std::string::npos) {
29 if(!is_extension_available(deviceID, "cl_khr_fp64")) {
30 log_info("Extension cl_khr_fp64 not supported; skipping double tests.\n");
31 return 0;
32 }
33 }
34
35 cl_int err = CL_SUCCESS;
36 int num = (int)h_lhs.size();
37 size_t lhs_bytes = num * sizeof(Tv);
38 size_t rhs_bytes = num * sizeof(Ts);
39 size_t res_bytes = lhs_bytes;
40 int vec_size = sizeof(Tv) / sizeof(Ts);
41
42 clMemWrapper lhs = clCreateBuffer(context, CL_MEM_READ_ONLY, lhs_bytes, NULL, &err);
43 SPIRV_CHECK_ERROR(err, "Failed to create lhs buffer");
44
45 err = clEnqueueWriteBuffer(queue, lhs, CL_TRUE, 0, lhs_bytes, &h_lhs[0], 0, NULL, NULL);
46 SPIRV_CHECK_ERROR(err, "Failed to copy to lhs buffer");
47
48 clMemWrapper rhs = clCreateBuffer(context, CL_MEM_READ_ONLY, rhs_bytes, NULL, &err);
49 SPIRV_CHECK_ERROR(err, "Failed to create rhs buffer");
50
51 err = clEnqueueWriteBuffer(queue, rhs, CL_TRUE, 0, rhs_bytes, &h_rhs[0], 0, NULL, NULL);
52 SPIRV_CHECK_ERROR(err, "Failed to copy to rhs buffer");
53
54 std::string kernelStr;
55
56 {
57 std::stringstream kernelStream;
58
59 if (is_double<Ts>::value) {
60 kernelStream << "#pragma OPENCL EXTENSION cl_khr_fp64 : enable\n";
61 } else if (sizeof(Ts) == sizeof(cl_half)) {
62 kernelStream << "#pragma OPENCL EXTENSION cl_khr_fp16 : enable\n";
63 }
64
65 kernelStream << "#define Ts " << Tname << "\n";
66 kernelStream << "#define Tv " << Tname << vec_size << "\n";
67 kernelStream << "__kernel void vector_times_scalar( \n";
68 kernelStream << " __global Tv *out, \n";
69 kernelStream << " const __global Tv *lhs,\n";
70 kernelStream << " const __global Ts *rhs)\n";
71 kernelStream << "{ \n";
72 kernelStream << " int id = get_global_id(0); \n";
73 kernelStream << " out[id] = lhs[id] * rhs[id]; \n";
74 kernelStream << "} \n";
75 kernelStr = kernelStream.str();
76 }
77
78 size_t kernelLen = kernelStr.size();
79 const char *kernelBuf = kernelStr.c_str();
80
81 std::vector<Tv> h_ref(num);
82 {
83 // Run the cl kernel for reference results
84 clProgramWrapper prog;
85 clKernelWrapper kernel;
86 err = create_single_kernel_helper(context, &prog, &kernel, 1,
87 &kernelBuf, "vector_times_scalar");
88 SPIRV_CHECK_ERROR(err, "Failed to create cl program");
89
90 clMemWrapper ref = clCreateBuffer(context, CL_MEM_READ_WRITE, res_bytes, NULL, &err);
91 SPIRV_CHECK_ERROR(err, "Failed to create ref buffer");
92
93 err = clSetKernelArg(kernel, 0, sizeof(cl_mem), &ref);
94 SPIRV_CHECK_ERROR(err, "Failed to set arg 0");
95
96 err = clSetKernelArg(kernel, 1, sizeof(cl_mem), &lhs);
97 SPIRV_CHECK_ERROR(err, "Failed to set arg 1");
98
99 err = clSetKernelArg(kernel, 2, sizeof(cl_mem), &rhs);
100 SPIRV_CHECK_ERROR(err, "Failed to set arg 2");
101
102 size_t global = num;
103 err = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &global, NULL, 0, NULL, NULL);
104 SPIRV_CHECK_ERROR(err, "Failed to enqueue cl kernel");
105
106 err = clEnqueueReadBuffer(queue, ref, CL_TRUE, 0, res_bytes, &h_ref[0], 0, NULL, NULL);
107 SPIRV_CHECK_ERROR(err, "Failed to read from ref");
108 }
109
110 cl_uint bits = sizeof(void *) * 8;
111 std::string ref = "vector_times_scalar_";
112 ref += Tname;
113 const char *spvName = ref.c_str();
114
115 clProgramWrapper prog;
116 err = get_program_with_il(prog, deviceID, context, spvName);
117 SPIRV_CHECK_ERROR(err, "Failed to build program");
118
119 clKernelWrapper kernel = clCreateKernel(prog, "vector_times_scalar", &err);
120 SPIRV_CHECK_ERROR(err, "Failed to create spv kernel");
121
122 clMemWrapper res = clCreateBuffer(context, CL_MEM_READ_WRITE, res_bytes, NULL, &err);
123 SPIRV_CHECK_ERROR(err, "Failed to create res buffer");
124
125 err = clSetKernelArg(kernel, 0, sizeof(cl_mem), &res);
126 SPIRV_CHECK_ERROR(err, "Failed to set arg 0");
127
128 err = clSetKernelArg(kernel, 1, sizeof(cl_mem), &lhs);
129 SPIRV_CHECK_ERROR(err, "Failed to set arg 1");
130
131 err = clSetKernelArg(kernel, 2, sizeof(cl_mem), &rhs);
132 SPIRV_CHECK_ERROR(err, "Failed to set arg 2");
133
134 size_t global = num;
135 err = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &global, NULL, 0, NULL, NULL);
136 SPIRV_CHECK_ERROR(err, "Failed to enqueue cl kernel");
137
138 std::vector<Tv> h_res(num);
139 err = clEnqueueReadBuffer(queue, res, CL_TRUE, 0, res_bytes, &h_res[0], 0, NULL, NULL);
140 SPIRV_CHECK_ERROR(err, "Failed to read from ref");
141
142 for (int i = 0; i < num; i++) {
143 if (h_res[i] != h_ref[i]) {
144 log_error("Values do not match at location %d\n", i);
145 return -1;
146 }
147 }
148 return 0;
149 }
150
151 #define TEST_VECTOR_TIMES_SCALAR(TYPE, N) \
152 TEST_SPIRV_FUNC(op_vector_times_scalar_##TYPE) \
153 { \
154 if (sizeof(cl_##TYPE) == 2) { \
155 PASSIVE_REQUIRE_FP16_SUPPORT(deviceID); \
156 } \
157 typedef cl_##TYPE##N Tv; \
158 typedef cl_##TYPE Ts; \
159 const int num = 1 << 20; \
160 std::vector<Tv> lhs(num); \
161 std::vector<Ts> rhs(num); \
162 \
163 RandomSeed seed(gRandomSeed); \
164 \
165 for (int i = 0; i < num; i++) { \
166 lhs[i] = genrandReal<cl_##TYPE##N>(seed); \
167 rhs[i] = genrandReal<cl_##TYPE>(seed); \
168 } \
169 \
170 return test_vector_times_scalar<Tv, Ts>(deviceID, \
171 context, queue, \
172 #TYPE, \
173 lhs, rhs); \
174 }
175
176 TEST_VECTOR_TIMES_SCALAR(float, 4)
177 TEST_VECTOR_TIMES_SCALAR(double, 4)
178