Une nouvelle version du portail de gestion des comptes externes sera mise en production lundi 09 août. Elle permettra d'allonger la validité d'un compte externe jusqu'à 3 ans. Pour plus de détails sur cette version consulter : https://doc-si.inria.fr/x/FCeS

Commit 5d8596f4 authored by BRAMAS Berenger's avatar BRAMAS Berenger
Browse files

update opencl starpu

parent 471bff07
......@@ -25,6 +25,14 @@ struct FEmptyOpenCLFilename{
template <class OriginalKernelClass, class KernelFilenameClass = FEmptyOpenCLFilename>
class FOpenCLDeviceWrapper {
protected:
struct Uptr9{
cl_mem ptrs[9];
};
struct size_t9{
size_t v[9];
};
static void SetKernelArgs(cl_kernel& kernel, const int pos){
}
template <class ParamClass, class... Args>
......@@ -83,7 +91,7 @@ public:
if(filename){
starpu_opencl_get_context (workerDevid, &context);
const int err = starpu_opencl_load_opencl_from_string(filename, &opencl_code, NULL);
const int err = starpu_opencl_load_opencl_from_string(filename, &opencl_code, "-cl-std=CL2.0 -cl-mad-enable -Werror");
if(err != CL_SUCCESS) STARPU_OPENCL_REPORT_ERROR(err);
FAssertLF( starpu_opencl_load_kernel(&kernel_bottomPassPerform, &queue_bottomPassPerform, &opencl_code, "FOpenCL__bottomPassPerform", workerDevid) == CL_SUCCESS);
......@@ -129,7 +137,12 @@ public:
void upwardPassPerform(cl_mem currentCellsPtr, size_t currentCellsSize, cl_mem subCellGroupsPtr[9], size_t subCellGroupsSize[9], int nbSubCellGroups, int idxLevel){
SetKernelArgs(kernel_upwardPassPerform, 0, &currentCellsPtr, &currentCellsSize, &subCellGroupsPtr, &subCellGroupsSize, &nbSubCellGroups, &idxLevel, &user_data);
Uptr9 ptrs;
memcpy(ptrs.ptrs, subCellGroupsPtr, sizeof(cl_mem)*9);
size_t9 sizes;
memcpy(sizes.v, subCellGroupsSize, sizeof(size_t)*9);
SetKernelArgs(kernel_upwardPassPerform, 0, &currentCellsPtr, &currentCellsSize, &ptrs, &sizes, &nbSubCellGroups, &idxLevel, &user_data);
size_t dim = 1;
const int err = clEnqueueNDRangeKernel(queue_upwardPassPerform, kernel_upwardPassPerform, 1, NULL, &dim, NULL, 0, NULL, NULL);
if (err != CL_SUCCESS) STARPU_OPENCL_REPORT_ERROR(err);
......@@ -164,8 +177,13 @@ public:
void downardPassPerform(cl_mem currentCellsPtr,
size_t currentCellsSize, cl_mem subCellGroupsPtr[9], size_t subCellGroupsSize[9], int nbSubCellGroups, int idxLevel){
Uptr9 ptrs;
memcpy(ptrs.ptrs, subCellGroupsPtr, sizeof(cl_mem)*9);
size_t9 sizes;
memcpy(sizes.v, subCellGroupsSize, sizeof(size_t)*9);
SetKernelArgs(kernel_downardPassPerform, 0, &currentCellsPtr,
&currentCellsSize, &subCellGroupsPtr, &subCellGroupsSize, &nbSubCellGroups, &idxLevel, &user_data);
&currentCellsSize, &ptrs, &sizes, &nbSubCellGroups, &idxLevel, &user_data);
size_t dim = 1;
const int err = clEnqueueNDRangeKernel(queue_downardPassPerform, kernel_downardPassPerform, 1, NULL, &dim, NULL, 0, NULL, NULL);
if (err != CL_SUCCESS) STARPU_OPENCL_REPORT_ERROR(err);
......
......@@ -44,7 +44,8 @@ int main(int argc, char* argv[]){
{"-bs"},
"The size of the block of the blocked tree"
};
FHelpDescribeAndExit(argc, argv, "Test the blocked tree by counting the particles.",
FHelpDescribeAndExit(argc, argv, "Test the blocked tree by counting the particles."
"Usually run with STARPU_NCPU=0 STARPU_NOPENCL=1 STARPU_OPENCL_ONLY_ON_CPUS=1 ./Tests/Release/testBlockedWithOpenCLAlgorithm",
FParameterDefinitions::OctreeHeight, FParameterDefinitions::NbThreads,
FParameterDefinitions::NbParticles, LocalOptionBlocSize);
// Initialize the types
......
Markdown is supported
0% or .
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment