I am trying to learn OpenCL basics.
I thought this code in a kernel:
out[ 1 & ((a+b)==(b+a)) ] = (char)1;
Would generate the same results as:
out[ 1 ] = (char)1;
Can anybody tell me why it generates different result?
Here is the full source code for my kernel:
#pragma OPENCL EXTENSION cl_khr_byte_addressable_store : enable
__kernel void hello(__global char * out)
{
size_t tid = get_global_id(0);
int a = tid & 0xff;
int b = (tid >> 8) & 0xff;
out[ 1 & ((a+b)==(b+a)) ] = (char)1;
}
If I replace that last statement with "out[1]=(char)1;" then "out[0]" will not be written. But the kernel above will write "1" to out[0].
EDIT: This is my C++ code:
#include <utility>
#define __NO_STD_VECTOR // Use cl::vector instead of STL version
#include <CL/cl.hpp>
#include <cstdio>
#include <cstdlib>
#include <fstream>
#include <iostream>
#include <string>
#include <iterator>
inline void checkErr(cl_int err, const char * name)
{
if (err != CL_SUCCESS) {
std::cerr << "ERROR: " << name
<< " (" << err << ")" << std::endl;
std::exit(EXIT_FAILURE);
}
}
int main()
{
cl_int err;
cl::vector< cl::Platform > platformList;
cl::Platform::get(&platformList);
checkErr(platformList.size()!=0 ? CL_SUCCESS : -1, "cl::Platform::get");
std::cerr << "Platform number is: " << platformList.size() << std::endl;
std::string platformVendor;
platformList[0].getInfo((cl_platform_info)CL_PLATFORM_VENDOR, &platformVendor);
std::cerr << "Platform is by: " << platformVendor << "\n";
cl_context_properties cprops[3] =
{CL_CONTEXT_PLATFORM, (cl_context_properties)(platformList[0])(), 0};
cl::Context context(
CL_DEVICE_TYPE_GPU,
cprops,
NULL,
NULL,
&err);
checkErr(err, "Conext::Context()");
unsigned char outH[2] = {0};
cl::Buffer outCL(
context,
CL_MEM_WRITE_ONLY | CL_MEM_USE_HOST_PTR,
sizeof(outH),
outH,
&err);
checkErr(err, "Buffer::Buffer()");
cl::vector<cl::Device> devices;
devices = context.getInfo<CL_CONTEXT_DEVICES>();
checkErr(
devices.size() > 0 ? CL_SUCCESS : -1, "devices.size() > 0");
std::ifstream file("condition1.cl");
checkErr(file.is_open() ? CL_SUCCESS:-1, "condition1.cl");
const std::string prog(std::istreambuf_iterator<char>(file), (std::istreambuf_iterator<char>()));
cl::Program::Sources source(1, std::make_pair(prog.c_str(), prog.length()+1));
cl::Program program(context, source);
err = program.build(devices,"");
checkErr(err, "Program::build()");
cl::Kernel kernel(program, "hello", &err);
checkErr(err, "Kernel::Kernel()");
err = kernel.setArg(0, outCL);
checkErr(err, "Kernel::setArg()");
cl::CommandQueue queue(context, devices[0], 0, &err);
checkErr(err, "CommandQueue::CommandQueue()");
cl::Event event;
err = queue.enqueueNDRangeKernel(
kernel,
cl::NullRange,
cl::NDRange(65536),
cl::NDRange(1, 1),
NULL,
&event);
checkErr(err, "ComamndQueue::enqueueNDRangeKernel()");
event.wait();
err = queue.enqueueReadBuffer(
outCL,
CL_TRUE,
0,
sizeof(outH),
outH);
checkErr(err, "ComamndQueue::enqueueReadBuffer()");
for (int i = 0; i < sizeof(outH); i++)
std::cout << (int)outH[i] << " ";
std::string str;
std::getline(std::cin, str);
return EXIT_SUCCESS;
}
EDIT2:
It is not caused by "undefined behaviour" at least. I get weird results for this kernel code too:
char result = 0;
for (int a = 0; a < 2; a++) {
for (int b = 0; b < 2; b++) {
if ((a+b) != (b+a))
result |= (1 << (a+2*b));
}
}
I would expect that "result" will get the value 0. But it gets the value 6. If I change the != to an == then result gets the value 9. So in this code, when "a" is not equal with "b" then (a+b) does not equal (b+a).
If I change the code and set a known value for either "a" or "b" then the result will be 0 as I expect. For example:
char result = 0;
int a = 1;
/*for (int a = 0; a < 2; a++)*/ {
for (int b = 0; b < 2; b++) {
if ((a+b) != (b+a))
result |= (1 << (a+2*b));
}
}