guih2 / repository

sashahush

por Sasha Rush - srush_nlp As arquiteturas de GPU são críticas para o aprendizado de máquina e parecem estar se tornando ainda mais importantes a cada dia. No entanto, você pode ser um especialista em aprendizado de máquina sem nunca tocar no código da GPU. É um pouco estranho trabalhar sempre através da abstração. Este notebook é uma tentativa de ensinar programação de GPU para iniciantes de uma forma completamente interativa. Em vez de fornecer texto com conceitos, ele o joga direto na codificação e na construção de kernels de GPU. Os exercícios usam NUMBA que mapeia diretamente o código Python para os kernels CUDA. Parece Python, mas é basicamente idêntico a escrever código CUDA de baixo nível. Em algumas horas, acho que você pode ir do básico ao entendimento dos algoritmos reais que alimentam 99% do aprendizado profundo hoje. Se você quiser ler o manual, está aqui: Guia NUMBA CUDA Eu recomendo fazer isso no Colab, pois é fácil começar. Certifique-se de fazer sua própria cópia, ative o modo GPU nas configurações ( Runtime / Change runtime type, defina Hardware acceleratorcomo GPU) e, em seguida, vá para a codificação. Abrir no Colab (Se você gosta desse estilo de quebra-cabeça, confira também meus Tensor Puzzles for PyTorch.) !pip install -qqq git+https://github.com/danoneata/chalk@srush-patch-1 !wget -q https://github.com/srush/GPU-Puzzles/raw/main/robot.png https://github.com/srush/GPU-Puzzles/raw/main/lib.py import numba import numpy as np import warnings from lib import CudaProblem, Coord warnings.filterwarnings( action="ignore", category=numba.NumbaPerformanceWarning, module="numba" ) Quebra-cabeça 1: Mapa Implemente um "kernel" (função GPU) que adicione 10 a cada posição do vetor a e armazene-o em vetor out. Você tem 1 thread por posição. Atenção Este código se parece com Python, mas na verdade é CUDA! Você não pode usar ferramentas padrão do python, como compreensão de lista ou solicitar propriedades do Numpy, como forma ou tamanho (se você precisar do tamanho, ele será fornecido como um argumento). Os quebra-cabeças exigem apenas operações simples, basicamente +, *, indexação de array simples, loops for e instruções if. Se você receber um erro, provavelmente é porque você fez algo extravagante :). Dica: Pense na função callcomo sendo executada 1 vez para cada thread. A única diferença é que cuda.threadIdx.xmuda a cada vez. def map_spec(a): return a + 10 def map_test(cuda): def call(out, a) -> None: local_i = cuda.threadIdx.x # FILL ME IN (roughly 1 lines) return call SIZE = 4 out = np.zeros((SIZE,)) a = np.arange(SIZE) problem = CudaProblem( "Map", map_test, [a], out, threadsperblock=Coord(SIZE, 1), spec=map_spec ) problem.show() # Map Score (Max Per Thread): | Global Reads | Global Writes | Shared Reads | Shared Writes | | 0 | 0 | 0 | 0 | svg problem.check() Failed Tests. Yours: [0. 0. 0. 0.] Spec : [10 11 12 13] Quebra-cabeça 2 - Zíper Implemente um kernel que some cada posição de aand be o armazene em out. Você tem 1 thread por posição. def zip_spec(a, b): return a + b def zip_test(cuda): def call(out, a, b) -> None: local_i = cuda.threadIdx.x # FILL ME IN (roughly 1 lines) return call SIZE = 4 out = np.zeros((SIZE,)) a = np.arange(SIZE) b = np.arange(SIZE) problem = CudaProblem( "Zip", zip_test, [a, b], out, threadsperblock=Coord(SIZE, 1), spec=zip_spec ) problem.show() # Zip Score (Max Per Thread): | Global Reads | Global Writes | Shared Reads | Shared Writes | | 0 | 0 | 0 | 0 | svg problem.check() Failed Tests. Yours: [0. 0. 0. 0.] Spec : [0 2 4 6] Quebra-cabeça 3 - Guardas Implemente um kernel que adicione 10 a cada posição ae o armazene em out. Você tem mais threads do que posições. def map_guard_test(cuda): def call(out, a, size) -> None: local_i = cuda.threadIdx.x # FILL ME IN (roughly 2 lines) return call SIZE = 4 out = np.zeros((SIZE,)) a = np.arange(SIZE) problem = CudaProblem( "Guard", map_guard_test, [a], out, [SIZE], threadsperblock=Coord(8, 1), spec=map_spec, ) problem.show() # Guard Score (Max Per Thread): | Global Reads | Global Writes | Shared Reads | Shared Writes | | 0 | 0 | 0 | 0 | svg problem.check() Failed Tests. Yours: [0. 0. 0. 0.] Spec : [10 11 12 13] Quebra-cabeça 4 - Mapa 2D Implemente um kernel que adicione 10 a cada posição ae o armazene em out. A entrada aé 2D e quadrada. Você tem mais threads do que posições. def map_2D_test(cuda): def call(out, a, size) -> None: local_i = cuda.threadIdx.x local_j = cuda.threadIdx.y # FILL ME IN (roughly 2 lines) return call SIZE = 2 out = np.zeros((SIZE, SIZE)) a = np.arange(SIZE * SIZE).reshape((SIZE, SIZE)) problem = CudaProblem( "Map 2D", map_2D_test, [a], out, [SIZE], threadsperblock=Coord(3, 3), spec=map_spec ) problem.show() # Map 2D Score (Max Per Thread): | Global Reads | Global Writes | Shared Reads | Shared Writes | | 0 | 0 | 0 | 0 | svg problem.check() Failed Tests. Yours: [[0. 0.] [0. 0.]] Spec : [[10 11] [12 13]] Quebra-cabeça 5 - Transmissão Implemente um kernel que adiciona aand be armazena em out. As entradas ae bsão vetores. Você tem mais threads do que posições. def broadcast_test(cuda): def call(out, a, b, size) -> None: local_i = cuda.threadIdx.x local_j = cuda.threadIdx.y # FILL ME IN (roughly 2 lines) return call SIZE = 2 out = np.zeros((SIZE, SIZE)) a = np.arange(SIZE).reshape(SIZE, 1) b = np.arange(SIZE).reshape(1, SIZE) problem = CudaProblem( "Broadcast", broadcast_test, [a, b], out, [SIZE], threadsperblock=Coord(3, 3), spec=zip_spec, ) problem.show() # Broadcast Score (Max Per Thread): | Global Reads | Global Writes | Shared Reads | Shared Writes | | 0 | 0 | 0 | 0 | svg problem.check() Failed Tests. Yours: [[0. 0.] [0. 0.]] Spec : [[0 1] [1 2]] Quebra-cabeça 6 - Blocos Implemente um kernel que adicione 10 a cada posição ae o armazene em out. Você tem menos threads por bloco do que o tamanho de a. Dica: Um bloco é um grupo de threads. O número de threads por bloco é limitado, mas podemos ter muitos blocos diferentes. A variável cuda.blockIdxnos diz em que bloco estamos. def map_block_test(cuda): def call(out, a, size) -> None: i = cuda.blockIdx.x * cuda.blockDim.x + cuda.threadIdx.x # FILL ME IN (roughly 2 lines) return call SIZE = 9 out = np.zeros((SIZE,)) a = np.arange(SIZE) problem = CudaProblem( "Blocks", map_block_test, [a], out, [SIZE], threadsperblock=Coord(4, 1), blockspergrid=Coord(3, 1), spec=map_spec, ) problem.show() # Blocks Score (Max Per Thread): | Global Reads | Global Writes | Shared Reads | Shared Writes | | 0 | 0 | 0 | 0 | svg problem.check() Failed Tests. Yours: [0. 0. 0. 0. 0. 0. 0. 0. 0.] Spec : [10 11 12 13 14 15 16 17 18] Quebra-cabeça 7 - Blocos 2D Implemente o mesmo kernel em 2D. Você tem menos threads por bloco do que o tamanho de aem ambas as direções. def map_block2D_test(cuda): def call(out, a, size) -> None: i = cuda.blockIdx.x * cuda.blockDim.x + cuda.threadIdx.x # FILL ME IN (roughly 4 lines) return call SIZE = 5 out = np.zeros((SIZE, SIZE)) a = np.ones((SIZE, SIZE)) problem = CudaProblem( "Blocks 2D", map_block2D_test, [a], out, [SIZE], threadsperblock=Coord(3, 3), blockspergrid=Coord(2, 2), spec=map_spec, ) problem.show() # Blocks 2D Score (Max Per Thread): | Global Reads | Global Writes | Shared Reads | Shared Writes | | 0 | 0 | 0 | 0 | svg problem.check() Failed Tests. Yours: [[0. 0. 0. 0. 0.] [0. 0. 0. 0. 0.] [0. 0. 0. 0. 0.] [0. 0. 0. 0. 0.] [0. 0. 0. 0. 0.]] Spec : [[11. 11. 11. 11. 11.] [11. 11. 11. 11. 11.] [11. 11. 11. 11. 11.] [11. 11. 11. 11. 11.] [11. 11. 11. 11. 11.]] Quebra-cabeça 8 - Compartilhado Implemente um kernel que adicione 10 a cada posição ae o armazene em out. Você tem menos threads por bloco do que o tamanho de a. Dica: Cada bloco pode ter uma quantidade constante de memória compartilhada que somente as threads desse bloco podem ler e escrever. Depois de escrever você deve usar cuda.syncthreadspara garantir que os threads não se cruzem. (Este exemplo realmente não precisa de memória compartilhada ou syncthreads, mas é uma demonstração.) TPB = 4 def shared_test(cuda): def call(out, a, size) -> None: shared = cuda.shared.array(TPB, numba.float32) i = cuda.blockIdx.x * cuda.blockDim.x + cuda.threadIdx.x local_i = cuda.threadIdx.x if i < size: shared[local_i] = a[i] cuda.syncthreads() # FILL ME IN (roughly 2 lines) return call SIZE = 8 out = np.zeros(SIZE) a = np.ones(SIZE) problem = CudaProblem( "Shared", shared_test, [a], out, [SIZE], threadsperblock=Coord(TPB, 1), blockspergrid=Coord(2, 1), spec=map_spec, ) problem.show() # Shared Score (Max Per Thread): | Global Reads | Global Writes | Shared Reads | Shared Writes | | 1 | 0 | 0 | 1 | svg problem.check() Failed Tests. Yours: [0. 0. 0. 0. 0. 0. 0. 0.] Spec : [11. 11. 11. 11. 11. 11. 11. 11.] Quebra-cabeça 9 - Agrupamento Implemente um kernel que some as 3 últimas posições ae as armazene em out. Você tem 1 thread por posição. Você só precisa de 1 leitura global e 1 gravação global por thread. Dica: Lembre-se de ter cuidado com a sincronização. def pool_spec(a): out = np.zeros(*a.shape) for i in range(a.shape[0]): out[i] = a[max(i - 2, 0) : i + 1].sum() return out TPB = 8 def pool_test(cuda): def call(out, a, size) -> None: shared = cuda.shared.array(TPB, numba.float32) i = cuda.blockIdx.x * cuda.blockDim.x + cuda.threadIdx.x local_i = cuda.threadIdx.x # FILL ME IN (roughly 8 lines) return call SIZE = 8 out = np.zeros(SIZE) a = np.arange(SIZE) problem = CudaProblem( "Pooling", pool_test, [a], out, [SIZE], threadsperblock=Coord(TPB, 1), blockspergrid=Coord(1, 1), spec=pool_spec, ) problem.show() # Pooling Score (Max Per Thread): | Global Reads | Global Writes | Shared Reads | Shared Writes | | 0 | 0 | 0 | 0 | svg problem.check() Failed Tests. Yours: [0. 0. 0. 0. 0. 0. 0. 0.] Spec : [ 0. 1. 3. 6. 9. 12. 15. 18.] Quebra-cabeça 10 - Produto escalar Implemente um kernel que calcule o produto escalar de aand be o armazene em out. Você tem 1 thread por posição. Você só precisa de 1 leitura global e 1 gravação global por thread. Observação: para esse problema, você não precisa se preocupar com o número de leituras compartilhadas. Vamos lidar com esse desafio mais tarde. def dot_spec(a, b): return a @ b TPB = 8 def dot_test(cuda): def call(out, a, b, size) -> None: a_shared = cuda.shared.array(TPB, numba.float32) b_shared = cuda.shared.array(TPB, numba.float32) i = cuda.blockIdx.x * cuda.blockDim.x + cuda.threadIdx.x local_i = cuda.threadIdx.x # FILL ME IN (roughly 9 lines) return call SIZE = 8 out = np.zeros(1) a = np.arange(SIZE) b = np.arange(SIZE) problem = CudaProblem( "Dot", dot_test, [a, b], out, [SIZE], threadsperblock=Coord(SIZE, 1), blockspergrid=Coord(1, 1), spec=dot_spec, ) problem.show() # Dot Score (Max Per Thread): | Global Reads | Global Writes | Shared Reads | Shared Writes | | 0 | 0 | 0 | 0 | svg problem.check() Failed Tests. Yours: [0.] Spec : 140 Quebra-cabeça 11 - Convolução 1D Implemente um kernel que calcula uma convolução 1D entre aand be armazena em out. Você precisa lidar com o caso geral. Você só precisa de 2 leituras globais e 1 gravação global por encadeamento. def conv_spec(a, b): out = np.zeros(*a.shape) len = b.shape[0] for i in range(a.shape[0]): out[i] = sum([a[i + j] * b[j] for j in range(len) if i + j < a.shape[0]]) return out MAX_CONV = 5 TPB = 8 TPB_MAX_CONV = TPB + MAX_CONV def conv_test(cuda): def call(out, a, b, a_size, b_size) -> None: i = cuda.blockIdx.x * cuda.blockDim.x + cuda.threadIdx.x local_i = cuda.threadIdx.x # FILL ME IN (roughly 17 lines) return call Teste 1 SIZE = 6 CONV = 3 out = np.zeros(SIZE) a = np.arange(SIZE) b = np.arange(CONV) problem = CudaProblem( "1D Conv (Simple)", conv_test, [a, b], out, [SIZE, CONV], Coord(1, 1), Coord(TPB, 1), spec=conv_spec, ) problem.show() # 1D Conv (Simple) Score (Max Per Thread): | Global Reads | Global Writes | Shared Reads | Shared Writes | | 0 | 0 | 0 | 0 | svg problem.check() Failed Tests. Yours: [0. 0. 0. 0. 0. 0.] Spec : [ 5. 8. 11. 14. 5. 0.] Teste 2 out = np.zeros(15) a = np.arange(15) b = np.arange(4) problem = CudaProblem( "1D Conv (Full)", conv_test, [a, b], out, [15, 4], Coord(2, 1), Coord(TPB, 1), spec=conv_spec, ) problem.show() # 1D Conv (Full) Score (Max Per Thread): | Global Reads | Global Writes | Shared Reads | Shared Writes | | 0 | 0 | 0 | 0 | svg problem.check() Failed Tests. Yours: [0. 0. 0. 0. 0. 0. 0. 0. 0. 0. 0. 0. 0. 0. 0.] Spec : [14. 20. 26. 32. 38. 44. 50. 56. 62. 68. 74. 80. 41. 14. 0.] Quebra-cabeça 12 - Soma do prefixo Implemente um kernel que calcula uma soma ae a armazena em out. Se o tamanho de afor maior que o tamanho do bloco, armazene apenas a soma de cada bloco. Faremos isso usando o algoritmo de soma de prefixo paralelo na memória compartilhada. Ou seja, cada passo do algoritmo deve somar metade dos números restantes. Siga este diagrama: TPB = 8 def sum_spec(a): out = np.zeros((a.shape[0] + TPB - 1) // TPB) for j, i in enumerate(range(0, a.shape[-1], TPB)): out[j] = a[i : i + TPB].sum() return out def sum_test(cuda): def call(out, a, size: int) -> None: cache = cuda.shared.array(TPB, numba.float32) i = cuda.blockIdx.x * cuda.blockDim.x + cuda.threadIdx.x local_i = cuda.threadIdx.x # FILL ME IN (roughly 12 lines) return call Teste 1 SIZE = 8 out = np.zeros(1) inp = np.arange(SIZE) problem = CudaProblem( "Sum (Simple)", sum_test, [inp], out, [SIZE], Coord(1, 1), Coord(TPB, 1), spec=sum_spec, ) problem.show() # Sum (Simple) Score (Max Per Thread): | Global Reads | Global Writes | Shared Reads | Shared Writes | | 0 | 0 | 0 | 0 | svg problem.check() Failed Tests. Yours: [0.] Spec : [28.] Teste 2 SIZE = 15 out = np.zeros(2) inp = np.arange(SIZE) problem = CudaProblem( "Sum (Full)", sum_test, [inp], out, [SIZE], Coord(2, 1), Coord(TPB, 1), spec=sum_spec, ) problem.show() # Sum (Full) Score (Max Per Thread): | Global Reads | Global Writes | Shared Reads | Shared Writes | | 0 | 0 | 0 | 0 | svg problem.check() Failed Tests. Yours: [0. 0.] Spec : [28. 77.] Quebra-cabeça 13 - Soma do Eixo Implemente um kernel que calcula uma soma em cada linha ae a armazena em out. TPB = 8 def sum_spec(a): out = np.zeros((a.shape[0], (a.shape[1] + TPB - 1) // TPB)) for j, i in enumerate(range(0, a.shape[-1], TPB)): out[..., j] = a[..., i : i + TPB].sum(-1) return out def axis_sum_test(cuda): def call(out, a, size: int) -> None: cache = cuda.shared.array(TPB, numba.float32) i = cuda.blockIdx.x * cuda.blockDim.x + cuda.threadIdx.x local_i = cuda.threadIdx.x batch = cuda.blockIdx.y # FILL ME IN (roughly 12 lines) return call BATCH = 4 SIZE = 6 out = np.zeros((BATCH, 1)) inp = np.arange(BATCH * SIZE).reshape((BATCH, SIZE)) problem = CudaProblem( "Axis Sum", axis_sum_test, [inp], out, [SIZE], Coord(1, BATCH), Coord(TPB, 1), spec=sum_spec, ) problem.show() # Axis Sum Score (Max Per Thread): | Global Reads | Global Writes | Shared Reads | Shared Writes | | 0 | 0 | 0 | 0 | svg problem.check() Failed Tests. Yours: [[0.] [0.] [0.] [0.]] Spec : [[ 15.] [ 51.] [ 87.] [123.]] Quebra-cabeça 14 - Multiplicação de Matrizes! Implemente um kernel que multiplique matrizes quadradas ae bo armazene em out. Dica: O algoritmo mais eficiente aqui copiará um bloco na memória compartilhada antes de computar cada um dos produtos escalares linha-coluna. Isso é fácil de fazer se a matriz couber na memória compartilhada. Faça esse caso primeiro. Em seguida, atualize seu código para calcular um produto escalar parcial e mova iterativamente a parte que você copiou para a memória compartilhada. Você deve ser capaz de fazer o hard case em 6 leituras globais. def matmul_spec(a, b): return a @ b TPB = 3 def mm_oneblock_test(cuda): def call(out, a, b, size: int) -> None: a_shared = cuda.shared.array((TPB, TPB), numba.float32) b_shared = cuda.shared.array((TPB, TPB), numba.float32) i = cuda.blockIdx.x * cuda.blockDim.x + cuda.threadIdx.x j = cuda.blockIdx.y * cuda.blockDim.y + cuda.threadIdx.y local_i = cuda.threadIdx.x local_j = cuda.threadIdx.y # FILL ME IN (roughly 14 lines) return call Teste 1 SIZE = 2 out = np.zeros((SIZE, SIZE)) inp1 = np.arange(SIZE * SIZE).reshape((SIZE, SIZE)) inp2 = np.arange(SIZE * SIZE).reshape((SIZE, SIZE)).T problem = CudaProblem( "Matmul (Simple)", mm_oneblock_test, [inp1, inp2], out, [SIZE], Coord(1, 1), Coord(TPB, TPB), spec=matmul_spec, ) problem.show(sparse=True) # Matmul (Simple) Score (Max Per Thread): | Global Reads | Global Writes | Shared Reads | Shared Writes | | 0 | 0 | 0 | 0 | svg problem.check() Failed Tests. Yours: [[0. 0.] [0. 0.]] Spec : [[ 1 3] [ 3 13]] # Test 2 SIZE = 8 out = np.zeros((SIZE, SIZE)) inp1 = np.arange(SIZE * SIZE).reshape((SIZE, SIZE)) inp2 = np.arange(SIZE * SIZE).reshape((SIZE, SIZE)).T problem = CudaProblem( "Matmul (Full)", mm_oneblock_test, [inp1, inp2], out, [SIZE], Coord(3, 3), Coord(TPB, TPB), spec=matmul_spec, ) problem.show(sparse=True) # Matmul (Full) Score (Max Per Thread): | Global Reads | Global Writes | Shared Reads | Shared Writes | | 0 | 0 | 0 | 0 | svg