strings_To_Colaboratory.ipynb

{
  "cells": [
    {
      "cell_type": "markdown",
      "metadata": {
        "id": "view-in-github",
        "colab_type": "text"
      },
      "source": [
        "<a href=\"https://colab.research.google.com/github/tonystz/gitpod/blob/main/strings_To_Colaboratory.ipynb\" target=\"_parent\"><img src=\"https://colab.research.google.com/assets/colab-badge.svg\" alt=\"Open In Colab\"/></a>"
      ]
    },
    {
      "cell_type": "code",
      "source": [
        "!pip install pycuda # install cuda\n",
        "import pycuda.driver as cuda\n",
        "import pycuda.autoinit\n",
        "from pycuda.compiler import SourceModule"
      ],
      "metadata": {
        "id": "6QAStcPlRxWx"
      },
      "execution_count": null,
      "outputs": []
    },
    {
      "cell_type": "code",
      "source": [
        "import numpy as np\n",
        "ttt = np.asarray([ \"stuff\" + str(i)  for i in range(0,2) ])\n",
        "\n",
        "print( ttt.dtype, type(ttt[0]) ) \n",
        "print(ttt)\n",
        "\n",
        "s= np.array(['abcdf'], dtype=object)\n",
        "print( s.dtype, type(s[0]) )\n",
        "\n",
        "#print(\"ctype:\",np.ctypeslib.as_ctypes_type(s.dtype))"
      ],
      "metadata": {
        "colab": {
          "base_uri": "https://localhost:8080/"
        },
        "id": "tIfpULJwlMWY",
        "outputId": "9aa8bf07-9695-43a1-bc6e-2faabecbcb1d"
      },
      "execution_count": 40,
      "outputs": [
        {
          "output_type": "stream",
          "name": "stdout",
          "text": [
            "<U6 <class 'numpy.str_'>\n",
            "['stuff0' 'stuff1']\n",
            "object <class 'str'>\n"
          ]
        }
      ]
    },
    {
      "cell_type": "markdown",
      "source": [
        "## string test"
      ],
      "metadata": {
        "id": "MkqQD552fAXv"
      }
    },
    {
      "cell_type": "code",
      "source": [
        "%%writefile s2.py\n",
        "import time\n",
        "import numpy as np\n",
        "from pycuda import driver, compiler, gpuarray, tools\n",
        "import math\n",
        "from sys import getsizeof\n",
        "\n",
        "import pycuda.autoinit\n",
        "\n",
        "kernel_code1 = \"\"\"\n",
        "__global__ void test1(char** d_wordList) {\n",
        "    (d_wordList[blockIdx.x][threadIdx.x])++;      \n",
        "}\n",
        "    \"\"\"\n",
        "\n",
        "kernel_code2 = \"\"\"\n",
        "__global__ void test2(char* d_wordList, size_t *offsets) {\n",
        "\n",
        "    int idx = threadIdx.x+blockDim.x*blockIdx.x;\n",
        "    printf(\"index: threadIdx.x=%d blockIdx.x=%d  blockDim.x=%d \\\\n\",threadIdx.x,blockIdx.x,blockDim.x);\n",
        "    printf(\"Hello from thread %d, my string is %s ,offsets:%u \\\\n\", idx, d_wordList+offsets[idx],offsets[blockIdx.x]);\n",
        "    (d_wordList[offsets[blockIdx.x] + threadIdx.x])++;\n",
        "}\n",
        "    \"\"\"\n",
        "\n",
        "\n",
        "\n",
        "\n",
        "mod = compiler.SourceModule(kernel_code1)\n",
        "ker_test1 = mod.get_function(\"test1\")\n",
        "\n",
        "\n",
        "\n",
        "wordList = ['asd','bsd','csd']\n",
        "\n",
        "d_words = []\n",
        "\n",
        "for word in wordList:\n",
        "    d_words.append(gpuarray.to_gpu(np.array(word, dtype=str)))\n",
        "\n",
        "d_wordList = gpuarray.to_gpu(np.array([word.ptr for word in d_words], dtype=np.uintp))\n",
        "\n",
        "ker_test1(d_wordList, block=(3,1,1), grid=(3,1,1))\n",
        "\n",
        "for word in d_words:\n",
        "  result = word.get()\n",
        "  print(result)\n",
        "\n",
        "mod2 = compiler.SourceModule(kernel_code2)\n",
        "ker_test2 = mod2.get_function(\"test2\")\n",
        "\n",
        "d_words2 = gpuarray.to_gpu(np.array(['asd','bsd','csd'], dtype=np.string_))\n",
        "offsets = gpuarray.to_gpu(np.array([0,3,6,9], dtype=np.uint64))\n",
        "ker_test2(d_words2, offsets, block=(3,1,1), grid=(1,1,1))\n",
        "h_words2 = d_words2.get()\n",
        "print(h_words2)"
      ],
      "metadata": {
        "colab": {
          "base_uri": "https://localhost:8080/"
        },
        "id": "jMJcWbEffD9I",
        "outputId": "c7eb6093-2611-4bd3-d606-70500ca792b3"
      },
      "execution_count": 189,
      "outputs": [
        {
          "output_type": "stream",
          "name": "stdout",
          "text": [
            "Overwriting s2.py\n"
          ]
        }
      ]
    },
    {
      "cell_type": "code",
      "source": [
        "!python s2.py"
      ],
      "metadata": {
        "colab": {
          "base_uri": "https://localhost:8080/"
        },
        "id": "tZPgiuMAfGdZ",
        "outputId": "6b01ccee-ec35-4b88-baa8-2b37ce6f2476"
      },
      "execution_count": 190,
      "outputs": [
        {
          "output_type": "stream",
          "name": "stdout",
          "text": [
            "𐅢sd\n",
            "𐅣sd\n",
            "𐅤sd\n",
            "index: threadIdx.x=0 blockIdx.x=0  blockDim.x=3 \n",
            "index: threadIdx.x=1 blockIdx.x=0  blockDim.x=3 \n",
            "index: threadIdx.x=2 blockIdx.x=0  blockDim.x=3 \n",
            "Hello from thread 0, my string is asdbsdcsd ,offsets:0 \n",
            "Hello from thread 1, my string is bsdcsd ,offsets:0 \n",
            "Hello from thread 2, my string is csd ,offsets:0 \n",
            "[b'bte' b'bsd' b'csd']\n"
          ]
        }
      ]
    },
    {
      "cell_type": "markdown",
      "source": [
        "## all string size must equal?"
      ],
      "metadata": {
        "id": "cWZj7ByhRDsb"
      }
    },
    {
      "cell_type": "code",
      "source": [
        "%%writefile s.py\n",
        "#!python \n",
        "import pycuda.driver as cuda\n",
        "import pycuda.autoinit\n",
        "from pycuda.compiler import SourceModule\n",
        "from pycuda import gpuarray\n",
        "import numpy as np\n",
        "\n",
        "mod = SourceModule(\"\"\"\n",
        "    #include <stdio.h>\n",
        "    const int strLen=6;\n",
        "    __global__ void say_hi(char*out_gpu, unsigned int *s_offset_gpu)\n",
        "    { \n",
        "      int idx = threadIdx.x + blockIdx.x * blockDim.x;\n",
        "\n",
        "      //printf(\"index: threadIdx.x=%d blockIdx.x=%d  blockDim.x=%d \\\\n\",threadIdx.x,blockIdx.x,blockDim.x);\n",
        "      //printf(\"thread id: [%d]  offset:%d \\\\n\",idx,s_offset_gpu[idx]);\n",
        "      printf(\"thread id: [%d]  offset:%d  get str:%s  strlen=%d \\\\n\",idx,s_offset_gpu[idx],out_gpu+s_offset_gpu[idx]+1, strLen);\n",
        "\n",
        "      // string len is 6\n",
        "      char sub[strLen+1];\n",
        "      for(int i = 0; i<strLen; i++)\n",
        "      {\n",
        "        sub[i]=out_gpu[idx*strLen+i];\n",
        "      }\n",
        "      printf(\"substring:%s\\\\n\",sub);\n",
        "    }\n",
        "    \"\"\")\n",
        "\n",
        "say_hi = mod.get_function(\"say_hi\")\n",
        "\n",
        "#padding string with \\0 char\n",
        "s= np.array(['abcmh\\0','edfg\\0\\0','s12wc\\0','h45q\\0\\0'], dtype=np.string_)\n",
        "\n",
        "offset=[0]\n",
        "for i in s:\n",
        "  offset.append(offset[-1]+len(i))\n",
        "s_offset=np.array(offset,dtype=np.uint32)\n",
        "print('shape:',s.shape)\n",
        "print(s,s.data)\n",
        "print('offset:',s_offset,s_offset.shape)\n",
        "\n",
        "s_gpu = gpuarray.to_gpu(s)\n",
        "s_offset_gpu = gpuarray.to_gpu(s_offset)\n",
        "\n",
        "say_hi(s_gpu,s_offset_gpu,block=(s.size,1,1),grid=(1,1,1))\n",
        "print('modify:',s_gpu.get(),s_offset_gpu.get())\n",
        "\n",
        "print(\"s_offset ctype:\",np.ctypeslib.as_ctypes_type(s_offset.dtype))"
      ],
      "metadata": {
        "colab": {
          "base_uri": "https://localhost:8080/"
        },
        "id": "C3vCpSIuULQt",
        "outputId": "257489db-fc01-4b31-c6cf-2b7c3c917f36"
      },
      "execution_count": 191,
      "outputs": [
        {
          "output_type": "stream",
          "name": "stdout",
          "text": [
            "Overwriting s.py\n"
          ]
        }
      ]
    },
    {
      "cell_type": "code",
      "source": [
        "!python s.py"
      ],
      "metadata": {
        "colab": {
          "base_uri": "https://localhost:8080/"
        },
        "id": "UNIVqz_mUS2u",
        "outputId": "45fe26b2-261f-4028-e834-29d372ff3129"
      },
      "execution_count": 192,
      "outputs": [
        {
          "output_type": "stream",
          "name": "stdout",
          "text": [
            "shape: (4,)\n",
            "[b'abcmh' b'edfg' b's12wc' b'h45q'] <memory at 0x7f56d7025580>\n",
            "offset: [ 0  5  9 14 18] (5,)\n",
            "thread id: [0]  offset:0  get str:bcmh  strlen=6 \n",
            "thread id: [1]  offset:5  get str:edfg  strlen=6 \n",
            "thread id: [2]  offset:9  get str:  strlen=6 \n",
            "thread id: [3]  offset:14  get str:wc  strlen=6 \n",
            "substring:abcmh\n",
            "substring:edfg\n",
            "substring:s12wc\n",
            "substring:h45q\n",
            "modify: [b'abcmh' b'edfg' b's12wc' b'h45q'] [ 0  5  9 14 18]\n",
            "s_offset ctype: <class 'ctypes.c_uint'>\n"
          ]
        }
      ]
    },
    {
      "cell_type": "markdown",
      "source": [
        "## test un-equal string process?"
      ],
      "metadata": {
        "id": "q0aII8SvieuC"
      }
    },
    {
      "cell_type": "code",
      "source": [
        "%%writefile su.py\n",
        "#!python \n",
        "import pycuda.driver as cuda\n",
        "import pycuda.autoinit\n",
        "from pycuda.compiler import SourceModule\n",
        "from pycuda import gpuarray\n",
        "import numpy as np\n",
        "\n",
        "mod = SourceModule(\"\"\"\n",
        "    #include <stdio.h>\n",
        "\n",
        "    __global__ void say_hi(char*out_gpu, unsigned int *s_offset_gpu)\n",
        "    {  //max string len is 13\n",
        "      const int strLen=13;\n",
        "      int idx = threadIdx.x + blockIdx.x * blockDim.x;\n",
        "      //printf(\"index: threadIdx.x=%d blockIdx.x=%d  blockDim.x=%d \\\\n\",threadIdx.x,blockIdx.x,blockDim.x);\n",
        "      //printf(\"thread id: [%d]  offset:%d \\\\n\",idx,s_offset_gpu[idx]);\n",
        "      printf(\"thread id: [%d]  offset:%d  get str:%s\\\\n\",idx,s_offset_gpu[idx],out_gpu+s_offset_gpu[idx]);\n",
        "  \n",
        "      char sub[strLen+1];\n",
        "      for(int i = 0; i<strLen; i++)\n",
        "      {\n",
        "        sub[i]=out_gpu[idx*strLen+i];\n",
        "      }\n",
        "      printf(\"thread[%d]:substring:%s\\\\n\",idx,sub);\n",
        "    }\n",
        "    \"\"\")\n",
        "\n",
        "say_hi = mod.get_function(\"say_hi\")\n",
        "\n",
        "\n",
        "s= np.array(['abc\\0','edfg\\0','stz12345678b\\0','short\\0'], dtype=np.string_)\n",
        "\n",
        "offset=[0]\n",
        "for i in s:\n",
        "  offset.append(offset[-1]+len(i))\n",
        "s_offset=np.array(offset,dtype=np.uint32)\n",
        "print('shape:',s.shape)\n",
        "print(s,s.data)\n",
        "print('offset:',s_offset,s_offset.shape)\n",
        "\n",
        "s_gpu = gpuarray.to_gpu(s)\n",
        "s_offset_gpu = gpuarray.to_gpu(s_offset)\n",
        "\n",
        "say_hi(s_gpu,s_offset_gpu,block=(s.size,1,1),grid=(1,1,1))\n",
        "print('modify:',s_gpu.get(),s_offset_gpu.get())\n",
        "\n",
        "print(\"s_offset ctype:\",np.ctypeslib.as_ctypes_type(s_offset.dtype))"
      ],
      "metadata": {
        "colab": {
          "base_uri": "https://localhost:8080/"
        },
        "id": "TaQ8-EtoieRb",
        "outputId": "521be165-b7a1-4a69-b8b2-49fbb9b913a0"
      },
      "execution_count": 193,
      "outputs": [
        {
          "output_type": "stream",
          "name": "stdout",
          "text": [
            "Overwriting su.py\n"
          ]
        }
      ]
    },
    {
      "cell_type": "code",
      "source": [
        "!python su.py"
      ],
      "metadata": {
        "colab": {
          "base_uri": "https://localhost:8080/"
        },
        "id": "gJFwn3gkiy7D",
        "outputId": "65cf029b-6e70-4075-e67f-4b6c59e0fb72"
      },
      "execution_count": 194,
      "outputs": [
        {
          "output_type": "stream",
          "name": "stdout",
          "text": [
            "shape: (4,)\n",
            "[b'abc' b'edfg' b'stz12345678b' b'short'] <memory at 0x7f4b43ab7580>\n",
            "offset: [ 0  3  7 19 24] (5,)\n",
            "thread id: [0]  offset:0  get str:abc\n",
            "thread id: [1]  offset:3  get str:\n",
            "thread id: [2]  offset:7  get str:\n",
            "thread id: [3]  offset:19  get str:\n",
            "thread[0]:substring:abc\n",
            "thread[1]:substring:edfg\n",
            "thread[2]:substring:stz12345678b\n",
            "thread[3]:substring:short\n",
            "modify: [b'abc' b'edfg' b'stz12345678b' b'short'] [ 0  3  7 19 24]\n",
            "s_offset ctype: <class 'ctypes.c_uint'>\n"
          ]
        }
      ]
    },
    {
      "cell_type": "code",
      "source": [
        "%%writefile t389.cu\n",
        "#include <stdio.h>\n",
        "#include <string.h>\n",
        "\n",
        "#define nTPB 256\n",
        "\n",
        "__global__ void kern_1D(char *data, unsigned *indices, unsigned num_strings){\n",
        "\n",
        "  int idx = threadIdx.x+blockDim.x*blockIdx.x;\n",
        "  if (idx < num_strings)\n",
        "    printf(\"Hello from thread %d, offset=%d, my string is %s\\n\", idx, indices[idx],data+indices[idx]);\n",
        "}\n",
        "\n",
        "__global__ void kern_2D(char **data, unsigned num_strings){\n",
        "\n",
        "  int idx = threadIdx.x+blockDim.x*blockIdx.x;\n",
        "  if (idx < num_strings)\n",
        "    printf(\"Hello from thread %d, my string is %s\\n\", idx, data[idx]);\n",
        "}\n",
        "\n",
        "int main(){\n",
        "\n",
        "  const int num_strings = 3;\n",
        "  const char s0[] = \"s1\\0\";\n",
        "  const char s1[] = \"s2\\0\";\n",
        "  const char s2[] = \"sstz3\";\n",
        "  int ds[num_strings];\n",
        "  ds[0] = sizeof(s0)/sizeof(char);\n",
        "  ds[1] = sizeof(s1)/sizeof(char);\n",
        "  ds[2] = sizeof(s2)/sizeof(char);\n",
        "  // pretend we have a dynamically allocated char** array\n",
        "  char **data;\n",
        "  data = (char **)malloc(num_strings*sizeof(char *));\n",
        "  data[0] = (char *)malloc(ds[0]*sizeof(char));\n",
        "  data[1] = (char *)malloc(ds[1]*sizeof(char));\n",
        "  data[2] = (char *)malloc(ds[2]*sizeof(char));\n",
        "  // initialize said array\n",
        "  strcpy(data[0], s0);\n",
        "  strcpy(data[1], s1);\n",
        "  strcpy(data[2], s2);\n",
        "  // method 1: \"flattening\"\n",
        "  char *fdata = (char *)malloc((ds[0]+ds[1]+ds[2])*sizeof(char));\n",
        "  unsigned *ind   = (unsigned *)malloc(num_strings*sizeof(unsigned));\n",
        "  unsigned next = 0;\n",
        "  for (int i = 0; i < num_strings; i++){\n",
        "    strcpy(fdata+next, data[i]);\n",
        "    ind[i] = next;\n",
        "    next += ds[i];}\n",
        "  //copy to device\n",
        "  char *d_fdata;\n",
        "  unsigned *d_ind;\n",
        "  cudaMalloc(&d_fdata, next*sizeof(char));\n",
        "  cudaMalloc(&d_ind, num_strings*sizeof(unsigned));\n",
        "  cudaMemcpy(d_fdata, fdata, next*sizeof(char), cudaMemcpyHostToDevice);\n",
        "  cudaMemcpy(d_ind, ind, num_strings*sizeof(unsigned), cudaMemcpyHostToDevice);\n",
        "  printf(\"method 1: %d  --> %d\\n\",(num_strings+nTPB-1)/nTPB,nTPB);\n",
        "  kern_1D<<<(num_strings+nTPB-1)/nTPB, nTPB>>>(d_fdata, d_ind, num_strings);\n",
        "  cudaDeviceSynchronize();\n",
        "  //method 2: \"2D\" (pointer-to-pointer) array\n",
        "  char **d_data;\n",
        "  cudaMalloc(&d_data, num_strings*sizeof(char *));\n",
        "  char **d_temp_data;\n",
        "  d_temp_data = (char **)malloc(num_strings*sizeof(char *));\n",
        "  for (int i = 0; i < num_strings; i++){\n",
        "    cudaMalloc(&(d_temp_data[i]), ds[i]*sizeof(char));\n",
        "    cudaMemcpy(d_temp_data[i], data[i], ds[i]*sizeof(char), cudaMemcpyHostToDevice);\n",
        "    cudaMemcpy(d_data+i, &(d_temp_data[i]), sizeof(char *), cudaMemcpyHostToDevice);}\n",
        "  printf(\"method 2:\\n\");\n",
        "  kern_2D<<<(num_strings+nTPB-1)/nTPB, nTPB>>>(d_data, num_strings);\n",
        "  cudaDeviceSynchronize();\n",
        "  // method 3: managed allocations\n",
        "  // start over with a managed char** array\n",
        "  char **m_data;\n",
        "  cudaMallocManaged(&m_data, num_strings*sizeof(char *));\n",
        "  cudaMallocManaged(&(m_data[0]), ds[0]*sizeof(char));\n",
        "  cudaMallocManaged(&(m_data[1]), ds[1]*sizeof(char));\n",
        "  cudaMallocManaged(&(m_data[2]), ds[2]*sizeof(char));\n",
        "  // initialize said array\n",
        "  strcpy(m_data[0], s0);\n",
        "  strcpy(m_data[1], s1);\n",
        "  strcpy(m_data[2], s2);\n",
        "  // call kernel directly on managed data\n",
        "  printf(\"method 3:\\n\");\n",
        "  kern_2D<<<(num_strings+nTPB-1)/nTPB, nTPB>>>(m_data, num_strings);\n",
        "  cudaDeviceSynchronize();\n",
        "\n",
        "  return 0;\n",
        "}"
      ],
      "metadata": {
        "colab": {
          "base_uri": "https://localhost:8080/"
        },
        "id": "-LixIpFBrd7x",
        "outputId": "9be27f49-55dc-4a7b-bdbd-41f82cbdc514"
      },
      "execution_count": 195,
      "outputs": [
        {
          "output_type": "stream",
          "name": "stdout",
          "text": [
            "Overwriting t389.cu\n"
          ]
        }
      ]
    },
    {
      "cell_type": "code",
      "source": [
        "!nvcc t389.cu -o t389\n",
        "!nvprof ./t389"
      ],
      "metadata": {
        "colab": {
          "base_uri": "https://localhost:8080/"
        },
        "id": "ffznc4yfbUXI",
        "outputId": "022f5c1f-05d5-4982-c9de-3cd858e730b3"
      },
      "execution_count": 196,
      "outputs": [
        {
          "output_type": "stream",
          "name": "stdout",
          "text": [
            "==46409== NVPROF is profiling process 46409, command: ./t389\n",
            "method 1: 1  --> 256\n",
            "Hello from thread 0, offset=0, my string is s1\n",
            "Hello from thread 1, offset=4, my string is s2\n",
            "Hello from thread 2, offset=8, my string is sstz3\n",
            "method 2:\n",
            "Hello from thread 0, my string is s1\n",
            "Hello from thread 1, my string is s2\n",
            "Hello from thread 2, my string is sstz3\n",
            "method 3:\n",
            "Hello from thread 0, my string is s1\n",
            "Hello from thread 1, my string is s2\n",
            "Hello from thread 2, my string is sstz3\n",
            "==46409== Profiling application: ./t389\n",
            "==46409== Profiling result:\n",
            "            Type  Time(%)      Time     Calls       Avg       Min       Max  Name\n",
            " GPU activities:   76.82%  419.57us         2  209.79us  97.565us  322.01us  kern_2D(char**, unsigned int)\n",
            "                   20.95%  114.40us         1  114.40us  114.40us  114.40us  kern_1D(char*, unsigned int*, unsigned int)\n",
            "                    2.23%  12.193us         8  1.5240us  1.4080us  1.9840us  [CUDA memcpy HtoD]\n",
            "      API calls:   90.84%  213.86ms         6  35.644ms  2.7100us  213.84ms  cudaMalloc\n",
            "                    8.75%  20.597ms         4  5.1492ms  6.0420us  20.489ms  cudaMallocManaged\n",
            "                    0.25%  585.89us         3  195.30us  111.83us  345.62us  cudaDeviceSynchronize\n",
            "                    0.06%  140.02us       101  1.3860us     140ns  57.490us  cuDeviceGetAttribute\n",
            "                    0.05%  124.19us         3  41.396us  13.566us  84.474us  cudaLaunchKernel\n",
            "                    0.03%  81.646us         8  10.205us  5.3110us  23.553us  cudaMemcpy\n",
            "                    0.01%  25.862us         1  25.862us  25.862us  25.862us  cuDeviceGetName\n",
            "                    0.00%  7.2580us         1  7.2580us  7.2580us  7.2580us  cuDeviceGetPCIBusId\n",
            "                    0.00%  2.3090us         3     769ns     250ns  1.7420us  cuDeviceGetCount\n",
            "                    0.00%     995ns         2     497ns     232ns     763ns  cuDeviceGet\n",
            "                    0.00%     540ns         1     540ns     540ns     540ns  cuModuleGetLoadingMode\n",
            "                    0.00%     457ns         1     457ns     457ns     457ns  cuDeviceTotalMem\n",
            "                    0.00%     245ns         1     245ns     245ns     245ns  cuDeviceGetUuid\n",
            "\n",
            "==46409== Unified Memory profiling result:\n",
            "Device \"Tesla T4 (0)\"\n",
            "   Count  Avg Size  Min Size  Max Size  Total Size  Total Time  Name\n",
            "       2  32.000KB  4.0000KB  60.000KB  64.00000KB  10.39900us  Host To Device\n",
            "       1         -         -         -           -  222.1390us  Gpu page fault groups\n",
            "Total CPU Page faults: 1\n"
          ]
        }
      ]
    },
    {
      "cell_type": "code",
      "source": [
        "!nvcc -o t389 t389.cu\n",
        "!compute-sanitizer ./t389"
      ],
      "metadata": {
        "colab": {
          "base_uri": "https://localhost:8080/"
        },
        "id": "9sSHVTI7rrBI",
        "outputId": "643e1bba-d176-4804-beaa-0eb8c338f619"
      },
      "execution_count": 197,
      "outputs": [
        {
          "output_type": "stream",
          "name": "stdout",
          "text": [
            "========= COMPUTE-SANITIZER\n",
            "method 1: 1  --> 256\n",
            "Hello from thread 0, offset=0, my string is s1\n",
            "Hello from thread 1, offset=4, my string is s2\n",
            "Hello from thread 2, offset=8, my string is sstz3\n",
            "method 2:\n",
            "Hello from thread 0, my string is s1\n",
            "Hello from thread 1, my string is s2\n",
            "Hello from thread 2, my string is sstz3\n",
            "method 3:\n",
            "Hello from thread 0, my string is s1\n",
            "Hello from thread 1, my string is s2\n",
            "Hello from thread 2, my string is sstz3\n",
            "========= ERROR SUMMARY: 0 errors\n"
          ]
        }
      ]
    },
    {
      "cell_type": "markdown",
      "source": [
        "## test string 1 dim-array fixed stringlen[one thread one string ]"
      ],
      "metadata": {
        "id": "CUZbprPagheb"
      }
    },
    {
      "cell_type": "code",
      "source": [
        "import numpy as np\n",
        "s=['abc\\0','edfg\\0','stz12345678b\\0','b1\\0','k2\\0','hellworld3\\0']\n",
        "print('maxlen:',max([len(i) for i in  s]))\n",
        "print('maxlen:',max([len(i.encode()) for i in  s]))\n",
        "s= np.array(s, dtype=np.string_)\n",
        "print(s)\n",
        "s2=s.reshape(2,3)\n",
        "print(s2)\n",
        "ttest=np.array([1,2],np.int32)\n",
        "print(\"ctype:\",np.ctypeslib.as_ctypes_type(ttest.dtype))"
      ],
      "metadata": {
        "colab": {
          "base_uri": "https://localhost:8080/"
        },
        "id": "6kqd2yEygpHr",
        "outputId": "75188b31-49ea-43ce-ee18-6cac2af7af7b"
      },
      "execution_count": 234,
      "outputs": [
        {
          "output_type": "stream",
          "name": "stdout",
          "text": [
            "maxlen: 13\n",
            "maxlen: 13\n",
            "[b'abc' b'edfg' b'stz12345678b' b'b1' b'k2' b'hellworld3']\n",
            "[[b'abc' b'edfg' b'stz12345678b']\n",
            " [b'b1' b'k2' b'hellworld3']]\n",
            "ctype: <class 'ctypes.c_int'>\n"
          ]
        }
      ]
    },
    {
      "cell_type": "code",
      "source": [
        "%%writefile sfixed.py\n",
        "#!python \n",
        "import pycuda.driver as cuda\n",
        "import pycuda.autoinit\n",
        "from pycuda.compiler import SourceModule\n",
        "from pycuda import gpuarray\n",
        "import numpy as np\n",
        "\n",
        "mod = SourceModule(\"\"\"\n",
        "    #include <stdio.h>\n",
        "    \n",
        "    __global__ void say_hi(char*out_gpu)\n",
        "    { \n",
        "      const int strLen=13;\n",
        "      int idx = threadIdx.x + blockIdx.x * blockDim.x;\n",
        "\n",
        "      //printf(\"index: threadIdx.x=%d blockIdx.x=%d  blockDim.x=%d \\\\n\",threadIdx.x,blockIdx.x,blockDim.x);\n",
        "      printf(\"thread id: [%d]  get str:%s  strlen=%d \\\\n\",idx,out_gpu+strLen+1, strLen);\n",
        "\n",
        "      // string len is 6\n",
        "      char sub[strLen+1];\n",
        "      for(int i = 0; i<strLen; i++)\n",
        "      {\n",
        "        sub[i]=out_gpu[idx*strLen+i];\n",
        "      }\n",
        "      printf(\"thread[%d] substring:%s\\\\n\",idx,sub);\n",
        "    }\n",
        "    \"\"\")\n",
        "\n",
        "say_hi = mod.get_function(\"say_hi\")\n",
        "\n",
        "#padding string with \\0 char\n",
        "s= np.array(['abc\\0','edfg\\0','stz12345678b\\0','b1\\0','k2\\0','hellworld3\\0'], dtype=np.string_)\n",
        "\n",
        "\n",
        "print('shape:',s.shape)\n",
        "print(s,s.data)\n",
        "\n",
        "s_gpu = gpuarray.to_gpu(s)\n",
        "say_hi(s_gpu,block=(s.size,1,1),grid=(1,1,1))\n",
        "print('modify:',s_gpu.get())"
      ],
      "metadata": {
        "colab": {
          "base_uri": "https://localhost:8080/"
        },
        "id": "gVKSc9b5gnUj",
        "outputId": "5934d15e-b655-402e-fe4e-4da1b6c0f9aa"
      },
      "execution_count": 235,
      "outputs": [
        {
          "output_type": "stream",
          "name": "stdout",
          "text": [
            "Overwriting sfixed.py\n"
          ]
        }
      ]
    },
    {
      "cell_type": "code",
      "source": [
        "!python sfixed.py"
      ],
      "metadata": {
        "colab": {
          "base_uri": "https://localhost:8080/"
        },
        "id": "21Cbfqwcho_t",
        "outputId": "481a5e31-80cb-4f58-c3a0-bbedbbb43cc5"
      },
      "execution_count": 236,
      "outputs": [
        {
          "output_type": "stream",
          "name": "stdout",
          "text": [
            "shape: (6,)\n",
            "[b'abc' b'edfg' b'stz12345678b' b'b1' b'k2' b'hellworld3'] <memory at 0x7f4dd2231580>\n",
            "thread id: [0]  get str:dfg  strlen=13 \n",
            "thread id: [1]  get str:dfg  strlen=13 \n",
            "thread id: [2]  get str:dfg  strlen=13 \n",
            "thread id: [3]  get str:dfg  strlen=13 \n",
            "thread id: [4]  get str:dfg  strlen=13 \n",
            "thread id: [5]  get str:dfg  strlen=13 \n",
            "thread[0] substring:abc\n",
            "thread[1] substring:edfg\n",
            "thread[2] substring:stz12345678b\n",
            "thread[3] substring:b1\n",
            "thread[4] substring:k2\n",
            "thread[5] substring:hellworld3\n",
            "modify: [b'abc' b'edfg' b'stz12345678b' b'b1' b'k2' b'hellworld3']\n"
          ]
        }
      ]
    },
    {
      "cell_type": "markdown",
      "source": [
        "##test string with 2 dim-array[one row, one thread to process array data]"
      ],
      "metadata": {
        "id": "2LpaUTgCAzF9"
      }
    },
    {
      "cell_type": "code",
      "source": [
        "%%writefile s2-1thread.py\n",
        "#!python \n",
        "import pycuda.driver as cuda\n",
        "import pycuda.autoinit\n",
        "from pycuda.compiler import SourceModule\n",
        "from pycuda import gpuarray\n",
        "import numpy as np\n",
        "\n",
        "mod = SourceModule(\"\"\"\n",
        "    #include <stdio.h>\n",
        "    __global__ void say_hi(char*out_gpu)\n",
        "    { \n",
        "      const int strLen=13;\n",
        "      const int colLen=3;\n",
        "      const int rowCnt=2;\n",
        "      int idx = threadIdx.x + blockIdx.x * blockDim.x;\n",
        "\n",
        "      //printf(\"index: threadIdx.x=%d blockIdx.x=%d  blockDim.x=%d \\\\n\",threadIdx.x,blockIdx.x,blockDim.x);\n",
        "      printf(\"thread id: [%d]  get str:%s  strlen=%d \\\\n\",idx,out_gpu+strLen+1, strLen);\n",
        "\n",
        "      // string len is 6\n",
        "      char sub[strLen+1];\n",
        "      int strStartIndex=0;//string start locaion\n",
        "      int rowOffset=colLen*strLen;\n",
        "      if(idx > rowCnt){\n",
        "        printf(\"[ERROR]idx=%d, rowCnt=%d, we want one thread process all one rows'data .please tune\");\n",
        "      }\n",
        "      \n",
        "      int rowStartIndex=rowOffset*idx;\n",
        "      for(int c=0;c<colLen;c++){\n",
        "        for(int i = 0; i<strLen; i++){ \n",
        "            strStartIndex=c*strLen;\n",
        "            sub[i]=out_gpu[strStartIndex+i+rowStartIndex];\n",
        "        }\n",
        "        printf(\"thread[%d] rowStartIndex=%d, substring:%s strStartIndex %d\\\\n\",idx,rowStartIndex,sub,strStartIndex);\n",
        "      }\n",
        "    }\n",
        "    \"\"\")\n",
        "\n",
        "say_hi = mod.get_function(\"say_hi\")\n",
        "\n",
        "#padding string with \\0 char\n",
        "s= np.array(['abc\\0','edfg\\0','stz12345678b\\0','b1\\0','k2\\0','hellworld3\\0'], dtype=np.string_)\n",
        "\n",
        "s=s.reshape(2,3)\n",
        "print('shape:',s.shape,s.size)\n",
        "print(s,s.data)\n",
        "\n",
        "s_gpu = gpuarray.to_gpu(s)\n",
        "say_hi(s_gpu,block=(2,1,1),grid=(1,1,1))\n",
        "print('modify:',s_gpu.get())"
      ],
      "metadata": {
        "colab": {
          "base_uri": "https://localhost:8080/"
        },
        "id": "5LLCos__BhuG",
        "outputId": "a3608e13-8bac-4dde-9e79-226e0865f433"
      },
      "execution_count": 237,
      "outputs": [
        {
          "output_type": "stream",
          "name": "stdout",
          "text": [
            "Overwriting s2-1thread.py\n"
          ]
        }
      ]
    },
    {
      "cell_type": "code",
      "source": [
        "!python s2-1thread.py"
      ],
      "metadata": {
        "colab": {
          "base_uri": "https://localhost:8080/"
        },
        "id": "CF8blkLVBuwe",
        "outputId": "a49d0f02-d6c6-4ea3-8b3c-a862c1c8c9b3"
      },
      "execution_count": 238,
      "outputs": [
        {
          "output_type": "stream",
          "name": "stdout",
          "text": [
            "shape: (2, 3) 6\n",
            "[[b'abc' b'edfg' b'stz12345678b']\n",
            " [b'b1' b'k2' b'hellworld3']] <memory at 0x7fa1ff91dc70>\n",
            "thread id: [0]  get str:dfg  strlen=13 \n",
            "thread id: [1]  get str:dfg  strlen=13 \n",
            "thread[0] rowStartIndex=0, substring:abc strStartIndex 0\n",
            "thread[1] rowStartIndex=39, substring:b1 strStartIndex 0\n",
            "thread[0] rowStartIndex=0, substring:edfg strStartIndex 13\n",
            "thread[1] rowStartIndex=39, substring:k2 strStartIndex 13\n",
            "thread[0] rowStartIndex=0, substring:stz12345678b strStartIndex 26\n",
            "thread[1] rowStartIndex=39, substring:hellworld3 strStartIndex 26\n",
            "modify: [[b'abc' b'edfg' b'stz12345678b']\n",
            " [b'b1' b'k2' b'hellworld3']]\n"
          ]
        }
      ]
    },
    {
      "cell_type": "markdown",
      "source": [
        "## test struct with number\n",
        "https://github.com/inducer/pycuda/blob/main/examples/demo_struct.py"
      ],
      "metadata": {
        "id": "UIXOg0vKeURO"
      }
    },
    {
      "cell_type": "code",
      "source": [
        "%%writefile demo_struct.py\n",
        "# prepared invocations and structures -----------------------------------------\n",
        "import pycuda.driver as cuda\n",
        "import pycuda.autoinit\n",
        "import numpy\n",
        "import struct\n",
        "from pycuda.compiler import SourceModule\n",
        "\n",
        "class DoubleOpStruct:\n",
        "    mem_size = 8 + numpy.uintp(0).nbytes\n",
        "    def __init__(self, array, struct_arr_ptr):\n",
        "        self.data = cuda.to_device(array)\n",
        "        self.shape, self.dtype = array.shape, array.dtype\n",
        "\n",
        "        packed_args = struct.pack(\"ixP\", array.size, numpy.uintp(self.data))\n",
        "        cuda.memcpy_htod(struct_arr_ptr, packed_args)\n",
        "\n",
        "    def __str__(self):\n",
        "        return str(cuda.from_device(self.data, self.shape, self.dtype))\n",
        "\n",
        "struct_arr = cuda.mem_alloc(2 * DoubleOpStruct.mem_size)\n",
        "do2_ptr = int(struct_arr) + DoubleOpStruct.mem_size\n",
        "\n",
        "array1 = DoubleOpStruct(numpy.array([1, 2, 3], dtype=numpy.float32), struct_arr)\n",
        "array2 = DoubleOpStruct(numpy.array([0, 4], dtype=numpy.float32), do2_ptr)\n",
        "\n",
        "print(\"original arrays\")\n",
        "print(array1)\n",
        "print(array2)\n",
        "\n",
        "mod = SourceModule(\"\"\"\n",
        "    struct DoubleOperation {\n",
        "        int datalen, __padding; // so 64-bit ptrs can be aligned\n",
        "        float *ptr;\n",
        "    };\n",
        "\n",
        "\n",
        "    __global__ void double_array(DoubleOperation *a)\n",
        "    {\n",
        "        a = a + blockIdx.x;\n",
        "        for (int idx = threadIdx.x; idx < a->datalen; idx += blockDim.x)\n",
        "        {\n",
        "            float *a_ptr = a->ptr;\n",
        "            a_ptr[idx] *= 2;\n",
        "        }\n",
        "    }\n",
        "    \"\"\")\n",
        "func = mod.get_function(\"double_array\")\n",
        "func(struct_arr, block=(32, 1, 1), grid=(2, 1))\n",
        "\n",
        "print(\"doubled arrays\")\n",
        "print(array1)\n",
        "print(array2)\n",
        "\n",
        "func(numpy.uintp(do2_ptr), block=(32, 1, 1), grid=(1, 1))\n",
        "print(\"doubled second only\")\n",
        "print(array1)\n",
        "print(array2)\n",
        "\n",
        "if cuda.get_version() < (4, ):\n",
        "    func.prepare(\"P\", block=(32, 1, 1))\n",
        "    func.prepared_call((2, 1), struct_arr)\n",
        "else:\n",
        "    func.prepare(\"P\")\n",
        "    block = (32, 1, 1)\n",
        "    func.prepared_call((2, 1), block, struct_arr)\n",
        "\n",
        "\n",
        "print(\"doubled again\")\n",
        "print(array1)\n",
        "print(array2)\n",
        "\n",
        "if cuda.get_version() < (4, ):\n",
        "    func.prepared_call((1, 1), do2_ptr)\n",
        "else:\n",
        "    func.prepared_call((1, 1), block, do2_ptr)\n",
        "\n",
        "\n",
        "print(\"doubled second only again\")\n",
        "print(array1)\n",
        "print(array2)\n"
      ],
      "metadata": {
        "colab": {
          "base_uri": "https://localhost:8080/"
        },
        "id": "Ru4QqnvFwuka",
        "outputId": "a0c11338-3f9a-4bd9-ac4f-5661da7b5764"
      },
      "execution_count": 239,
      "outputs": [
        {
          "output_type": "stream",
          "name": "stdout",
          "text": [
            "Overwriting demo_struct.py\n"
          ]
        }
      ]
    },
    {
      "cell_type": "code",
      "source": [
        "!python demo_struct.py"
      ],
      "metadata": {
        "colab": {
          "base_uri": "https://localhost:8080/"
        },
        "id": "TifAInm2w9lS",
        "outputId": "9d22fdf7-42b1-47d9-ab65-ec458246e0b9"
      },
      "execution_count": 240,
      "outputs": [
        {
          "output_type": "stream",
          "name": "stdout",
          "text": [
            "original arrays\n",
            "[1. 2. 3.]\n",
            "[0. 4.]\n",
            "doubled arrays\n",
            "[2. 4. 6.]\n",
            "[0. 8.]\n",
            "doubled second only\n",
            "[2. 4. 6.]\n",
            "[ 0. 16.]\n",
            "doubled again\n",
            "[ 4.  8. 12.]\n",
            "[ 0. 32.]\n",
            "doubled second only again\n",
            "[ 4.  8. 12.]\n",
            "[ 0. 64.]\n"
          ]
        }
      ]
    },
    {
      "cell_type": "markdown",
      "source": [
        "## test struct with string"
      ],
      "metadata": {
        "id": "o2nYLRLPxTwJ"
      }
    },
    {
      "cell_type": "markdown",
      "source": [
        "### reference:\n",
        "https://stackoverflow.com/questions/34700999/cuda-passing-char-to-kernel/34712905\n",
        "https://documen.tician.de/pycuda/tutorial.html#structures\n",
        "\n",
        "Here numpy uses a **special fixed length string** data type whose length is calculated from the input data. This is effective a C ordered array of **char**[7]. See more here. PyCUDA automagically understands how to handle this because of the buffer protocol and the underlying direct mapping to a native C type\n",
        "\n",
        "https://stackoverflow.com/questions/48038577/pycuda-numpy-and-working-with-strings-in-general\n",
        "\n"
      ],
      "metadata": {
        "id": "x_WJoSceKlMM"
      }
    }
  ],
  "metadata": {
    "colab": {
      "name": "Welcome To Colaboratory",
      "provenance": [],
      "include_colab_link": true
    },
    "kernelspec": {
      "display_name": "Python 3",
      "name": "python3"
    },
    "accelerator": "GPU",
    "gpuClass": "standard"
  },
  "nbformat": 4,
  "nbformat_minor": 0
}