该运行时API的作用
作为在访问共享内存时作为线程块内的同步机制出现,保证同一线程块内所有线程到程序运行到这个运行时API调用时都能运行完毕(注意,该API不能同步不同线程块内的线程),例如下列Cuda静态共享内存使用代码示例程序中的第23行所示:
/**************************************************************************************
* file name : static_shared_memory.cu
* author : 权 双
* data : 2023-12-26
* brief : 静态共享内存使用
***************************************************************************************/
#include <cuda_runtime.h>
#include <iostream>
#include "common.cuh"
__global__ void kernel_1(float* d_A, const int N)
{
const int tid = threadIdx.x;
const int bid = blockIdx.x;
const int n = bid * blockDim.x + tid;
__shared__ float s_array[32];
if (n < N)
{
s_array[tid] = d_A[n];
}
__syncthreads();
if (tid == 0)
{
for (int i = 0; i < 32; ++i)
{
printf("kernel_1: %f, blockIdx: %d\n", s_array[i], bid);
}
}
}
int main(int argc, char **argv)
{
int devID = 0;
cudaDeviceProp deviceProps;
CUDA_CHECK(cudaGetDeviceProperties(&deviceProps, devID));
std::cout << "运行GPU设备:" << deviceProps.name << std::endl;
int nElems = 64;
int nbytes = nElems * sizeof(float);
float* h_A = nullptr;
h_A = (float*)malloc(nbytes);
for (int i = 0; i < nElems; ++i)
{
h_A[i] = float(i);
}
float* d_A = nullptr;
CUDA_CHECK(cudaMalloc(&d_A, nbytes));
CUDA_CHECK(cudaMemcpy(d_A, h_A, nbytes,cudaMemcpyHostToDevice));
dim3 block(32);
dim3 grid(2);
kernel_1<<<grid, block>>>(d_A, nElems);
CUDA_CHECK(cudaFree(d_A));
free(h_A);
CUDA_CHECK(cudaDeviceReset());
}
GitHubID:shiyi23
践行 活在当下 的理念