2 Replies Latest reply on Feb 11, 2010 3:25 AM by genaganna

    Build-in "cross( float4, float4 )" function returns invalid values on K8 core kernel

    redditisgreat

      compiled on an Athlon64 X2

      <code>

      float4 n =  cross( ab, -ca ) ;

      n = normalize( n );  // result fucks up this line ==> n = (NaN,NaN,NaN,NaN)

      </code>

      n.xyz is computed corectly but n.w is not 0 as per spec. but NaN. But only some of the time.

      More specifically, for some builds of the same kernel it works and for some it doesn't.

      I build the same Cl program multiple times with different compile-time constants (subproblem size) as build options. The compiletime constant has no influence on the parameters of the "cross" call though, it is related to other parts of the code.

      There is one problematic subproblem size lets say it's 6.

      If I define  a

      testCase<N>()
      {

      clBuildProgram( ... , "-D SUBSIZE=N", ... )

      // setup Kernel parameter

      // run Kernel

      // evaluate results

      }

      like this.

      The undefined behavior happens whenever I run multiple testcases like

      testcase<4>(); testcase<6>();

      or

      testcase<6>(); testcase<4>();

      but not in

      testcase<6>(); testcase<6>();

      or

      testcase<6>();

       

      I know I can work arround te problem by simply setting the n.w = 0.0; but this undefined behavior makes me nervous and I'd like to know what's going on.

       

       

       

       

       

        • Build-in "cross( float4, float4 )" function returns invalid values on K8 core kernel
          omkaranathan

          Do you have a test case which reproduces the problem?

          • Build-in "cross( float4, float4 )" function returns invalid values on K8 core kernel
            genaganna

             

            Originally posted by: redditisgreat compiled on an Athlon64 X2

             

             

             

            float4 n =  cross( ab, -ca ) ;

             

            n = normalize( n );  // result fucks up this line ==> n = (NaN,NaN,NaN,NaN)

             

             

            n.xyz is computed corectly but n.w is not 0 as per spec. but NaN. But only some of the time.

             

            More specifically, for some builds of the same kernel it works and for some it doesn't.

             

            I build the same Cl program multiple times with different compile-time constants (subproblem size) as build options. The compiletime constant has no influence on the parameters of the "cross" call though, it is related to other parts of the code.

             

            There is one problematic subproblem size lets say it's 6.

             

            If I define  a

             

            testCase() {

             

            clBuildProgram( ... , "-D SUBSIZE=N", ... )

             

            // setup Kernel parameter

             

            // run Kernel

             

            // evaluate results

             

            }

             

            like this.

             

            The undefined behavior happens whenever I run multiple testcases like

             

            testcase<4>(); testcase<6>();ortestcase<6>(); testcase<4>();but not intestcase<6>(); testcase<6>();ortestcase<6>();

             

              I know I can work arround te problem by simply setting the n.w = 0.0; but this undefined behavior makes me nervous and I'd like to know what's going on.

             



            Redditisgreat,

                I am not able to reproduce your issue.  See attached test case. Let me know it reproduces at you end.

            System Information
            ==================
            OS : XP64, SP1
            GPU - RV770
            CPU - AMD Phenom Quad-Core
            OpenCL - Internal version
            CAL DRIVER : Interal version

            Header file ********* /* ============================================================ Copyright (c) 2009 Advanced Micro Devices, Inc. All rights reserved. Redistribution and use of this material is permitted under the following conditions: Redistributions must retain the above copyright notice and all terms of this license. In no event shall anyone redistributing or accessing or using this material commence or participate in any arbitration or legal action relating to this material against Advanced Micro Devices, Inc. or any copyright holders or contributors. The foregoing shall survive any expiration or termination of this license or any agreement or access or use related to this material. ANY BREACH OF ANY TERM OF THIS LICENSE SHALL RESULT IN THE IMMEDIATE REVOCATION OF ALL RIGHTS TO REDISTRIBUTE, ACCESS OR USE THIS MATERIAL. THIS MATERIAL IS PROVIDED BY ADVANCED MICRO DEVICES, INC. AND ANY COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" IN ITS CURRENT CONDITION AND WITHOUT ANY REPRESENTATIONS, GUARANTEE, OR WARRANTY OF ANY KIND OR IN ANY WAY RELATED TO SUPPORT, INDEMNITY, ERROR FREE OR UNINTERRUPTED OPERA TION, OR THAT IT IS FREE FROM DEFECTS OR VIRUSES. ALL OBLIGATIONS ARE HEREBY DISCLAIMED - WHETHER EXPRESS, IMPLIED, OR STATUTORY - INCLUDING, BUT NOT LIMITED TO, ANY IMPLIED WARRANTIES OF TITLE, MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE, ACCURACY, COMPLETENESS, OPERABILITY, QUALITY OF SERVICE, OR NON-INFRINGEMENT. IN NO EVENT SHALL ADVANCED MICRO DEVICES, INC. OR ANY COPYRIGHT HOLDERS OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, PUNITIVE, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, REVENUE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED OR BASED ON ANY THEORY OF LIABILITY ARISING IN ANY WAY RELATED TO THIS MATERIAL, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. THE ENTIRE AND AGGREGATE LIABILITY OF ADVANCED MICRO DEVICES, INC. AND ANY COPYRIGHT HOLDERS AND CONTRIBUTORS SHALL NOT EXCEED TEN DOLLARS (US $10.00). ANYONE REDISTRIBUTING OR ACCESSING OR USING THIS MATERIAL ACCEPTS THIS ALLOCATION OF RISK AND AGREES TO RELEASE ADVANCED MICRO DEVICES, INC. AND ANY COPYRIGHT HOLDERS AND CONTRIBUTORS FROM ANY AND ALL LIABILITIES, OBLIGATIONS, CLAIMS, OR DEMANDS IN EXCESS OF TEN DOLLARS (US $10.00). THE FOREGOING ARE ESSENTIAL TERMS OF THIS LICENSE AND, IF ANY OF THESE TERMS ARE CONSTRUED AS UNENFORCEABLE, FAIL IN ESSENTIAL PURPOSE, OR BECOME VOID OR DETRIMENTAL TO ADVANCED MICRO DEVICES, INC. OR ANY COPYRIGHT HOLDERS OR CONTRIBUTORS FOR ANY REASON, THEN ALL RIGHTS TO REDISTRIBUTE, ACCESS OR USE THIS MATERIAL SHALL TERMINATE IMMEDIATELY. MOREOVER, THE FOREGOING SHALL SURVIVE ANY EXPIRATION OR TERMINATION OF THIS LICENSE OR ANY AGREEMENT OR ACCESS OR USE RELATED TO THIS MATERIAL. NOTICE IS HEREBY PROVIDED, AND BY REDISTRIBUTING OR ACCESSING OR USING THIS MATERIAL SUCH NOTICE IS ACKNOWLEDGED, THAT THIS MATERIAL MAY BE SUBJECT TO RESTRICTIONS UNDER THE LAWS AND REGULATIONS OF THE UNITED STATES OR OTHER COUNTRIES, WHICH INCLUDE BUT ARE NOT LIMITED TO, U.S. EXPORT CONTROL LAWS SUCH AS THE EXPORT ADMINISTRATION REGULATIONS AND NATIONAL SECURITY CONTROLS AS DEFINED THEREUNDER, AS WELL AS STATE DEPARTMENT CONTROLS UNDER THE U.S. MUNITIONS LIST. THIS MATERIAL MAY NOT BE USED, RELEASED, TRANSFERRED, IMPORTED, EXPORTED AND/OR RE-EXPORTED IN ANY MANNER PROHIBITED UNDER ANY APPLICABLE LAWS, INCLUDING U.S. EXPORT CONTROL LAWS REGARDING SPECIFICALLY DESIGNATED PERSONS, COUNTRIES AND NATIONALS OF COUNTRIES SUBJECT TO NATIONAL SECURITY CONTROLS. MOREOVER, THE FOREGOING SHALL SURVIVE ANY EXPIRATION OR TERMINATION OF ANY LICENSE OR AGREEMENT OR ACCESS OR USE RELATED TO THIS MATERIAL. NOTICE REGARDING THE U.S. GOVERNMENT AND DOD AGENCIES: This material is provided with "RESTRICTED RIGHTS" and/or "LIMITED RIGHTS" as applicable to computer software and technical data, respectively. Use, duplication, distribution or disclosure by the U.S. Government and/or DOD agencies is subject to the full extent of restrictions in all applicable regulations, including those found at FAR52.227 and DFARS252.227 et seq. and any successor regulations thereof. Use of this material by the U.S. Government and/or DOD agencies is acknowledgment of the proprietary rights of any copyright holders and contributors, including those of Advanced Micro Devices, Inc., as well as the provisions of FAR52.227-14 through 23 regarding privately developed and/or commercial computer software. This license forms the entire agreement regarding the subject matter hereof and supersedes all proposals and prior discussions and writings between the parties with respect thereto. This license does not affect any ownership, rights, title, or interest in, or relating to, this material. No terms of this license can be modified or waived, and no breach of this license can be excused, unless done so in a writing signed by all affected parties. Each term of this license is separately enforceable. If any term of this license is determined to be or becomes unenforceable or illegal, such term shall be reformed to the minimum extent necessary in order for this license to remain in effect in accordance with its terms as modified by such reformation. This license shall be governed by and construed in accordance with the laws of the State of Texas without regard to rules on conflicts of law of any state or jurisdiction or the United Nations Convention on the International Sale of Goods. All disputes arising out of this license shall be subject to the jurisdiction of the federal and state courts in Austin, Texas, and all defenses are hereby waived concerning personal jurisdiction and venue of these courts. ============================================================ */ #ifndef REDUCTION_H_ #define REDUCTION_H_ #include <CL/cl.h> #include <stdio.h> #include <stdlib.h> #include <assert.h> #include <string.h> #include <iostream> #define SDK_SUCCESS 0 #define SDK_FAILURE 1 #define GROUP_SIZE 64 cl_device_type dType = CL_DEVICE_TYPE_GPU; /** * Reduction * Class implements OpenCL Reduction sample * Derived from SDKSample base class */ class Reduction { size_t maxWorkGroupSize; /**< Max allowed work-items in a group */ cl_uint maxDimensions; /**< Max group dimensions allowed */ size_t * maxWorkItemSizes; /**< Max work-items sizes in each dimensions */ cl_ulong totalLocalMemory; /**< Max local memory allowed */ cl_ulong usedLocalMemory; /**< Used local memory */ cl_int length; /**< length of the input array */ cl_float *input; /**< Input array */ cl_float *output; /**< Output result */ cl_context context; /**< CL context */ cl_device_id *devices; /**< CL device list */ cl_mem inputBuffer; /**< CL memory buffer */ cl_mem outputBuffer; /**< CL memory buffer */ cl_command_queue commandQueue; /**< CL command queue */ cl_program program; /**< CL program */ cl_kernel kernel; /**< CL kernel */ public: /** * Constructor * Initialize member variables * @param name name of sample (string) */ explicit Reduction(std::string name) : input(NULL), output(NULL), maxWorkItemSizes(NULL), devices(NULL) { length = 8192; } /** * Constructor * Initialize member variables * @param name name of sample (const char*) */ explicit Reduction(const char* name) : input(NULL), output(NULL), maxWorkItemSizes(NULL), devices(NULL) { length = 8192; } ~Reduction(); /** * Allocate and initialize host memory array with random values * @return 1 on success and 0 on failure */ int setupReduction(); /** * OpenCL related initialisations. * Set up Context, Device list, Command Queue, Memory buffers * Build CL kernel program executable * @return 1 on success and 0 on failure */ int setupCL(); /** * Set values for kernels' arguments, enqueue calls to the kernels * on to the command queue, wait till end of kernel execution. * Get kernel start and end time if timing is enabled * @return 1 on success and 0 on failure */ int runCLKernels(); /** * Reference CPU implementation of Reduction * for performance comparison * @param input the input array * @param length length of the array * @param output value */ void reductionCPUReference( cl_uint * input, const cl_uint length, cl_uint& output); /** * Override from SDKSample. Initialize * command line parser, add custom options */ int initialize(); /** * Override from SDKSample, adjust width and height * of execution domain, perform all sample setup */ int setup(); /** * Override from SDKSample * Run OpenCL Reduction */ int run(); /** * Override from SDKSample * Cleanup memory allocations */ int cleanup(); /** * Override from SDKSample * Verify against reference implementation */ int verifyResults(); char* getKernelSourceString(const char* clFileName, const char* prependStr, size_t* strSize); }; #endif // REDUCTION_H_ CPP file ******* /* ============================================================ Copyright (c) 2009 Advanced Micro Devices, Inc. All rights reserved. Redistribution and use of this material is permitted under the following conditions: Redistributions must retain the above copyright notice and all terms of this license. In no event shall anyone redistributing or accessing or using this material commence or participate in any arbitration or legal action relating to this material against Advanced Micro Devices, Inc. or any copyright holders or contributors. The foregoing shall survive any expiration or termination of this license or any agreement or access or use related to this material. ANY BREACH OF ANY TERM OF THIS LICENSE SHALL RESULT IN THE IMMEDIATE REVOCATION OF ALL RIGHTS TO REDISTRIBUTE, ACCESS OR USE THIS MATERIAL. THIS MATERIAL IS PROVIDED BY ADVANCED MICRO DEVICES, INC. AND ANY COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" IN ITS CURRENT CONDITION AND WITHOUT ANY REPRESENTATIONS, GUARANTEE, OR WARRANTY OF ANY KIND OR IN ANY WAY RELATED TO SUPPORT, INDEMNITY, ERROR FREE OR UNINTERRUPTED OPERA TION, OR THAT IT IS FREE FROM DEFECTS OR VIRUSES. ALL OBLIGATIONS ARE HEREBY DISCLAIMED - WHETHER EXPRESS, IMPLIED, OR STATUTORY - INCLUDING, BUT NOT LIMITED TO, ANY IMPLIED WARRANTIES OF TITLE, MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE, ACCURACY, COMPLETENESS, OPERABILITY, QUALITY OF SERVICE, OR NON-INFRINGEMENT. IN NO EVENT SHALL ADVANCED MICRO DEVICES, INC. OR ANY COPYRIGHT HOLDERS OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, PUNITIVE, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, REVENUE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED OR BASED ON ANY THEORY OF LIABILITY ARISING IN ANY WAY RELATED TO THIS MATERIAL, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. THE ENTIRE AND AGGREGATE LIABILITY OF ADVANCED MICRO DEVICES, INC. AND ANY COPYRIGHT HOLDERS AND CONTRIBUTORS SHALL NOT EXCEED TEN DOLLARS (US $10.00). ANYONE REDISTRIBUTING OR ACCESSING OR USING THIS MATERIAL ACCEPTS THIS ALLOCATION OF RISK AND AGREES TO RELEASE ADVANCED MICRO DEVICES, INC. AND ANY COPYRIGHT HOLDERS AND CONTRIBUTORS FROM ANY AND ALL LIABILITIES, OBLIGATIONS, CLAIMS, OR DEMANDS IN EXCESS OF TEN DOLLARS (US $10.00). THE FOREGOING ARE ESSENTIAL TERMS OF THIS LICENSE AND, IF ANY OF THESE TERMS ARE CONSTRUED AS UNENFORCEABLE, FAIL IN ESSENTIAL PURPOSE, OR BECOME VOID OR DETRIMENTAL TO ADVANCED MICRO DEVICES, INC. OR ANY COPYRIGHT HOLDERS OR CONTRIBUTORS FOR ANY REASON, THEN ALL RIGHTS TO REDISTRIBUTE, ACCESS OR USE THIS MATERIAL SHALL TERMINATE IMMEDIATELY. MOREOVER, THE FOREGOING SHALL SURVIVE ANY EXPIRATION OR TERMINATION OF THIS LICENSE OR ANY AGREEMENT OR ACCESS OR USE RELATED TO THIS MATERIAL. NOTICE IS HEREBY PROVIDED, AND BY REDISTRIBUTING OR ACCESSING OR USING THIS MATERIAL SUCH NOTICE IS ACKNOWLEDGED, THAT THIS MATERIAL MAY BE SUBJECT TO RESTRICTIONS UNDER THE LAWS AND REGULATIONS OF THE UNITED STATES OR OTHER COUNTRIES, WHICH INCLUDE BUT ARE NOT LIMITED TO, U.S. EXPORT CONTROL LAWS SUCH AS THE EXPORT ADMINISTRATION REGULATIONS AND NATIONAL SECURITY CONTROLS AS DEFINED THEREUNDER, AS WELL AS STATE DEPARTMENT CONTROLS UNDER THE U.S. MUNITIONS LIST. THIS MATERIAL MAY NOT BE USED, RELEASED, TRANSFERRED, IMPORTED, EXPORTED AND/OR RE-EXPORTED IN ANY MANNER PROHIBITED UNDER ANY APPLICABLE LAWS, INCLUDING U.S. EXPORT CONTROL LAWS REGARDING SPECIFICALLY DESIGNATED PERSONS, COUNTRIES AND NATIONALS OF COUNTRIES SUBJECT TO NATIONAL SECURITY CONTROLS. MOREOVER, THE FOREGOING SHALL SURVIVE ANY EXPIRATION OR TERMINATION OF ANY LICENSE OR AGREEMENT OR ACCESS OR USE RELATED TO THIS MATERIAL. NOTICE REGARDING THE U.S. GOVERNMENT AND DOD AGENCIES: This material is provided with "RESTRICTED RIGHTS" and/or "LIMITED RIGHTS" as applicable to computer software and technical data, respectively. Use, duplication, distribution or disclosure by the U.S. Government and/or DOD agencies is subject to the full extent of restrictions in all applicable regulations, including those found at FAR52.227 and DFARS252.227 et seq. and any successor regulations thereof. Use of this material by the U.S. Government and/or DOD agencies is acknowledgment of the proprietary rights of any copyright holders and contributors, including those of Advanced Micro Devices, Inc., as well as the provisions of FAR52.227-14 through 23 regarding privately developed and/or commercial computer software. This license forms the entire agreement regarding the subject matter hereof and supersedes all proposals and prior discussions and writings between the parties with respect thereto. This license does not affect any ownership, rights, title, or interest in, or relating to, this material. No terms of this license can be modified or waived, and no breach of this license can be excused, unless done so in a writing signed by all affected parties. Each term of this license is separately enforceable. If any term of this license is determined to be or becomes unenforceable or illegal, such term shall be reformed to the minimum extent necessary in order for this license to remain in effect in accordance with its terms as modified by such reformation. This license shall be governed by and construed in accordance with the laws of the State of Texas without regard to rules on conflicts of law of any state or jurisdiction or the United Nations Convention on the International Sale of Goods. All disputes arising out of this license shall be subject to the jurisdiction of the federal and state courts in Austin, Texas, and all defenses are hereby waived concerning personal jurisdiction and venue of these courts. ============================================================ */ #include "Reduction.hpp" #include <math.h> int Reduction::setupReduction() { //// make sure length is multiple of group size //length = (length < GROUP_SIZE) ? GROUP_SIZE : length; //length = (length / GROUP_SIZE) * GROUP_SIZE; //input = new cl_float[length]; ///* random initialisation of input */ //for(int i = 0; i < length * VECTOR_LENGHT; ++i) // input[i] = (cl_float)(rand() % 5); //output = new cl_float[length * VECTOR_LENGHT]; //memset(output, 0, length * VECTOR_LENGHT * sizeof(cl_float)); output = new cl_float[length]; memset(output, 0, length * sizeof(cl_float)); return SDK_SUCCESS; } char* Reduction::getKernelSourceString(const char* clFileName, const char* prependStr, size_t* strSize) { // locals FILE* filePtr = NULL; size_t srcSize; //printf("FileName = %s, prependStr = %s", clFileName, prependStr); // open the OpenCL source code file #ifdef _WIN32 // Windows version if(fopen_s(&filePtr, clFileName, "rb") != 0) { printf("Failed to open file : %s \n", clFileName); return NULL; } #else // Linux version filePtr = fopen(clFileName, "rb"); if(filePtr == 0) { printf("Failed to open file : %s \n", clFileName); return NULL; } #endif size_t prependStrSize = 0; if(prependStr != NULL) prependStrSize = strlen(prependStr); // get the length of the source code fseek(filePtr, 0, SEEK_END); srcSize = ftell(filePtr); fseek(filePtr, 0, SEEK_SET); // allocate a buffer for the source code string and read it in char* str = (char *)malloc(srcSize + prependStrSize + 1); memcpy(str, prependStr, prependStrSize); fread((str) + prependStrSize, srcSize, 1, filePtr); // close the file and return the total length of the combined (prepend + source) string fclose(filePtr); if(strSize != NULL) { *strSize = srcSize + prependStrSize; } str[srcSize + prependStrSize] = '\0'; return str; } int Reduction::setupCL() { cl_int status = CL_SUCCESS; size_t deviceListSize; /* * Have a look at the available platforms and pick either * the AMD one if available or a reasonable default. */ cl_uint numPlatforms; cl_platform_id platform = NULL; status = clGetPlatformIDs(0, NULL, &numPlatforms); if(status != CL_SUCCESS) { std::cout << "clGetPlatformIDs failed." << std::endl; return SDK_FAILURE; } if (0 < numPlatforms) { cl_platform_id* platforms = new cl_platform_id[numPlatforms]; status = clGetPlatformIDs(numPlatforms, platforms, NULL); if(status != CL_SUCCESS) { std::cout << "clGetPlatformIDs failed." << std::endl; return SDK_FAILURE; } for (unsigned i = 0; i < numPlatforms; ++i) { char pbuf[100]; status = clGetPlatformInfo(platforms[i], CL_PLATFORM_VENDOR, sizeof(pbuf), pbuf, NULL); if(status != CL_SUCCESS) { std::cout << "clGetPlatformInfo failed." << std::endl; return SDK_FAILURE; } platform = platforms[i]; if (!strcmp(pbuf, "Advanced Micro Devices, Inc.")) { break; } } delete[] platforms; } /* * If we could find our platform, use it. Otherwise pass a NULL and get whatever the * implementation thinks we should be using. */ cl_context_properties cps[3] = { CL_CONTEXT_PLATFORM, (cl_context_properties)platform, 0 }; /* Use NULL for backward compatibility */ cl_context_properties* cprops = (NULL == platform) ? NULL : cps; context = clCreateContextFromType(cprops, dType, NULL, NULL, &status); if(status != CL_SUCCESS) { std::cout << "clCreateContextFromType failed."; return SDK_FAILURE; } /* First, get the size of device list data */ status = clGetContextInfo( context, CL_CONTEXT_DEVICES, 0, NULL, &deviceListSize); if(status != CL_SUCCESS) { std::cout << "clGetContextInfo failed."; return SDK_FAILURE; } /* Now allocate memory for device list based on the size we got earlier */ devices = (cl_device_id *)malloc(deviceListSize); if(devices == NULL) { std::cout << "Failed to allocate memory"; return SDK_FAILURE; } /* Now, get the device list data */ status = clGetContextInfo( context, CL_CONTEXT_DEVICES, deviceListSize, devices, NULL); if(status != CL_SUCCESS) { std::cout << "clGetContextInfo failed."; return SDK_FAILURE; } /* Create command queue */ commandQueue = clCreateCommandQueue( context, devices[0], 0, &status); if(status != CL_SUCCESS) { std::cout << "clCreateCommandQueue failed."; return SDK_FAILURE; } /* Get Device specific Information */ status = clGetDeviceInfo( devices[0], CL_DEVICE_MAX_WORK_GROUP_SIZE, sizeof(size_t), (void*)&maxWorkGroupSize, NULL); if(status != CL_SUCCESS) { std::cout << "clGetDeviceInfo failed."; return SDK_FAILURE; } status = clGetDeviceInfo( devices[0], CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS, sizeof(cl_uint), (void*)&maxDimensions, NULL); if(status != CL_SUCCESS) { std::cout << "clGetDeviceInfo failed."; return SDK_FAILURE; } maxWorkItemSizes = (size_t*)malloc(maxDimensions * sizeof(size_t)); status = clGetDeviceInfo( devices[0], CL_DEVICE_MAX_WORK_ITEM_SIZES, sizeof(size_t) * maxDimensions, (void*)maxWorkItemSizes, NULL); if(status != CL_SUCCESS) { std::cout << "clGetDeviceInfo failed."; return SDK_FAILURE; } status = clGetDeviceInfo( devices[0], CL_DEVICE_LOCAL_MEM_SIZE, sizeof(cl_ulong), (void *)&totalLocalMemory, NULL); if(status != CL_SUCCESS) { std::cout << "clGetDeviceInfo failed."; return SDK_FAILURE; } /* * Create and initialize memory objects */ /* Create memory objects for input array */ // inputBuffer = clCreateBuffer( // context, // CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, // length * VECTOR_LENGHT * sizeof(cl_float), // input, // &status); // if(status != CL_SUCCESS) // { // std::cout << "clCreateBuffer failed(input)."; //return SDK_FAILURE; // } /* Create memory objects for input array */ outputBuffer = clCreateBuffer( context, CL_MEM_WRITE_ONLY | CL_MEM_COPY_HOST_PTR, length * sizeof(cl_float), output, &status); if(status != CL_SUCCESS) { std::cout << "clCreateBuffer failed(output)."; return SDK_FAILURE; } /* create a CL program using the kernel source */ size_t sourceSize = 0; const char * source = getKernelSourceString("Reduction_Kernels.cl", NULL, &sourceSize); program = clCreateProgramWithSource( context, 1, &source, &sourceSize, &status); if(status != CL_SUCCESS) { std::cout << "clCreateProgramWithSource failed."; return SDK_FAILURE; } /* create a cl program executable for all the devices specified */ status = clBuildProgram( program, 1, devices, NULL, NULL, NULL); if(status != CL_SUCCESS) { if(status == CL_BUILD_PROGRAM_FAILURE) { cl_int logStatus; char * buildLog = NULL; size_t buildLogSize = 0; logStatus = clGetProgramBuildInfo(program, devices[0], CL_PROGRAM_BUILD_LOG, buildLogSize, buildLog, &buildLogSize); if(logStatus !=CL_SUCCESS) { std::cout << "clGetProgramBuildInfo failed." << std::endl; return SDK_FAILURE; } buildLog = (char*)malloc(buildLogSize); if(buildLog == NULL) { std::cout << "Failed to allocate host memory.(buildLog)" << std::endl; return SDK_FAILURE; } memset(buildLog, 0, buildLogSize); logStatus = clGetProgramBuildInfo(program, devices[0], CL_PROGRAM_BUILD_LOG, buildLogSize, buildLog, NULL); if(logStatus !=CL_SUCCESS) { std::cout << "clGetProgramBuildInfo failed." << std::endl; free(buildLog); return SDK_FAILURE; } std::cout << " \n\t\t\tBUILD LOG\n"; std::cout << " ************************************************\n"; std::cout << buildLog << std::endl; std::cout << " ************************************************\n"; free(buildLog); } if(status != CL_SUCCESS) { std::cout << "clBuildProgram failed." << std::endl; return SDK_FAILURE; } } /* get a kernel object handle for a kernel with the given name */ kernel = clCreateKernel( program, "qsimKernel", &status); if(status != CL_SUCCESS) { std::cout << "clCreateKernel failed."; return SDK_FAILURE; } return SDK_SUCCESS; } int Reduction::runCLKernels() { cl_int status; cl_event events[1]; /* * This algorithm reduces each group of work-items to a single value * on OpenCL device and later each reduced items per group is further * reduced to a single value on CPU */ /*** Set appropriate arguments to the kernel ***/ // /* the input array */ // status = clSetKernelArg( // kernel, // 0, // sizeof(cl_mem), // (void *)&inputBuffer); // if(status != CL_SUCCESS) // { // std::cout << "clSetKernelArg failed."; //return SDK_FAILURE; // } // /* local memory */ // status = clSetKernelArg( // kernel, // 1, // GROUP_SIZE * VECTOR_LENGHT * sizeof(cl_float), // NULL); // if(status != CL_SUCCESS) // { // std::cout << "clSetKernelArg failed(Local memory)."; //return SDK_FAILURE; // } /* temporary output buffer */ status = clSetKernelArg( kernel, 0, sizeof(cl_mem), (void *)&outputBuffer); if(status != CL_SUCCESS) { std::cout << "clSetKernelArg failed."; return SDK_FAILURE; } /* * Enqueue a kernel run call. */ size_t globalThreads[] = {length / 4}; size_t localThreads[] = {GROUP_SIZE, 1}; if(localThreads[0] > maxWorkItemSizes[0] || localThreads[0] > maxWorkGroupSize) { std::cout<<"Unsupported: The current device does not support " <<"the execution domain specified. Try with a smaller size.\n"; return SDK_FAILURE; } status = clGetKernelWorkGroupInfo(kernel, devices[0], CL_KERNEL_LOCAL_MEM_SIZE, sizeof(cl_ulong), &usedLocalMemory, NULL); if(status != CL_SUCCESS) { std::cout << "clGetKernelWorkGroupInfo CL_KERNEL_LOCAL_MEM_SIZE failed."; return SDK_FAILURE; } std::cout << "Used local memory = " << usedLocalMemory << std::endl; std::cout << "availble local memory = " << totalLocalMemory << std::endl; status = clEnqueueNDRangeKernel( commandQueue, kernel, 1, NULL, globalThreads, localThreads, 0, NULL, NULL); if(status != CL_SUCCESS) { std::cout << "clEnqueueNDRangeKernel failed."; return SDK_FAILURE; } status = clFinish(commandQueue); if(status != CL_SUCCESS) { std::cout << "clFinish failed."; return SDK_FAILURE; } /* Enqueue readBuffer*/ status = clEnqueueReadBuffer( commandQueue, outputBuffer, CL_TRUE, 0, length * sizeof(cl_float), output, 0, NULL, &events[0]); if(status != CL_SUCCESS) { std::cout << "clEnqueueReadBuffer failed."; return SDK_FAILURE; } /* Wait for the read buffer to finish execution */ status = clWaitForEvents(1, &events[0]); if(status != CL_SUCCESS) { std::cout << "clWaitForEvents failed."; return SDK_FAILURE; } clReleaseEvent(events[0]); return SDK_SUCCESS; } int Reduction::setup() { if(setupReduction()!=SDK_SUCCESS) { return SDK_FAILURE; } if(setupCL()!=SDK_SUCCESS) { return SDK_FAILURE; } return SDK_SUCCESS; } int Reduction::run() { /* Arguments are set and execution call is enqueued on command buffer */ if(runCLKernels()!=SDK_SUCCESS) return SDK_FAILURE; return SDK_SUCCESS; } int Reduction::verifyResults() { if(true) { /* reference implementation * it overwrites the input array with the output */ bool passed = true; for(int i = 0; i < length / 4; i = i + 4) std::cout << output[i + 0] << " " << output[i + 1] << " " << output[i + 2] << " " << output[i + 3] << std::endl; } return SDK_SUCCESS; } int Reduction::cleanup() { /* Releases OpenCL resources (Context, Memory etc.) */ cl_int status; status = clReleaseKernel(kernel); if(status != CL_SUCCESS) { std::cout << "clReleaseKernel failed."; return SDK_FAILURE; } status = clReleaseProgram(program); if(status != CL_SUCCESS) { std::cout << "clReleaseProgram failed."; return SDK_FAILURE; } // status = clReleaseMemObject(inputBuffer); // if(status != CL_SUCCESS) // { // std::cout << "clReleaseMemObject failed."; //return SDK_FAILURE; // } status = clReleaseMemObject(outputBuffer); if(status != CL_SUCCESS) { std::cout << "clReleaseMemObject failed."; return SDK_FAILURE; } status = clReleaseCommandQueue(commandQueue); if(status != CL_SUCCESS) { std::cout << "clReleaseCommandQueue failed."; return SDK_FAILURE; } status = clReleaseContext(context); if(status != CL_SUCCESS) { std::cout << "clReleaseContext failed."; return SDK_FAILURE; } return SDK_SUCCESS; } Reduction::~Reduction() { /* release program resources (input memory etc.) */ //delete[] input; delete[] output; if(devices) { free(devices); devices = NULL; } if(maxWorkItemSizes) { free(maxWorkItemSizes); maxWorkItemSizes = NULL; } } int main(int argc, char * argv[]) { Reduction clReduction("OpenCL Reduction"); // Parse command line options // for(int i = 0; i < argc && argv; i++) { if(!argv[i]) continue; if(strstr(argv[i], "cpu")) { dType = CL_DEVICE_TYPE_CPU; std::cout << "Running on CL_DEVICE_TYPE_CPU" << std::endl; } else if(strstr(argv[i], "gpu")) { dType = CL_DEVICE_TYPE_GPU; std::cout << "Running on CL_DEVICE_TYPE_GPU" << std::endl; } } if(clReduction.setup()!=SDK_SUCCESS) return SDK_FAILURE; if(clReduction.run()!=SDK_SUCCESS) return SDK_FAILURE; if(clReduction.verifyResults()!=SDK_SUCCESS) return SDK_FAILURE; if(clReduction.cleanup()!=SDK_SUCCESS) return SDK_FAILURE; return SDK_SUCCESS; } Kernel file ******** __kernel void qsimKernel(__global float4* output) { size_t globalId = get_global_id(0); float4 a = (float4)(1.0f, 2.0f, 3.0f, 4.0f); float4 b = (float4)(4.0f, 3.0f, 2.0f, 1.0f); float4 n = cross(a, -b); //output[globalId] = n; output[globalId] = normalize(n); }