Skip to content

Fixes the out of host memory issue (#17) #68

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Open
vinays6 wants to merge 1 commit into
base: main
Choose a base branch
from
Open

Fixes the out of host memory issue (#17) #68

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
232 changes: 116 additions & 116 deletions AutoDock-Vina-GPU-2.1/lib/cache.cpp
View file
Open in desktop

Large diffs are not rendered by default.

406 changes: 205 additions & 201 deletions QuickVina-W-GPU-2.1/Kernel1_Opt.bin
View file
Open in desktop

Large diffs are not rendered by default.

76,513 changes: 38,278 additions & 38,235 deletions QuickVina-W-GPU-2.1/Kernel2_Opt.bin
View file
Open in desktop

Large diffs are not rendered by default.

6 changes: 3 additions & 3 deletions QuickVina-W-GPU-2.1/Makefile
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,7 +1,7 @@
# Need to be modified according to different users
WORK_DIR=/home/shidi/Vina-GPU-2.1/QuickVina-W-GPU-2.1
BOOST_LIB_PATH=/home/shidi/boost_1_84_0
OPENCL_LIB_PATH=/usr/local/cuda
WORK_DIR=/home/vinay/code/Vina-GPU-2.1/QuickVina-W-GPU-2.1
BOOST_LIB_PATH=/home/vinay/boost_1_77_0
OPENCL_LIB_PATH=/usr/include/CL
OPENCL_VERSION=-DOPENCL_3_0
GPU_PLATFORM=-DNVIDIA_PLATFORM
DOCKING_BOX_SIZE=-DSMALL_BOX
Expand Down
54 changes: 28 additions & 26 deletions QuickVina-W-GPU-2.1/OpenCL/src/kernels/kernel2.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -30,26 +30,27 @@ bool metropolis_accept(float old_f, float new_f, float temperature, float n) {

__kernel
void kernel2(
const __global output_type_cl* ric,
__global m_cl* mg,
__constant pre_cl* pre,
__constant grids_cl* grids,
__constant random_maps* random_maps,
__global ligand_atom_coords_cl* coords,
__global output_type_cl* results,
const __global mis_cl* mis,
const int torsion_size,
const int search_depth,
const int max_bfgs_steps,
float center_x,
float center_y,
float center_z,
float size_x,
float size_y,
float size_z,
__global ele_cl* global_ptr,
__global int* count_id,
const int rilc_bfgs_enable
const __global output_type_cl* ric,
__global m_cl* mg,
__constant pre_cl* pre,
__constant grids_cl* grids,
__constant random_maps* random_maps,
__global ligand_atom_coords_cl* coords,
__global output_type_cl* results,
const __global mis_cl* mis,
const int torsion_size,
const int search_depth,
const int max_bfgs_steps,
float center_x,
float center_y,
float center_z,
float size_x,
float size_y,
float size_z,
__global ele_cl* global_ptr,
__global int* count_id,
const int rilc_bfgs_enable,
__global individual_container* all_lists
) {
int gx = get_global_id(0);
int gy = get_global_id(1);
Expand All @@ -76,8 +77,9 @@ void kernel2(

output_type_cl candidate;

__private individual_container list;
circularvisited_init_cl(&list);
// Use global memory for list
__global individual_container* list = &all_lists[gl];
circularvisited_init_cl(list);

global_container g_container;
global_init_cl(&g_container);
Expand Down Expand Up @@ -111,7 +113,7 @@ void kernel2(
max_bfgs_steps,
true,
&g_container,
&list,
list,
search_depth,
&origin,
&boxsize,
Expand All @@ -129,7 +131,7 @@ void kernel2(
max_bfgs_steps,
true,
&g_container,
&list,
list,
search_depth,
&origin,
&boxsize,
Expand Down Expand Up @@ -159,7 +161,7 @@ void kernel2(
max_bfgs_steps,
false,
&g_container,
&list,
list,
search_depth,
&origin,
&boxsize,
Expand All @@ -177,7 +179,7 @@ void kernel2(
max_bfgs_steps,
false,
&g_container,
&list,
list,
search_depth,
&origin,
&boxsize,
Expand Down
4 changes: 2 additions & 2 deletions QuickVina-W-GPU-2.1/OpenCL/src/kernels/quasi_newton.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -672,7 +672,7 @@ void bfgs( output_type_cl* x,
const int max_bfgs_steps,
bool global_flag,
global_container* g_container,
__private individual_container* circularvisited,
__global individual_container* circularvisited,
int search_depth,
vec3_cl* origin,
vec3_cl* box_size,
Expand Down Expand Up @@ -903,7 +903,7 @@ void rilc_bfgs( output_type_cl* x,
const int max_bfgs_steps,
bool global_flag,
global_container* g_container,
__private individual_container* circularvisited,
__global individual_container* circularvisited,
int search_depth,
vec3_cl* origin,
vec3_cl* box_size,
Expand Down
18 changes: 10 additions & 8 deletions QuickVina-W-GPU-2.1/OpenCL/src/wrapcl.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -296,14 +296,16 @@ void SetupDevice(cl_platform_id* platforms, cl_device_id** devices, cl_int gpu_p
err = clGetDeviceInfo((*devices)[i], CL_DEVICE_NAME, 0, NULL, &device_name_size); checkErr(err);
char* device_name = (char*)malloc(sizeof(char) * device_name_size);
err = clGetDeviceInfo((*devices)[i], CL_DEVICE_NAME, device_name_size, device_name, NULL);
printf("\nGPU Device: %s", device_name);

#ifdef DISPLAY_ADDITION_INFO
err = clGetDeviceInfo((*devices)[i], CL_DEVICE_GLOBAL_MEM_SIZE, sizeof(cl_ulong), &mem_size, NULL);
printf("Platform %d global memory size:%f GB\n", N, (double)mem_size/1000000000);
err = clGetDeviceInfo((*devices)[i], CL_DEVICE_LOCAL_MEM_SIZE, sizeof(cl_ulong), &mem_size, NULL);
printf("Platform %d local memory size:%f KB\n", N, (double)mem_size / 1000);
#endif
cl_device_type dtype;
clGetDeviceInfo((*devices)[i], CL_DEVICE_TYPE, sizeof(dtype), &dtype, NULL);
cl_uint vendor_id;
clGetDeviceInfo((*devices)[i], CL_DEVICE_VENDOR_ID, sizeof(vendor_id), &vendor_id, NULL);
cl_ulong global_mem = 0;
clGetDeviceInfo((*devices)[i], CL_DEVICE_GLOBAL_MEM_SIZE, sizeof(global_mem), &global_mem, NULL);
printf("\nOpenCL Device %d: %s\n Type: %s\n Vendor ID: %u\n Global Mem: %.2f MB\n", i, device_name,
(dtype == CL_DEVICE_TYPE_GPU ? "GPU" : (dtype == CL_DEVICE_TYPE_CPU ? "CPU" : "Other")),
vendor_id, global_mem / (1024.0 * 1024.0));
free(device_name);
}
}

Expand Down
Binary file modified QuickVina-W-GPU-2.1/QuickVina-W-GPU-2-1
View file
Open in desktop
Binary file not shown.
2 changes: 1 addition & 1 deletion QuickVina-W-GPU-2.1/input_file_example/2bm2_config.txt
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,6 +1,6 @@
receptor = ./input_file_example/2bm2_protein.pdbqt
ligand_directory = ./test
opencl_binary_path = /home/shidi/Vina-GPU-2.1/QuickVina-W-GPU-2.1
opencl_binary_path = /home/vinay/code/Vina-GPU-2.1/QuickVina-W-GPU-2.1
center_x = 40.415
center_y = 110.986
center_z = 82.673
Expand Down
36 changes: 30 additions & 6 deletions QuickVina-W-GPU-2.1/lib/main_procedure_cl.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -78,6 +78,7 @@ std::vector<output_type> cl_to_vina(output_type_cl result_ptr[],
for (int j = 0; j < MAX_NUM_OF_ATOMS; j++) {
vec v_tmp(tmp_coords.coords[j][0], tmp_coords.coords[j][1], tmp_coords.coords[j][2]);
if ((v_tmp[0] != 0 || v_tmp[1] != 0) || (v_tmp[2] != 0)) tmp_vina.coords.push_back(v_tmp);
tmp_vina.coords.push_back(v_tmp);
}
results_vina.push_back(tmp_vina);
if (i == 0)num_atoms = tmp_vina.coords.size();
Expand All @@ -93,6 +94,7 @@ int get_n(int thread, int search_depth) {
}
}


void main_procedure_cl(cache& c, const std::vector<model>& ms, const precalculate& p, const parallel_mc par,
const vec& corner1, const vec& corner2, const int seed, std::vector<output_container>& outs, std::string opencl_binary_path,
const std::vector<std::vector<std::string>> ligand_names, const int rilc_bfgs) {
Expand Down Expand Up @@ -563,7 +565,7 @@ void main_procedure_cl(cache& c, const std::vector<model>& ms, const precalcula
}
};
tmp_struct ts;
for (int i = 0; i < m_ligand.children.size(); i++) {
for (int i = 0; i < m_ligand.children.size(); i++) {
ts.parent_index = 0; // Start a new branch, whose parent is 0
ts.store_node(m_ligand.children[i], m_ptr->ligand.rigid);
}
Expand All @@ -585,7 +587,7 @@ void main_procedure_cl(cache& c, const std::vector<model>& ms, const precalcula

// Init results
result_ptrs[ligand_count] = (output_type_cl*)malloc(par.mc.thread * sizeof(output_type_cl));

size_t global_buffer_size = pow(2, code_num) * sizeof(ele_cl);
//cl_mem ric_gpu;
CreateDeviceBuffer(&ric_gpus[ligand_count], CL_MEM_READ_ONLY, ric_size, context);
err = clEnqueueWriteBuffer(queue, ric_gpus[ligand_count], false, 0, ric_size, ric_ptr, 0, NULL, NULL); checkErr(err);
Expand All @@ -611,12 +613,17 @@ void main_procedure_cl(cache& c, const std::vector<model>& ms, const precalcula

//allocate global_buffer
ele_cl output_type_c;
size_t global_buffer_size = pow(2, code_num) * sizeof(ele_cl);
cl_mem global_buffer_gpu;
CreateDeviceBuffer(&global_buffer_gpu, CL_MEM_READ_WRITE, global_buffer_size, context);
size_t kernel2_global_size[2] = { 512, 32 };
size_t kernel2_local_size[2] = { 16, 2 };

clFinish(queue);
// Before kernel launch, allocate all_lists buffer for individual_container
size_t num_work_items = kernel2_global_size[0] * kernel2_global_size[1];
cl_mem all_lists_gpu;
CreateDeviceBuffer(&all_lists_gpu, CL_MEM_READ_WRITE, num_work_items * sizeof(individual_container), context);

clFinish(queue);
SetKernelArg(kernels[1], 0, sizeof(cl_mem), &ric_gpus[ligand_count]);
SetKernelArg(kernels[1], 1, sizeof(cl_mem), &m_gpus[ligand_count]);
SetKernelArg(kernels[1], 2, sizeof(cl_mem), &pre_gpu);
Expand All @@ -637,8 +644,7 @@ void main_procedure_cl(cache& c, const std::vector<model>& ms, const precalcula
SetKernelArg(kernels[1], 17, sizeof(cl_mem), &global_buffer_gpu);
SetKernelArg(kernels[1], 18, sizeof(cl_mem), &count_gpu);
SetKernelArg(kernels[1], 19, sizeof(int), &rilc_bfgs);
size_t kernel2_global_size[2] = { 512, 32 };
size_t kernel2_local_size[2] = { 16,2 };
SetKernelArg(kernels[1], 20, sizeof(cl_mem), &all_lists_gpu); // new argument

err = clEnqueueNDRangeKernel(queue, kernels[1], 2, 0, kernel2_global_size, kernel2_local_size,
0, NULL, &ligands_events[ligand_count]); checkErr(err);
Expand Down Expand Up @@ -686,11 +692,29 @@ void main_procedure_cl(cache& c, const std::vector<model>& ms, const precalcula
err = clReleaseMemObject(random_maps_gpus[ligand_count]); checkErr(err);
err = clReleaseMemObject(global_buffer_gpu); checkErr(err);
err = clReleaseMemObject(count_gpu); checkErr(err);
err = clReleaseMemObject(all_lists_gpu); checkErr(err);
#ifndef TIME_ANALYSIS
err = clReleaseEvent(ligands_events[ligand_count]); checkErr(err);
#endif // !TIME_ANALYSIS
}
catch(...){
// Robust cleanup for all per-ligand OpenCL resources
printf("[Cleanup] Exception caught for ligand %d. Releasing OpenCL resources...\n", ligand_count);
if (result_coords_gpus[ligand_count]) { clReleaseMemObject(result_coords_gpus[ligand_count]); printf("Released result_coords_gpu\n"); }
if (result_gpus[ligand_count]) { clReleaseMemObject(result_gpus[ligand_count]); printf("Released result_gpu\n"); }
if (ric_gpus[ligand_count]) { clReleaseMemObject(ric_gpus[ligand_count]); printf("Released ric_gpu\n"); }
if (m_gpus[ligand_count]) { clReleaseMemObject(m_gpus[ligand_count]); printf("Released m_gpu\n"); }
if (random_maps_gpus[ligand_count]) { clReleaseMemObject(random_maps_gpus[ligand_count]); printf("Released random_maps_gpu\n"); }
// if (all_lists_gpu) { clReleaseMemObject(all_lists_gpu); printf("Released all_lists_gpu\n"); }
#ifndef TIME_ANALYSIS
if (ligands_events[ligand_count]) { clReleaseEvent(ligands_events[ligand_count]); printf("Released ligand event\n"); }
#endif
// Free host-side allocations
if (rand_maps_ptrs[ligand_count]) { free(rand_maps_ptrs[ligand_count]); printf("Freed rand_maps_ptr\n"); }
if (ric_ptrs[ligand_count]) { free(ric_ptrs[ligand_count]); printf("Freed ric_ptr\n"); }
if (m_ptrs[ligand_count]) { free(m_ptrs[ligand_count]); printf("Freed m_ptr\n"); }
if (result_coords_ptrs[ligand_count]) { free(result_coords_ptrs[ligand_count]); printf("Freed result_coords_ptr\n"); }
if (result_ptrs[ligand_count]) { free(result_ptrs[ligand_count]); printf("Freed result_ptr\n"); }
continue;
}
}
Expand Down
Loading