Skip to content

Updates for cuda 13. #88

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
kaushikcfd wants to merge 1 commit into
base: master-2025
Choose a base branch
from
Open

Updates for cuda 13. #88

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
2 changes: 1 addition & 1 deletion config/find_cuda_libs.sh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -42,7 +42,7 @@ libspec=`$NVCC -dryrun bogus.cu 2>&1 | grep LIBRARIES | sed 's/^.*LIBRARIES=//'`
#echo "libspec=$libspec"
if [ $cudart_flag_supported -eq 1 ]
then
cudalibs=`$NVCC -dryrun bogus.cu 2>&1 | tail -1 | sed "s#^.*-o \"a.out\"##" | sed 's#"[a-zA-Z0-9./_-]*\.o"##g' | sed 's/-Wl,--start-group//' | sed 's/-Wl,--end-group//'`
cudalibs=`$NVCC -dryrun bogus.cu 2>&1 | tail -1 | grep -oE '(^|[[:space:]])"?-[lL][^[:space:]]*"?'`
else
cudalibs=$libspec
fi
Expand Down
2 changes: 1 addition & 1 deletion configure
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -5855,7 +5855,7 @@ fi

# get list of supported archs
supported_NVCC_archs=""
for i in `$NVCC --help | grep -o "compute_[0-9][0-9]" | sed 's,compute,,' | sort | uniq`
for i in `$NVCC --help | grep -o "compute_[0-9][0-9]*" | sed 's,compute,,' | sort | uniq`
do
supported_NVCC_archs+="-gencode=arch=compute$i,code=sm$i "
done
Expand Down
2 changes: 1 addition & 1 deletion configure.ac
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -199,7 +199,7 @@ AS_IF([test "x$with_cuda" != xno],

# get list of supported archs
supported_NVCC_archs=""
for i in `$NVCC --help | grep -o "compute_[[0-9]][[0-9]]" | sed 's,compute,,' | sort | uniq`
for i in `$NVCC --help | grep -o "compute_[[0-9]][[0-9]]*" | sed 's,compute,,' | sort | uniq`
do
supported_NVCC_archs+="-gencode=arch=compute$i,code=sm$i "
done
Expand Down
6 changes: 6 additions & 0 deletions src/cuda/common/main.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -100,7 +100,13 @@ void EnumerateDevicesAndChoose(int chooseDevice, bool verbose)
deviceProp.totalConstMem) << endl;
cout << " major (hw version) = " << deviceProp.major << endl;
cout << " minor (hw version) = " << deviceProp.minor << endl;
#if CUDART_VERSION >= 13000
int clockRate;
cudaDeviceGetAttribute(&clockRate, cudaDevAttrClockRate, device);
cout << " clockRate = " << clockRate << endl;
#else
cout << " clockRate = " << deviceProp.clockRate << endl;
#endif
cout << " textureAlignment = " << deviceProp.textureAlignment
<< endl;
}
Expand Down
2 changes: 1 addition & 1 deletion src/cuda/level0/BusSpeedReadback.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -122,7 +122,7 @@ void RunBenchmark(ResultDatabase &resultDB,

cudaMemcpy(device, hostMem1,
numMaxFloats*sizeof(float), cudaMemcpyHostToDevice);
cudaThreadSynchronize();
cudaDeviceSynchronize();
const unsigned int passes = op.getOptionInt("passes");

cudaEvent_t start, stop;
Expand Down
53 changes: 29 additions & 24 deletions src/cuda/level0/DeviceMemory.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -21,11 +21,9 @@ __global__ void writeGlobalMemoryCoalesced(float *output, int size, int repeat);
__global__ void writeGlobalMemoryUnit(float *output, int size, int repeat);
__global__ void writeLocalMemory(float *output, int size, int repeat);
__device__ int getRand(int seed, int mod);
__global__ void readTexels(int n, float *d_out, int width);
__global__ void readTexelsInCache(int n, float *d_out);
__global__ void readTexelsRandom(int n, float *d_out, int width, int height);
// Texture to use for the benchmarks
texture<float4, 2, cudaReadModeElementType> texA;
__global__ void readTexels(cudaTextureObject_t tex, int n, float *d_out, int width);
__global__ void readTexelsInCache(cudaTextureObject_t tex, int n, float *d_out);
__global__ void readTexelsRandom(cudaTextureObject_t tex, int n, float *d_out, int width, int height);

// ****************************************************************************
// Function: addBenchmarkSpecOptions
Expand Down Expand Up @@ -308,12 +306,6 @@ void TestTextureMem(ResultDatabase &resultDB, OptionParser &op, double scalet)
cudaEventCreate(&stop);
CHECK_CUDA_ERROR();

// make sure our texture behaves like we want....
texA.normalized = false;
texA.addressMode[0] = cudaAddressModeClamp;
texA.addressMode[1] = cudaAddressModeClamp;
texA.filterMode = cudaFilterModePoint;

for (int j = 0; j < nsizes; j++)
{
cout << "Benchmarking Texture Memory, Test Size: " << j+1 << " / 5\n";
Expand Down Expand Up @@ -350,15 +342,28 @@ void TestTextureMem(ResultDatabase &resultDB, OptionParser &op, double scalet)

// Allocate a cuda array
cudaArray* cuArray;
cudaMallocArray(&cuArray, &texA.channelDesc, width, height);
cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<float4>();
cudaMallocArray(&cuArray, &channelDesc, width, height);
CHECK_CUDA_ERROR();

// Copy in source data
cudaMemcpyToArray(cuArray, 0, 0, h_in, size, cudaMemcpyHostToDevice);
CHECK_CUDA_ERROR();

// Bind texture to the array
cudaBindTextureToArray(texA, cuArray);
cudaResourceDesc resDesc = {};
resDesc.resType = cudaResourceTypeArray;
resDesc.res.array.array = cuArray;

cudaTextureDesc texDesc = {};
texDesc.addressMode[0] = cudaAddressModeClamp;
texDesc.addressMode[1] = cudaAddressModeClamp;
texDesc.filterMode = cudaFilterModePoint;
texDesc.readMode = cudaReadModeElementType;
texDesc.normalizedCoords = false;

cudaTextureObject_t texA = 0;
cudaCreateTextureObject(&texA, &resDesc, &texDesc, NULL);
CHECK_CUDA_ERROR();

for (int p = 0; p < passes; p++)
Expand All @@ -370,8 +375,8 @@ void TestTextureMem(ResultDatabase &resultDB, OptionParser &op, double scalet)
// read texels from texture
for (int iter = 0; iter < iterations; iter++)
{
readTexels<<<gridSize, blockSize>>>(kernelRepFactor, d_out,
width);
readTexels<<<gridSize, blockSize>>>(texA, kernelRepFactor,
d_out, width);
}
cudaEventRecord(stop, 0);
CHECK_CUDA_ERROR();
Expand All @@ -398,7 +403,7 @@ void TestTextureMem(ResultDatabase &resultDB, OptionParser &op, double scalet)
for (int iter = 0; iter < iterations; iter++)
{
readTexelsInCache<<<gridSize, blockSize>>>
(kernelRepFactor, d_out);
(texA, kernelRepFactor, d_out);
}
cudaEventRecord(stop, 0);
cudaEventSynchronize(stop);
Expand All @@ -425,7 +430,7 @@ void TestTextureMem(ResultDatabase &resultDB, OptionParser &op, double scalet)
for (int iter = 0; iter < iterations; iter++)
{
readTexelsRandom<<<gridSize, blockSize>>>
(kernelRepFactor, d_out, width, height);
(texA, kernelRepFactor, d_out, width, height);
}

cudaEventRecord(stop, 0);
Expand All @@ -446,7 +451,7 @@ void TestTextureMem(ResultDatabase &resultDB, OptionParser &op, double scalet)
delete[] h_out;
cudaFree(d_out);
cudaFreeArray(cuArray);
cudaUnbindTexture(texA);
cudaDestroyTextureObject(texA);
}
cudaEventDestroy(start);
cudaEventDestroy(stop);
Expand Down Expand Up @@ -638,7 +643,7 @@ writeLocalMemory(float *output, int size, int repeat)
}

// Simple Repeated Linear Read from texture memory
__global__ void readTexels(int n, float *d_out, int width)
__global__ void readTexels(cudaTextureObject_t tex, int n, float *d_out, int width)
{
int idx_x = (blockIdx.x * blockDim.x) + threadIdx.x;
int idx_y = (blockIdx.y * blockDim.y) + threadIdx.y;
Expand All @@ -647,30 +652,30 @@ __global__ void readTexels(int n, float *d_out, int width)
int width_bits = width-1;
for (int i = 0; i < n; i++)
{
float4 v = tex2D(texA, float(idx_x), float(idx_y));
float4 v = tex2D<float4>(tex, float(idx_x), float(idx_y));
idx_x = (idx_x+1) & width_bits;
sum += v.x;
}
d_out[out_idx] = sum;
}

// Repeated read of only 4kb of texels (should fit in texture cache)
__global__ void readTexelsInCache(int n, float *d_out)
__global__ void readTexelsInCache(cudaTextureObject_t tex, int n, float *d_out)
{
int idx_x = (blockIdx.x * blockDim.x) + threadIdx.x;
int idx_y = (blockIdx.y * blockDim.y) + threadIdx.y;
int out_idx = idx_y * gridDim.x + idx_x;
float sum = 0.0f;
for (int i = 0; i < n; i++)
{
float4 v = tex2D(texA, float(idx_x), float(idx_y));
float4 v = tex2D<float4>(tex, float(idx_x), float(idx_y));
sum += v.x;
}
d_out[out_idx] = sum;
}

// Read "random" texels
__global__ void readTexelsRandom(int n, float *d_out, int width, int height)
__global__ void readTexelsRandom(cudaTextureObject_t tex, int n, float *d_out, int width, int height)
{
int idx_x = (blockIdx.x * blockDim.x) + threadIdx.x;
int idx_y = (blockIdx.y * blockDim.y) + threadIdx.y;
Expand All @@ -680,7 +685,7 @@ __global__ void readTexelsRandom(int n, float *d_out, int width, int height)
int height_bits = height-1;
for (int i = 0; i < n; i++)
{
float4 v = tex2D(texA, float(idx_x), float(idx_y));
float4 v = tex2D<float4>(tex, float(idx_x), float(idx_y));
idx_x = (idx_x*3+29)&(width_bits);
idx_y = (idx_y*5+11)&(height_bits);
sum += v.x;
Expand Down
8 changes: 4 additions & 4 deletions src/cuda/level1/fft/fftlib.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -163,7 +163,7 @@ forward(void* work, const int n_ffts)
(cufftComplex*)work, CUFFT_FORWARD);
}
printCUFFTError(res);
cudaThreadSynchronize();
cudaDeviceSynchronize();
CHECK_CUDA_ERROR();
#else
if (do_dp)
Expand All @@ -174,7 +174,7 @@ forward(void* work, const int n_ffts)
{
FFT512_device<float2, float><<<grid2D(n_ffts), 64>>>((float2*)work);
}
cudaThreadSynchronize();
cudaDeviceSynchronize();
CHECK_CUDA_ERROR();
#endif
}
Expand Down Expand Up @@ -206,7 +206,7 @@ inverse(void* work, const int n_ffts)
{
norm512_device<float2><<<grid2D(n_ffts), 64>>>((float2*)work);
}
cudaThreadSynchronize();
cudaDeviceSynchronize();
CHECK_CUDA_ERROR();
#else
if (do_dp)
Expand All @@ -217,7 +217,7 @@ inverse(void* work, const int n_ffts)
{
IFFT512_device<float2, float><<<grid2D(n_ffts), 64>>>((float2*)work);
}
cudaThreadSynchronize();
cudaDeviceSynchronize();
CHECK_CUDA_ERROR();
// normalization built in to inverse...
#endif
Expand Down
2 changes: 1 addition & 1 deletion src/cuda/level1/gemm/GEMM.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -228,7 +228,7 @@ void RunTest(string testName, ResultDatabase &resultDB, OptionParser &op)
// Warm Up
devGEMM<T>(transa, transb, m, n, k, alpha, dA, lda, dB, ldb, beta,
dC, ldc);
CUDA_SAFE_CALL(cudaThreadSynchronize());
CUDA_SAFE_CALL(cudaDeviceSynchronize());

double cublas_time;
float kernel_time = 0.0f;
Expand Down
56 changes: 30 additions & 26 deletions src/cuda/level1/md/MD.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -36,21 +36,19 @@ inline int populateNeighborList(std::list<T>& currDist,
std::list<int>& currList, const int j, const int nAtom,
int* neighborList);

// Texture caches for position info
texture<float4, 1, cudaReadModeElementType> posTexture;
texture<int4, 1, cudaReadModeElementType> posTexture_dp;

struct texReader_sp {
cudaTextureObject_t tex;
__device__ __forceinline__ float4 operator()(int idx) const
{
return tex1Dfetch(posTexture, idx);
return tex1Dfetch<float4>(tex, idx);
}
};

// CUDA doesn't support double4 textures, so we have to do some conversion
// here, resulting in a bit of overhead, but it's still faster than
// an uncoalesced read
struct texReader_dp {
cudaTextureObject_t tex;
__device__ __forceinline__ double4 operator()(int idx) const
{
#if (__CUDA_ARCH__ < 130)
Expand All @@ -60,11 +58,11 @@ struct texReader_dp {
// but since the arch doesn't support DP, it will never be called
return make_double4(0., 0., 0., 0.);
#else
int4 v = tex1Dfetch(posTexture_dp, idx*2);
int4 v = tex1Dfetch<int4>(tex, idx*2);
double2 a = make_double2(__hiloint2double(v.y, v.x),
__hiloint2double(v.w, v.z));

v = tex1Dfetch(posTexture_dp, idx*2 + 1);
v = tex1Dfetch<int4>(tex, idx*2 + 1);
double2 b = make_double2(__hiloint2double(v.y, v.x),
__hiloint2double(v.w, v.z));

Expand Down Expand Up @@ -105,7 +103,8 @@ __global__ void compute_lj_force(forceVecType* __restrict__ force3,
const T cutsq,
const T lj1,
const T lj2,
const int inum)
const int inum,
const texReader positionTexReader)
{
// Global ID - one thread per atom
int idx = blockIdx.x*blockDim.x + threadIdx.x;
Expand All @@ -116,8 +115,6 @@ __global__ void compute_lj_force(forceVecType* __restrict__ force3,
// Force accumulator
forceVecType f = {0.0f, 0.0f, 0.0f};

texReader positionTexReader;

int j = 0;
while (j < neighCount)
{
Expand Down Expand Up @@ -349,20 +346,24 @@ void runTest(const string& testName, ResultDatabase& resultDB, OptionParser& op)
position[i].z = (T)(drand48() * domainEdge);
}

texReader positionTexReader = {};
if (useTexture)
{
// Set up 1D texture to cache position info
cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<float4>();

// Bind a 1D texture to the position array
CUDA_SAFE_CALL(cudaBindTexture(0, posTexture, d_position, channelDesc,
nAtom*sizeof(float4)));

cudaChannelFormatDesc channelDesc2 = cudaCreateChannelDesc<int4>();

// Bind a 1D texture to the position array
CUDA_SAFE_CALL(cudaBindTexture(0, posTexture_dp, d_position,
channelDesc2, nAtom*sizeof(double4)));
// Set up a 1D texture object to cache position info
cudaResourceDesc resDesc = {};
resDesc.resType = cudaResourceTypeLinear;
resDesc.res.linear.devPtr = d_position;
// CUDA has no native double4 texture format, so texReader_dp
// reinterprets the data as int4 pairs instead.
resDesc.res.linear.desc = (sizeof(T) == sizeof(double)) ?
cudaCreateChannelDesc<int4>() : cudaCreateChannelDesc<float4>();
resDesc.res.linear.sizeInBytes = nAtom*sizeof(posVecType);

cudaTextureDesc texDesc = {};
texDesc.readMode = cudaReadModeElementType;

CUDA_SAFE_CALL(cudaCreateTextureObject(&positionTexReader.tex, &resDesc,
&texDesc, NULL));
}

// Keep track of how many atoms are within the cutoff distance to
Expand Down Expand Up @@ -403,8 +404,8 @@ void runTest(const string& testName, ResultDatabase& resultDB, OptionParser& op)
compute_lj_force<T, forceVecType, posVecType, useTexture, texReader>
<<<gridSize, blockSize>>>
(d_force, d_position, maxNeighbors, d_neighborList,
cutsq, lj1, lj2, nAtom);
CUDA_SAFE_CALL(cudaThreadSynchronize());
cutsq, lj1, lj2, nAtom, positionTexReader);
CUDA_SAFE_CALL(cudaDeviceSynchronize());

// Copy back forces
cudaEvent_t outputTransfer_start, outputTransfer_stop;
Expand Down Expand Up @@ -446,7 +447,7 @@ void runTest(const string& testName, ResultDatabase& resultDB, OptionParser& op)
compute_lj_force<T, forceVecType, posVecType, useTexture, texReader>
<<<gridSize, blockSize>>>
(d_force, d_position, maxNeighbors, d_neighborList, cutsq,
lj1, lj2, nAtom);
lj1, lj2, nAtom, positionTexReader);
}
cudaEventRecord(kernel_stop, 0);
CUDA_SAFE_CALL(cudaEventSynchronize(kernel_stop));
Expand Down Expand Up @@ -489,7 +490,10 @@ void runTest(const string& testName, ResultDatabase& resultDB, OptionParser& op)
CUDA_SAFE_CALL(cudaFreeHost(force));
CUDA_SAFE_CALL(cudaFreeHost(neighborList));
// Device
CUDA_SAFE_CALL(cudaUnbindTexture(posTexture));
if (useTexture)
{
CUDA_SAFE_CALL(cudaDestroyTextureObject(positionTexReader.tex));
}
CUDA_SAFE_CALL(cudaFree(d_position));
CUDA_SAFE_CALL(cudaFree(d_force));
CUDA_SAFE_CALL(cudaFree(d_neighborList));
Expand Down
Loading