From 59f3e5e335b0e7cc901406ccff309e4b4cce7d6f Mon Sep 17 00:00:00 2001 From: Kaushik Kulkarni Date: Sat, 18 Jul 2026 22:23:16 -0500 Subject: [PATCH] Updates for cuda 13. - Support 3-digits compute capabilities (e.g. sm_100) - cudaThreadSynchronize -> cudaDeviceSynchronize - Replace cudaUnbindTexture with newer API. - deviceProp.clockRate -> cudaDeviceGetAttribute(...) Assisted-by: Claude Sonnet 5 --- config/find_cuda_libs.sh | 2 +- configure | 2 +- configure.ac | 2 +- src/cuda/common/main.cpp | 6 ++ src/cuda/level0/BusSpeedReadback.cu | 2 +- src/cuda/level0/DeviceMemory.cu | 53 ++++++----- src/cuda/level1/fft/fftlib.cu | 8 +- src/cuda/level1/gemm/GEMM.cpp | 2 +- src/cuda/level1/md/MD.cu | 56 +++++------ src/cuda/level1/spmv/Spmv.cu | 93 ++++++++++--------- src/cuda/level1/triad/Triad.cu | 2 +- src/cuda/level2/qtclustering/QTC.cu | 40 ++++---- src/cuda/level2/qtclustering/kernels_common.h | 20 ++-- .../qtclustering/kernels_compact_storage.h | 4 +- .../qtclustering/kernels_full_storage.h | 10 +- src/mpi/contention-mt/cuda/CUDADriver.cpp | 6 ++ src/mpi/contention/cuda/CUDADriver.cpp | 6 ++ 17 files changed, 179 insertions(+), 135 deletions(-) diff --git a/config/find_cuda_libs.sh b/config/find_cuda_libs.sh index 7222dca3..1b828a0d 100755 --- a/config/find_cuda_libs.sh +++ b/config/find_cuda_libs.sh @@ -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 diff --git a/configure b/configure index 8341da58..9eadee38 100755 --- a/configure +++ b/configure @@ -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 diff --git a/configure.ac b/configure.ac index 3fec8427..26a85d4c 100644 --- a/configure.ac +++ b/configure.ac @@ -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 diff --git a/src/cuda/common/main.cpp b/src/cuda/common/main.cpp index 232afa1a..74a15a6d 100644 --- a/src/cuda/common/main.cpp +++ b/src/cuda/common/main.cpp @@ -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; } diff --git a/src/cuda/level0/BusSpeedReadback.cu b/src/cuda/level0/BusSpeedReadback.cu index 7b5a3bed..9380f8c3 100644 --- a/src/cuda/level0/BusSpeedReadback.cu +++ b/src/cuda/level0/BusSpeedReadback.cu @@ -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; diff --git a/src/cuda/level0/DeviceMemory.cu b/src/cuda/level0/DeviceMemory.cu index b637661d..0f16167e 100644 --- a/src/cuda/level0/DeviceMemory.cu +++ b/src/cuda/level0/DeviceMemory.cu @@ -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 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 @@ -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"; @@ -350,7 +342,8 @@ void TestTextureMem(ResultDatabase &resultDB, OptionParser &op, double scalet) // Allocate a cuda array cudaArray* cuArray; - cudaMallocArray(&cuArray, &texA.channelDesc, width, height); + cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc(); + cudaMallocArray(&cuArray, &channelDesc, width, height); CHECK_CUDA_ERROR(); // Copy in source data @@ -358,7 +351,19 @@ void TestTextureMem(ResultDatabase &resultDB, OptionParser &op, double scalet) 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++) @@ -370,8 +375,8 @@ void TestTextureMem(ResultDatabase &resultDB, OptionParser &op, double scalet) // read texels from texture for (int iter = 0; iter < iterations; iter++) { - readTexels<<>>(kernelRepFactor, d_out, - width); + readTexels<<>>(texA, kernelRepFactor, + d_out, width); } cudaEventRecord(stop, 0); CHECK_CUDA_ERROR(); @@ -398,7 +403,7 @@ void TestTextureMem(ResultDatabase &resultDB, OptionParser &op, double scalet) for (int iter = 0; iter < iterations; iter++) { readTexelsInCache<<>> - (kernelRepFactor, d_out); + (texA, kernelRepFactor, d_out); } cudaEventRecord(stop, 0); cudaEventSynchronize(stop); @@ -425,7 +430,7 @@ void TestTextureMem(ResultDatabase &resultDB, OptionParser &op, double scalet) for (int iter = 0; iter < iterations; iter++) { readTexelsRandom<<>> - (kernelRepFactor, d_out, width, height); + (texA, kernelRepFactor, d_out, width, height); } cudaEventRecord(stop, 0); @@ -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); @@ -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; @@ -647,7 +652,7 @@ __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(tex, float(idx_x), float(idx_y)); idx_x = (idx_x+1) & width_bits; sum += v.x; } @@ -655,7 +660,7 @@ __global__ void readTexels(int n, float *d_out, int width) } // 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; @@ -663,14 +668,14 @@ __global__ void readTexelsInCache(int n, float *d_out) float sum = 0.0f; for (int i = 0; i < n; i++) { - float4 v = tex2D(texA, float(idx_x), float(idx_y)); + float4 v = tex2D(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; @@ -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(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; diff --git a/src/cuda/level1/fft/fftlib.cu b/src/cuda/level1/fft/fftlib.cu index a7dc41e1..2d007294 100644 --- a/src/cuda/level1/fft/fftlib.cu +++ b/src/cuda/level1/fft/fftlib.cu @@ -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) @@ -174,7 +174,7 @@ forward(void* work, const int n_ffts) { FFT512_device<<>>((float2*)work); } - cudaThreadSynchronize(); + cudaDeviceSynchronize(); CHECK_CUDA_ERROR(); #endif } @@ -206,7 +206,7 @@ inverse(void* work, const int n_ffts) { norm512_device<<>>((float2*)work); } - cudaThreadSynchronize(); + cudaDeviceSynchronize(); CHECK_CUDA_ERROR(); #else if (do_dp) @@ -217,7 +217,7 @@ inverse(void* work, const int n_ffts) { IFFT512_device<<>>((float2*)work); } - cudaThreadSynchronize(); + cudaDeviceSynchronize(); CHECK_CUDA_ERROR(); // normalization built in to inverse... #endif diff --git a/src/cuda/level1/gemm/GEMM.cpp b/src/cuda/level1/gemm/GEMM.cpp index 00068e03..f2700ae1 100644 --- a/src/cuda/level1/gemm/GEMM.cpp +++ b/src/cuda/level1/gemm/GEMM.cpp @@ -228,7 +228,7 @@ void RunTest(string testName, ResultDatabase &resultDB, OptionParser &op) // Warm Up devGEMM(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; diff --git a/src/cuda/level1/md/MD.cu b/src/cuda/level1/md/MD.cu index e23e4b7a..390eaa0d 100644 --- a/src/cuda/level1/md/MD.cu +++ b/src/cuda/level1/md/MD.cu @@ -36,14 +36,11 @@ inline int populateNeighborList(std::list& currDist, std::list& currList, const int j, const int nAtom, int* neighborList); -// Texture caches for position info -texture posTexture; -texture posTexture_dp; - struct texReader_sp { + cudaTextureObject_t tex; __device__ __forceinline__ float4 operator()(int idx) const { - return tex1Dfetch(posTexture, idx); + return tex1Dfetch(tex, idx); } }; @@ -51,6 +48,7 @@ struct texReader_sp { // 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) @@ -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(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(tex, idx*2 + 1); double2 b = make_double2(__hiloint2double(v.y, v.x), __hiloint2double(v.w, v.z)); @@ -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; @@ -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) { @@ -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(); - - // Bind a 1D texture to the position array - CUDA_SAFE_CALL(cudaBindTexture(0, posTexture, d_position, channelDesc, - nAtom*sizeof(float4))); - - cudaChannelFormatDesc channelDesc2 = cudaCreateChannelDesc(); - - // 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() : cudaCreateChannelDesc(); + 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 @@ -403,8 +404,8 @@ void runTest(const string& testName, ResultDatabase& resultDB, OptionParser& op) compute_lj_force <<>> (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; @@ -446,7 +447,7 @@ void runTest(const string& testName, ResultDatabase& resultDB, OptionParser& op) compute_lj_force <<>> (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)); @@ -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)); diff --git a/src/cuda/level1/spmv/Spmv.cu b/src/cuda/level1/spmv/Spmv.cu index 6432e0f2..00e2662d 100644 --- a/src/cuda/level1/spmv/Spmv.cu +++ b/src/cuda/level1/spmv/Spmv.cu @@ -10,21 +10,20 @@ using namespace std; -texture vecTex; // vector textures -texture vecTexD; - // Texture Readers (used so kernels can be templated) struct texReaderSP { + cudaTextureObject_t tex; __device__ __forceinline__ float operator()(const int idx) const { - return tex1Dfetch(vecTex, idx); + return tex1Dfetch(tex, idx); } }; struct texReaderDP { + cudaTextureObject_t tex; __device__ __forceinline__ double operator()(const int idx) const { - int2 v = tex1Dfetch(vecTexD, idx); + int2 v = tex1Dfetch(tex, idx); #if (__CUDA_ARCH__ < 130) // Devices before arch 130 don't support DP, and having the // __hiloint2double() intrinsic will cause compilation to fail. @@ -43,21 +42,24 @@ __global__ void spmv_csr_scalar_kernel(const fpType * __restrict__ val, const int * __restrict__ cols, const int * __restrict__ rowDelimiters, - const int dim, fpType * __restrict__ out); + const int dim, fpType * __restrict__ out, + const texReader vecTexReader); template __global__ void spmv_csr_vector_kernel(const fpType * __restrict__ val, const int * __restrict__ cols, const int * __restrict__ rowDelimiters, - const int dim, fpType * __restrict__ out); + const int dim, fpType * __restrict__ out, + const texReader vecTexReader); template __global__ void spmv_ellpackr_kernel(const fpType * __restrict__ val, const int * __restrict__ cols, const int * __restrict__ rowLengths, - const int dim, fpType * __restrict__ out); + const int dim, fpType * __restrict__ out, + const texReader vecTexReader); template __global__ void @@ -216,17 +218,25 @@ void csrTest(ResultDatabase& resultDB, OptionParser& op, floatType* h_val, // Bind texture for position string suffix; + texReader vecTexReader = {}; + cudaResourceDesc resDesc = {}; + resDesc.resType = cudaResourceTypeLinear; + resDesc.res.linear.devPtr = d_vec; + cudaTextureDesc texDesc = {}; + texDesc.readMode = cudaReadModeElementType; if (sizeof(floatType) == sizeof(float)) { - cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc(); - CUDA_SAFE_CALL(cudaBindTexture(0, vecTex, d_vec, channelDesc, - numRows * sizeof(float))); + resDesc.res.linear.desc = cudaCreateChannelDesc(); + resDesc.res.linear.sizeInBytes = numRows * sizeof(float); + CUDA_SAFE_CALL(cudaCreateTextureObject(&vecTexReader.tex, &resDesc, + &texDesc, NULL)); suffix = "-SP"; } else { - cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc(); - CUDA_SAFE_CALL(cudaBindTexture(0, vecTexD, d_vec, channelDesc, - numRows * sizeof(int2))); + resDesc.res.linear.desc = cudaCreateChannelDesc(); + resDesc.res.linear.sizeInBytes = numRows * sizeof(int2); + CUDA_SAFE_CALL(cudaCreateTextureObject(&vecTexReader.tex, &resDesc, + &texDesc, NULL)); suffix = "-DP"; } @@ -252,7 +262,7 @@ void csrTest(ResultDatabase& resultDB, OptionParser& op, floatType* h_val, { spmv_csr_scalar_kernel <<>> - (d_val, d_cols, d_rowDelimiters, numRows, d_out); + (d_val, d_cols, d_rowDelimiters, numRows, d_out, vecTexReader); } CUDA_SAFE_CALL(cudaEventRecord(stop, 0)); CUDA_SAFE_CALL(cudaEventSynchronize(stop)); @@ -280,7 +290,7 @@ void csrTest(ResultDatabase& resultDB, OptionParser& op, floatType* h_val, gflop / (scalarKernelTime+totalTransfer)); } zero<<>>(d_out, numRows); - cudaThreadSynchronize(); + cudaDeviceSynchronize(); cout << "CSR Vector Kernel\n"; for (int k=0; k <<>> - (d_val, d_cols, d_rowDelimiters, numRows, d_out); + (d_val, d_cols, d_rowDelimiters, numRows, d_out, vecTexReader); } CUDA_SAFE_CALL(cudaEventRecord(stop, 0)); CUDA_SAFE_CALL(cudaEventSynchronize(stop)); @@ -299,7 +309,7 @@ void csrTest(ResultDatabase& resultDB, OptionParser& op, floatType* h_val, CUDA_SAFE_CALL(cudaEventElapsedTime(&vectorKernelTime, start, stop)); CUDA_SAFE_CALL(cudaMemcpy(h_out, d_out, numRows * sizeof(floatType), cudaMemcpyDeviceToHost)); - cudaThreadSynchronize(); + cudaDeviceSynchronize(); // Compare reference solution to GPU result if (! verifyResults(refOut, h_out, numRows, k)) { @@ -318,8 +328,7 @@ void csrTest(ResultDatabase& resultDB, OptionParser& op, floatType* h_val, CUDA_SAFE_CALL(cudaFree(d_out)); CUDA_SAFE_CALL(cudaFree(d_val)); CUDA_SAFE_CALL(cudaFree(d_cols)); - CUDA_SAFE_CALL(cudaUnbindTexture(vecTexD)); - CUDA_SAFE_CALL(cudaUnbindTexture(vecTex)); + CUDA_SAFE_CALL(cudaDestroyTextureObject(vecTexReader.tex)); CUDA_SAFE_CALL(cudaEventDestroy(start)); CUDA_SAFE_CALL(cudaEventDestroy(stop)); } @@ -377,17 +386,25 @@ void ellPackTest(ResultDatabase& resultDB, OptionParser& op, floatType* h_val, cmSize * sizeof(int), cudaMemcpyHostToDevice)); // Bind texture for position + texReader vecTexReader = {}; + cudaResourceDesc resDesc = {}; + resDesc.resType = cudaResourceTypeLinear; + resDesc.res.linear.devPtr = d_vec; + cudaTextureDesc texDesc = {}; + texDesc.readMode = cudaReadModeElementType; if (sizeof(floatType) == sizeof(float)) { - cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc(); - CUDA_SAFE_CALL(cudaBindTexture(0, vecTex, d_vec, channelDesc, - numRows * sizeof(float))); + resDesc.res.linear.desc = cudaCreateChannelDesc(); + resDesc.res.linear.sizeInBytes = numRows * sizeof(float); + CUDA_SAFE_CALL(cudaCreateTextureObject(&vecTexReader.tex, &resDesc, + &texDesc, NULL)); } else { - cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc(); - CUDA_SAFE_CALL(cudaBindTexture(0, vecTexD, d_vec, channelDesc, - numRows * sizeof(int2))); + resDesc.res.linear.desc = cudaCreateChannelDesc(); + resDesc.res.linear.sizeInBytes = numRows * sizeof(int2); + CUDA_SAFE_CALL(cudaCreateTextureObject(&vecTexReader.tex, &resDesc, + &texDesc, NULL)); } int nBlocks = (int) ceil((floatType) cmSize / BLOCK_SIZE); int passes = op.getOptionInt("passes"); @@ -401,7 +418,7 @@ void ellPackTest(ResultDatabase& resultDB, OptionParser& op, floatType* h_val, for (int j = 0; j < iters; j++) { spmv_ellpackr_kernel<<>> - (d_val, d_cols, d_rowLengths, cmSize, d_out); + (d_val, d_cols, d_rowLengths, cmSize, d_out, vecTexReader); } CUDA_SAFE_CALL(cudaEventRecord(stop, 0)); CUDA_SAFE_CALL(cudaEventSynchronize(stop)); @@ -433,14 +450,7 @@ void ellPackTest(ResultDatabase& resultDB, OptionParser& op, floatType* h_val, CUDA_SAFE_CALL(cudaFree(d_out)); CUDA_SAFE_CALL(cudaFree(d_val)); CUDA_SAFE_CALL(cudaFree(d_cols)); - if (sizeof(floatType) == sizeof(double)) - { - CUDA_SAFE_CALL(cudaUnbindTexture(vecTexD)); - } - else - { - CUDA_SAFE_CALL(cudaUnbindTexture(vecTex)); - } + CUDA_SAFE_CALL(cudaDestroyTextureObject(vecTexReader.tex)); CUDA_SAFE_CALL(cudaEventDestroy(start)); CUDA_SAFE_CALL(cudaEventDestroy(stop)); CUDA_SAFE_CALL(cudaFreeHost(h_rowLengths)); @@ -670,10 +680,10 @@ __global__ void spmv_csr_scalar_kernel(const fpType * __restrict__ val, const int * __restrict__ cols, const int * __restrict__ rowDelimiters, - const int dim, fpType * __restrict__ out) + const int dim, fpType * __restrict__ out, + const texReader vecTexReader) { int myRow = blockIdx.x * blockDim.x + threadIdx.x; - texReader vecTexReader; if (myRow < dim) { @@ -720,7 +730,8 @@ __global__ void spmv_csr_vector_kernel(const fpType * __restrict__ val, const int * __restrict__ cols, const int * __restrict__ rowDelimiters, - const int dim, fpType * __restrict__ out) + const int dim, fpType * __restrict__ out, + const texReader vecTexReader) { // Thread ID in block int t = threadIdx.x; @@ -729,8 +740,6 @@ spmv_csr_vector_kernel(const fpType * __restrict__ val, int warpsPerBlock = blockDim.x / warpSize; // One row per warp int myRow = (blockIdx.x * warpsPerBlock) + (t / warpSize); - // Texture reader for the dense vector - texReader vecTexReader; __shared__ volatile fpType partialSums[BLOCK_SIZE]; @@ -791,10 +800,10 @@ __global__ void spmv_ellpackr_kernel(const fpType * __restrict__ val, const int * __restrict__ cols, const int * __restrict__ rowLengths, - const int dim, fpType * __restrict__ out) + const int dim, fpType * __restrict__ out, + const texReader vecTexReader) { int t = blockIdx.x * blockDim.x + threadIdx.x; - texReader vecTexReader; if (t < dim) { diff --git a/src/cuda/level1/triad/Triad.cu b/src/cuda/level1/triad/Triad.cu index 3feaedea..7cfe59db 100644 --- a/src/cuda/level1/triad/Triad.cu +++ b/src/cuda/level1/triad/Triad.cu @@ -236,7 +236,7 @@ void RunBenchmark(ResultDatabase &resultDB, OptionParser &op) currStream = !currStream; } - cudaThreadSynchronize(); + cudaDeviceSynchronize(); double time = Timer::Stop(TH, "thread synchronize"); double triad = ((double)numMaxFloats * 2.0) / (time*1e9); diff --git a/src/cuda/level2/qtclustering/QTC.cu b/src/cuda/level2/qtclustering/QTC.cu index 441b1dd9..832c9d72 100644 --- a/src/cuda/level2/qtclustering/QTC.cu +++ b/src/cuda/level2/qtclustering/QTC.cu @@ -22,8 +22,6 @@ #include "comm.h" -texture texDistance; - using namespace std; #include "kernels_common.h" @@ -297,6 +295,7 @@ void QTC(const string& name, ResultDatabase &resultDB, OptionParser& op, int mat bool be_verbose = false; bool synthetic_data = true; cudaArray *distance_matrix_txt; + cudaTextureObject_t texDistance = 0; void *distance_matrix_gmem, *distance_matrix; float *dist_source, *pnts; float threshold; @@ -400,14 +399,11 @@ void QTC(const string& name, ResultDatabase &resultDB, OptionParser& op, int mat output = pmsAllocHostBuffer(max_degree); if( can_use_texture ){ - texDistance.addressMode[0] = cudaAddressModeClamp; - texDistance.addressMode[1] = cudaAddressModeClamp; - texDistance.filterMode = cudaFilterModePoint; - texDistance.normalized = false; // do not normalize coordinates // This is the actual distance matrix (dst_matrix_elems should be "point_count^2, or point_count*max_degree) printf("Allocating: %luMB (%lux%lux%lu) bytes in texture memory\n", dst_matrix_elems*sizeof(float)/(1024*1024), dst_matrix_elems/point_count, point_count, (long unsigned int)sizeof(float)); - cudaMallocArray(&distance_matrix_txt, &texDistance.channelDesc, dst_matrix_elems/point_count, point_count); + cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc(); + cudaMallocArray(&distance_matrix_txt, &channelDesc, dst_matrix_elems/point_count, point_count); }else{ allocDeviceBuffer(&distance_matrix_gmem, dst_matrix_elems*sizeof(float)); } @@ -430,7 +426,19 @@ void QTC(const string& name, ResultDatabase &resultDB, OptionParser& op, int mat if( can_use_texture ){ cudaMemcpyToArray(distance_matrix_txt, 0, 0, dist_source, dst_matrix_elems*sizeof(float), cudaMemcpyHostToDevice); CHECK_CUDA_ERROR(); - cudaBindTextureToArray(texDistance, distance_matrix_txt); + + cudaResourceDesc resDesc = {}; + resDesc.resType = cudaResourceTypeArray; + resDesc.res.array.array = distance_matrix_txt; + + cudaTextureDesc texDesc = {}; + texDesc.addressMode[0] = cudaAddressModeClamp; + texDesc.addressMode[1] = cudaAddressModeClamp; + texDesc.filterMode = cudaFilterModePoint; + texDesc.readMode = cudaReadModeElementType; + texDesc.normalizedCoords = false; // do not normalize coordinates + + cudaCreateTextureObject(&texDistance, &resDesc, &texDesc, NULL); }else{ copyToDevice(distance_matrix_gmem, dist_source, dst_matrix_elems*sizeof(float)); } @@ -444,7 +452,7 @@ void QTC(const string& name, ResultDatabase &resultDB, OptionParser& op, int mat tpb = ( point_count > THREADSPERBLOCK )? THREADSPERBLOCK : point_count; compute_degrees<<>>((int *)indr_mtrx, (int *)degrees, point_count, max_degree); - cudaThreadSynchronize(); + cudaDeviceSynchronize(); CHECK_CUDA_ERROR(); const char *sizeStr; @@ -512,10 +520,10 @@ void QTC(const string& name, ResultDatabase &resultDB, OptionParser& op, int mat (int *)indr_mtrx, (int *)cardnl, (int *)ungrpd_pnts_indr, (float *)dist_to_clust, (int *)degrees, point_count, max_point_count, max_degree, threshold, cwrank, active_node_count, - total_thread_block_count, matrix_type, can_use_texture); + total_thread_block_count, matrix_type, can_use_texture, texDistance); ///////// ----------------- Main kernel ----------------- ///////// //////////////////////////////////////////////////////////////////////////////////////////////// - cudaThreadSynchronize(); + cudaDeviceSynchronize(); CHECK_CUDA_ERROR(); t_krn += Timer::Stop(Tkernel, "Kernel Only"); @@ -523,7 +531,7 @@ void QTC(const string& name, ResultDatabase &resultDB, OptionParser& op, int mat if( thread_block_count > 1 ){ // We are reducing 128 numbers or less, so one thread should be sufficient. reduce_card_device<<>>((int *)cardnl, thread_block_count); - cudaThreadSynchronize(); + cudaDeviceSynchronize(); CHECK_CUDA_ERROR(); } @@ -548,8 +556,8 @@ void QTC(const string& name, ResultDatabase &resultDB, OptionParser& op, int mat trim_ungrouped_pnts_indr_array<<>>(winner_index, (int*)ungrpd_pnts_indr, (float*)distance_matrix, (int *)result, (char *)Ai_mask, (char *)clustered_pnts_mask, (int *)indr_mtrx, (int *)cardnl, (float *)dist_to_clust, (int *)degrees, - point_count, max_point_count, max_degree, threshold, matrix_type, can_use_texture ); - cudaThreadSynchronize(); + point_count, max_point_count, max_degree, threshold, matrix_type, can_use_texture, texDistance ); + cudaDeviceSynchronize(); CHECK_CUDA_ERROR(); t_trim += Timer::Stop(Ttrim, "Trim Only"); @@ -572,7 +580,7 @@ void QTC(const string& name, ResultDatabase &resultDB, OptionParser& op, int mat int Tupdt = Timer::Start(); update_clustered_pnts_mask<<>>((char *)clustered_pnts_mask, (char *)Ai_mask, max_point_count); - cudaThreadSynchronize(); + cudaDeviceSynchronize(); CHECK_CUDA_ERROR(); t_updt += Timer::Stop(Tupdt, "Update Only"); @@ -606,7 +614,7 @@ void QTC(const string& name, ResultDatabase &resultDB, OptionParser& op, int mat pmsFreeHostBuffer(indr_mtrx_host); if( can_use_texture ){ cudaFreeArray(distance_matrix_txt); - cudaUnbindTexture(texDistance); + cudaDestroyTextureObject(texDistance); }else{ freeDeviceBuffer(distance_matrix_gmem); } diff --git a/src/cuda/level2/qtclustering/kernels_common.h b/src/cuda/level2/qtclustering/kernels_common.h index 9a39916a..e77a7d97 100644 --- a/src/cuda/level2/qtclustering/kernels_common.h +++ b/src/cuda/level2/qtclustering/kernels_common.h @@ -12,13 +12,13 @@ #include "qtc_common.h" // Forward declarations -__global__ void QTC_device( float *dist_matrix, char *Ai_mask, char *clustered_pnts_mask, int *indr_mtrx, int *cluster_cardinalities, int *ungrpd_pnts_indr, float *dist_to_clust, int *degrees, int point_count, int N0, int max_degree, float threshold, int cwrank, int node_rank, int node_count, int total_thread_block_count, bool can_use_texture); +__global__ void QTC_device( float *dist_matrix, char *Ai_mask, char *clustered_pnts_mask, int *indr_mtrx, int *cluster_cardinalities, int *ungrpd_pnts_indr, float *dist_to_clust, int *degrees, int point_count, int N0, int max_degree, float threshold, int node_rank, int node_count, int total_thread_block_count, int matrix_type_mask, bool can_use_texture, cudaTextureObject_t texDistance); -__device__ int generate_candidate_cluster_compact_storage(int seed_point, int degree, char *Ai_mask, float *compact_storage_dist_matrix, char *clustered_pnts_mask, int *indr_mtrx, float *dist_to_clust, int point_count, int N0, int max_degree, int *candidate_cluster, float threshold, bool can_use_texture); +__device__ int generate_candidate_cluster_compact_storage(int seed_point, int degree, char *Ai_mask, float *compact_storage_dist_matrix, char *clustered_pnts_mask, int *indr_mtrx, float *dist_to_clust, int point_count, int N0, int max_degree, int *candidate_cluster, float threshold, bool can_use_texture, cudaTextureObject_t texDistance); -__device__ int generate_candidate_cluster_full_storage(int seed_point, int degree, char *Ai_mask, float *work, char *clustered_pnts_mask, int *indr_mtrx, float *dist_to_clust, int pointCount, int N0, int max_degree, int *candidate_cluster, float threshold, bool can_use_texture); +__device__ int generate_candidate_cluster_full_storage(int seed_point, int degree, char *Ai_mask, float *work, char *clustered_pnts_mask, int *indr_mtrx, float *dist_to_clust, int pointCount, int N0, int max_degree, int *candidate_cluster, float threshold, bool can_use_texture, cudaTextureObject_t texDistance); -__device__ int find_closest_point_to_cluster(int seed_point, int latest_point, char *Ai_mask, char *clustered_pnts_mask, float *work, int *indr_mtrx, float *dist_to_clust, int pointCount, int N0, int max_degree, float threshold); +__device__ int find_closest_point_to_cluster(int seed_point, int latest_point, char *Ai_mask, char *clustered_pnts_mask, float *work, int *indr_mtrx, float *dist_to_clust, int pointCount, int N0, int max_degree, float threshold, bool can_use_texture, cudaTextureObject_t texDistance); void QTC(const string& name, ResultDatabase &resultDB, OptionParser& op, int matrix_type); @@ -138,7 +138,7 @@ update_clustered_pnts_mask(char *clustered_pnts_mask, char *Ai_mask, int N0 ) { __global__ void -trim_ungrouped_pnts_indr_array(int seed_index, int *ungrpd_pnts_indr, float *dist_matrix, int *result_cluster, char *Ai_mask, char *clustered_pnts_mask, int *indr_mtrx, int *cluster_cardinalities, float *dist_to_clust, int *degrees, int point_count, int N0, int max_degree, float threshold, int matrix_type_mask, bool can_use_texture) { +trim_ungrouped_pnts_indr_array(int seed_index, int *ungrpd_pnts_indr, float *dist_matrix, int *result_cluster, char *Ai_mask, char *clustered_pnts_mask, int *indr_mtrx, int *cluster_cardinalities, float *dist_to_clust, int *degrees, int point_count, int N0, int max_degree, float threshold, int matrix_type_mask, bool can_use_texture, cudaTextureObject_t texDistance) { int cnt; int tid = threadIdx.x; int curThreadCount = blockDim.x*blockDim.y*blockDim.z; @@ -147,10 +147,10 @@ trim_ungrouped_pnts_indr_array(int seed_index, int *ungrpd_pnts_indr, float *dis int degree = degrees[seed_index]; if( matrix_type_mask & FULL_STORAGE_MATRIX ){ (void)generate_candidate_cluster_full_storage(seed_index, degree, Ai_mask, dist_matrix, clustered_pnts_mask, indr_mtrx, - dist_to_clust, point_count, N0, max_degree, result_cluster, threshold, can_use_texture); + dist_to_clust, point_count, N0, max_degree, result_cluster, threshold, can_use_texture, texDistance); }else{ (void)generate_candidate_cluster_compact_storage(seed_index, degree, Ai_mask, dist_matrix, clustered_pnts_mask, indr_mtrx, - dist_to_clust, point_count, N0, max_degree, result_cluster, threshold, can_use_texture); + dist_to_clust, point_count, N0, max_degree, result_cluster, threshold, can_use_texture, texDistance); } __shared__ int cnt_sh; @@ -201,7 +201,7 @@ trim_ungrouped_pnts_indr_array(int seed_index, int *ungrpd_pnts_indr, float *dis __global__ -void QTC_device( float *dist_matrix, char *Ai_mask, char *clustered_pnts_mask, int *indr_mtrx, int *cluster_cardinalities, int *ungrpd_pnts_indr, float *dist_to_clust, int *degrees, int point_count, int N0, int max_degree, float threshold, int node_rank, int node_count, int total_thread_block_count, int matrix_type_mask, bool can_use_texture) { +void QTC_device( float *dist_matrix, char *Ai_mask, char *clustered_pnts_mask, int *indr_mtrx, int *cluster_cardinalities, int *ungrpd_pnts_indr, float *dist_to_clust, int *degrees, int point_count, int N0, int max_degree, float threshold, int node_rank, int node_count, int total_thread_block_count, int matrix_type_mask, bool can_use_texture, cudaTextureObject_t texDistance) { int max_cardinality = -1; int max_cardinality_index; int i, tblock_id, tid, base_offset; @@ -224,12 +224,12 @@ void QTC_device( float *dist_matrix, char *Ai_mask, char *clustered_pnts_mask, i cnt = generate_candidate_cluster_full_storage(seed_index, degree, Ai_mask, dist_matrix, clustered_pnts_mask, indr_mtrx, dist_to_clust, point_count, N0, max_degree, NULL, threshold, - can_use_texture); + can_use_texture, texDistance); }else{ cnt = generate_candidate_cluster_compact_storage( seed_index, degree, Ai_mask, dist_matrix, clustered_pnts_mask, indr_mtrx, dist_to_clust, point_count, N0, max_degree, NULL, threshold, - can_use_texture); + can_use_texture, texDistance); } if( cnt > max_cardinality ){ max_cardinality = cnt; diff --git a/src/cuda/level2/qtclustering/kernels_compact_storage.h b/src/cuda/level2/qtclustering/kernels_compact_storage.h index 98e7d962..041d7ab9 100644 --- a/src/cuda/level2/qtclustering/kernels_compact_storage.h +++ b/src/cuda/level2/qtclustering/kernels_compact_storage.h @@ -26,7 +26,7 @@ do{\ }\ if( tmp_pnt == (_CAND_PNT_) ){\ if( can_use_texture ){ \ - dist_to_new_point = tex2D(texDistance, float(j)+0.5f, float(latest_point)+0.5f );\ + dist_to_new_point = tex2D(texDistance, float(j)+0.5f, float(latest_point)+0.5f );\ }else{\ dist_to_new_point = compact_storage_dist_matrix[ latest_p_off + j ];\ }\ @@ -65,7 +65,7 @@ do{\ } inline __device__ -int generate_candidate_cluster_compact_storage(int seed_point, int degree, char *Ai_mask, float *compact_storage_dist_matrix, char *clustered_pnts_mask, int *indr_mtrx, float *dist_to_clust, int point_count, int N0, int max_degree, int *candidate_cluster, float threshold, bool can_use_texture) +int generate_candidate_cluster_compact_storage(int seed_point, int degree, char *Ai_mask, float *compact_storage_dist_matrix, char *clustered_pnts_mask, int *indr_mtrx, float *dist_to_clust, int point_count, int N0, int max_degree, int *candidate_cluster, float threshold, bool can_use_texture, cudaTextureObject_t texDistance) { bool flag; diff --git a/src/cuda/level2/qtclustering/kernels_full_storage.h b/src/cuda/level2/qtclustering/kernels_full_storage.h index 22edb231..61d6158f 100644 --- a/src/cuda/level2/qtclustering/kernels_full_storage.h +++ b/src/cuda/level2/qtclustering/kernels_full_storage.h @@ -2,7 +2,7 @@ #define _KERNELS_FULL_STORAGE_H_ inline __device__ -int find_closest_point_to_cluster(int seed_point, int latest_point, char *Ai_mask, char *clustered_pnts_mask, float *work, int *indr_mtrx, float *dist_to_clust, int point_count, int N0, int max_degree, float threshold, bool can_use_texture){ +int find_closest_point_to_cluster(int seed_point, int latest_point, char *Ai_mask, char *clustered_pnts_mask, float *work, int *indr_mtrx, float *dist_to_clust, int point_count, int N0, int max_degree, float threshold, bool can_use_texture, cudaTextureObject_t texDistance){ int point_index = -1; float min_dist=2*threshold; @@ -21,7 +21,7 @@ int find_closest_point_to_cluster(int seed_point, int latest_point, char *Ai_mas float curr_dist_to_clust = dist_to_clust[i+tid]; if( can_use_texture ){ - dist_to_new_point = tex2D(texDistance, float(latest_point)+0.5f, float(pnt2_indx)+0.5f ); + dist_to_new_point = tex2D(texDistance, float(latest_point)+0.5f, float(pnt2_indx)+0.5f ); }else{ dist_to_new_point = work[pnt2_indx*N0 + latest_point]; } @@ -56,7 +56,7 @@ do{ \ \ float curr_dist_to_clust = _CAND_PNT_DIST_; \ if( can_use_texture ){\ - dist_to_new_point = tex2D(texDistance, float(latest_point)+0.5f, float( (_CAND_PNT_) )+0.5f ); \ + dist_to_new_point = tex2D(texDistance, float(latest_point)+0.5f, float( (_CAND_PNT_) )+0.5f ); \ }else{\ dist_to_new_point = work[ (_CAND_PNT_) * N0 + latest_point]; \ }\ @@ -79,7 +79,7 @@ do{ \ inline __device__ -int generate_candidate_cluster_full_storage(int seed_point, int degree, char *Ai_mask, float *work, char *clustered_pnts_mask, int *indr_mtrx, float *dist_to_clust, int point_count, int N0, int max_degree, int *candidate_cluster, float threshold, bool can_use_texture) +int generate_candidate_cluster_full_storage(int seed_point, int degree, char *Ai_mask, float *work, char *clustered_pnts_mask, int *indr_mtrx, float *dist_to_clust, int point_count, int N0, int max_degree, int *candidate_cluster, float threshold, bool can_use_texture, cudaTextureObject_t texDistance) { bool flag; int cnt, latest_point; @@ -194,7 +194,7 @@ int generate_candidate_cluster_full_storage(int seed_point, int degree, char *Ai while( (cnt < point_count) && flag ){ int min_G_index; min_G_index = find_closest_point_to_cluster(seed_point, latest_point, Ai_mask, clustered_pnts_mask, work, - indr_mtrx, dist_to_clust, point_count, N0, max_degree, threshold, can_use_texture); + indr_mtrx, dist_to_clust, point_count, N0, max_degree, threshold, can_use_texture, texDistance); if(min_G_index >= 0 ){ if( 0 == tid ){ Ai_mask[min_G_index] = 1; diff --git a/src/mpi/contention-mt/cuda/CUDADriver.cpp b/src/mpi/contention-mt/cuda/CUDADriver.cpp index 84f3f167..d60e22a1 100644 --- a/src/mpi/contention-mt/cuda/CUDADriver.cpp +++ b/src/mpi/contention-mt/cuda/CUDADriver.cpp @@ -51,7 +51,13 @@ void EnumerateDevicesAndChoose(int chooseDevice, bool verbose, const char* prefi cout << " totalConstMem = "<