Skip to content

Faster ptt #39

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
36000 wants to merge 12 commits into
base: master
Choose a base branch
from
Open

Faster ptt #39

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
12 commits
Select commit Hold shift + click to select a range
5f338dc
transition to tex
36000 Mar 17, 2026
e55bea2
more prog
36000 Mar 17, 2026
9f83d5d
in progress
36000 Mar 17, 2026
436854c
in progress
36000 Mar 17, 2026
11a5465
add rng seed for seeds
36000 Mar 17, 2026
5af7623
in progress
36000 Mar 17, 2026
bf705d9
in progress
36000 Mar 17, 2026
bbf107b
in progress
36000 Mar 17, 2026
c8299f3
WIP ptt tex memory
36000 Mar 17, 2026
8e3e550
return to st
36000 Mar 18, 2026
ad20879
remove this
36000 Mar 18, 2026
fc32030
massive improvement to PTT: ODF LUT
36000 Mar 18, 2026
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
23 changes: 14 additions & 9 deletions cuslines/cuda_c/generate_streamlines_cuda.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -36,6 +36,7 @@
#include "utils.cu"
#include "tracking_helpers.cu"
#include "boot.cu"
#include "ptt_init.cu"
#include "ptt.cu"

#define MAX_NUM_DIR (128)
Expand Down Expand Up @@ -292,6 +293,7 @@ __device__ int get_direction_prob_d(curandStatePhilox4_32_10_t *st,
template<int BDIM_X,
int BDIM_Y,
ModelType MODEL_T,
typename DATA_T,
typename REAL_T,
typename REAL3_T>
__device__ int tracker_d(curandStatePhilox4_32_10_t *st,
Expand All @@ -308,10 +310,10 @@ __device__ int tracker_d(curandStatePhilox4_32_10_t *st,
const int dimy,
const int dimz,
const int dimt,
const REAL_T *__restrict__ dataf,
DATA_T dataf,
const REAL_T *__restrict__ metric_map,
const int samplm_nr,
const REAL3_T *__restrict__ sphere_vertices,
DATA_T sphere_vertices,
const int2 *__restrict__ sphere_edges,
const int num_edges,
int *__restrict__ nsteps,
Expand Down Expand Up @@ -356,7 +358,7 @@ __device__ int tracker_d(curandStatePhilox4_32_10_t *st,
direction,
dimx, dimy, dimz, dimt,
point,
sphere_vertices,
(const REAL3_T *__restrict__) sphere_vertices,
sphere_edges,
num_edges,
__sh_new_dir + tidy);
Expand Down Expand Up @@ -511,6 +513,7 @@ __global__ void getNumStreamlinesProb_k(const REAL_T max_angle,
template<int BDIM_X,
int BDIM_Y,
ModelType MODEL_T,
typename DATA_T,
typename REAL_T,
typename REAL3_T>
__global__ void genStreamlinesMergeProb_k(
Expand All @@ -527,10 +530,10 @@ __global__ void genStreamlinesMergeProb_k(
const int dimy,
const int dimz,
const int dimt,
const REAL_T *__restrict__ dataf,
DATA_T dataf,
const REAL_T *__restrict__ metric_map,
const int samplm_nr,
const REAL3_T *__restrict__ sphere_vertices,
DATA_T sphere_vertices,
const int2 *__restrict__ sphere_edges,
const int num_edges,
const int *__restrict__ slineOutOff,
Expand Down Expand Up @@ -594,13 +597,13 @@ __global__ void genStreamlinesMergeProb_k(
if (MODEL_T == PTT) {
if (!init_frame_ptt_d<BDIM_X, BDIM_Y>(
&st,
dataf,
(cudaTextureObject_t*) dataf,
max_angle,
step_size,
first_step,
dimx, dimy, dimz, dimt,
seed,
sphere_vertices,
(cudaTextureObject_t*) sphere_vertices,
__ptt_frame
)) { // this fails rarely
if (tidx == 0) {
Expand All @@ -617,7 +620,8 @@ __global__ void genStreamlinesMergeProb_k(
int stepsB;
const int tissue_classB = tracker_d<BDIM_X,
BDIM_Y,
MODEL_T>(&st,
MODEL_T,
DATA_T>(&st,
max_angle,
tc_threshold,
step_size,
Expand Down Expand Up @@ -651,7 +655,8 @@ __global__ void genStreamlinesMergeProb_k(
int stepsF;
const int tissue_classF = tracker_d<BDIM_X,
BDIM_Y,
MODEL_T>(&st,
MODEL_T,
DATA_T>(&st,
max_angle,
tc_threshold,
step_size,
Expand Down
2 changes: 0 additions & 2 deletions cuslines/cuda_c/globals.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -85,8 +85,6 @@

#define EXCESS_ALLOC_FACT 2

#define NORM_EPS ((REAL)1e-8)

#if 0
#define DEBUG
#endif
Expand Down
Loading
Loading