lukasplevac/cuda-samples
1
0
Fork 0
mirror of https://github.com/NVIDIA/cuda-samples.git synced 2025-04-11 02:32:12 +01:00
Code Issues Packages Projects Releases Wiki Activity

Merge 6aac4717b878ce3a78c1159ff2794e681a914e5b into 8d564d5e3afdab5dda868f42a13d85f3d0e75bc9

Browse Source
This commit is contained in:
li wei 2025-02-20 17:35:32 +01:00 committed by GitHub
commit 601384af49
No known key found for this signature in database
GPG Key ID: B5690EEEBB952194
1 changed files with 3 additions and 3 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

6
Samples/3_CUDA_Features/globalToShmemAsyncCopy/globalToShmemAsyncCopy.cu
View File

@ -128,7 +128,7 @@ __global__ void MatrixMulAsyncCopyMultiStageLargeChunk(
a <= aEnd; a += aStep, b += bStep, ++i) {
// Load the matrices from device memory to shared memory; each thread loads
// one element of each matrix
for (; aStage <= a + aStep * maxPipelineStages;
for (; aStage < a + aStep * maxPipelineStages;
aStage += aStep, bStage += bStep, ++iStage) {
pipe.producer_acquire();
if (aStage <= aEnd && t4x < BLOCK_SIZE) {
@ -137,7 +137,7 @@ __global__ void MatrixMulAsyncCopyMultiStageLargeChunk(
cuda::memcpy_async(&As[j][threadIdx.y][t4x],
&A[aStage + wA * threadIdx.y + t4x], shape4, pipe);
cuda::memcpy_async(&Bs[j][threadIdx.y][t4x],
&B[aStage + wA * threadIdx.y + t4x], shape4, pipe);
&B[bStage + wB * threadIdx.y + t4x], shape4, pipe);
}
pipe.producer_commit();
}
@ -222,7 +222,7 @@ __global__ void MatrixMulAsyncCopyLargeChunk(float *__restrict__ C,
cuda::memcpy_async(&As[threadIdx.y][t4x], &A[a + wA * threadIdx.y + t4x],
shape4, pipe);
cuda::memcpy_async(&Bs[threadIdx.y][t4x], &B[a + wA * threadIdx.y + t4x],
cuda::memcpy_async(&Bs[threadIdx.y][t4x], &B[b + wB * threadIdx.y + t4x],
shape4, pipe);
pipe.producer_commit();