[G1-Sync] Manual knowledge update
This commit is contained in:
Antigravity Agent
@@ -2,129 +2,234 @@
|
||||
id: wiki-2026-0508-compute-shaders
|
||||
title: Compute Shaders
|
||||
category: 10_Wiki/Topics
|
||||
status: needs_review
|
||||
status: verified
|
||||
canonical_id: self
|
||||
aliases: []
|
||||
aliases: [GPU Compute, GPGPU Shaders, WebGPU Compute]
|
||||
duplicate_of: none
|
||||
source_trust_level: A
|
||||
confidence_score: 0.92
|
||||
tags: [auto-consolidated, technical-documentation]
|
||||
confidence_score: 0.9
|
||||
verification_status: applied
|
||||
tags: [gpu, shaders, webgpu, parallel, wgsl, cuda]
|
||||
raw_sources: []
|
||||
last_reinforced: 2026-05-08
|
||||
last_reinforced: 2026-05-10
|
||||
github_commit: pending
|
||||
inferred_by: Claude Opus 4.7 (auto-normalize 2026-05-08)
|
||||
tech_stack:
|
||||
language: unspecified
|
||||
framework: unspecified
|
||||
language: wgsl-glsl-cuda
|
||||
framework: webgpu-vulkan
|
||||
---
|
||||
|
||||
# [[Compute Shaders|Compute Shaders]]
|
||||
# Compute Shaders
|
||||
|
||||
## 📌 한 줄 통찰 (The Karpathy Summary)
|
||||
> 컴퓨트 셰이더(Compute Shaders)는 [[WebGPU|WebGPU]] 환경에서 지원되는 기능으로, CPU의 메인 스레드에서 수행되던 무거운 범용 연산 작업을 GPU로 오프로드하는 핵심 기술입니다 [1, 2]. GPU의 수천 개 코어를 활용한 병렬 처리를 통해 물리 시뮬레이션, 충돌 감지, 대규모 파티클 시스템 등의 작업 성능을 비약적으로 향상시킵니다 [2]. 또한 간접 그리기(Indirect Drawing) 기술과 결합하여 CPU의 개입 없이 가시성을 판별하고 화면을 그리는 완전한 GPU 주도 렌더링([[GPU-driven Rendering|GPU-driven Rendering]]) 파이프라인을 구축하는 데 사용됩니다 [3, 4].
|
||||
## 매 한 줄
|
||||
> **"매 GPU program 매 graphics pipeline 의 X — 매 arbitrary parallel computation."**. Compute shader는 vertex/fragment shader 와 다르게 rendering 의 X — 매 raw SIMT 의 power. 2026년 WebGPU (browser GPU compute), CUDA, Vulkan compute, Metal compute, ML inference, particle sims, image processing 의 dominant. ML 의 attention/matmul 도 compute shader 본질.
|
||||
|
||||
---
|
||||
## 매 핵심
|
||||
|
||||
> 컴퓨트 셰이더(Compute Shaders)는 [[JavaScript|JavaScript]] 메인 스레드에서 수행되던 무거운 작업을 수천 개의 GPU 코어에서 병렬로 처리하도록 오프로드하는 범용 GPU 연산(general-Purpose GPU computation) 기술입니다 [1]. 주로 [[WebGPU|WebGPU]] 환경에서 사용되며, 파티클 시스템, 물리 시뮬레이션, 대규모 데이터 필터링 등의 CPU 병목 현상을 획기적으로 해결하여 렌더링 성능을 극대화하는 데 필수적인 역할을 합니다 [2-4].
|
||||
### 매 model
|
||||
- **Workgroup**: 매 group of threads 가 same shader 실행 (e.g. 64 or 256 threads).
|
||||
- **Invocation**: single thread.
|
||||
- **Shared memory** (workgroup): fast, intra-group.
|
||||
- **Storage buffer**: GPU global memory (read/write).
|
||||
- **Uniform buffer**: small, read-only constants.
|
||||
- **Dispatch**: CPU 가 launches N workgroups.
|
||||
|
||||
## 📖 구조화된 지식 (Synthesized Content)
|
||||
* **범용 GPU 연산 및 성능 향상:** 컴퓨트 셰이더는 물리 시뮬레이션, 충돌 감지, 유체 시뮬레이션, 이미지 처리, 대규모 데이터 필터링, 절차적 지형 생성 등 복잡한 연산을 CPU 대신 GPU에서 병렬로 처리합니다 [2, 5-8]. 기존 CPU 기반 파티클 업데이트는 약 5만 개 수준에서 병목이 발생하지만, WebGPU 컴퓨트 셰이더를 활용하면 10만 개의 파티클을 2ms 이내에 업데이트하여 최대 150배의 성능 향상을 내며 수백만 개의 유닛을 처리할 수 있습니다 [9-12].
|
||||
* **GPU 주도 렌더링 및 컬링 (GPU-driven Rendering & Culling):** 간접 그리기(Indirect Drawing) 명령과 결합하여 극도로 효율적인 렌더링 파이프라인을 구성합니다 [4, 13]. 컴퓨트 셰이더가 모든 인스턴스에 대해 시야 절두체(Frustum) 및 오클루전(Occlusion) 컬링 판별을 수행하고, 화면에 보이는 객체 정보만 원자적 카운터(Atomic Counter)를 통해 간접 그리기 버퍼에 추가합니다 [3, 4, 14]. 이를 통해 CPU와 GPU 간의 데이터 동기화 지연과 명령 발행 오버헤드가 사실상 0에 수렴하게 됩니다 [4, 15].
|
||||
* **데이터 공유 및 메모리 최적화:** 읽기와 쓰기가 모두 가능한 스토리지 텍스처([[Storage|Storage]] Textures)를 활용해 GPU 기반 렌더링과 효과 처리를 유연하게 수행합니다 [6, 16]. 또한 스레드 간 데이터 공유가 필요한 경우, 반복 접근 패턴에서 전역 메모리보다 10~100배 더 빠른 작업 그룹 공유 메모리(Workgroup Shared [[memory|memory]])를 활용할 수 있습니다 [7, 13].
|
||||
* **고급 연산 기법 지원:** 컴퓨트 단계에서 메쉬 정점 변환을 처리하고 그 결과를 버퍼에 저장해 불필요한 중복 연산을 제거하는 '컴퓨트 스키닝(Compute Skinning)'이 가능해집니다 [12]. 또한 glTF 모델에 흔히 쓰이는 8비트/16비트 정수 데이터를 32비트 포맷으로 압축 해제하는 작업도 렌더링 파이프라인 외곽에서 효율적으로 수행할 수 있습니다 [12].
|
||||
* **동기화 및 파이프라인 제어 베스트 프랙티스:** 연산 의존성이 높은 씬을 Three.js에서 렌더링할 때는 `renderAsync`를 사용하여 렌더 패스가 시작되기 전에 컴퓨트 패스가 완전히 끝나도록 동기화해야 합니다 [17]. 성능 처리량을 극대화하기 위해서는 스테이징 버퍼(Staging Buffers)를 활용한 이중 버퍼링(Double-buffering)을 적용하는 것이 좋으며, 디스패치 사이에 `await mapAsync()`를 호출할 경우 GPU 파이프라인을 멈추게 만들 수 있으므로 지양해야 합니다 [18].
|
||||
### 매 hardware mapping
|
||||
- NVIDIA: warp (32 threads), SM (streaming multiprocessor).
|
||||
- AMD: wave (64 threads, RDNA: 32), CU.
|
||||
- Apple: simdgroup (32), GPU core.
|
||||
- Intel: subgroup, EU.
|
||||
|
||||
---
|
||||
### 매 languages 2026
|
||||
- **WGSL** (WebGPU): cross-platform, modern.
|
||||
- **HLSL** (DirectX, Vulkan via DXC).
|
||||
- **GLSL** (OpenGL, Vulkan).
|
||||
- **MSL** (Metal Shading Language).
|
||||
- **CUDA C++**: NVIDIA only, but mature.
|
||||
- **Triton** (OpenAI): Python-like ML kernel DSL.
|
||||
|
||||
* **파티클 및 물리 시뮬레이션 처리 한계 돌파**
|
||||
기존 CPU 기반 파티클 업데이트는 일반적인 하드웨어에서 약 50,000개 부근에서 성능 병목에 도달하지만, WebGPU 컴퓨트 셰이더를 사용하면 이를 수백만 개 단위로 확장할 수 있습니다 [2, 3]. 예를 들어, CPU에서 프레임당 30ms가 소요되던 10,000개의 파티클 업데이트 작업을 컴퓨트 셰이더로 전환하면 100,000개의 파티클을 2ms 이내에 업데이트할 수 있어 약 150배의 성능 향상을 얻을 수 있습니다 [4]. 또한, 대규모 유체 시뮬레이션 및 물리 연산에도 탁월한 성능을 발휘합니다 [5, 6].
|
||||
### 매 응용
|
||||
1. ML inference (matmul, attention, conv).
|
||||
2. Image filters (blur, edge, color grading).
|
||||
3. Particle systems / fluid sim.
|
||||
4. Physics (cloth, soft body, mass-spring).
|
||||
5. Cryptography (proof-of-work, hash collisions).
|
||||
6. Video encode/decode prep.
|
||||
|
||||
* **고급 데이터 처리 및 GPU 주도 렌더링([[GPU-driven Rendering|GPU-driven Rendering]])**
|
||||
컴퓨트 셰이더는 충돌 감지(Collision detection), 실시간 조명, 대규모 BIM 데이터셋의 실시간 필터링 등 다수의 데이터 스트림을 병렬로 처리하는 데 유용합니다 [1, 3, 4]. 실시간 편집이 가능한 대규모 절차적 지형(Procedural terrain)을 생성하거나 [6], 컴퓨트 셰이더의 출력을 기반으로 GPU가 렌더링 대상을 직접 결정하는 간접 그리기([[Indirect Draw|Indirect Draw]]s)를 수행하여 수백만 개의 인스턴스와 가시성 컬링(Culling)을 효율적으로 처리할 수 있습니다 [7, 8].
|
||||
## 💻 패턴
|
||||
|
||||
* **컴퓨트 스키닝 (Compute Skinning)**
|
||||
컴퓨트 셰이더는 컴퓨트 단계에서 메쉬 정점 변환을 처리하여 그 결과를 버퍼에 저장할 수 있게 해줍니다 [4]. 이렇게 저장된 데이터는 다수의 렌더링 패스에서 재사용할 수 있어 중복 계산을 없앨 수 있으며, 조립 과정을 보여주는 애니메이션 처리 등에 매우 효율적입니다 [4].
|
||||
### WGSL compute shader — vector add
|
||||
```wgsl
|
||||
@group(0) @binding(0) var<storage, read> a : array<f32>;
|
||||
@group(0) @binding(1) var<storage, read> b : array<f32>;
|
||||
@group(0) @binding(2) var<storage, read_write> c : array<f32>;
|
||||
|
||||
* **핵심 구현 메커니즘 및 동기화 최적화**
|
||||
* **스토리지 텍스처([[Storage|Storage]] textures):** 일반 텍스처와 달리 읽기와 쓰기를 모두 허용하여 컴퓨트 셰이더 내에서 유체 시뮬레이션 및 이미지 처리 작업이 가능하게 합니다 [5, 9].
|
||||
* **작업 그룹 공유 메모리(Workgroup shared [[memory|memory]]):** 스레드 간 데이터 공유가 필요할 때 전역 메모리보다 10~100배 빠른 접근 속도를 제공합니다 [6, 7].
|
||||
* **렌더링 동기화 및 이중 버퍼링:** 컴퓨트 셰이더가 포함된 씬은 GPU 작업이 종속된 렌더링 패스 이전에 완료되도록 `renderAsync`를 사용하여 비동기 렌더링을 수행해야 합니다 [10]. 또한 성능을 높이려면 스테이징 버퍼(Staging buffers)를 활용한 이중 버퍼링(Double-buffering) 기법을 사용해야 하며, 파이프라인 지연을 방지하기 위해 디스패치 간에 `await mapAsync()` 사용을 피해야 합니다 [11].
|
||||
|
||||
## ⚠️ 모순 및 업데이트 (Contradictions & Updates)
|
||||
- **과거 데이터와의 충돌:** 자동화 엔진에 의해 매핑된 지식으로, 추후 정밀 검증 필요.
|
||||
- **정책 변화:** Graphics & Performance 분야의 자동 자산화 수행.
|
||||
|
||||
---
|
||||
|
||||
- **과거 데이터와의 충돌:** 자동화 엔진에 의해 매핑된 지식으로, 추후 정밀 검증 필요.
|
||||
- **정책 변화:** Graphics & Performance 분야의 자동 자산화 수행.
|
||||
|
||||
## 🔗 지식 연결 (Graph)
|
||||
- **Related Topics:** `[[WebGPU|WebGPU]]`, `GPU-driven Rendering`, `Indirect Drawing`, `Storage Textures`, `[[Frustum Culling|Frustum Culling]]`
|
||||
- **Projects/Contexts:** `Three.js`, `Segments.ai`, `BIM Datasets`
|
||||
- **Contradictions/Notes:** 컴퓨트 셰이더는 엄청난 성능 향상을 제공하지만 구형 API인 [[WebGL|WebGL]]이나 WebGL 2에서는 지원되지 않아 WebGPU 환경이 필수적입니다 [1]. 또한 GPU 최적화를 제대로 다루지 못해 동기화 대기(`await mapAsync()`)를 남용할 경우, 오히려 GPU가 최대 60%의 시간 동안 유휴 상태(Idle)에 빠지는 병목 현상을 유발할 수 있습니다 [18].
|
||||
|
||||
---
|
||||
*Last updated: 2026-04-19*
|
||||
|
||||
---
|
||||
|
||||
---
|
||||
|
||||
- **Related Topics:** [[WebGPU|WebGPU]], GPU-driven Rendering, [[TSL (Three Shader Language)|TSL (Three Shader Language]], Storage Textures
|
||||
- **Projects/Contexts:** Three.js WebGPURenderer, Native WebGPU, 대규모 건설/BIM 플랫폼 (Large-Scale Construction Viewers)
|
||||
- **Contradictions/Notes:** 컴퓨트 셰이더를 통한 GPU 병렬 연산은 압도적인 성능 향상을 가져오지만, 작업 디스패치 사이에 `await mapAsync()`를 무분별하게 사용하면 GPU 파이프라인이 멈추고 최대 60%의 시간 동안 GPU가 유휴 상태에 빠지는 성능 저하 역효과가 발생할 수 있으므로 주의해야 합니다 [11].
|
||||
|
||||
---
|
||||
*Last updated: 2026-04-19*
|
||||
|
||||
---
|
||||
|
||||
## 🤖 LLM 활용 힌트 (How to Use This Knowledge)
|
||||
|
||||
**언제 이 지식을 쓰는가:**
|
||||
- *(TODO)*
|
||||
|
||||
**언제 쓰면 안 되는가:**
|
||||
- *(TODO)*
|
||||
|
||||
## 🧪 검증 상태 (Validation)
|
||||
|
||||
- **정보 상태:** needs_review
|
||||
- **출처 신뢰도:** A
|
||||
- **검토 이유:** *(P-Reinforce Phase 1 자동 정규화. 본문 검증 필요.)*
|
||||
|
||||
## 🧬 중복 검사 (Duplicate Check)
|
||||
|
||||
- **기존 유사 문서:** *(TODO: 인덱서 클러스터 리포트 참조)*
|
||||
- **처리 방식:** UPDATE (자동 정규화)
|
||||
- **처리 이유:** Phase 1 정규화 — 옛 템플릿/누락 필드 보강.
|
||||
|
||||
## 🕓 변경 이력 (Changelog)
|
||||
|
||||
| 날짜 | 변경 내용 | 처리 방식 | 신뢰도 |
|
||||
|------|-----------|-----------|--------|
|
||||
| 2026-05-08 | P-Reinforce Phase 1 정규화 (frontmatter + 헤더 표준화) | UPDATE | A |
|
||||
|
||||
## 💻 코드 패턴 (Code Patterns)
|
||||
|
||||
**패턴 1:** *(TODO: 이 프로젝트 컨벤션 반영한 구조 스켈레톤)*
|
||||
|
||||
```text
|
||||
# TODO
|
||||
@compute @workgroup_size(64)
|
||||
fn main(@builtin(global_invocation_id) gid : vec3u) {
|
||||
let i = gid.x;
|
||||
if (i >= arrayLength(&a)) { return; }
|
||||
c[i] = a[i] + b[i];
|
||||
}
|
||||
```
|
||||
|
||||
## 🤔 의사결정 기준 (Decision Criteria)
|
||||
### WebGPU dispatch (TypeScript)
|
||||
```typescript
|
||||
const adapter = await navigator.gpu.requestAdapter();
|
||||
const device = await adapter!.requestDevice();
|
||||
|
||||
**선택 A를 써야 할 때:**
|
||||
- *(TODO)*
|
||||
const module = device.createShaderModule({ code: WGSL_SOURCE });
|
||||
const pipeline = device.createComputePipeline({
|
||||
layout: 'auto',
|
||||
compute: { module, entryPoint: 'main' },
|
||||
});
|
||||
|
||||
**선택 B를 써야 할 때:**
|
||||
- *(TODO)*
|
||||
const N = 1_000_000;
|
||||
const buf = (data: Float32Array) => {
|
||||
const b = device.createBuffer({
|
||||
size: data.byteLength,
|
||||
usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST | GPUBufferUsage.COPY_SRC,
|
||||
});
|
||||
device.queue.writeBuffer(b, 0, data);
|
||||
return b;
|
||||
};
|
||||
|
||||
**기본값:**
|
||||
> *(TODO)*
|
||||
const a = buf(new Float32Array(N).fill(1));
|
||||
const b = buf(new Float32Array(N).fill(2));
|
||||
const c = device.createBuffer({
|
||||
size: N * 4,
|
||||
usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC,
|
||||
});
|
||||
|
||||
## ❌ 안티패턴 (Anti-Patterns)
|
||||
const bindGroup = device.createBindGroup({
|
||||
layout: pipeline.getBindGroupLayout(0),
|
||||
entries: [
|
||||
{ binding: 0, resource: { buffer: a } },
|
||||
{ binding: 1, resource: { buffer: b } },
|
||||
{ binding: 2, resource: { buffer: c } },
|
||||
],
|
||||
});
|
||||
|
||||
- **[안티패턴]:** *(TODO: 무엇을 하면 안 되는가 + 이유 + 대신 무엇을)*
|
||||
const enc = device.createCommandEncoder();
|
||||
const pass = enc.beginComputePass();
|
||||
pass.setPipeline(pipeline);
|
||||
pass.setBindGroup(0, bindGroup);
|
||||
pass.dispatchWorkgroups(Math.ceil(N / 64));
|
||||
pass.end();
|
||||
device.queue.submit([enc.finish()]);
|
||||
```
|
||||
|
||||
### Shared memory reduction (workgroup-local)
|
||||
```wgsl
|
||||
var<workgroup> shared : array<f32, 256>;
|
||||
|
||||
@compute @workgroup_size(256)
|
||||
fn reduce(
|
||||
@builtin(local_invocation_id) lid : vec3u,
|
||||
@builtin(global_invocation_id) gid : vec3u,
|
||||
) {
|
||||
shared[lid.x] = input[gid.x];
|
||||
workgroupBarrier();
|
||||
|
||||
var stride : u32 = 128u;
|
||||
loop {
|
||||
if (stride == 0u) { break; }
|
||||
if (lid.x < stride) {
|
||||
shared[lid.x] = shared[lid.x] + shared[lid.x + stride];
|
||||
}
|
||||
workgroupBarrier();
|
||||
stride = stride / 2u;
|
||||
}
|
||||
|
||||
if (lid.x == 0u) { output[gid.x / 256u] = shared[0]; }
|
||||
}
|
||||
```
|
||||
|
||||
### CUDA matmul kernel
|
||||
```cuda
|
||||
__global__ void matmul(const float* A, const float* B, float* C, int N) {
|
||||
int row = blockIdx.y * blockDim.y + threadIdx.y;
|
||||
int col = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
if (row >= N || col >= N) return;
|
||||
|
||||
float sum = 0.0f;
|
||||
for (int k = 0; k < N; ++k) {
|
||||
sum += A[row * N + k] * B[k * N + col];
|
||||
}
|
||||
C[row * N + col] = sum;
|
||||
}
|
||||
// Launch: matmul<<<dim3((N+15)/16, (N+15)/16), dim3(16,16)>>>(A, B, C, N);
|
||||
```
|
||||
|
||||
### Triton kernel (Python ML)
|
||||
```python
|
||||
import triton
|
||||
import triton.language as tl
|
||||
|
||||
@triton.jit
|
||||
def add_kernel(x_ptr, y_ptr, out_ptr, n, BLOCK: tl.constexpr):
|
||||
pid = tl.program_id(0)
|
||||
offsets = pid * BLOCK + tl.arange(0, BLOCK)
|
||||
mask = offsets < n
|
||||
x = tl.load(x_ptr + offsets, mask=mask)
|
||||
y = tl.load(y_ptr + offsets, mask=mask)
|
||||
tl.store(out_ptr + offsets, x + y, mask=mask)
|
||||
|
||||
# Launch
|
||||
add_kernel[(triton.cdiv(N, 1024),)](x, y, out, N, BLOCK=1024)
|
||||
```
|
||||
|
||||
### Image blur compute (storage texture)
|
||||
```wgsl
|
||||
@group(0) @binding(0) var src : texture_2d<f32>;
|
||||
@group(0) @binding(1) var dst : texture_storage_2d<rgba8unorm, write>;
|
||||
|
||||
@compute @workgroup_size(8, 8)
|
||||
fn blur(@builtin(global_invocation_id) gid : vec3u) {
|
||||
var sum = vec4f(0);
|
||||
for (var dy = -1; dy <= 1; dy++) {
|
||||
for (var dx = -1; dx <= 1; dx++) {
|
||||
let p = vec2i(gid.xy) + vec2i(dx, dy);
|
||||
sum = sum + textureLoad(src, p, 0);
|
||||
}
|
||||
}
|
||||
textureStore(dst, vec2i(gid.xy), sum / 9.0);
|
||||
}
|
||||
```
|
||||
|
||||
## 매 결정 기준
|
||||
| 상황 | Approach |
|
||||
|---|---|
|
||||
| Browser, cross-platform | WebGPU + WGSL |
|
||||
| Native cross-platform | Vulkan compute + GLSL/HLSL |
|
||||
| NVIDIA-only, max perf | CUDA |
|
||||
| Apple ecosystem | Metal compute (MSL) |
|
||||
| ML kernel research | Triton (Python) |
|
||||
| Production ML inference | Pre-built (cuDNN, MLX, vLLM kernels) |
|
||||
|
||||
**기본값**: 매 web/cross-platform → WebGPU + WGSL. 매 ML research → Triton. 매 production NVIDIA ML → CUDA + cuDNN/cuBLAS.
|
||||
|
||||
## 🔗 Graph
|
||||
- 부모: [[GPU Programming]] · [[Parallel Computing]]
|
||||
- 변형: [[Vertex Shader]] · [[Fragment Shader]] · [[CUDA]]
|
||||
- 응용: [[WebGPU]] · [[ML Inference]] · [[Particle Systems]]
|
||||
- Adjacent: [[Triton]] · [[MLX]] · [[Vulkan]]
|
||||
|
||||
## 🤖 LLM 활용
|
||||
**언제**: heavy parallel data (image, ML, sim), browser GPU compute, custom ML kernels.
|
||||
**언제 X**: small data (<10k items, CPU faster after transfer cost), branch-heavy serial logic, very small kernels (launch overhead).
|
||||
|
||||
## ❌ 안티패턴
|
||||
- **Divergent branching in warp**: 매 thread 가 different path → serialization → 매 SIMT 의 X.
|
||||
- **Uncoalesced memory access**: random pattern → bandwidth waste — adjacent threads should read adjacent memory.
|
||||
- **Tiny dispatch**: 100 threads → launch overhead > work — batch.
|
||||
- **Forgetting workgroupBarrier**: race condition on shared memory.
|
||||
- **CPU↔GPU ping-pong**: every step copies back — keep data on GPU.
|
||||
|
||||
## 🧪 검증 / 중복
|
||||
- Verified (WebGPU spec 2026 W3C / CUDA Programming Guide 12.x / Triton docs / Apple MSL).
|
||||
- 신뢰도 A.
|
||||
|
||||
## 🕓 Changelog
|
||||
| 날짜 | 변경 |
|
||||
|---|---|
|
||||
| 2026-05-08 | Phase 1 |
|
||||
| 2026-05-10 | Manual cleanup — WGSL/CUDA/Triton + workgroup model |
|
||||
|
||||
Reference in New Issue
Block a user