static void createFPFnStub(Function *F, Module *M, FPParamVariant PV,
const MipsSubtarget &Subtarget ) {
bool PicMode = Subtarget.getRelocationModel() == Reloc::PIC_;
bool LE = Subtarget.isLittle();
LLVMContext &Context = M->getContext();
std::string Name = F->getName();
std::string SectionName = ".mips16.fn." + Name;
std::string StubName = "__fn_stub_" + Name;
std::string LocalName = "$$__fn_local_" + Name;
Function *FStub = Function::Create
(F->getFunctionType(),
Function::InternalLinkage, StubName, M);
FStub->addFnAttr("mips16_fp_stub");
FStub->addFnAttr(llvm::Attribute::Naked);
FStub->addFnAttr(llvm::Attribute::NoUnwind);
FStub->addFnAttr(llvm::Attribute::NoInline);
FStub->addFnAttr("nomips16");
FStub->setSection(SectionName);
BasicBlock *BB = BasicBlock::Create(Context, "entry", FStub);
InlineAsmHelper IAH(Context, BB);
if (PicMode) {
IAH.Out(".set noreorder");
IAH.Out(".cpload $$25");
IAH.Out(".set reorder");
IAH.Out(".reloc 0,R_MIPS_NONE," + Name);
IAH.Out("la $$25," + LocalName);
}
else {
IAH.Out("la $$25," + Name);
}
swapFPIntParams(PV, M, IAH, LE, false);
IAH.Out("jr $$25");
IAH.Out(LocalName + " = " + Name);
new UnreachableInst(FStub->getContext(), BB);
}
//
// Make sure that we know we already need a stub for this function.
// Having called needsFPHelperFromSig
//
static void assureFPCallStub(Function &F, Module *M,
const MipsSubtarget &Subtarget){
// for now we only need them for static relocation
if (Subtarget.getRelocationModel() == Reloc::PIC_)
return;
LLVMContext &Context = M->getContext();
bool LE = Subtarget.isLittle();
std::string Name = F.getName();
std::string SectionName = ".mips16.call.fp." + Name;
std::string StubName = "__call_stub_fp_" + Name;
//
// see if we already have the stub
//
Function *FStub = M->getFunction(StubName);
if (FStub && !FStub->isDeclaration()) return;
FStub = Function::Create(F.getFunctionType(),
Function::InternalLinkage, StubName, M);
FStub->addFnAttr("mips16_fp_stub");
FStub->addFnAttr(llvm::Attribute::Naked);
FStub->addFnAttr(llvm::Attribute::NoInline);
FStub->addFnAttr(llvm::Attribute::NoUnwind);
FStub->addFnAttr("nomips16");
FStub->setSection(SectionName);
BasicBlock *BB = BasicBlock::Create(Context, "entry", FStub);
InlineAsmHelper IAH(Context, BB);
IAH.Out(".set reorder");
FPReturnVariant RV = whichFPReturnVariant(FStub->getReturnType());
FPParamVariant PV = whichFPParamVariantNeeded(F);
swapFPIntParams(PV, M, IAH, LE, true);
if (RV != NoFPRet) {
IAH.Out("move $$18, $$31");
IAH.Out("jal " + Name);
} else {
IAH.Out("lui $$25,%hi(" + Name + ")");
IAH.Out("addiu $$25,$$25,%lo(" + Name + ")" );
}
switch (RV) {
case FRet:
IAH.Out("mfc1 $$2,$$f0");
break;
case DRet:
if (LE) {
IAH.Out("mfc1 $$2,$$f0");
IAH.Out("mfc1 $$3,$$f1");
} else {
IAH.Out("mfc1 $$3,$$f0");
IAH.Out("mfc1 $$2,$$f1");
}
break;
case CFRet:
if (LE) {
IAH.Out("mfc1 $$2,$$f0");
IAH.Out("mfc1 $$3,$$f2");
} else {
IAH.Out("mfc1 $$3,$$f0");
IAH.Out("mfc1 $$3,$$f2");
}
break;
case CDRet:
if (LE) {
IAH.Out("mfc1 $$4,$$f2");
IAH.Out("mfc1 $$5,$$f3");
IAH.Out("mfc1 $$2,$$f0");
IAH.Out("mfc1 $$3,$$f1");
} else {
IAH.Out("mfc1 $$5,$$f2");
IAH.Out("mfc1 $$4,$$f3");
IAH.Out("mfc1 $$3,$$f0");
IAH.Out("mfc1 $$2,$$f1");
}
break;
case NoFPRet:
break;
}
if (RV != NoFPRet)
IAH.Out("jr $$18");
else
IAH.Out("jr $$25");
new UnreachableInst(Context, BB);
}
int main(int argc, char* argv[])
{
/* glFinish(); */
/* status = clEnqueueAcquireGLObjects(commandQueue, 1, &cl_tex_mem, */
/* 0,NULL,NULL ); */
/* status = clEnqueueNDRangeKernel(commandQueue, tex_kernel, 2, NULL, */
/* tex_globalWorkSize, */
/* tex_localWorkSize, */
/* 0, NULL, NULL); */
/* clFinish(commandQueue); */
/* status = clEnqueueReleaseGLObjects(commandQueue, 1, &cl_tex_mem, 0, NULL, NULL ); */
/*Step1: Getting platforms and choose an available one.*/
initFns();
/* printf("HELLO\n"); */
cl_uint numPlatforms;//the NO. of platforms
cl_platform_id platform = NULL;//the chosen platform
IAH();
PP(clGetPlatformIDs);
cl_int status = clGetPlatformIDs(0, NULL, &numPlatforms);
if (status != CL_SUCCESS)
{
printf("Error: Getting platforms!\n");
return 1;
}
/*For clarity, choose the first available platform. */
if(numPlatforms > 0)
{
cl_platform_id* platforms = (cl_platform_id* )malloc(numPlatforms* sizeof(cl_platform_id));
IAH();
status = clGetPlatformIDs(numPlatforms, platforms, NULL);
platform = platforms[0];
free(platforms);
}
/*Step 2:Query the platform and choose the first GPU device if has one.Otherwise use the CPU as device.*/
cl_uint numDevices = 0;
cl_device_id *devices;
IAH();
PP(clGetDeviceIDs);
status = clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, 0, NULL, &numDevices);
if (numDevices == 0) //no GPU available.
{
printf("No GPU device available.\n");
printf("Choose CPU as default device.\n");
IAH();
status = clGetDeviceIDs(platform, CL_DEVICE_TYPE_CPU, 0, NULL, &numDevices);
CHECK_STATUS(status, "clGetDeviceIDs");
devices = (cl_device_id*)malloc(numDevices * sizeof(cl_device_id));
IAH();
status = clGetDeviceIDs(platform, CL_DEVICE_TYPE_CPU, numDevices, devices, NULL);
CHECK_STATUS(status, "clGetDeviceIDs");
}
else
{
devices = (cl_device_id*)malloc(numDevices * sizeof(cl_device_id));
IAH();
status = clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, numDevices, devices, NULL);
CHECK_STATUS(status, "clGetDeviceIDs");
}
/*Step 3: Create context.*/
IAH();
cl_context context = clCreateContext(NULL,1, devices,NULL,NULL,NULL);
CHECK_STATUS(status, "clCreateContext");
/*Step 4: Creating command queue associate with the context.*/
IAH();
cl_command_queue commandQueue = clCreateCommandQueue(context, devices[0], 0, &status);
CHECK_STATUS(status, "clCreateCommandQueue");
/*Step 5: Create program object */
//const char *filename = "HelloWorld_Kernel.cl";
//string sourceStr;
//status = convertToString(filename, sourceStr);
/* const char *source = KERNEL_SRC;//sourceStr.c_str(); */
const char *source = KERNEL_SRC2;//sourceStr.c_str();
size_t sourceSize[] = {strlen(source)};
IAH();
cl_program program = clCreateProgramWithSource(context, 1, &source, sourceSize, &status);
CHECK_STATUS(status, "clCreateProgramWithSource");
/*Step 6: Build program. */
IAH();
status=clBuildProgram(program, 1,devices,NULL,NULL,NULL);
CHECK_STATUS(status, "clBuildProgram");
/* printf("HELLO\n"); */
/*Step 7: Initial input,output for the host and create memory objects for the kernel*/
/* const char* input = "GdkknVnqkc"; */
char input[ARRAY_SIZE - 1];
/* size_t strlength = strlen(input); */
size_t strlength = ARRAY_SIZE - 1;
int i;
for (i = 0; i < strlength - 1; i++) {
input[i] = '1';
}
input[strlength - 1] = '\0';
printf("input string: %s\n",input);
char *output = (char*) malloc(strlength + 1);
IAH();
cl_mem inputBuffer = clCreateBuffer(context, CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR, (strlength + 1) * sizeof(char),(void *) input, &status);
CHECK_STATUS(status, "clCreateBuffer");
IAH();
cl_mem outputBuffer = clCreateBuffer(context, CL_MEM_WRITE_ONLY , (strlength + 1) * sizeof(char), NULL, &status);
CHECK_STATUS(status, "clCreateBuffer");
/*Step 8: Create kernel object */
IAH();
cl_kernel kernel = clCreateKernel(program,"helloworld", &status);
CHECK_STATUS(status, "clCreateKernel");
/*Step 9: Sets Kernel arguments.*/
/* IAH(); */
/* status = clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *)&inputBuffer); */
/* IAH(); */
/* status = clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *)&outputBuffer); */
/* ciErr1 |= clSetKernelArg(ckKernel, 3, sizeof(cl_int), (void*)&iNumElements); */
/* printf("HELLO\n",input); */
/*Step 9: Sets Kernel arguments.*/
IAH();
int strsize = strlength;
status = clSetKernelArg(kernel, 0, sizeof(cl_int), (void *)&strsize);
CHECK_STATUS(status, "clSetKernelArg");
IAH();
status = clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *)&inputBuffer);
CHECK_STATUS(status, "clSetKernelArg");
IAH();
status = clSetKernelArg(kernel, 2, sizeof(cl_mem), (void *)&outputBuffer);
CHECK_STATUS(status, "clSetKernelArg");
/*Step 10: Running the kernel.*/
/* size_t global_work_size[1] = {strlength}; */
size_t global_work_size[1] = {NUM_WORKERS};
IAH();
status = clEnqueueNDRangeKernel(commandQueue, kernel, 1, NULL, global_work_size, NULL, 0, NULL, NULL);
CHECK_STATUS(status, "clEnqueueNDRangeKernel");
/*Step 11: Read the cout put back to host memory.*/
IAH();
status = clEnqueueReadBuffer(commandQueue, outputBuffer, CL_TRUE, 0, strlength * sizeof(char), output, 0, NULL, NULL);
CHECK_STATUS(status, "clEnqueueReadBuffer");
output[strlength] = '\0';//Add the terminal character to the end of output.
printf("output string: %s\n",output);
for (i = 0; i < strlength - 1; i++) {
if (output[i] != '2') {
printf("ERROR: expected '2' but saw %c at output[%i]\n", output[i], i);
exit(EXIT_FAILURE);
}
input[i] = '1';
}
/* Print the maximum allocatable memory size for the device.
*/
cl_bool device_available = CL_FALSE;
status = clGetDeviceInfo(devices[0], CL_DEVICE_AVAILABLE, sizeof(cl_bool), &device_available, NULL);
CHECK_STATUS(status, "clGetDeviceInfo");
if (device_available != CL_TRUE)
{
printf("Error: Device %i is not available\n", 0);
return EXIT_FAILURE;
}
cl_ulong device_max_mem_alloc_size = 0;
status = clGetDeviceInfo(devices[0], CL_DEVICE_MAX_MEM_ALLOC_SIZE, sizeof(cl_ulong), &device_max_mem_alloc_size, NULL);
CHECK_STATUS(status, "clGetDeviceInfo");
cl_ulong device_global_mem_size = 0;
status = clGetDeviceInfo(devices[0], CL_DEVICE_GLOBAL_MEM_SIZE, sizeof(cl_ulong), &device_max_mem_alloc_size, NULL);
CHECK_STATUS(status, "clGetDeviceInfo");
printf("The max allocateable memory size is %i\n", device_max_mem_alloc_size);
printf("The size of global device memory is %i\n", device_global_mem_size);
/*Step 12: Clean the resources.*/
IAH();
status = clReleaseKernel(kernel);//*Release kernel.
CHECK_STATUS(status, "clReleaseKernel");
IAH();
status = clReleaseProgram(program); //Release the program object.
CHECK_STATUS(status, "clReleaseProgram");
IAH();
status = clReleaseMemObject(inputBuffer);//Release mem object.
CHECK_STATUS(status, "clReleaseMemObject");
IAH();
status = clReleaseMemObject(outputBuffer);
CHECK_STATUS(status, "clReleaseMemObject");
IAH();
status = clReleaseCommandQueue(commandQueue);//Release Command queue.
CHECK_STATUS(status, "clReleaseCommandQueue");
IAH();
status = clReleaseContext(context);//Release context.
CHECK_STATUS(status, "clReleaseContext");
IAH();
if (output != NULL)
{
IAH();
free(output);
output = NULL;
}
if (devices != NULL)
{
IAH();
free(devices);
devices = NULL;
}
IAH();
return 0;
}
