| // |
| // Copyright (c) 2017 The Khronos Group Inc. |
| // |
| // Licensed under the Apache License, Version 2.0 (the "License"); |
| // you may not use this file except in compliance with the License. |
| // You may obtain a copy of the License at |
| // |
| // http://www.apache.org/licenses/LICENSE-2.0 |
| // |
| // Unless required by applicable law or agreed to in writing, software |
| // distributed under the License is distributed on an "AS IS" BASIS, |
| // WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. |
| // See the License for the specific language governing permissions and |
| // limitations under the License. |
| // |
| #include "common.h" |
| |
| const char *SVMPointerPassing_test_kernel[] = { |
| "__kernel void verify_char(__global uchar* pChar, volatile __global uint* num_correct, uchar expected)\n" |
| "{\n" |
| " if(0 == get_global_id(0))\n" |
| " {\n" |
| " *num_correct = 0;\n" |
| " if(*pChar == expected)\n" |
| " {\n" |
| " *num_correct=1;\n" |
| " }\n" |
| " }\n" |
| "}\n" |
| }; |
| |
| |
| // Test that arbitrarily aligned char pointers into shared buffers can be passed directly to a kernel. |
| // This iterates through a buffer passing a pointer to each location to the kernel. |
| // The buffer is initialized to known values at each location. |
| // The kernel checks that it finds the expected value at each location. |
| // TODO: possibly make this work across all base types (including typeN?), also check ptr arithmetic ++,--. |
| int test_svm_pointer_passing(cl_device_id deviceID, cl_context context2, cl_command_queue queue, int num_elements) |
| { |
| clContextWrapper context = NULL; |
| clProgramWrapper program = NULL; |
| cl_uint num_devices = 0; |
| cl_int error = CL_SUCCESS; |
| clCommandQueueWrapper queues[MAXQ]; |
| |
| error = create_cl_objects(deviceID, &SVMPointerPassing_test_kernel[0], &context, &program, &queues[0], &num_devices, CL_DEVICE_SVM_COARSE_GRAIN_BUFFER); |
| if(error) return -1; |
| |
| clKernelWrapper kernel_verify_char = clCreateKernel(program, "verify_char", &error); |
| test_error(error,"clCreateKernel failed"); |
| |
| size_t bufSize = 256; |
| cl_uchar *pbuf_svm_alloc = (cl_uchar*) clSVMAlloc(context, CL_MEM_READ_WRITE, sizeof(cl_uchar)*bufSize, 0); |
| |
| cl_int *pNumCorrect = NULL; |
| pNumCorrect = (cl_int*) clSVMAlloc(context, CL_MEM_READ_WRITE, sizeof(cl_int), 0); |
| |
| { |
| clMemWrapper buf = clCreateBuffer(context, CL_MEM_USE_HOST_PTR, sizeof(cl_uchar)*bufSize, pbuf_svm_alloc, &error); |
| test_error(error, "clCreateBuffer failed."); |
| |
| clMemWrapper num_correct = clCreateBuffer(context, CL_MEM_USE_HOST_PTR, sizeof(cl_int), pNumCorrect, &error); |
| test_error(error, "clCreateBuffer failed."); |
| |
| error = clSetKernelArg(kernel_verify_char, 1, sizeof(void*), (void *) &num_correct); |
| test_error(error, "clSetKernelArg failed"); |
| |
| // put values into buf so that we can expect to see these values in the kernel when we pass a pointer to them. |
| cl_command_queue cmdq = queues[0]; |
| cl_uchar* pbuf_map_buffer = (cl_uchar*) clEnqueueMapBuffer(cmdq, buf, CL_TRUE, CL_MAP_READ | CL_MAP_WRITE, 0, sizeof(cl_uchar)*bufSize, 0, NULL,NULL, &error); |
| test_error2(error, pbuf_map_buffer, "clEnqueueMapBuffer failed"); |
| for(int i = 0; i<(int)bufSize; i++) |
| { |
| pbuf_map_buffer[i]= (cl_uchar)i; |
| } |
| error = clEnqueueUnmapMemObject(cmdq, buf, pbuf_map_buffer, 0,NULL,NULL); |
| test_error(error, "clEnqueueUnmapMemObject failed."); |
| |
| for (cl_uint ii = 0; ii<num_devices; ++ii) // iterate over all devices in the platform. |
| { |
| cmdq = queues[ii]; |
| for(int i = 0; i<(int)bufSize; i++) |
| { |
| cl_uchar* pChar = &pbuf_svm_alloc[i]; |
| error = clSetKernelArgSVMPointer(kernel_verify_char, 0, pChar); // pass a pointer to a location within the buffer |
| test_error(error, "clSetKernelArg failed"); |
| error = clSetKernelArg(kernel_verify_char, 2, sizeof(cl_uchar), (void *) &i ); // pass the expected value at the above location. |
| test_error(error, "clSetKernelArg failed"); |
| error = clEnqueueNDRangeKernel(cmdq, kernel_verify_char, 1, NULL, &bufSize, NULL, 0, NULL, NULL); |
| test_error(error,"clEnqueueNDRangeKernel failed"); |
| |
| pNumCorrect = (cl_int*) clEnqueueMapBuffer(cmdq, num_correct, CL_TRUE, CL_MAP_READ | CL_MAP_WRITE, 0, sizeof(cl_int), 0, NULL,NULL, &error); |
| test_error2(error, pNumCorrect, "clEnqueueMapBuffer failed"); |
| cl_int correct_count = *pNumCorrect; |
| error = clEnqueueUnmapMemObject(cmdq, num_correct, pNumCorrect, 0,NULL,NULL); |
| test_error(error, "clEnqueueUnmapMemObject failed."); |
| |
| if(correct_count != 1) |
| { |
| log_error("Passing pointer directly to kernel for byte #%d failed on device %d\n", i, ii); |
| return -1; |
| } |
| } |
| } |
| |
| error = clFinish(cmdq); |
| test_error(error, "clFinish failed"); |
| } |
| |
| |
| clSVMFree(context, pbuf_svm_alloc); |
| clSVMFree(context, pNumCorrect); |
| |
| return 0; |
| } |