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 T>
test_fmath(cl_device_id deviceID,cl_context context,cl_command_queue queue,const char * spvName,const char * funcName,const char * Tname,bool fast_math,std::vector<T> & h_lhs,std::vector<T> & h_rhs)21 int test_fmath(cl_device_id deviceID,
22 cl_context context,
23 cl_command_queue queue,
24 const char *spvName,
25 const char *funcName,
26 const char *Tname,
27 bool fast_math,
28 std::vector<T> &h_lhs,
29 std::vector<T> &h_rhs)
30 {
31
32 if(std::string(Tname).find("double") != std::string::npos) {
33 if(!is_extension_available(deviceID, "cl_khr_fp64")) {
34 log_info("Extension cl_khr_fp64 not supported; skipping double tests.\n");
35 return 0;
36 }
37 }
38 cl_int err = CL_SUCCESS;
39 int num = (int)h_lhs.size();
40 size_t bytes = num * sizeof(T);
41
42 clMemWrapper lhs = clCreateBuffer(context, CL_MEM_READ_ONLY, bytes, NULL, &err);
43 SPIRV_CHECK_ERROR(err, "Failed to create lhs buffer");
44
45 err = clEnqueueWriteBuffer(queue, lhs, CL_TRUE, 0, 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, bytes, NULL, &err);
49 SPIRV_CHECK_ERROR(err, "Failed to create rhs buffer");
50
51 err = clEnqueueWriteBuffer(queue, rhs, CL_TRUE, 0, 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<T>::value) {
60 kernelStream << "#pragma OPENCL EXTENSION cl_khr_fp64 : enable\n";
61 } else if (sizeof(T) == sizeof(cl_half)) {
62 kernelStream << "#pragma OPENCL EXTENSION cl_khr_fp16 : enable\n";
63 }
64
65 kernelStream << "#define spirv_fadd(a, b) (a) + (b) \n";
66 kernelStream << "#define spirv_fsub(a, b) (a) - (b) \n";
67 kernelStream << "#define spirv_fmul(a, b) (a) * (b) \n";
68 kernelStream << "#define spirv_fdiv(a, b) (a) / (b) \n";
69 kernelStream << "#define spirv_frem(a, b) fmod(a, b) \n";
70 kernelStream << "#define spirv_fmod(a, b) copysign(fmod(a,b),b) \n";
71 kernelStream << "#define T " << Tname << "\n";
72 kernelStream << "#define FUNC spirv_" << funcName << "\n";
73 kernelStream << "__kernel void fmath_cl(__global T *out, \n";
74 kernelStream << "const __global T *lhs, const __global T *rhs) \n";
75 kernelStream << "{ \n";
76 kernelStream << " int id = get_global_id(0); \n";
77 kernelStream << " out[id] = FUNC(lhs[id], rhs[id]); \n";
78 kernelStream << "} \n";
79 kernelStr = kernelStream.str();
80 }
81
82 const char *kernelBuf = kernelStr.c_str();
83
84 std::vector<T> h_ref(num);
85
86 {
87 // Run the cl kernel for reference results
88 clProgramWrapper prog;
89 clKernelWrapper kernel;
90 err = create_single_kernel_helper(context, &prog, &kernel, 1,
91 &kernelBuf, "fmath_cl");
92 SPIRV_CHECK_ERROR(err, "Failed to create cl kernel");
93
94 clMemWrapper ref = clCreateBuffer(context, CL_MEM_READ_WRITE, bytes, NULL, &err);
95 SPIRV_CHECK_ERROR(err, "Failed to create ref buffer");
96
97 err = clSetKernelArg(kernel, 0, sizeof(cl_mem), &ref);
98 SPIRV_CHECK_ERROR(err, "Failed to set arg 0");
99
100 err = clSetKernelArg(kernel, 1, sizeof(cl_mem), &lhs);
101 SPIRV_CHECK_ERROR(err, "Failed to set arg 1");
102
103 err = clSetKernelArg(kernel, 2, sizeof(cl_mem), &rhs);
104 SPIRV_CHECK_ERROR(err, "Failed to set arg 2");
105
106 size_t global = num;
107 err = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &global, NULL, 0, NULL, NULL);
108 SPIRV_CHECK_ERROR(err, "Failed to enqueue cl kernel");
109
110 err = clEnqueueReadBuffer(queue, ref, CL_TRUE, 0, bytes, &h_ref[0], 0, NULL, NULL);
111 SPIRV_CHECK_ERROR(err, "Failed to read from ref");
112 }
113
114 clProgramWrapper prog;
115 err = get_program_with_il(prog, deviceID, context, spvName);
116 SPIRV_CHECK_ERROR(err, "Failed to build program");
117
118 clKernelWrapper kernel = clCreateKernel(prog, "fmath_spv", &err);
119 SPIRV_CHECK_ERROR(err, "Failed to create spv kernel");
120
121 clMemWrapper res = clCreateBuffer(context, CL_MEM_READ_WRITE, bytes, NULL, &err);
122 SPIRV_CHECK_ERROR(err, "Failed to create res buffer");
123
124 err = clSetKernelArg(kernel, 0, sizeof(cl_mem), &res);
125 SPIRV_CHECK_ERROR(err, "Failed to set arg 0");
126
127 err = clSetKernelArg(kernel, 1, sizeof(cl_mem), &lhs);
128 SPIRV_CHECK_ERROR(err, "Failed to set arg 1");
129
130 err = clSetKernelArg(kernel, 2, sizeof(cl_mem), &rhs);
131 SPIRV_CHECK_ERROR(err, "Failed to set arg 2");
132
133 size_t global = num;
134 err = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &global, NULL, 0, NULL, NULL);
135 SPIRV_CHECK_ERROR(err, "Failed to enqueue cl kernel");
136
137 std::vector<T> h_res(num);
138 err = clEnqueueReadBuffer(queue, res, CL_TRUE, 0, bytes, &h_res[0], 0, NULL, NULL);
139 SPIRV_CHECK_ERROR(err, "Failed to read from ref");
140
141 for (int i = 0; i < num; i++) {
142 if (h_res[i] != h_ref[i]) {
143 log_error("Values do not match at location %d\n", i);
144 return -1;
145 }
146 }
147 return 0;
148 }
149
150 #define TEST_FMATH_FUNC(TYPE, FUNC, MODE) \
151 TEST_SPIRV_FUNC(op_##FUNC##_##TYPE##_##MODE) \
152 { \
153 if (sizeof(cl_##TYPE) == 2) { \
154 PASSIVE_REQUIRE_FP16_SUPPORT(deviceID); \
155 } \
156 const int num = 1 << 20; \
157 std::vector<cl_##TYPE> lhs(num); \
158 std::vector<cl_##TYPE> rhs(num); \
159 \
160 RandomSeed seed(gRandomSeed); \
161 \
162 for (int i = 0; i < num; i++) { \
163 lhs[i] = genrandReal<cl_##TYPE>(seed); \
164 rhs[i] = genrandReal<cl_##TYPE>(seed); \
165 } \
166 \
167 const char *mode = #MODE; \
168 return test_fmath(deviceID, context, queue, \
169 #FUNC "_" #TYPE, \
170 #FUNC, \
171 #TYPE, \
172 mode[0] == 'f', \
173 lhs, rhs); \
174 }
175
176 #define TEST_FMATH_MODE(TYPE, MODE) \
177 TEST_FMATH_FUNC(TYPE, fadd, MODE) \
178 TEST_FMATH_FUNC(TYPE, fsub, MODE) \
179 TEST_FMATH_FUNC(TYPE, fmul, MODE) \
180 TEST_FMATH_FUNC(TYPE, fdiv, MODE) \
181 TEST_FMATH_FUNC(TYPE, frem, MODE) \
182 TEST_FMATH_FUNC(TYPE, fmod, MODE) \
183
184 #define TEST_FMATH_TYPE(TYPE) \
185 TEST_FMATH_MODE(TYPE, regular) \
186 TEST_FMATH_MODE(TYPE, fast) \
187
188 TEST_FMATH_TYPE(float)
189 TEST_FMATH_TYPE(double)
190
191 TEST_FMATH_TYPE(float4)
192 TEST_FMATH_TYPE(double2)
193
194 TEST_FMATH_TYPE(half)
195