{"id":218,"date":"2024-10-17T23:08:55","date_gmt":"2024-10-17T15:08:55","guid":{"rendered":"https:\/\/www.qibinliang.com\/?p=218"},"modified":"2024-10-19T15:22:24","modified_gmt":"2024-10-19T07:22:24","slug":"cuda%e7%bc%96%e7%a8%8b%e5%ad%a6%e4%b9%a0%e7%ac%94%e8%ae%b0%ef%bc%88%e4%b8%80%ef%bc%89","status":"publish","type":"post","link":"https:\/\/www.qibinliang.com\/index.php\/2024\/10\/17\/cuda%e7%bc%96%e7%a8%8b%e5%ad%a6%e4%b9%a0%e7%ac%94%e8%ae%b0%ef%bc%88%e4%b8%80%ef%bc%89\/","title":{"rendered":"Cuda\u7f16\u7a0b\u5b66\u4e60\u7b14\u8bb0\uff08\u4e00\uff09"},"content":{"rendered":"\n

\u6700\u8fd1\u95f2\u6687\u65f6\u95f4\u60f3\u627e\u70b9\u4e8b\u60c5\u505a\uff0c\u5c31\u60f3\u8981\u4e0d\u5b66\u4e60Cuda\u7f16\u7a0b\u5427\uff0c\u5305\u597d\u73a9\u7684\u3002\u7136\u540e\uff0c\u5728Google\u4e0a\u627eCuda\u5b66\u4e60\u8d44\u6599\u65f6\u53d1\u73b0\u4e86\u5f88\u591a\u4eba\u90fd\u63a8\u8350Oakland National Laboratory\u7684\u89c6\u9891\u4ee5\u53ca\u7ec3\u4e60\u3002\u4e8e\u662f\uff0c\u51b3\u5b9a\u5728\u535a\u5ba2\u4e0a\u9762\u505a\u4e2a\u5b66\u4e60\u7b14\u8bb0\uff0c\u6253\u53d1\u4e0b\u65e0\u804a\u65f6\u5149\uff0c\u987a\u4fbf\u7ec3\u4e60\u4e0b\u8bb8\u4e45\u672a\u7528\u7684\u82f1\u8bed\u3002<\/p>\n\n\n\n

\u4ee5\u4e0b\u662f\u8bfe\u7a0b\u94fe\u63a5\uff1a<\/p>\n\n\n\n

\nhttps:\/\/www.olcf.ornl.gov\/cuda-training-series\n<\/div><\/figure>\n\n\n\n

\u8bfe\u7a0b\u5185\u5bb9(Introduction to CUDA C++)<\/h2>\n\n\n\n

1. \u540d\u79f0\u5b9a\u4e49<\/h3>\n\n\n\n

\u5728CUDA\u7f16\u7a0b\u91cc\u901a\u5e38\u628aGPU\u53eb\u505aDevice\uff0cCPU\u53eb\u505aHost\uff0c\u56e0\u6b64GPU\u7684memory\u4e5f\u53eb\u505aDevice memroy\uff0cCPU\u7684memory\u4e5f\u53eb\u505aHost memory\u3002\u800cDevice Code\u548cHost Code\u5206\u522b\u662f\u8fd0\u884c\u5728GPU\u4ee5\u53caCPU\u4e0a\u7684\u4ee3\u7801\u3002\u4f46\u8fd9\u5e76\u4e0d\u610f\u5473\u7740GPU\u7684\u4ee3\u7801\u548cCPU\u7684\u662f\u5206\u79bb\u7684\uff0c\u4ed6\u4eec\u901a\u5e38\u90fd\u662f\u5728\u4e00\u8d77\u8fd0\u4f5c\u7684\u3002GPU\u4e3b\u8981\u8868\u73b0\u5728\u5f3a\u5927\u7684\u8ba1\u7b97\u80fd\u529b\uff0cCPU\u9700\u8981\u8d1f\u8d23\u4f8b\u5982\u78c1\u76d8IO\uff0c\u7528\u6237IO\uff0c\u7f51\u7edc\u7b49\uff0c\u4f46GPU\u5e76\u4e0d\u8d1f\u8d23\u8fd9\u4e9b\u3002<\/p>\n\n\n

\n
\"\"<\/div>
\u56fe1. CPU\u548cGPU\u7684\u522b\u79f0<\/figcaption><\/figure>\n<\/div>\n\n\n

2. \u4e09\u6b65\u6d41\u7a0b<\/h3>\n\n\n\n

\u7531\u4e8eCPU\u548cGPU\u662f\u5b8c\u5168\u4e0d\u77e5\u9053\u5bf9\u65b9\u5728\u5e72\u4ec0\u4e48\u7684\uff0c\u6211\u4eec\u5728CPU\u4ee3\u7801\u4e0a\u5b9a\u4e49\u7684\u6570\u636eGPU\u662f\u4e0d\u53ef\u89c1\u7684\uff0c\u56e0\u6b64\u6211\u4eec\u9700\u8981\u4e00\u4e9b\u64cd\u4f5c\u53bb\u540c\u6b65\u4e24\u4e2a\u8bbe\u5907\u7684\u6570\u636e\u3002\u5bf9\u6b64\uff0c\u6211\u4eec\u53ef\u4ee5\u7528\u4e00\u4e2a\u7b80\u5355\u7684\u4e09\u6b65\u6d41\u7a0b\u53bb\u63cf\u8ff0CUDA\u7f16\u7a0b\u6240\u9700\u8981\u7684\u6b65\u9aa4\uff1a<\/p>\n\n\n\n

    \n
  1. \u5c06\u6570\u636e\u4eceCPU\u5185\u5b58\u590d\u5236\u5230GPU\u5185\u5b58\uff08\u56fe2\uff09<\/li>\n\n\n\n
  2. \u52a0\u8f7dGPU\u7a0b\u5e8f\u5e76\u6267\u884c\uff08\u56fe3\uff09<\/li>\n\n\n\n
  3. \u5c06\u7ed3\u679c\u4eceGPU\u5185\u5b58\u590d\u5236\u5230CPU\u5185\u5b58\uff08\u56fe4\uff09<\/li>\n<\/ol>\n\n\n
    \n
    \"\"<\/div>
    \u56fe2. \u4e09\u6b65\u6d41\u7a0b\u5176\u4e00<\/figcaption><\/figure>\n<\/div>\n\n
    \n
    \"\"<\/div>
    \u56fe3. \u4e09\u6b65\u6d41\u7a0b\u5176\u4e8c<\/figcaption><\/figure>\n<\/div>\n\n
    \n
    \"\"<\/div>
    \u56fe4. \u4e09\u6b65\u6d41\u7a0b\u5176\u4e09<\/figcaption><\/figure>\n<\/div>\n\n\n

    3. \u51fd\u6570\u5b9a\u4e49\u4ee5\u53ca\u8c03\u7528<\/h3>\n\n\n\n

    \u5728CUDA\u7f16\u7a0b\u4e2d\u901a\u5e38\u51fd\u6570\u90fd\u4f1a\u5e26\u6709__global__<\/code>\u3001__host__<\/code>\u6216__device__<\/code>\u7684\u5173\u952e\u5b57\uff0c\u5e76\u4e14\u8fd4\u56de\u503c\u4e3avoid\uff08\u56fe5\uff09\u3002\u901a\u5e38\u6211\u4eec\u4f1a\u4f7f\u7528nvcc\u6765\u7f16\u8bd1\u6211\u4eec\u7684\u4ee3\u7801\uff0c\u800cnvcc\u4f1a\u628a\u6e90\u7801\u5206\u6210host\u548cdevice\u4e24\u4e2a\u90e8\u5206\uff0c\u5206\u522b\u5728\u82f1\u4f1f\u8fbe\u7684\u7f16\u8bd1\u5668\u548cc\/c++\u7f16\u8bd1\u5668\u4e0a\u5904\u7406\u3002\u5728\u8c03\u7528\u51fd\u6570\u7684\u65f6\u5019\u6211\u4eec\u9700\u8981\u7528\u5230\u4e00\u4e2aCUDA\u4e2d\u7684\u8bed\u6cd5\u2014\u2014\u201c\u4e09\u91cd\u5c16\u62ec\u53f7\u201d\uff0c\u4e5f\u53eb\u505a\u5185\u6838\u542f\u52a8\uff08kernel launch\uff09\uff08\u56fe6\uff09\uff0c\u5c16\u62ec\u53f7\u5185\u7684\u6570\u5b57\u4ee3\u8868\u4e86\u542f\u52a8\u7684\u6838\u5fc3\u6570\u4ee5\u53ca\u7ebf\u7a0b\u6570\uff0c\u4ee3\u8868\u7740\u6211\u4eec\u7684\u7a0b\u5e8f\u5e76\u884c\u5730\u8fd0\u884c\u5728\u4e86\u591a\u5c11\u4e2a\u6838\u5fc3\u5185\u3002\u5982\u679c\u6211\u4eec\u6ca1\u6709\u4f7f\u7528\u8fd9\u4e2a\u8bed\u6cd5\uff0c\u800c\u662f\u50cfC\/C++\u4e00\u6837\u76f4\u63a5\u8c03\u7528\u51fd\u6570\uff0c\u90a3\u4e48nvcc\u5c31\u4f1a\u7f16\u8bd1\u62a5\u9519\u3002<\/p>\n\n\n

    \n
    \"\"<\/div>
    \u56fe5. mykernel\u51fd\u6570\u5b9a\u4e49<\/figcaption><\/figure>\n<\/div>\n\n
    \n
    \"\"<\/div>
    \u56fe6. mykernel\u51fd\u6570\u8c03\u7528<\/figcaption><\/figure>\n<\/div>\n\n\n

    4. \u5185\u5efa\u53d8\u91cf\uff08Built-in Variable\uff09<\/h3>\n\n\n\n

    \u9664\u4e86__global__<\/code>\uff0c__host__<\/code>\uff0c__device__<\/code>\u8fd9\u51e0\u4e2a\u5185\u5efa\u53d8\u91cf\uff0cCUDA C++\u63a5\u53e3\u8fd8\u7ed9\u6211\u4eec\u63d0\u4f9b\u4e86\u5176\u4ed6\u7684\u5185\u5efa\u53d8\u91cf\u3002\u5728\u5e76\u884c\u8ba1\u7b97\u7684\u65f6\u5019\u6211\u4eec\u9700\u8981\u77e5\u9053\u6709\u591a\u5c11\u7684\u7ebf\u7a0b\u5728\u5de5\u4f5c\u4ee5\u53ca\u7ebf\u7a0b\u7684id\uff0c\u8fd9\u6837\u53ef\u4ee5\u65b9\u4fbf\u6211\u4eec\u5206\u914d\u6570\u636e\u4ee5\u53ca\u8bbe\u8ba1\u7b97\u6cd5\u3002\u90e8\u5206\u89e3\u91ca\u53c2\u8003\u4e86\u5b98\u65b9\u6587\u6863<\/mark><\/strong><\/a>\u3002<\/p>\n\n\n\n

    threadIdx<\/summary>\n

    threadIdx<\/strong><\/a>\u662f\u4e00\u4e2a\u4e09\u5206\u91cfvector\uff0c\u6240\u4ee5threadIdx\u53ef\u4ee5\u88ab\u7528\u4f5c\u4e0e\u4e00\u7ef4\u3001\u4e8c\u7ef4\u4ee5\u53ca\u4e09\u7ef4\u7ebf\u7a0b\u5757\u7d22\u5f15\uff0c\u901a\u8fc7threadIdx\u6211\u4eec\u83b7\u53d6\u5f53\u524d\u7ebf\u7a0b\u5757\u5185\u8fd0\u884c\u7684\u7ebf\u7a0bid\u3002\u8b6c\u5982\u4ee5\u4e0b\u6837\u4f8b\uff0c\u4e8c\u7ef4\u77e9\u9635\u7684\u52a0\u6cd5\uff1a<\/p>\n\n\n\n

    __global__ void MatAdd(float A[N][N], float B[N][N],\n                       float C[N][N])\n{\n    int i = threadIdx.x;\n    int j = threadIdx.y;\n    C[i][j] = A[i][j] + B[i][j];\n}<\/code><\/code><\/pre>\n<\/details>\n\n\n\n
    blockIdx<\/summary>\n
    blockIdx\u4e0ethreadIdx\u7c7b\u4f3c\u4e5f\u662f\u7528\u6765\u5bf9\u7ebf\u7a0b\u5757\u7684\u7d22\u5f15\u7684\uff0c\u540c\u6837\u7684blockIdx\u4e5f\u662f\u4e00\u4e2a\u4e09\u5206\u91cfvector\uff0c\u53ef\u4ee5\u8868\u793a\u4e00\u5230\u4e09\u7ef4\u7684\u7f51\u683c\uff08grid\uff09<\/p>\n<\/details>\n\n\n\n
    \n
    blockDim<\/summary>\n
    blockDim\u53ef\u4ee5\u83b7\u53d6\u7ebf\u7a0b\u5757\u7684\u7ef4\u5ea6\uff0c\u5373\u4e00\u4e2a\u7ebf\u7a0b\u5757\u6709\u591a\u5c11\u6761\u7ebf\u7a0b\uff0c\u7ed3\u5408blockIdx\uff0cblockDim\u4ee5\u53cathreadIdx\uff0c\u6211\u4eec\u53ef\u4ee5\u5199\u4e00\u4e2a\u66f4\u52a0\u201c\u5e76\u884c\u201d\u7684\u77e9\u9635\u52a0\u6cd5\uff1a<\/p>\n\n\n\n
    \/\/ Kernel definition\n__global__ void MatAdd(float A[N][N], float B[N][N],\nfloat C[N][N])\n{\n    int i = blockIdx.x * blockDim.x + threadIdx.x;\n    int j = blockIdx.y * blockDim.y + threadIdx.y;\n    if (i < N && j < N)\n        C[i][j] = A[i][j] + B[i][j];\n}\n\nint main()\n{\n    ...\n    \/\/ Kernel invocation\n    dim3 threadsPerBlock(16, 16);\n    dim3 numBlocks(N \/ threadsPerBlock.x, N \/ threadsPerBlock.y);\n    MatAdd<<<numBlocks, threadsPerBlock>>>(A, B, C);\n    ...\n}<\/code><\/pre>\n<\/details>\n<\/div><\/div>\n\n\n\n
    \u8bfe\u7a0b\u7ec3\u4e60<\/h2>\n\n\n\n
    Oakland National Laboratory\u7684\u8bfe\u7a0b\u8fd8\u63d0\u4f9b\u4e86\u76f8\u5e94\u7684\u7ec3\u4e60\uff0c\u4ee3\u7801\u5728github<\/mark><\/strong><\/a>\u4e0a\u53ef\u4ee5\u627e\u5230\u3002\u6211\u6240\u7528\u7684\u662fAMD\u7684\u663e\u5361\uff0c\u867d\u7136\u786c\u4ef6\u67b6\u6784\u4e0d\u540c\uff0c\u4f46\u662fAMD\u5bf9\u81ea\u5bb6\u7684\u5f02\u6784\u8ba1\u7b97\u63a5\u53e3\uff08hip\uff09\u4e0eNvidia\u7684CUDA\u5e93\u505a\u4e86\u5927\u9762\u79ef\u7684\u63a5\u53e3\u6620\u5c04\uff0c\u6240\u4ee5\u662f\u9700\u8981\u7a0d\u4f5c\u4fee\u6539\u5373\u53ef\u5728AMD GPU\u4e0a\u8fd0\u884c\u3002<\/p>\n\n\n\n
    hw1 – hello<\/summary>\n
    \u8c03\u7528GPU\u5b9e\u73b0\u6253\u5370\u4e00\u4e2a\u7b80\u5355\u7684hello world\uff0c\u8fd9\u4e2a\u6837\u4f8b\u53ef\u4ee5\u6d4b\u8bd5\u4e0d\u540c\u5927\u5c0f\u7684block\u4ee5\u53cathread\u4f1a\u5bf9\u7ed3\u679c\u6709\u4ec0\u4e48\u5f71\u54cd\u3002<\/p>\n\n\n\n
    #include <cstdlib>\n#include <stdio.h>\n#include \"hip\/hip_runtime.h\"\n\n__global__ void hello(){\n  printf(\"Hello from block: %u, thread: %u\\n\", blockIdx.x, threadIdx.x);\n}\n\nint main(int argc, char *argv[]){\n  int griddim = atoi(argv[1]);\n  int blockdim = atoi(argv[2]);\n  hello<<<griddim, blockdim>>>();\n}<\/code><\/pre>\n\n\n\n
    \u4f7f\u7528AMD\u7684hipcc\u7f16\u8bd1\u7a0b\u5e8f\uff0c\u5e76\u4f20\u5165\u4e0d\u540c\u53c2\u6570\u8fd0\u884c\uff1a<\/p>\n\n\n\n
    (base) qibin:~\/projects\/cuda-training-series\/exercises\/hw1$ hipcc -o hello  hello.cu <\/mark>\n(base) qibin:~\/projects\/cuda-training-series\/exercises\/hw1$ .\/hello 1 1<\/mark>\nHello from block: 0, thread: 0<\/mark>\n(base) qibin:~\/projects\/cuda-training-series\/exercises\/hw1$ .\/hello 2 1<\/mark>\nHello from block: 1, thread: 0\nHello from block: 0, thread: 0<\/mark>\n(base) qibin:~\/projects\/cuda-training-series\/exercises\/hw1$ .\/hello 2 2<\/mark>\nHello from block: 0, thread: 0\nHello from block: 0, thread: 1\nHello from block: 1, thread: 0\nHello from block: 1, thread: 1<\/mark><\/code><\/pre>\n\n\n\n
    \u4ece\u7ed3\u679c\u770b\u4f20\u5165\u4e0d\u540c\u7684blockdim*griddim<\/strong>\u5c31\u662f\u603b\u5171\u5728GPU\u4e0a\u8fd0\u884c\u7684\u7ebf\u7a0b\u6570<\/p>\n<\/details>\n\n\n\n
    hw1 – vector_add<\/summary>\n
    \u4f7f\u7528\u672c\u7ae0\u8bfe\u7a0b\u8bb2\u5230\u7684\u5185\u5bb9\u5b9e\u73b0\u7b80\u5355\u7684\u5411\u91cf\u52a0\u6cd5\u3002hipMalloc<\/code>\uff0chipMemcpy<\/code>\u76f8\u5f53\u4e8eC\/C++\u4e2d\u7684Malloc<\/code>\u548cMemcpy<\/code>\uff0c\u4e0d\u540c\u7684\u7684\u662f\u5206\u914d\u5185\u5b58\u4ee5\u53ca\u62f7\u8d1d\u5185\u5b58\u662f\u5728GPU\u4e0a\u64cd\u4f5c\u7684\uff0c\u800chipMemcpyHostToDevice\u4ee5\u53cahipMemcpyDeviceToHost\u5219\u662fCUDA\/hip\u4e2d\u7684\u4e00\u4e2a\u679a\u4e3e\u7c7b\uff0c\u5b9a\u4e49\u4e86\u6570\u636e\u7684\u4f20\u8f93\u65b9\u5411\uff0c\u4ececpu\u5230gpu\u6216\u8005\u76f8\u53cd\u3002\u8be6\u60c5\u53ef\u89c1\u5b98\u65b9\u6587\u6863cudaMalloc<\/a>\/cudaMemcpy<\/a>\/cudaMemcpyKind<\/a>\u3002
    \u4ee5\u4e0b\u4ee3\u7801\u5219\u4e09\u6b65\u6d41\u7a0b<\/strong>\u4e2d\u6240\u8bf4\u7684\u4e09\u6b65\u5904\u7406\u6d41\u7a0b\u5b8c\u5168\u76f8\u540c\u3002\u6a59\u8272<\/mark>\u9ad8\u4eae\u4ee3\u7801\u5bf9\u5e94\u7b2c\u4e00\u6b65\uff0cgpu\u4e0a\u5206\u914d\u6570\u636e\u5185\u5b58\u5e76\u5c06cpu\u4e0a\u7684\u6570\u636e\u590d\u5236\u5230gpu\u4e0a\u3002\u7b2c\u4e8c\u6b65\u5219\u662f\u9ec4\u8272<\/mark>\u9ad8\u4eae\u90e8\u5206\uff0c\u5728gpu\u4e0a\u8fdb\u884c\u6570\u636e\u5904\u7406\u3002\u6700\u540e\u84dd\u8272<\/mark>\u9ad8\u4eae\u4ee3\u7801\u4e3a\u7b2c\u4e09\u6b65\uff0c\u5c06\u6570\u636e\u4ecegpu\u632a\u56decpu\u3002<\/p>\n\n\n\n
    #include <stdio.h>\n#include \"hip\/hip_runtime.h\"\n\n\/\/ error checking macro\n#define hipCheckErrors(msg) \\\n    do { \\\n       hipError_t __err =hipGetLastError(); \\\n        if (__err !=hipSuccess) { \\\n            fprintf(stderr, \"Fatal error: %s (%s at %s:%d)\\n\", \\\n                msg,hipGetErrorString(__err), \\\n                __FILE__, __LINE__); \\\n            fprintf(stderr, \"*** FAILED - ABORTING\\n\"); \\\n            exit(1); \\\n        } \\\n    } while (0)\n\n\nconst int DSIZE = 4096;\nconst int block_size = 256;  \/\/hip maximum is 1024\n\/\/ vector add kernel: C = A + B\n__global__ void vadd(const float *A, const float *B, float *C, int ds){\n\n  int idx = threadIdx.x+blockIdx.x*blockDim.x; \/\/ create typical 1D thread index from built-in variables\n  if (idx < ds)\n    C[idx] = A[idx] + B[idx];        \/\/ do the vector (element) add here\n}\n\nint main(){\n\n  float *h_A, *h_B, *h_C, *d_A, *d_B, *d_C;\n  h_A = new float[DSIZE];  \/\/ allocate space for vectors in host memory\n  h_B = new float[DSIZE];\n  h_C = new float[DSIZE];\n  for (int i = 0; i < DSIZE; i++){  \/\/ initialize vectors in host memory\n    h_A[i] = rand()\/(float)RAND_MAX;\n    h_B[i] = rand()\/(float)RAND_MAX;\n    h_C[i] = 0;}\n hipMalloc(&d_A, DSIZE*sizeof(float));  \/\/ allocate device space for vector A\n hipMalloc(&d_B, DSIZE*sizeof(float)); \/\/ allocate device space for vector B\n hipMalloc(&d_C, DSIZE*sizeof(float)); \/\/ allocate device space for vector C<\/mark>\n hipCheckErrors(\"cudaMalloc failure\"); \/\/ error checking\n  \/\/ copy vector A to device:\n hipMemcpy(d_A, h_A, DSIZE*sizeof(float), hipMemcpyHostToDevice);<\/mark>\n  \/\/ copy vector B to device:\n hipMemcpy(d_B, h_B, DSIZE*sizeof(float), hipMemcpyHostToDevice);<\/mark>\n hipCheckErrors(\"cudaMemcpy H2D failure\");\n  \/\/cuda processing sequence step 1 is complete\n  vadd<<<(DSIZE+block_size-1)\/block_size, block_size>>>(d_A, d_B, d_C, DSIZE);<\/mark>\n hipCheckErrors(\"kernel launch failure\");\n  \/\/cuda processing sequence step 2 is complete\n  \/\/ copy vector C from device to host:\n hipMemcpy(h_C, d_C, DSIZE*sizeof(float), hipMemcpyDeviceToHost);<\/mark>\n  \/\/cuda processing sequence step 3 is complete\n hipCheckErrors(\"kernel execution failure orhipMemcpy H2D failure\");\n  printf(\"A[0] = %f\\n\", h_A[0]);\n  printf(\"B[0] = %f\\n\", h_B[0]);\n  printf(\"C[0] = %f\\n\", h_C[0]);\n  return 0;\n}<\/code><\/pre>\n\n\n\n
    \u6700\u540e\u7ed3\u679c\u5982\u4e0b\uff1a<\/p>\n\n\n\n
    (base) qibin:~\/projects\/cuda-training-series\/exercises\/hw1$ hipcc -o vector_add vector_add.cu<\/mark> \nA[0] = 0.840188\nB[0] = 0.394383\nC[0] = 1.234571<\/mark><\/code><\/pre>\n<\/details>\n\n\n\n
    hw1 – matrix_mul<\/summary>\n
    \u6700\u540e\u8fd9\u4e2a\u7ec3\u4e60\u9898\u7528\u4e00\u7ef4\u6570\u7ec4\u6a21\u62df\u4e86\u77e9\u9635\u4e58\u6cd5\uff0c\u5176\u5b9e\u6838\u5fc3\u5185\u5bb9\u4ee5\u524d\u9762\u5927\u5dee\u4e0d\u5dee\uff0c\u4f46\u662f\u6211\u8ba4\u4e3a\u8fd9\u4e2a\u7ec3\u4e60\u7ed9\u6211\u7684\u6536\u83b7\u53cd\u800c\u662f\u6700\u6df1\u523b\u7684\u3002\u9996\u5148\u5b83\u8ba9\u6211\u4eec\u66f4\u52a0\u4e60\u60ef\u4e86\u5e76\u884c\u8ba1\u7b97\u7684\u5199\u6cd5\uff0c\u5176\u6b21\u5219\u662f\u6211\u4e00\u4e2a\u65e0\u610f\u95f4\u7684\u6539\u52a8\uff0c\u8ba9\u6211\u53c8\u6536\u83b7\u5230\u4e86\u989d\u5916\u7684\u77e5\u8bc6\u3002<\/p>\n\n\n\n
    \u4ee3\u7801\u4e2d\u7684\u6838\u5fc3\u90e8\u5206mmul\u5982\u4e0b\uff1a<\/p>\n\n\n\n
    __global__ void mmul(const float *A, const float *B, float *C, int ds) {\n\n  int idx = threadIdx.x+blockDim.x*blockIdx.x; \/\/ create thread x index\n  int idy = threadIdx.y+blockDim.y*blockIdx.y; \/\/ create thread y index\n\n  if ((idx < ds) && (idy < ds)){\n    float temp = 0;\n    for (int i = 0; i < ds; i++)\n      temp += A[idy*ds+i] * B[i*ds+idx];   \/\/ dot product of row and column\n    C[idy*ds+idx] = temp;<\/mark>\n  }\n}<\/code><\/pre>\n\n\n\n
    \u76f8\u53cd\u7684\u6211\u4eec\u5982\u679c\u5c06\u9ec4\u8272\u9ad8\u4eae\u90e8\u5206\u4fee\u6539\u4e3a\uff1a<\/p>\n\n\n\n
        for (int i = 0; i < ds; i++)\n      temp += A[<\/mark>idx<\/mark>*ds+i] * B[i*ds+<\/mark>idy<\/mark>];   \/\/ dot product of row and column\n    C[<\/mark>idx<\/mark>*ds+<\/mark>idy<\/mark>] = temp;<\/mark><\/code><\/pre>\n\n\n\n
    \u8fd9\u6837\u7c7b\u4f3c\u4e8e\u5c06\u65b9\u9635\u77e9\u9635\u7684\u884c\u548c\u5217\u8c03\u6362\uff0c\u4f46\u4e0d\u5f71\u54cd\u8ba1\u7b97\u7ed3\u679c\uff0c\u4f46\u662f\u8fd9\u4e24\u79cd\u5199\u6cd5\u7f16\u8bd1\u540e\u7684\u8fd0\u884c\u901f\u5ea6\u5374\u5927\u4e0d\u76f8\u540c\u3002\u4ece\u4ee5\u4e0b\u7ed3\u679c\u770b\uff0c\u5bf9\u6bd4\u4e8e\u4fee\u6539\u524d\u7684\u901f\u5ea6\uff0c\u4fee\u6539\u540e\u8fd0\u884c\u901f\u5ea6\u8db3\u8db3\u591a\u4e86\u4e00\u500d\u3002\u8fd9\u4e2a\u7ed3\u679c\u8ba9\u6211\u5341\u5206\u7684\u56f0\u60d1\uff0c\u56e0\u4e3a\u5728\u6211\u7684\u7406\u89e3\u91cc\u9020\u6210\u8fd9\u79cd\u73b0\u8c61\u5f88\u6709\u53ef\u80fd\u662f\u8bbf\u95ee\u975e\u8fde\u7eed\u5185\u5b58\u9020\u6210\u7684\u3002\u4f46\u662f\uff0c\u4ece\u4e0b\u6807\u8ba1\u7b97\u6765\u770bidy*ds+i<\/strong>\u4e0eidx*ds+i<\/strong>\u5e94\u8be5\u90fd\u662f\u8bbf\u95ee\u8fde\u7eed\u7684\u5185\u5b58\u7a7a\u95f4\uff0c\u800ci*ds+idx<\/strong>\u4e0ei*ds+idy<\/strong><\/mark>\u770b\u8d77\u6765\u867d\u7136\u4e0d\u662f\u8bbf\u95ee\u8fde\u7eed\u7684\u5185\u5b58\uff0c\u4f46\u662f\u4e8c\u8005\u7684\u901f\u5ea6\u5e94\u8be5\u4e5f\u662f\u76f8\u540c\u7684\uff0c\u90a3\u7a76\u7adf\u662f\u4ec0\u4e48\u56e0\u7d20\u9020\u6210\u4e86\u8fd9\u4e00\u73b0\u8c61\uff1f<\/p>\n\n\n\n
    \u4fee\u6539\u524d\u7ed3\u679c<\/strong>\uff1a\n(base) qibin:~\/projects\/cuda-training-series\/exercises\/hw1$ .\/matrix_mul \nInit took 0.037945 seconds.  Begin compute\nDone. Compute took 0.653286<\/mark> seconds\nSuccess!\n\n\u4fee\u6539\u540e\u7ed3\u679c<\/strong>\uff1a\n(base) qibin:~\/projects\/cuda-training-series\/exercises\/hw1$ .\/matrix_mul \nInit took 0.040502 seconds.  Begin compute\nDone. Compute took 1.278184 <\/mark>seconds\nSuccess!<\/code><\/pre>\n\n\n\n
    \u901a\u8fc7\u5728Stackoverflow\u4e0a\u67e5\u627e\uff0c\u6211\u627e\u5230\u4e86\u4e00\u4e2a\u6211\u8ba4\u4e3a\u5408\u7406\u7684\u89e3\u91ca\u90a3\u5c31\u662f\u5408\u5171\u8bbf\u95ee\uff08coalesced access\uff09\uff0c\u4ee5\u4e0b\u662f\u539f\u56de\u7b54\u7684\u8fde\u63a5\uff1a<\/p>\n\n\n\n
    https:\/\/stackoverflow.com\/a\/5044424<\/a><\/p>\n\n\n\n
    \u5176\u4e2d\u4e00\u53e5\u8bdd\uff1a<\/p>\n\n\n\n
    \n
    the correct way to do it is just have consecutive threads access consecutive memory addresses
    \u8ba9\u8fde\u7eed\u7684\u7ebf\u7a0b\u8bbf\u95ee\u8fde\u7eed\u7684\u5185\u5b58\u5730\u5740<\/p>\n<\/blockquote>\n\n\n\n
    \u4ee5\u53ca\u7ed3\u5408\u67e5\u9605\u5230\u7684\u5173\u4e8e\u884c\u4f18\u5148\uff08row-major order\uff09\u8d44\u6599\uff1a<\/p>\n\n\n\n
    https:\/\/engineering.purdue.edu\/~smidkiff\/ece563\/NVidiaGPUTeachingToolkit\/Mod5\/Lecture-5-1-warps-simd.pdf<\/a><\/p>\n\n\n\n
    https:\/\/lumetta.web.engr.illinois.edu\/408-S19\/slide-copies\/ece408-lecture22-S19-ZJUI.pdf<\/a><\/p>\n\n\n\n
    \n
    \n
    \n
    <\/p>\n<\/blockquote>\n<\/blockquote>\n<\/blockquote>\n\n\n\n
    \u6700\u7ec8\u6211\u8ba4\u4e3a\u5408\u7406\u7684\u89e3\u91ca\u5e94\u8be5\u662f\u8fd9\u6837\u7684\uff1a<\/p>\n\n\n\n
    \u5047\u8bbeblockIdx.x=0, blockIdx.y=0, blockDim=3, gridDim=3<\/code>\uff0c\u5373\u5f53\u524d\u4f4d\u4e8eblock(0, 0)<\/code>\uff0cblock\u5185\u67099\u6761\u7ebf\u7a0b\uff083*3\uff09\uff0c\u6839\u636erow-major order<\/code>\u6211\u4eec\u53ef\u4ee5\u5f97\u5230\u88681\u3002\u5f53idx=0\uff0cidy=0,<\/code>\u90a3\u4e48\u4e0b\u6807\u8ba1\u7b97\u7b49\u5f0f\u4e3aidx*ds+i,i*ds+idy<\/code>\uff0c<\/mark>\u5219(0,0), (1,0), (2,0)<\/code>\u4e3a\u76f8\u90bb\u7ebf\u7a0b\u4f46\u5b58\u50a8\u6570\u636e0\uff0c3\uff0c6<\/code>\u5e76\u4e0d\u662f\u8fde\u7eed\u7684\uff0c\u4e0d\u6ee1\u8db3coalesced access\uff1b\u53cd\u4e4b\u5982\u679c\u662fidy*ds+i,i*ds+idx<\/code>\uff0c\u5219(0,0), (0,1), (0,2)<\/code>\u4e3a\u76f8\u90bb\u7ebf\u7a0b\u4e14\u5b58\u50a8\u6570\u636e0\uff0c1\uff0c2<\/code>\u662f\u8fde\u7eed\u7684\uff0c\u6ee1\u8db3coalesced access\u3002\u6700\u7ec8\u4f1a\u56e0\u4e3acoalesced access\u9020\u6210\u8fd0\u884c\u901f\u5ea6\u4e0a\u7684\u5dee\u5f02\u3002<\/p>\n\n\n\n
    \u5143\u7d20idx<\/td>0<\/td>1<\/td>2<\/td>3<\/td>4<\/td>5<\/td>6<\/td>7<\/td>8<\/td><\/tr>
    threadidx(y,x)<\/td>(0,0)<\/td>(0,1)<\/td>(0,2)<\/td>(1,0)<\/td>(1,1)<\/td>(1,2)<\/td>(2,0)<\/td>(2,1)<\/td>(2,2)<\/td><\/tr><\/tbody><\/table>
    \u88681.\u6570\u7ec4\u5143\u7d20\u5bf9\u5e94cuda block<\/figcaption><\/figure>\n<\/details>\n\n\n\n
    \u603b\u7ed3<\/h2>\n\n\n\n
    \u672c\u7ae0\u8282\u7b80\u5355\u5730\u4ecb\u7ecd\u4e86CUDA\u4e2d\u4e00\u4e9b\u7b80\u5355\u7684\u6982\u5ff5\uff0c\u5982\u547d\u540d\uff0c\u5185\u5efa\u53d8\u91cf\uff0c\u540c\u65f6\u4e5f\u7b80\u8981\u4ecb\u7ecd\u8c03\u7528CUDA\u7684\u6d41\u7a0b\u4ee5\u53ca\u6837\u4f8b\u3002\u6700\u540e\u8bfe\u540e\u7ec3\u4e60\u4e5f\u6a21\u62df\u4e86\u77e9\u9635\u52a0\u6cd5\u3001\u4e58\u6cd5\uff0c\u4ee5\u53ca\u6211\u81ea\u5df1\u7684\u4e00\u70b9\u5c1d\u8bd5\u3002<\/p>\n","protected":false},"excerpt":{"rendered":"
    \u6700\u8fd1\u95f2\u6687\u65f6\u95f4\u60f3\u627e\u70b9\u4e8b\u60c5\u505a\uff0c\u5c31\u60f3\u8981\u4e0d\u5b66\u4e60Cuda\u7f16\u7a0b\u5427\uff0c\u5305\u597d\u73a9\u7684\u3002\u7136\u540e\uff0c\u5728Google\u4e0a\u627eCuda\u5b66\u4e60\u8d44\u6599\u65f6\u53d1\u73b0 […]<\/p>\n","protected":false},"author":1,"featured_media":0,"comment_status":"open","ping_status":"open","sticky":false,"template":"","format":"standard","meta":{"footnotes":""},"categories":[35],"tags":[33,27,31,29],"class_list":["post-218","post","type-post","status-publish","format-standard","hentry","category-35","tag-c","tag-cuda","tag-hip","tag-29"],"_links":{"self":[{"href":"https:\/\/www.qibinliang.com\/index.php\/wp-json\/wp\/v2\/posts\/218","targetHints":{"allow":["GET"]}}],"collection":[{"href":"https:\/\/www.qibinliang.com\/index.php\/wp-json\/wp\/v2\/posts"}],"about":[{"href":"https:\/\/www.qibinliang.com\/index.php\/wp-json\/wp\/v2\/types\/post"}],"author":[{"embeddable":true,"href":"https:\/\/www.qibinliang.com\/index.php\/wp-json\/wp\/v2\/users\/1"}],"replies":[{"embeddable":true,"href":"https:\/\/www.qibinliang.com\/index.php\/wp-json\/wp\/v2\/comments?post=218"}],"version-history":[{"count":61,"href":"https:\/\/www.qibinliang.com\/index.php\/wp-json\/wp\/v2\/posts\/218\/revisions"}],"predecessor-version":[{"id":317,"href":"https:\/\/www.qibinliang.com\/index.php\/wp-json\/wp\/v2\/posts\/218\/revisions\/317"}],"wp:attachment":[{"href":"https:\/\/www.qibinliang.com\/index.php\/wp-json\/wp\/v2\/media?parent=218"}],"wp:term":[{"taxonomy":"category","embeddable":true,"href":"https:\/\/www.qibinliang.com\/index.php\/wp-json\/wp\/v2\/categories?post=218"},{"taxonomy":"post_tag","embeddable":true,"href":"https:\/\/www.qibinliang.com\/index.php\/wp-json\/wp\/v2\/tags?post=218"}],"curies":[{"name":"wp","href":"https:\/\/api.w.org\/{rel}","templated":true}]}}