Tests for the program compilation

It took one day to have this going on, but there were bugs in my
openSUSE LLVM setup. The code is fairly clean and straightforward,
thanks to LLVM and Clang which are really well done libraries.

9 files changed, 165 insertions, 45 deletions

diff --git a/cmake/modules/FindClang.cmake b/cmake/modules/FindClang.cmake
index 6d6beb6..b740937 100644
--- a/cmake/modules/FindClang.cmake
+++ b/cmake/modules/FindClang.cmake
@@ -13,6 +13,7 @@ ENDMACRO(FIND_AND_ADD_CLANG_LIB)
 set(CLANG_INCLUDE_DIRS ${CLANG_INCLUDE_DIRS} ${LLVM_INCLUDE_DIR})
 set(CLANG_INCLUDE_DIRS ${CLANG_INCLUDE_DIRS} ${CLANG_INCLUDE_DIR})
 
+FIND_AND_ADD_CLANG_LIB(clang)
 FIND_AND_ADD_CLANG_LIB(clangFrontend)
 FIND_AND_ADD_CLANG_LIB(clangDriver)
 FIND_AND_ADD_CLANG_LIB(clangCodeGen)
diff --git a/src/api/api_program.cpp b/src/api/api_program.cpp
index 2964bcf..af6209c 100644
--- a/src/api/api_program.cpp
+++ b/src/api/api_program.cpp
@@ -84,7 +84,7 @@ clCreateProgramWithBinary(cl_context            context,
     
     *errcode_ret = context->info(CL_CONTEXT_DEVICES, 
                                  context_num_devices * sizeof(cl_device_id), 
-                                 &context_devices, 0);
+                                 context_devices, 0);
     
     if (*errcode_ret != CL_SUCCESS)
         return 0;
@@ -93,7 +93,7 @@ clCreateProgramWithBinary(cl_context            context,
     {
         bool found = false;
         
-        for (int j=0; j<context_num_devices; ++i)
+        for (int j=0; j<context_num_devices; ++j)
         {
             if (device_list[i] == context_devices[j])
             {
@@ -193,7 +193,7 @@ clBuildProgram(cl_program           program,
     
     result = context->info(CL_CONTEXT_DEVICES, 
                                  context_num_devices * sizeof(cl_device_id), 
-                                 &context_devices, 0);
+                                 context_devices, 0);
     
     if (result != CL_SUCCESS)
         return result;
@@ -202,7 +202,7 @@ clBuildProgram(cl_program           program,
     {
         bool found = false;
         
-        for (int j=0; j<context_num_devices; ++i)
+        for (int j=0; j<context_num_devices; ++j)
         {
             if (device_list[i] == context_devices[j])
             {
diff --git a/src/core/compiler.cpp b/src/core/compiler.cpp
index afe005e..4c4bf35 100644
--- a/src/core/compiler.cpp
+++ b/src/core/compiler.cpp
@@ -8,9 +8,12 @@
 #include <clang/Frontend/LangStandard.h>
 #include <clang/CodeGen/CodeGenAction.h>
 #include <llvm/ADT/SmallVector.h>
+#include <llvm/Support/Host.h>
 #include <llvm/Module.h>
 #include <llvm/LLVMContext.h>
 
+#include <stdio.h>
+
 using namespace Coal;
 
 Compiler::Compiler(const char *options)
@@ -18,38 +21,13 @@ Compiler::Compiler(const char *options)
 {
     size_t len = std::strlen(options);
     
-    // Parse the user options
-    std::istringstream options_stream(options);
-    std::string token;
-    llvm::SmallVector<const char *, 16> argv;
-    
-    while (options_stream >> token)
-    {
-        char *arg = new char[token.size() + 1];
-        
-        std::strcpy(arg, token.c_str());
-        argv.push_back(arg);
-    }
-    
-    argv.push_back("program.cl");
-    
-    // Create the diagnostics engine
-    p_compiler.createDiagnostics(0 ,NULL);
-    
-    if (!p_compiler.hasDiagnostics())
-        return;
+    // Set codegen options
+    clang::CodeGenOptions &codegen_opts = p_compiler.getCodeGenOpts();
+    codegen_opts.DebugInfo = false;
+    codegen_opts.AsmVerbose = true;
     
-    // Create a compiler invocation
-    clang::CompilerInvocation &invocation = p_compiler.getInvocation();
-    invocation.setLangDefaults(clang::IK_OpenCL);
-    
-    clang::CompilerInvocation::CreateFromArgs(invocation, 
-                                              argv.data(), 
-                                              argv.data() + argv.size(), 
-                                              p_compiler.getDiagnostics());
-
     // Set diagnostic options
-    clang::DiagnosticOptions &diag_opts = invocation.getDiagnosticOpts();
+    clang::DiagnosticOptions &diag_opts = p_compiler.getDiagnosticOpts();
     diag_opts.Pedantic = true;
     diag_opts.ShowColumn = true;
     diag_opts.ShowLocation = true;
@@ -58,29 +36,54 @@ Compiler::Compiler(const char *options)
     diag_opts.ShowColors = false;
     diag_opts.ErrorLimit = 19;
     diag_opts.MessageLength = 80;
+    diag_opts.DumpBuildInformation = std::string();
+    diag_opts.DiagnosticLogFile = std::string();
     
     // Set frontend options
-    clang::FrontendOptions &frontend_opts = invocation.getFrontendOpts();
+    clang::FrontendOptions &frontend_opts = p_compiler.getFrontendOpts();
     frontend_opts.ProgramAction = clang::frontend::EmitLLVMOnly;
     frontend_opts.DisableFree = true;
+    frontend_opts.Inputs.push_back(std::make_pair(clang::IK_OpenCL, "program.cl"));
     
     // Set header search options
-    clang::HeaderSearchOptions &header_opts = invocation.getHeaderSearchOpts();
+    clang::HeaderSearchOptions &header_opts = p_compiler.getHeaderSearchOpts();
     header_opts.Verbose = true;
     header_opts.UseBuiltinIncludes = false;
     header_opts.UseStandardIncludes = false;
     header_opts.UseStandardCXXIncludes = false;
     
     // Set lang options
-    clang::LangOptions &lang_opts = invocation.getLangOpts();
+    clang::LangOptions &lang_opts = p_compiler.getLangOpts();
     lang_opts.NoBuiltin = true;
     lang_opts.OpenCL = true;
     
+    // Set target options
+    clang::TargetOptions &target_opts = p_compiler.getTargetOpts();
+    target_opts.Triple = llvm::sys::getHostTriple();
+    
     // Set preprocessor options
-    clang::PreprocessorOptions &prep_opts = invocation.getPreprocessorOpts();
+    clang::PreprocessorOptions &prep_opts = p_compiler.getPreprocessorOpts();
     //prep_opts.Includes.push_back("stdlib.h");
     //prep_opts.addRemappedFile("stdlib.h", ...);
     
+    clang::CompilerInvocation &invocation = p_compiler.getInvocation();
+    invocation.setLangDefaults(clang::IK_OpenCL);
+    
+    // Parse the user options
+    std::istringstream options_stream(options);
+    std::string token;
+    
+    while (options_stream >> token)
+    {
+        
+    }
+    
+    // Create the diagnostics engine
+    p_compiler.createDiagnostics(0, NULL);
+    
+    if (!p_compiler.hasDiagnostics())
+        return;
+    
     p_valid = true;
 }
 
@@ -89,6 +92,11 @@ Compiler::~Compiler()
     
 }
 
+bool Compiler::valid() const
+{
+    return p_valid;
+}
+
 llvm::Module *Compiler::compile(llvm::MemoryBuffer *source)
 {
     // Feed the compiler with source
@@ -103,13 +111,15 @@ llvm::Module *Compiler::compile(llvm::MemoryBuffer *source)
         new clang::EmitLLVMOnlyAction(new llvm::LLVMContext)
     );
     
-   if (!p_compiler.ExecuteAction(*act))
-      return 0;
+    if (!p_compiler.ExecuteAction(*act))   
+    {
+        return 0;
+    }
 
     module = act->takeModule();
     
     // Cleanup
     prep_opts.eraseRemappedFile(prep_opts.remapped_file_buffer_end());
    
-   return module;
+    return module;
 }
\ No newline at end of file
diff --git a/src/core/program.cpp b/src/core/program.cpp
index c9a7208..0ab1625 100644
--- a/src/core/program.cpp
+++ b/src/core/program.cpp
@@ -113,6 +113,14 @@ cl_int Program::build(const char *options,
     {
         Compiler *compiler = new Compiler(options);
         
+        if (!compiler->valid())
+        {
+            if (pfn_notify)
+                pfn_notify((cl_program)this, user_data);
+            
+            return CL_BUILD_PROGRAM_FAILURE;
+        }
+        
         const llvm::StringRef s_data(p_source);
         const llvm::StringRef s_name("<source>");
         
diff --git a/tests/CMakeLists.txt b/tests/CMakeLists.txt
index ce238e1..3cd4a58 100644
--- a/tests/CMakeLists.txt
+++ b/tests/CMakeLists.txt
@@ -9,6 +9,7 @@ set(OPENCL_TESTS_SOURCE
     test_commandqueue.cpp
     test_mem.cpp
     test_kernel.cpp
+    test_program.cpp
     )
 
 add_executable(tests ${OPENCL_TESTS_SOURCE})
@@ -24,3 +25,4 @@ OPENCL_TEST(tests context)
 OPENCL_TEST(tests commandqueue)
 OPENCL_TEST(tests mem)
 OPENCL_TEST(tests kernel)
+OPENCL_TEST(tests program)
diff --git a/tests/test_mem.cpp b/tests/test_mem.cpp
index 9a965fa..a15f2d9 100644
--- a/tests/test_mem.cpp
+++ b/tests/test_mem.cpp
@@ -162,10 +162,10 @@ START_TEST (test_read_write_subbuf)
     cl_int result;
     char s[] = "Hello, Denis !";
     
-    cl_buffer_region create_info = {    // "Hello, [denis] !"
-        .origin = 7,
-        .size = 5
-    };
+    cl_buffer_region create_info;
+    
+    create_info.origin = 7;      // "Hello, [denis] !"
+    create_info.size = 5;
     
     result = clGetDeviceIDs(0, CL_DEVICE_TYPE_CPU, 1, &device, 0);
     fail_if(
diff --git a/tests/test_program.cpp b/tests/test_program.cpp
new file mode 100644
index 0000000..30ee8c9
--- /dev/null
+++ b/tests/test_program.cpp
@@ -0,0 +1,80 @@
+#include "test_program.h"
+#include "CL/cl.h"
+
+#include <stdio.h>
+
+const char program_source[] =
+    "#define __global __attribute__((address_space(1)))\n"
+    "\n"
+    "__kernel void test(__global float *a, __global float *b, int n) {\n"
+    "   int i;\n"
+    "\n"
+    "   for (i=0; i<n; i++) {\n"
+    "       a[i] = 3.1415926f * b[0] * b[0];\n"
+    "   }\n"
+    "}\n";
+
+START_TEST (test_create_program)
+{
+    cl_platform_id platform = 0;
+    cl_device_id device;
+    cl_context ctx;
+    cl_program program;
+    cl_int result;
+    
+    const char *src = program_source;
+    size_t program_len = sizeof(program_source);
+    
+    result = clGetDeviceIDs(platform, CL_DEVICE_TYPE_DEFAULT, 1, &device, 0);
+    fail_if(
+        result != CL_SUCCESS,
+        "unable to get the default device"
+    );
+    
+    ctx = clCreateContext(0, 1, &device, 0, 0, &result);
+    fail_if(
+        result != CL_SUCCESS || ctx == 0,
+        "unable to create a valid context"
+    );
+    
+    program = clCreateProgramWithSource(ctx, 0, &src, 0, &result);
+    fail_if(
+        result != CL_INVALID_VALUE,
+        "count cannot be 0"
+    );
+    
+    program = clCreateProgramWithSource(ctx, 1, 0, 0, &result);
+    fail_if(
+        result != CL_INVALID_VALUE,
+        "strings cannot be NULL"
+    );
+    
+    program = clCreateProgramWithSource(ctx, 1, &src, &program_len,
+                                        &result);
+    fail_if(
+        result != CL_SUCCESS,
+        "cannot create a program from source with sane arguments"
+    );
+    
+    clReleaseProgram(program); // Sorry
+    
+    program = clCreateProgramWithSource(ctx, 1, &src, 0, &result);
+    fail_if(
+        result != CL_SUCCESS,
+        "lengths can be NULL and it must work"
+    );
+    
+    result = clBuildProgram(program, 1, &device, "", 0, 0);
+    printf("result : %i\n", result);
+    
+    clReleaseContext(ctx);
+}
+END_TEST
+
+TCase *cl_program_tcase_create(void)
+{
+    TCase *tc = NULL;
+    tc = tcase_create("program");
+    tcase_add_test(tc, test_create_program);
+    return tc;
+}
diff --git a/tests/test_program.h b/tests/test_program.h
new file mode 100644
index 0000000..6a62329
--- /dev/null
+++ b/tests/test_program.h
@@ -0,0 +1,17 @@
+#ifndef __UTEST_PROGRAM__
+#define __UTEST_PROGRAM__
+
+#include <check.h>
+
+#ifdef __cplusplus
+extern "C" {
+#endif
+
+TCase *cl_program_tcase_create(void);
+
+#ifdef __cplusplus
+}
+#endif
+
+#endif
+
diff --git a/tests/tests.c b/tests/tests.c
index 45d3441..3a37f39 100644
--- a/tests/tests.c
+++ b/tests/tests.c
@@ -4,6 +4,7 @@
 #include "test_commandqueue.h"
 #include "test_mem.h"
 #include "test_kernel.h"
+#include "test_program.h"
 
 #include <stdlib.h>
 #include <stdio.h>
@@ -30,6 +31,7 @@ int main(int argc, char **argv)
     TESTSUITE(commandqueue, "commandqueue");
     TESTSUITE(mem, "mem");
     TESTSUITE(kernel, "kernel");
+    TESTSUITE(program, "program");
 
     if (s == NULL) {
         printf("test case %s does not exist", argv[1]);