diff --git a/repos/dde_uos-intel-gpgpu/src/uos-intel-gpgpu b/repos/dde_uos-intel-gpgpu/src/uos-intel-gpgpu index 113e28173f..06c1b3f205 160000 --- a/repos/dde_uos-intel-gpgpu/src/uos-intel-gpgpu +++ b/repos/dde_uos-intel-gpgpu/src/uos-intel-gpgpu @@ -1 +1 @@ -Subproject commit 113e28173fa9d3b0d3a40be3678e2d2c6f495f3e +Subproject commit 06c1b3f205ae5806b21d7a3ae795d723c9e8e46b diff --git a/repos/hello_gpgpu/src/hello_gpgpu/CL/cl.cc b/repos/hello_gpgpu/src/hello_gpgpu/CL/cl.cc index 4429d55eb1..ae88ee6b59 100644 --- a/repos/hello_gpgpu/src/hello_gpgpu/CL/cl.cc +++ b/repos/hello_gpgpu/src/hello_gpgpu/CL/cl.cc @@ -622,7 +622,7 @@ clCreateProgramWithBinary(cl_context context, return NULL; } - cl_program p = (cl_program)g_cl_genode->alloc(sizeof(_cl_program)); + cl_program p = (cl_program)g_cl_genode->alloc(sizeof(struct _cl_program)); p->binary = (uint8_t*)binaries[0]; p->size = lengths[0]; @@ -1428,7 +1428,7 @@ clEnqueueNDRangeKernel(cl_command_queue command_queue, } // create copy of kernel (vm and driver should not modify the same kernel) - kernel_config* kcopy = new(g_cl_genode->getAlloc()) kernel_config(); + struct kernel_config* kcopy = new(g_cl_genode->getAlloc()) kernel_config(); *kcopy = *kc; // also copy buff configs diff --git a/repos/hello_gpgpu/src/hello_gpgpu/benchmark/doitgen/doitgen.cl b/repos/hello_gpgpu/src/hello_gpgpu/benchmark/doitgen/doitgen.cl index 1a377952cf..3c239072b9 100644 --- a/repos/hello_gpgpu/src/hello_gpgpu/benchmark/doitgen/doitgen.cl +++ b/repos/hello_gpgpu/src/hello_gpgpu/benchmark/doitgen/doitgen.cl @@ -30,7 +30,6 @@ __kernel void doitgen_kernel1(int nr, int nq, int np, __global DATA_TYPE *A, __g for (int s = 0; s < np; s++) { sum[r * (nq * np) + q * np + p] = sum[r * (nq * np) + q * np + p] + A[r * (nq * np) + q * np + s] * C4[s * np + p]; - //sum[r * (nq * np) + q * np + p] = (float)q; } } } diff --git a/repos/hello_gpgpu/src/hello_gpgpu/main.cc b/repos/hello_gpgpu/src/hello_gpgpu/main.cc index 235bcee4e8..423014521e 100644 --- a/repos/hello_gpgpu/src/hello_gpgpu/main.cc +++ b/repos/hello_gpgpu/src/hello_gpgpu/main.cc @@ -50,7 +50,7 @@ void testvm_construct(Genode::Env &env) Genode::log("===Run bicg==="); ns_bicg::main(0, 0); Genode::log("===Run doitgen==="); - ns_doitgen::main(0, 0); // Number of misses: 2080768 + GPU kernel1 takes very long + ns_doitgen::main(0, 0); Genode::log("===Run gemm==="); ns_gemm::main(0, 0); Genode::log("===Run gemver==="); @@ -70,7 +70,7 @@ void testvm_construct(Genode::Env &env) ns_lu::main(0, 0); // Non-Matching CPU-GPU Outputs Beyond Error Threshold of 0.05 Percent: 516 Genode::log("===Run correlation==="); - ns_correlation::main(0, 0); // Non-Matching CPU-GPU Outputs Beyond Error Threshold of 1.05 Percent: 1 + ns_correlation::main(0, 0); Genode::log("===Run covariance==="); ns_covariance::main(0, 0); diff --git a/repos/hello_gpgpu/src/hello_gpgpu/polybench.cc b/repos/hello_gpgpu/src/hello_gpgpu/polybench.cc index 1e8c16350d..88b5c3833a 100644 --- a/repos/hello_gpgpu/src/hello_gpgpu/polybench.cc +++ b/repos/hello_gpgpu/src/hello_gpgpu/polybench.cc @@ -30,6 +30,7 @@ #endif #define POLYBENCH_TIME 1 +#define POLYBENCH_CYCLE_ACCURATE_TIMER int polybench_papi_counters_threadid = POLYBENCH_THREAD_MONITOR; double polybench_program_total_flops = 0;