update plscore

Minimap2onGPU · Feb 28, 2024 · 94f333a · 94f333a
1 parent 6836cb7
commit 94f333a
Showing 4 changed files with 25 additions and 10 deletions.
diff --git a/aac_config.json b/aac_config.json
@@ -21,7 +21,7 @@
         "mid_blockdim": 64,
         "//blockdim config": "options are not used: static config specified at compile time (make ... LONG_BLOCK_SIZE=1024)",
         "short_griddim": 16128,
-        "long_griddim": 50,
+        "long_griddim": 100,
         "mid_griddim": 16128
     }
 }
diff --git a/gpu/plscore.cu b/gpu/plscore.cu
@@ -395,7 +395,7 @@ __global__ void score_generation_long_map(int32_t* anchors_x, int32_t* anchors_y
     if (tid == 0) {
         segid = bid;
     }
-    #ifdef DEBUG_VERBOSE
+    #ifdef DEBUG_CHECK
     auto start = clock64();
     #endif
     __syncthreads();
@@ -413,7 +413,7 @@ __global__ void score_generation_long_map(int32_t* anchors_x, int32_t* anchors_y
     //     // seg_t seg = long_seg[segid]; // unsorted
     //     compute_sc_seg_multi_wf(anchors_x, anchors_y, sid, range, seg.start_idx, seg.end_idx, f, p);
     // }
-    #ifdef DEBUG_VERBOSE
+    #ifdef DEBUG_CHECK
     auto end = clock64();
     if (threadIdx.x == 0) {
         printf("bid: %d, long kernel time: %lu, process %u segs\n", bid, end - start, seg_count);
@@ -612,6 +612,20 @@ void plscore_async_long_forward_dp(deviceMemPtr* dev_mem, cudaStream_t* stream)
     #endif
     cudaCheck();
 
+
+    // #ifdef __LONG_BLOCK_SIZE__
+    // // fprintf(stderr, "long block size: %d\n", __LONG_BLOCK_SIZE__);
+    // score_generation_long<__LONG_BLOCK_SIZE__><<<longDimGrid, dim3(__LONG_BLOCK_SIZE__, 1, 1), 0, *stream>>>(
+    //     dev_mem->d_ax_long, dev_mem->d_ay_long, dev_mem->d_sid_long, dev_mem->d_range_long, dev_mem->d_long_seg,
+    //     dev_mem->d_long_seg_count, dev_mem->d_f_long, dev_mem->d_p_long);
+    // #else
+    // dim3 longDimBlock(score_kernel_config.long_blockdim, 1, 1);
+    // score_generation_long<<<longDimGrid, longDimBlock, 0, *stream>>>(
+    //     dev_mem->d_ax_long, dev_mem->d_ay_long, dev_mem->d_sid_long, dev_mem->d_range_long, dev_mem->d_long_seg,
+    //     dev_mem->d_long_seg_count, dev_mem->d_f_long, dev_mem->d_p_long);
+    // #endif
+    // cudaCheck();
+
 #ifdef DEBUG_PRINT
     // fprintf(stderr, "[Info] %s (%s:%d) long score generation launched\n", __func__, __FILE__, __LINE__);
 #endif

diff --git a/scripts/aac_omnitrace.slurm b/scripts/aac_omnitrace.slurm
@@ -28,9 +28,10 @@ module load rocm-5.7.1
 # export AMD_LOG_LEVEL=4
 # Replace the following line with the actual command(s) you want to run
 cd $MM2_ROOT
-make clean
-make MICRO_BATCH=2 GPU_CONFIG=aac_config.json SHORT_BLOCK_SIZE=64 LONG_BLOCK_SIZE=1024 MID_BLOCK_SIZE=512 MID_CUT=1 LONG_CUT=50 DEBUG=1 DEBUG_ANALYSIS=1
-omnitrace-sample -PTDH -E all -I rocm-smi -I roctracer -I rocprofiler -I roctx -o omni_output -- ./minimap2 -K 2500000000 -t 1 --max-chain-skip=2147483647 --gpu-chain /shareddata/umich_folder/data/ONT/hg38.mmi /shareddata/umich_folder/data/ONT/random_1GBases_100kto300k.fa
+# make clean
+# make MICRO_BATCH=2 GPU_CONFIG=aac_config.json SHORT_BLOCK_SIZE=64 LONG_BLOCK_SIZE=1024 MID_BLOCK_SIZE=512 MID_CUT=1 LONG_CUT=50 DEBUG=1 DEBUG_ANALYSIS=1
+# omnitrace-sample -PTDH -E all -I rocm-smi -I roctracer -I rocprofiler -I roctx -o omni_output -- ./minimap2 -K 2500000000 -t 1 --max-chain-skip=2147483647 --gpu-chain /shareddata/umich_folder/data/ONT/hg38.mmi /shareddata/umich_folder/data/ONT/random_1GBases_100kto300k.fa
+omnitrace-sample -PTDH -E all -I rocm-smi -I roctracer -I rocprofiler -I roctx -o omni_output -- ./minimap2 -K 2000000000 -t 1 --max-chain-skip=2147483647 --gpu-chain /shareddata/umich_folder/data/ONT/hg38.mmi /shareddata/umich_folder/data/ONT/random_5GBases_90kto100k.fa
 
 # ./minimap2 -t 1 --max-chain-skip=2147483647 --gpu-chain /shareddata/umich_folder/data/ONT/hg38.mmi /shareddata/umich_folder/data/ONT/reads_4f452f4a-d82a-4580-981b-32d14b997217.fa
 # ./minimap2 -t 1 --max-chain-skip=2147483647 --gpu-chain /shareddata/umich_folder/data/ONT/hg38.mmi /shareddata/umich_folder/data/ONT/random_500MBases_200kto300k.fa

diff --git a/scripts/acc_integrated.slurm b/scripts/acc_integrated.slurm
@@ -37,13 +37,13 @@ cd $MM2_ROOT
 
 # random data test
 # export AMD_LOG_LEVEL=4
-make clean
-make MICRO_BATCH=6 GPU_CONFIG=aac_config.json SHORT_BLOCK_SIZE=64 LONG_BLOCK_SIZE=1024 MID_BLOCK_SIZE=512 MID_CUT=1 LONG_CUT=40 DEBUG=1 DEBUG_ANALYSIS=1
-./minimap2 -K 2000000000 -t 1 --max-chain-skip=2147483647 --gpu-chain /shareddata/umich_folder/data/ONT/hg38.mmi /shareddata/umich_folder/data/ONT/random_5GBases_90kto100k.fa
+# make clean
+# make MICRO_BATCH=6 GPU_CONFIG=aac_config.json SHORT_BLOCK_SIZE=64 LONG_BLOCK_SIZE=1024 MID_BLOCK_SIZE=512 MID_CUT=1 LONG_CUT=40 DEBUG=1 DEBUG_ANALYSIS=1
+# ./minimap2 -K 2000000000 -t 1 --max-chain-skip=2147483647 --gpu-chain /shareddata/umich_folder/data/ONT/hg38.mmi /shareddata/umich_folder/data/ONT/random_5GBases_90kto100k.fa
 # ./minimap2 -K 2000000000 -t 1 --max-chain-skip=2147483647 --gpu-chain /shareddata/umich_folder/data/ONT/hg38.mmi /shareddata/umich_folder/data/ONT/random_4GBases_10kto300k.fa 
 
 
-# ./minimap2 -K 2000000000 -t 1 --max-chain-skip=2147483647 --gpu-chain /shareddata/umich_folder/data/ONT/hg38.mmi /shareddata/umich_folder/data/ONT/long_read_3G.fa 
+./minimap2 -K 2000000000 -t 1 --max-chain-skip=2147483647 --gpu-chain /shareddata/umich_folder/data/ONT/hg38.mmi /shareddata/umich_folder/data/ONT/long_read_3G.fa 
 # ./minimap2 -K 2000000000 -t 1 --max-chain-skip=2147483647 --gpu-chain /shareddata/umich_folder/data/ONT/hg38.mmi /shareddata/umich_folder/data/ONT/long_read_600M.fa 
 # ./minimap2 -K 2000000000 -t 1 --max-chain-skip=2147483647 --gpu-chain /shareddata/umich_folder/data/ONT/hg38.mmi /shareddata/umich_folder/data/ONT/random_5GBases_90kto100k.fa
 # ./minimap2 -t 1 --max-chain-skip=2147483647 --gpu-chain /shareddata/umich_folder/data/ONT/hg38.mmi /shareddata/umich_folder/data/ONT/random_500MBases_90kto100k.fa