Skip to content

Explicitly template and instantiate d_ValueFill function #437

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

Merged
dtegunov merged 1 commit into from
Apr 14, 2026
Merged

Explicitly template and instantiate d_ValueFill function #437

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
4 changes: 2 additions & 2 deletions NativeAcceleration/gtom/src/Correlation/SimilarityMatrix.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -38,7 +38,7 @@ namespace gtom

d_imagesft += ElementsFFT2(dimsimage) * startvalid;
d_similarity += startvalid;
d_ValueFill(d_similarity, nvalid, (tfloat)-1);
d_ValueFill<tfloat>(d_similarity, nvalid, (tfloat)-1);

for (int i = 0; i < anglesteps; i++)
{
Expand Down Expand Up @@ -73,7 +73,7 @@ namespace gtom
tcomplex* d_target = d_linesft + ElementsFFT2(dimslines) * target;
d_linesft += ElementsFFT2(dimslines) * startvalid;
d_similarity += startvalid;
d_ValueFill(d_similarity, nvalid, (tfloat)-1);
d_ValueFill<tfloat>(d_similarity, nvalid, (tfloat)-1);

for (int i = 0; i < anglesteps; i++)
{
Expand Down
4 changes: 2 additions & 2 deletions NativeAcceleration/gtom/src/Correlation/SubTomograms.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -39,8 +39,8 @@ namespace gtom
/*if (nvolumes > batchsize)
throw;*/

d_ValueFill(d_bestcorrelation, Elements(dimsvolume) * nvolumes, (tfloat)-1e30);
d_ValueFill(d_bestangle, Elements(dimsvolume) * nvolumes, (float)0);
d_ValueFill<tfloat>(d_bestcorrelation, Elements(dimsvolume) * nvolumes, (tfloat)-1e30);
d_ValueFill<float>(d_bestangle, Elements(dimsvolume) * nvolumes, (float)0);

tcomplex* d_projectedftctf;
cudaMalloc((void**)&d_projectedftctf, ElementsFFT(dimsvolume) * batchsize * sizeof(tcomplex));
Expand Down
8 changes: 7 additions & 1 deletion NativeAcceleration/gtom/src/Helpers/Memory.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -284,12 +284,18 @@ namespace gtom
return CudaMallocValueFilled<tfloat>(elements, (tfloat)0.0);
}

// Explicitly instantiate d_ValueFill to prevent linker errors
template void d_ValueFill<float>(float* d_array, size_t elements, float value);
template void d_ValueFill<float2> (float2* d_array, size_t elements, float2 value);
template void d_ValueFill<double>(double* d_array, size_t elements, double value);


template <class T> T* CudaMallocValueFilled(size_t elements, T value)
{
T* d_array;
cudaMalloc((void**)&d_array, elements * sizeof(T));

d_ValueFill(d_array, elements, value);
d_ValueFill<T>(d_array, elements, value);

return d_array;
}
Expand Down
5 changes: 2 additions & 3 deletions NativeAcceleration/gtom/src/ImageManipulation/LocalLowpass.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -21,7 +21,7 @@ namespace gtom
{
tcomplex* d_inputft;
cudaMalloc((void**)&d_inputft, ElementsFFT(dims) * sizeof(tcomplex));
d_ValueFill(d_output, Elements(dims), (tfloat)0);
d_ValueFill<tfloat>(d_output, Elements(dims), (tfloat)0);
tcomplex* d_maskedft;
cudaMalloc((void**)&d_maskedft, ElementsFFT(dims) * sizeof(tcomplex));
tfloat* d_cleanresolution;
Expand Down Expand Up @@ -63,7 +63,7 @@ namespace gtom
tfloat res = (tfloat)b * binsize + minval;
tfloat freq = (tfloat)dims.x / res;

d_ValueFill(d_mask, Elements(dims), (tfloat)1);
d_ValueFill<tfloat>(d_mask, Elements(dims), (tfloat)1);
d_SphereMask(d_mask, d_mask, dims, &freq, 0, NULL, false);
d_RemapFull2HalfFFT(d_mask, d_maskhalf, dims);

Expand Down Expand Up @@ -93,5 +93,4 @@ namespace gtom
////////////////
//CUDA kernels//
////////////////

}
6 changes: 3 additions & 3 deletions NativeAcceleration/src/CTF.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -49,7 +49,7 @@ __declspec(dllexport) void CreateSpectra(float* d_frame,
int pertimegroup = nframes / ctfgrid.z;

// Temp spectra will be summed up to be averaged later
d_ValueFill(d_outputall, ElementsFFT2(dimsregionscaled) * nspectra, 0.0f);
d_ValueFill<float>(d_outputall, ElementsFFT2(dimsregionscaled) * nspectra, 0.0f);

cufftHandle ownplanforw = planforw > 0 ? planforw : d_FFTR2CGetPlan(2, toInt3(dimsregion), norigins);
cufftHandle ownplanback = planback > 0 ? planback : d_IFFTC2RGetPlan(2, toInt3(dimsregionscaled), norigins);
Expand Down Expand Up @@ -97,7 +97,7 @@ __declspec(dllexport) void CreateSpectra(float* d_frame,
{
tfloat* d_tempspectraaccumulated;
cudaMalloc((void**)&d_tempspectraaccumulated, norigins * ElementsFFT2(dimsregion) * sizeof(tfloat));
d_ValueFill(d_tempspectraaccumulated, ElementsFFT2(dimsregion) * norigins, (tfloat)0);
d_ValueFill<tfloat>(d_tempspectraaccumulated, ElementsFFT2(dimsregion) * norigins, (tfloat)0);

for (size_t z = 0; z < nframes; z++)
{
Expand Down Expand Up @@ -247,4 +247,4 @@ __global__ void ScaleNormCorrSumKernel(float2* d_simcoords, float* d_sim, float*

d_scores[blockIdx.x] = sum1 / (float)length;
}
}
}
4 changes: 2 additions & 2 deletions NativeAcceleration/src/RealspaceProjection.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -12,7 +12,7 @@ __declspec(dllexport) void RealspaceProjectForward(float* d_volume,
float3* h_angles,
int batch)
{
d_ValueFill(d_projections, Elements2(dimsproj) * batch, 0.0f);
d_ValueFill<float>(d_projections, Elements2(dimsproj) * batch, 0.0f);

glm::mat3* d_matrices;

Expand Down Expand Up @@ -183,4 +183,4 @@ __global__ void RealspaceProjectBackwardKernel(float* d_volume, int3 dimsvolume,
else
d_volume[id] = sum;
}
}
}
8 changes: 4 additions & 4 deletions NativeAcceleration/src/Tools.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -651,7 +651,7 @@ __declspec(dllexport) void ProjectNMAPseudoAtoms(float3* d_positions,
int2 dimsproj,
uint batch)
{
d_ValueFill(d_proj, Elements2(dimsproj) * batch, 0.0f);
d_ValueFill<float>(d_proj, Elements2(dimsproj) * batch, 0.0f);

d_ProjNMAPseudoAtoms(d_positions,
d_intensities,
Expand Down Expand Up @@ -691,7 +691,7 @@ __declspec(dllexport) void ProjectSoftPseudoAtoms(float3* d_positions,
int2 dimsproj,
uint batch)
{
d_ValueFill(d_proj, Elements2(dimsproj) * batch, 0.0f);
d_ValueFill<float>(d_proj, Elements2(dimsproj) * batch, 0.0f);

d_ProjSoftPseudoAtoms(d_positions,
d_intensities,
Expand Down Expand Up @@ -759,12 +759,12 @@ __declspec(dllexport) void DestroyTexture(unsigned long long textureid, unsigned

__declspec(dllexport) void ValueFill(float* d_input, size_t elements, float value)
{
d_ValueFill(d_input, elements, value);
d_ValueFill<float>(d_input, elements, value);
}

__declspec(dllexport) void ValueFillComplex(float2* d_input, size_t elements, float2 value)
{
d_ValueFill(d_input, elements, value);
d_ValueFill<float2>(d_input, elements, value);
}

__declspec(dllexport) void Real(float2* d_input, float* d_output, size_t elements)
Expand Down