Publish tokenizer_kernel kernel and performance card

Browse files

CUDA source, short description, performance table, and plot.

Files changed (3) hide show

README.md +40 -0
kernel.cu +248 -0
performance.svg +8 -0

README.md ADDED Viewed

	@@ -0,0 +1,40 @@

+---
+license: apache-2.0
+tags:
+- cuda
+- kernel
+- gpu-optimization
+- hpc
+---
+# tokenizer_kernel
+Experimental tokenizer/matrix CUDA kernel file.
+This repository contains the standalone CUDA source for the `tokenizer_kernel` lane from
+ the PyC kernel lab. It is a source artifact for inspection and benchmarking;
+it is not a precompiled binary and the result below is not a universal ranking.
+## Performance
+| Kernel | GPU / architecture | Shape | Best recorded result | Evidence |
+|---|---|---|---|---|
+| `tokenizer_kernel` | not recorded | not recorded | Not measured in the published campaign | No published performance receipt was found for this lane. |
+![Performance plot](performance.svg)
+The result is reported with the original campaign's timing and correctness
+context. Compare kernels only when GPU, CUDA version, matrix shape, warmup,
+repeats, and reference/correctness mode match.
+## Source
+- `kernel.cu` — copied from `kernels/prototypes/experimental/tokenizer_matmul/kernel.cu`.
+- Original lane tags: `cuda, tokenizer, experimental`.
+## Build/run contract
+```text
+{nvcc} -O2 -c {source} -o {build_dir}/{name}.o
+(compile-only)
+```

kernel.cu ADDED Viewed

	@@ -0,0 +1,248 @@

+#include <cuda_runtime.h>
+#include <stdio.h>
+#include <ctype.h>
+#define MAX_TOKENS 1024
+// Enhanced token types
+typedef enum {
+    TOKEN_IDENTIFIER = 0,
+    TOKEN_NUMBER = 1,
+    TOKEN_OPERATOR = 2,
+    TOKEN_KEYWORD = 3,
+    TOKEN_STRING = 4,
+    TOKEN_COMMENT = 5,
+    TOKEN_PREPROCESSOR = 6,
+    TOKEN_PUNCTUATION = 7
+} TokenType;
+// Add token metadata
+typedef struct {
+    TokenType type;
+    int start_pos;
+    int end_pos;
+    int length;
+    int line;
+    int column;
+    char lexeme[256];
+    unsigned int hash;
+} EnhancedTokenGPU;
+typedef struct {
+    int type; // 0: identifier, 1: number, 2: operator
+    int start_pos;
+    int end_pos;
+    int length;
+} TokenGPU;
+// Add shared memory optimization
+__shared__ char shared_input[1024];
+__shared__ int shared_token_count;
+// Enhanced tokenization kernel with better pattern matching
+__global__ void enhanced_tokenize_kernel(const char* input, size_t input_length,
+                                       EnhancedTokenGPU* tokens, int* token_count,
+                                       bool enable_comments, bool enable_preprocessing) {
+    int idx = blockIdx.x * blockDim.x + threadIdx.x;
+    if (idx >= input_length) return;
+    // Load chunk into shared memory
+    int local_idx = threadIdx.x;
+    if (local_idx < 1024 && idx < input_length) {
+        shared_input[local_idx] = input[idx];
+    }
+    __syncthreads();
+    // Enhanced token detection with more patterns
+    if (idx > 0 && (isalnum(shared_input[local_idx-1]) && isalnum(shared_input[local_idx]))) return;
+    int tcount = atomicAdd(token_count, 0);
+    if (tcount >= MAX_TOKENS) return;
+    EnhancedTokenGPU token;
+    token.start_pos = idx;
+    token.hash = 0;
+    // Calculate line and column
+    int line = 1, column = 1;
+    for (int i = 0; i < idx; i++) {
+        if (input[i] == '\n') {
+            line++;
+            column = 1;
+        } else {
+            column++;
+        }
+    }
+    token.line = line;
+    token.column = column;
+    // Enhanced pattern matching
+    if (isalpha(shared_input[local_idx]) || shared_input[local_idx] == '_') {
+        // Handle identifiers and keywords
+        int end = local_idx;
+        while (end < 1024 && (isalnum(shared_input[end]) || shared_input[end] == '_')) {
+            token.hash = token.hash * 31 + shared_input[end];
+            end++;
+        }
+        token.type = TOKEN_IDENTIFIER;
+        token.end_pos = idx + (end - local_idx) - 1;
+        token.length = end - local_idx;
+    }
+    // ... Add more token pattern matching ...
+    // Store token if valid
+    if (token.length > 0) {
+        int new_count = atomicAdd(token_count, 1);
+        if (new_count < MAX_TOKENS) {
+            tokens[new_count] = token;
+        }
+    }
+}
+__global__ void tokenize_kernel(const char* input, size_t input_length, TokenGPU* tokens, int* token_count) {
+    int idx = blockIdx.x * blockDim.x + threadIdx.x;
+    if (idx >= input_length) return;
+    // Skip if not at token boundary
+    if (idx > 0 && (isalnum(input[idx-1]) && isalnum(input[idx]))) return;
+    int tcount = *token_count;
+    if (tcount >= MAX_TOKENS) return;
+    if (isalpha(input[idx])) {
+        int end = idx;
+        while (end < input_length && isalnum(input[end])) end++;
+        int new_count = atomicAdd(token_count, 1);
+        if (new_count < MAX_TOKENS) {
+            tokens[new_count].type = 0;
+            tokens[new_count].start_pos = idx;
+            tokens[new_count].end_pos = end - 1;
+            tokens[new_count].length = end - idx;
+        }
+    } else if (isdigit(input[idx])) {
+        int end = idx;
+        while (end < input_length && isdigit(input[end])) end++;
+        int new_count = atomicAdd(token_count, 1);
+        if (new_count < MAX_TOKENS) {
+            tokens[new_count].type = 1;
+            tokens[new_count].start_pos = idx;
+            tokens[new_count].end_pos = end - 1;
+            tokens[new_count].length = end - idx;
+        }
+    } else if (input[idx] == '+' || input[idx] == '-' || input[idx] == '*' || input[idx] == '/') {
+        int new_count = atomicAdd(token_count, 1);
+        if (new_count < MAX_TOKENS) {
+            tokens[new_count].type = 2;
+            tokens[new_count].start_pos = idx;
+            tokens[new_count].end_pos = idx;
+            tokens[new_count].length = 1;
+        }
+    }
+}
+__global__ void matrix_mult_kernel(float* a, float* b, float* c, int m, int n, int k) {
+    int row = blockIdx.y * blockDim.y + threadIdx.y;
+    int col = blockIdx.x * blockDim.x + threadIdx.x;
+    if (row < m && col < n) {
+        float sum = 0.0f;
+        for (int i = 0; i < k; i++) {
+            sum += a[row * k + i] * b[i * n + col];
+        }
+        c[row * n + col] = sum;
+    }
+}
+// Add parallel matrix operations
+__global__ void enhanced_matrix_mult_kernel(float* a, float* b, float* c,
+                                          int m, int n, int k,
+                                          bool use_shared_memory) {
+    // ... existing matrix multiplication code ...
+    // Add shared memory optimization
+    __shared__ float shared_a[16][16];
+    __shared__ float shared_b[16][16];
+    // ... implement block matrix multiplication ...
+}
+// Add new CUDA utilities
+void initialize_cuda_context(void) {
+    cudaFree(0); // Force context initialization
+}
+void optimize_kernel_launch(dim3* blocks, dim3* threads, size_t shared_memory_size) {
+    int device;
+    cudaGetDevice(&device);
+    cudaDeviceProp props;
+    cudaGetDeviceProperties(&props, device);
+    // Optimize launch configuration based on device properties
+    // ... implementation ...
+}
+void cuda_tokenize(const char* input, TokenGPU* tokens, int* token_count) {
+    size_t input_length = strlen(input);
+    char* d_input;
+    TokenGPU* d_tokens;
+    int* d_token_count;
+    cudaMalloc(&d_input, input_length + 1);
+    cudaMalloc(&d_tokens, MAX_TOKENS * sizeof(TokenGPU));
+    cudaMalloc(&d_token_count, sizeof(int));
+    cudaMemcpy(d_input, input, input_length + 1, cudaMemcpyHostToDevice);
+    cudaMemset(d_token_count, 0, sizeof(int));
+    int threads = 256;
+    int blocks = (input_length + threads - 1) / threads;
+    tokenize_kernel<<<blocks, threads>>>(d_input, input_length, d_tokens, d_token_count);
+    cudaMemcpy(token_count, d_token_count, sizeof(int), cudaMemcpyDeviceToHost);
+    cudaMemcpy(tokens, d_tokens, *token_count * sizeof(TokenGPU), cudaMemcpyDeviceToHost);
+    cudaFree(d_input);
+    cudaFree(d_tokens);
+    cudaFree(d_token_count);
+}
+void cuda_matrix_mult(float* a, float* b, float* c, int m, int n, int k) {
+    float *d_a, *d_b, *d_c;
+    cudaMalloc(&d_a, m * k * sizeof(float));
+    cudaMalloc(&d_b, k * n * sizeof(float));
+    cudaMalloc(&d_c, m * n * sizeof(float));
+    cudaMemcpy(d_a, a, m * k * sizeof(float), cudaMemcpyHostToDevice);
+    cudaMemcpy(d_b, b, k * n * sizeof(float), cudaMemcpyHostToDevice);
+    dim3 threads(16, 16);
+    dim3 blocks((n + threads.x - 1) / threads.x, (m + threads.y - 1) / threads.y);
+    matrix_mult_kernel<<<blocks, threads>>>(d_a, d_b, d_c, m, n, k);
+    cudaMemcpy(c, d_c, m * n * sizeof(float), cudaMemcpyDeviceToHost);
+    cudaFree(d_a); cudaFree(d_b); cudaFree(d_c);
+}
+void print_tokens(TokenGPU* tokens, int token_count) {
+    for (int i = 0; i < token_count; i++) {
+        printf("Token %d: Type %d, Start %d, End %d, Length %d\n",
+               i, tokens[i].type, tokens[i].start_pos, tokens[i].end_pos, tokens[i].length);
+    }
+}
+int main() {
+    const char* input = "int a = 5 + 3;";
+    TokenGPU tokens[MAX_TOKENS];
+    int token_count;
+    cuda_tokenize(input, tokens, &token_count);
+    print_tokens(tokens, token_count);
+    float a[6] = {1, 2, 3, 4, 5, 6};
+    float b[6] = {7, 8, 9, 10, 11, 12};
+    float c[4] = {0};
+    cuda_matrix_mult(a, b, c, 2, 3, 2);
+    for (int i = 0; i < 4; i++) {
+        printf("%f ", c[i]);
+    }
+    printf("\n");
+    return 0;
+}
+// Compile with nvcc -o kernel kernel.cu

performance.svg ADDED Viewed